nsight-compute(ncu)比nvprof更适合现代CUDA性能分析,因其支持Turing及更新架构、提供kernel级底层指标,而nvprof在CUDA 11.0+中已被弃用且不支持Ampere等新架构。

为什么 nsight-compute 比 nvprof 更适合现代 CUDA 性能分析
因为 nvprof 已在 CUDA 11.0+ 中被标记为 deprecated,且不支持 Ampere 及更新架构(如 A100、RTX 4090)的完整指标采集。实际 profiling 时若看到 Unsupported GPU architecture 或关键指标(如 s__inst_executed、dram__bytes_read)为空,基本就是 nvprof 失效了。
推荐直接用 nsight-compute(命令行工具 ncu),它原生支持所有 Turing 及之后架构,并能精确到 kernel launch 级别采集指令吞吐、内存带宽、warp 指令分发等底层数据。
-
ncu默认只采样单个 kernel;加--set full才能拿到完整指标集(含 L2、DRAM、Tensor Core 利用率) - 若程序含多个 kernel,用
--kernel-id all或--kernel-name ".*copy.*"过滤目标 kernel - 避免在 X11 图形会话下运行 —— 容易触发
NCU: Error: CUresult driver error: 999 (CUDA_ERROR_UNKNOWN),改用 TTY 或ssh -X后加export DISPLAY=显式指定
如何用 ncu 快速定位 memory-bound kernel
大多数性能瓶颈不在计算,而在显存带宽或延迟。用 ncu 跑一次 baseline,重点看三组比率:
-
lts__t_sectors.sum.average.pct_of_peak_sustained_elapsed:L2 缓存扇区读写占峰值带宽百分比 -
dram__bytes.sum.per_second:实际 DRAM 带宽(GB/s),对比卡标称值(如 A100 PCIe 是 1555 GB/s) -
sm__sass_thread_inst_executed_op_ld.sum和sm__sass_thread_inst_executed_op_st.sum:load/store 指令数,比值明显偏离 1:1 说明访存模式不均衡
如果 dram__bytes.sum.per_second 接近硬件上限但算力利用率(sm__inst_executed)很低,大概率是 global memory 访问未合并 —— 检查 kernel 中数组索引是否满足 threadIdx.x * sizeof(float) 对齐,以及是否用了 __ldg() 提升只读缓存命中。
立即学习“C++免费学习笔记(深入)”;
ncu --set full --kernel-id all --metrics sm__inst_executed,sm__sass_thread_inst_executed_op_ld,sm__sass_thread_inst_executed_op_st,dram__bytes.sum.per_second,lts__t_sectors.sum.average.pct_of_peak_sustained_elapsed ./my_cuda_app
为什么 nsight-systems(nsys)不能替代 ncu
nsys 是 timeline 工具,擅长看 CPU-GPU 协同、kernel 启动间隔、memory copy 与 compute 的重叠程度;但它不提供每个 kernel 的 micro-architectural 指标。两者不是互斥,而是互补。
- 先用
nsys record -t cuda,nvtx ./my_app看整体 timeline,确认是否存在 host 等待、kernel launch 频次过高、或 memcpy 占用过多时间 - 再用
ncu针对 timeline 中耗时最长的 1–2 个 kernel 做深度分析 - 若
nsys显示 kernel duration 很短(ncu 报告高 latency,可能是 kernel 内部有 divergent branch 或大量 __syncthreads(),需结合--metrics sms__warps_launched,sms__inst_executed看 warp 利用率
常见 ncu 错误及绕过方法
实际跑 ncu 时最常遇到三类失败:
-
NCU: Error: CUresult driver error: 700 (CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES):kernel 占用太多寄存器或 shared memory,导致无法启动 profiling context。加--launch-skip跳过首次 launch,或用--unified-memory-profiling off关闭 unified memory tracking -
NCU: Error: Profiling is not supported on this device:当前 GPU 不在ncu --list-gpus输出中,检查是否用的是 vGPU、Tesla 模式未启用(nvidia-smi -i 0 -c 3)、或驱动版本太旧(至少需要 R460+) - 输出中大量指标为
-或0:默认采样集太轻量,必须显式指定--set full或自定义--metrics,不能依赖默认行为
真正难的从来不是跑出数据,而是理解 sm__inst_executed_op_fadd_pred_on.sum 和 sm__inst_executed_op_fmul_pred_on.sum 的比值为何偏离理论 FMA 比例 —— 这往往指向 kernel 中隐式的类型转换或编译器未展开的循环。









