第3章 CUDA编程入门

第3章：CUDA编程入门：向量加法与编译调试

好，咱们直接进入正题。这一章我带你写第一个CUDA程序——向量加法。别小看这个例子，它麻雀虽小五脏俱全。我在带新人时，第一课永远是它。你把它吃透了，后面那些花里胡哨的算子优化，本质上都是这个框架的变种。

3.1 向量加法的GPU实现

先说说向量加法要干什么。两个数组A和B，长度都是N，我们要算C[i] = A[i] + B[i]。CPU上写个for循环就完事了。但GPU上不一样，我们要让每个线程算一个元素。

核心思路就一句话：一个线程负责一个元素。你想想看，如果N=1024，我们就启动1024个线程，每个线程只做一次加法。这就是所谓的「数据并行」。

来看代码：

// 向量加法的CUDA kernel
__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 N = 1024;
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

这里有个细节我当年踩过坑：if (i < N)这个边界检查。为什么要有它？因为N不一定能被线程块大小整除。比如N=1000，threadsPerBlock=256，那blocksPerGrid=4，总线程数1024，多出来的24个线程就得跳过。嗯，这个习惯一定要养成。

关键概念：

threadIdx.x：线程在块内的索引
blockIdx.x：块在网格内的索引
blockDim.x：每个块的线程数
全局线程ID = blockIdx.x * blockDim.x + threadIdx.x

3.2 nvcc编译器与编译流程

写好了代码怎么跑？用nvcc。它是NVIDIA的CUDA编译器，专门处理.cu文件。

我个人习惯把编译命令写成这样：

nvcc -o vec_add vec_add.cu -arch=sm_75

-arch=sm_75指定了目标架构。sm_75对应Turing架构（RTX 20系列）。你用的是什么显卡？可以用nvidia-smi查一下。我建议编译时加上-arch=sm_XX，别偷懒用默认值，否则性能可能打折扣。

nvcc的编译流程分两步：

分离设备代码和主机代码：nvcc把.cu文件拆成两部分。设备代码（kernel函数）被编译成PTX（Parallel Thread Execution）中间代码，主机代码（main函数等）交给宿主编译器（gcc/clang）处理。
生成cubin：PTX再被编译成实际的GPU机器码，打包成cubin文件。运行时由CUDA驱动加载。

为什么会这样设计？说白了就是为了兼容性。PTX是中间表示，可以针对不同架构的GPU生成不同的cubin。你写一次代码，跑在不同显卡上，nvcc帮你搞定底层差异。

小技巧： 用nvcc --ptx可以只生成PTX代码，方便你查看编译器做了什么优化。我在调试性能问题时经常用这招。

3.3 CUDA错误检查与调试

写CUDA程序最烦什么？运行时错误。CPU上你还能打日志，GPU上出错了直接崩，连个提示都没有。我刚开始做项目时，经常花半天时间找bug，后来学乖了——每个CUDA API调用都要检查返回值。

我的做法是封装一个宏：

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

// 使用示例
CUDA_CHECK(cudaMalloc(&d_A, N * sizeof(float)));
CUDA_CHECK(cudaMemcpy(d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice));
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
CUDA_CHECK(cudaGetLastError());  // 检查kernel启动是否成功
CUDA_CHECK(cudaDeviceSynchronize());  // 等待kernel执行完毕

注意最后两行。kernel启动是异步的，错误可能发生在执行过程中。cudaGetLastError()捕获启动时的错误，cudaDeviceSynchronize()等待执行完成并捕获运行时错误。我曾经漏掉cudaDeviceSynchronize()，结果一个越界访问的bug查了两天——教训深刻。

避坑指南： 我曾经在项目里遇到一个诡异问题：kernel执行结果偶尔正确偶尔错误。排查后发现是cudaMemcpy的拷贝方向写反了——从设备拷到主机写成了从主机拷到设备。用了CUDA_CHECK宏后，这种低级错误一眼就能发现。

3.4 完整示例与调试技巧

把上面所有东西串起来，一个完整的向量加法程序长这样：

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

#define CUDA_CHECK(call) /* 上面定义的宏 */

__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 = 1024;
    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;
    CUDA_CHECK(cudaMalloc(&d_A, size));
    CUDA_CHECK(cudaMalloc(&d_B, size));
    CUDA_CHECK(cudaMalloc(&d_C, size));

    // 拷贝数据到设备
    CUDA_CHECK(cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice));
    CUDA_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);
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());

    // 拷贝结果回主机
    CUDA_CHECK(cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost));

    // 验证结果
    for (int i = 0; i < N; i++) {
        if (h_C[i] != h_A[i] + h_B[i]) {
            printf("Error at index %d\n", i);
            break;
        }
    }
    printf("Success!\n");

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

    return 0;
}

调试时我常用cuda-gdb，它跟gdb用法差不多，但能调试GPU上的代码。设置断点、单步执行、查看变量，都支持。不过要注意，编译时要加-G选项：

nvcc -G -o vec_add vec_add.cu

-G会生成调试信息，但会关闭优化，所以调试完记得去掉。

3.5 知识体系图

下面这张图总结了本章的核心逻辑，从代码编写到编译执行再到错误处理，一条线串起来：

这张图把本章内容串起来了。你写代码时，脑子里要有这个流程：先写kernel，再编译，最后执行并检查错误。每一步都有坑，但每一步也都有对应的工具和方法。

好了，向量加法这个例子虽然简单，但它包含了CUDA编程的所有核心要素：线程模型、内存管理、编译流程、错误检查。你把这个例子跑通了，后面学矩阵乘法、归约这些复杂算子，就会轻松很多。

公众号:蓝海资料掘金营，微信deep3321

📚 课程目录