cudaEvent_t start, stop;
ErrorCheck(cudaEventCreate(&start), __FILE__, __LINE__);
ErrorCheck(cudaEventCreate(&stop), __FILE__, __LINE__);
ErrorCheck(cudaEventRecord(start), __FILE__, __LINE__);
cudaEventQuery(start); //此处不可用错误检测函数
/************************************************************
需要记时间的代码
************************************************************/
ErrorCheck(cudaEventRecord(stop), __FILE__, __LINE__);
ErrorCheck(cudaEventSynchronize(stop), __FILE__, __LINE__);
float elapsed_time;
ErrorCheck(cudaEventElapsedTime(&elapsed_time, start, stop), __FILE__, __LINE__);
printf("Time = %g ms.\n", elapsed_time);
ErrorCheck(cudaEventDestroy(start), __FILE__, __LINE__);
ErrorCheck(cudaEventDestroy(stop), __FILE__, __LINE__);
代码解析:
第1行cudaEvent_t start, stop:定义两个CUDA事件类型(cudaEvent_t)的变量;
第2、3行cudaEventCreate函数初始化定义的cudaEvent_t变量;
第4行通过cudaEventRecord函数,在需要记时的代码块之前记录代表时间开始的事件;
第5行cudaEventQuery函数在TCC驱动模式的GPU下可省略,但在处于WDDM驱动模式的GPU必须保留,因此,我们就一直保留这句函数即可。注意:cudaEventQuery函数不可使用错误检测函数;
第8行是需要记时的代码块;
第11行在需要记时的代码块之后记录代表时间结束的事件;
第12行cudaEventSynchronize函数作用是让主机等待事件stop被记录完毕;
第13~15行cudaEventElapsedTime函数的调用作用是计算cudaEvent_t变量start和stop时间差,记录在float变量elapsed_time中,并输出打印到屏幕上;
第17、18行调用cudaEventDestroy函数销毁start和stop这两个类型为cudaEvent_t的CUDA事件。
核函数计时示例:
此代码计算运行核函数10次的平均时间,核函数实际运行11次,由于第一次调用核函数,往往会花费更多的时间,如果将第一次记录进去,可能导致记录的时间不准确,因此忽略第一次调用核函数的时间,取10次平均值
#include <stdio.h>
#include "../tools/common.cuh"
#define NUM_REPEATS 10
__device__ float add(const float x, const float y)
{
return x + y;
}
__global__ void addFrom