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 不连续或未合并。
立即学习“C++免费学习笔记（深入）”；
- 避免只看
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，其中 --source-base . 告诉工具从当前目录找源文件。
遇到内联 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 定论，得按典型输入尺寸分段测。
