GPU高性能计算CUDA编程：PixCopy()核函数中线程的执行情况-CSDN博客

本文链接：https://blog.csdn.net/m0_51165837/article/details/142026063

GPU高性能计算CUDA编程：PixCopy()核函数中线程的执行情况

声明：本文不做商用

PixCopy()核函数的每个线程复制一个“字节”而非3个字节，因此当用户在命令行中输人C选项(复制)时，imflipG.cu必须启动三倍的线程。在此例中，121MB的图像将启动121M个线程，对应于gridDim.x=498834。每个块的ID范围为 blockIDx.x-0…498 833。我们来看看线程执行的各个步骤。
1.我是谁?下面是代码6.9中显示的 PixCopy()核函数。与其他两个核函数一样，也需要计算全局索引。

// 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;

2.我的任务是什么?该步骤首先进行范围检查，只是稍有不同。当全局线程ID大于图像的大小时，该线程是无用的。只有当图像的大小不能被每块线程数整除时，才会发生这种情况。例如，astronaut.bmp为7918x5376x3=127701504字节，如果我们使用每块1024个线程，则需要124708.5个线程块。因此，将启动124709个块，浪费了半个块(即512个线程)。尽管与121MB相比，这是一个微不足道的数字，但我们仍然必须用该行代码来检查全局线程ID是否超出了范围。参见【0voice C++】

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

3.执行……因为每个线程只复制一个像素，所以全局线程ID(MYgtid)与源和目标GPU内存地址直接相关。实际上，不仅如此:它与数组索引完全一样，不需要计算任何其他索引。因此，执行实际工作只需要一行代码。

	ImgDst[MYgtid] = ImgSrc[MYgtid];
}

这些代码所做的是将一个字节从GPU全局内存的一个地方(由* ImgSrc指向)复制到GPU全局内存的另一地方(由* ImgDst指向)。核函数 PixCopy()比其他两个更简单，这也许会让你挠头。