Kernel函数是执行在GPU上的函数，并且是异步执行，所以统计其执行时间与CPU函数有不同之处。通常有三种方式可以用于kernel函数的执行时间统计：
-采用CPU Timer
-采用GPU Timer，也就是CUDA提供的API
-采用NVIDIA Profiler，一种是命令行工具nvprof，另一种是图形化界面的Visual profiler
下面一一介绍。
CPU Timer
CPU Timer可以通过调用系统函数实现，比如在Linux系统上：
#include <sys/time.h>
double cpuSecond()
{
struct timeval tp;
gettimeofday(&tp,NULL);
return ((double)tp.tv_sec + (double)tp.tv_usec*1.e-6);
}
你可以使用cpuSecond()采用如下的方式来统计kernel函数的执行时间：
double iStart = cpuSecond();
kernel_name<<<grid, block>>>(argument list);
cudaDeviceSynchronize(); //wait for all GPU threads to complete
double iElaps = cpuSecond() - iStart;
由于kernel函数是异步执行的，所以在kernel函数后面统计时间戳前面加上了CUDADeviceSynchronize()以等待kerne函数的完成。为了保险起见，在获取iStart之前也应该调用一下同步函数，这是因为前面可能会有没有执行完成的kernel函数或者是调用了异步的CUDA API。这将导致错误的时间统计。
GPU Timer
CUDA Event 的 API也具有时间统计的功能。cudaEventElapsedTime () 函数可以计算两个事件（event）的间隔时间。返回的结果是以ms为单位，并且其时间分辨率大约为0.5us。其具体的使用方法如下：
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
...kernel functions...
cudaEventRecord(stop);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&time, start, stop);
cudaEventDestroy(start);
cudaEventDestroy(stop);
首先创建两个时间start与stop，然后采用cudaEventRecord记录一个事件，默认情况下是在default Stream （sstream=0）上。在kernel函数执行之后，再重新记录一个事件stop。cudaEventSynchronize()的功能是等待时间stop完成。最后采用cudaEventElapsedTime计算两个时间start与stop之间的时间间隔。
nvprof
采用nvprof可以直接打印每个kernel函数被调用的次数以及执行的时间，另外它还统计了每个CUDA API的调用情况。下面是一个示例：
==9703== Profiling application: ./a.out
==9703== Profiling result:
Type Time(%) Time Calls Avg Min Max Name
GPU activities: 100.00% 111.00ms 3 37.000ms 36.950ms 37.094ms testKernel(int*)
0.00% 2.6560us 3 885ns 736ns 1.1520us [CUDA memset]
API calls: 78.69% 418.81ms 2 209.41ms 271.66us 418.54ms cudaFree
13.88% 73.896ms 2 36.948ms 36.943ms 36.953ms cudaDeviceSynchronize
6.97% 37.086ms 1 37.086ms 37.086ms 37.086ms cudaEventSynchronize
0.16% 868.42us 1 868.42us 868.42us 868.42us cuDeviceTotalMem
0.14% 723.05us 96 7.5310us 183ns 329.77us cuDeviceGetAttribute
0.05% 273.57us 1 273.57us 273.57us 273.57us cudaMalloc
0.04% 221.74us 3 73.911us 42.784us 132.13us cudaMemset
0.03% 141.77us 3 47.256us 33.979us 67.575us cudaLaunchKernel
三种统计方式的比较
很显然，不论是CPU Timer还是GPU Timer我们都需要修改源代码，所以nvprof的优势就是使用更加简单便捷。不需要修改任何代码就可以获得kernel函数的执行时间。当然在源代码中使用Timer更加的灵活，我们可以仅打印我们所关心的代码段的执行时间。所以具体应该使用哪一个，就要看自己的需求了。下面的一段代码展示了CPU Timer与GPU Timer的使用。
#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#include <sys/time.h>
//in ms
double cpuMillisecond()
{
struct timeval tp;
gettimeofday(&tp,NULL);
return ((double)tp.tv_sec * 1000 + (double)tp.tv_usec*1.e-3);
}
__global__ void testKernel(int *a)
{
for(size_t i = 0; i < 1024 * 1024; i++)
{
a[0] += a[i];
}
}
int main() {
int *d_a;
int N = 2 << 20; //1M
cudaFree(0);
cudaMalloc(&d_a, N * sizeof(int));
cudaMemset(d_a, 0, N * sizeof(int));
//warp up
testKernel<<<1,1>>>(d_a);
cudaMemset(d_a, 0, N * sizeof(int));
//need to be added befor cpu timer
//you can guess what will happen when omitting this synchronization
cudaDeviceSynchronize();
//cpuTimer
double iStart = cpuMillisecond();
testKernel<<<1,1>>>(d_a);
cudaDeviceSynchronize();
double tcpu = cpuMillisecond() - iStart;
printf("cpuTimer: %.3f ms\n", tcpu);
//gpuTimer
float tgpu = 0.f;
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaMemset(d_a, 0, N * sizeof(int));
cudaEventRecord(start);
testKernel<<<1,1>>>(d_a);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&tgpu, start, stop);
printf("gpuTimer: %.3f ms\n", tgpu);
cudaEventDestroy(start);
cudaEventDestroy(stop);
cudaFree(d_a);
}
在一台P5000的显卡上执行的结果是：
cpuTimer: 37.241 ms
gpuTimer: 36.999 ms
采用nvprof的统计结果是：
Type Time(%) Time Calls Avg Min Max Name
GPU activities: 100.00% 110.79ms 3 36.931ms 36.773ms 37.081ms testKernel(int*)
