Nsight 调优工作流：Compute 与 Systems 怎么读

Roofline 篇 给了优化方向，但要把瓶颈定位到具体的 kernel、具体的指令或访存行为，得靠 profiler。NVIDIA 的两个工具分工明确：Nsight Systems 看全局时间线，Nsight Compute 看单个 kernel 的细节。这一篇讲它们各管什么、关键指标怎么读，以及没有完整工具链时怎么用 CUDA event 做轻量测量。

一、两个工具的分工

正确顺序是先 Systems 后 Compute：先用 Systems 找出真正占时间的 kernel（优化一个只占总时间 2% 的 kernel 没意义），再用 Compute 深挖那个 kernel。这对应一条通用原则——先定位热点，再优化热点。

二、Nsight Systems：时间线视角

nsys 抓取程序运行的时间线，把 CUDA API 调用、kernel 执行、内存拷贝、CPU 线程活动画在同一条时间轴上。典型用法：

nsys profile -o report ./my_app

时间线上要找的几类问题：

• kernel 之间的空隙：GPU 在等 CPU 提交下一个 kernel，说明 CPU 侧有瓶颈（如 Python 开销、同步过多），GPU 没喂饱。
• H2D/D2H 拷贝占比过大：数据在 host 和 device 之间反复搬，应该让数据留在 device（第 04 篇 提过）。
• 串行执行本可并发的操作：拷贝和计算没重叠、多个独立 kernel 没用多 stream，可以用异步拷贝 + stream 重叠（通信重叠篇）。
• 同步点：频繁的 cudaDeviceSynchronize 会制造气泡。

Systems 的产出是一份”时间预算表”：告诉你接下来该把精力放在哪个 kernel 上。

三、Nsight Compute：单 kernel 深挖

锁定热点 kernel 后，用 ncu 对它做详细分析：

ncu --set full -o kernel_report ./my_app
# 或只看某个 kernel、某些指标
ncu -k my_kernel --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed,\
dram__throughput.avg.pct_of_peak_sustained_elapsed ./my_app

Compute 会重放 kernel 多次采集硬件计数器。要看的核心指标分几组：

1. 两个吞吐率：先判断算力还是访存

• sm__throughput.avg.pct_of_peak_sustained_elapsed：SM（计算）吞吐占峰值百分比。
• dram__throughput.avg.pct_of_peak_sustained_elapsed：DRAM（访存）吞吐占峰值百分比。

哪个接近 100%，瓶颈就在哪。这等价于把 kernel 定位到 Roofline 的哪条屋顶上：DRAM 吞吐打满 → memory-bound；SM 吞吐打满 → compute-bound；两个都不高 → 是延迟暴露或并行不足的问题，往下看 stall。

2. Warp stall reasons：warp 在等什么

当两个吞吐都不高，说明 SM 大量周期里没指令可发，warp 都在 stall。Compute 会归因每个周期 warp 卡在什么上，常见几种：

• Long Scoreboard：等 global/local memory 返回。占比高说明访存延迟没被隐藏——提 occupancy、改善合并、减少访存。
• MIO Throttle / LSU：访存指令排队，访存单元过载——典型是访存太密或有 bank conflict。
• Barrier：等 __syncthreads()，block 内负载不均或同步太频繁。
• Math Pipe Throttle：算术单元打满——这其实是好事，说明 compute-bound。
• Not Selected：有就绪 warp 但调度器没选它，通常意味着资源充足，没问题。

stall reason 是 Compute 最有价值的输出：它把”慢”翻译成具体原因。

3. occupancy 与资源

Compute 报告 achieved occupancy（实际驻留 warp）和 theoretical occupancy（受寄存器/shared 限制的上限），以及每线程寄存器、每 block shared memory。两者差距大说明负载不均或尾部效应；理论 occupancy 低说明被资源卡住（第 06 篇）。注意结合 stall 判断 occupancy 够不够——如果 Long Scoreboard 占比低，occupancy 即使不高也够用。

4. 访存细节

l1tex__ 和 lts__（L2）相关指标能看缓存命中率；l1tex__data_bank_conflicts_pipe_lsu_mem_shared 直接报 shared memory 的 bank conflict 数。怀疑访问模式问题时看这些。

四、一个标准排查流程

flowchart TD
  A["nsys: 找占时间最多的 kernel"] --> B["ncu: 看 SM/DRAM 吞吐"]
  B -->|"DRAM 接近峰值"| C["memory-bound:<br/>减少访存/融合/合并"]
  B -->|"SM 接近峰值"| D["compute-bound:<br/>换更高吞吐指令/Tensor Core"]
  B -->|"两者都低"| E["看 stall reasons"]
  E -->|"Long Scoreboard"| F["访存延迟暴露:<br/>提 occupancy/改善合并"]
  E -->|"Barrier"| G["同步/负载不均"]
  E -->|"bank conflict 高"| H["加 padding 改 shared 布局"]

这个流程把前面几篇的优化手段串成了可执行的决策树：profiler 给证据，优化手段对症下药。

五、没有完整工具链时的轻量替代

Nsight Compute/Systems 依赖 CUDA 工具链和相应权限，在受限环境（如部分容器、本系列使用的 WSL2 + NVRTC 环境）里不一定能跑——在 WSL2 上 ncu --set full 常因权限不足或硬件计数器不可用而直接失败，这正是本系列用 CUDA event 计时兜底的原因。这种情况下，CUDA event 计时是最小可用的测量手段，也是本系列所有性能数字的来源：

cudaEvent_t start, stop;
cudaEventCreate(&start); cudaEventCreate(&stop);
// 预热若干次后：
cudaEventRecord(start);
my_kernel<<<grid, block>>>(...);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
float ms;
cudaEventElapsedTime(&ms, start, stop);  // GPU 侧耗时，毫秒

要点：

• 必须预热：第一次启动包含 JIT 编译、上下文初始化等一次性开销，会严重偏高。本系列统一预热 10 次、取 100 次的中位数。
• event 计的是 GPU 侧时间，比 CPU 端 clock() 准确，不含 host 提交开销。
• 算有效带宽和 GFLOP/s：用耗时反推 \(\text{带宽} = \text{字节数}/\text{时间}\)、\(\text{GFLOP/s} = \text{FLOP}/\text{时间}\)，和 Roofline 的屋顶比较，就能近似判断 memory-bound / compute-bound——不需要 ncu 也能做基本定位。

event 计时给不了 stall reason、缓存命中这些细节，但”这个 kernel 多快、贴没贴到带宽/算力屋顶”这类一阶问题它完全够用。能用 ncu 时优先 ncu，受限时用 event 计时兜底。

六、小结与下一步

• 先用 Nsight Systems 找时间热点（哪个 kernel/拷贝值得优化），再用 Nsight Compute 深挖单 kernel。
• Compute 先看 SM/DRAM 吞吐定位 compute/memory-bound，两者都低时看 warp stall reasons 找根因，再看 occupancy 和 bank conflict 细节。
• 没有完整工具链时，CUDA event 计时（预热 + 取中位数 + 反推带宽/GFLOP/s）是可靠的轻量替代，也是本系列的测量口径。

工具和方法论备齐，从下一篇开始进入真正的算子实现。先从最基础的协作模式——Reduction 与 Scan 开始，它是理解 warp 级原语和 shared memory 协作的入口。