第4章 CUDA编程模型

4、CUDA编程模型：线程层次结构（Grid、Block、Thread）、核函数定义、<<<>>>语法

好，咱们今天聊聊CUDA编程里最核心、也最绕不开的一个概念——线程层次结构。

说实话，我刚开始学CUDA的时候，最头疼的就是搞不清Grid、Block、Thread这三者到底怎么分工。你想想看，CPU那边就几个核，顶多十几个线程，到了GPU这边，一下子冒出成千上万个线程，怎么管理？怎么调度？

嗯，CUDA的设计者想了个很聪明的办法——分层管理。

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

说白了，CUDA把所有的并行线程组织成了三层结构：

Thread（线程）：最底层的工作单元，每个线程执行核函数中的一条指令。
Block（线程块）：一组线程的集合，同一个Block内的线程可以共享数据（通过共享内存），也可以同步。
Grid（网格）：一组Block的集合，一个Grid对应一个核函数的执行。

我打个比方你就明白了。假设你要给一张1000x1000的图片做滤镜处理：

每个像素点需要一个线程来处理——这就是Thread。
把图片切成若干个小块，比如16x16的小方块，每个小方块内的像素由同一个Block里的线程处理——这就是Block。
整张图片就是Grid，它包含了所有的小方块——这就是Grid。

为什么会这样设计？

我在项目中遇到过一个问题：如果所有线程都直接暴露在同一个层级，那共享内存和同步的开销会变得极其复杂。分层之后，Block内部的线程可以高效协作，而Block之间则相对独立，这样GPU硬件就能轻松调度了。

关键点：

一个Grid可以有多个Block
一个Block可以有多个Thread
Block内部的线程可以通过 __syncthreads() 同步
不同Block之间的线程无法直接同步

4.2 内置变量：threadIdx、blockIdx、blockDim、gridDim

每个线程在运行时，都能通过一些内置变量知道自己是谁、在哪干活。这些变量是CUDA自动提供的，你直接拿来用就行。

变量名	类型	含义
`threadIdx`	`dim3`	当前线程在Block内的索引（x, y, z）
`blockIdx`	`dim3`	当前Block在Grid内的索引（x, y, z）
`blockDim`	`dim3`	Block的维度，即每个Block有多少个线程
`gridDim`	`dim3`	Grid的维度，即每个Grid有多少个Block

举个例子，如果你启动了一个Grid，里面有2x3个Block，每个Block是4x4个线程，那么：

gridDim.x = 2, gridDim.y = 3
blockDim.x = 4, blockDim.y = 4
某个线程的 blockIdx.x 可能是0或1，threadIdx.x 可能是0到3

我个人习惯用一维索引来处理数组，这样代码更简洁。计算全局线程ID的公式是：

int tid = blockIdx.x * blockDim.x + threadIdx.x;

如果是二维的，那就稍微复杂一点：

int idx = blockIdx.x * blockDim.x + threadIdx.x;
int idy = blockIdx.y * blockDim.y + threadIdx.y;
int tid = idy * width + idx;  // 假设width是每行的元素数

4.3 核函数定义：global 关键字

核函数，就是跑在GPU上的函数。定义的时候要在前面加上 __global__ 修饰符。

语法规则很简单：

返回类型必须是 void
不能使用可变参数（比如 ...）
不能是类的成员函数（静态成员函数可以）
在主机端调用，在设备端执行

来看一个最经典的例子——向量加法：

// 核函数定义
__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];
    }
}

注意那个 if (i < N) 的判断。为什么要加这个？

嗯，这里有个坑。我曾经在项目里忘了加边界检查，结果数组越界，读到了脏数据，调试了一下午才发现。你想想看，我们启动的线程数往往是N的整数倍，比如N=1000，我们启动了1024个线程，那多出来的24个线程就会访问到数组外面。所以边界检查是必须的。

注意：核函数内部不能调用 printf（调试时可以，但正式代码不要用），不能使用C++标准库（比如 std::vector），也不能使用递归。

4.4 <<<>>> 语法：启动核函数

核函数写好了，怎么让它跑起来？用 <<<grid, block>>> 语法。

这个语法看起来有点奇怪，我第一次看到的时候也愣了一下。它的作用是指定：

第一个参数：Grid的维度，即有多少个Block
第二个参数：Block的维度，即每个Block有多少个线程

两个参数都是 dim3 类型。如果你只传一个整数，CUDA会默认另外两个维度是1。

调用方式：

// 假设我们要处理N个元素
int N = 1000;
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

// 分配设备内存并拷贝数据（略）

// 启动核函数
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

// 等待GPU执行完成
cudaDeviceSynchronize();

这里有个小技巧：(N + threadsPerBlock - 1) / threadsPerBlock 这个公式是向上取整，保证即使N不能被threadsPerBlock整除，也能覆盖所有元素。我每次写CUDA代码都会用这个公式，已经形成肌肉记忆了。

另外，<<<>>> 还可以传第三个参数——共享内存大小（以字节为单位），以及第四个参数——流（stream）。不过初学者先掌握前两个就够了。

个人建议：Block的线程数最好是32的倍数。为什么？因为GPU的warp大小是32，一个warp里的线程是同时执行的。如果线程数不是32的倍数，就会有空闲的CUDA核心，造成浪费。我一般用128、256或512，这几个值在大多数GPU上表现都不错。

4.5 完整的例子：向量加法

把上面所有的知识点串起来，写一个完整的向量加法程序：

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

// 核函数定义
__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 = 1000;
    size_t size = N * sizeof(float);

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

    // 初始化数据（略）

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

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

    // 启动核函数
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

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

    // 验证结果（略）

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

    return 0;
}

这个例子虽然简单，但麻雀虽小五脏俱全。你把这个流程跑通了，CUDA编程的基本功就算打牢了。

4.6 避坑指南

最后，分享几个我踩过的坑：

忘记同步：我曾经在核函数后面忘了加 cudaDeviceSynchronize()，结果主机端代码先执行了，读到的数据全是错的。记住，核函数是异步的，主机不会等它执行完。
Block太大：每个Block的线程数不能超过设备限制（一般是1024）。我刚开始写代码时设了2048，结果程序直接崩溃。
Grid太大：Grid的维度也有上限，一般是65535（老设备）或2^31-1（新设备）。不过一般用不了那么多。
忘记释放内存：cudaMalloc 分配的内存一定要用 cudaFree 释放，不然会内存泄漏。这个和 malloc/free 是一个道理。

好了，这一章的内容就到这里。线程层次结构是CUDA的基石，搞懂了它，后面的共享内存、同步、性能优化才能顺利展开。下一章我们聊聊共享内存和线程同步，那才是真正体现CUDA威力的地方。

📚 课程目录

4、CUDA编程模型：线程层次结构（Grid、Block、Thread）、核函数定义、<<<>>>语法

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

4.2 内置变量：threadIdx、blockIdx、blockDim、gridDim

4.3 核函数定义：__global__ 关键字

4.4 <<<>>> 语法：启动核函数

4.5 完整的例子：向量加法

4.6 避坑指南

4.3 核函数定义：global 关键字