因此,我正在研究CUDA程序,并且在索引块和线程时遇到一些问题。基本上,我正在尝试在CUDA中实现Pixel Sort算法。 (通过一种修改,我们只处理行或列,而不是同时处理)

我想像的方式是简单地运行N个块,每个块具有1个线程(用于行或列的数量),并使每个块彼此独立地处理该行/列。

因此,如果我们想对列进行排序,我们将像这样启动内核(有几个额外的参数仅与我们的特定处理相关,因此为简单起见,我将其省略了)

pixel_sort<<<cols, 1>>>(d_image, d_imageSort, rows, cols, ...);

然后在内核中,我使用
int tid = blockIdx.x;

这使我可以在每个块中处理一个行/列数据,但是存在一些问题。它的运行速度比我们针对较小图像的算法的串行实现慢,并且当图像尺寸太大时会直线崩溃。

我正在考虑的另一种线程方案是将图像的每个像素映射到一个线程,但是我对此有两个疑问。
  • 如果我们要启动具有M个线程的N个块,代表具有M行的N个列,那么如何避免每个块的线程数限制为512(或1024?)。在这种情况下,我们可以让每个线程处理列中的多个像素吗?索引在内核中的外观如何?
  • 该算法基本上要求我们在整个列上工作,因此每个线程不能仅在该像素上做一些工作,它们必须进行通信,大概使用共享内存。每个块都具有一个“主”线程(进行实际排序计算),然后让所有其他线程参与共享内存,这将是一种有效的策略吗?

  • 其他说明:
  • 我们的图像数据通过OpenCV读取,并将RGBA值存储在uchar4数组
  • 最佳答案

    如果每个块只有一个线程,则会很快遇到线程占用问题。如果您的目标是进行全行排序(对于列,您可以在将图像发送到GPU之前先进行转置,以利用全局合并),那么获得不错结果的最快方法可能是对数进行基数或合并排序每行,基本上从http://mgarland.org/files/papers/nvr-2008-001.pdf复制步骤。您可以为每行分配k个m线程块,以使km> =图像宽度。然后,您将启动k *(图像高度)块。这样,您的网格的大小将为(k,height,1)。

    至于您的具体问题:

  • 您无法绕过512/1024每块线程数限制,因此必须重新构造算法。
  • “主”线程通常设计不佳,会导致停顿,开销,并且无法充分利用许多内核。有时您可能需要利用单个线程,例如输出/广播结果,但是大多数情况下您都希望避免使用它。请参阅链接的文章,以获取大部分可以避免这种情况的示例算法。
  • 关于c++ - CUDA中图像行/列的线程索引,我们在Stack Overflow上找到一个类似的问题:https://stackoverflow.com/questions/40983372/

    10-11 22:46
    查看更多