CUDA计算π程序出现段错误，求原因及正确内存分配方式

段错误核心原因 #

• 主机直接访问设备内存：你用cudaMalloc给data分配了内存，但cudaMalloc是在GPU设备内存中分配空间，主机CPU不能直接读写设备内存地址。后面cout << data[i]尝试访问这块设备内存，直接触发段错误。正确做法是主机内存用malloc或普通数组分配，设备内存单独用cudaMalloc。

其他代码问题 #

• 变量未初始化：核函数里的y0没有初始化，第一次循环计算(y0 + y1)时会使用垃圾值，导致积分结果错误。
• 网格大小计算错误：GRID_SIZE*THREAD_SIZE/BLOCK_SIZE计算结果为1*1/16=0，网格大小为0的话，核函数根本不会执行。
• 积分逻辑问题：dx设置为0.1步长过大，精度极低；初始x0的计算方式不符合积分区间要求，最后也没有将四分之一圆面积乘4得到π。
• 未检查CUDA API错误：所有CUDA操作（cudaMalloc、cudaMemcpy、核函数执行）都应该检查返回值，方便快速定位问题。

修正后的代码 #

// System includes
#include <stdio.h>
#include <assert.h>
#include <cuda.h>
#include <iostream>
#include <cmath>

using namespace std;

#define CUDA_FLOAT float
#define BLOCK_SIZE 256  // 采用CUDA常用的块大小
#define GRID_SIZE 16    // 根据总线程数调整网格大小
#define NUM_STEPS 100000  // 积分总步数，替代固定dx提升精度

__global__ void pi_kern(CUDA_FLOAT *res)
{
    // 计算当前线程的全局ID
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    // 获取总线程数
    int total_threads = blockDim.x * gridDim.x;
    
    CUDA_FLOAT s = 0.0f;
    // 每个线程分配部分积分任务，实现负载均衡
    for (int i = tid; i < NUM_STEPS; i += total_threads) {
        CUDA_FLOAT x0 = static_cast<CUDA_FLOAT>(i) / NUM_STEPS;
        CUDA_FLOAT x1 = static_cast<CUDA_FLOAT>(i + 1) / NUM_STEPS;
        CUDA_FLOAT y0 = sqrtf(1.0f - x0 * x0);
        CUDA_FLOAT y1 = sqrtf(1.0f - x1 * x1);
        // 梯形法计算单步积分
        s += (y0 + y1) * (x1 - x0) / 2.0f;
    }
    res[tid] = s;
}

int main(int argc, char **argv)
{   
    printf("[pi-calculation] - Starting\n");
    
    // 主机内存：用malloc分配，确保是主机可访问的内存空间
    int total_threads = GRID_SIZE * BLOCK_SIZE;
    CUDA_FLOAT* data = (CUDA_FLOAT*)malloc(total_threads * sizeof(CUDA_FLOAT));
    if (!data) {
        cerr << "主机内存分配失败!" << endl;
        return 1;
    }
    
    // 设备内存分配
    CUDA_FLOAT* d_data;
    cudaError_t err = cudaMalloc((void **)&d_data, total_threads * sizeof(CUDA_FLOAT));
    if (err != cudaSuccess) {
        cerr << "cudaMalloc失败: " << cudaGetErrorString(err) << endl;
        free(data);
        return 1;
    }
    
    // 启动核函数
    dim3 block(BLOCK_SIZE);
    dim3 grid(GRID_SIZE);
    pi_kern<<<grid, block>>>(d_data);
    
    // 检查核函数启动错误
    err = cudaGetLastError();
    if (err != cudaSuccess) {
        cerr << "核函数启动失败: " << cudaGetErrorString(err) << endl;
        cudaFree(d_data);
        free(data);
        return 1;
    }
    // 等待核函数执行完成
    cudaDeviceSynchronize();
    
    // 将设备结果拷贝到主机
    err = cudaMemcpy(data, d_data, total_threads * sizeof(CUDA_FLOAT), cudaMemcpyDeviceToHost);
    if (err != cudaSuccess) {
        cerr << "cudaMemcpy失败: " << cudaGetErrorString(err) << endl;
        cudaFree(d_data);
        free(data);
        return 1;
    }
    
    // 汇总所有线程的积分结果，乘4得到完整π值
    CUDA_FLOAT pi = 0.0f;
    for (int i = 0; i < total_threads; i++) {
        pi += data[i];
    }
    pi *= 4.0f;
    cout << "计算得到的π值: " << pi << endl;
    
    // 释放内存
    cudaFree(d_data);
    free(data);
    
    return 0;
}

关键修正说明 #

• 内存分配规范：主机内存用malloc分配，设备内存用cudaMalloc，主机只访问自身内存空间，彻底避免段错误。
• 核函数逻辑优化：用总步数替代固定dx，每个线程分配部分积分任务，提升计算精度和并行效率。
• 错误检查完善：给所有CUDA操作添加错误检查，能快速定位执行过程中的问题。
• 结果正确汇总：将所有线程的积分结果相加后乘4，得到完整的圆周率π值。

内容的提问来源于stack exchange，提问作者Николай Козлов