GPU系统拆解-05-CUDA 执行映射：抽象模型如何落到真实硬件

本文是「GPU系统拆解」系列第 05 篇。
系列导读：GPU系统拆解-00-导读：从架构认知到推理系统的学习路线
上一篇：GPU系统拆解-04-数据中心架构：为什么大厂更偏好 Hopper、Blackwell 与 MI300
下一篇：GPU系统拆解-06-内存层次与访存主线：GPU 性能为什么常常卡在数据

这一篇要解决的是一个基础但很容易混乱的问题：你写下的 thread / block / grid，到底是怎么在真实 GPU 上跑起来的。学完这一篇之后，你应该能把 thread、warp、block、SM 放到同一张图里理解，并知道它们为什么会直接影响 kernel 性能和推理系统设计。

1. 先给结论

• thread 是最小逻辑工作单位，但不是 GPU 最重要的硬件调度单位。
• GPU 真正高频调度和发射指令时，更关键的粒度是 warp。
• block 不是多余中间层，它决定了调度、局部同步、shared memory 归属和可扩展性。
• SM 不是“一个大核”，而是能同时维护很多 warp 上下文的吞吐型执行资源池。
• GPU 隐藏延迟的核心不是把单线程做得像 CPU 一样快，而是让一个 SM 上同时挂很多可切换的 warp。
• occupancy 很重要，但只是手段，不是目标；真正目标是整体吞吐更高。
• 对 AI infra / 推理 来说，很多 kernel 设计问题，本质上都是在选择合适的 thread / warp / block 粒度。

2. 为什么 CUDA 不让你直接写“SM 程序”

一个很自然的问题是：既然 kernel 最终在 SM 上执行，为什么 CUDA 不让你直接指定“这个任务给哪个 SM”？

原因很简单：那样会让代码几乎不可移植，也不利于扩展。

不同 GPU 之间会变化的东西很多：

• SM 数量
• 每个 SM 的寄存器和 shared memory 容量
• L1 / L2 组织方式
• Tensor Core 路径和执行细节
• 最优 block size 和最优 tile 形状

如果 CUDA 让程序员直接绑定具体硬件细节，那么一份代码很难跨设备复用。

所以 CUDA 的设计思路是：

• 给程序员稳定的并行抽象
• 让编译器、runtime、driver 去做映射
• 保留足够多的优化空间，但不把你绑死在某代 GPU 上

可以把它理解成一种平衡：

程序员写的是“逻辑并行”，系统再把它翻译成“硬件执行”。

3. 从一个最常见的 kernel 开始

最常见的 CUDA kernel 会长这样：

__global__ void vec_add(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];
    }
}

如果只停留在语法层面，你会说：

• 有很多 thread
• 每个 thread 算一个元素
• blockDim.x 决定每个 block 有多少线程
• gridDim.x 决定总共有多少 block

这没错，但还不够。

真正重要的是：这段代码在 GPU 上并不是“所有线程同时独立起飞”，而是经历了一层层映射。

4. 从 kernel launch 到 GPU 执行，实际发生了什么

当你写下：

vec_add<<<grid_size, block_size>>>(d_a, d_b, d_c, n);

背后至少有四步。

4.1 Host 提交工作描述

CPU 侧通过 CUDA Runtime 发起 launch，请求里会包含：

• kernel 是谁
• grid 维度
• block 维度
• 参数
• stream 信息

这一步不是 CPU 在执行 kernel，而是 CPU 在向 GPU 提交一份工作描述。

4.2 Runtime / Driver 做准备

接下来 runtime 和 driver 会负责：

• 找到对应的 device code
• 打包参数
• 检查 launch 配置
• 处理 stream 依赖
• 管理执行上下文

你平时最常写的是 Runtime API，但更底层的系统框架经常会直接接触 Driver API，因为它控制更细。

4.3 Grid 不会被“一次性全部展开”

很多人刚学 CUDA 时会误以为：我有一百万个线程，GPU 就会把这一百万个线程一次性全跑起来。

真实情况不是这样。

更接近实际的描述是：

1. grid 里有很多 block
2. GPU 把这些 block 分批装载到各个 SM
3. 一个 SM 同时驻留若干个 block
4. 每个 block 再被切成多个 warp
5. warp scheduler 从可执行的 warp 中挑一个发射指令

所以真正的执行链条是：

grid -> block -> warp -> instruction issue on SM

4.4 真正“活起来”的不是孤立 thread

这句话要单独记住：

线程是你写代码时最直观的抽象，但硬件真正推进执行时，看的不是孤立 thread，而是 block 和 warp。

5. 为什么 block 这一层必不可少

很多人会问：既然硬件最终按 warp 调度，为什么 CUDA 还要多设计一层 block？

因为 block 解决的是和 warp 完全不同的问题。

5.1 block 是调度装载单位

GPU 不会把单个线程随意扔到不同 SM。真正被装到某个 SM 上执行的是 block。

这样做的好处是：

• 调度开销可控
• 一组线程天然归属于同一个局部协作域
• 资源边界清晰
• 任务容易拆分成大量独立工作包

5.2 block 是局部同步边界

__syncthreads() 只在 block 内有效，这不是限制太多，而是 CUDA 明确给出的设计边界。

这意味着：

• block 内线程可以同步
• block 内线程可以协作
• block 间默认独立

这样调度器才有足够自由度去分批执行大 grid。

5.3 block 定义了 shared memory 的归属

shared memory 不是整张卡共享，而是一个 block 在某个 SM 上执行时使用的片上共享存储。

所以很多需要局部复用的 kernel，天然会围绕 block 设计：

• GEMM
• 卷积
• attention tile
• block reduction

5.4 block 让程序天然可扩展

同一份代码里，如果任务被切成很多 block：

• 小 GPU 可以一次只跑少量 block
• 大 GPU 可以同时装更多 block

这就是 CUDA 跨设备扩展能力的根基之一。

6. 为什么硬件最终又回到 warp 这一层

如果 block 这么重要，为什么执行粒度又不是 block，而是 warp？

因为 block 太大，不适合作为最小指令发射单位。

假设一个 block 有 256 个线程。如果硬件每次都要“让 256 个线程一起执行一条指令”，会遇到很多问题：

• 发射粒度太粗
• 分支控制代价高
• 数据依赖管理复杂
• 流水线压力大

于是硬件采用更细的粒度：warp。

在 NVIDIA 体系里，一个 warp 通常是 32 个线程。你可以先把它理解成：

一组在执行路径和调度上强相关的线程集合。

所以：

• block 负责任务切分和局部协作
• warp 负责高吞吐执行和硬件调度

这两层都不可少。

7. thread、warp、block、SM 到底怎么对应

这一部分要建立成脑图。

7.1 thread：逻辑工作单位

thread 是最直观的编程单位。你会用它来映射数据空间，比如：

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

每个 thread 都有自己的：

• threadIdx
• 局部变量
• 寄存器上下文

但它更多是编程视角的最小单位，不是最经济的硬件调度粒度。

7.2 warp：执行和调度的重要粒度

一个 block 会按线程 ID 连续切成若干个 warp。

例如 block 大小为 256 时，会形成 8 个 warp：

• warp 0: thread 0-31
• warp 1: thread 32-63
• …
• warp 7: thread 224-255

很多性能问题都要放到 warp 这一层看：

• 分支发散
• warp-level primitive
• 访存合并
• 活跃 warp 数量

7.3 block：资源装载单位

一个 block 被放到某个 SM 上时，会消耗这台 SM 的一部分资源，例如：

• 寄存器
• shared memory
• warp slot
• block slot

只有等这个 block 执行结束，这些资源才会释放。

所以：

• block 太大可能压低同时驻留数量
• 寄存器用太多可能压低驻留数量
• shared memory 用太多也会压低驻留数量

7.4 SM：吞吐型执行资源池

不要把 SM 简化成“一个大核”。

更准确的理解是：

SM 是一个能同时维护很多 warp 上下文、拥有寄存器文件、shared memory、调度器和执行单元的吞吐型资源池。

这也是 GPU 和 CPU 在执行哲学上的关键差异：

• CPU 更强调少量线程的低延迟
• GPU 更强调大量 warp 的整体吞吐

8. 一个 block 是如何在 SM 上“活起来”的

继续用 block_size = 256 举例。

当一个 256-thread 的 block 被放到某个 SM 上时，大致会发生：

1. 它被切成 8 个 warp
2. 这些 warp 共享这个 block 的 shared memory
3. 每个 thread 使用自己的寄存器
4. warp scheduler 在很多就绪 warp 之间轮流发射指令

这里最关键的一点是：

GPU 的高吞吐不主要来自“单个 warp 很快”，而主要来自“一个 SM 上同时挂着很多 warp，可以来回切换”。

9. latency hiding：GPU 为什么没那么怕内存延迟

CPU 通常会用 cache、乱序执行、分支预测等手段，让少量线程尽量低延迟执行。

GPU 的思路更像是：

• 某个 warp 去等内存
• 调度器先不等它
• 切换到另一个 ready warp

这就是 latency hiding。

所以 GPU 喜欢一个 SM 上有足够多的 in-flight warps。这样当某个 warp 因为下面这些原因暂时不能前进时：

• global memory 访问还没返回
• 数据依赖还没满足
• 某类执行单元暂时繁忙

调度器还能继续推进别的 warp。

这就是为什么 GPU 不是靠“单线程做快”吃饭，而是靠“让很多 warp 在途”。

10. occupancy 到底是什么

occupancy 可以粗略理解为：

某个 SM 上当前活跃 warp 数，占该 SM 理论最大活跃 warp 数的比例。

它回答的是一个很实际的问题：

当前这个 kernel，在一个 SM 上到底留下了多少可供调度器切换的 warp？

如果 occupancy 太低，容易出现这种情况：

• 某个 warp 去等内存
• 结果没别的 warp 可切
• SM 空转

所以 occupancy 对隐藏延迟确实重要。

11. 为什么 occupancy 不是越高越好

这是非常高频的误区。

很多人知道 occupancy 之后，会不自觉地把目标变成“把 occupancy 拉满”。这不对。

因为提高 occupancy 往往有代价：

• block 变小，数据复用可能变差
• 每线程寄存器预算更紧，可能发生 spill
• shared memory tile 变小，访存次数变多
• 指令级并行度下降

所以正确理解应该是：

occupancy 是隐藏延迟的重要手段，但最终目标始终是整体吞吐更好。

更直白一点说：

• occupancy 太低，通常有风险
• occupancy 很高，不代表一定最快
• 还是要和寄存器、shared memory、访存模式、计算强度一起看

12. 为什么 block 之间默认独立

很多初学者会问：为什么 CUDA 不允许我随便在 grid 中间做一次全局同步？

因为一旦允许 block 之间自由同步，调度自由度会被严重破坏。

现在这套设计的好处是：

• block A 可以先在 SM0 跑
• block B 可以同时在 SM1 跑
• block C 暂时不跑，等资源空出来再上

也就是说，grid 可以远大于当前硬件能同时容纳的 block 数量，GPU 只需要分批处理。

如果中途强制所有 block 同步，就会变成：

• 你得保证所有 block 都同时有位置
• 调度器自由度下降
• 很多大规模问题会失去可扩展性

所以 CUDA 的默认设计是：

• block 内强协作
• block 间弱耦合
• 全局同步通常交给 kernel 边界

这也是为什么很多算法会拆成多个 kernel。不是程序员爱麻烦，而是 kernel 边界天然就是一个清晰的全局同步点。

13. 这套模型为什么特别适合深度学习和推理

深度学习里的很多操作都有类似结构：

• 输出空间很大
• 相邻输出往往依赖相邻输入块
• 同类运算重复很多次
• 局部复用很重要

这和 CUDA 的层次化抽象刚好很契合：

• grid：覆盖整个问题空间
• block：处理一个 tile、一段 token 或一块输出区域
• warp：推进更细粒度协作
• thread：处理局部元素

例如：

• GEMM 里，一个 block 常常负责输出矩阵的一块 tile
• attention 里，一个 block 或 warp 可能负责部分 token/head 区域
• decode kernel 里，粒度选择会直接影响并行度和缓存行为

所以 CUDA 长期成为 AI 计算主战场，不只是生态问题，也因为这套抽象和张量计算的结构高度匹配。

14. 为什么很多推理 kernel 的设计都绕不开“粒度”

当你开始看：

• FlashAttention
• PagedAttention
• layernorm
• RMSNorm
• TopK
• sampling kernel

会发现大家总在讨论：

• 一个 block 处理多少元素
• 一个 warp 处理几行或几列
• 一个 thread 一次处理几个值
• 归约放在 warp 内还是 block 内

原因很简单：

GPU 性能优化，本质上就是在选择最贴合硬件的执行粒度。

粒度太粗，容易出现：

• block 太大
• 资源占用过高
• 同时驻留数量下降
• occupancy 下降

粒度太细，容易出现：

• 调度和 launch 开销占比升高
• 数据复用变差
• shared memory 的价值下降
• 带宽压力变大

所以好的 kernel 不是“线程越多越好”，而是 thread / warp / block 三层粒度和数据组织方式刚好贴合问题结构。

15. Volta 之后，不要乱用“warp 天然同步”的假设

老资料里常见一种说法：同一个 warp 内线程天然同步，所以很多事情可以不写显式同步。

在现代 GPU 上，这种说法要非常谨慎。

从 Volta 开始，NVIDIA 引入了 Independent Thread Scheduling。这意味着你不能再随意依赖一些没有被正式保证的隐式行为。

更稳妥的原则是：

• 需要同步时，用明确同步原语
• 需要 warp 内通信时，用官方支持的 warp primitive
• 不要依赖某一代 GPU 的偶然执行细节

这不是“代码风格”问题，而是你是否把 CUDA 当成抽象模型来写的问题。

16. 一个简单例子：为什么归约会暴露执行粒度问题

先看一个简化的 block 级归约：

__global__ void block_sum(const float* x, float* out, int n) {
    __shared__ float smem[256];

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

    float val = 0.0f;
    if (idx < n) val = x[idx];

    smem[tid] = val;
    __syncthreads();

    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride) {
            smem[tid] += smem[tid + stride];
        }
        __syncthreads();
    }

    if (tid == 0) {
        out[blockIdx.x] = smem[0];
    }
}

这个例子虽然简单，但已经暴露了几个关键点：

• 为什么要有 block：因为归约需要局部协作
• 为什么 block 内线程必须在同一个 SM：否则 shared memory 没法成立
• 为什么后半段效率会下降：活跃线程越来越少
• 为什么很多高性能实现后半段会切到 warp primitive：因为协作粒度在变化

这个思路不只适用于 reduction，也适用于 softmax、layernorm、sampling 等推理 kernel。

17. 这一篇必须记住的几句话

• thread 是逻辑单位，warp 是关键执行粒度，block 是关键资源和协作粒度。
• SM 是吞吐型执行资源池，不是简单意义上的“一个核”。
• grid 中的 block 会被分批装载到 SM，不是一次性全部同时执行。
• GPU 的高吞吐依赖大量 in-flight warps 来隐藏延迟。
• occupancy 重要，但不是越高越好。
• block 间默认独立，是 CUDA 保持可扩展性的关键设计。
• 很多 AI 推理 kernel 的核心问题，其实都是执行粒度和数据组织是否贴合硬件。

18. 精简版面试表达

如果面试官问 CUDA 的 thread / block / warp / SM 是怎么对应的，可以这样答：

thread 是编程时最直观的逻辑工作单位，但 GPU 不会按单线程高成本调度。实际执行时，一个 block 会被装到某个 SM 上，再切成多个 warp，由 warp scheduler 选择可执行 warp 发射指令。block 负责资源分配、shared memory 和局部同步，warp 负责更细粒度的执行推进，SM 则是承载这些资源和上下文的吞吐型执行单元。很多 kernel 性能问题，本质上都来自这几层粒度与数据组织之间是否匹配。

---

系列导航

• 导读：GPU系统拆解-00-导读：从架构认知到推理系统的学习路线
• 上一篇：GPU系统拆解-04-数据中心架构：为什么大厂更偏好 Hopper、Blackwell 与 MI300
• 下一篇：GPU系统拆解-06-内存层次与访存主线：GPU 性能为什么常常卡在数据
• 延伸阅读：CUDA C编程权威指南：导读
• 延伸阅读：CUDA实践-01-基础概念与分块