CUDA 中图像行/列的线程索引

Question

所以我正在开发一个 CUDA 程序，但在索引块和线程时遇到了一些问题。基本上，我正在尝试在 CUDA 中实现 Pixel Sort 算法。 （通过一项修改，我们只处理任一行或列，而不是同时处理两者）

我将其可视化的方式是简单地运行 N 个块，每个块有 1 个线程（用于行数或列数），并让每个块彼此独立地处理该行/列。

因此，如果我们想对列进行排序，我们会像这样启动内核（有几个额外的参数只与我们的特定处理相关，所以为了简单起见，我将它们省略了）

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

然后在内核中，我用

访问块索引

int tid = blockIdx.x;

这允许我在每个块中处理一行/列数据，但它有一些问题。它的运行速度比我们针对较小图像的算法的串行实现要慢，并且当图像尺寸变得太大时会直接崩溃。

我正在考虑的另一种线程方案是将图像的每个像素映射到一个线程，但是我对此有几个问题。

1. 如果我们用 M 个线程启动 N 个块，用 M 行表示 N 个列，我们如何避免每个块的 512（或 1024 个？）线程限制。在这种情况下，我们可以让每个线程处理列中的多个像素吗？内核中的索引是什么样子的？
2. 该算法基本上要求我们处理整个列，因此每个线程不能只对那个像素做一些工作，它们必须进行通信，大概使用共享内存。每个块有一个“主”线程来进行实际的排序计算，然后让所有其他线程参与共享内存是否是一种有效的策略？

其他说明：

• 我们的图像数据通过 OpenCV 读取，并将 RGBA 值存储在 uchar4 数组中

Answer 1

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

至于你的具体问题：

1. 您无法绕过每块 512/1024 个线程的限制，您必须重新构建算法。
2. “主”线程通常设计不佳，会导致停顿、开销以及无法充分利用许多内核。您有时可能需要使用单个线程，例如输出/广播结果，但大多数情况下您想避免它。请参阅链接文章，了解主要避免这种情况的示例算法。