GPU高性能计算CUDA编程：线程与GPU核函数的执行

GPU高性能计算CUDA编程：线程与GPU核函数的执行

声明：本文不做商用

现在我们想知道的是，Vflp()函数如何计算它是谁以及它该负责处理哪个部分。我们先来回忆一下CPU代码:

void *MTFlipV(void* tid)
{
	...
	long ts = *((int *) tid);       	// My thread ID is stored here
	ts *= ip.Hbytes/NumThreads;			// start index
	long te = ts+ip.Hbytes/NumThreads-1; 	// end index

	for(col=ts; col<=te; col+=3)
	{
		row=0;
		while(row<ip.Vpixels/2)
		{
			pix.B = TheImage[row][col];
			pix.G = TheImage[row][col+1];
			pix.R = TheImage[row][col+2];
			...
}

此处的变量ts和te分别是图像中的起始和结束行号。通过两层循环实现垂直翻转，层是列扫描，另一层是行扫描。现在，将该代码与代码6.7中的Vflip()函数相比较:

ui ThrPerBlk = blockDim.x;
ui MYbid = blockIdx.x;
ui MYtid = threadIdx.x;
ui MYgtid = ThrPerBlk * MYbid + MYtid;

ui BlkPerRow = (Hpixels + ThrPerBlk - 1) / ThrPerBlk;  // ceil
ui RowBytes = (Hpixels * 3 + 3) & (~3);
...

可以看到CPU和GPU代码之间既有很多相似之处，也存在不小的差异。由于线程块和块内线程数的出现，GPU函数中的任务分配与CPU中的完全不同。所以，尽管GPU也需要计算一些索引值，但与CPU函数中的实现完全不同。GPU首先想知道每个块有多少个线程。答案在一个名为blockDim.x的GPU系统变量中。在本例中，这个值是256，因为我们指定了每个线程块启动256个线程(Vflip<<<…，256 >>>)。因此，每个块包含 256个线程，线程ID分别为0…255。线程的ID值在threadIDx.x中。我们还需要知道在166656个块中，每个线程块的ID是多少，这个值在另一个名为blockIdx.x的GPU系统变量中令人惊讶的是，本例并不关心线程块的总数是多少(166656)，而有的程序可能需要知道。

线程块的ID和线程ID被分别存放在变量bid和tid中。使用这两者的组合可以计算出一个全局线程ID(gtid)。gtid是每个启动的GPU线程的唯一ID(共有166656x256~41M个线程)，也就是线程ID的线性化。这与我们根据公式6.2线性化磁盘上的像素存储单元非常相似。然而，因为每行中都存在无用的线程，此时GPU线程的线性ID与像素存储单元的线性地址之间的对应关系并没那么容易地获得。接下来，需要计算每行的线程块数(BIkPerRow)，在本例中为31。最后，由于水平像素的个数(7918)作为第三个参数被传递给该函数，因此可以计算每行图像的总字节数(3x7918=23754字节)，并确定每个像素的字节索引。参见【0voice C++】
完成这些计算后，核函数计算需要复制的像素的行索引和列索引，如下所示:

ui MYrow = MYbid / BlkPerRow;
ui MYcol = MYgtid - MYrow*BlkPerRow*ThrPerBlk;
if (MYcol >= Hpixels) return;			// col out of range
ui MYmirrorrow = Vpixels - 1 - MYrow;
ui MYsrcOffset = MYrow       * RowBytes;
ui MYdstOffset = MYmirrorrow * RowBytes;
ui MYsrcIndex = MYsrcOffset + 3 * MYcol;
ui MYdstIndex = MYdstOffset + 3 * MYcol;

这些代码执行后，源像素的存储地址位于MYsrcIndex中，目的存储地址位于MYdstIndex中。由于每个像素都包含从该地址开始的3个字节(RGB)，因此核函数会从该地址开始复制3个连续的字节，如下所示:

// swap pixels RGB   @MYcol , @MYmirrorcol
ImgDst[MYdstIndex] = ImgSrc[MYsrcIndex];
ImgDst[MYdstIndex + 1] = ImgSrc[MYsrcIndex + 1];
ImgDst[MYdstIndex + 2] = ImgSrc[MYsrcIndex + 2];

将它与CPU代码2.7进行比较。在CPU中只能启动4~8个线程，而我们刚刚目睹了41M的巨量线程。GPU核函数的一个惊人之处是for循环消失了!换句话说，CPU函数必须显式地对列和行进行循环扫描，而GPU不需要。CPU函数中循环的目的是用row和col变量形成的某种顺序依次访问所有像素。但是，在GPU核函数中，可以通过tid和bid来实现该功能，因为我们知道像素坐标与变量tid和bid之间的准确关系。