针对CUDA中的图像行/列的线索索引

Question

因此，我正在研究CUDA程序，并且在索引块和线程时遇到了一些问题。基本上，我试图在CUDA中实现Pixel Sort算法。

我被可视化它是简单地与每个1级的线程中运行的N个块（对于行数的方式（有一个修改中，我们只，而不是两个同时处理要么行或列），或列），并让每个块独​​立处理该行/列。

所以，如果我们要进行排序列，我们推出这样的内核（也有一对额外的参数仅仅是我们的具体处理相关的，所以我离开了出来为简单起见）

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

然后在内核中，我访问与

int tid = blockIdx.x; 

块索引这使我每块一个行/列中的数据的工作，但仍存在一些问题。它的运行速度比我们对较小图像的算法的串行执行速度要慢，并且在图像尺寸变得过大时直接崩溃。

我正在考虑的替代线程方案是将每个图像的像素映射到一个线程，但是我有几个关于这个问题的问题。

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

其他注释：

• 我们的图像数据通过OpenCV的读入，并具有存储在uchar4阵列

Answer 1

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

至于你具体的问题：

1. 不能围绕512/1024线每块限制得到，你必须调整你的算法。
2. “主”线程通常设计不佳，造成停顿，开销，并没有充分利用多核。您有时可能需要使用单个线程，比如输出/广播结果，但大多数情况下您希望避免它。请参阅链接的文章以了解大多数避免这种情况的示例算法