GPU高性能计算CUDA编程：设计GPU线程（CPU线程）的原则及使用示例

最新推荐文章于 2024-10-08 00:02:21 发布

阿里猫的编程小辈

最新推荐文章于 2024-10-08 00:02:21 发布

阅读量502

点赞数 4

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

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

版权

007 专栏收录该内容

37 篇文章 0 订阅

订阅专栏

GPU高性能计算CUDA编程：设计GPU线程（CPU线程）的原则及使用示例

声明：本文不做商用

以下是设计GPU(或CPU)线程的一般性指导原则。我们将根据这些指导原则来检查GPU核函数的细节:
多线程CPU和GPU代码中的每个线程都会经历三个阶段的操作。这三个阶段是:
1.我是谁?核函数获得自己的 ID。
2.我的任务是什么?核函数根据ID来确定它应该处理哪部分数据,
3.执行……完成它应该做的事情。

让我们将这些准则应用于Vflip()中的每个步骤:
1.我是谁?下面是代码6.7中Vflip()核函数的第一部分:

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

在这里核函数Vflip()提取自己的块ID、线程ID和ThrPerBlk值。清楚起见，该核函数使用了与main()相同的变量名称ThrPerBlk，但它是该核函数的局部变量，因此名称其实可以随意。因为本例启动了41M个核函数，所以上面的Vflip()函数只代表这41M个核函数中的一个线程。因此，核函数的第一个任务是计算它是这41M个核函数中的哪一个。全局线程ID位于名为MYgtid的变量中。该步计算操作根据线程块ID(在MYbid变量中)和线程ID(在MYtid变量中)实现了线程索引(或者线程ID)的“线性化”。参见【0voice C++】

2.我的任务是什么?我用术语线性化线程ID来表示MYgtid。确定了MYgtid后Vflip()将继续执行如下的代码:

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;

线性化的概念是指将二维索引转换为一维索引，就像我们在公式6.2中看到的那样，根据像素的x和y坐标计算像素的线性存储器地址。这里的相似之处(以MYgtid为例)是41M个线程的启动模式实际上是2D的，块实际上就是x维，每个块中的线程就是y维。因此，在这种情况下，线性化允许一个线程确定其在所有41M个线程中的全局唯一DD-MYgtid，这是无法由MYbid或MYtid单独确定的。

在决定自己的任务是什么时，Vflip()首先确定它需要复制的像素位于哪一行(MYrow)和哪一列(MYcol)，以及将要被复制到镜像行的行索引(MYmirrorrow)。计算了需要处理的列索引后，一个线程一旦意识到自己是一个无用线程后就会退出，正如6.4.14节中所描述的那样。接下来，Vflip()将行索引和列索引转换为源和目标GPU内存地址(MYsrcIndex和 MYdstIndex)。请注意，它使用了在main()内通过cudaMalloc()函数分配后传递给该核函数的ImgSrc和ImgDst指针。
3.执行……计算好源和目标内存地址后，剩下的就是将该像素连续的三个字节从GPU内存中的源地址复制到目的地址。

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