4. 线程层次结构：Grid、Block、Thread

好，咱们直接进入正题。GPU 编程里最核心的概念，就是线程怎么组织。说白了，就是你怎么告诉 GPU：「这里有 100 万个任务，你帮我分一下，每人干一小块。」

我个人习惯把 GPU 的线程层次想象成一个军队：Thread 是士兵，Block 是班排，Grid 是整个军团。每个士兵干自己的活，但大家协同作战。

4.1 三个基本概念

• Thread（线程）：最小的执行单元。每个线程执行同一个 kernel 函数，但处理不同的数据。
• Block（线程块）：一组线程的集合。同一个 block 内的线程可以共享内存（shared memory），也能同步（__syncthreads()）。
• Grid（线程网格）：所有 block 的集合。一个 kernel 启动时，就对应一个 grid。

4.2 一维/二维/三维索引计算

你想想看，我们怎么知道「我是谁」？GPU 提供了几个内置变量，帮你定位：

• threadIdx：线程在 block 内的索引
• blockIdx：block 在 grid 内的索引
• blockDim：block 的维度大小（每个 block 有多少线程）
• gridDim：grid 的维度大小（有多少个 block）

嗯，这里要注意：这些变量可以是三维的（x, y, z）。但实际硬件上，z 维度用得少，x 和 y 最常见。

一维索引

最简单的情况。假设你有一个数组，长度是 N。每个线程处理一个元素：

int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < N) {
    data[tid] = data[tid] * 2;
}

这个公式我建议你背下来。我在项目中遇到过很多新手，上来就写 threadIdx.x 当全局索引，结果数据全乱了。

二维索引

处理图像、矩阵时，二维索引更自然。比如一张 1024x768 的图片：

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

if (row < height && col < width) {
    output[idx] = input[idx] * 2.0f;
}

说白了，二维就是把一维的公式在 y 方向再复制一份。每个线程负责一个像素点。

三维索引

三维用得少，但处理体数据（比如 CT 扫描）时会用到。原理一样：

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 idx = z * width * height + y * width + x;

4.3 内置变量详解

变量	类型	说明
threadIdx	dim3	线程在 block 内的索引（x, y, z）
blockIdx	dim3	block 在 grid 内的索引（x, y, z）
blockDim	dim3	block 的尺寸，即每个 block 有多少线程
gridDim	dim3	grid 的尺寸，即有多少个 block

注意：dim3 是个结构体，有 .x, .y, .z 三个成员。如果你只设了一维，那 .y 和 .z 默认是 1。

4.4 线程束（Warp）概念

这是 GPU 最底层的执行单位。一个 warp 包含 32 个线程。硬件以 warp 为单位调度指令。

为什么会这样？因为 GPU 是 SIMT（单指令多线程）架构。说白了，32 个线程同时执行同一条指令，只是处理的数据不同。

我曾经踩过一个坑：在 warp 内部做 if-else 分支。如果 32 个线程走不同的分支，那性能会急剧下降。因为硬件只能串行执行这些分支，最后再合并结果。这叫「warp divergence」。

嗯，这里有个实用技巧：block 的线程数最好是 32 的倍数。比如 128、256、512。这样不会浪费 warp 的槽位。我个人习惯用 256，因为大多数 GPU 上这个值效率最高。

4.5 知识结构图

下面这张图帮你理清线程层次的关系：

这张图里，grid 包含 2 个 block，每个 block 包含若干线程。每个线程通过 blockIdx 和 threadIdx 的组合，就能算出自己在全局数据中的位置。

4.6 总结

线程层次结构是 GPU 编程的基石。记住三点：

1. Thread → Block → Grid，层层嵌套
2. 全局索引 = blockIdx × blockDim + threadIdx，这个公式要刻在脑子里
3. Warp 是 32 个线程一组，写代码时要考虑 warp 对齐和分支问题

我个人觉得，理解了这个层次结构，你就掌握了 GPU 编程的一半。剩下的就是各种优化技巧和 API 调用了。嗯，下一节我们会聊内存层次结构，那又是另一个有趣的话题。

---