Nsight Systems CLI (nsys) 快速入门
基于 A100-SXM4-80GB + CUDA 13.1 实测。
nsys是 NVIDIA 系统级性能分析器,从时间线视角展示 CPU-GPU 交互、kernel 执行时长和 stream 并发效果。它与ncu(kernel 内部)和nvbandwidth(裸带宽)形成三件套工作流。
1. nsys 在三件套中的位置
06_nsight_compute_cli.md 的定位表在此展开:
| 工具 | 分析层级 | 核心问题 | 输出粒度 |
|---|---|---|---|
nvbandwidth |
裸硬件 | PCIe/HBM 带宽是否达标? | 一个数字 (GB/s) |
ncu |
单 kernel | Memory bound 还是 Compute bound? | 每个 kernel 的 200+ metrics |
nsys |
全系统时间线 | CPU 在等 GPU 吗?stream 重叠了吗?第一个 cudaMalloc 为什么花了 200ms? | 每个 CUDA API 调用的时间戳 |
ncu 告诉你”kernel 为什么慢”,nsys 告诉你”谁在等谁”。
对于以下问题,nsys 是首选工具:
- 为什么程序启动这么慢?→
cuda_api_sum看 cudaMalloc 和 context 初始化 - 我的 stream 重叠真的发生了吗?→ 时间线直接可视化
- CPU 侧是 poll/ioctl 还是计算逻辑在拖慢整体?→
osrt_sum看 OS runtime
2. 快速上手
2.1 安装
nsys 通常随 CUDA Toolkit 安装在独立路径。本环境:
/usr/local/bin/nsys --version
# NVIDIA Nsight Systems version 2025.5.2.266-255236693005v0
注意:
nsys不在<CUDA_PATH>/bin下,而是有独立的安装路径。在服务器上位于/usr/local/bin/nsys。
2.2 基本用法
# 基础 profiling(生成 .nsys-rep 报告文件)
nsys profile -o my_report ./my_app
# 自动输出统计摘要(最常用)
nsys profile --stats=true -o my_report ./my_app
# 限制 trace 类型(减少 overhead 和文件体积)
nsys profile --trace=cuda,osrt -o my_report ./my_app
# 查看已有报告的统计摘要
nsys stats my_report.nsys-rep
2.3 默认报告类型
nsys 默认生成 6 个统计报告(即使不使用 --stats=true,nsys stats 也会输出它们):
| 报告 | 文件 | 回答的问题 |
|---|---|---|
| osrt_sum | OS Runtime Summary | CPU 侧时间花在哪了?poll/ioctl/mmap? |
| cuda_api_sum | CUDA API Summary | cudaMalloc 花了多久?cudaMemcpy 多少次? |
| cuda_gpu_kern_sum | GPU Kernel Summary | 每个 kernel 跑了多久、跑了多少次? |
| cuda_gpu_mem_time_sum | GPU Memory Time Summary | H2D/D2H 总时间? |
| cuda_gpu_mem_size_sum | GPU Memory Size Summary | 传输了多少数据? |
| nvtx_sum | NVTX Summary | 自定义 annotation 区域的时间(需代码内嵌 NVTX) |
3. A100 实测:vectorAdd
以官方 cuda-samples 中的 vectorAdd 为目标程序:
cd cuda-samples/Samples/0_Introduction/vectorAdd
nvcc -arch=sm_80 -lineinfo -I../../../Common -o vectorAdd vectorAdd.cu -lcudart
nsys profile --stats=true -o vectorAdd_report ./vectorAdd
3.1 CUDA API Summary — CPU 侧 API 开销
Time (%) Total Time (ns) Num Calls Avg (ns) Med (ns) Name
-------- --------------- --------- ---------- -------- --------------
99.7 209,435,653 3 69,811,884 4,990 cudaMalloc
0.1 233,650 3 77,883 37,340 cudaMemcpy
0.1 194,810 3 64,936 17,670 cudaFree
0.0 75,690 1 75,690 75,690 cudaLaunchKernel
解读:
cudaMalloc占 99.7% 的 CUDA API 时间——首次调用触发 CUDA context 初始化 + 物理显存分配,属于一次性开销。Median 仅 4,990 ns 说明后续 2 次分配是快的cudaMemcpy(3 次,avg 77.9 μs)和cudaLaunchKernel(1 次,75.7 μs)是常态开销- 如果看到
cudaMemcpy的 avg » med(如上表:77.9 μs vs 37.3 μs),说明某次 memcpy 特别大——查cuda_gpu_mem_size_sum
3.2 GPU Kernel Summary — GPU 侧执行
Time (%) Total Time (ns) Instances Avg (ns) Name
-------- --------------- --------- -------- --------------------
100.0 3,040 1 3,040.0 vectorAdd(...)
vectorAdd kernel 在 GPU 上仅执行了 3.04 μs。对比前面 cudaLaunchKernel 的 75.7 μs,launch 开销是 kernel 执行时间的 25 倍——这就是为什么小 kernel 应该用 CUDA Graph(见 09_cuda_graphs.md)。
3.3 GPU Memory Time Summary — H2D/D2H 时间线
Time (%) Total Time (ns) Count Avg (ns) Operation
-------- --------------- ----- -------- ----------------------------
71.7 30,944 2 15,472 [CUDA memcpy Host-to-Device]
28.3 12,192 1 12,192 [CUDA memcpy Device-to-Host]
2 次 H2D(input 数组拷贝,15.5 μs 每次),1 次 D2H(结果拷回,12.2 μs)。总计 memcpy 43.1 μs,占比远超 kernel 执行时间(3.0 μs)。
3.4 GPU Memory Size Summary — 数据传输量
Total (MB) Count Avg (MB) Operation
---------- ----- -------- ----------------------------
0.400 2 0.200 [CUDA memcpy Host-to-Device]
0.200 1 0.200 [CUDA memcpy Device-to-Host]
2 次 H2D 每次 0.2 MB(50,000 floats × 2 arrays = 400 KB → 每个数组 200 KB),1 次 D2H 0.2 MB。总共 0.6 MB。
3.5 OS Runtime Summary — CPU 侧在干什么
Time (%) Total Time (ns) Num Calls Avg (ns) Name
-------- --------------- --------- ------------ --------------
65.3 688,931,159 14 49,209,368 poll
33.9 357,524,255 1,089 328,305 ioctl
0.4 4,371,933 53 82,489 mmap64
解读:
poll(65.3%, 689 ms)——程序在等 GPU 完成(cudaDeviceSynchronize或隐式同步),这是正常的ioctl(33.9%, 358 ms)——CUDA driver 通信开销,1089 次 ioctl 调用均值 328 μs,中位数 28 μs,说明大部分很快但少数很大mmap64/fopen(< 0.5%)——程序加载和数据文件 I/O
经验法则:如果
poll占比 > 50%,说明程序是 GPU-bound(CPU 在等 GPU)。如果ioctl占比高且次数多,考虑 CUDA Graph 减少 driver 往返。
4. 常用分析场景
4.1 确认 stream 重叠
# 采集时间线
nsys profile --trace=cuda -o timeline ./stream_overlap
# CLI 无法直接看 Gantt 图,但可以从 kernel 时间戳推断
nsys stats --report cuda_gpu_kern_sum timeline.nsys-rep
CLI 的局限:时间线图需要 GUI。但通过对比 cuda_gpu_kern_sum 中各 kernel 的总执行时间和 wall-clock 时间,可以推断重叠效果。
4.2 找启动瓶颈
nsys profile --stats=true -o startup ./my_app
# 重点看 cuda_api_sum 的第一个 cudaMalloc 时间
首次 cudaMalloc 超 100ms → CUDA context 初始化。解决方案:预先 warmup 或使用 cuInit 提前初始化。
4.3 批量分析多个程序的对比
nsys profile -o baseline ./baseline_app
nsys profile -o optimized ./optimized_app
nsys stats baseline.nsys-rep > baseline.txt
nsys stats optimized.nsys-rep > optimized.txt
diff baseline.txt optimized.txt
5. 三件套工作流
nvbandwidth
┌──────────┐
│ 硬件基线 │ ← PCIe/HBM 带宽是否正常?
└────┬─────┘
│
┌────▼─────┐
│ nsys │ ← 全局时间线:CPU 在等什么?stream 重叠了吗?
└────┬─────┘
│ 定位到具体的慢 kernel
┌────▼─────┐
│ ncu │ ← 进入 kernel:Memory bound?Occupancy 低?
└──────────┘
| 步骤 | 工具 | 问题 | 行动 |
|---|---|---|---|
| 1 | nvbandwidth |
硬件基线正常吗? | 若否 → 查 PCIe 插槽/链路状态 |
| 2 | nsys |
CPU-GPU 时间线合理吗?stream 重叠了吗? | 定位异常的 API 调用或 kernel 序列 |
| 3 | ncu |
单个 kernel 内部瓶颈是什么? | Memory → 查 coalescing;Compute → 查指令效率;Occupancy → 调 grid/block |