生成数据
为了简便期间,生成不重复的数据
#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__的话,就会创建在寄存器中,寄存器空间不够则会创