在GPU上进行并行冒泡排序

Question

我正在使用CUDA实现简单的冒泡排序算法，有一个问题。 我执行以下代码以交换数组中相邻的2个元素：

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

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

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

- scatman

Answer 1

我假设：

1. 你想要排序的数组很小（＜100个元素）
2. 它是某个更大的GPU算法的一部分
3. 该数组驻留在共享内存空间中，或者可以复制到共享内存空间中

如果这些条件不符合，请勿使用冒泡排序！

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

这是合理的。当warp中出现分歧分支时，所有线程都会完全同步地执行所有分支，只是一些线程的标志“禁用”设置。这样，每个分支只执行一次。唯一的例外是当来自warp的没有线程采取分支时，该分支被简单地跳过。

BUG！

然而，在你的代码中，我看到了一个问题。如果你想让一个线程处理数组的两个元素，请让它们独占处理，也就是说：

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

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

另一个错误！

如果您的块大于warp大小，请记得在需要时放置__syncthreads()。即使您的块较小（不应该），您也应该检查__threadfence_block()以确保共享内存中的写操作对同一块的其他线程可见。否则，编译器可能过于激进地进行优化，使您的代码无效。

另一个问题

如果您修复了第一个错误，则在共享内存中会有2路银行冲突。这并不是非常重要，但您可能希望重新组织数组中的数据以避免这种情况，例如按以下顺序排列连续元素：

[1, 3, 5, 7, 9, ..., 29, 31, 2, 4, 6, 8, ... , 30, 32]

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

- CygnusX1

Answer 2

很高兴你意识到在GPU上使用冒泡排序可能会表现得非常糟糕！我正在努力想出如何获得足够的并行性而不必启动许多内核。此外，你可能会难以确定何时完成。

无论如何，回答你的具体问题：是的，在这种情况下，你很可能会有warp分歧。然而，考虑到“else”分支实际上是空的，这不会减慢你的速度。平均而言（直到这个列表被排序），一个warp中大约一半的线程将进入“if”分支，其他线程将等待，然后当“if”分支完成时，warp线程可以重新保持同步。这远非你最大的问题 :)

- Edric