GPU高性能计算CUDA编程：PixCopy()复制图像的GPU核函数

GPU高性能计算CUDA编程：PixCopy()复制图像的GPU核函数

声明：本文不做商用

以下示例代码所示为如何完整地复制一幅图像。将示例代码中的 PixCopy() 函数与前面文章示例代码中的其他两个函数进行比较，可以看到主要区别在于PixCopy() 中使用了一维像素索引。换句话说，PixCopy() 没有计算行号和列号。PixCopy() 中的每个线程只复制一个字节。这消除了每行末尾的无用线程，但在图像的末尾还是会有无用线程。如果启动程序时，选择了C(copy) 选项，将以每块256个线程启动它。但是，由于复制过程不是按照二维索引，而是按照线性顺序进行的，因此需要使用另一种块的索引方式来处理。做一个简单的数学变换:图像共有127712256字节。每个线程块有256个线程，我们总共需要启动127712256/256]=498876个线程块。没用无用线程，运气不错。但是，如果文件再多20个字节，最后一个块的256个线程中会有236个被浪费掉。这就是为什么在核函数中仍必须输入以下if语句来检查这种情况:

if (MYgtid > FS) return;            // 超出内存分配范围

示例代码中的if语句与之前的示例代码中的if语句非常类似，用于检查“是否是个无用线程”，如果是则不做任何事情。该行代码对性能的影响与前两个核函数类似:尽管这个条件判断式只会在很少的线程中为真，但是所有线程仍然必须在复制每个字节时执行它。我们可以对它进行改进，所有这些改进方法将在后续的章节中给出。就目前而言，你需要意识到的是核函数PixCopy()中的if语句对性能的影响比其他两个核函数中的严重很多。 PixCopy()核函数的并行粒度更细，因为它只复制单个字节。正因为如此，PixCopy()中只有6行C代码，其中之一就是if语句。虽然“代码行数”并不等同于“GPU 核心需要执行的指令的周期数”，但仍然可以大致反映问题的严重程度。参见【0voice C++】

// Kernel that copies an image from one part of the
// GPU memory (ImgSrc) to another (ImgDst)
__global__
void PixCopy(uch *ImgDst, uch *ImgSrc, ui FS)
{
	ui ThrPerBlk = blockDim.x;
	ui MYbid = blockIdx.x;
	ui MYtid = threadIdx.x;
	ui MYgtid = ThrPerBlk * MYbid + MYtid;

	if (MYgtid > FS) return;				// outside the allocated memory
	ImgDst[MYgtid] = ImgSrc[MYgtid];
}