线程层次结构：Grid、Block、Thread 的概念与配置

聊到 CUDA 编程，线程层次结构是绕不开的核心概念。说白了，你得先搞清楚你的「工人」是怎么组织起来的，才能指挥他们干活。

我个人习惯把 CUDA 的线程层次想象成一个公司架构：Grid 是总公司，Block 是部门，Thread 就是每个员工。员工（Thread）在部门（Block）里干活，部门又在总公司（Grid）的管辖下。每个员工都有自己独立的工位编号，部门也有自己的编号，这样你就能精准地找到任何一个员工。

一、Grid、Block、Thread 是什么？

先看官方定义：

• Thread（线程）：CUDA 执行的最小单位。每个线程执行相同的 kernel 函数，但处理不同的数据。
• Block（线程块）：一组线程的集合。Block 内的线程可以通过共享内存通信，也能同步执行。
• Grid（线程网格）：一组 Block 的集合。一个 kernel 启动时，就对应一个 Grid。

嗯，这里要注意：Block 内的线程数量是有限制的。我记得在 Fermi 架构时代，每个 Block 最多 512 个线程；到了 Kepler 之后，上限提升到了 1024。你写代码时别超了，否则 kernel 启动会失败。

二、一维/二维/三维线程索引计算

你想想看，如果只有一维索引，处理一张 1920×1080 的图片得多麻烦？你得手动把 (x, y) 坐标映射成一个线性索引。CUDA 直接支持多维索引，省心多了。

1. 一维索引

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

// 一维 Grid，一维 Block
int tid = blockIdx.x * blockDim.x + threadIdx.x;

// 示例：处理 1024 个元素
// 配置：gridDim.x = 4, blockDim.x = 256
// 线程 0 到 1023 分别对应 tid = 0 到 1023

我在项目中遇到过一个问题：数组长度不是 blockDim.x 的整数倍怎么办？比如 N=1000，blockDim.x=256，那最后一个 Block 只有 232 个线程在工作。你需要在 kernel 里加一个边界检查：

if (tid < N) {
    // 处理数据
}

2. 二维索引

处理图像或矩阵时，二维索引更直观。每个线程对应一个像素或矩阵元素：

// 二维 Grid，二维 Block
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int tid = y * width + x;  // 线性化索引

// 示例：处理 1920×1080 图像
// 配置：gridDim = (1920/32, 1080/32), blockDim = (32, 32)
// 每个线程处理一个像素 (x, y)

为什么用二维？因为图像本身就是二维的。你用一维索引也能算，但代码可读性差，还容易算错。我个人习惯：数据是几维的，线程就用几维。

3. 三维索引

处理体素数据（比如 CT 扫描、3D 卷积）时，三维索引就派上用场了：

// 三维 Grid，三维 Block
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 * width * height + y * width + x;  // 线性化索引

// 示例：处理 256×256×256 体素
// 配置：gridDim = (8, 8, 8), blockDim = (32, 32, 32)
// 每个线程处理一个体素 (x, y, z)

三、如何配置 Grid 和 Block？

配置没有固定公式，但有几条经验法则：

举个例子：你要处理 100 万个元素。如果每个 Block 用 256 个线程，那需要 1000000 / 256 ≈ 3907 个 Block。你可以配置 gridDim.x = 3907，blockDim.x = 256。但 3907 不是 SM 数量的整数倍（比如 80 个 SM），那可以向上取整到 4000，然后 kernel 里加边界检查。

四、知识体系结构图

下面这张图帮你理清线程层次的核心逻辑：

五、避坑指南与个人经验

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

• Block 内线程数别超 1024：这是硬件限制。我曾经试图用 blockDim = (32, 32, 2) 共 2048 个线程，结果 kernel 启动直接报错。
• Grid 维度别太大：gridDim.x 最大是 2^31 - 1，但实际中别超过几百万，否则调度开销会很大。
• 二维/三维索引的线性化：如果你需要把多维索引转成一维，注意行优先还是列优先。CUDA 默认是行优先（row-major）。
• 调试时先用小数据：我习惯先用 16×16 的 Block 和 2×2 的 Grid 测试，确认索引计算正确后再放大。

好了，线程层次结构就聊到这儿。记住那个公司架构的比喻，写代码时多想想你的「员工」是怎么被组织起来的，就不会迷路。