4. 线程组织与索引映射

在 CUDA 中，核函数并不是以“一个线程”为单位被调用的，而是以线程块（block）为调度单位、以网格（grid）为整体范围被启动。真正决定每个线程处理哪一份数据的，是线程在网格中的坐标以及由这些坐标计算出的全局线程 ID。理解这套索引映射机制，是写出正确且高效 CUDA 程序的基础。

一、从线程坐标到全局线程 ID

CUDA 为每个线程提供了一组内建变量：threadIdx 表示线程在块内的坐标，blockIdx 表示线程块在网格中的坐标，blockDim 表示每个块的维度大小，gridDim 表示网格的维度大小。它们都是 dim3 类型，可以只使用 .x，也可以使用 .x/.y/.z 三个分量。

对于一维线程组织，全局线程 ID 的计算公式为：

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

这个公式的含义是：先跳过当前块之前的所有块所包含的线程，再加上当前线程在块内的偏移。例如 blockDim.x = 256，blockIdx.x = 3，threadIdx.x = 17，则 gid = 3 * 256 + 17 = 785。只要网格覆盖的线程总数不小于数据元素个数，每个元素就能被唯一一个线程处理。

二、多维网格与块

当处理矩阵、图像或三维体数据时，使用二维或三维线程组织更自然。二维情况下，全局坐标分别为：

int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;

这里约定：x 方向对应列（连续内存方向），y 方向对应行。这个约定必须全文一致，否则访问顺序会与内存布局错位。对于按行优先存储的矩阵，元素 (row, col) 的线性下标为 row * width + col。三维情况同理，增加 z 分量即可。

选择块维度时，通常让 blockDim.x 为 32 的倍数，以匹配一个 warp 的线程数。例如二维块常用 dim3 block(32, 8)，共 256 个线程；网格维度则按数据规模向上取整：

dim3 block(32, 8);
dim3 grid((width + block.x - 1) / block.x,
          (height + block.y - 1) / block.y);

三、边界处理：向上取整带来的越界风险

由于网格维度必须覆盖数据范围，而数据尺寸往往不是块大小的整数倍，最后一行的块中会存在“多余线程”。这些线程的全局坐标已经超出数据边界，如果直接访问就会越界。标准做法是在核函数入口处做边界判断：

if (row < height && col < width) {
    int idx = row * width + col;
    out[idx] = a[idx] + b[idx];
}

边界判断必须放在任何内存访问之前。对于只读输入，越界读取可能返回垃圾值；对于写入，越界写会破坏其他数据甚至导致程序崩溃。因此“先判断、后访问”是必须遵守的纪律。

四、矩阵与网格数据的线程映射策略

线程映射的核心目标是让相邻线程访问相邻内存，从而获得合并访问（coalesced access）。对于行优先矩阵，应让 threadIdx.x 沿列方向变化，这样同一 warp 内的线程访问的是连续地址。若把 threadIdx.x 用作行方向，则相邻线程访问的地址间隔为 width，合并效率会显著下降。

对于二维网格类数据（如有限差分、模板计算），每个线程通常需要读取自身及上下左右邻居。此时除了全局索引映射，还应结合共享内存做分块（tiling）：先把一块数据连同边界 halo 区载入共享内存，再在共享内存上完成邻域计算。这样既减少了全局内存访问次数，也避免了同一数据被多个线程重复读取。

二维线程映射与边界判断示意 数据矩阵 (width=10, height=6) 有效数据区域：row < 6 且 col < 10 线程块覆盖 (block=4×4) 越界线程 越界线程 映射 索引计算与边界判断 col = blockIdx.x * blockDim.x + threadIdx.x row = blockIdx.y * blockDim.y + threadIdx.y if (row < height && col < width) { ... } // 先判断后访问

五、避免分支发散

边界判断本身会引入分支：满足条件的线程执行计算，不满足的线程跳过。在一个 warp 内，如果部分线程走 if 分支、部分线程走 else 分支，硬件会串行执行两条路径，这就是分支发散（branch divergence）。边界判断造成的发散通常只出现在最后一列或最后一行的块中，影响有限，属于可接受的开销。

真正需要警惕的是在循环内部、且条件与线程索引强相关的分支。例如在模板计算中，若对每个线程都判断“是否处于边界”，则内部线程与边界线程会走不同路径，导致整个 warp 反复串行。更好的做法是：把边界区域单独处理，或者用共享内存 halo 区把边界条件“填平”，让主体循环对所有线程走同一条路径。

此外，应避免用线程索引作为分支条件去决定是否执行主要计算。如果确实需要按数据特征分支，尽量让同一 warp 内的线程具有相同的数据特征，例如按行分块、按区域划分线程块，使分支在 warp 粒度上保持一致。

六、一个可复用的映射检查清单

在编写核函数前，可以按以下要点逐项确认：

  • 数据尺寸与块尺寸是否已知，网格维度是否按向上取整计算；
  • 全局索引公式是否与内存布局一致（行优先时 x 对应列）；
  • 所有全局内存访问前是否都有边界判断；
  • 相邻线程是否访问相邻地址，是否满足合并访问；
  • 是否存在与线程索引强相关的分支，能否通过分块或 halo 区消除；
  • 共享内存分块时，halo 区是否已正确加载并同步。
提示：当数据尺寸恰好是块大小的整数倍时，边界判断仍然建议保留，因为它不会带来明显开销，却能让核函数在尺寸变化时依然正确。

线程组织与索引映射看似只是几行坐标计算，但它直接决定了程序的正确性、内存访问效率和分支行为。把全局 ID 公式、边界判断和合并访问这三件事做扎实，后续的共享内存优化、流水线并行和性能调优才有稳固的基础。

← 上一章 📖 返回目录 下一章 →