Nsight Compute 实操梳理：理论占用率 vs 硬件实测占用率

核心主旨：
用寄存器、共享内存算出来的理论 SM 占用率只是硬件并行上限，只代表资源够不够；访存延迟、Warp 分支发散、指令依赖这类动态运行问题，哪怕资源完全空闲，照样会让 GPU 算力大量浪费，出现「理论数值很高，实际跑很慢」。Nsight Compute 通过硬件计数器拿到的实测 Achieved Occupancy，才是 Kernel 真实运行状态。
一、两个核心概念区分
理论 SM 占用率
输入：每个线程寄存器用量、Block 共享内存大小。
输出：SM 最多能够容纳多少个活跃 Warp，算出并行上限。
只评估静态硬件资源：寄存器、共享内存。
不考虑：显存访问、分支跳转、指令等待、流水线阻塞。
含义：硬件允许最多同时跑多少 Warp，不等于 Warp 真的一直在干活。
实测 SM 占用率（Achieved Occupancy，Nsight 硬件计数器采集）
统计整个 Kernel 生命周期，处于就绪 / 执行状态 Warp 的时间占比。
Warp 大量时间处于 Stalled 停滞：等 HBM 显存、等共享内存、分支掩码、指令数据依赖，就算寄存器、共享内存还有空余，Warp 也会挂起不发射指令。
二者大小关系：实测占用率 ≤ 理论占用率
两者数值接近：瓶颈偏向计算本身
两者差距巨大：瓶颈来自访存、分支发散、指令依赖等动态阻塞
经典坑例：理论占用率 75%，Nsight 实测只有 35%，硬件资源充足，但线程大部分时间在等待。
二、Nsight Compute 实操流程
1、编译配置
nvcc 编译参数
–generate-line-info -Xptxas=-v -O3
-Xptxas=-v输出regPerThread，用来和理论占用率 Python 脚本做对照
被测 Kernel 输入规模要足够，运行时长 > 几十 ms，硬件计数器采样才可信；小 Kernel 会出现统计失真。
2、命令行采集（无需 GUI）

ncu -o kernel_report ./your_program

输出.ncu‑rep报告文件，Nsight Compute GUI 打开分析。
3、重点观测指标页：Kernel‑Level → Solver → Occupancy
Theoretical Occupancy：Nsight 计算的理论占用率，应和 Python 脚本结果对齐
Achieved Occupancy：硬件统计真实实测占用率
Warp Stall Reasons：调优最核心，看 Warp 停滞原因占比
判断逻辑：
Theoretical ≈ Achieved：Warp 多数时间执行指令，瓶颈偏向计算
Theoretical >> Achieved：大量 Warp 停滞，依据 Stall 分布定位访存 / 分支 / 指令依赖问题
4、常见 Warp Stall 类型
Stall Memory Dependence：等待全局显存 / 共享内存，最常见
Stall Branch Divergence：Warp 内部线程分支发散，线程掩码造成额外开销
Stall Instruction Dependence：前后指令存在数据依赖，等待前一条指令输出
Stall Resource Conflict：ALU、浮点执行单元硬件排队冲突
Stall Not Selected：存在就绪 Warp，但调度器未选中，出现概率低
Stall 不是报错，是 GPU 固有机制；GPU 靠切换就绪 Warp 掩盖延迟，但 Stall 占比过高直接性能暴跌。
三、两大性能杀手：访存延迟、Warp 分支发散
3.1 访存延迟
HBM 显存访问延迟数百时钟周期。Warp 发起全局内存读取后直接进入 Stall，让出硬件单元；GPU 靠切换其他就绪 Warp 掩盖等待延迟。
掩盖延迟的前提：SM 内部要有足够多就绪活跃 Warp。
三种典型场景
场景 A：理论占用率高，SM 大量活跃 Warp。一部分 Warp 等显存，调度器切其他 Warp 干活，延迟被掩盖，实测接近理论值。
场景 B：理论占用率本身很低。少量 Warp 同时发起访存，没有多余 Warp 切换，SM 空转；静态资源瓶颈叠加访存 Stall，双重性能打击。
场景 C：理论占用率很高，但全部 Warp 同一时刻发起大规模全局访存，全体集体进入 Memory Stall，无工作线程可切换。现象：理论 80%，实测掉到 30‑40%，硬件资源空闲，但全部线程在等内存返回。
访存问题特征：
Stall Memory Dependence占比＞40%；显存带宽已经打满，但 SM 利用率不高；扩大 batch，性能提升不成线性，被 HBM 带宽锁死。
优化方向：
拉高理论占用率，保证有充足 Warp 用来掩盖访存延迟
使用 shared memory 做 Tiling 分块缓存，减少全局显存访问次数
保证全局内存合并访问 coalesced，杜绝碎片化访存
只读数据使用__ldg__、const修饰，启用只读缓存
3.2 Warp 分支发散 Divergence
GPU SIMT 机制：1 个 Warp 固定 32 线程。
同一个 Warp 内，一部分线程走 if、一部分走 else，硬件会完整执行两条分支，依靠掩码屏蔽不需要执行的线程，实际工作量直接翻倍。
关键：分支发散不消耗寄存器、不占用共享内存。理论占用率的静态计算完全感知不到这个问题。哪怕理论占用率 100%，分支发散严重也会性能腰斩。
Nsight 特征：Stall Branch Divergence占比显著抬高。
高发业务：图像 Kernel 边界判断、稀疏算子、Attention Mask 掩码、Padding 较多的大模型推理。
优化方向：
判断逻辑提升至 Warp 粒度，让整个 Warp 所有线程走同一条分支路径
边界逻辑拆分：绝大多数 Block 走无边界校验快速路径，少量边缘 Block 单独处理
减少 Warp 内部多层 if‑else 嵌套
大模型推理：Token 重排，把 Padding 集中到少数 Warp，避免大量 Warp 出现分支发散
四、实战案例：理论占用率 96.9%，实测仅 38%
硬件：RTX5090
Kernel：读写 + 简单计算，Block=256，寄存器、共享内存开销很小
脚本计算：MaxWarps_reg=62，MaxWarps_smem=64，ActiveWarps=62，理论占用率 96.9%
Nsight 采集结果：
Theoretical Occupancy：96.9%
Achieved Occupancy：38%
Warp Stall：Memory Dependence 占比 62%
根因：内存访问模式差，非合并访存；大量 Warp 同时等待显存，没有足够可调度工作线程掩盖访存延迟。静态资源充足，但动态访存拖垮性能。
优化动作：调整内存访问顺序实现合并访问，增加 shared memory tile 缓存。
优化后：Achieved Occupancy 提升 79%，整体吞吐提升 2.1 倍。
启示：只看寄存器、共享内存做调优，会漏掉访存、分支发散这类重大性能缺陷。
五、标准化调优排查流程
Python 脚本算出 Kernel 理论 SM 占用率。Nsight 采集报告，拿到理论占用率、实测占用率、全套 Warp Stall 分布。
如果理论占用率本身偏低：优先优化寄存器、共享内存，解决静态资源瓶颈。
如果理论占用率高，实测占用率低，看 Stall 占比：
Memory Dependence 为主 → 访存问题：合并访问、shared tile、__ldg__只读缓存
Branch Divergence 为主 → 分支发散：重构条件逻辑，Warp 粒度判断
Instruction Dependence 为主 → 指令数据依赖，调整指令调度、复用中间变量
如果实测占用率已经很高，但整体性能依旧上不去：
基本触达硬件物理上限：HBM 显存带宽、FP 算力峰值，Kernel 层面调优收益很小，需要上层业务逻辑优化。
重要提醒：高实测占用率不等于高性能。Warp 虽然一直在跑，但如果执行碎片化低效访存，依旧打不满硬件峰值；占用率只是观测指标，不能作为唯一评判标准。
六、复现 Demo 要点（divergent_demo.cu）
两个 Kernel 寄存器、共享内存开销极低，理论 SM 占用率接近 100%。
divergent_kernel：线程粒度 if‑else，threadIdx.x & 1，同一个 Warp 内部奇偶线程分走不同分支，强分支发散。
Nsight 现象：理论占用率接近 100%；Achieved Occupancy 大幅下跌；Stall Branch Divergence 占比很高；ALU 有效利用率下降。
no_divergent_kernel：判断提升至 Warp 粒度warp_id_in_block &1，整个 Warp 全部线程共用一条分支路径，消除发散。
现象：理论占用率不变，Branch Stall 几乎消失；实测占用率上涨，运行速度接近快一倍。
编译 & 采集命令

nvcc divergent_demo.cu -o divergent_demo --generate-line-info -Xptxas=-v -O3
ncu -o divergent_report ./divergent_demo

注意输入数据量足够，保证 Kernel 运行几十 ms 以上，硬件计数器采样有效。
最终总结
理论占用率：寄存器 + 共享内存算出的硬件并行上限，静态模型，看不见访存、分支