GPU高性能计算CUDA编程：Hflip()水平翻转的GPU核函数

最新推荐文章于 2024-10-05 12:07:50 发布

阿里猫的编程小辈

最新推荐文章于 2024-10-05 12:07:50 发布

阅读量315

点赞数 5

分类专栏： 007 文章标签： c++ 开发语言人工智能 gpu算力

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

版权

007 专栏收录该内容

37 篇文章 0 订阅

订阅专栏

GPU高性能计算CUDA编程：Hflip()水平翻转的GPU核函数

声明：本文不做商用

以下示例代码显示了在水平方向上翻转图像的GPU核函数。尽管计算索引的方式不同，代码6.8与代码6.7几乎相同，当然，代码6.7完成的是像素的垂直翻转。

// Kernel that flips the given image horizontally
// each thread only flips a single pixel (R,G,B)
__global__
void Hflip(uch *ImgDst, uch *ImgSrc, ui Hpixels)
{
	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 MYmirrorcol = Hpixels - 1 - MYcol;
	ui MYoffset = MYrow * RowBytes;
	ui MYsrcIndex = MYoffset + 3 * MYcol;
	ui MYdstIndex = MYoffset + 3 * MYmirrorcol;

	// 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位操作系统才能有效地使用。

代码6.7和代码6.8有一个共同点:两者都没有显式的循环，因为Nvidia硬件负责向每个启动的线程提供threadIDx.x、blockIdx.x和blockDim.x变量。每个启动的线程都知道它的线程ID和线程块ID各是什么，以及每个块中有多少个线程。所以，通过巧妙的索引,每个线程都可以获得类似于for 循环的控制，正如我们在代码6.7和代码6.8中看到的那样。由于GPU核函数每个线程的工作量都很小，去掉循环操作所节约的开销可能会非常巨大。