2. CUDA编程基础与调试入口

各位同学，今天我们来聊聊CUDA编程最基础、也最容易被忽视的部分。说实话，我见过太多人一上来就写核函数，跑出错误结果后一脸懵。其实，很多问题在入口处就能避免。

2.1 CUDA核函数编写要点

核函数，就是跑在GPU上的那个函数。用 __global__ 修饰，返回类型必须是 void。嗯，这是硬性规定。

__global__ void myKernel(float *d_data, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        // 你的计算逻辑
        d_data[idx] = d_data[idx] * 2.0f;
    }
}

我个人习惯，在核函数里一定会加边界检查。你想想看，如果线程总数超过数据长度，不检查的话就会访问非法内存。我在项目中遇到过好几次，就是因为少写了这个 if 判断，结果数据被莫名其妙地改写了。

几个关键点，我列出来：

• 线程索引计算：threadIdx.x + blockIdx.x * blockDim.x 是最常用的方式。多维网格的话，要分别算好 x、y、z 方向。
• 内存访问模式：尽量让相邻线程访问连续地址。这叫合并访问，能充分利用显存带宽。
• 避免分支发散：同一个 warp 里的线程，如果走不同的分支，性能会下降。我建议尽量让 warp 内的线程执行路径一致。

2.2 cudaError_t 错误处理

CUDA 的 API 函数几乎都返回 cudaError_t 类型。说白了，这就是一个枚举值，告诉你调用成功还是失败。但很多人不检查返回值，这是大忌。

我曾经接手过一个项目，代码跑了一整天，结果全是错的。查了半天，发现是 cudaMalloc 返回了 cudaErrorMemoryAllocation，但代码里根本没检查。你说冤不冤？

正确的做法是：

cudaError_t err = cudaMalloc(&d_data, size);
if (err != cudaSuccess) {
    fprintf(stderr, "cudaMalloc failed: %s\n", cudaGetErrorString(err));
    // 处理错误，比如退出或回滚
}

常用的错误类型，我整理了一张表：

2.3 cudaGetLastError 的使用

这个函数，可以说是调试核函数的利器。为什么？因为核函数是异步执行的，你调用它之后，错误不会立刻返回。等你调用 cudaGetLastError 时，它才会把之前积压的错误告诉你。

我一般这样用：

myKernel<<<grid, block>>>(d_data, N);
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
    printf("Kernel launch failed: %s\n", cudaGetErrorString(err));
}

// 或者，在同步之后检查
cudaDeviceSynchronize();
err = cudaGetLastError();
if (err != cudaSuccess) {
    printf("Kernel execution failed: %s\n", cudaGetErrorString(err));
}

这里有个坑，我踩过好几次。你想想看，如果之前有某个 API 调用已经报错了，但你没处理，那么 cudaGetLastError 返回的是那个旧错误，而不是核函数的错误。所以，我建议在每个可能出错的地方都检查一下，或者至少在执行核函数前调用一次 cudaGetLastError 清空错误状态。

#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            fprintf(stderr, "CUDA error at %s:%d - %s\n", \
                    __FILE__, __LINE__, cudaGetErrorString(err)); \
            exit(EXIT_FAILURE); \
        } \
    } while(0)

知识体系总览

下面这张图，是我梳理的本章核心逻辑。你可以看到，从核函数编写到错误处理，再到调试入口，是一条完整的链路。

说白了，这三块内容是一个整体。核函数写得再好，不检查错误也是白搭。而 cudaGetLastError 就是你调试核函数的最后一道防线。我建议你从现在开始，养成每个 CUDA 调用都检查返回值的习惯。刚开始可能觉得麻烦，但等你遇到几次莫名其妙的 bug 后，就知道这有多重要了。