17370845950

ncu 命令行分析需确保同步、多次运行取中位数，并按瓶颈类型关注关键指标：compute-bound 核函数应重点考察 fadd/fmul/ffma 指令数及 IPC。

nsight-compute 命令行基本分析流程

直接运行 ncu 即可对 CUDA 可执行文件做基础 profiling，无需额外编译标记（但需确保程序含 cudaDeviceSynchronize() 或等效同步点，否则核函数可能未完成就被采样中断）。

常见误操作是只跑一次就下结论——ncu 默认仅采集单次 kernel launch，而 GPU 频率、缓存预热、内存碎片都会影响结果。实际应使用 --set full 或指定关键 metric，并配合 --repeat 3 多次运行取中位数。

• ncu --set full --repeat 3 ./my_cuda_app：启用全指标集并重复三次
• ncu -k my_kernel_name ./my_cuda_app：只 profiling 名为 my_kernel_name 的核函数（注意名称需与 nvcc -lineinfo 编译后符号一致）
• 若报错 No kernels were profiled，大概率是程序没触发 kernel launch，或 launch 后未同步（cudaStreamSynchronize(0) 或 cudaDeviceSynchronize() 缺失）

定位瓶颈：看哪些 metric 最关键

不是所有指标都同等重要。对 compute-bound 核函数，优先盯紧 sm__inst_executed_op_fadd_pred_on.sum、sm__inst_executed_op_fmul_pred_on.sum 和 sm__inst_executed_op_ffma_pred_on.sum 这三类指令计数，它们反映实际浮点计算吞吐；再对比 sm__cycles_elapsed 算出 IPC（每周期指令数），IPC

对 memory-bound 场景，重点看 l1tex__t_bytes.sum（L1/纹理单元总字节数）、lts__t_sectors.sum（L2 扇区数）和 sys__bytes.sum（显存带宽占用）。如果 l1tex__t_bytes.sum / sm__cycles_elapsed 显著低于硬件峰值（如 A100 的 ~1.2 TB/s），说明 cache 利用率低，大概率是访存 pattern 不连续或未合并。

• 避免只看 achieved_occupancy：它高不代表快，比如大量空循环也能占满 occupancy
• inst_per_warp 值过低（
• 若 gpu__time_duration 远大于 sm__cycles_elapsed * gpu_clock，说明 kernel 被 preempted 或有隐式同步开销

源码关联与 inline assembly 分析

要让 ncu 显示源码行号甚至 SASS 指令，编译时必须加 -g -lineinfo（nvcc -g -lineinfo -O3 kernel.cu），且不能 strip 符号表。运行时用 ncu --source-base . --set full ./app

遇到内联 PTX（如 asm volatile("shfl.sync.down.b32 ...")）或模板展开过深的 kernel，ncu 可能无法准确定位到 C++ 行。此时可先用 cuobjdump -sass ./app 提取 SASS，再对照 ncu 输出中的 pc（program counter）列，人工匹配热点指令位置。

• 若 ncu 报 Source not found，检查路径是否含空格或软链接——ncu 不解析 symlink，需用真实路径
• --unified-memory-profiling on 可捕获 Unified Memory page fault，但会显著拖慢 profiling 速度，仅在怀疑缺页导致 stall 时开启
• 使用 ncu --csv -f report.csv ... 导出 CSV 后，可用 Python 快速统计各 kernel 的 sm__warps_launched 与 sm__inst_executed_op_fadd_pred_on.sum 比值，识别算力浪费最严重的 kernel

常见陷阱：profiling 结果失真的几个原因

GPU clock 动态调频会让 sm__cycles_elapsed 在不同 run 间波动，尤其在笔记本或共享服务器上。默认 ncu 不锁频，导致 latency 类指标不可比。真正做深度对比前，务必先用 nvidia-smi -r 重置 GPU，再执行 nvidia-smi -lgc 1200（设为固定 1200 MHz）——注意 A100/A800 等数据中心卡需 root 权限且支持 persistence mode。

另一个高频问题是 kernel 被拆成多个 sub-launch（如由 cuBLAS 或 Thrust 触发），而 ncu 默认只显示 top N 个，漏掉真正耗时的子核。此时要用 --kernel-id all 强制采集全部，或先用 nsys profile ./app 做 trace 级概览，再根据 timeline 定位具体 kernel ID 后单独分析。

• 在 WSL2 下无法使用 ncu：Nsight Compute 不支持 Windows 子系统，必须在原生 Linux 环境运行
• 容器内 profiling 需挂载 /dev/nvidiactl、/dev/nvidia-uvm 和 /dev/nvidia0，且镜像内需安装匹配版本的 nvidia-cuda-toolkit
• 若 kernel 运行时间 ncu 可能因采样精度不足返回全零指标——改用 nsys profile --trace=cuda,nvtx --force-overwrite true 获取更细粒度时间戳

ncu --set full --kernel-id all --duration 10000 --timeout 30000 \
  --metrics sm__inst_executed_op_fadd_pred_on.sum,sm__inst_executed_op_fmul_pred_on.sum,sm__inst_executed_op_ffma_pred_on.sum \
  ./my_cuda_app

复杂点在于，同一 kernel 在不同数据规模下瓶颈可能完全不同：小 batch 时 register pressure 主导，大 batch 时 L2 bandwidth 成瓶颈。别依赖单次 profiling 定论，得按典型输入尺寸分段测。