我们固然可以选择通过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个参数,第一个参数是用来返回时间差结果的,第二个参数是起始时刻的事件,第三个参数则是结束时刻的事件。