CUDA实践（1）--性能分析工具

本文记录几种CUDA实践中常用的运行计时和性能分析工具。

1. 运行计时

虽然标准C语言也有相关计时方法，但是由于CPU与GPU之间的同步问题可能造成测时不准确，这里分别介绍这两种测试方法：

（1）标准C语言计时函数

C语言当前版本中包含一个头文件<time.h>，该文件中定义了一个时间变量clock_t，一个获取到目前为止的运行时间的函数clock()，以及一个将clock()函数结果转换为以秒为单位的常量CLOCK_PER_SEC。下面是一个简单的例子：

clock_t start = clock();
myKernel<<<GridNum, BlockNum, size_smem, stream>>>(parameter list);
clock_t end = clock();
double time = ((double)(end - start))/CLOCK_PER_SEC； 

这样计时得到的结果远小于正确结果，是因为核函数的执行是异步的。核函数一启动，CPU就会继续执行其他任务，它会在内核函数执行完毕之前就读取时钟的值，并将其赋给end。但是这个计时方法可以用于cudaMemcpy()函数的计时，这是因为数据复制默认是同步的，所以当cudaMemcpy()执行完毕时CPU才允许继续执行其他操作。

• 同步操作：对计算流进行阻塞，防止其他操作的执行
• 异步操作：在异步操作执行的同时允许其他操作的执行

（2）CUDA事件计时

CUDA拥有一套自己的计时机制，以防止CPU与GPU之间同步导致的问题，并提供更精准的测量。在CUDA的API中，有一个特殊的数据类型cudaEvent_t，以及几个与CUDA事件相关的关键函数：

• cudaEventCreate()与cudaEventDestroy()负责创建和销毁事件。
• cudaEventRecord()负责记录时间。
• cudaEventSynchronize()用于确保异步函数全部执行完毕。
• cudaEventElapsedTime()负责将两个事件记录转成运行时间（ms）。

cudaEvent_t start, stop; //事件对象
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, stream);

myKernel<<<GridNum, BlockNum, size_smem, stream>>>(parameters list);

cudaEventRecord(stop, stream);
cudaEventSynchronize(stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);

2. 性能分析

（1）NVIDIA Visual Profiler (NVVP)

NVVP是一款跨平台的可视化性能分析工具，其包含在CUDA Toolkit中。NVVP执行程序病生成一个主机端和设备端的时间轴。在NVVP下方还有各种标签页，如Analysis、Details、Console以及Settings。查看Analysis标签页，可以看到有两个图标，一个是数字列表图标，一个是无序符号列表图标，分别表示guided analysis与unguided analysis。用户可以在guided模式下逐步获取不同的性能分析结果，也可以再unguided模式下根据自己的需要获取指定的性能分析结果。

（2）nvprof

nvprof和NVVP本质上是一样的，它能从命令行收集和查看分析数据。

• 时间分析：使用nvprof命令启动编译后的可执行文件，除了输出程序的结果，还输出分析结果。

  nvprof ./a.out

  profiling result：是GPU（kernel函数）上运行的时间；API calls：在CPU上测量的程序调用API的时间。

• 内核可实现占用率：

  nvprof --metrics achieved_occupancy ./a.out

   

• 内核的内存操作效率：

  nvprof --metrics gld_throughput ./a.out

   

• 全局加载效率：

  nvprof --metrics gld_efficiency ./a.out