GPU 占用率 occupancy：寄存器压力、shared memory 与 warp 驻留的三方账

更新时间：2026-09-25。本文是「集成电路大类 · 芯片设计」进阶层（intermediate） 的一篇。写 CUDA 的人几乎都被问过一个问题：你的 kernel 占用率（occupancy）是多少？这个数字经常被当成 GPU 程序的「跑分」，仿佛越高越好。但真打开 profiler 会发现，有些 kernel 占用率不到 30% 依然跑满频率，有些冲到 100% 却在显存带宽上堵成一团。占用率到底衡量什么、被什么限制、为什么不是越高越好，是本篇要算清的三方账。warp 调度器如何用切换隐藏延迟见 GPU 架构深入：SIMT、warp 调度与内存层级，warp 锁步执行与分支发散见 SIMT 分支发散。

本文要回答的问题

• occupancy 的准确定义是什么？SM 上「活跃 warp」和「驻留 warp」是一回事吗？
• 高占用为什么能隐藏内存延迟？要藏住 400 到 800 周期的显存访问，大概需要多少 warp？
• 一个 SM 有 65536 个寄存器，每线程多用几个寄存器，驻留 warp 数具体怎么掉？
• shared memory 和 block 数这另外两道闸门是怎么参与算账的？
• 寄存器溢出（spilling）到 local memory 为什么是最糟的一种「省占用率」方式？
• 为什么计算密集型 kernel 低占用也能满速？occupancy calculator 怎么用？

一、占用率是什么：SM 上驻留 warp 数与硬件上限的比值

先给定义：占用率（occupancy）= SM 实际驻留的 warp 数 / 该 SM 最多能驻留的 warp 数。注意用词是「驻留（resident）」而不是「正在执行」。一个 SM 一般只有 4 个 warp 调度器、每拍最多发射 4 条指令，任何瞬间真正在执行的 warp 只有几个；但 SM 上可以同时驻留几十个 warp，它们各自的寄存器和线程状态都实实在在占在硬件里，等待调度器挑选。占用率描述的是池子的充盈程度，不是此刻有多少 warp 在干活。

典型数字先给出来，后文反复用到：

• 一个 SM 有 65536 个 32 位寄存器（约 256KB 寄存器堆），这是所有驻留线程瓜分的总盘子；
• SM 的 warp 驻留上限随架构不同，常见为 48 或 64 个 warp，对应每 SM 1536 或 2048 个线程；
• block 维度也有上限，常见每 SM 最多 16 或 32 个 block 同时驻留；
• shared memory（共享内存）每 SM 容量从几十 KB 到上百 KB，例如 A100 每 SM 有 164KB。

线程如何组织成 grid/block/warp、一个 block 又怎么被切成 warp，见 CUDA 线程层次：grid、block 与 warp。这里要抓住一个推论：block 是调度的不可分割单位——一个 block 的资源一旦分配，它的全部 warp 必须整体驻留直到执行结束，SM 不能只要一个 block 的一半。所以占用率的计算从「单个 block 消耗多少资源」开始，而不是从单线程开始。

还要区分两个在 profiler 里常见的指标：理论占用率（theoretical occupancy） 由资源配置算出，kernel 启动前就能确定，是上限；实测占用率（achieved occupancy） 是运行中活跃 warp 对时间的平均，除了资源限制，还受 grid 里 block 总数的影响——如果整个 kernel 只启动了几个 block，分摊到几十个 SM 上，每个 SM 可能只分到一个 block，理论占用率再高，实测也很低。这种尾部效应在数据量小、grid 形状狭长时最常见，属于启动规模问题，不是寄存器问题。

二、高占用的意义：让调度器手里始终有就绪 warp

GPU 为什么需要池子里有几十个 warp？因为它对抗长延迟的办法不是 CPU 那套乱序执行和大缓存，而是用并行度换延迟。一次全局显存访问的延迟约 400 到 800 个周期，warp 一旦发出访存指令就停住等数据。此时调度器若手里还有其他就绪 warp，立刻切过去发射——warp 的寄存器常驻寄存器堆，切换没有保存恢复成本，发射槽因此始终被填满。SM/CU 的调度器与发射结构见 SM/CU 内部结构，显存带宽与各层延迟见 GPU 内存层级与带宽。

读图抓一个点：占用率低首先伤的不是并行度，是「延迟隐藏」。一个简化估算（这就是经典的 Little 定律在 GPU 上的用法）：若 warp 平均每 10 个周期需要一次切换、每次停顿 400 周期，那么想让发射槽不停摆，池子里至少要有 400 / 10 = 40 个处于不同等待阶段的 warp。这解释了为什么经验上内存密集型 kernel 追求 50% 以上的占用率（64 warp 上限时约 32 个），而不是非要 100%。访存越密集、每次访存间隔的指令越少，需要的 warp 越多；反过来一个 warp 若在两次访存之间能垫上几十条算术指令，它自己就盖住了大半延迟。

还要注意「驻留」不等于「就绪」：被 __syncthreads() 屏障挡住、等数据返回的 warp 占着名额却不能被调度，实际可切换的 warp 比占用率显示的更少。所以同一个占用率数字，在 barrier 频繁的 kernel 里和在无同步的 elementwise kernel 里，含金量并不相同——这也是为什么理论占用率只能当起点，最终要看 achieved occupancy 和 stall reason（停顿原因）的实测分布。

三、三道资源闸门：寄存器、shared memory、block 数

一个 block 能不能驻留、SM 上能同时放几个 block，由硬件对三类资源分别做除法，取最小值。这就是占用率的三道闸门。

闸门一：每线程寄存器数。 编译器为每个线程分配固定数量的寄存器，记为 R。SM 的 65536 个寄存器按驻留线程数瓜分：每 SM 最多线程数 = 65536 / R，再换算成 warp（除以 32）。账非常直观：

每线程寄存器 R	寄存器允许的线程数/SM	折合 warp（64 上限）	寄存器口径占用率
32	2048	64	100%
48	约 1365	42	约 66%
64	1024	32	50%
128	512	16	25%

实际硬件还有分配粒度：寄存器以每个 warp 256 个（折合每线程 8 个）为单位分配，R 会向上取整到 8 的倍数，所以 R 从 33 增到 40 的效果一样，跨过 40 才真正掉一档。这意味着 R 卡在 41、49、57 这类数字上最不划算——多出来的几个寄存器没带来实质能力，却掉了一整档驻留名额；反过来，把 R 从 41 压回 40，往往是零成本的一档占用率。编译器报告的 R 可以用 --ptxas-options=-v 编译选项查看，输出里每个 kernel 的寄存器数、shared memory 字节数都会列出，调优时建议先把这份编译日志存底保留，作为后续每轮对比的基线。

闸门二：每 block 的 shared memory。 若 block 声明使用 S 字节共享内存，SM 最多放 shared memory 总量 / S 个 block。一个用 48KB shared memory 的 block，在 164KB/SM 的 A100 上只能放 3 个；即使寄存器还很宽裕，驻留 warp 数也被这 3 个 block 锁死。

闸门三：每 SM 的 block 数上限。 即使每个 block 只用 1 个 warp、几乎不耗 shared memory，每 SM 也最多驻留 16 到 32 个 block。这惩罚的是 block 尺寸过小：16 个 block × 32 线程只有 512 个线程，占用率天然只有 25%。反过来把一个完整算例串成三个小节并走三道闸门：256 线程的 block、每线程 40 个寄存器，寄存器闸门允许约 6 个 block；若该 block 用 40KB shared memory，共享内存闸门只允许 4 个；block 上限是 16——最终驻留 4 个 block、共 32 个 warp，占用率 50%，瓶颈清清楚楚在 shared memory。这种逐道代入的习惯比盯着计算器的输出更能指导优化。

上图是硬件分配 block 的真实流程：三道闸门各算各的，谁最小谁说了算。优化占用率的第一步永远是先找出当前 kernel 卡在哪一道——盲目减寄存器，结果卡在 shared memory 上，纯属白费功夫。

四、寄存器溢出：降到 local memory 的「伪优化」

寄存器不够会怎样？注意硬件不会因此拒绝启动 kernel：当编译器发现一个线程需要的变量超过可分配寄存器，它会把放不下的变量溢出（spill）到 local memory。名字叫 local，物理上却挂在显存上，走的是全局内存的完整延迟路径，只是地址对该线程私有、不与别人共享。

溢出的代价是双重的。其一，被溢出的变量每次读写都变成一次显存访问，原本单周期的寄存器操作变成 400 周期以上的停顿，编译器还会生成 st.local/ld.local 存取指令把指令序列拉长；其二，溢出让 warp 更频繁地进入等待，于是需要更高的占用率才能藏住它自己制造的延迟——为了一点占用率指标把 kernel 改成大量溢出，是方向性错误：占用率数字可能变好看了，实测延迟却全面恶化。

举个具体例子：kernel 里声明 float tmp[64] 并按下标变量访问，编译器无法把它固定映射到寄存器，整个数组被放进 local memory——看似只是 64 个 float 的临时空间，实际每次循环访问都走一遍显存路径。常见的溢出诱因包括：kernel 里大的局部数组（尤其是下标不可预测的访问）、过深的调用链、循环内大量活跃变量，以及用 __launch_bounds__ 强行压低寄存器预算。排查上，ptxas 的编译日志会直接输出 X bytes spill stores, Y bytes spill loads，任何非零数字都值得警惕；profiler 里 local_load、local_store 指令占比、显存事务量以及 L1 命中率也能看到溢出的实际流量。顺带澄清一个误解：local memory 的数据通常经过 L1/L2 缓存，命中时延迟没有完整显存访问那么夸张，但它的容量和命中率都不可指望，绝不能当寄存器用。

反过来，这也暴露了「为占用率而占用率」的矛盾：寄存器用得少 → 能驻留更多 warp，但编译器可能因寄存器紧张而溢出；显式放宽寄存器（如降低占用率目标）让数据留在寄存器里，常常更快。寄存器是 GPU 上最快、离 ALU 最近的存储，没有之一。

五、高占用不是越高越好：一张权衡图与三类场景

占用率和性能的关系不是线性的，而是一条边际递减的饱和曲线：占用率从低往上升时，气泡被消除，性能快速改善；越过「足够藏住延迟」的拐点后，剩余 warp 基本只是备胎，性能趋于平台；再继续堆高占用，可能要逼迫每个线程少用寄存器、牺牲循环展开和变量常驻，性能甚至回落。

拐点的位置取决于 kernel 的「算数密度」——每单位访存做多少计算。计算越密集，每条访存指令前后垫的计算越多，单个 warp 自己就能把延迟填平，需要的备胎 warp 越少。

占用率水平	典型适用场景	该场景下的取舍要点
高占用（75% 到 100%）	内存密集、随机访存、单次访存间计算少：稀疏矩阵、图遍历、轻量 elementwise 算子、带宽型 reduction	warp 等待频繁，需要大量就绪 warp 轮流填槽；注意别把寄存器压到溢出
中占用（40% 到 75%）	规则访存加一定计算密度：tile 化矩阵乘、卷积、stencil、带 shared memory 复用的 kernel	延迟大体可藏，富余寄存器留给循环展开和 tile 内复用；多数 kernel 的甜区
低占用（25% 到 40%）	计算密集、寄存器重度使用：大 tile GEMM、寄存器分块、寄存器内多级展开	每线程用上百个寄存器换数据零访存，单 warp 指令级并行即可喂满流水线；低占用是主动选择

表里最反直觉的是最后一行：高性能矩阵乘的 kernel 常常只跑到 25% 到 50% 的占用率，因为每个线程持有数十个累加器和分块数据。算一笔账：一个线程若负责 8×8 的输出子块，仅累加器就要 64 个寄存器，再加地址、指针和预取的操作数，R 轻松超过 80，按前文的寄存器除法，驻留 warp 数自然压到 24 个以下。但这 80 个寄存器换来了每个输出元素只在显存与寄存器之间搬一次数据，显存流量下降一个数量级——对 GEMM 这种算术密度极高的 kernel，这笔买卖非常划算。

占用率是输入约束，不是优化目标；目标是让发射槽和显存带宽都不空着。也可以换个说法：占用率回答的是「硬件还剩多少床位」，而性能问题问的是「病人为什么没出院」——床位全满不代表周转快，床位空闲也可能是病人都在手术中。把这两个问题分开，调优时才不会南辕北辙。

六、occupancy calculator 与工程实操

把上面的算账固化下来的工具就是 occupancy calculator（占用率计算器）：选定 GPU 架构，输入每线程寄存器数、每 block 线程数、每 block shared memory，它直接给出可驻留 block 数和理论占用率，并标出卡在哪道闸门。CUDA Toolkit 里有电子表格版，Nsight Compute profiler 则在每次 profile 后直接给出 achieved occupancy（实测占用率）与理论上限的对比，还会列出限制因素。

写 kernel 时常用的三个调节手段：

• __launch_bounds__：在 kernel 定义上声明每 block 最大线程数和期望最小驻留 block 数，编译器据此约束寄存器分配。注意它是把双刃剑——逼编译器少用寄存器可能引发溢出，只在确认卡在寄存器闸门时使用。
• 调 block 尺寸：block 线程数取 32 的倍数（通常 128 或 256），同时避开 block 数上限闸门；用两三种 block 尺寸实测，比纸面计算可靠。
• 管 shared memory：能复用的数据尽量留在寄存器；shared memory 用量按驻留 block 数倒推，宁可每 block 少用一点换两个额外驻留 block。同时留意 bank conflict——shared memory 的容量闸门过了，访问冲突仍会拖慢，这是另一类问题。

最后仍是那条原则：先测量，再动手。先看 achieved occupancy 与 SM 活跃度，再看 warp 停在哪里：Nsight Compute 会把停顿原因分成 long scoreboard（等显存）、short scoreboard（等 math/L1）、barrier、wait（等指令依赖）等类别。若大头是 long scoreboard 且占用率低，提升占用率、合并访存才有效；若停顿以 short scoreboard 和依赖为主，瓶颈在指令流水和 ILP，堆 warp 无济于事；若显存带宽已经饱和，正确方向是减少访存量（分块、复用、压缩），而不是让更多 warp 去抢同一条总线。占用率只是这些诊断中的一个输入量。

这一篇在知识体系里的位置

本篇属于芯片设计进阶层，回答「SM 上为什么只能驻留这么多 warp，以及占用率与性能之间的真实关系」。向前补硬件模型：GPU 为什么走众核吞吐路线见 GPU 与 CPU：为什么 GPU 走向了众核吞吐，SM 内部的调度器与执行通路见 SM/CU 内部结构；横向配套：grid/block/thread 的组织见 CUDA 线程层次：grid、block 与 warp，warp 切换隐藏延迟与内存合并的全景见 GPU 架构深入：SIMT、warp 调度与内存层级，warp 内分支对利用率的另一种折损见 SIMT 分支发散，各层存储的带宽与延迟见 GPU 内存层级与带宽。再往后，内存合并与 bank conflict、roofline 性能建模是这条 GPU 线的后续专题。

总结

一句话总结：GPU 占用率 occupancy 是 SM 驻留 warp 数与上限（典型 32/64 warp）的比值，衡量的是调度器手里有多少备胎 warp——高占用的唯一意义是在 warp 等待 400 到 800 周期显存延迟时有就绪 warp 可零开销切换、把发射槽填满；驻留数由三道闸门各做除法取最小值：每线程寄存器数瓜分 65536 个寄存器（32 个/线程可满 64 warp，128 个只剩 16 个）、每 block shared memory 瓜分 SM 共享内存容量、每 SM block 数上限（16 到 32 个）；寄存器不够时变量溢出到挂在显存上的 local memory，单周期操作变成数百周期停顿，是最糟的伪优化；占用率与性能是边际递减关系，拐点由算数密度决定——内存密集型要 50% 以上，计算密集型用大量寄存器分块、25% 到 40% 占用即可满速；实操用 occupancy calculator 和 Nsight Compute 找限制闸门，用 __launch_bounds__、block 尺寸和 shared memory 用量调节，但永远先测量真实停顿与带宽，别把占用率数字本身当目标。