使用Nsight Compute生成Roofline分析失败问题求助

核心排查步骤 #

1. 确认内核被正确捕获 #

• 先执行ncu --list-kernels ./my_program，检查是否能识别到你认为包含浮点运算的目标内核。如果列表里没有，说明ncu没attach到目标进程/内核，可尝试添加--target-processes all参数确保捕获所有进程的内核。
• 去掉--launch-count 50参数，避免只采样前50次启动（若目标内核在后期才执行）。如果内核启动次数过多，改用--launch-skip 10 --launch-count 10跳过初始化阶段的小内核，采样稳定运行的批次。

2. 确保采集Roofline所需的Metrics #

• --set full理论上包含所有Metrics，但部分场景下可能遗漏Roofline计算依赖的指标。可直接指定Roofline专用集合：

  ncu -f --set roofline --export profile_rep ./my_program
  

• 若仍有问题，手动指定关键Metrics覆盖浮点运算和内存访问统计：

  ncu -f --metrics smsp__sass_thread_inst_executed_op_fadd_pred_on.sum,smsp__sass_thread_inst_executed_op_fmul_pred_on.sum,smsp__sass_thread_inst_executed_op_ffma_pred_on.sum,smsp__sass_thread_inst_executed_op_dadd_pred_on.sum,smsp__sass_thread_inst_executed_op_dmul_pred_on.sum,smsp__sass_thread_inst_executed_op_dfma_pred_on.sum,dram__bytes.sum,l1tex__t_bytes.sum,l2__t_bytes.sum --export profile_rep ./my_program
  

3. 验证浮点运算未被编译器优化消除 #

• 用nvcc -Xptxas -v your_kernel.cu查看编译时的指令统计，确认有浮点指令生成。
• 用cuobjdump -sass ./my_program解析生成的SASS代码，搜索FADD/FMUL/FFMA（FP32）或DADD/DMUL/DFMA（FP64）指令，确认浮点运算真实存在。若没有，说明编译器因死代码消除（比如运算结果未被使用）移除了浮点逻辑，需修改代码确保结果被有效利用。

4. 检查版本与架构兼容性 #

• 2023.3.0.0版本的ncu对部分新架构（如Ada Lovelace）可能存在Roofline分析的bug，尝试升级到最新版ncu（如2024.x版本）。
• 用nvidia-smi确认GPU型号，若为Kepler及更早架构，Roofline分析支持有限，建议更换较新的GPU测试。

5. 调整采样触发方式 #

• 添加--profile-from-start off参数，然后在代码中插入cudaProfilerStart()和cudaProfilerStop()，精准控制只采样目标内核的执行阶段：

  // 初始化代码...
  cudaProfilerStart();
  your_kernel<<<...>>>(...); // 目标内核
  cudaDeviceSynchronize();
  cudaProfilerStop();
  // 后续代码...
  

• 重新执行ncu命令，确保采样聚焦在目标内核上。

6. 正确查看报告 #

• 必须用ncu-ui ./profile_rep.ncu-rep打开报告，在左侧面板切换到Roofline Chart模块查看分析结果。若导出为CSV等非原生格式，将无法显示Roofline图表。

内容的提问来源于stack exchange，提问作者Severus Snape