sorting - GPU 上的并行冒泡排序

Question

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

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

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

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

Answer 1

我假设:

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

如果其中任何一个不正确，请不要进行冒泡排序!

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

这是有道理的。当 warp 中出现发散分支时，所有线程都以完美同步的方式执行所有分支，只是一些线程将其标志设置为“禁用”。这样，每个分支只执行一次。唯一的异常(exception)——当没有来自 warp 的线程采用分支时——那么该分支将被简单地跳过。

错误!

但是我在您的代码中发现了一个问题。如果想让一个线程对数组的两个元素进行操作，就让它们独占处理，即:

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

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

另一个错误!

如果您的 block 大于 warp 大小，请记住在需要时放置 __syncthreads()。即使您的 block 较小(不应该)，您也应该检查 __threadfence_block() 以确保同一 block 的其他线程可以看到对共享内存的写入。否则，编译器可能会过度优化并使您的代码无效。

另一个问题

如果您修复了第一个错误，您的共享内存将出现 2-way bank 冲突。它不是非常重要，但您可能希望重新组织数组中的数据以避免它们，例如具有以下顺序的连续元素:

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

这样，元素 1 和 2 在共享内存中属于同一组。