CUDA进阶第三篇：CUDA计时方式

写在前面

内容分为两部分，第一部分为翻译《Professional CUDA C Programming》 Section 2. CUDA Programming Model中的TIMING YOUR KERNEL；第二部分为自己的经验。经验不足，欢迎各位大大补充。

写CUDA，追求的就是加速比，想要得到准确的时间，计时函数就是必不可少
计时通常分为两种情况，（1）直接得到接口函数的时间，一般用于得到加速比；（2）获得接口函数内核函数、内存拷贝函数等所耗时间，一般用于优化代码时。
情况（1）方法有两种，CPU计时函数和GPU计时函数。
情况（2）有三种工具nsight，nvvp，nvprof

本博客会详细介绍情况（1）的两种方法；情况（2），nsight不会用，简单介绍一下nvvp和nvprof的用法。

CPU计时函数

在利用CPU计时函数时，要考虑的一个问题是：核函数的执行是异步执行的，所以必须加上核函数同步函数，才能得到准确的时间。
示例代码如下：

double cpuSecond() {
    struct timeval tp;
    gettimeofday(&tp,NULL);
    return ((double)tp.tv_sec + (double)tp.tv_usec*1.e-6);
}
double iStart = cpuSecond();
function(argument list);
cudaDeviceSynchronize();  // 同步函数
double iElaps = cpuSecond() - iStart;

GPU计时函数

GPU计时函数就不需要考虑同步问题，直接用计时事件函数就可以了，示例代码如下：

cudaEvent_t start, stop;
float elapsedTime = 0.0;

cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);

function(argument list);;

cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);

cudaEventElapsedTime(&elapsedTime, start, stop);

cout << elapsedTime << endl;

cudaEventDestroy(start);
cudaEventDestroy(stop);

如何获得精确的计时

正常情况下，第一次执行核函数的时间会比第二次慢一些。这是因为GPU在第一次计算时需要warmup。所以想要第一次核函数的执行时间是不精确的。获得精确的计时我总结为三种，如下

1. 循环执行一百次所需要计时的部分，求平均值，将第一次的误差缩小100倍。这种方法的优点是简单粗暴。但缺点也很明显：(1)程序的执行时间大大增长，特别是比较大的程序(2)要考虑内存溢出问题，C++的内存需要程序猿自己手动管理。写出执行一次不出内存溢出问题的程序很容易，但是写出循环执行一百次而不出内存溢出问题的代码就有一定难度了(3)计时不是特别准确，虽然误差已经被缩小了100倍。
2. 在计时之前先执行一个warmup函数，warmup函数随便写，比如我从cuda sample里的vectoradd作为warmup函数。这种方法的优点是程序执行时间缩短；缺点是需要在程序中添加一个函数，而且因为GPU乱序并行的执行方式，核函数的两次执行时间并不能完全保持一样。所以推荐使用方法3.
3. 先执行warmup函数，在循环10遍计时部分。

nvvp和nvprof的用法

nvprof是自cuda5.0开始存在的一个命令行Profiler，你可以只用nvprof来你代码的一些执行细节。简单用法如下：

$ nvprof ./sumArraysOnGPU-timer

你就可以得到如下：

./sumArraysOnGPU-timer Starting...
Using Device 0: Tesla M2070
==17770== NVPROF is profiling process 17770, command: ./sumArraysOnGPU-timer
Vector size 16777216
sumArraysOnGPU <<<16384, 1024>>> Time elapsed 0.003266 sec
Arrays match.
......

关于nvprof的更多参数信息，可以使用帮助命令：

$ nvprof --help

The NVIDIA Visual Profiler（nvvp）是一款图形化界面的Profiler，也是我一直在用的Profiler。
简单图文教程见链接

写在后面

OpenCUDA：CUDA图像算法开源项目，算法内都有详细的注释，大家一起学习。

---

私人接各种CUDA相关外包（调试、优化、开发图像算法等），有意向请联系，加好友时请注明。