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

第 19 章 Nsight Compute / Systems 实战

作者 杨艺韬 · 3,109 字 · 发布于 · 更新于

前四篇的每一个优化，都需要一个能证明它有效的读数。 这一章讲 Nsight Systems 与 Nsight Compute 的分工，以及一份 kernel 报告该按什么顺序读下去。

19.1 两个工具的分工

NVIDIA 面向 CUDA 的主力 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 是这些库的 API 调用。
  • --capture-range=cudaProfilerApi：只在 cudaProfilerStart() / cudaProfilerStop() 之间录制，避免初始化阶段干扰。程序里必须真的调用这对 API（PyTorch 里是 torch.cuda.profiler.start() / stop()），否则什么也录不到。

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

19.2.2 看 timeline 的关键点

打开 timeline 后，关注几条 row：

  1. CUDA API row：CPU 上的 CUDA API 调用（cudaMemcpyAsync、cudaLaunchKernel）。每个 API call 有 latency；如果某个 API 阻塞太久（比如 cudaStreamSynchronize），可能是同步逻辑写错。
  2. GPU Kernels row：实际跑在 SM 上的 kernel。看 kernel 之间有没有 gap。
  3. NVLink / PCIe Throughput：多卡推理时跨卡通信带宽。这几行来自 GPU Metrics 采样，默认不采（--gpu-metrics-devices 默认为 none），要在上面的命令里加 --gpu-metrics-devices=all 才会出现。注意 H100 那个 900 GB/s 是双向聚合口径（18 条链路 × 25 GB/s 单向 = 450 GB/s 单向），all-reduce 这种双向都跑满的负载才可能贴近 900；单向流量看到 450 就已经是打满了，别拿 900 去对。
  4. 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）
  • 上下游依赖经过 host 中转（例如 kernel2 的启动参数要等 kernel1 的结果拷回 CPU 才能算出来）；同一 stream 里的普通依赖由 GPU 按序执行，本身不产生 gap

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

Pattern 2：CPU 串行调度

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

可能原因：CPU 单线程串行调度太慢，GPU 把队列里的活干完了还在等下一批。

修复：减少要派发的次数（kernel fusion、CUDA Graph 把整段 forward 一次提交），或者让派发本身并行（多 stream + 多个 host 线程各自派发）。注意 cudaLaunchCooperativeKernel 解决的是 grid 级同步，不是派发速度，别拿它治这个病。

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 \
    -k regex:gemm \
    --launch-skip 5 --launch-count 1 \
    --target-processes all \
    -o my_kernel_profile \
    python my_inference.py
  • --set full：抓完整的 section 集合。嫌慢可以换成更小的集合——用 ncu --list-sets 看当前版本支持哪些，别照抄别人博客里的集合名，Nsight Compute 各版本之间改过：不加 --set 时采集的默认集合在 2023.1 及以前叫 default，2023.2 起改名 basic；roofline 集合是 2021.3 才加的。
  • -k regex:gemm：只选名字匹配 gemm 的 kernel。
  • --launch-skip 5 --launch-count 1：在匹配 -k 的 launch 里跳过前 5 次（warm-up），抓第 6 次。不加 -k 时计数的是进程里所有 kernel 的 launch，前 5 次未必是你关心的那个 kernel。
  • -o：输出文件名。

运行时 ncu 需要把同一个 kernel 重放（replay）多次才能凑齐一整套 metric，还会串行化所有 kernel launch，整个程序的墙钟时间会比正常执行慢很多（metric 越多、replay 次数越多越慢）。报告里的 Duration 也不等于生产环境的耗时：ncu 默认锁定 GPU 时钟（--clock-control），并在每次 replay 前清空各级缓存（--cache-control all），kernel 之间也不再并发。所以 ncu 报告里的耗时不能当性能数据用，它只用来诊断。要测真实延迟，请用 nsys 或自己插 cudaEvent。

19.3.2 ncu 的关键指标

打开 .ncu-rep 文件，最重要的几个指标（下面各段的读数都是示意，不是实测）：

A. SM utilization

sm__cycles_active.avg.pct_of_peak_sustained_elapsed : 96.3%   ← SM 平均 96.3% 的周期处于活跃
Issue Slots Busy                                    : 78.5%   ← 78.5% 的 issue slot 发出了指令

第一行是 metric 名（SM 活跃率），第二行是 Compute Workload Analysis 里的显示名，两者反映 SM 的"忙碌程度"。如果 SM 活跃率很低（比如不到一半，这是经验阈值），说明某些 SM 闲置——可能是 grid 太小或负载不均。

B. Compute / Memory Throughput

Compute (SM) Throughput   : 72.1%
Memory Throughput          : 35.4%
  • Compute Throughput 高 + Memory Throughput 低 → 算力 bound（某条计算流水接近饱和，未必是 Tensor Core，要展开 breakdown 看是哪条）。
  • Compute Throughput 低 + Memory Throughput 高 → 访存 bound（某一级存储是瓶颈，是 HBM 还是 L1/L2/shared memory，要看 breakdown 和 Memory Workload）。
  • 两个都高 → 接近最优。
  • 两个都低 → kernel 有其他瓶颈（latency hiding 不够、寄存器 spill、warp stall）。

C. Roofline Plot

ncu 直接给出 Roofline 图（需要采集 roofline 相关 section，--set full 已包含），自动标出当前 kernel 的位置。看到一个点在带宽屋顶下方很远，说明带宽利用率低；在算力屋顶下方，说明算力利用率低。

D. Stall Reason

ncu 把每个 warp 在执行时 stall 的原因分解成几类（Warp State Statistics / Source Counters 里的名字）：

Warp Stall Reasons:
  Stall Long Scoreboard       : 23.5%   ← 等 L1TEX 访存（global/local 等）的结果
  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 高 → 固定延迟的依赖链长（例如 FFMA 接 FFMA），需要重排指令、展开循环或增加 warp；官方文档提醒它通常只在已经高度优化的 kernel 里才排到前面。

Stall Math Pipe Throttle 高 → 数学流水已经被喂饱、发不进去了。对 GEMM/FA 这类 compute-bound kernel 来说这恰恰是好消息（说明瓶颈真的在算力）；但如果这个 kernel 本该是带宽 bound 的，那多半是某类低吞吐的运算（__expf 这类走 MUFU 的指令、会展开成一长串指令的整数除法、FP64）被误用在了热路径上。

另外，只有当调度器不是每个周期都能发出指令时，stall 分解才值得细看（官方 Source Counters 说明原话：Only focus on stalls if the schedulers fail to issue every cycle）。

E. Memory Workload Analysis

L1/TEX Hit Rate              : 92.1%
L2 Hit Rate                  : 65.3%
DRAM Throughput              : 78.5%   ← 这一项在 Speed of Light 表里
  • L1 命中率 高 → 数据访问局部性好。GEMM 这类靠 shared memory 显式复用的 kernel，这个数参考意义不大。
  • L2 命中率 低 → 跨 block 数据复用差。
  • DRAM Throughput 78.5% → 已经把 HBM 用得差不多。

F. Source-level Counters

ncu 最强大的功能之一：把 SASS（以及 PTX）指令和 CUDA 源码对应，每行看 stall 占比、access 数量。前提是编译时加 -lineinfo（例如 nvcc -arch=sm_90a -O3 -lineinfo），否则 Source 页只能看 SASS，对不上源码行。

打开 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 = 88% → 带宽 bound。（注意 ncu 的 Memory Throughput 是整个访存子系统各级利用率的最大值，所以它不会低于下一步看到的 DRAM Throughput。）

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 复用同一份 A/B 分块。
  • 用 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 加法指令数 在 LLM kernel 里应该是 0——H100 SXM5 的 FP64（非 Tensor Core）峰值只有 FP32 的一半（34 vs 67 TFLOPs），误用 double 字面量就会悄悄掉进来
dram__throughput.avg.pct_of_peak_sustained_elapsed HBM 利用率 > 85%
lts__t_sector_op_read_hit_rate.pct L2 读命中率（要看 miss 的扇区数用 lts__t_sectors_op_read_lookup_miss.sum） > 70%
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum SMEM 读的 bank conflict（写用 ..._op_st.sum） 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 发射密度（Hopper 上一条 WGMMA 按 4 条 warp 指令计） 与同尺寸 cuBLAS 对比比看绝对值有意义；看 Tensor Core 忙不忙，更直接的是 sm__pipe_tensor_cycles_active

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_trace_handler 实际写出的是 Chrome trace 格式的 *.pt.trace.json，可以直接用 Perfetto 打开；在 TensorBoard 里看要另装 torch-tb-profiler 插件，该插件已被 PyTorch 标为弃用、不再维护。

19.7 这一章的小结与下一章

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

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

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

本章动手练习：

  1. 用 nsys 抓一段 LLM 推理 trace，找出占时最长的 5 个 kernel。
  2. 用 ncu 分析其中一个 kernel，记录 SM 活跃率、Compute / Memory Throughput、Top 3 Stall Reason。
  3. 给你的 kernel 加 NVTX marker（attention、FFN、layernorm 各一个），看 timeline 区分。