本文记录几种CUDA实践中常用的运行计时和性能分析工具。1. 运行计时虽然标准C语言也有相关计时方法但是由于CPU与GPU之间的同步问题可能造成测时不准确这里分别介绍这两种测试方法1标准C语言计时函数C语言当前版本中包含一个头文件time.h该文件中定义了一个时间变量clock_t一个获取到目前为止的运行时间的函数clock()以及一个将clock()函数结果转换为以秒为单位的常量CLOCKS_PER_SEC。下面是一个简单的例子clock_t start clock(); myKernelGridNum, BlockNum, size_smem, stream(parameter list); clock_t end clock(); double time ((double)(end - start))/CLOCKS_PER_SEC这样计时得到的结果远小于正确结果是因为核函数的执行是异步的。核函数一启动CPU就会继续执行其他任务它会在内核函数执行完毕之前就读取时钟的值并将其赋给end。但是这个计时方法可以用于cudaMemcpy()函数的计时这是因为数据复制默认是同步的所以当cudaMemcpy()执行完毕时CPU才允许继续执行其他操作。同步操作对计算流进行阻塞防止其他操作的执行异步操作在异步操作执行的同时允许其他操作的执行2CUDA事件计时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); myKernelGridNum, BlockNum, size_smem, stream(parameters list); cudaEventRecord(stop, stream); cudaEventSynchronize(stop); float elapsedTime; cudaEventElapsedTime(elapsedTime, start, stop);2. 性能分析1NVIDIA Visual Profiler (NVVP)NVVP是一款跨平台的可视化性能分析工具其包含在CUDA Toolkit中。NVVP执行程序病生成一个主机端和设备端的时间轴。在NVVP下方还有各种标签页如Analysis、Details、Console以及Settings。查看Analysis标签页可以看到有两个图标一个是数字列表图标一个是无序符号列表图标分别表示guided analysis与unguided analysis。用户可以在guided模式下逐步获取不同的性能分析结果也可以再unguided模式下根据自己的需要获取指定的性能分析结果。2nvprofnvprof和NVVP本质上是一样的它能从命令行收集和查看分析数据。时间分析使用nvprof命令启动编译后的可执行文件除了输出程序的结果还输出分析结果。nvprof ./a.outprofiling result是GPUkernel函数上运行的时间API calls在CPU上测量的程序调用API的时间。内核可实现占用率nvprof --metrics achieved_occupancy ./a.out内核的内存操作效率nvprof --metrics gld_throughput ./a.out全局加载效率nvprof --metrics gld_efficiency ./a.out