阶段 5 · 工程化与 AI13 / 15约 24 分钟
性能分析：用数据代替猜测
Nsight Compute 与 Nsight Systems 的分工、必看指标、roofline 判读
学完这一课你会
- 分清 nsys 和 ncu 各自解决什么问题
- 记住一次 kernel 分析必看的五个指标
- 会读 roofline 图并据此决定优化方向
- 用 compute-sanitizer 定位越界和竞争
CUDA 优化最常见的失败模式是优化了不是瓶颈的地方。你花两天把一个 kernel 的指令数砍掉 30%，结果整体没变——因为它本来就卡在显存带宽上。所以规矩只有一条：先测，再改，再测。NVIDIA 的工具链分工很清楚，用对了能省下大量时间。
| 工具 | 看什么 | 什么时候用 |
|---|---|---|
Nsight Systems（nsys） | 整个应用的时间线：kernel、拷贝、CPU、API 调用 | 第一步。定位是谁慢、有没有空隙、重叠有没有生效 |
Nsight Compute（ncu） | 单个 kernel 的硬件计数器细节 | 第二步。已知某 kernel 是瓶颈，要挖它内部的问题 |
compute-sanitizer | 越界、未初始化、竞争、泄漏 | 正确性可疑时，或结果不稳定时 |
nvidia-smi dmon | 利用率、功耗、温度、时钟 | 怀疑降频或多进程争抢时 |
第一步：nsys 看全局
shell
1nsys profile -o report --stats=true --force-overwrite=true ./app2 3# 命令行直接看统计摘要4nsys stats report.nsys-rep5 6# 图形界面看时间线（能看到流之间的重叠情况）7nsys-ui report.nsys-rep- ▸kernel 之间有大片空隙 → CPU 侧成了瓶颈，考虑 CUDA Graph 或减少同步。
- ▸拷贝和计算完全不重叠 → 检查是不是用了默认流，或 host 内存没 pin。
- ▸某个 kernel 占了 80% 时间 → 目标锁定，进入第二步。
- ▸大量细碎的小 kernel → 考虑算子融合。
第二步：ncu 挖细节
shell
1# 完整分析（较慢，会 replay kernel 多次）2ncu --set full -o profile ./app3 4# 只分析指定 kernel 的前 3 次调用5ncu --kernel-name matmulTiled --launch-count 3 ./app6 7# 只取关心的指标，速度快很多8ncu --metrics \9 sm__throughput.avg.pct_of_peak_sustained_elapsed,\10 gpu__dram_throughput.avg.pct_of_peak_sustained_elapsed,\11 sm__warps_active.avg.pct_of_peak_sustained_active,\12 l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,\13 l1tex__data_bank_conflicts_pipe_lsu_mem_shared.sum \14 ./app15 16# 直接让工具给出优化建议17ncu --set full --section SpeedOfLight --section Occupancy ./app| 指标 | 含义 | 怎么判读 |
|---|---|---|
| Compute Throughput | SM 算力利用率 | > 60% 说明算力受限，该优化指令和算法 |
| Memory Throughput | 显存带宽利用率 | > 70% 说明带宽已近打满，只能减少数据搬运 |
| Achieved Occupancy | 实际驻留 warp 比例 | 远低于理论值 → 尾部效应或负载不均 |
| sectors / requests | 平均每次请求的 sector 数 | 理想 4，接近 32 说明访存完全没合并 |
| Bank Conflicts | 共享内存冲突次数 | 非零就值得看看能不能 padding 掉 |
| Warp Stall Reasons | warp 卡在哪 | 见下表，这是最有信息量的一项 |
| Stall 原因 | 含义 | 对策 |
|---|---|---|
Long Scoreboard | 等全局内存返回 | 提高占用率、改善合并、增加 ILP |
Short Scoreboard | 等共享内存返回 | 消除 bank conflict |
Barrier | 卡在 __syncthreads() | 减少屏障、平衡 block 内负载 |
MIO Throttle | 访存指令队列满 | 用向量化访存减少指令条数 |
Math Pipe Throttle | 算力单元排队 | 好现象，说明确实在算 |
Not Selected | 有别的 warp 被优先调度 | 好现象，说明并行度充足 |
Roofline：把 kernel 放进坐标系
性能 (GFLOP/s)
▲
19500┤ ┌────────────────── 峰值算力（算力屋顶）
│ ╱
│ ╱ ← 斜线区：受带宽限制
│ ╱ 性能 = 带宽 × 计算强度
│ ╱
│ ● ╱ ● tiled matmul（还有上升空间）
│ ╱
│ ● ╱ ● 向量加（贴在斜线上 = 已打满带宽）
└──────┴──────────────────────────▶ 计算强度 (FLOP/Byte)
12.5
拐点 = 峰值算力 / 峰值带宽
点在斜线上 → 带宽已打满，只能提高计算强度（融合、tiling、复用）
点在斜线下 → 有优化空间：访存没合并 / 占用率不足 / 有冲突
点在平台上 → 算力已打满，考虑 Tensor Core 或换算法shell
1ncu --set roofline -o roofline_report ./app2ncu-ui roofline_report.ncu-rep # 图形界面里直接能看到 kernel 落点compute-sanitizer：CUDA 版的 valgrind
shell
1# 越界访问、非法地址（最常用）2compute-sanitizer --tool memcheck ./app3 4# 共享内存 / 全局内存的数据竞争5compute-sanitizer --tool racecheck ./app6 7# 读取未初始化的内存8compute-sanitizer --tool initcheck ./app9 10# 显存泄漏11compute-sanitizer --tool synccheck ./app12 13# 想看到出错的源码行，编译时加 -lineinfo（比 -G 轻量，不影响优化）14nvcc -O3 -lineinfo -arch=sm_80 -o app main.cu一个可复用的优化循环
- 1.建立 baseline：用 cudaEvent 测出稳定的时间和有效带宽，写进脚本。
- 2.算理论上限：数据量 / 显存带宽 = 理论最短时间。差距多大，一目了然。
- 3.nsys 定位：哪个 kernel 占大头？有没有空隙和串行化？
- 4.ncu 诊断：看 SOL 定性，看 stall 原因定位，看 sectors/requests 查合并。
- 5.只改一处：一次只动一个变量，改完立刻重测。同时改三处的话，你永远不知道是哪一处起了作用（或抵消了效果）。
- 6.验证正确性：每轮优化后跑一次结果对拍 + compute-sanitizer。快而错的 kernel 毫无价值。
自测
先自己在心里回答一遍，再展开对照。答不上来的说明这一段值得重读。