上周刷B站和知乎，我发现一个很反常的现象：短短一周之内，四五个不同的作者像约好了一样，密集更新同一类内容——Nsight Compute（下称NCU）性能分析教程。
数据也支持这个判断。8月12日，noinfra 的《以GEMM为例，学会使用Nsight Compute》在B站和小红书同步上线，要把 Nsight Compute 的每个面板彻底算清楚哔哩哔哩小红书。这条视频播放量不到两千，收藏却有380，收藏率接近20%。同一作者的续作《内存分析详解：Request、Sector、Wavefront》上线两天又被收了30多次哔哩哔哩。知乎上 Leezhou 的CUDA系列写到第三篇，主题也正好是从向量加法切到NCU性能分析知乎。"弹幕教我写代码"的NCU系列每期稳定上百收藏哔哩哔哩。
教程扎堆只有一个解释：大量学CUDA的人卡在了同一个地方。
入门教程会教你 thread、block、grid，教你写一个正确的向量加法和矩阵乘法。但没人教你下一步——你写的 kernel 只有理论峰值的零头，慢在哪？怎么查？网上的优化技巧文章倒是很多，共享内存、合并访存、双缓冲，每个都说得头头是道，但落到你自己的 kernel 上，第一个问题永远是：我这到底是缺计算还是缺带宽？
这篇就把这周社区里被反复验证的排查路线整理出来：3步定位瓶颈，外加4个让过来人栽了跟头的坑。适合已经能写出正确kernel、正在被性能折磨的中阶同学。
第一步：先搞清楚NCU是什么、不测什么
NCU 是 NVIDIA 官方的 kernel 级性能分析器，随 CUDA Toolkit 完整安装自带，命令行工具叫 `ncu`，图形界面叫 `ncu-ui`。从 CUDA 12 开始，老的 nvprof 已经退役，kernel 级分析的主力就是它。
两个边界先划清楚，能少走弯路：
NCU 只测单个 kernel 内部，不测 kernel 启动开销、不测 H2D/D2H 拷贝、不测多 kernel 之间的空隙。端到端时间线是 Nsight Systems（nsys）的活。程序整体慢，先用 nsys 看时间线，确认瓶颈真在某个 kernel 里，再对它开 NCU。
Windows 用户不用纠结环境。有同学在评论区问是不是必须装 Linux——不用，NCU 在 Windows 上有原生版本可以直接装，WSL2 里也能用哔哩哔哩。
最简命令就一行：
```
ncu ./your_app
```
这会跑一遍基于规则的默认分析。要看全量指标，用 `ncu --set full ./your_app`——注意 full 档要采集几百个计数器，kernel 会被反复重放很多轮，慢是正常的，后面会讲这个机制带来的坑。
第二步：看SOL面板，先给瓶颈定性
打开报告，第一个要看的是 GPU Speed Of Light（SOL） 面板，它就两根条：
Compute (SM) Throughput：计算单元忙到什么程度
Memory Throughput：存储子系统忙到什么程度
两根条的相对高度，直接决定你后面往哪个方向挖。知乎上 Leezhou 用向量加法做了个标准示范：跑出来 DRAM Throughput 82.95%，Memory Throughput 82.95%，而 Compute (SM) Throughput 只有 6.16%，L1/TEX 只有 7.92%——计算单元基本在干等数据，这是教科书级的 memory-bound（访存瓶颈）知乎。这也符合向量加法的本质：每个元素只做1次加法，却要读2次写1次，算术强度太低，带宽先被吃满。
定性规则可以总结成一张表：
SOL表现 | 瓶颈类型 | 下一步看哪 |
|---|---|---|
Memory高、SM低 | 访存瓶颈 | Memory Workload Analysis |
SM高、Memory低 | 计算瓶颈 | Compute Workload Analysis / Instruction Statistics |
两个都不高 | 延迟瓶颈（最难啃） | Warp State Statistics + Occupancy |
前两种好办，第三种最磨人：两个都不高，说明 GPU 既没算满也没搬满，时间都花在"等"上——等数据到达、等同步、等调度。这时候要换工具面板，下面说。
第三步：三个分支，往下钻一层
分支一：访存瓶颈 → 看合并程度
打开 Memory Workload Analysis，核心看一件事：你的访存合并得好不好。官方文档里这张 Memory Chart 把存储层级的流量画得很直观：每一级的命中率、吞吐占峰值的百分比一目了然，哪一层在流血一眼就能看到NVIDIA官方文档。
这里需要对齐三个概念，也是 noinfra 那期《内存分析详解》讲透的内容哔哩哔哩。
Request：一个 warp 执行一条访存指令发出的一次请求
Sector：显存系统处理数据的最小粒度，32字节
Wavefront：实际派发到显存子系统的一批 sector
理想情况下，一个 warp 的32个线程访问连续对齐的128字节，正好合成4个 sector——这是最省的形态。Leenzhou 的向量加法实测就是完美的 1/8：按元素算"应该"有 2,097,152 次 sector 级访问，实际只有 262,144 次，和理论推导完全吻合知乎。
反过来，如果线程访问的地址零散跨行，一次请求就要扇出成大量 sector。报告里会直接标出 excessive sectors（超额扇区）。暮易在知乎连载的 FlashAttention 优化实录里有个震撼的数字：他通过调整 stmatrix 的布局加 padding，把 excessive sectors 从 600 多万直接干到 0知乎。
顺便说一句 roofline 视角：如果 SOL 显示访存瓶颈，而你的算法算术强度本来就低（比如逐元素操作），那带宽打满就是终点，别折腾了；如果算术强度不低却打不满带宽，问题大概率出在合并和缓存利用上。
分支二：计算瓶颈 → 看指令构成
打开 Compute Workload Analysis 和 Instruction Statistics，看是哪条流水线在满负荷：FMA、ALU、Tensor Core 还是 LSU。如果目标吞吐对应的管道确实打满了，说明优化方向是减少指令数或换更强的指令。
这里埋着一个真实悬案，下面坑点部分细说。
分支三：延迟瓶颈 → 看 warp 在等什么
Warp State Statistics 面板会把每个周期的 warp 状态拆开，重点看停顿原因：
Long Scoreboard：在等全局内存的数据回来——典型延迟瓶颈，考虑预取、加大每线程工作量或提高并行度
Short Scoreboard：在等共享内存/L1 相关操作——查 bank conflict
Barrier：在等 `__syncthreads()`——同步点太密或block内负载不均
Math Pipe Throttle：数学流水线排队——这其实是好消息，说明你就是计算瓶颈，回到分支二
再配合 Occupancy 面板：Achieved Occupancy 离 Theoretical Occupancy 差多远？差得远的话，看是寄存器、共享内存还是 block 尺寸把占用率卡住了。
4个坑，都是过来人用小时为单位换来的
坑一：开着 debug 模式测性能。 知乎上 wild pointer 记录了一桩完整的"性能悬案"：Softmax 加了 warp shuffle 优化，理论上寄存器路径更短，结果耗时反而暴增。他一层层查：NCU 里发现 Local Memory 有流量→怀疑函数调用压栈，加 `forceinline`→流量还在→用 `–ptxas-options=-v` 查寄存器溢出，是0→最后在 Instruction Statistics 里发现 MOV 指令比优化前多了近一倍，SASS 里 `__shfl_down_sync` 竟然变成了 CALL 调用。真相是 CUDA 编译选项里的 Generate GPU Debug Information（也就是 `-G`）没关——debug 模式下内联和 peephole 优化全被禁用。他原话：排查耗时数小时知乎。结论：测性能永远用 Release 配置，别开 `-G`。
坑二：拿 ncu 跑出来的时间当真实性能。 有同学发现，加了 ncu 之后 cudaEvent 测出的 kernel 时间从 2ms 变成 1803ms，以为程序写崩了知乎。其实这是 NCU 的工作机制：为了采集几百个硬件计数器，同一个 kernel 要被重放很多轮，每轮重放前都要把显存状态恢复一遍，cudaEvent 计的是全部重放的总开销NVIDIA官方文档。真实性能看报告里的 Duration，或者干脆在不开 ncu 时单独计时。
坑三：数据太小，带宽虚高。 还是 Leenzhou 的实验：他按实测数据反推带宽，发现算出来的值超过了 A6000 Ada 的理论峰值 960GB/s。不是超频了，是 NCU 多轮重放期间，他那点测试数据（N=1M 的 float 数组）早就整个住进了 96MB 的 L2 缓存，DRAM 流量被严重低估知乎。结论：测带宽时数据量要明显大于 L2 容量，或者用 `–cache-control` 选项留意缓存行为。
坑四：把 occupancy 当越高越好的 KPI。 这周还有个反常识视频值得看：DeeparchWorks 的《GPU核越多越快吗？占用率与寄存器压力的反转》——核心更多、利用率满格、occupancy 更高，kernel 反而可能更慢哔哩哔哩。占用率只保证"有 warp 可调度"，真正决定性能的是数据复用和流水线效率。暮易那个 FA 优化系列里也有同样的体会：从算术指令到 occupancy 挣扎的那一轮（G2），收益是所有轮次里最小的，真正的跃升来自 pipeline 重构（一轮从221干到727 TFLOPS）知乎知乎。NCU 给的规则提示也要带着怀疑看，他的原话是"ncu 的 FMA 提示是幻觉"——只要 kernel 还是 latency-bound，FMA 提示就没有意义知乎。
实操清单，存一下
场景 | 命令 |
|---|---|
快速看一眼 | `ncu ./app` |
全量指标 | `ncu --set full ./app` |
只看某个面板 | `ncu --section MemoryWorkloadAnalysis ./app` |
指定指标 | `ncu --metrics dram__bytes_read.sum ./app` |
跳过预热kernel | `ncu --launch-skip 4 --launch-count 1 ./app` |
图形界面 | `ncu-ui` |
两条补充建议：程序里有多个 kernel 时，务必用 `–launch-skip` 和 `–launch-count` 圈定要分析的那一个，否则报告会混在一起；优化是迭代的活，把每轮的 `.ncu-rep` 存下来，在界面里做 baseline 对比，比凭记忆靠谱得多。
接下来学什么
按这周的更新节奏，给个顺手的学习序列：先看 noinfra 的《以GEMM为例，学会使用Nsight Compute》把面板过一遍哔哩哔哩。知乎上 Leezhou 的系列适合对照着读文字版推导知乎。再跟"弹幕教我写代码"的系列补合并访存和共享内存哔哩哔哩。想啃进阶案例，暮易的 FlashAttention 优化七连（56→986 TFLOPS）是目前中文社区少见的完整实录，每一轮改了什么、收益多少都记得很清楚知乎。
最后一个提醒：Blackwell 这一代的工具链更新很快，CUDA 13.x 还在高频迭代，NCU 的面板和指标名也会跟着变，建议收藏官方 Profiling Guide 跟着版本看NVIDIA官方文档。路线本身——先定性、再钻一层、用证据说话——倒是不会过时的。
