【cuda编程】CUDA中计算程序耗时

在我们日常使用CUDA优化程序时，为了方便分析，经常会统计自己写的程序耗时情况。一般情况下，可以使用C/C++的CPU方式，另外一种则选择则是使用CUDA的事件方式。

1. CPU计算方式

配合cuda的同步方法，加上cpu的计时方式，可以实现GPU端的耗时。

    double cpuSecond()
    {
        struct timeval tp;
        gettimeofday(&tp, NULL);
        return tp.tv_sec * 1000000 + tp.tv_usec; //微秒
    }
        double iStart = cpuSecond();
        
        kernel_function << <1, 10>> > ();
        cudaDeviceSynchronize(); // synchronzie
        
        double iElaps = cpuSecond() - iStart;
        std::cout << "耗时： " << iElaps << "us"<< std::endl;

2. CUDA计算方式

本节介绍一下CUDA统计耗时的方式cudaEventElapsedTime，它的功能为记录两次事件之间相差的时间，单位为毫秒，精度为0.5微妙。

主要涉及以下几个函数：

//创建事件对象
cudaError_t cudaEventCreate( cudaEvent_t* event );
//记录事件
cudaError_t cudaEventRecord( cudaEvent_t event，CUstream stream );
//计算两次事件之间相差的时间
cudaError_t cudaEventElapsedTime( float* time，cudaEvent_t start，cudaEvent_t end );
//销毁事件对象
cudaError_t cudaEventDestroy( cudaEvent_t event );

注意：CUDA事件是直接在GPU上实现的，因此它们不适用于对同时包含设备代码和主机代码的混合代码计时。也就是说，如果你试图通过CUDA事件对Kernel和设备内存复制之外的代码进行计时，将得到不可靠的结果。

代码举例：

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

#define BASE_CUDA_CHECK(condition) { GPUAssert((condition), __FILE__, __LINE__); }

inline void GPUAssert(cudaError_t code, const char *file, int line, bool abort = true) {
    if (code != cudaSuccess) {
        fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
        if (abort) {
            exit(code);
        }
    }
}

//产生大量0-9之间的随机数
void GenerateNumbers(int *number, const int size)
{
    for (int i = 0; i < size; i++)
    {
        number[i] = rand() % 10;
    }
}

__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = threadIdx.x;
    c[i] = a[i] + b[i];
}

int main()
{
	// const int arraySize = 104857;
	const int arraySize = 1000;
	int host_a[arraySize];
	int host_b[arraySize];
	//生成数据
	GenerateNumbers(host_a, arraySize);
	GenerateNumbers(host_b, arraySize);
	int host_c[arraySize] = {0};
	int *dev_a = 0;
    int *dev_b = 0;
	int *dev_c = 0;

	BASE_CUDA_CHECK(cudaMalloc((void**)&dev_a, arraySize * sizeof(int)));
	BASE_CUDA_CHECK(cudaMalloc((void**)&dev_b, arraySize * sizeof(int)));
	BASE_CUDA_CHECK(cudaMalloc((void**)&dev_c, arraySize * sizeof(int)));
	BASE_CUDA_CHECK(cudaMemcpy(dev_a, host_a, arraySize * sizeof(int), cudaMemcpyHostToDevice));
	BASE_CUDA_CHECK(cudaMemcpy(dev_b, host_b, arraySize * sizeof(int), cudaMemcpyHostToDevice));

	//一些cuda事件定义
	cudaEvent_t start, stop;
	float time;
	cudaEventCreate(&start); //创建计时
	cudaEventCreate(&stop); //创建计时
	cudaEventRecord(start, 0); // 开始计时

	// cuda调用
	addKernel<<<1, arraySize>>>(dev_c, dev_a, dev_b);
	BASE_CUDA_CHECK(cudaDeviceSynchronize());

	cudaEventRecord(stop, 0);
	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&time, start, stop);
	cudaEventDestroy(start); //销毁
	cudaEventDestroy(stop);
	std::cout << "耗时: " << time << "ms" << std::endl;

	//取回数据
	BASE_CUDA_CHECK(cudaMemcpy(host_c, dev_c, arraySize * sizeof(int), cudaMemcpyDeviceToHost));

	std::cout<<std::endl;
	return 0;
}

输出：

耗时: 0.269248ms