GPU上的并行冒泡排序

Question

我正在使用CUDA实现简单的冒泡排序算法，我有一个问题。

为了交换数组中的两个连续元素，我执行了以下代码：

if(a[threadIdx.x]>a[threadIdx.x + 1])
    Swap(a[threadIdx.x] , a[threadIdx.x + 1]);

请注意，块中的线程数是数组大小的一半。这是一个好的实现吗？即使有分支，单个warp中的线程也会并行执行吗？因此，实际上需要N次迭代才能对数组进行排序？

另请注意，我知道我可以实现更好的排序算法，我可以使用推力、CUDPP或SDK中的示例排序算法，但在我的情况下，我只需要一个简单的算法来实现。

Answer 1

我很高兴你意识到GPU上的冒泡排序可能会表现得非常糟糕！我正在努力找出如何在不启动许多内核的情况下获得足够的并行性。而且，当你完成的时候，你可能很难找到答案。

无论如何，为了回答你的具体问题:是的，在这种情况下，你很可能会有曲折发散。但是，考虑到"else“分支实际上是空的，这不会减慢您的速度。平均而言(直到这个列表被排序)，一个warp中大约一半的线程将采用"if“分支，其他线程将等待，然后当"if”分支完成时，warp线程可以回到同步状态。这远不是您最大的问题:)

Answer 2

我假设：

1. 您要排序的数组很小(<100个元素)
2. 它是某个更大的GPU算法的一部分
3. 数组驻留在共享内存空间中，或者可以复制到那里

如果这些都不是真的，就不要做冒泡排序！

块中的线程数是数组大小的一半。这是一个好的实现吗？

这是合理的。当warp中出现分支分支时，所有线程都会完全同步地执行所有分支，只是一些线程将其标志设置为"disabled“。这样，每个分支只执行一次。唯一的例外是-当没有来自经线的线采用分支时-那么分支被简单地跳过。

虫子！

然而，在您的代码中，我看到了一个问题。如果您希望一个线程对数组的两个元素进行操作，请让它们以独占方式处理它，即：

if(a[2*threadIdx.x]>a[2*threadIdx.x + 1])
    Swap(a[2*threadIdx.x] , a[2*threadIdx.x + 1]);

否则，如果Swap由两个相邻的线程执行，一些值可能会消失，而其他一些值可能会在数组中重复。

又一个bug！

如果你的块大于一个warp大小，记得在需要的时候放入__syncthreads()。即使您的块较小(不应该是)，您也应该检查__threadfence_block()，以确保对共享内存的写入对于同一个块的其他线程是可见的。否则，编译器可能在优化方面过于激进，从而使您的代码无效。

另一个问题

如果你修复了第一个错误，你将在你的共享内存上有双向存储体冲突。这不是很重要，但您可能希望重新组织数组中的数据以避免它们，例如，将连续的元素按以下顺序排列：

1、3、5、7、9、...、29、31、2、4、6、8、...、30、32

这样，元素1和2属于共享内存中的同一存储体。