CUDA 算子工程：手写 FlashAttention v2 之路

第 19 章 Nsight Compute / Systems 实战

作者杨艺韬 · 2,121 字

第 19 章 Nsight Compute / Systems 实战

"Premature optimization is the root of all evil — except in CUDA, where every cycle counts and the only sin is optimizing without a profiler." ——一句被 CUDA 团队改写过的 Knuth

19.1 两个工具的分工

NVIDIA 的 profiler 有两个：

flowchart LR
  subgraph NSYS [Nsight Systems · nsys]
    NS1[系统级 timeline]
    NS2[CPU + GPU 行为]
    NS3[Kernel launch 顺序]
    NS4[CUDA API call latency]
    NS5[NVLink / PCIe 流量]
  end
  subgraph NCU [Nsight Compute · ncu]
    NC1[Kernel 内部细节]
    NC2[每个 SM 的 metrics]
    NC3[Memory 层级 hit rate]
    NC4[Stall reason 分解]
    NC5[Source-level annotation]
  end
  NSYS -->|"发现『哪个 kernel 慢』"| NCU
  NCU -->|"发现『为什么这个 kernel 慢』"| FIX[修代码]

Nsight Systems（nsys）回答"哪个慢"：拿到一段完整的 LLM 推理 trace，看哪些 kernel 占了大头、kernel 之间的 gap 是不是有问题、是否有 host 端阻塞。

Nsight Compute（ncu）回答"为什么慢"：选定一个 kernel 后，深入到 SM 内部——algorithm level 是带宽 bound 还是算力 bound？memory level 是 L1 miss 还是 L2 miss？warp 是否在等内存？

新人常见的错误是直接用 ncu 看一切。但 ncu 的 metrics 很多很复杂，没有 nsys 提供的"宏观视角"，容易陷入"调一个不是瓶颈的 kernel"的陷阱。

19.2 Nsight Systems 实战

19.2.1 抓 trace

最简单的命令：

nsys profile \
    --output=my_trace \
    --trace=cuda,nvtx,cudnn,cublas \
    --capture-range=cudaProfilerApi \
    python my_inference.py

参数说明：

--trace：跟踪哪些层。cuda 是 CUDA API 和 kernel；nvtx 是用户自定义 marker；cudnn/cublas 是这些库的内部细节。
--capture-range=cudaProfilerApi：只在 cudaProfilerStart() / cudaProfilerStop() 之间录制，避免初始化阶段干扰。

跑完得到一个 .nsys-rep 文件，用 Nsight Systems GUI 打开看 timeline。

19.2.2 看 timeline 的关键点

打开 timeline 后，关注几条 row：

CUDA API row：CPU 上的 CUDA API 调用（cudaMemcpyAsync、cudaLaunchKernel）。每个 API call 有 latency；如果某个 API 阻塞太久（比如 cudaStreamSynchronize），可能是同步逻辑写错。
GPU Kernels row：实际跑在 SM 上的 kernel。看 kernel 之间有没有 gap。
NVLink / PCIe Throughput：多卡推理时跨卡通信带宽。如果跨卡是瓶颈，会看到 NVLink 带宽贴近 900 GB/s 上限。
NVTX ranges（如果你打了 marker）：用户标注的逻辑段，方便定位代码位置。

19.2.3 常见 timeline pattern

Pattern 1：Kernel 之间有大 gap

GPU:  [kernel1]                [kernel2]                  [kernel3]
              ←   gap 5μs   ←        ←   gap 8μs   ←

可能原因：

CPU 端有同步等待（cudaDeviceSynchronize）
Kernel launch 慢（参数太多、没有 graph capture）
上下游依赖（kernel2 需要 kernel1 的输出）

修复：把多个 kernel 用 CUDA Graph 录制，整体提交。

Pattern 2：CPU 串行调度

CPU: [launch1] [launch2] [launch3] [launch4] ...
GPU: [k1][k2][k3][k4][...]
              ↑ GPU 已经空闲, CPU 还在派发

可能原因：CPU 单线程串行调度太慢。

修复：用 multi-stream + cudaLaunchCooperativeKernel + Graph。

Pattern 3：内存拷贝阻塞

GPU:  [kernel1]            [memcpy H2D]            [kernel2]
                                ↑ 5ms 拷贝

修复：用 page-locked memory + async memcpy + 用 stream 把 memcpy 和 kernel 重叠。

19.3 Nsight Compute 实战

19.3.1 抓 metrics

最常用的命令：

ncu --set full \
    --launch-skip 5 --launch-count 1 \
    --target-processes all \
    -o my_kernel_profile \
    python my_inference.py

--set full：抓所有 metrics。可以用 --set basic 或 --set roofline 限制范围（更快）。
--launch-skip 5 --launch-count 1：跳过前 5 次 launch（warm-up），抓第 6 次。
-o：输出文件名。

运行时 ncu 会插入 instrumentation，会让 kernel 跑慢 10-100×——所以不要用 ncu 测最终性能，只用来诊断。

19.3.2 ncu 的关键指标

打开 .ncu-rep 文件，最重要的几个指标：

A. SM utilization

SM Active                 : 96.3%   ← 几乎所有 SM 都在算
SM Issue Active           : 78.5%   ← 78.5% 周期发出指令
SM Issued Slot Utilization: 65.2%   ← 65.2% issue slot 在用

这三个数字反映 SM 的"忙碌程度"。如果 SM Active 低（< 50%），说明某些 SM 闲置——可能是 grid 太小或负载不均。

B. Compute / Memory Throughput

Compute (SM) Throughput   : 72.1%
Memory Throughput          : 35.4%

Compute Throughput 高 + Memory Throughput 低 → 算力 bound（已经把 Tensor Core 喂饱）。
Compute Throughput 低 + Memory Throughput 高 → 带宽 bound（HBM 是瓶颈）。
两个都高 → 接近最优。
两个都低 → kernel 有其他瓶颈（latency hiding 不够、寄存器 spill、warp stall）。

C. Roofline Plot

ncu 直接给出 Roofline 图，自动标出当前 kernel 的位置。看到一个点在带宽屋顶下方很远，说明带宽利用率低；在算力屋顶下方，说明算力利用率低。

D. Stall Reason

ncu 把每个 warp 在执行时 stall 的原因分解成几类：

Warp Stall Reasons:
  Stall Long Scoreboard       : 23.5%   ← 等内存
  Stall MIO Throttle          :  8.2%   ← MIO 队列满
  Stall Wait                  : 12.1%   ← 等指令依赖
  Stall Tex Throttle          :  3.4%
  Stall Math Pipe Throttle    :  5.1%
  ...

Stall Long Scoreboard 高 → 内存延迟没掩盖好，需要更多 warp（提高 occupancy）或更深 pipeline。 Stall Wait 高 → 指令依赖链长，需要重排指令或减少串行依赖。 Stall Math Pipe Throttle 高 → ALU 队列满，可能是 Tensor Core 用不上需要降低 issue rate。

E. Memory Workload Analysis

L1/TEX Hit Rate              : 92.1%
L2 Hit Rate                  : 65.3%
DRAM Throughput              : 78.5%

L1 命中率高 → 数据访问局部性好（典型 GEMM）。
L2 命中率低 → 跨 block 数据复用差。
DRAM Throughput 78.5% → 已经把 HBM 用得差不多。

F. Source-level Counters

ncu 最强大的功能之一：把 PTX 指令和 CUDA 源码对应，每行看 stall 占比、access 数量。

打开 source view，能看到类似：

// CUDA C source           |  Counter (Stall)
for (int k = 0; k < K; ++k) {     // 0.1%
    sum += A[i*K + k] * B[k*N + j]; // 23.5%   ← 这一行卡了
}                                  // 0.0%

可以精确定位到"哪一行代码导致最多 stall"。

19.3.3 一个完整案例

假设一个 GEMM kernel 实测 30% 算力，预期 80%。怎么诊断？

Step 1：跑 nsys 确认 GEMM kernel 是否真的是热点。

时间占总推理 80%? → 是热点。
时间占 5%? → 不是瓶颈，去优化别的。

Step 2：跑 ncu 看 GEMM kernel 的 metrics。

Compute Throughput = 30%, Memory Throughput = 65% → 带宽 bound。

Step 3：检查 Memory Workload。

L2 Hit Rate = 30% → L2 miss 很多。
DRAM Throughput = 88% → HBM 已经接近峰值。

结论：当前 kernel 已经把 HBM 带宽用到 88%，但因为 L2 miss 导致从 HBM 拉数据多于必要——可能是访存模式让 L2 cache 频繁驱逐。

修复方向：

调 tile 大小，让一个 SM 上的连续 tile 复用同一份 K/V data。
用 L2 Persistence API 锁定热数据。
重新排列 GEMM 的 swizzle，让 L2 cache 友好。

这种诊断流程在生产环境每天发生几十次。熟练掌握 ncu 是 CUDA 工程师的核心技能。

19.4 一组高频指标速查

常用指标速查表：

指标	意义	健康值
`sm__cycles_active.avg.pct_of_peak_sustained_elapsed`	SM 活跃率	> 90%
`sm__sass_thread_inst_executed_op_dadd_pred_on.sum`	FP64 加法数	量化算法用量
`dram__throughput.avg.pct_of_peak_sustained_elapsed`	HBM 利用率	> 85%
`lts__t_sectors_op_read_lookup_hit.sum` / `lts__t_sectors_op_read_lookup_miss.sum`	L2 hit/miss	hit rate > 70%
`smsp__sass_l1tex_data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum`	SMEM bank conflict	0 是理想
`sm__warps_active.avg.pct_of_peak_sustained_active`	Warp occupancy	看具体场景
`smsp__sass_average_branch_targets_threads_uniform.pct`	分支一致性	> 95% (避免 divergence)
`sm__inst_executed_pipe_tensor.sum.per_cycle_active`	Tensor Core 利用	接近 1 是峰值

19.5 NVTX：让 nsys timeline 更可读

默认情况下 nsys timeline 上每个 kernel 是它的真实 C++ 名字（很长很难读）。可以用 NVTX 给关键代码段打 marker：

#include <nvtx3/nvtx3.hpp>

void llm_forward(...) {
    nvtx3::scoped_range layer_range{"layer-12"};

    {
        nvtx3::scoped_range attn_range{"attention"};
        attention_kernel<<<...>>>(...);
    }
    {
        nvtx3::scoped_range ffn_range{"ffn"};
        gemm_kernel<<<...>>>(...);
        gemm_kernel<<<...>>>(...);
    }
}

跑 nsys 时加 --trace=nvtx，timeline 上会显示这些 marker——一眼看出每一层、每个子模块的耗时。

19.6 自动化 profiling：CUPTI 与 PyTorch profiler

工业级生产环境的 profiling 通常是自动化的：

CUPTI (CUDA Profiling Tools Interface)：底层 C API，自定义抓 metrics。
PyTorch Profiler：PyTorch 内置的 profiler，可以一键导出 Chrome trace 格式。
NVIDIA DCGM + Nsight Systems：服务器级监控。

常用工作流：

import torch.profiler as profiler

with profiler.profile(
    activities=[profiler.ProfilerActivity.CPU,
                profiler.ProfilerActivity.CUDA],
    schedule=profiler.schedule(wait=1, warmup=1, active=3),
    on_trace_ready=profiler.tensorboard_trace_handler('./logs'),
    record_shapes=True,
    with_stack=True,
) as prof:
    for step, batch in enumerate(loader):
        outputs = model(batch)
        prof.step()
        if step >= 5:
            break

输出可以在 TensorBoard 里看，或者导出 Chrome trace 用 Perfetto 看。

19.7 这一章的小结与下一章

性能调优的工具链是日常必备：

nsys 看宏观，ncu 看微观：先 nsys 定位热点 kernel，再 ncu 深入分析。
关键指标三件套：Compute Throughput、Memory Throughput、Stall Reason。
Source view 是诊断神器：把 stall 精确归到代码行。
NVTX marker 让 timeline 可读：花几分钟标记关键代码段，省几小时盯着名字猜。
生产环境用 PyTorch Profiler / DCGM：自动化抓 trace，不用每次手动跑。

第 20 章我们继续往低层走——讲 PTX 与 SASS。当 ncu 已经告诉你"这一行慢"，下一步可能需要看编译出的 PTX 中间代码、甚至 SASS 机器码，才能搞清楚"编译器到底干了什么"。读完第 20 章，读者掌握读 PTX/SASS 的能力，能诊断那些 ncu 都看不出来的"指令级"问题。

本章动手练习：

用 nsys 抓一段 LLM 推理 trace，找出占时最长的 5 个 kernel。

用 ncu 分析其中一个 kernel，记录 SM Active、Compute / Memory Throughput、Top 3 Stall Reason。

给你的 kernel 加 NVTX marker（attention、FFN、layernorm 各一个），看 timeline 区分。