GPU 执行模型：SM、warp、线程层次与 occupancy

上一篇 说算子工程发生在框架与硬件之间。要写快算子，先得知道硬件怎么执行你的线程。这一篇建立 GPU 执行模型的硬件直觉：线程的层次结构、它们如何映射到物理单元、warp 为什么是 32、分支发散为什么慢。这些是后面所有优化的地基。

一、线程层次：grid / block / thread

CUDA 的并行模型是三层：你启动一个 grid，grid 由若干 block（线程块）组成，每个 block 由若干 thread 组成。

// grid 有 blk 个 block，每个 block 有 256 个线程
kernel<<<blk, 256>>>(...);

block 和 grid 都可以是一维到三维。线程通过内建变量定位自己：threadIdx（block 内的线程坐标）、blockIdx（grid 内的 block 坐标）、blockDim（每个 block 的线程数）、gridDim（grid 的 block 数）。一维场景下全局索引就是：

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

这套抽象的关键性质：block 之间相互独立。不同 block 不能假设彼此的执行顺序，也不能直接同步（同一 grid 内）。这正是 GPU 可扩展性的来源——同一份 kernel，在 38 个 SM 的卡和 132 个 SM 的卡上都能跑，只是 block 被分到更多物理单元上。

二、物理映射：block 到 SM，thread 到 warp

逻辑层次要落到物理单元才有意义。GPU 的核心计算单元是 SM（Streaming Multiprocessor）。本系列测试卡 RTX 3060 Ti 有 38 个 SM。

映射规则：

• 一个 block 被整体调度到某一个 SM 上，不会跨 SM。block 的全部线程共享这个 SM 的 shared memory 和寄存器文件。
• 一个 SM 可以同时驻留多个 block，只要寄存器和 shared memory 够分。RTX 3060 Ti 每个 SM 最多驻留 1536 个线程、16 个 block。
• SM 内部以 warp 为单位执行。warp 是 32 个连续线程（threadIdx 相邻）的固定分组。一个 256 线程的 block 被切成 8 个 warp。

warp 是硬件真正调度的单位。SM 上有多个 warp 调度器（scheduler），每个周期从就绪的 warp 中挑一个发射指令。这是 GPU 隐藏延迟的核心机制：当一个 warp 在等显存返回时，调度器切到另一个就绪 warp，访存延迟就被计算”盖住”了。这一点在 occupancy 篇 展开。

flowchart TD
  G["Grid"] --> B0["Block 0"] & B1["Block 1"] & Bn["Block ..."]
  B0 -->|"整块调度到一个 SM"| SM["SM (38 个之一)"]
  SM --> W0["Warp 0 (lane 0-31)"] & W1["Warp 1 (lane 32-63)"] & Wn["Warp ..."]
  W0 -->|"warp 调度器每周期选一个就绪 warp"| EX["发射指令到执行单元"]

三、SIMT：warp 内 32 线程锁步执行

warp 的执行模型叫 SIMT（Single Instruction, Multiple Threads）：同一个 warp 内的 32 个线程在同一时刻执行同一条指令，只是各自作用在不同数据上。这和 CPU 的 SIMD 类似，但 SIMT 让每个线程看起来有独立的控制流——代价藏在分支里。

warp 为什么是 32？这是 NVIDIA 架构的固定设计（warpSize == 32，至今所有 NVIDIA GPU 都如此）。很多优化都以 32 为粒度：访存合并按 warp 发起、warp 级原语（__shfl_*、__ballot_*）在 32 个 lane 间通信、归约先在 warp 内做。这些会在 访存优化篇 和 reduction 篇 反复出现。

四、分支发散：SIMT 的代价

既然一个 warp 同一时刻只能执行一条指令，那么当 warp 内的线程走向不同分支时会发生什么？

答案是串行化。如果一个 warp 里一半线程满足 if、一半不满足，硬件会先执行 if 分支（让不满足的线程闲置/被掩码屏蔽），再执行 else 分支（让满足的线程闲置）。两条路径的时间相加，而不是取最大——这就是分支发散（warp divergence）。

关键在于发散只发生在 warp 内部。如果分支条件按 warp 对齐（整个 warp 走同一条路），就没有发散代价。下面用一个对照实验量化：同样的总计算量，一种让分支按 warp 划分（threadIdx.x >> 5，warp 内一致），一种按 lane 划分（threadIdx.x & 1，warp 内交替发散）。

int cond = (mode == 0) ? ((threadIdx.x >> 5) & 1)   // warp-uniform
                       : (threadIdx.x & 1);          // intra-warp divergent
if (cond) { for (int k=0;k<2048;++k) x = x*1.0001f + 0.5f; }
else      { for (int k=0;k<2048;++k) x = x*0.9999f - 0.5f; }

在 RTX 3060 Ti 上，\(n=2^{22}\)、block 256 线程，CUDA event 计时 100 次取中位数：

发散版本慢约 1.74 倍。没有到理论上限 2 倍，是因为访存、循环开销等部分在两种模式下相同，只有 if/else 计算体被串行化。结论很实用：让分支条件尽量按 warp 对齐，例如用 threadIdx.x / 32 而非 threadIdx.x % 2 做粗粒度分流；数据相关的分支无法消除时，考虑把数据按条件重排（访存优化篇 会谈到布局重排）。

需要补充一点架构背景：从 Volta 开始，NVIDIA 引入了独立线程调度（independent thread scheduling），warp 内线程各自维护程序计数器，使得发散后的线程可以交错前进、支持更复杂的同步模式。但这改变的是正确性与灵活性，不改变”发散路径要分别执行”的性能本质。一个直接后果是：warp 内线程不再保证隐式锁步重新汇聚，依赖 warp 内线程一致的归约、shuffle 等操作必须显式调用 __syncwarp() 同步，这一点在 first-kernel 篇 和 reduction 篇 会再次出现。

五、occupancy：SM 上住了多少 warp

occupancy（占用率） 指一个 SM 上实际驻留的 warp 数与硬件上限的比值。RTX 3060 Ti 每个 SM 最多 1536 线程，即 48 个 warp；如果你的 kernel 配置只让 SM 驻留 24 个 warp，occupancy 就是 50%。

occupancy 由三个资源中最紧的那个决定：

• 每线程寄存器数：寄存器文件每 SM 65536 个。若每线程用 64 个寄存器，则最多 \(65536/64 = 1024\) 线程，occupancy 上限 \(1024/1536 \approx 67\%\)。
• 每 block 的 shared memory：每 SM 100 KB。用得多，能并存的 block 就少。
• block 大小与 block 数上限：每 SM 最多 16 block、1536 线程。

occupancy 高意味着 SM 有更多就绪 warp 可供调度，更容易隐藏延迟。但高 occupancy 不等于高性能：寄存器用得少虽然能提高 occupancy，却可能因为变量溢出到显存（register spilling）反而变慢。occupancy 只是延迟隐藏的必要条件之一，完整的权衡留给 第 06 篇。

六、小结与下一步

• 逻辑层次 grid/block/thread 映射到物理 SM/warp/lane；block 整块进一个 SM，warp 是 32 线程的调度单位。
• SIMT 让 warp 内 32 线程锁步执行；warp 内分支发散会把多条路径串行化，本卡实测约 1.7 倍代价，按 warp 对齐分支可避免。
• occupancy 是 SM 驻留 warp 与上限之比，由寄存器、shared memory、block 限制中最紧的决定，是延迟隐藏的基础但不是性能的全部。

有了执行模型，下一篇看数据住在哪——内存层次：global / L2 / shared / register 的带宽与延迟，这是大多数算子真正的瓶颈所在。