第四章：CUDA编程模型——线程层次结构与核函数实战
各位同学，今天我们来啃一块硬骨头——CUDA的线程层次结构。说实话，我刚接触CUDA时，也被Grid、Block、Thread这三层结构搞得晕头转向。但等你真正理解了，就会发现这设计其实非常巧妙。
4.1 线程层次结构：从宏观到微观
CUDA的线程组织方式，说白了就是三级管理：
- Grid（网格）：最顶层，对应一个核函数启动的所有线程
- Block（线程块）：中间层，一组可以协作的线程
- Thread（线程）：最底层，实际执行计算的最小单元
我习惯把Grid想象成一个工厂，Block是车间，Thread就是工人。每个工人干自己的活，同一个车间的工人可以互相借工具（共享内存），不同车间之间就不能直接借了。
关键点：一个Block内的线程可以同步，可以通过共享内存通信。不同Block的线程，理论上完全独立。
4.2 线程索引与维度
每个线程都有唯一的ID。CUDA提供了几个内置变量：
| 变量 | 含义 | 示例 |
|---|---|---|
| threadIdx | 线程在Block内的索引 | threadIdx.x, threadIdx.y |
| blockIdx | Block在Grid内的索引 | blockIdx.x, blockIdx.y |
| blockDim | Block的维度大小 | blockDim.x = 256 |
| gridDim | Grid的维度大小 | gridDim.x = 1024 |
计算全局线程ID的公式，我建议你直接背下来：
int tid = blockIdx.x * blockDim.x + threadIdx.x;为什么会这样？你想想看，每个Block有blockDim.x个线程，第blockIdx.x个Block的起始位置就是blockIdx.x * blockDim.x，再加上线程在Block内的偏移threadIdx.x，就是全局ID了。
我的习惯：处理一维数据时，我通常只用x维度。二维或三维数据时，再用y和z。别把问题搞复杂了。
4.3 内存模型：搞清楚数据在哪
CUDA的内存模型，嗯，这里要注意，它和CPU的内存模型差别很大。我刚开始写CUDA程序时，经常因为数据放错地方导致性能惨不忍睹。
主要的内存类型：
- 全局内存（Global Memory）：所有线程都能访问，但速度最慢。我把它比作「仓库」。
- 共享内存（Shared Memory）：同一个Block内的线程共享，速度很快。相当于「车间里的工具台」。
- 寄存器（Register）：每个线程私有，速度最快。就是「工人手里的螺丝刀」。
- 常量内存（Constant Memory）：只读，有缓存。适合存放查找表之类的数据。
- 本地内存（Local Memory）：线程私有，但实际存在全局内存里。寄存器不够用时编译器会自动用。
我曾经踩过的坑：以为共享内存是全局的，结果不同Block之间想共享数据，折腾了半天才发现根本行不通。记住：共享内存只在Block内可见！
4.4 核函数（Kernel）编写实战
核函数就是用__global__修饰的函数，在GPU上执行。写核函数有几个铁律：
- 返回类型必须是
void - 不能使用可变参数
- 不能是类的成员函数（静态成员函数可以）
- 调用的设备函数要用
__device__修饰
来看一个最简单的向量加法：
__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)的边界检查。我见过太多新手忘记这个，结果数组越界，程序崩溃还找不到原因。
4.5 <<<>>>语法详解
这个语法是CUDA特有的，用来配置核函数的执行参数。格式是：
kernel_name<<<grid_dim, block_dim, shared_mem, stream>>>(args);四个参数：
| 参数 | 含义 | 示例 |
|---|---|---|
| grid_dim | Grid的维度（dim3类型） | dim3 grid(1024, 1, 1) |
| block_dim | Block的维度（dim3类型） | dim3 block(256, 1, 1) |
| shared_mem | 动态共享内存大小（字节），默认0 | sizeof(float) * 256 |
| stream | CUDA流，默认0（默认流） | stream1 |
实际调用：
int N = 1 << 20; // 约100万个元素
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);这里有个小技巧：(N + threadsPerBlock - 1) / threadsPerBlock这个公式，保证即使N不能被threadsPerBlock整除，也能启动足够的Block覆盖所有数据。
我个人建议：Block大小选128或256通常比较稳妥。太小了硬件利用率低，太大了可能受限于寄存器数量。我一般从256开始调优。
4.6 实战避坑指南
最后分享几个我这些年积累的经验：
- 线程同步：同一个Block内用
__syncthreads()同步。不同Block之间别想同步，那是设计上不允许的。 - 内存对齐：全局内存访问尽量对齐，128字节对齐最好。不对齐的话，性能可能掉一个数量级。
- Bank Conflict：共享内存有32个Bank，同时访问同一个Bank会导致冲突。我当年调优时，光解决这个问题就花了两天。
- Occupancy：不是Block越大越好。要平衡每个SM上活跃的线程数和每个线程使用的资源。
好了，这一章的内容就到这里。记住，CUDA编程的核心就是理解线程怎么组织、数据怎么流动。把这些搞清楚了，后面写优化代码就会顺手很多。