使用 CUDA 并行化冒泡排序

Question

我被分配了一项任务，即并行化冒泡排序并使用 CUDA 实现它。
我看不出冒泡排序是如何并行化的。我认为它本质上是顺序的。因为，它比较两个连续的元素并在条件分支之后交换它们。
有什么想法吗？

Answer 1

老实说，我也很难考虑并行化冒泡排序的方法。我最初想到了一种混合排序，您可以在其中平铺，对每个平铺进行冒泡排序，然后合并（如果可以使其工作，可能仍会提高性能）。但是，我浏览了“并行冒泡排序”，发现this page。如果向下滚动，您会发现以下并行冒泡排序算法：

For k = 0 to n-2
If k is even then
    for i = 0 to (n/2)-1 do in parallel
        If A[2i] > A[2i+1] then
            Exchange A[2i] ↔ A[2i+1]
Else
    for i = 0 to (n/2)-2 do in parallel
        If A[2i+1] > A[2i+2] then
            Exchange A[2i+1] ↔ A[2i+2]
Next k

您可以在 CPU 中运行 for 循环，然后为每个 do in parallels 使用内核。这对于大型阵列来说似乎很有效，但对于小型阵列来说可能开销太大。如果您正在编写 CUDA 实现，则假定使用大型数组。由于这些内核中的交换是相邻的元素对，因此您应该能够相应地平铺。我搜索了通用的、非 GPU 特定的并行冒泡排序，这是我唯一能找到的。

我确实找到了一个（非常轻微的）helpful visualization here，如下所示。我很乐意在 cmets 中对此进行更多讨论。

编辑：我发现了另一个并行版本的冒泡排序，称为Cocktail Shaker Sort。这是伪代码：

procedure cocktailShakerSort( A : list of sortable items ) defined as:
  do
    swapped := false
    for each i in 0 to length( A ) - 2 do:
      if A[ i ] > A[ i + 1 ] then // test whether the two elements are in the wrong order
        swap( A[ i ], A[ i + 1 ] ) // let the two elements change places
        swapped := true
      end if
    end for
    if not swapped then
      // we can exit the outer loop here if no swaps occurred.
      break do-while loop
    end if
    swapped := false
    for each i in length( A ) - 2 to 0 do:
      if A[ i ] > A[ i + 1 ] then
        swap( A[ i ], A[ i + 1 ] )
        swapped := true
      end if
    end for
  while swapped // if no elements have been swapped, then the list is sorted
end procedure

看起来这也有两个比较相邻元素的for循环bubbly..这些算法看起来有点相似的对立面，因为第一个算法（我现在学到的叫做odd-even sort）假设已排序并让 for 循环指定 false，而鸡尾酒摇酒器排序有条件地检查每个循环中的排序。

这篇文章中包含的 odd-even sort 代码似乎只是运行了足够的 while 循环来保证排序，维基百科伪代码检查的地方。潜在的第一步可能是实现本文的算法，然后使用检查进行优化，尽管使用 CUDA 进行检查实际上可能会更慢。

无论如何排序都会很慢。这是related SO question 仅供参考，但没有太多帮助。他们同意它对小型阵列无效，并真正强调它的失败。

您是在寻找特定的 CUDA 代码还是已经足够了？您似乎想要了解可能的选项并了解 CUDA 实施。

Answer 2

TL;DR;

如需通用并行冒泡排序的完整实现，请查看generic-bubble-sort.cu。 *这里的“通用”是指算法对任何类型的元素进行排序，只要您提供比较器。

充其量

使用与N 线程成线性比例的数字（比如N/2），您可以获得O(N) 时间复杂度的并行冒泡排序（其中N 是数组的大小你想排序）。

提示

这可能不是微不足道的，但是当您仔细观察时，您会意识到 顺序冒泡排序 所做的所有事情都是在没有正确排序的情况下交换一对元素一次一对！。

由于对可以独立排序，并行冒泡排序可以利用排序对独立的特性。

一种方法

假设我们要对以下array 进行升序排序：

#   [7][1][3][2][0] 

1. 我们首先采用初始未排序的array 并将每两个元素array[i] 和array [i+1] 视为一个独立的对。对于第一次迭代，i 将是一个EVEN 索引，所以我们的巴黎是{ {array[0], array[1]} , {array[2],array[3]}, ...}。

#    [7][1][3][2][0]   <-- Unsorted array of 5 elements
# [7][1]  [3][2]  [0] <-- A set of independent pairs.

1. 然后，如果每对中的每两个元素未按所需顺序排列，我们将交换它们。

# [7][1]  [3][2]  [0] --┑ Sorting first set of pairs
#                       |
# [1][7]  [2][3]  [0] <-┛ starting from an even idx

1. 我们的array 在第一次迭代后将如下所示：

# [1][7][2][3][0]   <-- Result after first iteration

4. 我们现在将再次重申，但与以前不同的是，我们现在将从 ODD 索引 { {array[1], array[2]} , {array[3],array[4]}, ...} 开始对对进行排序。值得一提的是，没有同行的元素不会被考虑。

#   [1][7][2][3][0]   <-- Result after first iteration

# [1]  [7][2]  [3][0] --┑ Sorting second set of pairs
                        |
# [1]  [2][7]  [0][3] <-┛ starting from an odd index

#   [1][2][7][0][3]   <-- Result after second iteration

5. 在N EVEN/ODD 对排序迭代之后，我们将有一个sorted array。

# [1][2]  [7][0]  [3] --┑
# [1][2]  [0][7]  [3]   |
#                       |
#   [1][2][0][7][3]     |
#                       | The whole parallel sorting
# [1]  [2][0]  [7][3]   | will converge after N iterations
# [1]  [0][2]  [3][7]   | So we keep sorting pairs for 3 more
#                       | iterations.
#   [1][0][2][3][7]     |
#                       |
# [1][0]  [2][3]  [7]   |
# [0][1]  [2][3]  [7] <-┛
# 
#   [0][1][2][3][7]   <-- Sorted array!

使用 CUDA 进行并行冒泡排序

上述方法的 CUDA 程序的直接实现将按如下方式完成：

• 每个线程负责对单独的对进行排序

• 你需要N/2线程

• 因为warp发散是我们需要关心同步我们的线程的事情

  • 使用单个块：如果我们的线程适合单个块，我们只会在每次迭代后使用__synchronize()，并且我们可以通过拥有所有线程来利用共享内存我们的数组在那里。

  • 使用多个块：我们必须确保内核中所有线程的线程同步。每次内核启动我们只能执行一次迭代，并启动我们的内核 N 次。坏消息是我们只能使用全局内存来处理我们的数组，因为共享内存只有内核生命周期。

一些代码

这是上面解释的一个简单的实现，只考虑一个块。完整代码可在in this repo获取。

template<typename T>
__global__
void bubbleSort(T* v, const unsigned int n, ShouldSwap<T> shouldSwap) {
    const unsigned int tIdx = threadIdx.x;

    for (unsigned int i = 0; i < n; i++) {

        unsigned int offset = i % 2;
        unsigned int leftIndex = 2 * tIdx + offset;
        unsigned int rightIndex = leftIndex + 1;

        if (rightIndex < n) {
            if (shouldSwap(v[leftIndex ], v[rightIndex ])) {
                swap<T>(&v[leftIndex ], &v[rightIndex ]);
            }
        }
        __syncthreads();
    }
}

如果您想了解 ShouldSwap 和 swap 的实现，请查看以下代码：

swap

用于交换元素的设备函数。

template<typename T>
__host__ __device__ __inline__
void swap (T* a, T* b) {
    T tmp = *a;
    *a = *b;
    *b = tmp;
}

ShouldSwap

用作通用比较器的 C++ Functor。

template<typename T>
__host__ __device__
bool ShouldSwap<T>::operator() (const T left, const T right) const {
    return left > right;
}