GPU高性能计算CUDA编程：Vflip()垂直翻转的GPU核函数

GPU高性能计算CUDA编程：Vflip()垂直翻转的GPU核函数

声明：本文不做商用

以下示例代码所示为GPU核函数Vflp()。该核函数是每个线程都需要执行的功能。正如我们在前面计算的那样，假设我们以维度 Vflip <<< 166656,256 >>>(…) 启动该核函数，将会启动4000多万个线程运行该核函数。如果以更高的维度启动它，启动的线程也会更多。到目前为止，我们看到的每一行代码都是CPU代码。代码6.7是我们看到的第一个GPU代码。代码6.7和其他代码都是用C语言编写的，但我们关心的是代码6.7经过编译后得到的CPU指令。尽管本章之前的所有代码都将被编译为x64指令—Intel 64位指令集架构(ISA)，代码6.7将被编译为Nvidia的GPU指令集架构，称为并行线程执行(PTX)。就像英特尔和AMD在新一代的CPU中都会扩展指令集架构一样，Nvidia在每一代GPU中也是如此。

// Kernel that flips the given image vertically
// each thread only flips a single pixel (R,G,B)
__global__
void Vflip(uch *ImgDst, uch *ImgSrc, ui Hpixels, ui Vpixels)
{
	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);
	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;

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

CPU指令集架构和GPU指令集架构之间的一大区别是，采用x64编译的CPU指令集架构的输出是在运行时依次执行的x86指令(即完全编译)，但PTX实际上只是一种中间表示(IR)。这意味着它是“半编译的”，就像Java的字节码。参见【0voice C++】Nvidia运行时引擎输入 PTX 指令，并在运行时进一步半编译它，并将完全编译后的指令提供给GPU核心。在Windows中，实现这种“分次编译”的“Nvidia代码”都在名为cudart(CUDA运行时)的动态链接库(DLL)中。有两种配置:在较新的x64操作系统中称为cudart64，在传统的32位操作系统中称为cudart32。当然不应该再继续使用后者，因为所有新一代的Nvidia GPU都需要64位操作系统才能有效地使用。