3 CUDA 学习笔记—— 性能分析、计时与设备号读取

限制内核性能因素：

• 存储带宽
• 计算资源
• 指令和内存延迟

应用性能分析工具检测内核性能：

• nvvp, 独立可视化分析器，显示CPU与GPU上的程序活动的时间表；
• nvprof, 命令行分析器， 在命令行上收集和显示分析数据，获得CPU与GPU上CUDA关联活动的时间表，包括内核执行、内存传输和CUDA的API调用，以及硬件计数器和CUDA内核的性能指标；

nvprof [nvprof_args] <application> [application_args]

 

并行性表现

• 用nvprof检测活跃的线程束:一个内核的可实现占用率被定义为,每周期内活跃线程束的平均数量与一个SM支持的线程束最大数量的比值。

nvprof --metrics achieved_occupancy <application>

• 用nvprof检测内存操作:即检查内核的内存读取效率。

nvprof --metrics gld_throughput <application>

• 用nvprof检测全局加载效率:全局加载效率,即被请求的全局加载吞吐量占所需的全局加载吞吐量的比值，其衡量了应用程序的加载操作利用设备内存带宽的程度。

nvprof --metrics gld_efficiency <application>

注：增大并行性是性能优化的一个重要因数，值得注意的是，最好的执行配置既不具有最高的可实现占用率，也不具有最高的加载吞吐量，没有一个单独的指标能直接优化性能。因此需要在几个相关的指标间寻找一个恰当的平衡来达到最佳的总体性能。

 

使用系统时间计时（添加头文件sys/time.h）

#include<sys/time.h>

double cpuSecond()

{

    struct timeval tp;

    gettimeofday(&tp,NULL);

    return ((double)tp.tv_sec + (double)tp.tv_usec*1.e-6);

}

调用方式如下，由于核函数调用与主机端程序是异步的，因此用cudaDeviceSynchronize()函数来等待所有的GPU线程运行结束

double iStart = cpuSecond();

kernel_name<<<grid,block>>>(argument list);

cudaDeviceSynchronize();

double iEnd = cpuSecond() - iStart;

读取显卡设备号:

printf("%s Starting...\n",argv[0]);

//set up device

int dev = 0;

cudaDeviceProp deviceProp;

CHECK(cudaGetDeviceProperties(&deviceProp,dev));

printf("Using Device %d:%s\n",dev,deviceProp.name);

CHECK(cudaSetDevice(dev));