【CUDA编程】CPU计时与GPU计时

使用CUDA进行编程，主要目的就是时间上加速。为此，如何计时必不可少。在CUDA中，我们可以使用CPU计时函数和GPU计时函数。对于CPU计时，我们在之前的文章(精确系统计时：秒、毫秒、微秒)中已经介绍在一般的C/C++编程中的计时方法。下面我们介绍在CUDA中如何计时：

---

CPU计时

CUDA中的核函数是异步执行的，即调用核函数后（而非等待其运行结束）就继续执行后面的语句。因此，使用CPU计时的时候，我们需要加上同步函数，这样才能得到核函数的运行时间，否则就是调用时间。下面给出一个简单的实例，由于代码简单，这里不再过多说明：

#include"cuda_runtime.h"
#include"device_launch_parameters.h"
#include<iostream>
#include<random>
#include<ctime>


using namespace std;

__global__ void kernel_function()
{
	printf("Hello From GPU\n");

}


//Timing using CPU 
int main()
{
	float esp_time_cpu;
	clock_t start_cpu, stop_cpu;

	kernel_function << <1, 10>> > ();// warming up

	start_cpu = clock();// start timing
	kernel_function<<<1,10>>> ();
	cudaDeviceSynchronize(); // synchronzie
	stop_cpu = clock();// end timing

	esp_time_cpu = (float)(stop_cpu - start_cpu) / CLOCKS_PER_SEC;

	printf("The time by host:\t%f(ms)\n", esp_time_cpu);


	cudaDeviceReset();
	return 0;
}

需要注意的是：在计时前最好先warming up一下，即先把要计时的函数运行一遍。

---

GPU计时

这里我们可以使用CUDA提供的事件管理API来实现计时，具体可以参考NVIDIA官方的文档。具体实例如下：

#include"cuda_runtime.h"
#include"device_launch_parameters.h"
#include<iostream>
#include<random>
#include<ctime>


using namespace std;


__global__ void kernel_function()
{
	printf("Hello From GPU\n");

}


 Timing using GPU
int main()
{
	cudaEvent_t start, stop;
	float esp_time_gpu;
	cudaEventCreate(&start);
	cudaEventCreate(&stop);

	kernel_function << <1, 10 >> > ();// warming up

	cudaEventRecord(start, 0);// start

	kernel_function << <1, 10 >> > ();

	cudaEventRecord(stop, 0);// stop

	cudaEventSynchronize(stop);


	cudaEventElapsedTime(&esp_time_gpu, start, stop);
	printf("Time for the kernel: %f ms\n", esp_time_gpu);


	cudaDeviceReset();
	return 0;
}

---

CPU和GPU计时的结果分别如下：

注意：上述的结果不同，是因为有误差，我们可以通过多运行几次取平均值。

另外，上述CPU和GPU计时上的主要区别在于是否需要同步：CPU计时需要在核函数后调用同步函数。为此，当需要同步时，我们可以使用CPU计时；当不能同步时，我们使用GPU计时。当然，GPU计时也能用于需要同步的场景，此时的同步函数cudaDeviceSynchronize()需要放置在计时模块的后面。