CUDA Micro-benchmark 实战（上）：从源码、SASS 到 Nsight Compute 证据链

2026-04-26 约 3629 字预计阅读 8 分钟

GPU 性能分析需要一条完整证据链：

1
2
3
4
5
6
7


CUDA source
    ↓ compiler
PTX / SASS
    ↓ Warp Scheduler
issued instructions
    ↓ Nsight Compute
metrics and source correlation

如果只看源码，我们不知道编译器最终生成了什么；只看 stall reason，又容易把正常等待误判成瓶颈。本篇从编译产物出发，建立 Warp 调度模型，再给出一套 Nsight Compute 的 top-down 分析流程。下篇则在这条证据链之上设计具体 micro-benchmark。

1. CUDA 源码如何变成 GPU 指令？

.cu 文件同时包含 Host 和 Device 代码。nvcc 是编译驱动：它拆分两部分代码，把 Host 部分交给平台 C++ 编译器，把 Device 部分编译成 PTX 或特定架构的 cubin，最后把 Device image 打包进 fat binary 并参与 Host 链接。

NVCC 与 PTX 编译流程 — NVCC 编译流程：Host/Device 拆分、PTX、ptxas、cubin、fatbinary 与链接

这张图信息很多，可以分三条路径读：

Host 路径：预处理后的 Host C++ 交给系统编译器，生成 .o/.obj；
Device 路径：Device code 先形成虚拟 ISA PTX，再由 ptxas 生成面向 sm_xy 的 cubin；
打包与链接：fatbinary 可包含一个或多个 cubin，也可保留 PTX 供 Driver 在没有匹配 cubin 时 JIT。

PTX 是虚拟指令集，适合观察编译器的中间表达；SASS 是目标 GPU 实际执行的机器指令。性能结论最终必须落到 SASS。

1.1 正确生成和提取 PTX/SASS

生成独立 PTX：

1

nvcc --ptx -arch=compute_90 kernel.cu -o kernel.ptx

生成面向目标 GPU 的可执行文件并保留中间产物：

1

nvcc -O3 -lineinfo -arch=sm_90 --keep kernel.cu -o kernel

从可执行文件或库中提取：

1
2


cuobjdump --dump-ptx  kernel > kernel.ptx.txt
cuobjdump --dump-sass kernel > kernel.sass.txt

对独立 cubin 做更详细的控制流和机器码分析：

1

nvdisasm -g kernel.cubin

cuobjdump and nvdisasm — cuobjdump 适合检查容器中的 Device image；nvdisasm 适合深入分析 cubin/SASS

其他常用工具：

cu++filt：还原被 C++ name mangling 修饰的符号；
nvprune：裁剪 object/library 中不需要的 GPU 架构；
NVBit：NVLabs 的独立 SASS 二进制插桩框架，不是 CUDA Toolkit 内置命令。

1.2 Micro-benchmark 不应默认关闭优化

性能 benchmark 应尽量测量生产环境会执行的优化代码。完全关闭优化会改变寄存器分配、指令选择、循环结构和访存行为，测到的已不是目标 kernel。

防止 dead-code elimination 的正确方法是：

让结果写回可观察内存；
在 Host 端验证结果；
必要时用 asm volatile 固定短指令序列；
检查 PTX/SASS，确认目标指令和循环确实存在。

-O3 主要控制 Host 优化；Device optimization 在未使用 -G 时默认开启。Source-level profiling 应使用 -lineinfo，不要用 -G 代替：-G 会显著改变 Device code。

2. 从 Launch 到 Resident Warps

Kernel launch 后，Thread Blocks 被分配到 SM。每 32 个线程组成一个 Warp。Block 能否同时驻留，受以下资源共同限制：

每 SM 最大 blocks/warps/threads；
每 block 线程数；
每线程寄存器数；
每 block Shared Memory；
架构限制。

由 launch 配置和资源使用推导出的上限叫 Theoretical Occupancy；真正执行时驻留的 Warp 是 Active Warps。

Warp Stall 基本概念 — Device limit、theoretical occupancy、active/eligible/selected warps 的层级关系

图中的 Fermi/Kepler 数字属于历史架构示例，不应当作现代 GPU 固定上限。应关注左侧的包含关系：

1
2
3
4


Device limit
  ⊇ theoretical occupancy
      ⊇ active warps
          = eligible + stalled

Occupancy 只表示潜在并发量，不直接等于性能。若一条计算 pipeline 已经饱和，继续增加 Active Warps 不会提升吞吐；若 Warp 大量等待长延迟依赖，更高并发才可能帮助隐藏延迟。

3. Warp Scheduler 实际在做什么？

3.1 Active、Eligible、Selected 与 Stalled

驻留在 SM 上的 Active Warp，在每个调度周期可分为：

Eligible：下一条指令的输入、执行资源和同步条件都已满足，可以发射；
Selected：Eligible Warp 中本周期被 Scheduler 选中发射的 Warp；
Stalled：下一条指令仍在等待依赖、内存、barrier 或执行资源。

Active Warp 状态转换 — Warp slot 中 stalled、eligible 和 selected 的概念示意

Selected 不是长期状态，而是“本周期实际发射”。若有多个 Eligible Warps，硬件按架构相关策略选择；不能依赖某个固定 Round-Robin 顺序。

3.2 Issue Slot 与延迟隐藏

Warp 指令发射周期示例 — 单 issue-slot 示意：有 Eligible Warp 时发射，全部 Stalled 时出现空 issue cycle

图中前三个有颜色的 cycle 分别选择一个 Eligible Warp；Cycle 3 没有绿色 Eligible Warp，因此 issue slot 空闲。性能分析的关键问题不是“有没有 stall”，而是：

Stall 是否让 Scheduler 找不到任何可发射 Warp，从而浪费 issue slot？

当一个 Warp 等待 Global Memory 或数据依赖时，它的寄存器、PC 和执行状态已经驻留在 SM，不需要 CPU/OS 式保存恢复。Scheduler 可以选择另一个 Eligible Warp，因此切换开销很低；但如果所有 Active Warps 都 Stall，延迟仍无法隐藏。

3.3 分支发散的现代模型

Warp 执行一条指令时，只有 active mask 中的 lanes 参与。分支发散会让不同 lane 路径分时执行，降低有效线程利用率。Volta 之后支持 Independent Thread Scheduling，旧的“固定 branch reconvergence stack”只是教学简化，不能完整描述现代实现。

因此 Warp 内通信应使用 __syncwarp(mask)、Cooperative Groups 等显式同步机制，不要假设分支后的线程会自动在某个源码位置重新同步。

Warp 执行基本步骤 — 从 Thread grouping、fetch/decode 到 issue 与 divergence 的概念流程

4. Nsight Compute：从高层瓶颈逐步下钻

Nsight Compute 是 Kernel 级 profiler，但一次采集所有指标可能需要多次 replay，显著增加开销。建议先收集高层 section，再按问题追加指标。

4.1 稳定采集与 Source Mapping

编译优化版本并保留行号：

1

nvcc -O3 -lineinfo kernel.cu -o benchmark

采集指定 Kernel launch：

1
2
3
4
5
6
7
8
9


ncu \
  --launch-skip 1 \
  --launch-count 3 \
  --section SpeedOfLight \
  --section Occupancy \
  --section SchedulerStats \
  --section WarpStateStats \
  -o report \
  ./benchmark

只有在需要全面指标时再使用 --set full。由于 metric replay 可能多次执行 Kernel，被测程序必须是可重复的，并避免其他进程共享 GPU。

NCU SSH 远程连接 — 远程运行 ncu 并用 GUI 打开 `.ncu-rep` 报告

-lineinfo 让 NCU 将 SASS 映射回 CUDA C++/PTX，而不关闭优化：

NCU 获得代码热点 — Source 页面把 CUDA C++、PTX、SASS 与采样指标关联起来

5. 一套可复用的 Top-down 分析顺序

5.1 第一步：确认时间与 Launch

先检查：

Kernel duration 是否稳定；
grid 是否足够大，能否覆盖全部 SM；
是否存在 wave quantization、tail effect；
Block size、register、Shared Memory 是否限制并发。

如果 workload 太小，后面的 cache hit rate 和 stall 统计都可能缺乏代表性。

5.2 第二步：判断 Compute 还是 Memory

查看 Speed of Light / Roofline：

DRAM/L2 吞吐接近峰值，通常优先调查内存流量和访问模式；
某计算 pipeline 接近峰值，可能已经 compute-bound；
两者都低，则更可能是延迟、依赖、同步、指令供给或 workload 不足。

不要使用旧架构的静态吞吐表替代当前 GPU 实测。原稿的 pipe_utilization.png 只覆盖 Compute Capability 2.0–3.5，因此不再作为现代 GPU 结论。

5.3 第三步：看 Scheduler 是否真的缺指令

关注每个 Scheduler 的：

Active Warps；
Eligible Warps；
Issued Warps / issue active；
skipped issue slots。

若 Scheduler 基本每周期都能发射，Stall 百分比再高也未必是性能问题。若 issue slots 经常空闲，才继续追查为什么没有 Eligible Warp。

Warp Stall Statistics 概览 — Warp State Statistics：每条 issued instruction 对应的 selected、not selected 与 stall cycles

这张图的横轴不是“时间占比”，而是 Warp Cycles per Issued Instruction。Not Selected 表示 Warp 已 Eligible，但本周期其他 Warp 被选中；当 issue slot 已经繁忙时，它通常不是坏事。

5.4 第四步：将 Stall Reason 定位到指令

Stall reason 是症状，不是自动优化建议。名称和细节还会随架构、NCU 版本变化，应以报告内 tooltip 和当前 Nsight Compute Profiling Guide 为准。

常见类别可这样理解：

Long Scoreboard：等待与 L1TEX 路径相关的长延迟依赖，例如 Global/Local/Texture load；
Short Scoreboard：等待较短的 scoreboard 依赖，来源可能包括 MIO、Shared/Constant 路径，具体看架构和生产指令；
LG Throttle：Local/Global 指令相关队列无法继续接收，需结合 load/store 指令频率、spilling 和访问模式分析；
MIO Throttle：MIO 相关指令队列繁忙，可能涉及 Shared Memory、特殊函数或其他 MIO 路径；
Math Pipe Throttle：目标数学 pipeline 繁忙；只有结合高 pipeline utilization 才能判断是否已接近 compute ceiling；
Barrier：Warp 在执行 barrier 等待其他 Warp，常与 block 内 workload 不均衡有关；
Membar：等待 memory barrier 及相关未完成内存操作，不等同于 __syncthreads()；
Branch Resolving：等待 branch target / program counter 更新，不等同于所有 branch divergence；
No Instruction：前端暂时没有可发射指令，可能涉及 instruction fetch/cache 或控制流。

随后进入 Source 页面，找到贡献最大的 SASS 指令，再反查对应 CUDA/PTX。不要看到 long_scoreboard 就直接加 Shared Memory；如果数据只使用一次，额外搬运和同步可能更慢。

6. 三种典型诊断路径

6.1 Issue 低 + Long Scoreboard 高

检查：

Source/SASS 中是哪条 load 的 consumer 在等待；
L1/L2 hit rate 与 DRAM sectors；
访问是否合并、是否存在 pointer chasing；
是否有足够独立工作隐藏 latency；
预取或 Shared Memory staging 是否能复用数据。

6.2 Issue 低 + Barrier 高

检查 barrier 前各 Warp 的工作量是否均衡、Block 是否过大、是否能减少同步次数或改成更细粒度 Cooperative Groups。单纯提高 occupancy 无法解决同一个 Block 内的 workload imbalance。

6.3 Issue 高 + Compute Pipeline 高

这可能已经接近目标 pipeline 的吞吐上限。此时 Stall 不一定需要“修复”；应比较实测 FLOPS、指令吞吐和理论上限，确认 benchmark 是否达到了设计目标。

7. 从上篇到下篇

本篇建立了完整证据链：

1
2
3
4
5
6


Source
  → PTX/SASS
  → Warp readiness and issue
  → NCU high-level bottleneck
  → Stall reason
  → Source/SASS instruction

下一篇文章将在这条证据链上进一步讨论如何测量 ALU throughput、DRAM bandwidth，以及如何用 Inline PTX、WGMMA 与 TMA 构造 Hopper 异步流水线。

目录