第3章：CUDA编程基础：CUDA C/C++基础、Kernel函数编写、线程层次结构

各位同学，欢迎来到CUDA编程的实战环节。说实话，很多初学者一上来就被各种并行概念吓住了，觉得GPU编程高不可攀。其实没那么玄乎——你想想看，CUDA C/C++本质上就是在C语言的基础上，加了一些“小插件”，让你能指挥成千上万个线程同时干活。

我个人习惯把CUDA编程比作“包工头带团队”。CPU是那个包工头，负责发号施令；GPU就是那帮干活麻利的工人，一拥而上把活干完。今天我们就来聊聊，怎么当好这个包工头。

3.1 CUDA C/C++基础：从Hello World到设备代码

先看一个最基础的例子。我当年第一次写CUDA程序时，就卡在“怎么让GPU干活”这个问题上。其实答案很简单：用__global__关键字声明一个函数，这就是你的Kernel（内核函数）。

// 一个简单的Kernel函数
__global__ void hello_from_gpu() {
    printf("Hello from GPU thread %d\n", threadIdx.x);
}

int main() {
    // 启动Kernel，<<<1, 10>>>表示1个Block，10个线程
    hello_from_gpu<<<1, 10>>>();
    // 同步，确保GPU执行完毕
    cudaDeviceSynchronize();
    return 0;
}

这里有几个关键点：

• __global__：告诉编译器，这个函数要在GPU上跑，并且能被CPU调用。
• <<<grid, block>>>：这是CUDA的“启动配置”语法，指定你要用多少线程来干活。
• cudaDeviceSynchronize()：CPU等GPU干完活再继续。我曾经忘了加这行，结果CPU跑完了，GPU还在排队，输出全丢了。

3.2 Kernel函数编写：线程的“工作手册”

Kernel函数就是每个线程要执行的代码。你写一份代码，成千上万个线程同时跑。但每个线程怎么知道“我该处理哪块数据”？靠的是内置变量：threadIdx、blockIdx、blockDim。

举个例子，我要把两个数组相加：

__global__ void vector_add(float *A, float *B, float *C, int n) {
    // 计算当前线程的全局索引
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 边界检查，防止越界
    if (idx < n) {
        C[idx] = A[idx] + B[idx];
    }
}

int main() {
    // 假设数组大小是1000
    int n = 1000;
    int block_size = 256;
    int grid_size = (n + block_size - 1) / block_size;  // 向上取整
    
    // ... 分配内存、拷贝数据 ...
    vector_add<<<grid_size, block_size>>>(d_A, d_B, d_C, n);
    // ... 拷贝结果回CPU ...
}

这里有个坑，我踩过好几次：边界检查。因为grid_size是向上取整的，最后一个Block可能只有部分线程在工作。如果不加if (idx < n)，那些多余的线程会访问非法内存，轻则数据错误，重则程序崩溃。

3.3 线程层次结构：Grid、Block、Thread

这是CUDA最核心的概念。我画了一张图，帮你理清关系：

这张图展示了三层结构：

• Grid（网格）：一个Kernel启动时，对应一个Grid。Grid里包含多个Block。
• Block（线程块）：一组线程的集合。同一个Block内的线程可以共享数据（通过共享内存），也能同步。
• Thread（线程）：最小的执行单元。每个线程执行相同的Kernel代码，但处理不同的数据。

每个线程都有一个唯一的ID，计算公式是：

全局线程ID = blockIdx.x * blockDim.x + threadIdx.x

这里blockIdx是Block在Grid中的索引，threadIdx是线程在Block中的索引，blockDim是Block的尺寸（即包含多少个线程）。

3.4 实战经验：如何选择Block大小？

这个问题我当年纠结了很久。其实没有标准答案，但有经验法则：

记住一个原则：Block大小最好是32的倍数。为什么？因为GPU的线程调度是以32个线程为一组（称为Warp）进行的。如果Block大小不是32的倍数，最后一个Warp就会“浪费”一些线程，降低利用率。

3.5 总结：记住这三句话

1. Kernel是GPU的“工作函数”，用__global__声明，用<<<grid, block>>>启动。
2. 线程ID是定位数据的“坐标”，通过blockIdx、threadIdx、blockDim计算。
3. 层次结构是性能的“钥匙”，Block大小选得好，性能翻倍；选得差，GPU空转。

嗯，今天就先聊到这儿。这些基础概念虽然简单，但它们是后续所有优化的基石。你想想看，如果连线程怎么编号都搞不清楚，后面那些共享内存、Warp优化根本无从谈起。所以，多写几个小例子练练手，把这三个内置变量玩熟，后面就顺了。