如何使用NVIDIA Nsight分析c++ CUDA程序的性能？ (GPU Profiling)

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 中隐式的类型转换或编译器未展开的循环。