CUDA程序计时函数总结

在CUDA中统计运算时间，大致有三种方法：

<1>使用cutil.h中的函数
unsigned int timer=0;
//创建计时器
cutCreateTimer(&timer);
//开始计时
cutStartTimer(timer);
{
  //统计的代码段
  …………
}
//停止计时
cutStopTimer(timer);
//获得从开始计时到停止之间的时间
cutGetTimerValue( timer);
//删除timer值
cutDeleteTimer( timer);
 

不知道在这种情况下，统计精度。

 

<2>time.h中的clock函数
clock_t start, finish;
float costtime;
start = clock(); 
{
  //统计的代码段
  …………
}
finish = clock();
//得到两次记录之间的时间差
costtime = (float)(finish - start) / CLOCKS_PER_SEC; 
时钟计时单元的长度为1毫秒，那么计时的精度也为1毫秒。

 

<3>事件event
cudaEvent_t start,stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecend(start,0);
{
  //统计的代码段
   …………
}
cudaEventRecord(stop,0);
float costtime;
cudaEventElapsedTime(&costtime,start,stop);
 
cudaError_t cudaEventCreate( cudaEvent_t* event )---创建事件对象；
cudaError_t cudaEventRecord( cudaEvent_t event，CUstream stream )--- 记录事件；
cudaError_t cudaEventElapsedTime( float* time，cudaEvent_t start，cudaEvent_t end )---计算两次事件之间相差的时间；
cudaError_t cudaEventDestroy( cudaEvent_t event )---销毁事件对象。
计算两次事件之间相差的时间（以毫秒为单位，精度为0.5微秒）。如果尚未记录其中任何一个事件，此函数将返回cudaErrorInvalidValue。如果记录其中任何一个事件使用了非零流，则结果不确定。

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遍计时部分。