GPU 性能剖析:Nsight Systems 与 Nsight Compute 实战

GPU 跑不满、Kernel 慢却不知卡在哪,是推理优化的日常困境。本文拆解 NVIDIA 剖析双工具的分工:Nsight Systems 看时间线(谁在等谁、Kernel 间空隙、数据传输与同步),Nsight Compute 看单个 Kernel 内部(占用率、访存、算力与带宽利用率)。给出采集命令、关键指标解读、典型瓶颈模式与从剖析到优化的闭环流程。

推理服务里最常见的场景是:nvidia-smi 显示 GPU 利用率只有 40%,token 吞吐却上不去;或者一个 Kernel 比别人慢 3 倍,但代码看起来一样。这类问题靠「看代码猜」几乎无解,必须用剖析工具把 GPU 上真实发生的事摊开来看。NVIDIA 的剖析工具链分两层:Nsight Systems(系统级时间线)和 Nsight Compute(Kernel 级计数器)。搞清两者的分工,才能从「哪里慢」一路定位到「为什么慢」。

前置:CUDA 内存管理与性能优化 、算子融合与 Kernel 优化 、推理性能调优检查清单 。

一、两把刀:先看时间线,再看 Kernel

剖析的常见错误是「一上来就 profile 单个 Kernel」——但如果瓶颈根本不在 Kernel 里(而是在数据拷贝、同步或启动开销上),Kernel 剖析再细也没用。正确顺序是自上而下:

工具命令粒度回答的问题
Nsight Systemsnsys进程 / 时间线时间花在哪个阶段?Kernel 之间有空隙吗?谁在等谁?
Nsight Computencu单个 Kernel这个 Kernel 内部是算力受限还是访存受限?占用率够吗?
定位顺序(自顶向下):
1. nsys 看全局:CPU 是否在等 GPU?GPU 是否有大段空闲?
   → 有空闲 = 调度/拷贝/同步问题,先解决
2. nsys 看热点:哪个 Kernel 占的时间最多?占比多少?
   → 找出 Top-N Kernel,才算进入 Kernel 优化
3. ncu 看内部:热点 Kernel 为何慢?算力/带宽/延迟哪个受限?
   → 拿到具体指标,才知道该改什么

工程要点:nsys 负责「找嫌疑人」,ncu 负责「审问嫌疑人」。跳过 nsys 直接 ncu,很可能把一个只占 5% 时间的 Kernel 优化到极致,而整体性能纹丝不动。先用 nsys 确认瓶颈在哪一层,再决定要不要 ncu。

二、Nsight Systems:把时间线摊开

nsys 采集的是一段时间内 CPU 线程、CUDA API、Kernel、内存拷贝、同步事件的完整时间线。

2.1 采集命令

# 基本采集:整个进程
nsys profile -o report_name --force-overwrite true python infer.py

# 只采 CUDA/NVTX 相关,减少噪音
nsys profile -t cuda,nvtx,osrt -o report_name ./app

# 延迟采集:跳过预热阶段,只采稳态(避免冷启动干扰)
nsys profile --delay 10 --duration 20 -o report_name ./app

# 加上 GPU 计数器采样(功耗/利用率/温度)
nsys profile --gpu-metrics-device=0 -o report_name ./app

采集产物是 .nsys-rep,用 GUI 打开(nsys-ui report.nsys-rep),或直接命令行导出统计:

# 命令行汇总:Kernel 耗时排行 + CUDA API 耗时排行
nsys stats --report cuda_gpu_kern_sum,cuda_api_sum report_name.nsys-rep

2.2 时间线上要看的四件事

① GPU 空闲间隙(Gap)
   Kernel A 结束到 Kernel B 开始之间有没有空隙?
   → 有空隙 = CPU 提交慢 / 同步阻塞 / 启动开销大
   → 小 Kernel 密集场景,空隙可能吃掉 30%+ 时间

② CPU 是否在等 GPU
   看 CPU 线程状态:是否长时间停在 cudaStreamSynchronize / cudaMemcpy
   → 说明 CPU 被同步拖住,应改异步(async copy + stream)

③ 数据传输占比
   H2D / D2H 拷贝是否与计算重叠?
   → 未重叠的拷贝是纯浪费,应上 pinned memory + 多流

④ 同步点
   cudaDeviceSynchronize 是否频繁?每次同步都会打断流水线
典型症状 → 结论对照:
GPU 利用率 30%,时间线大段空白   → 启动/调度瓶颈(CPU 提交慢或同步多)
Kernel 之间有细密空隙            → 小 Kernel 启动开销(考虑融合/CUDA Graph)
D2H 拷贝与 Kernel 串行不重叠     → 拷贝未异步化(pinned + stream)
CPU 长期停在 Sync 上             → 同步点过多(去同步化)

2.3 用 NVTX 给时间线加标注

裸时间线上只有 Kernel 名和 API 名,很难对应到业务逻辑。NVTX(NVIDIA Tools Extension) 允许在代码里插入带颜色的区间标记,让时间线直接显示「检索阶段」「生成阶段」这样的语义标签:

import torch.cuda.nvtx as nvtx

def generate(prompt):
    nvtx.range_push("prefill")
    # ... 预填充计算
    nvtx.range_pop()

    nvtx.range_push("decode")
    for _ in range(max_tokens):
        # ... 逐 token 解码
        pass
    nvtx.range_pop()

# 上下文管理器写法(更安全,异常也会 pop)
with nvtx.range("embedding"):
    vec = embed(prompt)
# 采集时带上 nvtx,时间线会显示自定义区间
nsys profile -t cuda,nvtx -o report ./app

有了 NVTX 标注,就能一眼看出「prefill 占了 60% 时间」「decode 里有大段空白」,把 Kernel 层面的数据直接映射到业务阶段,定位效率提升明显。

工程要点:nsys 时间线的价值在于**「看见空隙」**。LLM 推理里,解码阶段每 token 会启动几十上百个小 Kernel,如果没开 CUDA Graph,启动开销能占到总时间的可观比例。时间线上那些「看起来不多」的细密空隙,累计起来往往就是 20%~40% 的损失。相关内容见 算子融合与 Kernel 优化 。

三、Nsight Compute:钻进 Kernel 内部

确认了热点 Kernel,才轮到 ncu。它用硬件计数器回答「这个 Kernel 为什么慢」。

3.1 采集命令

# 剖析指定 Kernel(按名字匹配),采集完整指标集
ncu --kernel-name regex:flash --set full -o k_report ./app

# 快速概览:只看 roofline 与瓶颈提示
ncu --set default --launch-count 5 ./app

# 指定指标集合
ncu --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed,\
gpu__compute_memory_throughput.avg.pct_of_peak_sustained_elapsed ./app

# 只剖析第 N 次调用(跳过预热)
ncu --launch-skip 20 --launch-count 1 ./app

⚠️ --set full 会多次重放 Kernel(每类计数器一轮),对长时间 Kernel 非常慢。生产剖析优先用 --set default 或 --section 精确选择需要的部分。

3.2 关键指标解读

① 占用率(Occupancy)
   Achieved Occupancy:实际活跃 warp / 理论最大 warp
   → 低占用率常见原因:寄存器压力、共享内存超限、block 太小
   → 但低占用不一定是问题:访存受限的 Kernel 不需要高占用

② 计算吞吐(Compute Throughput)
   sm__throughput.avg.pct_of_peak_sustained_elapsed
   → 接近 100% = 算力受限(compute-bound)

③ 内存吞吐(Memory Throughput)
   gpu__compute_memory_throughput.avg.pct_of_peak_sustained_elapsed
   → 接近 100% = 带宽受限(memory-bound)

④ 停顿原因(Warp Stall Reasons)
   Stall Long Scoreboard → 等全局内存(访存延迟)
   Stall Barrier        → 等 __syncthreads
   Stall MIO Throttle   → 共享内存/特殊功能单元排队
   Stall Not Selected   → 可发射的 warp 太多(其实是好事)
指标含义受限类型
sm__throughput 高 + memory 低算力打满compute-bound
memory 高 + sm__throughput 低带宽打满memory-bound
两者都低 + 占用率低延迟受限latency-bound
两者都低 + 占用率高有同步/分支/发散其它

3.3 Roofline:一眼判断受限类型

ncu 的 Roofline 图把 Kernel 的「算术强度(Arithmetic Intensity,FLOP/Byte)」与硬件峰值对比,直观判断离哪条屋顶线更近:

算术强度低(< 硬件拐点)→ 落在带宽屋顶线下 → memory-bound
   → 优化方向:减少访存字节(量化、融合、共享内存复用)
算术强度高(> 硬件拐点)→ 落在算力屋顶线下 → compute-bound
   → 优化方向:用 Tensor Core、减少冗余计算、提高并行度
离两条屋顶线都很远 → 延迟/占用率问题 → 调 block/寄存器/提高并行

3.4 剖析开销与 Kernel 重放

ncu 采集计数器的方式是重放 Kernel——每采集一组硬件计数器,就把 Kernel 重新执行一遍。这带来两个必须知道的后果:

① 耗时放大
   --set full 可能重放几十次,一个 1ms 的 Kernel 要花几十 ms 采完
   → 长 Kernel / 大模型场景,优先用 --section 精确采集

② 状态污染
   重放会重复执行 Kernel 的副作用(如原子加、写全局内存)
   → 对「会改状态」的 Kernel,采集前后要做隔离或用 --replay-mode

③ 只适合稳态 Kernel
   依赖外部随机状态、时间戳的 Kernel 采集结果会失真
# 只采关键 section,避免全量重放
ncu --section SpeedOfLight --section Occupancy \
    --kernel-name regex:"attention" ./app

# 应用重放模式(快照/恢复内存状态,副作用更小)
ncu --replay-mode application --kernel-name regex:"gemm" ./app

工程要点:ncu 的核心是**「用计数器替代猜测」**。同一个 Kernel 在不同输入形状下可能是完全不同的受限类型,别用经验硬套。先看 sm__throughput 与 memory 两个百分比,基本就能判定算力受限还是带宽受限;再用 Warp Stall Reasons 与占用率定位更深层的原因。记住 ncu 会重放 Kernel,采集本身有开销,别拿 ncu 下的耗时当性能基线(基线要用 nsys 或无剖析器时测)。

四、典型瓶颈模式与对策

把常见模式整理成一张速查表,剖析时对号入座:

现象(nsys/ncu)根因对策
GPU 大段空闲CPU 提交慢 / 同步阻塞异步化、多流、CUDA Graph
Kernel 间细密空隙小 Kernel 启动开销算子融合、CUDA Graph 捕获
memory 吞吐 ~100%带宽受限量化(FP8/INT8)、融合减少读写、共享内存复用
sm__throughput ~100%算力受限Tensor Core、提高并行度、减少冗余
占用率低 + 寄存器多寄存器压力限制 __launch_bounds__、拆 Kernel
占用率低 + 共享内存大共享内存超限减小 tile、动态共享内存
Stall Long Scoreboard 高访存延迟未被掩盖增大并行、预取、提高占用率
Stall Barrier 高同步等待减少 __syncthreads、调整 tile
数据传输不重叠拷贝未异步pinned memory + 多流 + cudaMemcpyAsync
同一 Kernel 时快时慢输入形状差异动态 Shape 优化、按形状分桶
一个高频误判:
「占用率低 → 一定是问题」
错。访存受限的 Kernel 靠少量 warp 打满带宽是正常的,
盲目提高占用率反而可能因寄存器/共享内存挤占而变慢。
判断标准:占用率低 + 吞吐也低,才是真问题。

4.1 多流与多卡剖析

单流场景看完了,生产推理往往涉及多流、多卡(TP/PP)甚至多机。这时剖析要额外关注:

多流场景:
□ 各 stream 是否真正并行?还是串行排队(false dependency)
□ 事件(event)等待是否成为隐式同步点
□ 拷贝流与计算流是否重叠

多卡场景(TP/PP):
□ NCCL 通信 Kernel 占比多少?是否与计算重叠?
□ AllReduce/AllGather 是否落在关键路径上
□ 是否有卡在「等别的卡」(负载不均)
# 多进程剖析:每个 rank 一份报告
nsys profile -o rank_%p --force-overwrite true torchrun --nproc_per_node=4 app.py

# 只看 NCCL 通信占比
nsys stats --report cuda_gpu_kern_sum rank_*.nsys-rep | grep -i nccl

多卡推理里,通信开销常被低估——TP 越大,AllReduce 越频繁,如果通信没和计算重叠,扩卡反而变慢。这部分机制可参考同专题的 NCCL 集合通信 与 分布式推理与 GPU 集群调度 。

工程要点:剖析的目标不是「把所有指标调好看」,而是**「找到当前的第一瓶颈并消除它」**。GPU 优化是串行游戏:解决了带宽瓶颈,下一个瓶颈可能变成算力或延迟。每次只优化一个瓶颈,用 nsys + ncu 复测确认收益,别同时改多个地方(否则分不清谁起作用)。

五、从剖析到优化的闭环

剖析不是一次性动作,而是一个「采集 → 定位 → 优化 → 复测」的循环:

Step 1  建立基线
        nsys 采集稳态时间线 + ncu 采集 Top-3 Kernel 指标
        记录:端到端延迟、Kernel 总耗时、GPU 利用率

Step 2  定位第一瓶颈
        自顶向下:先看 nsys 找空隙/热点,再看 ncu 判受限类型

Step 3  单点优化
        只改一个地方(融合 / 量化 / 占用率 / 异步化)

Step 4  复测对比
        同一输入、同一预热,重跑 nsys/ncu,对比指标
        端到端是否变快?Kernel 指标是否改善?

Step 5  记录并迭代
        写进「调优日志」:改动 + 指标前后 + 结论
        当前瓶颈消除后,回到 Step 2 找下一个
# 最小化的剖析脚本(把两步串起来)
nsys profile -t cuda,nvtx --delay 5 --duration 15 -o base ./serve.sh
nsys stats --report cuda_gpu_kern_sum base.nsys-rep | head -20
# 挑出 Top-1 Kernel 名,再用 ncu 深挖
ncu --kernel-name regex:"<TopKernel>" --set default --launch-count 3 ./serve.sh

工程要点:剖析要可复现——固定输入形状、固定预热轮次、固定采集时长,否则指标抖动会让你误判。把每次优化的前后指标记成表(对应 推理基准与压测 的方法论),形成团队的「Kernel 性能档案」。连续剖析(continuous profiling)思路在 持续剖析与性能归因 中有更系统的讨论。

六、速查表与一句话记忆

问题一句话答案
先看什么先 nsys 看时间线(找空隙/热点),再 ncu 看 Kernel 内部
GPU 大段空闲启动/同步/拷贝瓶颈,不是 Kernel 问题
Kernel 间细密空隙小 Kernel 启动开销,融合或 CUDA Graph
判断受限类型看 sm__throughput vs memory__throughput 两个百分比
占用率低不一定有问题;占用率低 + 吞吐低才是真问题
访存慢Stall Long Scoreboard 高 → 访存延迟未掩盖
同步多Stall Barrier 高 → 减少 __syncthreads
采集技巧--delay 跳预热,--launch-skip 跳前 N 次
优化原则一次只改一个瓶颈,复测确认收益
剖析纪律固定输入/预热/时长,指标才可复现

一句话记忆:GPU 剖析 = 先 nsys 看时间线(空隙/热点/拷贝/同步)+ 再 ncu 看 Kernel(占用率/算力/带宽/停顿)+ 用 Roofline 判受限类型 + 对号入座找瓶颈 + 单点优化后复测闭环——「自顶向下定位,用计数器说话,一次改一个」。

延伸阅读

继续阅读

探索更多技术文章

浏览归档,发现更多关于系统设计、工具链和工程实践的内容。

全部文章 返回首页

「ai」更多文章

  1. 排序学习与搜索召回排序系统
  2. 数据版本控制与血缘:DVC 与 LakeFS
  3. 模型可解释性:SHAP、LIME 与注意力归因