Occupancy 与延迟隐藏：寄存器、shared memory 的取舍

第 02 篇 提过 occupancy 是 SM 上驻留 warp 与硬件上限的比值，也提过”高 occupancy 不等于高性能”。这一篇把它讲透：occupancy 到底为什么重要、由什么决定、多高才够，以及为什么一味追求高 occupancy 会适得其反。

一、occupancy 为什么重要：延迟隐藏

GPU 访问 global memory 要几百个周期。如果 SM 上只有一个 warp，它发出访存请求后只能干等，几百周期都在空转。GPU 的对策是用并发掩盖延迟：SM 上同时驻留很多 warp，一个 warp 卡在访存时，调度器立刻切到另一个就绪 warp 发指令。只要就绪 warp 足够多，访存延迟就被持续的计算”盖住”，执行单元不空转。

occupancy 衡量的就是”手头有多少 warp 可切换”。RTX 3060 Ti 每个 SM 最多驻留 48 个 warp（1536 线程），occupancy 100% 即 48 个。occupancy 越高，可调度的 warp 越多，越容易隐藏延迟。

二、occupancy 由什么决定：三个资源上限

一个 SM 能驻留多少 warp，由三类资源中最紧的那个决定：

1. 寄存器：每 SM 65536 个。每线程用 \(R\) 个寄存器，则最多 \(\lfloor 65536 / (R \times 32) \rfloor\) 个 warp。例如每线程 32 个寄存器 → 最多 64 warp，但被 48 的硬上限截到 48；每线程 64 个寄存器 → 最多 32 warp，occupancy 上限 \(32/48 \approx 67\%\)。
2. shared memory：每 SM 100 KB。每 block 用得越多，能并存的 block 越少。
3. block 数 / 线程数硬上限：每 SM 最多 16 个 block、1536 个线程、48 个 warp。

nvcc --ptxas-options=-v 会打印每个 kernel 的寄存器和 shared memory 用量，配合 NVIDIA 的 occupancy 计算器或 cudaOccupancyMaxActiveBlocksPerMultiprocessor 可以算出理论 occupancy。

三、多高才够：实测延迟隐藏的饱和点

关键问题是 occupancy 要追到多高。用一个访存密集的流式 copy kernel，通过控制 grid 大小调节每 SM 的 warp 数——把 grid 的 block 总数设为”SM 数（38）× 目标每 SM block 数”，让每个 SM 恰好分到 1/2/4 个 block——实测有效带宽（RTX 3060 Ti，256 线程/block）：

结论很有意思：occupancy 从 17% 提到 33% 时带宽明显上升（363 → 399），但从 33% 再往上就基本持平。对这个访存密集 kernel，16 个 warp 已经足够隐藏 DRAM 延迟，再多的 warp 没有额外收益。

这推翻了一个常见误区——“occupancy 越高越好”。实际上 occupancy 只要过了延迟隐藏的饱和点就够了，那个点常常远低于 100%。不同 kernel 的饱和点不同：访存延迟长、每 warp 计算少的 kernel 需要更高 occupancy；计算密集、每 warp 指令多的 kernel 用很低的 occupancy 也能跑满。

四、高 occupancy 的代价：寄存器压力

既然高 occupancy 不总是有用，为什么不能无脑追求？因为提高 occupancy 通常要压低每线程寄存器用量，而这有副作用。

寄存器是有限的。要让更多 warp 驻留，编译器必须把每线程的寄存器控制在更低水平。当 kernel 逻辑复杂、需要的寄存器超过限额时，编译器会把一部分变量溢出（spill）到 local memory——而 local memory 物理上在 global 显存里，慢得多。于是出现一种典型的反噬：

• 强行限制寄存器（如用 __launch_bounds__ 或 -maxrregcount）→ occupancy 上去了；
• 但寄存器不够用导致 spilling → 每次访问溢出变量都打显存 → kernel 反而变慢。

正确做法是把 occupancy 当成一个需要平衡的旋钮，而不是越大越好的目标：

// 提示编译器：每 block 最多 256 线程，至少保证 4 个 block/SM 驻留
// 编译器据此约束寄存器分配；设得过激会引发 spilling
__global__ void __launch_bounds__(256, 4) my_kernel(...) { ... }

__launch_bounds__(maxThreadsPerBlock, minBlocksPerSM) 告诉编译器你的启动意图，让它在寄存器分配上做相应权衡。用它之前先用 profiler 确认 kernel 是不是真的受 occupancy 限制——很多时候瓶颈在别处。

五、计算密集 kernel：低 occupancy 也能很快

GEMM 这类计算密集 kernel 是反例。它们每个线程持有大量寄存器（用于寄存器分块，GEMM 篇 详述）来积累结果，每线程寄存器用量很高、occupancy 很低，但因为每个 warp 有大量独立的算术指令可以填满流水线，照样能逼近算力峰值。

这里体现了 ILP（指令级并行）与 TLP（线程级并行）的互补：延迟既可以靠多 warp（高 occupancy/TLP）隐藏，也可以靠单 warp 内大量独立指令（高 ILP）隐藏。寄存器分块用更多寄存器换更高 ILP，是 GEMM 不依赖高 occupancy 也快的原因。

六、决策清单

调 occupancy 前后问自己：

1. 这个 kernel 受 occupancy 限制吗？ 用 profiler 看是访存延迟暴露（occupancy 太低）还是别的瓶颈。compute-bound 或已饱和带宽的 kernel 调 occupancy 没用。
2. 当前 occupancy 离饱和点多远？ 访存密集 kernel 常在 30–50% 就饱和，盲目冲 100% 无益。
3. 提 occupancy 会引发 spilling 吗？ 看 -v 输出的 spill stores/loads；一旦 spill，得不偿失。
4. 能否靠 ILP 替代？ 计算密集 kernel 用寄存器分块提 ILP，比堆 occupancy 更有效。

七、小结与下一步

• occupancy 是延迟隐藏的基础：足够多的就绪 warp 才能掩盖访存延迟。
• 它由寄存器、shared memory、block/线程硬上限中最紧的决定。
• 实测访存密集 kernel 在约 33% occupancy 就饱和，更高无益；高 occupancy 不是目标。
• 强行提 occupancy 可能引发寄存器溢出反而变慢；计算密集 kernel 靠 ILP（寄存器分块）也能不依赖高 occupancy。

到这里，影响性能的两大因素（访存模式、延迟隐藏）都齐了。下一篇把它们统一进一个判断框架——Roofline 模型，回答”这个算子到底受算力还是受带宽限制”。