当你写下这行代码：
myKernel<<<1000, 256>>>(...);
你其实是在说：“给我 25.6 万个线程，去把这件事干了。”
问题来了：这 25.6 万个线程，具体在哪里跑？谁决定谁先谁后？有多少个是同时真的在跑的？
上一篇我们讲了线程与数据如何对位，这一篇向下钻一层，去看 GPU 这台机器本身。PMPP 第四章给出的答案，会颠覆不少从 CPU 带过来的直觉。
1. 车间俯视图：SM 与 GPC
先看一张"厂区平面图"。
GPU 的基本构成单位是 SM（Streaming Multiprocessor，流多处理器）。每个 SM 内部有一批 streaming processor（流处理器），它们共享控制逻辑和内存资源。
以 Hopper 架构的 H100 为例：
| 层级 | 数量 |
|---|---|
| GPC（GPU Processing Cluster） | 8 个 |
| SM | 132 个 |
| 每个 SM 的流处理器 | 128 个 |
| 流处理器总数 | 132 × 128 = 16,896 个 |
从 Hopper 开始，SM 被进一步组织成 GPC。132 个 SM 分布在 8 个 GPC 里。
一个类比：SM 是车间，流处理器是工位，GPC 是厂区里的楼。 而接下来要说清楚的，是"班组"——warp。
2. 块调度的"任意性"：透明可伸缩性的代价
你启动一个网格，里面有 1000 个块。它们是怎么被分配到 132 个 SM 上的？
答案是：以任意顺序。
CUDA 编程模型不承诺任何分配顺序，也不承诺哪些块会同时运行。这件事带来一个巨大的好处——透明可伸缩性（transparent scalability）：
同一份二进制代码，不需要任何修改，就能在只有 1 个 SM 的低端显卡上跑，也能在 132 个 SM 的 H100 上跑，性能随硬件规模自然放大。
但这是一笔交易。自由的代价是约束：
- 不能假设块的执行顺序；
- 不能假设同时运行的块的数量；
- 线程块之间不能互相同步。
第三条是硬性规定，不是性能建议。如果两个块互相等待，而它们恰好没有被同时调度上来，程序就会直接死锁——而且这种死锁在低端卡上更容易复现，因为 SM 少，块之间更容易排不开。
Hopper 起给了这个问题一个受控的例外：线程块簇（thread block cluster）。同一个簇里的块会被放在同一个 GPC 上，可以互相同步，还能访问彼此的分布式共享内存。但这是显式声明的、有边界的合作，不是"随便跨块同步"。
3. warp：32 个线程的联合命运
块被送上 SM 之后，会立刻被切成 warp：每组 32 个线程。
warp 是 GPU 真正执行的最小单位，它的执行模型是 SIMD（单指令多数据）：一条指令，32 个线程一起执行。
这里有个细节值得记住：块的 x 维度是划分 warp 时变化最快的维度。也就是说，threadIdx.x 相邻的线程，最有可能落在同一个 warp 里，一起走完整个生命周期。
4. 控制发散：半个 warp 摸鱼，整个 warp 买单
SIMD 的麻烦之处在于条件分支。
如果同一个 warp 里，有一部分线程走 if 分支、另一部分走 else，硬件不会把它们分成两组并行跑。它会：
- 先执行
if分支，只让走这条路的线程处于活跃状态，其余线程空转； - 再执行
else分支，同样只让另一批线程活跃； - 两条路径串行走完。
这就是控制发散（control divergence）。它的代价正比于一个 warp 里出现的不同路径数量——两条路径就是两倍时间，三条就是三倍。
一句话总结这一节的直觉：
同一个 warp 里的线程，是命运共同体。只要有一个人走岔路，全组陪他多走一遍。
这也解释了一个后面会反复用到的技巧：在归约（reduction）这类操作里，为什么要把线程下标倒着安排（从高往低数），让活跃线程尽量连成一片——正是为了让"空转的人"整整齐齐地聚在几个 warp 里，而不是散落在每个 warp 中。
5. 延迟容忍：零开销切换的秘密
GPU 的设计哲学里，有一个和 CPU 截然不同的选择。
CPU 拼命降低每一次操作的延迟：大缓存、分支预测、乱序执行，都是为了"让单个线程跑得快"。
GPU 的选择是：不消灭延迟，而是容忍它、覆盖它。
机制是这样的：SM 上常驻的线程数，远远多于它同一时刻真正能执行的线程数。任何时刻，SM 只挑一部分 warp 发射指令，其余的都在等——等一个长延迟的内存访问返回。当某个 warp 卡住时，调度器立刻切到另一个准备好的 warp，而这次切换不需要保存和恢复现场，因为每个 warp 的状态本来就常驻在硬件里。
这就是零开销线程调度（zero-overhead thread scheduling）。
换个比方：这不是让一个工人跑得更快，而是雇足够多的人，让每个人在等电梯的时候，别人正好在干活。老派的 CPU 是让一个博士生又快又聪明地做完一切；GPU 是请一万个普通人，用排班把空闲时间填满。
6. 占用率：不是分数，是"待办工作池"
于是我们引出这一章最重要的概念：占用率（occupancy）。
$$\text{Occupancy} = \frac{\text{resident threads per SM}}{\text{max threads per SM}}$$它的意义不是"跑分"，而是：SM 手里有多少可切换的工作可以做。这个池子越大，就越能把长延迟操作藏起来。
关键在于，这个池子被四项资源同时限制，取其中最小的那个：
| 限制项 | 说明 |
|---|---|
| 线程数上限 | SM 最多能容纳多少线程 |
| 块数上限 | SM 最多能容纳多少个块 |
| 寄存器上限 | 每个线程吃的寄存器 × 线程数 |
| 共享内存上限 | 每个块吃的共享内存 × 块数 |
所以调优时的第一个问题永远不是"怎么提高占用率"，而是：
到底是哪一项在卡我？
修错了资源上限，等于白干。
7. 一个算例：谁才是瓶颈？
书上给了一组典型的设备参数，我们照抄一遍算：
| 资源 | 设备上限 |
|---|---|
| 线程 | 2048 / SM |
| 块 | 32 / SM |
| 寄存器 | 64K / SM |
| 共享内存 | 96 KB / SM |
假设你的 kernel 用 128 线程的块、每线程 32 个寄存器、每块 8 KB 共享内存：
| 约束 | 计算 | 允许的块数 |
|---|---|---|
| 线程 | 2048 ÷ 128 | 16 |
| 块数 | 32 | 32（不构成限制） |
| 寄存器 | 65536 ÷ (128 × 32) | 16 |
| 共享内存 | 96 KB ÷ 8 KB | 12 |
共享内存把 SM 卡在 12 个块 = 1536 线程，也就是 75% 的占用率。线程和寄存器明明都够 16 个块，但决定权在最小的那个数手里。
这个算例顺便预告了第五章的核心矛盾：共享内存既是提速的关键，也是占用率的杀手。
8. 一个反直觉的伏笔：低占用率有时是对的
读到这儿很容易得出一个结论：“占用率当然是越高越好。”
先别急着下结论。
后面的第十五章会把矩阵乘法优化到接近峰值，而那个版本恰恰是寄存器吃得极多、占用率很低的。原因是：当每个线程手里有足够多互相独立的指令（比如展开的乘加、向量化的访存）时，单个线程自己就能把延迟填满，不需要那么多邻居陪跑。
所以占用率的正确读法是：
它描述的是隐藏延迟的手段之一，而不是目的本身。手段可以换：更多 warp，或者更多 ILP（指令级并行）。
记住这句话，第五篇之后你会反复用到它。
9. 顺手带走一段代码
这些限制不必去查手册硬记，运行时就能问出来：
int dev = 0;
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, dev);
printf("%s: %d SMs, %d threads/SM, %d regs/SM, %zu bytes shared/SM\n",
prop.name, prop.multiProcessorCount, prop.maxThreadsPerMultiProcessor,
prop.regsPerMultiprocessor, prop.sharedMemPerMultiprocessor);
把这段输出记在你的笔记本里，它就是你在做占用率算术时的"设备参数表"。
10. 小结
- GPU 是 SM 的阵列，Hopper 起 SM 又被组织成 GPC；这也是"车间—工位—班组"的物理基础。
- 线程块被任意分派给 SM，这份自由换来透明可伸缩性，代价是跨块不能同步（除非用 Hopper 的线程块簇）。
- warp 是 32 个线程的共同体，SIMD 执行；控制发散让整个 warp 陪跑。
- GPU 不消灭延迟，而是用零开销的 warp 切换去覆盖延迟。
- 占用率是"待办工作池"的大小，被线程、块、寄存器、共享内存四项取最小所限制。
- 低占用率不必然是坏事——只要每个线程自己有足够的独立指令。
现在我们已经知道了车间长什么样、班组怎么排班。那么下一个问题就变得非常尖锐：
既然机器这么能跑，为什么很多 kernel 还是慢得可怜？
下一篇，我们去撞那堵墙——内存墙。
参考链接
Programming Massively Parallel Processors: A Hands-on Approach (5th Edition) — Wen-mei W. Hwu, David B. Kirk, Izzat El Hajj. 第 4 章 Compute architecture and scheduling.