线程与块配置：Grid与Block维度设计、线程索引计算、Occupancy优化基础

好，咱们进入正题。今天聊的是CUDA编程里最基础、也最容易被忽视的一环——线程与块的配置。说实话，我见过太多人一上来就写 dim3 block(256, 1, 1)，然后就不管了。嗯，这样跑小数据还行，一旦数据量上去，性能立马崩盘。

我个人习惯是，写任何kernel之前，先花5分钟想清楚三个问题：数据怎么分？线程怎么排？SM能塞下多少？ 这三个问题搞定了，后面调优就顺风顺水。

1. Grid与Block的维度设计

先看个最简单的例子。假设你有1024个元素要做向量加法。你会怎么分？

// 方案A：一维网格，一维块
dim3 block(256, 1, 1);
dim3 grid((1024 + 255) / 256, 1, 1);

// 方案B：二维网格，二维块
dim3 block(16, 16, 1);
dim3 grid(64, 1, 1);  // 16*64 = 1024

两种方案都能跑。但性能一样吗？不一样。方案A的block是256个线程，方案B是256个线程。看起来一样，但内存访问模式完全不同。

我建议你记住一个原则：Block的维度尽量匹配数据的内存布局。比如处理二维图像，用二维block；处理一维数组，用一维block。别为了炫技搞花里胡哨的维度。

2. 线程索引计算——别搞错了

线程索引计算是新手最容易翻车的地方。我刚开始学CUDA时，就因为在二维block里算错了索引，debug了整整一个下午。

先看标准公式：

// 一维情况
int tid = blockIdx.x * blockDim.x + threadIdx.x;

// 二维情况
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
int tid = row * width + col;  // 行优先

// 三维情况
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
int tid = z * height * width + y * width + x;

这里有个坑：千万别把row和col搞反了。C/C++是行优先存储，所以row对应y方向，col对应x方向。我见过有人写成 tid = col * width + row，结果数据全乱套了。

3. Occupancy优化基础——别让SM闲着

Occupancy，说白了就是SM上活跃的warp数量占最大容量的比例。你想想看，SM就像个工厂，warp就是工人。工人越多，工厂越忙，产出越高。但工人太多，资源（寄存器、共享内存）就不够分了。

Occupancy的计算公式：

Occupancy = 活跃warp数 / 最大warp数

影响Occupancy的因素有三个：

• 每个线程使用的寄存器数量——寄存器越多，能塞的线程越少
• 每个block使用的共享内存大小——共享内存越大，能放的block越少
• block的大小——block太小，warp数上不去

我举个例子。假设你的SM最多能放64个warp（2048个线程），每个block有256个线程（8个warp）。那么最多能放8个block（64/8=8）。但如果每个block用了太多寄存器，可能只能放4个block，Occupancy就掉到50%了。

4. 实战：如何选择block大小

选block大小，我有个三步法：

1. 先看数据量：数据量大，block可以大一点（512或1024）；数据量小，block小一点（128或256）
2. 再看资源需求：如果每个线程用很多寄存器或共享内存，block要小一点
3. 最后看warp效率：block大小最好是32的倍数，避免warp中有空闲线程

下面这张图展示了不同block大小对Occupancy的影响：

从图上能看出来，block为256或512时，Occupancy通常最高。但如果你每个线程用了很多寄存器（橙色柱子），那block为1024反而会降低Occupancy，因为寄存器不够用了。

5. 一个完整的例子

最后，我写个完整的向量加法kernel，把上面说的都串起来：

// 向量加法kernel
__global__ void vecAdd(float* A, float* B, float* C, int N) {
    // 计算全局线程索引
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 边界检查
    if (tid < N) {
        C[tid] = A[tid] + B[tid];
    }
}

// 调用时
int N = 1 << 20;  // 1M个元素
int blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize;

vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);

这里blockSize选256，是因为：

• 256是32的倍数，warp不会浪费
• 每个线程只用几个寄存器，Occupancy能到100%
• gridSize是4096，远小于设备限制

好了，线程与块配置这块就聊这么多。记住：没有银弹。每个kernel的最优配置都不一样，多试几次，用profiler看看，你就能找到最适合你的方案。

想要更多资料，联系可可