CUDA 算法之 奇偶排序

我们常用的排序算法基本上都是单线程的，要在CUDA运用并行进行排序，就需要新的排序算法。这篇就介绍CUDA适用的排序方法：奇偶排序（odd even sort），它是基于“老朋友”冒泡排序衍生出来的算法。

算法示意图

CUDA的算法复杂度：O(n),稳定排序

由于每个“小块”进行交换的时候都是互不相关的，所以使用CUDA进行并行计算非常合适。

CUDA代码实现

核函数：

__global__ void cudaOddEvenSort(int* data, __uint32_t len){
    unsigned int idx = blockIdx.x*blockDim.x + threadIdx.x;
    __uint8_t isOdd = 0;
    __uint32_t counter = len;

    do{
        isOdd = (++isOdd)%2;
        int d0 = data[isOdd+2*idx];
        if(isOdd+2*idx+1>=len) {continue;}    //边界检查
        int d1 = data[isOdd+2*idx+1];

        if(d0>d1) {
            data[isOdd+2*idx] = d1;
            data[isOdd+2*idx+1] = d0;
        }
    }while(counter--);
}

主函数调用：

#include <cuda_runtime.h>

#define LEN 127

int main(int argc,char **argv) {

    int *sortdata ;
    cudaMallocManaged(&sortdata, LEN);
    for(int i=0; i<LEN; i++) sortdata[i]=LEN-i;
    cudaOddEvenSort<<< 1, LEN/2>>>(sortdata, LEN);
    cudaSafeCall(cudaDeviceSynchronize());
    for(int i=0; i<LEN; i++) printf("%d ", sortdata[i]);
    cudaFree(sortdata);

    return 0;
}

假设数据集长度为N=2n（偶数），需要使用的threads则为n。当数据N=2n+1（奇数）时则需要n个threads。核函数中的需要做数据边界处理。循环次数为数据集长度：N。

优化

我们使用CUDA是意在加速，上述程序能否继续追求更高的性能？其实是可以的，后续优化可以加入L1共享内存功能，将数据集放入L1缓存中进行操作，可以加速IO速度，缺点就是数据集不能超过48KB（这是我们能编程的最大L1缓存了）。

核函数：

__global__ void cudaOddEvenSortAccShared(int* data, __uint32_t len){
    unsigned int idx = blockIdx.x*blockDim.x + threadIdx.x;
    __uint8_t isOdd = 0;
    __uint32_t counter = len;
    //copy2L1
    extern __shared__ int s[];
    for(int i=0; i<len; i++) s[i] = data[i];

    do{
        isOdd = (++isOdd)%2;
        int d0 = s[isOdd+2*idx];
        if(isOdd+2*idx+1>=len) {continue;}    //boundary
        int d1 = s[isOdd+2*idx+1];

        if(d0>d1) {
            s[isOdd+2*idx] = d1;
            s[isOdd+2*idx+1] = d0;
            __syncthreads();
        }
    }while(counter--);

    //copy back
    for(int i=0; i<len; i++) data[i] = s[i];
}

共享内存使用有两种方式：动态和静态。这里使用的是动态分配。

主函数调用：

#include <cuda_runtime.h>

#define LEN 127

int main(int argc,char **argv) {

    int *sortdata ;
    cudaMallocManaged(&sortdata, LEN);
    for(int i=0; i<LEN; i++) sortdata[i]=LEN-i;
    cudaOddEvenSortAccShared<<< 1, LEN/2, LEN*sizeof(int)>>>(sortdata, LEN);
    cudaSafeCall(cudaDeviceSynchronize());
    for(int i=0; i<LEN; i++) printf("%d ", sortdata[i]);
    cudaFree(sortdata);

    return 0;
}