第四节：第一个CUDA程序——向量加法实战

从这一节开始，我们正式进入CUDA编程的实操环节。向量加法是CUDA世界里的“Hello World”，它虽然简单，却完整地展示了GPU编程的核心流程：分配显存、拷贝数据、启动内核、同步结果、释放资源。理解这个最小闭环，后续的矩阵乘法、卷积、归约等复杂内核都只是在这个骨架上添加血肉。本节将带你从零编写、编译、运行并调试一个向量加法程序，同时剖析初学者最容易踩中的“雷区”。

上图展示了CUDA向量加法的完整数据流。请注意，内核启动是异步的，CPU不会等待GPU执行完毕，而是立即返回继续执行后续代码。因此，在主机端读取结果之前，必须使用cudaMemcpy（D2H方向）或cudaDeviceSynchronize()进行同步。这是初学者最容易忽略的时序问题。

4.1 核心原理：GPU线程层次与内存模型

在编写第一行CUDA代码之前，必须理解GPU的执行模型。CUDA将线程组织为三层结构：网格（Grid）→ 线程块（Block）→ 线程（Thread）。内核启动时，我们通过<<<gridDim, blockDim>>>指定网格和块的大小。

• 线程（Thread）：最小的执行单元，每个线程执行相同的内核代码，但通过内置变量threadIdx.x区分自己的身份。
• 线程块（Block）：一组线程的集合，块内线程可以通过共享内存（Shared Memory）协作，并通过__syncthreads()同步。
• 网格（Grid）：内核启动时创建的所有线程块的集合。网格和块都可以是一维、二维或三维的。

对于向量加法，我们通常使用一维布局。假设向量长度为N，每个块包含256个线程，则需要的块数为(N + 255) / 256。每个线程的全局索引计算公式为：

int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < N) {  // 边界检查，防止越界
    c[tid] = a[tid] + b[tid];
}

内存模型方面，CUDA将内存分为主机端（Host）和设备端（Device）。设备端内存包括全局内存（Global Memory）、共享内存（Shared Memory）、常量内存（Constant Memory）和纹理内存（Texture Memory）。全局内存是所有线程都可访问的，但延迟较高（数百个时钟周期）；共享内存是块内线程共享的，延迟极低（约1-2个时钟周期），但容量有限（通常几十KB）。

4.2 操作流程：编写、编译与运行

下面我们逐步完成一个完整的向量加法程序。假设我们要计算两个长度为N=1024的浮点向量相加。

步骤1：编写内核函数（Kernel）

__global__ void vecAdd(float *a, float *b, float *c, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        c[tid] = a[tid] + b[tid];
    }
}

步骤2：编写主函数（Host Code）

int main() {
    int n = 1024;
    size_t size = n * sizeof(float);
    float *h_a, *h_b, *h_c;
    float *d_a, *d_b, *d_c;

    // 1. 分配主机内存并初始化
    h_a = (float*)malloc(size);
    h_b = (float*)malloc(size);
    h_c = (float*)malloc(size);
    for (int i = 0; i < n; i++) { h_a[i] = i * 1.0f; h_b[i] = i * 2.0f; }

    // 2. 分配设备显存
    cudaMalloc(&d_a, size);
    cudaMalloc(&d_b, size);
    cudaMalloc(&d_c, size);

    // 3. 将输入数据从主机拷贝到设备
    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);

    // 4. 配置内核启动参数并启动
    int blockSize = 256;
    int gridSize = (n + blockSize - 1) / blockSize;
    vecAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);

    // 5. 将结果从设备拷贝回主机（隐式同步）
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);

    // 6. 验证结果
    bool ok = true;
    for (int i = 0; i < n; i++) {
        if (fabs(h_c[i] - (h_a[i] + h_b[i])) > 1e-5) { ok = false; break; }
    }
    printf("Result: %s\n", ok ? "PASS" : "FAIL");

    // 7. 释放资源
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    free(h_a); free(h_b); free(h_c);
    return 0;
}

步骤3：编译与运行

nvcc -o vecAdd vecAdd.cu
./vecAdd

如果一切正常，输出应为Result: PASS。这里的关键点在于：cudaMemcpy（D2H方向）会阻塞CPU，直到内核执行完成，因此不需要显式调用cudaDeviceSynchronize()。但如果你在拷贝之前需要查看中间状态，就必须手动同步。

4.3 案例拆解：一个典型的“未检查错误码”事故

假设你运行上面的程序，发现输出是Result: FAIL，或者干脆没有任何输出就崩溃了。你可能会怀疑算法逻辑，但更常见的原因是CUDA API调用失败。例如，显存不足、设备不支持某个特性、或者内核启动配置错误。

让我们模拟一个场景：你的GPU是较老的型号，最大线程块大小限制为512。如果你不小心设置了blockSize = 1024，内核启动会失败，但cudaMemcpy会返回一个错误码。如果你没有检查这个错误码，程序会继续执行，最终得到错误结果或崩溃。

初始现象：程序输出Result: FAIL，但没有任何报错信息。

分析过程：在vecAdd<<<gridSize, blockSize>>>()之后，立即添加错误检查代码：

cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
    printf("Kernel launch failed: %s\n", cudaGetErrorString(err));
    return -1;
}

重新编译运行，这次你会看到：Kernel launch failed: invalid configuration argument。问题定位到blockSize过大。

修复与验证：将blockSize改为256，重新运行，程序通过。这个案例告诉我们：CUDA编程中，错误检查不是可选项，而是必须项。因为GPU执行是异步的，很多错误不会立即抛出，而是延迟到下一个同步点。

#define CHECK(call) { \
    cudaError_t err = call; \
    if (err != cudaSuccess) { \
        printf("Error %s at %s:%d\n", cudaGetErrorString(err), __FILE__, __LINE__); \
        exit(1); \
    } \
}

4.4 调试技巧与常见误区

除了错误码检查，还有几个调试技巧能显著提升效率：

• 使用cuda-memcheck工具：在命令行运行cuda-memcheck ./vecAdd，可以检测越界访问、未初始化内存等错误。这是排查“程序崩溃但无报错”的利器。
• 打印调试信息：在内核中直接使用printf()是可行的（计算能力3.5+），但会严重拖慢性能。建议只在调试时使用，且只打印前几个线程的值。
• 逐步缩小规模：如果结果不对，先将N改为16，blockSize改为4，手动计算期望值，对比输出。
• 检查边界条件：当N不能被blockSize整除时，最后一个块的线程数会少于blockSize。务必使用if (tid < n)进行保护。

常见误区总结：