CUDA进阶补充篇：详析各种CUDA函数计时函数 |
您所在的位置：网站首页 › c语言clock函数头文件 › CUDA进阶补充篇：详析各种CUDA函数计时函数 |
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(); 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 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 (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的浮动cudaEvent328.312744ms上下3ms左右的浮动clock()330ms很稳定chrono324ms上下3ms左右的浮动从实验结果可以看出，后三种计时方式都是比较稳定，可以放心使用。 坑细心的人可能会发现，我在代码最前面加了一个空的warmup函数。这个在精确统计时间是非常重要的！！！因为GPU第一次被调用时会消耗不定的时间来预热。 如果把预热那行注释掉，得到的计时结果如下： 计时方式Time(ms)gettimeofday()535.159826mscudaEvent346.573151msclock()440.000000mschrono470ms可以看出，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架构，所以这里的分析并不保证是完全正确的。 
|
【本文地址】
今日新闻 |
推荐新闻 |