推理服务里最常见的场景是:nvidia-smi 显示 GPU 利用率只有 40%,token 吞吐却上不去;或者一个 Kernel 比别人慢 3 倍,但代码看起来一样。这类问题靠「看代码猜」几乎无解,必须用剖析工具把 GPU 上真实发生的事摊开来看。NVIDIA 的剖析工具链分两层:Nsight Systems(系统级时间线)和 Nsight Compute(Kernel 级计数器)。搞清两者的分工,才能从「哪里慢」一路定位到「为什么慢」。
一、两把刀:先看时间线,再看 Kernel
剖析的常见错误是「一上来就 profile 单个 Kernel」——但如果瓶颈根本不在 Kernel 里(而是在数据拷贝、同步或启动开销上),Kernel 剖析再细也没用。正确顺序是自上而下:
| 工具 | 命令 | 粒度 | 回答的问题 |
|---|---|---|---|
| Nsight Systems | nsys | 进程 / 时间线 | 时间花在哪个阶段?Kernel 之间有空隙吗?谁在等谁? |
| Nsight Compute | ncu | 单个 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 判受限类型 + 对号入座找瓶颈 + 单点优化后复测闭环——「自顶向下定位,用计数器说话,一次改一个」。
延伸阅读
- CUDA 内存管理与性能优化
- 算子融合与 Kernel 优化
- FlashAttention 与高效注意力内核
- 推理性能调优检查清单
- 推理基准与压测
- 持续剖析与性能归因 — 生产环境的持续剖析实践
- GPU 与加速器可观测性 — 监控侧的 GPU 指标采集
继续阅读
探索更多技术文章
浏览归档,发现更多关于系统设计、工具链和工程实践的内容。