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__ void
vectorAdd(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
 */
int
main(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

时间统计结果如下：

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

坑

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

可以看出，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架构，所以这里的分析并不保证是完全正确的。