GPU高性能计算CUDA编程：GPU核函数的执行

阿里猫的编程小辈

已于 2024-09-09 00:39:07 修改

阅读量673

点赞数 9

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

于 2024-09-07 23:37:23 首次发布

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

版权

007 专栏收录该内容

37 篇文章 0 订阅

订阅专栏

GPU高性能计算CUDA编程：GPU核函数的执行

声明：本文不做商用

完成CPU到GPU的数据传输后，现在我们可以在GPU端执行GPU核函数了。代码6.3中与 GPU 端代码执行有关的代码行如下所示:

int IPH = ip.Hpixels;
int IPV = ip.Vpixels;
...
BlkPerRow = (IPH + ThrPerBlk -1 ) / ThrPerBlk;
NumBlocks = IPV*BlkPerRow; 
switch (Flip){
    case 'H': Hflip <<< NumBlocks, ThrPerBlk >>> (GPUCopyImg, GPUImg, IPH);
        GPUResult = GPUCopyImg;
        GPUDataTransfer = 2*IMAGESIZE;
        break;
    case 'V': Vflip <<< NumBlocks, ThrPerBlk >>> (GPUCopyImg, GPUImg, IPH, IPV);
        GPUResult = GPUCopyImg;
        GPUDataTransfer = 2*IMAGESIZE;
        break;
    ...
}

Flip 参数是根据用户输人的命令行参数设置的。当用户选择H选项时，Hflip()的GPU端函数被调用，并将三个指定的参数(GPUCopyImg、GPUImg和IPH)从CPU端传递给Hflip()。V选项会启动包含四个参数(GPUCopyImg、GPUImg、IPH和IPV)的 Vflip()核函数，而不是Hflip()核函数的三个参数。当看到两个核函数的具体细节后，你就会明白为什么 Vflip()需要一个额外的参数。

下面几行代码显示的是当用户在命令行中选择T(转置)或C(复制)选项时会发生什么。可以通过编写一个核函数来以更高效的方式实现转置操作。但我们的目标是展示如何依次地执行两个核函数。因此，在实现T操作时，我们首先启动Hflip，然后是Vflip高效地对图像进行了转置。当然，在实现C选项时，我们设计了一个完全不同的核函数PixCopy().

switch (Flip){
    ...
    case 'T': Hflip <<< NumBlocks, ThrPerBlk >>> (GPUCopyImg, GPUImg, IPH);
        Vflip <<< NumBlocks, ThrPerBlk >>> (GPUImg, GPUCopyImg, IPH, IPV);
        GPUResult = GPUImg;
        GPUDataTransfer = 4*IMAGESIZE;
        break;
    case 'C': NumBlocks = (IMAGESIZE+ThrPerBlk-1) / ThrPerBlk;
        PixCopy <<< NumBlocks, ThrPerBlk >>> (GPUCopyImg, GPUImg, IMAGESIZE);
        GPUResult = GPUCopyImg;
        GPUDataTransfer = 2*IMAGESIZE;
        break;
}

当用户选择日选项时，下面这一行代码的执行由Nvidia运行时引擎处理，包括启动Hflip()核函数并从CPU端向它传递上述的三个参数。

Hflip <<< NumBlocks,ThrPerBlk >>> (GPUCopyImg，GPUImg， IPH);

自此开始，我们将使用术语启动GPU核函数。这与术语调用CPU函数形成了对比。参见【0voice C++】根据类比6.2，CPU是在自己星球(地球)的内部调用一个函数，这对于GPU核函数来说可能并不准确。GPU实际上是一个协处理器，使用比CPU自身的内部总线慢得多的连接方式与CPU相连，因此调用一个位于像月亮这样遥远的位置处的函数需要一个更恰当的术语，比如启动。在上面启动GPU核函数的代码中，Hflip()是GPU核函数名，在<<<和>>>符号之间的两个参数(NumBlocks和ThrPerBlk)告诉Nvidia 运行时引擎以什么维度运行该核函数。第一个参数(NumBlocks)表示要启动多少个线程块，第二个参数(ThrPerBlk)表示在每个线程块中启动了多少个线程。请记住在类比6.2中，这两个数字就是cudaTown人想知道的参数:盒子的数量(NumBlocks)和每个盒子中的椰子数量(ThrPerBIk)。通用的核函数启动代码如下所示:

GPU Kernel Name <<< dimension,dimension >>> (arg1,arg2,...);

其中argl、arg2、……是从CPU端传递到GPU核函数的参数。在代码6.3中，它们是两个指针和IPH。在GPU内存区域创建存储图像存储空间时，由两个指针(GPUCopyImg和GPUImg)cudaMalloc()返回。IPH是一个变量，存放图像水平方向上的像素数(ipHpixels)。GPU的核函数Hfip()在执行期间需要这三个参数，如果在核函数启动时没有被传递进来，核函数将无法获得它们。回忆一下，类比6.2中的两个维度分别是166656和256，实际上对应于下面的启动代码:

Hflip <<< 166,656，256 >>> (GPUCopyImg，GPUImg， IPH);

这告诉Nvidia运行时引擎为Hflip()核函数启动166656个线程块，并将这三个参数传递给每一个线程块。也就是启动:线程块0、线程块1、线程块2、……、线程块166655。每一…、tid=255)，与我们在本书第一部分中个线程块都将启动256个线程(tid=0、tid=1、看到的 pthread示例相同。我们真正想说的是，该行代码总共启动了166656x256~41M个线程。