理解并优化 CUDA Occupancy

关键词：CUDA Occupancy（占用率）、Streaming Multiprocessor (SM)、CUDA Optimization、Register/Shared Memory Management、Latency Hiding

• Learn HPC with me: Understanding and Optimizing CUDA Occupancy
• https://medium.com/@manisharadwad/unlocking-gpu-potential-understanding-and-optimizing-cuda-occupancy-2f43ee01ad7e
• 6000 字，阅读 20 分钟，播客 13 分钟

• 首个硬件感知智能体 GPU Kernel 优化框架 SwizzlePerf：让 L2 命中率飙升 70%，最高提速 2.06倍！
• ISCA'25 建模GPU LLM推理性能AMALI: 借助张量核心优化与warp分布，推理误差低于GCoM的GPU分析模型
• 94% GPU 利用率！Llama 模型推测解码 10%+ 提速与 2.0× 批量性能跃升！从训练到推理的生产级全栈优化

本文聚焦 CUDA 占用率这一 GPU 性能优化关键概念，阐述其定义、影响因素及优化方法。CUDA 占用率指流多处理器（SM）上活跃 warp 数与最大支持数的比值，核心作用是通过多活跃 warp 切换隐藏内存访问等延迟，最大化 SM 利用率与吞吐量，低占用率会导致设备闲置。

SM 的有限资源（最大线程、最大block数、最大寄存器、最大共享内存，这种“最大”在本文中也被称为“槽”，slot）及动态分区机制是占用率的核心制约因素，线程/block 对同一资源的冲突、block 大小与 SM 最大线程数不可整除、寄存器或共享内存压力等，都可能降低占用率，甚至引发性能骤降的“性能悬崖”。

而占用率优化需从四方面入手：选择 32 倍数的合理块大小；精简局部变量、利用编译器工具管控寄存器使用；按需分配共享内存；借助 CUDA 占用率计算器预判限制因素，NeuralTalk 再补充一点：Grid 设置为 SM 个数的倍数^[1]。

需注意，占用率是性能优化手段而非目标，需通过代码剖析验证优化效果，最终解锁 GPU 计算潜力。

unsetunset本文目录unsetunset

• 本文目录
• 关键问题
• 问题 1：A100 SM 双重限制下，优先调整块大小还是优化寄存器 usage 以最低成本提升占用率？
• 问题 2：占用率提升但执行时间增加时，如何量化判断瓶颈并修正？
• 一、什么是占用率？为什么它很重要？
• 1.1 隐藏访存需要多少线程：Little's Law
• 二、GPU 的核心：流多处理器（SM）的资源
• 三、动态分配：兼具灵活性与复杂性
• 四、资源限制如何制约占用率
• 4.1 线程槽与线程块槽的限制
• 4.2 线程块大小的“不可整除”问题
• 4.3 寄存器压力（Register Pressure）
• 4.4 共享内存限制（Shared Memory Limits）
• 五、“性能悬崖”（Performance Cliff）
• 六、占用率优化：开发人员可以做些什么？
• 6.1 合理选择线程块大小
• 6.2 管理寄存器使用
• 6.3 管理共享内存使用
• 6.4 使用 CUDA 占用率计算器
• 总结

交流加群请在 NeuralTalk 公众号后台回复：加群

unsetunset关键问题unsetunset

问题 1：A100 SM 双重限制下，优先调整块大小还是优化寄存器 usage 以最低成本提升占用率？

在 A100 SM（2048 线程槽、32 块槽、65536 寄存器）场景下，若内核同时面临“块槽不足（如 32 线程/块时仅能跑 32 块）”与“寄存器压力（如单线程需 40 寄存器）”的双重限制，应优先调整块大小还是优化寄存器 usage，才能以【最低开发】成本实现占用率最大提升？

从 A100 SM 硬件参数看，其支持 2048 线程槽、32 块槽及 65536 寄存器：

• 若块大小为 32 线程/块，会因块槽上限 32 仅能激活 1024 线程（50%占用率）
• 此时将块大小调整为 64 线程/块（32 倍数，符合作者推荐的“合理块大小原则”），32 块可恰好填满 2048 线程槽

问题 2：占用率提升但执行时间增加时，如何量化判断瓶颈并修正？

作者提到“寄存器溢出可能提升占用率却增加内存访问延迟”，且“占用率是手段而非目标”，那么当优化后占用率从 75%提升至 100%但内核执行时间反而增加时，如何量化判断是“寄存器溢出的延迟代价”还是“其他资源瓶颈”导致的性能反向，且该如何针对性修正？

需通过“工具检测+指标对比”量化定位，再针对性修正：

1. 判断是否为寄存器溢出代价：用编译器 flag -Xptxas -v 查看寄存器溢出情况，若输出“spilled to local memory”，说明溢出导致额外内存访问延迟（与作者提到的“寄存器溢出虽可能提占用率但增延迟”一致）；
2. 排查其他资源瓶颈：
• 用 CUDA Occupancy Calculator 输入内核参数（线程/块、寄存器/线程等），查看是否存在共享内存超限或线程槽未充分利用；
• 同时通过 profiling 工具对比优化前后的“全局内存访问耗时”，若耗时显著增加，可能是块大小调整导致内存访问模式恶化如合并访问被破坏。

unsetunset一、什么是占用率？为什么它很重要？unsetunset

我们可以把流多处理器（SM）想象成一个能同时容纳许多“工人”（线程）工作的工厂车间。

• 在 CUDA 中，执行调度的基本单位是线程束（warp），通常包含 32 个执行相同指令的线程（线程束是 GPU 调度的最小单元，同一线程束内的线程会同步执行）。
• 一个流多处理器（SM）可以同时处理多个线程束。

占用率的定义是：当前分配给某个流多处理器（SM）的活跃线程束数量，与该 SM 在物理上可支持的最大线程束数量的比值。NeuralTalk 注：具体来说，这里对占用率的定义，是表达实际 GPU运行的 Achieved Occupancy，不是理论占用率。

为什么要追求高占用率？最主要的原因是为了隐藏延迟。

• 高占用率意味着有更多“就绪”的线程束可供调度，能让 SM 的执行单元（如 CUDA Core）始终保持忙碌，从而有效隐藏长耗时操作的延迟，最大化吞吐量（throughput，单位时间内完成的任务数量）。
• 低占用率则意味着 SM 可能会频繁处于空闲状态——因为没有可执行的就绪线程束，只能等待停滞的线程束恢复。

1.1 隐藏访存需要多少线程：Little's Law

Little's Law^[2]确定了通过吞吐量完全隐藏延迟所需的并发量。

该定律是被 Lazowska 等人所著的经典定量系统教科书^[3]中描述为分析学的“基本定律中最重要的一条”。

• 如果一款 GPU 的峰值吞吐量为每周期 1 条指令，内存访问延迟为 400 个周期，那么程序中所有活跃线程束需要执行 400 次并发内存操作。
• 如果吞吐量提升至每周期 10 条指令，那么为了充分利用这一提升，程序需要执行 4000 次并发内存操作。

对于 Little's Law 的一个重要应用，我们来看看 Vasily Volkov  关于“延迟隐藏”的博士论文^[4]第 4.3 节中的观察结果：隐藏纯内存访问延迟所需的线程束数量并不比隐藏纯算术延迟所需的多太多（在他的实验中，前者为30，后者为24）。

直觉上，内存访问的延迟更长，似乎需要更多的并发。但并发不仅由延迟决定，还由吞吐量决定。而且，由于“内存带宽”远低于“算术带宽”，所需的并发量结果大致相同——这对于以“延迟隐藏”为导向、将算术操作和内存操作混合进行的系统来说，是一种有用的平衡形式。

unsetunset二、GPU 的核心：流多处理器（SM）的资源unsetunset

要理解是什么限制了占用率，我们需要先了解每个流多处理器（SM）内部可用的资源。这些资源是有限的，必须在该 SM 上运行的所有线程和线程块（thread block）之间共享。

关键资源包括：

1. 线程槽（Thread Slots）：一个 SM 可同时管理的最大线程数量。
2. 线程块槽（Thread Block Slots）：一个 SM 可同时容纳的最大线程块数量（线程块是线程的集合，多个线程块会分配到不同 SM 上执行）。
3. 寄存器（Registers）：用于存储每个线程局部变量的片上快速内存（访问速度极快，是线程存储临时数据的主要位置）。一个 SM 中的寄存器总量是所有线程共享的。
4. 共享内存（Shared Memory）：片上快速内存区域，仅供同一线程块内的线程共享访问（速度和寄存器相差 1 个数量级，常用于线程块内线程间的数据交换）。

unsetunset三、动态分配：兼具灵活性与复杂性unsetunset

像 NVIDIA 安培（Ampere）架构 A100 这样的现代 GPU，采用动态分配方式管理资源。这意味着 SM 的资源（线程槽、寄存器、共享内存）不会预先划分为固定大小的块分配给每个线程块，而是根据当前分配到该 SM 的线程块的需求，动态进行资源分配。

例如，安培架构 A100 的一个 SM 支持最多 2048 个线程（对应 64 个线程束，因为每个线程束含 32 个线程：2048÷32=64），同时最多可容纳 32 个线程块。

这种动态分配方式具有很强的灵活性，能让 SM 高效运行线程块大小差异极大的核函数。

相比之下，固定分配方案的效率更低：

• 如果线程块所需资源少于固定分配的量，就会造成资源浪费；
• 如果所需资源超过固定分配的量，线程块则无法运行。

然而，这种动态性也导致了资源限制之间的复杂交互，可能使 SM 无法达到理论上的最大占用率（100%）。

unsetunset四、资源限制如何制约占用率unsetunset

由于资源有限且采用动态分配，多种因素会限制 SM 上的活跃线程束数量：

4.1 线程槽与线程块槽的限制

你可能会在达到线程槽上限之前，先达到线程块槽的上限。

基于线程数计算的占用率为 1024÷2048=50%。要在该 SM 上实现 100%的线程占用率，你选择的线程块大小至少需要为 64 个线程（2048 线程 ÷ 32 个线程块 = 64 线程/线程块）。

4.2 线程块大小的“不可整除”问题

如果 SM 的最大线程数不能被你选择的线程块大小整除，那么部分线程槽将不可避免地处于闲置状态。

4.3 寄存器压力（Register Pressure）

核函数中的每个线程都会使用寄存器存储局部变量。SM 拥有容量较大但有限的寄存器池（例如，A100 的一个 SM 拥有 65536 个寄存器），这个寄存器池需要在所有活跃线程之间分配。

• 如果你的核函数每个线程需要 40 个寄存器，那么寄存器池可支持的最大活跃线程数为 65536÷40≈1638 个。即使根据线程槽/线程块槽的限制，你的线程块大小允许 2048 个线程，但寄存器限制会将活跃线程数上限压低到 1638 个。此时占用率被限制为 1638÷2048≈80%。
• 如果核函数每个线程需要 64 个寄存器，那么该 SM 仅能支持 65536÷64=1024 个活跃线程。此时最大可实现占用率为 1024÷2048=50%，无论线程块大小如何设置，都无法突破这个限制。

4.4 共享内存限制（Shared Memory Limits）

与寄存器类似，SM 拥有固定大小的共享内存。如果你的核函数每个线程块需要大量共享内存，会限制该 SM 上可同时容纳的线程块数量，进而可能降低占用率。

unsetunset五、“性能悬崖”（Performance Cliff）unsetunset

上述资源限制之间的相互作用，可能导致即使是微小的代码改动，也会引发占用率和性能的急剧下降——这种现象有时被称为“性能悬崖”。

示例（基于 A100 GPU）：假设你的核函数每个线程使用 31 个寄存器，且你设置的线程块大小为 512 个线程。

• 基于线程数限制的 SM 最大线程块数：2048 个线程 ÷ 512 线程/线程块 = 4 个线程块。
• 所需寄存器总量：2048 个线程 × 31 个寄存器/线程 = 63488 个寄存器，这小于 A100 的寄存器池容量（65536 个）。这种情况下，你可以实现 100%的占用率（暂时不考虑共享内存的影响）。

现在，假设你添加了几个局部变量，导致每个线程的寄存器使用量增加到 33 个。

• 4 个线程块（2048 个线程）所需的寄存器总量变为：2048 个线程 × 33 个寄存器/线程 = 67584 个寄存器，这超过了 65536 个的寄存器池上限！
• CUDA 运行时系统无法再调度 4 个线程块，可能会将活跃线程块数量减少到 3 个。
• 此时活跃线程数变为：3 个线程块 × 512 线程/线程块 = 1536 个。
• 实际使用的寄存器数量为：1536 个线程 × 33 个寄存器/线程 = 50688 个，远低于寄存器池上限。

unsetunset六、占用率优化：开发人员可以做些什么？unsetunset

理解上述限制后，开发人员就能做出更合理的决策来优化占用率：

6.1 合理选择线程块大小

当一个内核执行的块数量较少时，可能会导致尾部效应^[5]。所以，在可能的情况下，尽可能让网格启动尽可能夺得块，让他们占满 SM。

• 如果有成百上千个波，末尾部分波的影响会大大降低，相对占比百分之一或者千分之一。
• 选择的线程块大小（每个线程块的线程数）通常应为线程束大小（32）的整数倍。目标是在充分利用线程槽的同时，避免过早触及线程块槽的限制。
• 还要避免选择会导致 SM 最大线程数无法整除的线程块大小。常见的线程块大小选择包括 128、256、512，但最佳选择需根据具体硬件和核函数特性而定。
• NeuralTalk 认为：让 block 的个数（即 grid）是 SM 个数的整数倍，因为在 SM 资源用满的情况下，打满的单个波的大小就是 SM 的个数，仅供参考。因为，实际情况可能会由于资源分配的 SM 运行时的动态性、GPU 不只是单个 Stream 独占、使用 NCCL 占用部分 SM 等等情况，让我们预设的 grid 达不到期望的一个 wave 打满 SM 的情况。

6.2 管理寄存器使用

寄存器数量是主要的占用率限制因素^[6]

• 注意声明的局部变量数量，尽量使用最小的合适数据类型（例如，能用int就不用long，能用float就不用double）。
• 通过编译器标志（如-Xptxas -v）查看每个线程的实际寄存器使用量。该标志会在编译时输出寄存器和共享内存的使用信息。
• 考虑使用启动边界（__launch_bounds__(MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_SM)）为编译器提供预期的并发信息，帮助编译器优化寄存器分配（MAX_THREADS_PER_BLOCK表示每个线程块的最大线程数，MIN_BLOCKS_PER_SM表示每个 SM 上至少要容纳的线程块数）。
• 如果单个线程使用的寄存器过多，可能导致“寄存器溢出”（register spilling）——将部分变量从寄存器转移到速度较慢的全局内存（访问延迟远高于寄存器）。这种操作可能会提高占用率，但会因额外的内存访问延迟导致性能下降，因此需要在占用率和性能之间权衡。

6.3 管理共享内存使用

每个线程块仅分配必要的共享内存，因为共享内存用量也会限制 SM 上可同时运行的线程块数量。

6.4 使用 CUDA 占用率计算器

NVIDIA 提供了一个基于电子表格的工具——CUDA 占用率计算器（CUDA Occupancy Calculator，NeuralTalk注：电子表格这个就过时了，最新可以使用 Nsight 来计算）。

你只需输入目标 GPU 架构、核函数的资源使用情况（每个线程块的线程数、每个线程的寄存器数、每个线程块的共享内存量），该工具就能计算出理论占用率，并识别出限制占用率的关键因素。

unsetunset总结unsetunset

占用率是衡量 CUDA 性能的关键指标，直接影响 GPU 隐藏内存延迟的能力。

占用率由以下因素共同决定：核函数的配置（线程块大小）、核函数的资源需求（每个线程的寄存器数、每个线程块的共享内存量），以及 GPU 流多处理器（SM）的特定硬件限制（最大线程数、最大线程块数、寄存器总量、共享内存总量）。

因此，【务必】通过性能分析工具（如 NVIDIA Nsight）验证：为提高占用率而进行的代码修改，是否真的能缩短执行时间。使用 CUDA 占用率计算器这类工具，可以为优化工作提供指导，帮助你充分释放 GPU 的计算潜力。