CUDA 归并排序的GPU加速

原创

已于 2022-02-10 15:13:26 修改 · 4.2k 阅读

标签

#cuda #c++ #排序 #算法 #数据结构

于 2022-02-02 22:19:09 首次发布

本文探讨了如何使用CUDA实现归并排序的GPU加速，通过CUDA编程将归并过程并行化，对比了GPU与CPU串行版本的时间效率，并讨论了并行策略优化和GPU潜力。

CUDA实现归并排序 (merge sort) 的GPU加速

最近在学习CUDA, 寻找一个可以练练手的问题，正好想到了一个问题特别适合用GPU加速，那就是归并排序。
归并排序使用分治的思想，化解成多个互不相关的子问题，因此特别适合用并行化来解决。

用GPU进行归并排序的思路：

把数组分成多段, 把相邻两段的有序数组合并，由于有很多段，因此可以并行的方式同时合并。合并完之后段数减半(每段长度加倍)，然后继续上一过程进行并行化地合并，直到最后只剩下一段数组，那就是整个数组都排好序。

注意点：

为了实现并行，需要把归并排序中原本的递归转化为迭代。需要从最底层（比如分解到只有一个数的数组）逐层往上进行合并，与通常的归并排序使用的递归的思路有所不同。（即把递归的从上层往下层分解的过程反过来，从下层逐渐整合到上层，类似的思想在FFT中把递归优化为迭代的算法中也出现过）
为了使迭代过程中的合并更好实现，数组长度在合并过程中最好是1变2, 2变4, 4变8, 8变16……，这样我们每次合并都能确切地知道左右两边的数组长度是多少。如果最终的数组长度不是2的n次方，我们可以通过数组长度的二进制表示，把把其分解为多段长度为 2 ** n 的数组（比如，如果长度是26，二进制是11010，即 26 = 2 + 8 + 16，于是可以拆解成长度分别是2, 8, 16的数组，分别进行GPU版本的归并排序，最后从短数组向长数组依次把这几段合并起来）

代码如下：

作为对比，我写了一个CPU串行版本的归并排序，与GPU并行版本的排序对比运行时间。（其中，函数mergeSort_gpu调用了GPU上运行的kernel核函数mergeVec_half；另外手写了一个CPU串行版本的归并排序的函数mergeSort_cpu作为对比）

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <time.h> 
#include <math.h>
#include <vector>
#include <memory>
#include <iostream>
#include <algorithm>
#define BIG (1e7)
// #define DEBUG
using namespace std;
template<typename theIterator> void print(theIterator begin, theIterator end);


template<typename T> __global__ void
mergeVec_half(T *A, T *tmp, const int64_t vSize) {
   
   

    /* splict the vector A into two halfs
     * merge these two half together
     *
     * tmp is a temporary vector to 
     * receive the merge result
     */

    int64_t left = blockIdx.x * vSize;
    int64_t right = left + vSize - 1;
    int64_t mid = (left + right) / 2;

    int64_t i = left, j = mid + 1, k = left;  // index of left half, right half, and the mergeVec
    while ((i <= mid) && (j <= right)) {
   
   
        if (A[i] <= A[j]) {
   
   
            tmp[k] = A[i];
            ++i; ++k;
        } else {
   
   
            tmp[k] = A[j];
            ++j; ++k;
        }
    }
    if (i > mid) {
   
   
        for (; j <= right; ++j, ++k) {
   
   
            tmp[k] = A[j];
        }
    } else {
   
   
        for (; i <= mid; ++i, ++k) {
   
   
            tmp[k] = A[i];
        }
    }
    /// copy tmp to A
    for (k = left; k <= right; ++k) {
   
   
        A[k] = tmp[k];
    }
}


template<typename theIterator, typename T> void 
mergeSort_power2n(theIterator begin, theIterator end, T args) {
   
   
    /* 
        sort a vector with size of power(2, n)
    */
    clock_t begT, endT;

    T *dataA, *dataTmp;
    int64_t vSize = end - begin;
    cudaMalloc((void**)&dataA, sizeof(*begin) * vSize);
    cudaMalloc((void**)&dataTmp, sizeof(*begin) * vSize);