CUDA事件--时间性能测量

我们固然可以选择通过CPU上的计时器来测量CUDA应用的性能。但这种方式并不会给出很精确的结果。CPU上的时间测量还需要CPU端具有高精度的定时器。很多时候，在GPU内核异步运行的同时，主机上正在执行着计算，所以CPU端的计时器可能无法给出正确的内核执行时间。
CUDA事件等于是在你的CUDA应用运行的特定时刻被记录的时间戳。通过使用CUDA事件API，由GPU来记录这个时间戳，因此消除了CPU端的计时器测量性能时所会受到的影响。使用CUDA测量时间需要两个步骤:创建事件和记录事件。我们将会记录两个事件:一个记录我们的代码运行开始的时刻，另一个记录我们的代码运行结束的时刻。接着我们会通过两个事件记录的时刻相减来计算出运行时间，这将会给出我们的代码整体性能的参考信息。
在CUDA代码中可以使用CUDA事件API导入以下代码来测量性能:

#include "stdio.h"
#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>

//Defining number of elements in Array
#define N	50000

//Defining Kernel function for vector addition
__global__ void gpuAdd(int *d_a, int *d_b, int *d_c) 
{
	//Getting Thread index of current kernel

	int tid = threadIdx.x + blockIdx.x * blockDim.x;
	while (tid < N)
	{
		d_c[tid] = d_a[tid] + d_b[tid];
		tid += blockDim.x * gridDim.x;
	}
}

int main() 
{
	//Defining host arrays
	int h_a[N], h_b[N], h_c[N];
	//Defining device pointers
	int *d_a, *d_b, *d_c;
	cudaEvent_t e_start, e_stop;
	cudaEventCreate(&e_start);
	cudaEventCreate(&e_stop);
	cudaEventRecord(e_start, 0);
	// allocate the memory
	cudaMalloc((void**)&d_a, N * sizeof(int));
	cudaMalloc((void**)&d_b, N * sizeof(int));
	cudaMalloc((void**)&d_c, N * sizeof(int));
	//Initializing Arrays
	for (int i = 0; i < N; i++)
	{
		h_a[i] = 2 * i*i;
		h_b[i] = i;
	}
	// Copy input arrays from host to device memory
	cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice);
	//Calling kernels passing device pointers as parameters
	gpuAdd << <512, 512 >> >(d_a, d_b, d_c);
	//Copy result back to host memory from device memory
	cudaMemcpy(h_c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);
	cudaDeviceSynchronize();
	cudaEventRecord(e_stop, 0);
	cudaEventSynchronize(e_stop);
	float elapsedTime;
	cudaEventElapsedTime(&elapsedTime, e_start, e_stop);
	printf("Time to add %d numbers: %3.1f ms\n",N, elapsedTime);

	int Correct = 1;
	printf("Vector addition on GPU \n");
	//Printing result on console
	for (int i = 0; i < N; i++) 
	{
		if ((h_a[i] + h_b[i] != h_c[i]))
		{
			Correct = 0;
		}
	}
	if (Correct == 1)
	{
		printf("GPU has computed Sum Correctly\n");
	}
	else
	{
		printf("There is an Error in GPU Computation\n");
	}
	//Free up memory
	cudaFree(d_a);
	cudaFree(d_b);
	cudaFree(d_c);
	return 0;
}

我们将建立两个事件对象，e_start和 e_stop，用来测量一段代码的开始和结束区间(的执行性能)。我们通过cudaEvent_t类型来定义相关的变量，然后通过cudaEventCreate建立事件，保存在刚才定义好的变量上。我们通过将cudaEvent_t类型的变量地址传递给cudaEventCreate来得到对应的事件对象，而非 cudaEventCreate通过返回值直接返回它。在代码的开始，我们通过调用cudaEventRecord，用刚才建立的第一个事件对象e_start记录一次GPU时间戳。第二个参数我们指定为0，这里代表了CUDA流。
在开头的代码发出了时刻记录命令后，你可以继续写你的其他GPU代码。当这些代码结束后，我们会再记录一次时刻，这次是记录在事件对象e_stop里。这次通过cudaEventRecord(e_stop,0)这行代码进行记录。一旦我们记录了开始的时刻和结束的时刻，两者的差值将会给出用来衡量你的代码性能的时间。但直接计算两个事件对象的差值时间的话，还存在问题。
如同我们之前讨论过的那样，CUDA C(代码）的执行可能是异步的。当GPU正在执行内核的时候，没等它执行完毕，CPU可能就已经继续执行后续的代码行了。同样的，类似内核，事件的记录命令从它在CPU上发出到它实际在GPU上执行也是异步的。所以，CPU和 GPU不进行同步，直接认为发出了事件的记录命令后就可以立刻使用测量值的想法，可能会给出错误的结果。
因为cudaEventRecord()作用的事件对象将会在它之前所有发出GPU命令都执行完毕后才会记录一个时刻值。我们必须等待事件对象e_stop实际的已经完成了时刻值记录，也就是它之前的所有GPU工作都完成后才应该试图读取使用它。这样，对于异步的cudaEventRecord()操作,等它完成后再访问才是安全的。
我们通过CUDA提供的cudaEventElapsedTime这个API函数来计算两个时间戳之间的差值。该API函数具有3个参数，第一个参数是用来返回时间差结果的，第二个参数是起始时刻的事件，第三个参数则是结束时刻的事件。