cuda练习(三)：使用gpu进行排序

生成数据

为了简便期间，生成不重复的数据

#define NUM_ELEMENT 4096
#define NUM_LISTS   32

template<class T> void c_swap(T &x, T &y){ T tmp = x; x = y; y = tmp; }

unsigned int srcData[NUM_ELEMENT];

void gen_and_shuffle(unsigned int * const srcData)
{
    for(int i = 0; i < NUM_ELEMENT; i++)    //生成不重复的数据
        srcData[i] = i;
    for(int i = 0; i < NUM_ELEMENT; i++)
        c_swap(srcData[rand()%NUM_ELEMENT], srcData[i]);
    return;
}

void print_data(unsigned int * const srcData)
{
    for(int i = 0; i < NUM_ELEMENT; i++)
    {
        printf("%4u", srcData[i]);
        if((i+1)%32 == 0)
            printf("\n");
    }
}

cpu排序

sort(srcData, srcData+NUM_ELEMENT);

gpu排序

归并排序

归并排序因为可以分解，很适合并行计算；但是传统的归并排序在合并时，到最后只有少数线程在操作，效率比较低。因此先使用桶排序让数据在列方向有序，然后所有列一起归并，而不是两两归并。

__global__ void sortincuda(unsigned int * const data, gpu_calc_type type)
{
    const unsigned int tid = threadIdx.x;
    __shared__ unsigned int sort_tmp[NUM_ELEMENT], sort_tmp_1[NUM_ELEMENT];

    copy_data_in_gpu(data, sort_tmp, tid);  //因为数据要经常被读取写入，因此拷贝数据到共享内存中以加速

    radix_sort2(sort_tmp, sort_tmp_1, tid); //桶排序

    switch(type)
    {
        case cpu_sort:  break;
        case gpu_merge_1: merge_1(sort_tmp, data, tid); break;  //单线程合并
        case gpu_merge_all: merge_atomicMin(sort_tmp, data, tid); break;    //多线程合并
        case gpu_merge_reduction: merge_two(sort_tmp, data, tid); break;    //两两归约合并
        case gpu_merge_reduction_modified: merge_final(sort_tmp, data, tid); break; //分块归约合并
        default: break;
    }
}

上述内核函数只在一个线程块运行，线程数为列数

sortincuda<< <1, NUM_LISTS>> >(gpu_srcData, type);

桶排序

首先可以进行按列的桶排序，让数据在每一列有序。这里桶排序为0/1桶排序。

__device__ void radix_sort2(unsigned int * const sort_tmp, 
                            unsigned int * const sort_tmp_1,
                            const unsigned int tid) //桶排序
{
    for(unsigned int bit_mask = 1; bit_mask > 0; bit_mask <<= 1)    //32位
    {
        unsigned int base_cnt_0 = 0;
        unsigned int base_cnt_1 = 0;

        for (unsigned int i = 0; i < NUM_ELEMENT; i+=NUM_LISTS) 
        {
            if(sort_tmp[i+tid] & bit_mask)  //该位是1，放到sort_tmp_1中
            {
                sort_tmp_1[base_cnt_1+tid] = sort_tmp[i+tid];
                base_cnt_1 += NUM_LISTS;
            }
            else    //该位是0，放到sort_tmp的前面的
            {
                sort_tmp[base_cnt_0+tid] = sort_tmp[i+tid];
                base_cnt_0 += NUM_LISTS;
            }
        }

        for (unsigned int i = 0; i < base_cnt_1; i+=NUM_LISTS)  //将sort_tmp_1的数据放到sort_tmp后面
        {
            sort_tmp[base_cnt_0+i+tid] = sort_tmp_1[i+tid];
        }
        __syncthreads();
    }
}

单线程合并

桶排序后，数据在每一列方向上是有序的，现在需要将所有列归并。
单线程归并指的是使用gpu中的一个线程进行归并操作。由于没有利用到多线程，速度较慢。

__device__ void merge_1(unsigned int * const srcData, 
                        unsigned int * const dstData,
                        const unsigned int tid) //单线程合并
{
    __shared__ unsigned int list_index[NUM_LISTS];  //不使用__shared__的话，就会创建在寄存器中，寄存器空间不够则会创