GPU高性能计算CUDA编程：GPU端的时间戳

GPU高性能计算CUDA编程：GPU端的时间戳

声明：本文不做商用

当我们进入CUDA代码执行的细节后，你很快就会知道线程块是什么意思。现在，让我们把注意力集中在CPU-GPU的交互上。你可以在CUDA代码的开头部分查询GPU，并根据查询结果用不同的参数执行CUDA代码，以获得最佳效率。

GPU端的一个函数(如Vflip())被称为一个核函数。代码6.3显示了CPU如何启动一个GPU的核函数，并将GPU用作协处理器。这是代码中需要理解的最重要的部分，所以我将详细解释它的每一部分。一旦你理解了这部分代码，下面的章节就会毫无困难。首先，与CPU代码的计时操作非常相似，我们也希望统计完成CPU→GPU的传输和GPU→CPU的传输需要花费多长时间。另外，我们希望统计在这两个操作之间执行GPU代码需要多长时间。当查看代码6.3时，我们发现以下几行代码实现了GPU代码的计时:

cudaEvent_t time1, time2，time3, time4;
...
cudaEventCreate(&time1);
cudaEventCreate(&time2);
cudaEventCreate(&time3);
cudaEventCreate(&time3);
...
// 将数据从CPU复制到GPU。
cudaEventRecord(time1，0);  // 开始进行GPU传输的时间戳
...
// 将数据从CPU传输到GPU
cudaEventRecord(time2，0);  // CPU到GPU的传输结束后的时间戳
...
//执行GPU代码
cudaEventRecord(time3，0);  // GPU代码执行完成后的时间戳

变量 time1、time2、time3 和 time4 都是 CPU端变量，存储了CPU 和 GPU 之间进行数据传输时的时间戳，以及GPU代码在设备端执行过程中的时间戳。上面代码中一个奇怪的地方是，我们只用Nvidia的API为与GPU相关的事件添加时间戳。任何涉及GPU的东西都必须用Nvidia的API来得到时间戳，此处是cudaEventRecord()。参见【0voice C++】但为什么?为什么我们不能简单地使用表现优异的gettimeofday()函数?我们在CPU代码清单中看到过它。

uch *TheImg, *CopyImg;					// Where images are stored in CPU
uch *GPUImg, *GPUCopyImg, *GPUResult;	// Where images are stored in GPU
...
#define	IPHB		ip.Hbytes
#define	IPH			ip.Hpixels
#define	IPV			ip.Vpixels
#define	IMAGESIZE	(IPHB*IPV)
...
int main(int argc, char **argv)
{
	cudaError_t cudaStatus, cudaStatus2;
	cudaEvent_t time1, time2, time3, time4;
	
	ui BlkPerRow, ThrPerBlk = 256, NumBlocks, GPUDataTransfer;
	...

	cudaEventCreate(&time1);
	cudaEventCreate(&time2);
	cudaEventCreate(&time3);
	cudaEventCreate(&time4);

	cudaEventRecord(time1, 0);		// Time stamp at the start of the GPU transfer
	// Allocate GPU buffer for the input and output images
	cudaStatus = cudaMalloc((void**)&GPUImg, IMAGESIZE);
	cudaStatus2 = cudaMalloc((void**)&GPUCopyImg, IMAGESIZE);
	if ((cudaStatus != cudaSuccess) || (cudaStatus2 != cudaSuccess)){
		fprintf(stderr, "cudaMalloc failed! Can't allocate GPU memory");
		exit(EXIT_FAILURE);
	}

	// Copy input vectors from host memory to GPU buffers.
	cudaStatus = cudaMemcpy(GPUImg, TheImg, IMAGESIZE, cudaMemcpyHostToDevice);
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaMemcpy  CPU to GPU  failed!");
		exit(EXIT_FAILURE);
	}

	cudaEventRecord(time2, 0);		// Time stamp after the CPU --> GPU tfr is done
	
	//dim3 dimBlock(ThrPerBlk);
	//dim3 dimGrid(ip.Hpixels*BlkPerRow);

	BlkPerRow = (IPH + ThrPerBlk -1 ) / ThrPerBlk;
	NumBlocks = IPV*BlkPerRow; 
	switch (Flip){
		case 'H': Hflip <<< NumBlocks, ThrPerBlk >>> (GPUCopyImg, GPUImg, IPH);
				  GPUResult = GPUCopyImg;
				  GPUDataTransfer = 2*IMAGESIZE;
				  break;
		case 'V': Vflip <<< NumBlocks, ThrPerBlk >>> (GPUCopyImg, GPUImg, IPH, IPV);
				  GPUResult = GPUCopyImg;
				  GPUDataTransfer = 2*IMAGESIZE;
				  break;
		case 'T': Hflip <<< NumBlocks, ThrPerBlk >>> (GPUCopyImg, GPUImg, IPH);
				  Vflip <<< NumBlocks, ThrPerBlk >>> (GPUImg, GPUCopyImg, IPH, IPV);
				  GPUResult = GPUImg;
				  GPUDataTransfer = 4*IMAGESIZE;
				  break;
		case 'C': NumBlocks = (IMAGESIZE+ThrPerBlk-1) / ThrPerBlk;
				  PixCopy <<< NumBlocks, ThrPerBlk >>> (GPUCopyImg, GPUImg, IMAGESIZE);
				  GPUResult = GPUCopyImg;
				  GPUDataTransfer = 2*IMAGESIZE;
				  break;
	}
	...
}

我们完全依赖Nvidia的API(来自月球的人)来计算与GPU端有关的任何事情。如果这样做，不妨也让他们对太空旅行进行计时，无论是出发还是回家。数据传输的开始和结束或执行一个GPU核函数的开始和结束都被记录为某种事件，这使我们能够用 Nvidia事件计时API对它们进行计时，例如cudaEventRecord()。要在API中使用一个事件，必须先用cudaEventCreate()创建该事件。由于事件记录机制内嵌在Nvidia的API中，我们可以很方便地用它们来统计执行GPU核函数或在CPU与GPU之间进行传输所需要的时间，这和CPU代码非常相似。
在代码6.3中，在代码刚开始的地方我们用time1系加时间，并在CPU到GPU传输完成时将时间戳保存在time2中。类似地，time3是GPU代码执行完成时的时间戳而time4是结果完全到达CPU端时的时间戳。相邻两个时间戳之间的差值告诉我们完成这些事件中的每一个需要多长时间。毫不奇怪，差值也必须通过调用CUDAAPI库中的cudaEventElapsedTime()API来计算(如代码6.4所示)，因为时间戳的存储格式是Nvidia API的一部分，并不是普通的变量。