第三章 Nsight Compute深度分析：Kernel分析、Occupancy分析、Memory分析、指令分析
各位同学，欢迎来到第三章。这一章我们直接切入核心——Nsight Compute。说实话，这工具我用了快十年了，每次新项目上线前，我都要跟它「深度交流」几个小时。它就像GPU的CT机，能把你的kernel里里外外扫描个遍。
很多同学问我：「老师，我写了个kernel，跑起来感觉慢，但不知道慢在哪。」嗯，这就是典型的「黑盒焦虑」。Nsight Compute就是来帮你打开这个黑盒的。今天我们就从四个维度来解剖它：Kernel分析、Occupancy分析、Memory分析、指令分析。
3.1 Kernel分析：先看整体，再抠细节
打开Nsight Compute，第一个映入眼帘的就是Kernel Summary。我个人习惯先看这里，就像医生看病先看体温一样。
核心指标一览：
- Duration：kernel总耗时，单位微秒
- Grid Size & Block Size：线程网格和线程块配置
- Registers Per Thread：每个线程使用的寄存器数
- Shared Memory Per Block：每个block使用的共享内存
我记得有一次，一个学生跑过来跟我说他的kernel耗时200微秒，优化了一周还是200微秒。我让他打开Nsight Compute一看——好家伙，寄存器用了64个，occupancy直接掉到25%。这就是典型的「寄存器压力」问题。
为什么会这样？你想想看，每个SM的寄存器总量是固定的（比如V100是65536个）。如果每个线程用太多寄存器，能同时驻留的线程数就少了。说白了，就是「人多了住不下」。
我的经验：先看Duration，如果超过100微秒，大概率有优化空间。然后看Registers Per Thread，超过32就要警惕了。
3.2 Occupancy分析：别被数字骗了
Occupancy，中文叫「占用率」。很多教程告诉你occupancy越高越好，但我在项目中吃过亏——有一次我把occupancy从33%优化到了100%，结果性能反而下降了15%。
为什么会这样？因为occupancy只是「理论上的并行度」，不代表实际效率。你想想看，如果每个线程都在等内存，occupancy再高也没用。
| Occupancy区间 | 典型瓶颈 | 我的建议 |
|---|---|---|
| 0% - 33% | 寄存器或共享内存限制 | 减少寄存器使用，或拆分kernel |
| 33% - 66% | 可能受计算或内存限制 | 结合其他指标判断 |
| 66% - 100% | 通常不是occupancy问题 | 关注指令级并行度 |
Nsight Compute里有个「Occupancy Analysis」页面，它会告诉你限制因素是什么。我曾经遇到一个案例，限制因素是shared memory——每个block用了48KB，而V100每个SM只有96KB，所以最多同时跑2个block。解决方案？把shared memory砍到32KB，occupancy从25%跳到了50%。
避坑指南：我曾经为了追求高occupancy，把一个大kernel拆成三个小kernel。结果呢？kernel launch overhead增加了，总耗时反而多了。所以，occupancy不是唯一指标，要结合latency一起看。
3.3 Memory分析：带宽利用率是王道
好，接下来是重头戏——Memory分析。说白了，GPU计算大部分时候都是「内存墙」问题。计算再快，数据搬不过来也是白搭。
Nsight Compute的Memory Workload Analysis页面会给你几个关键数字：
- Memory Throughput：实际内存吞吐量
- L1/L2 Cache Hit Rate：缓存命中率
- Sector Utilization：扇区利用率（这个很重要！）
我记得有一次优化一个矩阵乘法kernel，理论带宽是900GB/s，实际只跑了200GB/s。打开Sector Utilization一看——只有40%。这意味着每次内存访问，有60%的带宽被浪费了。
为什么会这样？因为GPU内存访问是以128字节为单位的（一个cache line）。如果你的数据布局不连续，比如按列访问行主序的数组，那每次只用到了一小部分数据。
解决方案：
- 确保全局内存访问是合并访问的（coalesced access）
- 使用共享内存做数据重排
- 考虑使用向量化加载（float4/int4）
这里有个小技巧：在Nsight Compute的「Source」页面，它会用颜色标记每一行代码的内存访问效率。绿色表示高效，红色表示低效。我一般先扫一眼红色区域，那里就是优化的突破口。
3.4 指令分析：看看GPU在忙什么
最后一个维度——指令分析。这个很多人会忽略，但恰恰是「压榨最后一滴性能」的关键。
Nsight Compute的Instruction Analysis页面会告诉你：
- IPC (Instructions Per Cycle)：每个时钟周期执行的指令数
- Warp Stall Reasons：warp停顿的原因
- Compute Workload Analysis：计算指令的分布
IPC的理想值是多少？对于V100，理论峰值是4（每个时钟周期发射4条指令）。但实际项目中，能跑到2.5就算不错了。我见过最差的kernel，IPC只有0.3——那基本就是在「空转」。
Warp Stall Reasons是我最常看的页面。它会告诉你warp为什么停顿：
| 停顿原因 | 含义 | 优化方向 |
|---|---|---|
| Long Scoreboard | 等待全局内存加载 | 提高数据局部性，使用缓存 |
| Short Scoreboard | 等待共享内存或寄存器 | 减少bank conflict，优化寄存器使用 |
| Not Selected | warp调度器选择了其他warp | 提高occupancy，增加并行度 |
| Sync | 等待__syncthreads() | 减少同步点，或使用更细粒度的同步 |
我曾经优化过一个图像处理kernel，IPC从0.8提升到了2.1。怎么做到的？其实就是把一些重复计算从运行时搬到了编译时——用常量内存存储查找表，用__constant__修饰。指令数没变，但停顿时间大大减少了。
我的习惯：先看IPC，如果低于1.0，说明有严重的停顿问题。然后看Warp Stall Reasons，找到占比最高的停顿原因，针对性优化。
3.5 实战：一个完整的分析流程
好了，理论讲完了。我给大家梳理一个我常用的分析流程：
- 第一步：跑Nsight Compute的默认Profile，看Kernel Summary
- 第二步：看Duration，如果超过预期，进入下一步
- 第三步：看Occupancy，如果低于50%，检查寄存器/共享内存限制
- 第四步：看Memory Throughput，如果低于理论带宽的50%，检查Sector Utilization
- 第五步：看IPC，如果低于1.5，检查Warp Stall Reasons
- 第六步：定位到具体代码行，在Source页面看颜色标记
这个流程我用了好几年，基本能覆盖90%的性能问题。当然，有些特殊情况需要单独处理，比如原子操作冲突、动态并行等，那些我们后面章节再讲。
记住一句话：Nsight Compute不是用来「证明你的kernel很快」的，而是用来「发现你的kernel哪里慢」的。心态摆正了，工具才能用好。
好，这一章就到这里。下一章我们会深入TensorRT的模型优化，到时候会用到今天讲的这些分析方法。大家先把Nsight Compute装好，跑几个自己的kernel试试。有问题随时问我。
公众号:蓝海资料掘金营，微信deep3321