极客工场

Appearance

共 4,980 字预计阅读约 17 分钟

Nsight Systems 与 Nsight Compute：GPU 性能分析双工具工作流 ​

更新时间：2026-09-25。上一篇 GPU 显存层级与带宽 讲清了数据从 HBM 到寄存器要过哪几关，但一个真实的 CUDA 程序跑慢了，问题往往不止于「某个 kernel 访存不合并」：可能是 CPU 端根本没及时把 kernel 发射出去，可能是 H2D 拷贝占了一半时间，也可能是时间线上挤满了上千个几微秒的小 kernel。NVIDIA 提供的 Nsight Systems Compute 这对工具正是为这种分层排查设计的——前者看程序的「一生」，后者看一个 kernel 的「内脏」。本篇讲两者的分工与配合流程，瓶颈的理论判据对照 Roofline 模型：用算术强度定位 GPU 性能瓶颈，访存与占用率的深入解释分别见 GPU 访存合并与缓存行利用 和 Occupancy 与寄存器压力。

本文要回答的问题 ​

  • Nsight Systems 和 Nsight Compute 各看什么层级？为什么不能只靠其中一个下结论？
  • 在 Systems 时间线上，GPU 空转、H2D 拷贝阻塞、小 kernel 过多这三类典型问题分别长什么样？
  • Nsight Compute 的 SOL roofline、SM 吞吐、内存吞吐、warp stall 采样各回答什么问题？gld_efficiency 等指标怎么读？
  • 怎样按 warp stall 的主要原因，把一个慢 kernel 判为 compute-bound、memory-bound 还是 latency-bound，然后对症开方？
  • 命令行下 nsys profile 和 ncu 最常用的参数与指标名有哪些？无图形界面的服务器上怎么产出报告？

一、分工：一个看时间线，一个看 kernel 内部 ​

新手最常见的误区是打开 Nsight Compute（以下简称 ncu）逐个 kernel 翻指标，翻了半天却没发现 GPU 有 60% 的时间根本没在执行 kernel——单 kernel 内部的指标再好，也回答不了「GPU 为什么闲着」。Nsight Systems Compute 的分工首先是观测尺度的分工。

Nsight Systems（命令行 nsys）是系统级的时间线采样器：它以很低的开销记录 CUDA Runtime/Driver API 调用、kernel 起止、内存拷贝（H2D/D2H/D2D）、CPU 线程调度，乃至 MPI、NVTX 标注、OS 运行时事件，然后把所有轨道对齐到同一条时间轴上。它回答的是宏观问题：GPU 忙不忙、谁让它等、kernel 与拷贝如何排布。它不深入 kernel 内部，不收集 SM 级硬件计数器。

Nsight Compute（命令行 ncu）是 kernel 级的性能剖析器：它通过重放（kernel replay）多次启动同一个 kernel，每次采集一组硬件计数器，最终给出该 kernel 的 SOL（speed-of-light，光速）roofline、各子系统吞吐占峰值的百分比、warp stall 原因分布、occupancy、访存效率等。它回答的是微观问题：这个 kernel 跑不快，时间到底耗在哪个执行单元或哪次等待上。它一次只看一个 kernel，且采集开销大，绝不能对整段程序全量跑。补充一点工作环境上的现实：训练或生产服务器通常没有图形界面，标准做法是在服务器上用命令行产出 .nsys-rep / .ncu-rep 报告，再 scp 到本地工作站用 GUI 打开深入分析；报告文件自包含全部数据，不需要在本机重跑程序。因此命令行参数不是可选项，而是这套流程的入口；报告体积可能不小，长任务采集时注意用时长与过滤参数控制文件大小。

对比项Nsight Systems（nsys）Nsight Compute（ncu）
观测层级系统级全局时间线单个 kernel 内部
数据来源API 跟踪 + 时间线事件硬件计数器 + kernel 重放
典型开销个位数百分比，可全程开数倍到数十倍，只对目标 kernel
核心视图CPU 线程 / CUDA API / 拷贝 / kernel 轨道SOL roofline / Section 指标 / Source 页
关键指标GPU 活跃占比、API 间隙、拷贝时长、kernel 间隔sm__throughput、gld_efficiency、warp stall、occupancy
回答的问题时间花在谁身上、GPU 为何空转这个 kernel 内部被什么卡住
使用顺序先用，定位嫌疑 kernel后用，解剖 Systems 点名的 kernel

二、Nsight Systems：在时间线上抓三类宏观病灶 ​

典型的命令行用法是先采一份报告，再拷回本地用 GUI 打开：

bash
nsys profile -o report -t cuda,nvtx,osrt --stats=true ./your_app

-t 指定跟踪的事件域（cuda 包含 API、kernel 和拷贝，nvtx 是用户代码里手写的区间标注，osrt 是系统调用）；--stats=true 在 .nsys-rep 之外额外打印一份汇总，包括 CUDA API 耗时排名、kernel 耗时排名和 GPU 核利用率，没有 GUI 时靠它就能做第一轮判断。在长时间运行的服务上用 --duration=30 限制采集时长，用 -y 10 丢弃前 10 秒的预热段。

打开时间线后，第一类要抓的病灶是 GPU 空转（空隙）。正常的满负荷程序里，GPU 轨道上的 kernel 色块应当首尾相接；如果色块之间有大片空白，沿空白向上对齐到 CPU 线程轨道和 CUDA API 轨道，看空窗里 CPU 在做什么。最常见的发现是 cudaMemcpy 之外夹杂着大量 host 端计算、磁盘 IO 或锁等待——kernel 是 CPU 串行发射的，host 没走到 launch 这一行，GPU 就只能干等。另一高频原因是隐式同步：cudaMemcpy 默认同步、cudaDeviceSynchronize、任何紧挨着 launch 的内存分配（cudaMalloc 触发上下文级同步），都会在 API 轨道上显示为一条横跨整个 GPU 空隙的长条。修法是异步流水化：预分配显存、用 stream + event 把拷贝和计算重叠、用 NVTX 区间把业务阶段标出来便于定位。

第二类病灶是 拷贝与计算不重叠。H2D/D2H 拷贝在独立轨道上显示，对照 kernel 轨道看两者是否并行。若拷贝块与 kernel 块在时间上严格错开、且拷贝总时长不可忽略，说明程序在「先整块搬入 → 再整体计算 → 再整块搬出」的单流串行模式里。按 batch 切块、多 stream 双缓冲后，时间线上应出现 H2D(k+1)、kernel(k)、D2H(k-1) 三层重叠的阶梯形状。判断是否值得做这一步，直接看拷贝占 GPU 活跃时间之外的墙钟比例即可。

第三类病灶是 小 kernel 过多。表现为 GPU 轨道上密密麻麻的短小色块，单个 kernel 只有几微秒，汇总表里 kernel 数量成千上万、且平均时长极短。每个 kernel 都有固定的 launch 与硬件调度开销，kernel 短到与开销同量级时，实际算力被 launch 间隙吃掉。典型修法是 kernel 融合（把相邻的逐元素操作并进一个 kernel）、用 CUDA Graph 捕获整串 launch 然后一次 cudaGraphLaunch 重放、或调整问题切分粒度减少发射次数。判别融合是否可行，可以先在时间线上看相邻小 kernel 之间有没有 D2H/H2D 或 host 介入：纯 device 侧首尾相接的 kernel 串最适合融合或 Graph，一旦中间夹着 host 决策，则要先消除中间同步。

chip-design/intermediate/nsight-performance-analysis.md D2 示意图

补一个定位技巧：框架代码的时间线往往很长、阶段很多，肉眼对不上业务语义，这时在代码里插入 NVTX 区间（C++ 用 nvtxRangePush("stage-name") / nvtxRangePop()，Python 用 torch.cuda.nvtx.range_push），时间线上就会出现带名字的彩色泳道，能立刻把「数据预处理」「前向」「优化器更新」等阶段与底层 API 对齐。对于 PyTorch 类程序，还要注意时间线上的 cudaMemcpyAsync 可能来自框架的隐式拷贝与缓存分配器行为，先排除框架噪声再改业务代码。--stats=true 的汇总文本里，CUDA API 统计按累计耗时排序，若 cudaDeviceSynchronize、cudaFree、cudaMalloc 排在前列，即使不看图形界面也能直接判定存在频繁同步或分配抖动。

走完时间线这一层，产出不是结论而是一张短名单：哪几个 kernel 占了大头、它们各自可能带着什么宏观嫌疑。这一步决定了昂贵的 ncu 该花在谁身上。一个纪律是不要被「单次最慢」的 kernel 带偏——汇总表里同时看总耗时和调用次数，一个中等时长但被调用上万次的 kernel，优化收益可能远大于只跑几次的最慢 kernel。

三、Nsight Compute：SOL roofline 与各子系统吞吐 ​

对短名单上的 kernel，用 ncu 定点采集。最常用的调用方式：

bash
ncu --set full -f -o kernel_report \
    --kernel-name regex:gemm_kernel \
    --launch-skip 10 --launch-count 1 \
    ./your_app

--set full 采集全部 section（粗略摸底时用 --set basic 或只取 roofline、SpeedOfLight）；--kernel-name 用正则过滤目标 kernel；--launch-skip 跳过预热、--launch-count 只采一次，因为每次采集都要重放 kernel、开销很大。服务器上无 GUI 时加 --csv --log-file metrics.csv 导出表格，或只打印需要的指标：

bash
ncu --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed,\
dram__throughput.avg.pct_of_peak_sustained_elapsed,\
l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,\
smsp__warp_issue_stalled_long_scoreboard_per_warp_active.pct,\
sm__warps_active.avg.pct_of_peak_sustained_active \
    --launch-skip 10 --launch-count 1 ./your_app

先看的是 SOL（SpeedOfLight）页。它把计算吞吐（SM Activity / Compute SM Throughput）和内存吞吐（Memory Throughput，细分为 DRAM 和 L1/L2）统一表示成占硬件峰值的百分比，本质上就是一张表格化的 Roofline 模型体检表：两个百分比都低，说明既不是算力墙也不是带宽墙，问题在并行度或停顿；计算高、内存低，偏向 compute-bound；反之偏向 memory-bound。Roofline 页则直接画出该 kernel 的实测算术强度和标点，标点落在斜线段还是平顶段一目了然。

其次看 访存效率指标。gld_efficiency（global load efficiency，新版指标名对应 smsp__inst_executed_pipe_lsu... 与请求 sector 数的比值口径）表示全局加载请求实际用到的字节占搬运字节的比例：100% 意味着完全合并访存，低于 60% 通常说明 warp 内线程访问地址离散、大量缓存行只被用了一小部分。配套看 l1tex__t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio（每次加载平均消耗的 sector 数，理想为 1，恶化可达 8 甚至 32），这是诊断非合并访问最直接的数字，优化原理见 GPU 访存合并与缓存行利用。共享内存侧则看 shared memory bank conflict 相关指标。

然后是 occupancy 页。sm__warps_active.avg.pct_of_peak_sustained_active 给出实测活跃 warp 占峰值的比例，并展开限制因素：是每线程寄存器数过多（launch__registers_per_thread），还是每 block 共享内存过大，还是 block 内线程数配置不当。注意 occupancy 低是「病因不明时的嫌疑犯」而非自动的优化目标——延迟低、吞吐已饱和的 kernel 在低 occupancy 下照样跑满，盲目压寄存器反而可能拖慢，判别方法见 Occupancy 与寄存器压力。

读 ncu 报告还有两个页面值得一提。其一是 Source/Source Counters 页：ncu 把执行计数、各 warp stall 的采样直接标到每行 CUDA C 源码和对应 SASS 指令上，当指标指向明确时，可以顺着它找到究竟是哪一条 load 指令贡献了大部分 long scoreboard，而不必通读整个 kernel。其二是 指令混合统计（Instruction Statistics）：inst_executed 按指令类型拆分，若 Load/Store 与整数指令的条数远超浮点指令，说明大量发射槽被地址计算和搬运占据——这正是 Roofline 公式「悄悄乐观」的部分，重写索引计算、让一个线程处理多个元素往往比抠浮点本身更见效。看指标时还要区分 avg（平均值）与 max（最大值）口径：SM 间负载不均时，平均吞吐尚可但个别 SM 已近饱和，max 口径会先暴露热点，常见于 block 数不足或网格划分与问题规模不对齐的情形。

四、warp stall 采样：把等待时间归因 ​

ncu 最有诊断价值的一页是 Warp State（warp stall 原因分布）。warp 被发射后并非一直占着执行单元，大量时间处于「就绪但发射不出去」的停顿状态，硬件按原因对每个停顿周期记账。读懂几个主要 stall 原因，就拿到了对症的钥匙。

  • long scoreboard：等待全局内存加载返回，是 memory latency 的典型信号。占比高说明 warp 发出的 global load 迟迟未到，常见于非合并访问、缓存命中率低或活跃 warp 太少不足以掩盖延迟。
  • short scoreboard：等待片上快速操作（共享内存、数学单元流水）返回，常对应 bank conflict 或较短流水依赖。
  • stall barrier：等待 __syncthreads()，说明 block 内线程负载不均、有人先到栅栏空等。
  • stall wait：线程内指令间的数据依赖（如用上一条 load 的结果做地址计算），指令级并行不足。
  • not selected：warp 已就绪却没被选中发射，往往意味着 SM 上就绪 warp 不足——occupancy 过低或其他 stall 的次生结果。
  • imc miss / mio 等：指令缓存未命中、数学流水线排队等，在规整 kernel 中通常占比很小。

读这页时有两个要点。第一，stall 百分比是「停顿周期」的分布，不是时间分布，各原因之间也不是互斥排队，而是 warp 每次无法发射时记录的首要原因，所以看到 long scoreboard 占 60% 的含义是：全部停顿机会里六成在等全局访存。第二，要结合「每指令平均停顿周期」看绝对量级——有的 kernel stall 分布很均匀，但活跃 warp 多、发射槽始终被别的 warp 填满，停顿并不构成瓶颈；反过来，not selected 占比高时要向上追问是 occupancy 不足还是其他 stall 饿死了就绪队列，单看一页容易因果倒置。一个有用的核对动作：把 stall 结论与 SOL 页交叉验证，long scoreboard 居首却看不到可观的 DRAM 流量，通常意味着访问全部打在 L2/L1 上、问题其实是缓存层级的延迟而非带宽，优化手段也就不同。

五、决策流程：从 stall 分布到处方 ​

把 SOL 百分比、stall 分布和访存指标合起来，得到一套可重复的判定流程。全程不用猜，每一步都有数字支撑。

chip-design/intermediate/nsight-performance-analysis.md D2 示意图

三条分支的典型数字画像值得记住。compute-bound：Compute SM Throughput 接近或超过 80%，stall 中数学流水线（mio / pipe busy）类占比高，roofline 标点在平顶段；处方是换 Tensor Core 指令与更低精度、消除地址计算和边界判断等非主体指令。memory-bound：DRAM Throughput 逼近峰值且 long scoreboard 居首；若同时 gld_efficiency 很低，先修访存合并，因为此时带宽数字是被浪费撑高的；若效率已高，则走分块复用、压缩精度、双缓冲。latency/并行度 bound：两项吞吐都不高、stall 以 not selected 或 stall wait 为主，说明 SM 缺少足够的就绪 warp 来填空，处方是提高 occupancy、用更多独立访存增加指令级并行、均衡线程负载消除 barrier 空等。

还有两个容易误判的边界情况值得交代。第一，roofline 标点远低于屋顶、SM 和 DRAM 吞吐都低时，别急着在三类处方里挑一个——先查问题规模是否真的能填满 GPU：grid 尺寸小于 SM 数、每个 SM 摊不到一个 block 时，所有百分比都会难看，这是并行规模问题而非 kernel 实现问题，先增大批量或合并工作量。第二，ncu 的重放机制会改变时序，它报告的绝对时长不能代表真实运行时间，判断优化有没有生效，真实墙钟要用 nsys 或应用自身的计时为准；ncu 只负责回答「为什么慢」，不负责回答「慢了多少」。采集权限方面，较新的驱动上普通用户默认只能采有限计数器，遇到指标显示 N/A 或报权限错误时，需要管理员把 NVreg_RestrictProfilingToAdminUsers=0 写入 nvidia 内核模块参数并重新加载模块（容器内则要确认没有禁用性能计数器的安全策略），这类环境问题在多租户服务器上很常见，提前确认能省掉一轮无效采集。

一个完整的实战回合是这样的：nsys profile --stats=true 跑完发现 GPU 利用率只有 35%，时间线显示 host 线程在每个 batch 之间有约 20ms 的空白且夹着 cudaMalloc——先把分配移出循环、改为内存池，利用率升到 70%；剩余时间里一个 reduce_kernel 占用最长，ncu 采一次，SOL 显示 DRAM 75%、SM 12%，stall 中 long scoreboard 占六成，gld_efficiency 只有 40%——确诊为非合并访问导致的 memory-bound；改输入布局使 warp 内访问连续后，效率回到 90%+，kernel 时长缩短近一半。注意每一轮只改一类因素、再用同一套命令复测，避免多处改动互相掩盖效果。

这一篇在知识体系里的位置 ​

总结 ​

Nsight Systems Compute 是按观测尺度分工的一对工具：先用 nsys 低开销采全局时间线，沿 CPU/API/H2D/kernel 轨道抓 GPU 空转、拷贝不重叠、小 kernel 过多三类宏观病灶，产出 1-3 个嫌疑 kernel；再用 ncu 重放采集硬件计数器，SOL 页给出 SM/DRAM 吞吐占峰值的百分比（即表格化 roofline），Warp State 页把停顿归因为 long scoreboard、barrier、not selected 等，配合 gld_efficiency 与 occupancy 判定三类瓶颈——SM 高内存低为 compute-bound（上 Tensor Core、向量化），DRAM 高且 long scoreboard 居首为 memory-bound（先修访存合并，再分块复用、降精度），两项都低且 not selected 居首为 latency/并行度 bound（提 occupancy、增 ILP、均衡负载）；命令行上 nsys profile -t cuda,nvtx,osrt 摸底、ncu --set full --kernel-name --launch-skip/--launch-count 定点，每轮只改一类因素并用同一命令复测，是这套双工具工作流不变的纪律。