Occupancy 受什么限制？
是不是越高越好？

英文问法What is occupancy and what limits it? · Is higher occupancy always better? · How do you check how many registers a kernel uses, and how do you tune occupancy?

考过别家面经原题：摩尔线程 AI 算子库及高性能计算工程师一面 2026-10-15「你知道在（CUDA 和 HIP）这两个平台上怎么去看某一个内核的寄存器用量吗？」「有没有学习过、或者看过一些关于（访存吞吐、峰值浮点计算利用率、occupancy）这些比较重要的量的分析，然后运用到算法里面来？」（牛客）· NVIDIA AI Infra 校招「Occupancy 受哪些因素影响？如何进行调控？」（AIInfraGuide 的二手整理，收录 2026-04-17，GitHub）· 两份 AI infra 面经样本里 occupancy 这个关键词只占 1% 和 0%，偏底层 · Unity 一侧没有直接依据，也没有 Unity 自己的面经原题

先弄清楚：每个概念是什么

Occupancy
SM 上活跃的 warp 数，除以这个 SM 最多能放的 warp 数。
Latency hiding
一个 warp 在等 memory 时，SM 切到别的 warp 继续发指令。warp 越多，越不容易全都在等。
Limiter
register、shared memory、thread 数三个上限里最紧的那个，occupancy 由它决定。
Register spill
register 不够用时，compiler 把值放到 local memory，它在 HBM 上，和 global memory 一样慢。

看图：三个上限，取最紧的

一个 SM 最多 64 个 warp。每个 block 8 个 warp，能放几个 block？ thread 数 2,048 ÷ 256 8 个 register 65,536 ÷ 12,288 5 个 shared memory 228 KB ÷ 48 KB 4 个，最紧 occupancy = 4 × 8 ÷ 64 = 50%，limiter 是 shared memory
图 1：示意的 kernel，每个 block 256 个 thread、48 KB shared memory，每个 thread 48 个 register，上限按 H100 算。虚线格是放不进去的 block，三个上限取最小的 4 个。

30 秒回答

Occupancy is the number of active warps on an SM divided by the maximum it can hold. When one warp waits on memory, the SM switches to another, hiding the latency. It's capped by whichever runs out first: registers per thread, shared memory per block, or the block size. Higher isn't always better. Once there are enough warps to keep the SM busy, more don't help, and squeezing registers can cause spills. A GEMM that gives each thread more registers to reuse data often runs faster at lower occupancy. So low occupancy is a warning sign, but I judge by measured time.

三个上限，手算一遍

H100 的一个 SM 最多放 2,048 个 thread（64 个 warp），有 65,536 个 register 和最多 228 KB shared memory（见 thread 上限）。图 1 的 kernel 被 shared memory 限制在 4 个 block，occupancy 是 50%。

把每个 block 的 shared memory 减到 32 KB，这一项放宽到 7 个 block，limiter 换成 register，能放 5 个 block，occupancy 是 62.5%。要到 100%，每个 thread 最多用 65,536 ÷ 2,048 = 32 个 register。

是不是越高越好

太低一定不好，但不是越高越好。

  1. warp 已经多到 SM 每个 cycle 都有指令可发时，再加也不会更快。NVIDIA 的文档说从 66% 提到 100% 通常得不到成比例的加速
  2. 提高它要压 register 或 shared memory。register 压过头会 spill，shared memory 小了 tile 就小，数据复用变少
  3. GEMM 反过来做，每个 thread 多占 register，occupancy 低了，实际更快（见 matmul 的 register tiling）
要做的事办法
看 occupancy 和 limiterncu 的 Occupancy 一节，报 theoretical 和 achieved occupancy，以及 limiter 是哪个
看 register 用量编译时加 -Xptxas -v；ncu 的 Launch Statistics 也报每个 thread 的 register 数
调__launch_bounds__ 限制 register；block size 取 32 的倍数，从 128 到 256 试起；每个 block 少用 shared memory

必答点

追问

怎么知道 occupancy 是不是这个 kernel 的瓶颈？

ncu 里 memory 和 compute throughput 都低时才怀疑它（见 Nsight profiling）。NVIDIA 文档的试法是把 dynamic shared memory 调大来压低 occupancy，看时间变不变。

__launch_bounds__ 是什么意思？

kernel 上的声明，告诉 compiler 每个 block 最多多少 thread、每个 SM 至少放几个 block，compiler 据此限制 register 用量。快手校招问过，出处是 AIInfraGuide 的二手整理。

block size 越大，occupancy 越高吗？

不是，block 是整个放进 SM 的。只看 register 这一项，图里 256 个 thread 的 block 能放 5 个，是 40 个 warp。512 个 thread 的 block 只放得下 2 个，是 32 个 warp。

kernel agent 会用到 occupancy 吗？

会。Meta 的 KernelAgent 在 Diagnose 一步用 ncu 指标把 kernel 判成 memory、compute 或 occupancy bound，再选对应的优化（见 kernel agent 设计）。

来源CUDA C++ Best Practices Guide §11.1、§11.1.1、§11.3、§11.4 · Nsight Compute Profiling Guide（Occupancy、Launch Statistics 两节）· KernelAgent · ml-八股 08 第 20 题 · 面经和频率出处见 面试题与 JD 3.1、3.6 节和 NVIDIA 题库笔记 1.8 节、Q2