CUDA进阶补充篇：详析各种CUDA函数计时函数

写在前面：之前写了一篇CUDA进阶第三篇：CUDA计时方式，列出了几种当时遇到的CUDA计时方式，只是个教程式的东西，没有太多技术含量，也不太全面。前几天在CUDA Professional（45157483）群里和大佬们讨论到CUDA官方event函数在计时cpu和cpu混合代码时有问题，虎躯一震，tm这么多年要是一直用的都是错的就瞎了。今天特花时间实验探究一番。有不足之处还望各位前辈指点。

概要

本文分为两部分，前半部分为测验不同计时函数在计时CUDA函数的表现以及分析出的一个坑。后半部分为分析了GPGPU-sim仿真器中cudaevent计时函数的源码。

不同计时函数在计时CUDA函数的表现

实验设计

选取一段cpu和gpu混合代码（这里选择的cuda samples里的vectorAdd），分别用四种不同的计时函数统计程序运行时间进行对比。四种计时函数如下：

• gettimeofday（）
• 官方推荐的cudaEvent方式
• clock()函数
• c++中的chrono库

代码简单思路

vectorAdd的代码比较简单，只有一个核函数global void vectorAdd()和一个main()函数。
main()函数内如下：

__global__ voidvectorAdd(const float *A, const float *B, float *C, int numElements){/*节省空间删掉*/}__global__ void warmup(){/*预热GPU，调用一个空的核函数*/}double cpuSecond() {    struct timeval tp;    gettimeofday(&tp,NULL);    return ((double)tp.tv_sec + (double)tp.tv_usec*1.e-6);}/** * Host main routine */intmain(int argc, char **argv){    // 预热GPU    warmup<<<1, 1>>>();    cudaDeviceSynchronize();    // 变量申请    double start = 0.0f, end = 0.0f;    float elapsedTime = 0.0;    cudaEvent_t event_start, event_stop;    clock_t clock_start;    clock_t clock_end;    std::chrono::time_point<std::chrono::system_clock> c11_start, c11_end;    // 四种计时方式    if(atoi(argv[1]) == 1) {        start = cpuSecond();    } else if(atoi(argv[1]) == 2) {        cudaEventCreate(&event_start);        cudaEventCreate(&event_stop);        cudaEventRecord(event_start, 0);    } else if(atoi(argv[1]) == 3) {        clock_start = clock();    } else if(atoi(argv[1]) == 4) {        c11_start = system_clock::now();    }    /*vectorAdd代码，包含内存申请、初始化、拷贝、计算、拷回、释放。数据量大小为5000000*/    if(atoi(argv[1]) == 1) {        // 如果使用CPU计时方式，一定要加同步函数！！        cudaDeviceSynchronize();        end = cpuSecond();        printf("gettimeofday time = %lfms\n", (end - start) * 1000);    } else if(atoi(argv[1]) == 2) {        cudaEventRecord(event_stop, 0);        cudaEventSynchronize(event_stop);        cudaEventElapsedTime(&elapsedTime, event_start, event_stop);        printf("cudaevent time = %lfms\n", elapsedTime);    } else if(atoi(argv[1]) == 3) {        cudaDeviceSynchronize();        clock_end= clock();        double clock_diff_sec = ((double)(clock_end- clock_start) / CLOCKS_PER_SEC);        printf("clock_ time: %lfms.\n", clock_diff_sec * 1000);    }else if(atoi(argv[1]) == 4) {        cudaDeviceSynchronize();        c11_end = system_clock::now();        int elapsed_seconds = std::chrono::duration_cast<std::chrono::milliseconds>                             (c11_end-c11_start).count();        printf("chrono time: %dms.\n", elapsed_seconds);    }}

实验结果

实验平台：
GPU ： Tesla K80
系统 ： Centos 6
gcc ： 4.7.2

时间统计结果如下：

计时方式 Time(ms) 评价 gettimeofday() 326.971769ms 不太稳定，上下有大概20ms的浮动 cudaEvent 328.312744ms 上下3ms左右的浮动 clock() 330ms 很稳定 chrono 324ms 上下3ms左右的浮动

---

从实验结果可以看出，后三种计时方式都是比较稳定，可以放心使用。

坑

细心的人可能会发现，我在代码最前面加了一个空的warmup函数。这个在精确统计时间是非常重要的！！！因为GPU第一次被调用时会消耗不定的时间来预热。
如果把预热那行注释掉，得到的计时结果如下：

计时方式 Time(ms) gettimeofday() 535.159826ms cudaEvent 346.573151ms clock() 440.000000ms chrono 470ms

可以看出，1,3,4三种CPU计时方式结果与真实结果大相径庭，cudaEvent还算比较接近。
所以个人比较推荐的精确计时方式为：（1）前面加warmup函数；（2）循环N(比如100次)然后求平均；（3）针对某个kernel函数，用nvvp或者nvprof看精准的时间。

cudaevent计时函数源码分析

GPGPU-Sim是一款cycle级别的GPU仿真器。我之前也写过几篇介绍GPGPU-sim的博客。我从GPGPU-Sim的源码中找到了cudaEvent计时方式的源码，简单分析了一下。
cudaEvent计时方式的流程如下，核心函数为cudaEventRecord(）和cudaEventElapsedTime()

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);cudaEventDestroy(start);cudaEventDestroy(stop)

其中cudaEventRecord()函数的源码如下：

__host__ cudaError_t CUDARTAPI cudaEventRecord(cudaEvent_t event, cudaStream_t stream){    CUevent_st *e = get_event(event);    if( !e ) return g_last_cudaError = cudaErrorUnknown;    struct CUstream_st *s = (struct CUstream_st *)stream;    stream_operation op(e,s);    g_stream_manager->push(op);    return g_last_cudaError = cudaSuccess;}

其中cudaEventElapsedTime()函数的源码如下：

__host__ cudaError_t CUDARTAPI cudaEventElapsedTime(float *ms, cudaEvent_t start, cudaEvent_t end){    time_t elapsed_time;    CUevent_st *s = get_event(start);    CUevent_st *e = get_event(end);    if( s==NULL || e==NULL )        return g_last_cudaError = cudaErrorUnknown;    elapsed_time = e->clock() - s->clock();    *ms = 1000*elapsed_time;    return g_last_cudaError = cudaSuccess;}

可以看出cudaEventRecord()函数其实就是一个流操作，并压入一个栈中g_stream_manager。
在cudaEventElapsedTime函数中，在栈中找到两次压入的cudaEvent，然后再调用clock()函数计算时间差。而这个clock()函数本质是一个time_t类型。所以归根到底还是调用CPU端的计时函数进行计时。

PS:因为GPU是商业软件，不公开内部实现细节。所以GPGPU-Sim只能模拟非常老的Fermi架构，所以这里的分析并不保证是完全正确的。