4、第一个CUDA程序：编写向量加法Kernel、nvcc编译流程、运行时API、错误检查与调试

好，咱们终于要动手写第一个CUDA程序了。

说实话，我当年第一次写CUDA的时候，心里也没底。毕竟CPU编程和GPU编程，思维方式完全不同。但别担心，向量加法就是CUDA世界的"Hello World"，简单、直观，却能让你把整个流程跑通一遍。

4.1 向量加法的Kernel怎么写？

先看一个最简单的例子。假设有两个数组A和B，长度都是N。我们要把A[i] + B[i]的结果存到C[i]里。

CPU版本大家都会写：

for (int i = 0; i < N; i++) {
    C[i] = A[i] + B[i];
}

GPU版本呢？说白了，就是把那个for循环拆开，让每个线程处理一个元素。

__global__ void vecAdd(float *A, float *B, float *C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

注意几个关键点：

  • __global__ 修饰符告诉编译器，这个函数是在GPU上跑的，但可以从CPU调用
  • blockIdx.x、blockDim.x、threadIdx.x 是CUDA内置变量，用来计算线程的全局ID
  • 那个 if (i < N) 是边界检查。为什么需要？因为线程总数可能比N大，多出来的线程不能越界访问
我的习惯：写Kernel时，我总会在第一行就把线程索引算出来。这样后面写起来清爽，不容易出错。

4.2 nvcc编译流程——到底发生了什么？

写好了Kernel，怎么编译？用nvcc。

nvcc不是普通的编译器。它是个"二传手"。你想想看，代码里既有CPU代码（比如main函数），又有GPU代码（Kernel）。nvcc会把它们分开处理：

  1. 分离阶段：nvcc把.cu文件拆成两部分——CPU代码交给主机编译器（gcc或cl.exe），GPU代码自己处理
  2. GPU编译：GPU代码被编译成PTX（一种中间表示），再转成cubin（二进制）
  3. 链接阶段：最后把两部分链接到一起，生成可执行文件

编译命令很简单：

nvcc -o vecAdd vecAdd.cu

如果你想看看中间产物，加个 --keep 选项：

nvcc --keep -o vecAdd vecAdd.cu

这样会生成一堆 .ptx、.cubin 文件。我建议你试试看，能帮你理解编译流程。

注意：我曾经遇到过一个问题——在Windows上用nvcc编译，默认主机编译器是cl.exe。如果没装Visual Studio，会报错。解决办法是装好VS，或者在Linux上用gcc。

4.3 运行时API——怎么把数据搬上GPU？

Kernel写好了，编译通过了。但数据还在CPU上呢。怎么搬？

CUDA运行时API提供了几个关键函数：

函数 作用
cudaMalloc 在GPU上分配显存
cudaMemcpy 在CPU和GPU之间拷贝数据
cudaFree 释放显存

完整的流程是这样的：

// 1. 分配主机内存
float *h_A, *h_B, *h_C;
h_A = (float*)malloc(N * sizeof(float));
h_B = (float*)malloc(N * sizeof(float));
h_C = (float*)malloc(N * sizeof(float));

// 2. 初始化数据（略）

// 3. 分配设备内存
float *d_A, *d_B, *d_C;
cudaMalloc(&d_A, N * sizeof(float));
cudaMalloc(&d_B, N * sizeof(float));
cudaMalloc(&d_C, N * sizeof(float));

// 4. 把数据从CPU拷贝到GPU
cudaMemcpy(d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, N * sizeof(float), cudaMemcpyHostToDevice);

// 5. 调用Kernel
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

// 6. 把结果拷回CPU
cudaMemcpy(h_C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);

// 7. 释放内存
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
free(h_A);
free(h_B);
free(h_C);

这里有个细节：cudaMemcpy的最后一个参数指定方向。有四种：

  • cudaMemcpyHostToDevice：CPU → GPU
  • cudaMemcpyDeviceToHost：GPU → CPU
  • cudaMemcpyDeviceToDevice：GPU内部拷贝
  • cudaMemcpyHostToHost：CPU内部拷贝（基本不用）
避坑指南：我刚开始写的时候，经常忘记调用cudaMemcpy把数据拷回CPU。结果在CPU上打印h_C，全是0。嗯，检查了半天才发现数据根本没回来。

4.4 错误检查与调试——别让bug偷偷溜过去

CUDA程序有个特点：很多错误不会直接报出来。比如你访问了越界内存，程序可能继续跑，但结果不对。或者Kernel根本没执行，但CPU代码照常运行。

所以，错误检查是必须的。

最简单的办法：检查每个CUDA API调用的返回值。

cudaError_t err = cudaMalloc(&d_A, N * sizeof(float));
if (err != cudaSuccess) {
    printf("cudaMalloc failed: %s\n", cudaGetErrorString(err));
    exit(1);
}

但这样写太啰嗦了。我一般会封装一个宏：

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

然后这样用：

CHECK(cudaMalloc(&d_A, N * sizeof(float)));
CHECK(cudaMemcpy(d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice));

对于Kernel调用，错误检查要分两步：

vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
CHECK(cudaGetLastError());  // 检查Kernel启动有没有错误
CHECK(cudaDeviceSynchronize());  // 等待Kernel执行完，检查执行错误

为什么需要两步？因为Kernel启动是异步的。你调用 <<<>>> 时，CPU不会等GPU执行完。所以第一步只能检查启动参数对不对（比如block数太多），第二步才能检查执行时的错误（比如越界）。

我曾经踩过的坑：有一次Kernel里写了个死循环，但没加cudaDeviceSynchronize。结果CPU代码跑完了，程序退出，GPU还在跑。最后系统报了个"GPU hung"的错误。从那以后，我每个Kernel后面都加同步检查。

4.5 完整示例——把一切串起来

好了，咱们把上面所有内容拼成一个完整的程序：

#include <stdio.h>
#include <cuda_runtime.h>

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

__global__ void vecAdd(float *A, float *B, float *C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

int main() {
    int N = 1 << 20;  // 约100万个元素
    size_t size = N * sizeof(float);

    // 分配主机内存
    float *h_A = (float*)malloc(size);
    float *h_B = (float*)malloc(size);
    float *h_C = (float*)malloc(size);

    // 初始化数据
    for (int i = 0; i < N; i++) {
        h_A[i] = i * 1.0f;
        h_B[i] = i * 2.0f;
    }

    // 分配设备内存
    float *d_A, *d_B, *d_C;
    CHECK(cudaMalloc(&d_A, size));
    CHECK(cudaMalloc(&d_B, size));
    CHECK(cudaMalloc(&d_C, size));

    // 拷贝数据到GPU
    CHECK(cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice));

    // 启动Kernel
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    CHECK(cudaGetLastError());
    CHECK(cudaDeviceSynchronize());

    // 拷贝结果回CPU
    CHECK(cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost));

    // 验证结果
    for (int i = 0; i < 10; i++) {
        printf("h_C[%d] = %f (expected %f)\n", i, h_C[i], h_A[i] + h_B[i]);
    }

    // 清理
    CHECK(cudaFree(d_A));
    CHECK(cudaFree(d_B));
    CHECK(cudaFree(d_C));
    free(h_A);
    free(h_B);
    free(h_C);

    return 0;
}

编译运行：

nvcc -o vecAdd vecAdd.cu
./vecAdd

如果一切正常，你会看到前10个元素的结果，每个都等于 i + 2*i = 3*i。

总结一下：写第一个CUDA程序，核心就四步——写Kernel、编译、搬数据、检查错误。把这四步练熟了，后面的路就好走了。

嗯，下一章咱们会聊聊线程、块、网格的配置策略。你会发现，同样的Kernel，不同的配置，性能能差好几倍。到时候再细说。

← 上一章 📖 返回目录 下一章 →