如何用ncu命令行统计重复10次的CUDA核函数平均耗时/使用率？

问题描述 #

我有一个包含5个CUDA核函数的测试程序，每个核函数先运行10次预热，再运行10次用于性能测量。我需要获取这10次测量的平均耗时、资源使用率等统计数据。请问如何通过ncu命令行优雅实现？是否需要借助cudaProfilerStart()/End()？另外我希望将结果导出至Excel文件，我是初学者，恳请帮助。测试程序示例如下：

int main()
{
    for (int i = 0; i < 10; i++){
        kernel_1<<<...>>>(...); // warm up
    }
    for (int i = 0; i < 10; i++){
        kernel_1<<<...>>>(...); // to be measured
    }
    ...
    for (int i = 0; i < 10; i++){
        kernel_5<<<...>>>(...); // warm up
    }
    for (int i = 0; i < 10; i++){
        kernel_5<<<...>>>(...); // to be measured
    }
    return 0;
}

解决方案 #

1. 是否需要cudaProfilerStart()/cudaProfilerStop()？ #

可选，但推荐使用。不加的话ncu会记录所有核函数调用（包括预热的10次），后续需手动筛选有效数据；加上后可精准控制只采集测量阶段的运行，避免无效数据干扰。

2. 带代码标记的测量方案（推荐） #

先修改测试代码，在每个核的测量阶段前后添加Profiler开关：

#include <cuda_profiler_api.h> // 必须包含该头文件

int main()
{
    // kernel_1 预热
    for (int i = 0; i < 10; i++){
        kernel_1<<<...>>>(...);
    }
    cudaDeviceSynchronize(); // 确保预热完全完成
    cudaProfilerStart(); // 启动性能采集
    for (int i = 0; i < 10; i++){
        kernel_1<<<...>>>(...);
        cudaDeviceSynchronize(); // 确保单个核运行完成，避免重叠干扰
    }
    cudaProfilerStop(); // 停止性能采集
    cudaDeviceSynchronize();

    // kernel_2 重复上述逻辑
    for (int i = 0; i < 10; i++){
        kernel_2<<<...>>>(...);
    }
    cudaDeviceSynchronize();
    cudaProfilerStart();
    for (int i = 0; i < 10; i++){
        kernel_2<<<...>>>(...);
        cudaDeviceSynchronize();
    }
    cudaProfilerStop();
    cudaDeviceSynchronize();

    // kernel_3、kernel_4、kernel_5 同理执行

    return 0;
}

执行以下ncu命令完成测量：

ncu --profile-from-start off \
    --kernel-regex "kernel_[1-5]" \
    --metrics elapsed_time,sm_efficiency,global_memory_throughput \
    --statistics \
    --csv \
    --output cuda_perf_results.csv \
    ./your_compiled_test_program

参数说明：

• --profile-from-start off：默认不启动采集，仅响应cudaProfilerStart()/Stop()的控制
• --kernel-regex "kernel_[1-5]"：精准匹配需要测量的5个核函数
• --metrics：指定采集的性能指标，这里选了耗时、SM利用率、全局内存吞吐量，可按需添加（如shared_memory_throughput、ipc等）
• --statistics：让ncu自动计算每个核多次运行的平均、最大、最小值
• --csv：输出CSV格式，方便直接导入Excel
• --output：指定结果保存的文件名

3. 无代码修改的测量方案 #

若不想修改代码，直接采集所有核调用后手动筛选：

ncu --kernel-regex "kernel_[1-5]" \
    --metrics elapsed_time,sm_efficiency \
    --csv \
    --output all_results.csv \
    ./your_compiled_test_program

生成的CSV会包含每个核的20次调用记录（10次预热+10次测量），可在Excel中按核函数名称分组，取每组后10条数据计算平均值。

4. 导入Excel并分析 #

1. 打开生成的.csv文件：直接双击用Excel打开，或通过Excel「数据」→「自文本/CSV」导入
2. 数据分析：
   • 若使用了--statistics参数，结果中会直接显示Average列，可直接查看平均耗时和资源使用率
   • 若手动筛选，选中对应核的10次数据，用AVERAGE()函数计算平均值

注意事项 #

• 编译程序时需添加优化选项（如nvcc -O3），模拟真实运行环境
• 测量前关闭其他占用GPU的程序，避免干扰结果准确性
• 若核函数运行时间极短（<1ms），可增加测量次数（如20次）以提高统计可靠性

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