CUDA系统拆解-03-线程组织模型：grid、block、thread 到底在表达什么

本文是「CUDA系统拆解」系列第 03 篇。
系列导读：CUDA系统拆解-00-导读：从编程模型到 AI 推理系统的学习路线
上一篇：CUDA系统拆解-02-第一个CUDA程序：最小闭环与代码执行路径
下一篇：CUDA系统拆解-04-warp、SIMT 与 SM：真实执行不是“线程各跑各的”

1. 这篇解决什么问题

CUDA 为什么不是“直接启动很多线程”，而是要分成 thread / block / grid 三层。
thread、block、grid 各自到底解决什么问题。
为什么几乎所有基础 kernel 都要先算全局索引，再做边界判断。
什么时候适合用 1D、2D、3D 线程映射。
这种线程组织方式为什么会直接影响后面的 shared memory、同步、GEMM、attention 和推理算子设计。

2. 先记住的核心结论

CUDA 的并行不是“你写很多函数”，而是“你定义一个线程做什么，然后复制成很多线程并行执行”。
thread 是最小计算单位，block 是局部协作单位，grid 是整体任务覆盖范围。
block 的关键价值不只是分组，而是提供同步和 shared memory 的作用边界。
不同 block 默认独立，这个约束既限制了算法写法，也换来了更好的调度和可扩展性。
一维索引公式和边界判断不是模板技巧，而是线程空间映射到数据空间的直接结果。
后面几乎所有高性能 kernel，本质上都要先回答：一个线程处理什么、一个 block 为什么这样协作、grid 如何覆盖整个问题空间。

3. 正文讲解

3.1 CUDA 的并行思维是什么

在 CPU 上，你更常见的是这种思路：

写一个 for 循环
每次迭代处理一个元素
让同一个线程把全部数据做完

在 CUDA 上，思路变成了：

写一个 kernel
定义“一个线程该做什么”
启动大量线程
让每个线程通过自己的线程 ID 去处理不同的数据位置

所以 CUDA 的并行不是“手工写很多份逻辑”，而是：

写一份线程逻辑，然后让大量线程并行执行这份逻辑。

这就是为什么你会不断看到这类代码：

int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
    C[idx] = A[idx] + B[idx];
}

这里真正重要的不是公式本身，而是背后的思维：

先定义线程怎么找到自己的数据
再让很多线程并行处理这些数据

3.2 为什么不是只有线程，而是 `thread / block / grid`

如果 CUDA 只有“线程”这一层，看起来会更简单，但很多关键问题就没法解决。

三层结构各自承担不同职责。

thread

最小计算单位
负责一份局部工作
例如处理一个向量元素、一个矩阵元素、一个像素或一个 tile 内的一小部分数据

block

一组线程组成的局部协作单位
block 内线程可以同步
block 内线程可以共享 shared memory
block 也是资源分配和调度中的重要粒度

grid

一次 kernel launch 的整体工作空间
负责用很多个 block 去覆盖整个问题规模

如果只保留线程，没有 block 这一层，会遇到两个大问题：

没有清晰的局部协作边界
shared memory 和同步语义会变得很混乱

如果只有 block，没有 grid，又很难把任务扩展到更大规模。

所以这三层不是语法拆分，而是把“计算粒度”“局部协作”“整体覆盖”分开表达。

3.3 `thread`、`block`、`grid` 分别解决什么问题

可以把这三层压缩成一张表：

层级	本质角色	主要解决的问题
`thread`	计算单位	单个线程做什么
`block`	协作单位	哪些线程放在一起同步和共享数据
`grid`	覆盖单位	如何把全部数据空间都覆盖到

最重要的一句是：

block 内可以协作，不同 block 默认独立。

这句话后面会不断出现，因为它决定了：

shared memory 的作用范围
__syncthreads() 的有效范围
为什么很多全局问题要拆成多个 kernel

3.4 最基础的 1D 线程映射

最常见的启动方式是：

1
2
3

int blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize;
kernel<<<gridSize, blockSize>>>(...);

这里：

blockSize 是每个 block 的线程数
gridSize 是 block 的数量

总线程数通常是：

1	gridSize * blockSize

这通常会大于等于 N，而不是严格等于 N。

原因很简单：按 block 对齐去组织线程，更符合 GPU 的调度方式，也更方便统一 launch 配置。

这里要额外记住一层：<<<gridSize, blockSize>>> 的本质不是“把两个数字传给 kernel”，而是在声明怎么把数据空间切成很多块去并行覆盖。也就是说，launch 参数本质上是在做数据分解。

线程如何找到自己的全局编号？最经典的 1D 公式就是：

1	int idx = blockIdx.x * blockDim.x + threadIdx.x;

其中：

threadIdx.x：线程在当前 block 内的局部位置
blockIdx.x：当前 block 在整个 grid 中的位置
blockDim.x：每个 block 的线程数

这行代码的意义就是：

把线程空间映射到一维数据空间。

3.5 为什么几乎总要写边界判断

边界判断是和上面的 launch 方式配套出现的：

1	if (idx < N) { ... }

原因是总线程数经常会向上取整。

例如：

N = 1000
blockSize = 256
gridSize = (1000 + 255) / 256 = 4
总线程数 = 1024

这意味着最后会多出 24 个线程。
如果这些线程继续访问 A[idx]、B[idx]、C[idx]，就会越界。

所以边界判断的本质不是“语法习惯”，而是：

允许线程组织按硬件友好的方式对齐
用很小的代价裁掉尾部多出来的线程

这体现了 CUDA 很典型的工程思路：

接受少量冗余线程，换取更统一、更可扩展的线程组织方式。

3.6 一个最小 1D 骨架

下面这段代码已经足够体现 1D 线程组织的核心：

__global__ void vecAdd(const float* A, const float* B, float* C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

从线程组织的角度看，这段代码只回答了两件事：

当前线程处理哪个元素
如果线程越界，该如何安全退出

而真正的并行规模，是由 launch 时的 grid 和 block 决定的。

3.7 为什么很多问题更适合 2D 或 3D 映射

不是所有问题都天然是一维的。

很多问题本身就是二维甚至三维结构，例如：

图像：height × width
矩阵：rows × cols
feature map：H × W
部分张量计算：batch × seq × hidden

这时继续硬用 1D 当然也能写，但代码结构会不自然，后续优化也会更绕。

2D 映射的典型写法是：

1 2	int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y;

例如做矩阵加法时，线程坐标可以直接对应到 (row, col)。

它的好处是：

代码和数据结构更一致
更容易构造 tile
更方便后面配合 shared memory 做局部优化

3D 映射进一步扩展到 z 维，常见于体数据、三维网格，或者把 batch / head / spatial 某一维塞进 z 方向。

不过要记住一点：

2D / 3D launch 的意义主要是更自然地表达问题结构，而不是硬件真的按“二维”或“三维”特殊加速。

3.8 `block` 为什么是协作边界

block 最重要的价值，不是把线程分组，而是让局部线程协作有了明确边界。

block 内线程通常可以做两件事：

用 __syncthreads() 做同步
共享 shared memory

这意味着一个 block 常常对应“一个局部工作块”。

例如在矩阵乘法里：

一个 block 可能负责输出矩阵中的一个 tile
block 内线程一起把输入 tile 搬进 shared memory
然后再在 block 内协作完成计算

如果没有 block 这一层，这种局部协作会非常难组织。

3.9 为什么不同 `block` 默认独立

这也是 CUDA 非常关键的设计取舍。

不同 block 默认独立，有两个明显好处：

第一，调度更灵活。
GPU 可以把不同 block 分发到不同 SM 上执行，不需要让它们彼此强依赖。

第二，可扩展性更强。
同一个 kernel 可以在不同 GPU 上运行，只要 block 之间独立即可，不需要高度绑定具体硬件规模。

代价也很清楚：

block 间默认不能直接做像 block 内那样的同步
很多全局问题不能在一个 kernel 里一次性完成

例如对一个大数组做全局归约，常见做法往往是：

第一个 kernel：每个 block 先做局部归约
第二个 kernel：再把各 block 的部分结果继续归约

这不是写法笨，而是线程组织模型决定的。

3.10 `block size` 为什么不能随便取

block size 既不能太小，也不能太大。

太小的问题：

并行度不够
很难把 warp 级执行吃满
调度粒度不划算

太大的问题：

会吃掉更多寄存器和 shared memory
可能降低一个 SM 上能同时驻留的 block 数
影响 occupancy 和调度灵活性

所以 block size 本质上是在平衡：

并行度
协作粒度
资源占用
调度灵活性

为什么常见值经常是 128 / 256 / 512？

它们通常是 32 的倍数，便于和 warp 粒度对齐
对很多 kernel 来说是比较实用的经验起点

但要记住：

block size 不是越大越好，也不是固定答案。它最终要和 kernel 的寄存器、shared memory、访存模式一起看。

4. 和 AI 推理的关系

线程组织模型和 AI 推理的关系非常直接，因为后面你看到的几乎所有推理 kernel，都先要回答这个问题：

一份数据由哪个线程处理
一个 block 为什么这样分工
哪些线程需要局部协作
grid 如何覆盖完整问题空间

例如：

GEMM

grid 覆盖输出矩阵的所有 tile
block 负责一个 tile
thread 负责 tile 内的一部分结果

attention

某些维度会映射到 batch、head、query block、key block
block 内线程共同处理一个 attention tile

decode 阶段的部分 kernel

可能按 token、head、hidden slice 或 cache block 做映射

所以这篇不是在学“基础语法”，而是在学：

以后所有 CUDA 算子和推理 kernel 的最底层骨架。

5. 常见误区

thread / block / grid 只是语法层次。不是，它们分别对应计算粒度、协作边界和整体覆盖范围。
block 只是为了把线程分组。不是，block 最关键的是同步和 shared memory 的作用边界。
if (idx < N) 只是模板。不是，它是向上取整 launch 配置的自然结果。
只会 1D 索引就够了。不是，矩阵、图像、张量问题经常更适合 2D 或 3D 映射。
不同 block 之间也能像 block 内一样直接协作。默认不行，这正是很多算法需要拆成多个 kernel 的原因。

6. 复习自测

为什么 CUDA 不是只有线程，而是要分成 thread / block / grid 三层？
thread、block、grid 各自解决什么问题？
一维索引公式为什么写成 blockIdx.x * blockDim.x + threadIdx.x？
为什么边界判断几乎总会出现？
为什么 block 既是组织单位，又是协作单位？
为什么不同 block 默认独立，这种设计带来了什么好处和限制？
什么时候 2D 映射会比 1D 更自然？
如果把一个 GEMM 或 attention kernel 拿出来，你会先从哪些线程组织问题开始分析？

系列导航

导读：CUDA系统拆解-00-导读：从编程模型到 AI 推理系统的学习路线
上一篇：CUDA系统拆解-02-第一个CUDA程序：最小闭环与代码执行路径
下一篇：CUDA系统拆解-04-warp、SIMT 与 SM：真实执行不是“线程各跑各的”
延伸阅读：GPU系统拆解-05-CUDA 执行映射：抽象模型如何落到真实硬件
延伸阅读：CUDA C编程权威指南：导读