想象你是一家高端餐厅的总厨。餐厅后厨里有 8 个灶台组，每个灶台组一排能同时放 32 口锅。客人点了一大堆菜，你要把订单按“桌”分组交给各灶台。问题来了——菜谱有的需要大案板，有的需要很多调料，但灶台旁的案板和调料架是有限的。如果每道菜都豪横地占一大块案板，一桌只能坐 2 道菜，灶台空着一大半；如果大家挤一挤，所有灶台全满，但厨师可能因为案板太小而手忙脚乱，反而更慢。GPU 里的“占用率”就是这个灶台的实际开工率。这篇实战课帮你弄清：占用率到底怎么算、什么在限制它、如何像老总厨一样精打细算地排布订单和案板，让 GPU 这台“算力餐厅”每一口锅都烧得恰到好处，不浪费一寸硬件。

岗位关联

对 AI 从业者：GPU 是深度学习最贵的资源，占用率优化能在不加卡的情况下榨出 30% 以上性能，是算子优化、大模型推理部署、AI 基础设施团队的硬核技能。

目录
  1. 可视化演示
  2. 占用率定义
  3. 占用率计算
  4. 资源限制
  5. Block调优
  6. 寄存器优化
  7. 共享内存
  8. Nsight分析
  9. 实战策略
  10. AI 应用场景
  11. Python 演示
  12. 常见误区
  13. 面试题

可视化演示

展示系统处理流水线、实时性能指标曲线及多方案对比分析，包含数据流转、吞吐延迟监控与方案评估。

本节配有交互动效（进入视口后自动加载，约 3KB）。

系统可视化 处理流程：输入→处理→模型→输出 性能监控：实时指标 P=L×T

占用率定义

走进后厨，你一眼就能看出餐厅生意好不好——如果 8 组灶台只有 2 组在冒火，说明“开台率”很低。GPU 里的占用率（Occupancy）衡量的就是每个 SM（流式多处理器，相当于一个灶台组）上有多少线程束（warp，相当于一位掌勺大厨，一次同时处理 32 道菜）正处于活跃执行状态。灶台空着，算力就空转，这是 GPU 最隐蔽的浪费。

为什么 GPU 需要这么多“大厨”同时开工？因为访存太慢。GPU 从显存取数据往往要等几百个周期，如果 SM 上只有零星几个 warp，数据没到的时候整个灶台组只能干等。而如果驻留了大量 warp，一个 warp 在等数据，调度器会立刻把另一个准备好的 warp 换上来炒菜——这就是“隐藏延迟”。warp 数量越多，延迟被盖住的可能性越大。

一个 SM 能容纳的 warp 数量有硬件上限（比如现代架构通常是 64 个，即 2048 个线程 ÷ 32）。占用率这个指标，就是把当前“实际能驻留的 warp 数”和这个上限做对比：

占用率不是“灶台火力的大小”，而是“灶台有没有被占满”。哪怕每道菜做得再快，只要桌上空着一半灶台，翻台率就上不去。所以在看 GPU 性能报告时，第一眼就要问：我的厨房里，有多少口锅正在同时冒火？

Occupancy=WactiveWmax×100%\text{Occupancy} = \frac{W_{\text{active}}}{W_{\text{max}}} \times 100\%
这是“灶台开工率”：活跃线程束数除以最大容量，得到 0%~100% 的占用率。100% 意味着灶台全满。
  1. W_active（活跃线程束数）当前 SM 上“已经挂号、随时能开工”的 warp 数，它取决于 kernel 的资源占用和硬件调度，是一个动态结果。
    WactiveW_{\text{active}}
  2. W_max（最大线程束数）硬件规定的每个 SM 的 warp 容量上限，例如 64。这是分母，代表“灶台全满”的理想状态。
    WmaxW_{\text{max}}
  3. 除法并转成百分比活跃数除以容量再乘 100，得到 0%~100% 的直观指标，方便工程师横向比较不同配置。
    WactiveWmax×100%\frac{W_{\text{active}}}{W_{\text{max}}} \times 100\%

占用率计算

老总厨排班时要算三笔账：一共能接几桌？每桌坐下来占几个灶位？共用调料架够不够分？三笔账都算完，取最小的那个数字，才是真正能安排的单量。GPU 算占用率也一样——要分别从线程数、寄存器、共享内存三个维度推导上限，再取最小值。

为什么是“取最小值”而不是“取平均值”？因为这三样资源是并联约束：任何一个容量爆了，block 就放不进去。好比厨房案板再大，如果调料架只能放 3 桌的料，第 4 桌就只能排队等着。三者必须同时满足。

以普遍参数为例：每 SM 最大线程数 2048，寄存器文件 65536 个 32 位寄存器，共享内存 96KB。假设你的 block 大小是 256 线程（即 8 个 warp），每线程用 32 个寄存器，每块不申请共享内存。先算线程维度：2048 ÷ 256 = 8 个 block，即 8 × 8 = 64 个 warp；再算寄存器维度：65536 ÷ (32 × 32) = 64 个 warp；共享内存维度不限。三者取最小，活跃 warp 数为 64，占用率 100%。

这三笔账就像三张草稿纸。总厨先把“一桌坐几人”定下来，再算每张草稿纸能接几桌，最后取最紧的那张。后面你会发现，寄存器往往是那只最紧的手——这也是为什么它是调优的第一靶点。

Wactive=min⁡(Wthread,Wreg,Wsmem)W_{\text{active}} = \min\left(W_{\text{thread}}, W_{\text{reg}}, W_{\text{smem}}\right)
Wthread=⌊TSM-maxTblock⌋×Tblock32W_{\text{thread}} = \left\lfloor \frac{T_{\text{SM-max}}}{T_{\text{block}}} \right\rfloor \times \frac{T_{\text{block}}}{32}
第一行是“三笔账取最小”的总公式；第二行是其中第一笔账：每 SM 最大线程数除以块大小，向下取整得到块数，再换算成 warp 数。
  1. T_SM-max（每 SM 最大线程数）硬件规格，比如 2048，是第一个限制因素的分子，代表灶台组最多能站多少个“厨师位”。
    TSM-maxT_{\text{SM-max}}
  2. T_block（每块线程数）程序员自己定的块大小，比如 128 或 256，它决定每桌团餐占多少个厨师位。
    TblockT_{\text{block}}
  3. 整除取整（向下）最大线程数除以块大小，向下取整，得到线程维度最多能放几个 block。整除是因为 block 不能切半个。
    ⌊TSM-maxTblock⌋\left\lfloor \frac{T_{\text{SM-max}}}{T_{\text{block}}} \right\rfloor
  4. 换算成 warpblock 数乘以“每块包含的 warp 数”（线程数除以 32），得到线程维度支撑的活跃 warp 数，参与最终取最小。
    ×Tblock32\times \frac{T_{\text{block}}}{32}

资源限制

为什么不能无限接单？因为厨房有三样东西是焊死的：灶台位（线程数上限）、案板（寄存器文件）、共用调料架（共享内存）。任何一样不够用，单子就接不了那么多。CUDA 里的三样对应资源，每一样都能独立卡死你的占用率，形成“木桶效应”。

寄存器是每个线程私有的“案板”，存取速度极快，但总量有限。编译器为了减少访存，倾向于给每个线程多分几块案板。共享内存则是 block 内所有线程公用的“调料架”，速度快但同样容量有限。这两者一旦被某个 kernel“大手大脚”占用，SM 能同时容纳的 block 数就会直线下降。

寄存器维度能容纳多少 warp，计算公式如下。假设每线程用 64 个寄存器，一个 warp 要 32 份案板：

共享内存维度则这样算——每块申请 8KB，96KB 容量只能放 12 个 block：

资源限制的本质是“硬件物理容量”对并行度的天花板。案板给每道菜分得多，能站的厨师就少；调料架被每桌占得多，能接的桌数同步缩水。三样资源里只要有一块短板，整个厨房的接待能力就被它锁死——优化占用率的第一步，永远是找出最短的那块板。

Wreg=⌊RSM⌈Rper-thread×32⌉⌋W_{\text{reg}} = \left\lfloor \frac{R_{\text{SM}}}{\left\lceil R_{\text{per-thread}} \times 32 \right\rceil} \right\rfloor
Wsmem=⌊SSMSper-block⌋×Tblock32W_{\text{smem}} = \left\lfloor \frac{S_{\text{SM}}}{S_{\text{per-block}}} \right\rfloor \times \frac{T_{\text{block}}}{32}
第一个公式是“案板账”：寄存器总量除以每个 warp 的案板需求。第二个是“调料架账”：共享内存总量除以每块申请量，再乘每块 warp 数，换算成 warp 容量。
  1. R_SM（寄存器文件容量）每 SM 的寄存器总数，例如 65536 个 32 位寄存器，这是所有厨师共享的私人案板总面积。
    RSMR_{\text{SM}}
  2. R_per-thread（每线程寄存器数）编译后每个线程实际占用的寄存器数，由编译器分配，也可以通过编译选项或 __launch_bounds__ 干预。
    Rper-threadR_{\text{per-thread}}
  3. 乘 32 并向上取整一个 warp 有 32 个线程，案板需求按整 warp 计算，向上取整保证不会出现“半个 warp 的案板”。
    ⌈Rper-thread×32⌉\left\lceil R_{\text{per-thread}} \times 32 \right\rceil
  4. 共享内存块数转 warp 数共享内存容量除以每块申请量得到块数，再乘以每块线程数除以 32，换算成 warp 维度，参与取最小。
    ⌊SSMSper-block⌋×Tblock32\left\lfloor \frac{S_{\text{SM}}}{S_{\text{per-block}}} \right\rfloor \times \frac{T_{\text{block}}}{32}

Block调优

订单一桌 4 人和一桌 12 人，对厨房的安排完全不同——桌数少了大厨闲着，桌数多了又坐不下。block（线程块）就是 GPU 里的“一桌团餐订单”：同一 block 内的线程共享共享内存、可以同步协作，它是调度和资源共享的基本单位。

block 大小直接决定前面三笔账的每一项。如果 block 太小，比如只有 32 个线程（1 个 warp），虽然块数很多，但每 SM 有 block 数量上限（例如 32 个），这反而成了新瓶颈——32 块 × 1 warp = 32 warp，占用率只有 50%。如果 block 太大，比如 1024 线程，每 SM 最多 2 块（2048 ÷ 1024），虽然也能到 64 warp，但块内同步开销巨大，且块数太少会让调度器没有回旋余地。

因此 block 大小有一条铁律：必须是 32 的整数倍。因为 warp 是调度单位，非整数倍会让最后一个 warp 空一半座位，白白浪费。经验上最常用的是 128 和 256——既保证每 SM 有足够多的 block 供调度，又不至于让块内协作负担过重。块数上限带来的影响公式如下：

团餐不是人越多越好——一桌 12 人把灶台占满，另一桌只能排队；一桌 2 人，又浪费调度弹性。业界默认“一桌 4 到 8 个厨师”（128/256 线程）是黄金区间：灶台满转、桌数充足、块内同步成本可控。选定 block 大小后，再回头看三笔账，哪笔最紧就动哪笔。

Wblock-limit=Bmax×Tblock32W_{\text{block-limit}} = B_{\text{max}} \times \frac{T_{\text{block}}}{32}
这是“桌数账”：每 SM 最大块数乘以每块 warp 数，得到块数上限维度能支撑的 warp 总数。如果它比线程账小，说明块太小导致被桌数卡住。
  1. B_max（每 SM 最大块数）硬件规定的 SM 能同时容纳的 block 个数上限，比如 32。这是硬件给“桌数”设的硬顶。
    BmaxB_{\text{max}}
  2. 乘每块 warp 数块数上限乘以每块线程数除以 32，算出在“只受桌数限制”时最多有多少个 warp 在工作。
    ×Tblock32\times \frac{T_{\text{block}}}{32}

寄存器优化

案板是每个厨师私人的操作台。菜谱越来越复杂，需要的案板就越多——如果案板占得太多，整个厨房只能站 8 个厨师，剩下 6 个灶台空着。把案板压缩到刚好够用，就能多站几位厨师，灶台利用率立刻上去了。在 CUDA 里，这个“案板”就是寄存器。

每线程寄存器数是影响占用率的最大变量。编译器默认会尽量多给线程分配寄存器，把变量留在最快的存储里；但这会挤占“厨师名额”。CUDA 提供了 __launch_bounds__ 来干预编译器：__launch_bounds__(256, 8) 表示“这个 kernel 每块最多 256 线程，并且我希望每个 SM 至少驻留 8 个块”。

编译器收到这个“目标桌数”后，会倒推每线程最多能用多少寄存器。比如寄存器总量 65536，希望每 SM 至少 8 块 × 256 线程 = 2048 线程，那么每线程最多 65536 ÷ 2048 = 32 个寄存器。但如果 kernel 真的需要 40 个寄存器，编译器只能把多出的变量“溢出”（spill）到 local memory——那是全局内存的一部分，速度可能慢十倍。这就是寄存器优化最核心的权衡：

压缩案板是门手艺活。老总厨熟记每道菜最少需要几块案板，宁可让厨师稍微挤一点也要保证 8 个灶台全开；但他也绝不会把案板砍到低于最低需求——否则厨师只能频繁跑去仓库取料（local memory），单道菜耗时暴增。寄存器优化的艺术，就是在这条“案板 vs 跑腿”的曲线上找到甜点。

Nblock-by-reg=⌊RSMRper-thread×Tblock⌋N_{\text{block-by-reg}} = \left\lfloor \frac{R_{\text{SM}}}{R_{\text{per-thread}} \times T_{\text{block}}} \right\rfloor
这是“案板账”的另一个写法：每 SM 寄存器总量除以“每块线程数 × 每线程寄存器数”，得到寄存器维度最多能放几个 block。分子分母一目了然。
  1. 每 block 寄存器总需求每线程寄存器数乘以每块线程数，得到这一整桌厨师总共要占多少块案板。
    Rper-thread×TblockR_{\text{per-thread}} \times T_{\text{block}}
  2. 除法取整SM 寄存器总量除以每块总需求，向下取整，得到案板维度能放的 block 数。
    ⌊RSMRper-thread×Tblock⌋\left\lfloor \frac{R_{\text{SM}}}{R_{\text{per-thread}} \times T_{\text{block}}} \right\rfloor

共享内存

共享内存是厨房中央的“公用调料架”：同一个 block 的厨师们用它传递半成品、共享热数据，避免每个人都跑去大仓库（全局内存）拿同一味调料。不过调料架总共就那么宽——这桌多占几格，厨房里能站下的桌数就少几个。

共享内存的声明有两种：静态 __shared__（编译期写死大小）和动态 extern __shared__（启动 kernel 时通过第三个配置参数指定字节数）。两种加起来就是每 block 的总申请量 S_per-block。如果每 block 申请 16KB，而 SM 共享内存是 96KB，那共享内存维度最多放下 6 个块——即使线程和寄存器都允许 8 个块，实际也只能驻留 6 个。

不过，“省共享内存”并不总是对的。共享内存的价值在于消除全局内存的重复访问：如果一个 block 反复用到同一批数据，放调料架上取用比每次去仓库快几十倍。省下共享内存导致访存暴涨，新的瓶颈会变成 DRAM 带宽。合理思路是：先算清这块数据会被复用多少次，再决定申请量，并观察它对占用率的挤压：

共用调料架的哲学是“按需取用，但别贪婪”。多放一点调料能省下跑远路的功夫，但架子堆得太满，厨房就站不下更多桌。真实的优化顺序是：先用共享内存把访存瓶颈吸掉，再回头调 block 大小和寄存器，让“省下的访存时间”和“多出来的灶台数”彼此平衡。

Sper-block=Sstatic+SdynamicS_{\text{per-block}} = S_{\text{static}} + S_{\text{dynamic}}
每块共享内存的总需求 = 静态声明的部分 + 动态声明的部分。它直接参与“调料架账”的计算，决定了 SM 能驻留多少块。
  1. S_static（静态共享内存）代码里用 __shared__ 声明的大小，编译期就确定，写在可执行文件里。
    SstaticS_{\text{static}}
  2. S_dynamic（动态共享内存）kernel 启动时由 host 端额外指定的共享内存字节数，灵活但需要程序员手动管理。
    SdynamicS_{\text{dynamic}}

Nsight分析

想成为老总厨，光靠感觉不够，得装监控：每个灶台多少锅在冒火、案板够不够用、哪道菜在拖时间，一眼看清。NVIDIA Nsight Compute（简称 ncu）就是 CUDA 工程师的“厨房监控仪表盘”。

在命令行执行 ncu --set full ./mykernel，它会输出每个 kernel 的详细报告。最核心的两行是：Theoretical Occupancy（理论占用率，由资源静态计算得出）和 Achieved Occupancy（实际占用率，运行中采样得出）。如果两者接近，说明调度充分；如果实际远低于理论，说明有分支发散、__syncthreads 等待或内存停顿在拖后腿。

报告中还会直接告诉你“是哪只桶最短”：Launch Stats 页面会列出每线程寄存器数、每块共享内存数、block 大小，以及当时的限制因素（Limit Registers / Limit Shared Memory / Limit Warps）。这就省去了手工计算的麻烦——但看懂原理才能根据建议做正确取舍。实际占用率与理论值的关系如下：

仪表盘告诉你两件事：理论上能开几桌（Theoretical），实际开了几桌（Achieved）。两者之间的差距，就是后厨流程里的“摸鱼时间”——可能是等食材、可能是分工不均、可能是灶台间互相等待。监控数据会像侦探一样，替你指出这些时间藏在哪一行代码里。

Achieved=W‾residentWmax×100%\text{Achieved} = \frac{\overline{W}_{\text{resident}}}{W_{\text{max}}} \times 100\%
实际占用率 = 采样周期内平均驻留 warp 数 ÷ 最大 warp 数。它反映的是 GPU 真实工作状态，而不是资源静态推算值。
  1. 平均驻留 warp 数Profiler 在运行期间周期采样 SM 上“此刻正在运行”的 warp 数并取平均，它会因为分支、同步和访存停顿而低于理论值。
    W‾resident\overline{W}_{\text{resident}}
  2. 除以最大 warp 数与硬件上限比较，得到真实利用率。这个数字才是你优化后真正想拉高的指标。
    WmaxW_{\text{max}}

实战策略

优化不是把每个旋钮都拧到最大，而是先判断餐厅到底是“等食材”（访存瓶颈）还是“炒菜慢”（计算瓶颈）。前者多开灶台能藏住等待；后者灶台全开也救不了火——因为计算单元已经满载，再加线程只是增加排队。

先看 Nsight 报告里的 Memory Workload 与 Compute Workload：如果访存单元利用率远高于计算单元，这是 memory-bound，优先提高占用率来隐藏延迟，手段包括压寄存器、调 block 大小；如果计算单元已经 90% 以上，这是 compute-bound，占用率再高也白搭，应该优化算法本身，比如减少冗余计算、提高指令级并行。

如果要判断“到底需要多少个 warp 才够隐藏延迟”，可以用一个粗略公式估算：

实战调优建议按这个次序走：先定位瓶颈类型，再查理论占用率，然后每次只改一个变量（比如只调 __launch_bounds__ 或只换 block 大小），用端到端耗时做 A/B 对比。记住：占用率是手段，不是目的。最终指标永远是“这道菜从下单到上桌的总时间”——也就是 kernel 的墙钟时间。

Nwarp-needed=LmemoryIIN_{\text{warp-needed}} = \frac{L_{\text{memory}}}{\mathit{II}}
需要的 warp 数 ≈ 访存延迟周期 ÷ warp 平均发射间隔。比如访存要 400 周期、每个 warp 每 8 周期发一条指令，就需要约 50 个 warp 轮换，才能保证计算单元不出现空白。
  1. L_memory（访存延迟）一次全局内存访问等待的时钟周期数，通常几百个 cycle，是 GPU 隐藏延迟的目标对象。
    LmemoryL_{\text{memory}}
  2. II（发射间隔）一个 warp 连续发射两条独立指令的平均间隔周期数。间隔越短，单个 warp 越“忙”，需要的替补 warp 就越少。
    II\mathit{II}
  3. 两者相除用延迟除以间隔，得到填满等待期所需的最小并行 warp 数。如果硬件上限小于这个数，说明从根上就无法完全隐藏延迟。
    LmemoryII\frac{L_{\text{memory}}}{\mathit{II}}

AI 应用场景

大模型推理的 batch size 调度

阶段：模型部署与推理优化

解决的问题：LLM 推理时，小 batch 场景下 GPU 的 SM 占用率经常只有 20%~30%，好比大厨房只接了两桌客人。传统做法是把 batch 撑大，但显存不够时反而 OOM。用占用率分析指导 batch 与 kernel 配置，可以在显存约束内让 SM 满转。

具体案例：某 7B 模型在 A100 上做 decode 推理，原本 batch=8 时每个 GEMM kernel 的占用率只有 30%。通过分析发现寄存器用量过高导致 block 数受限，用 __launch_bounds__ 压到每线程 40 寄存器后，占用率提升到 70%，单卡吞吐提高到 1.8 倍。

训练框架中的算子融合

阶段：训练性能优化

解决的问题：PyTorch 默认逐个算子启动 kernel，每次启动都有开销且资源分配不均，就像每道菜都单独开火。把多个算子融合成一个 kernel，再为融合后的 kernel 精细调 block 大小和寄存器上限，能显著提升 SM 利用率。

具体案例：Megatron 训练框架中，把 LayerNorm + Residual Add + 激活函数融合成单个 kernel，融合后结合 __launch_bounds__ 将每线程寄存器从 56 压到 32，kernel 耗时降低 37%，端到端训练速度提升约 12%。

Nsight Compute 日常性能巡检

阶段：性能分析定位

解决的问题：程序跑得慢，但不知道是哪个 kernel 慢、为什么慢。传统方式是靠猜，逐个改参数反复试。用 ncu 一键 Profile，按占用率和 stall 原因给 kernel 排序，直接定位“案板不够”还是“等食材太久”。

具体案例：某推荐系统训练任务中，ncu --set full 报告显示 top-1 热点 kernel 是 embedding 查找，占用率只有 25%，原因是每线程寄存器高达 128 导致 warp 数不足。将 block 从 128 调到 256 并配合 __launch_bounds__ 后，kernel 时间从 4.2ms 降到 1.9ms。

Python 演示

python42 行
def compute_occupancy(block_size, regs_per_thread, smem_per_block,                      max_threads=2048, max_warps=64, max_blocks=32,                      reg_file=65536, smem_capacity=98304):    # 每个 block 对应的 warp 数    warp_size = 32    warp_per_block = block_size // warp_size     # 第一笔账: 线程数上限    blocks_by_thread = max_threads // block_size    warps_by_thread = blocks_by_thread * warp_per_block     # 第二笔账: 寄存器文件上限    warps_by_regs = reg_file // (regs_per_thread * warp_size)     # 第三笔账: 共享内存上限    if smem_per_block > 0:        blocks_by_smem = smem_capacity // smem_per_block    else:        blocks_by_smem = blocks_by_thread    warps_by_smem = blocks_by_smem * warp_per_block     # 额外: 硬件块数上限    blocks_by_hw = min(blocks_by_thread, max_blocks)    warps_by_blocks = blocks_by_hw * warp_per_block     # 三笔账取最小, 合成最终活跃 warp 数    warps_active = min(warps_by_thread, warps_by_regs,                       warps_by_smem, warps_by_blocks, max_warps)    occupancy = warps_active / max_warps * 100    return occupancy, warps_by_thread, warps_by_regs, warps_by_smem  def main():    print('block_size regs smem occupancy')    for bs in (128, 256, 512):        for regs in (32, 64, 128):            occ, wt, wr, ws = compute_occupancy(bs, regs, 0)            print('%9d %4d %4d %8.1f%%' % (bs, regs, 0, occ))  if __name__ == '__main__':    main()
目的：用 Python 模拟 CUDA 占用率的“三笔账”计算，输入 block 大小、每线程寄存器数、每块共享内存申请量，输出理论占用率，方便你在改写 kernel 前快速做估算。
  1. warp_per_block = block_size // warp_size先求每桌团餐对应几桌灶台（几个 warp）。
  2. warps_by_thread = blocks_by_thread * warp_per_block第一笔账：线程数上限能支撑多少个 warp。
  3. warps_by_regs = reg_file // (regs_per_thread * warp_size)第二笔账：寄存器总量除以“每 warp 案板需求”。
  4. warps_by_smem = blocks_by_smem * warp_per_block第三笔账：共享内存能放的块数换算成 warp。
  5. warps_active = min(warps_by_thread, ...)木桶效应：三笔账取最小值才是真实上限。
  6. occupancy = warps_active / max_warps * 100除以硬件最大 warp 数，得到 0~100% 的占用率。

常见误区

误区我把占用率从 50% 拉到 100% 后，kernel 怎么反而变慢了？不应该越快吗？

正解占用率只是“并行度足够让硬件忙起来”的一个度量。如果 kernel 是计算密集且计算单元已经满转，提高占用率只会增加线程间的资源竞争，甚至引发寄存器溢出。好比 8 个灶台全开但厨师只有 2 个，火再旺也没人炒菜。占用率高还要配合足够的访存压力才有效：memory-bound 时提高占用率收益明显，compute-bound 时几乎没有收益。

为何常见：教程常说“占用率低就优化”，让人误把指标当成目标。实际上要看瓶颈类型，占用率只是手段。

误区我把所有 kernel 的 block 都设成 256，这个值到处都在用，应该没问题吧？

正解256 只是通用起点，不是万能解。如果 block 内每个线程需要大量寄存器，或者频繁执行 __syncthreads() 屏障，块太大会导致驻留块数过少、同步等待占比高。合理的块大小应保证每 SM 至少能驻留 4~8 个 block，给调度器足够的“换手”弹性。

为何常见：模板代码和示例普遍从 256 起步，缺少对不同资源压力的理解，导致照搬。

误区我加了 `__launch_bounds__(256, 8)`，占用率确实上去了，但 kernel 慢了一半，编译器是不是有问题？

正解编译器为了满足 minBlocksPerMultiprocessor，会强行把寄存器压到目标值。但局部变量放不下时会被 spill 到 local memory，而 local memory 实际在全局内存里，一次 spill 可能多出几百周期延迟。正确做法是先用 ncu 确认寄存器确实是瓶颈，再用 __launch_bounds__ 适度压缩，比如从 40 压到 32，而不是一步压到 16。

为何常见：看到占用率低就下意识想用 __launch_bounds__，忽略了 spill 的代价。

误区共享内存是限制因素，所以我把每块共享内存减到最小，占用率上去了，数据复用我不管了。

正解共享内存的价值是消除全局内存重复访问。省 smem 导致全局内存访问量暴涨，DRAM 带宽成为新瓶颈，整体反而更慢。优化第一优先级是让算法先在 SM 内完成任务（复用数据），第二才是调占用率。先保复用，再压占用，顺序不能反。

为何常见：人们把“资源限制”理解为“资源越少越好”，忽略资源是拿来换取访存效率的。

误区Nsight 报告里 Theoretical Occupancy 是 100%，那我就已经榨干 GPU 了。

正解理论占用率只反映静态资源能装下多少 warp；实际运行中分支发散、__syncthreads() 等待、访存延迟都会让 Achieved Occupancy 明显低于理论值。如果实际只有 60%，要去看 warp stall 的原因，而不是满足于理论数字。

为何常见：报告列太多，只认最显眼的百分比，忽略了“理论”和“实际”两个词的本质区别。

面试题

Q1. GPU 占用率最接近下列哪个比喻？

单选题 · easy

  • A.单核 CPU 的频率高低
  • B.显存总容量的大小
  • C.灶台被利用的占比← 正确答案
  • D.网络带宽的吞吐率

答案：C。占用率衡量的是 SM 上活跃 warp 占最大容量的比例，好比同一时间灶台被使用了几成。CPU 频率、显存容量、网络带宽描述的是“速度”或“容量”，与“利用率”是不同维度。

Q2. CUDA 中一个线程束 warp 包含多少线程？

单选题 · easy

  • A.64 个线程
  • B.16 个线程
  • C.32 个线程← 正确答案
  • D.8 个线程

答案：C。warp 是 GPU 调度的基本单位，固定 32 个线程，32 个线程同时执行同一条指令。这就像一个大厨同时管 32 口锅。其他数字是常见的误导项。

Q3. 某 SM 寄存器文件共 65536 个寄存器，每线程使用 64 个寄存器，每块 256 线程。寄存器维度下每 SM 最多驻留几个 block？

单选题 · mid

  • A.2 个块
  • B.4 个块← 正确答案
  • C.8 个块
  • D.16 个块

答案：B。65536 ÷ (64 × 256) = 4。案板总量除以每桌（每块）的总需求。注意单位要一致：寄存器总个数除以“每线程寄存器数 × 每块线程数”，得到的是块数。

Q4. 下列哪些属于限制 GPU 占用率的硬件资源？

多选题 · mid

  • A.每 SM 最大线程数← 正确答案
  • B.寄存器文件大小← 正确答案
  • C.共享内存容量← 正确答案
  • D.全局显存带宽

答案：ABC。占用率的三笔账来自线程数、寄存器、共享内存三个容量型资源。全局显存带宽是速度型指标，它决定访存快慢，但不限制“能驻留多少 warp”，所以不属于占用率的直接限制因素。

Q5. __launch_bounds__(256, 8) 中第二个参数 8 的含义是什么？

单选题 · mid

  • A.每个块最多 8 个线程
  • B.每个 SM 至少驻留 8 个块← 正确答案
  • C.每个线程最多用 8 个寄存器
  • D.每个 SM 最多容纳 8 个块

答案：B。__launch_bounds__ 的第一个参数是每块最大线程数，第二个参数 minBlocksPerMultiprocessor 是“期望每个 SM 至少驻留多少块”。编译器据此反推每线程寄存器上限。注意是“至少”而不是“最多”。

Q6. 某个访存密集 kernel，提高占用率后性能反而下降，最可能的原因是什么？

单选题 · hard

  • A.共享内存被耗尽
  • B.缓存命中率下降
  • C.寄存器溢出局存← 正确答案
  • D.带宽已近饱和

答案：C。提高占用率最常见的做法是压寄存器。压过头导致局部变量溢出到 local memory（实际在全局内存），单次访问延迟暴涨，性能必然下降。共享内存耗尽会导致启动失败而不是变慢；带宽饱和只会让提升不明显，不会“下降”。

Q7. 哪些情况下低占用率也能获得接近峰值性能？

多选题 · mid

  • A.计算密集且指令级并行高← 正确答案
  • B.访存延迟极高的长尾访问
  • C.内核极短启动开销主导← 正确答案
  • D.分支发散极其严重的任务

答案：AC。高指令级并行（ILP）时，单个 warp 自身就能连续发射指令填满流水线，不需要大量 warp 换手。kernel 极短时启动开销占主导，占用率提升的收益被开销淹没。而高延迟访存和严重分支发散恰恰需要更多 warp 来填补等待空隙，低占用率会更糟。

Q8. Nsight Compute 中 Achieved Occupancy 表示什么？

单选题 · easy

  • A.计算单元的平均利用率
  • B.理论上的最大占用率
  • C.显存带宽的实际利用率
  • D.实际采样到的驻留 warp 占比← 正确答案

答案：D。Achieved Occupancy 是 profiler 在运行期间采样到的平均驻留 warp 占比，反映真实运行状态。理论占用率是另一个指标，两者不可混淆。

Q9. 关于 block size 的选择，下列说法正确的有哪些？

多选题 · hard

  • A.最好是 32 的整数倍← 正确答案
  • B.越大占用率必然越高
  • C.应兼顾资源与调度弹性← 正确答案
  • D.固定 1024 一定最优

答案：AC。block 是 warp 的整数倍才不会浪费半个 warp 的槽位；同时要兼顾寄存器、共享内存的占用和每 SM 的块数上限，保证调度弹性。“越大越高”和“固定最优”都是绝对化错误认知。

Q10. 某 SM 最大支持 64 个 warp，访存延迟 500 周期，每个 warp 平均每 8 周期发射一条指令。要完全隐藏访存延迟，至少需要多少活跃 warp？

单选题 · extreme

  • A.32 个 warp
  • B.63 个 warp← 正确答案
  • C.64 个 warp
  • D.不需要这么多

答案：B。隐藏延迟需要 warp 数 = 延迟 ÷ 发射间隔 = 500 ÷ 8 = 62.5，向上取整为 63 个。少于 63 个时，某些周期会出现计算单元无指令可发的空白；63 个刚好在硬件 64 上限内，理论上能做到完全隐藏。