CUDA编程练习（十三） 样本排序

咱就是说，不必每期都写肉麻的话

1 想法

这个算法就是为了避免合并排序而提出的。合并意味着大量的跨block交互，大量的非合并访存。这是万万使不得的。但是这样的方法也有弊端：不是均匀分布的乱序数组，整的桶子大小不一，影响效率。

它包含了以下步骤：

1. 提取样本；
2. 样本排序；
3. 装桶（提取+前缀和）；
4. 桶内排序；

乍一看，合理、有创意，实现起来会发现步骤一个又一个，需要注意的点也很多。本篇文章以CPU+GPU代码一起写的方式来整。参考了书，没办法，它真的很麻烦，如果我自己写会更好，但是考虑的多了有时也会没必要。反正原理很简单，有的步骤就按照书上的全局内存写法来。问题不大。

本篇只讨论升序，因为二分排序要改降序会死脑细胞

2 代码

2.1 提取样本

CPU代码

__host__ double select_sample_cpu(const ELEM_TYPE* const src_data, ELEM_TYPE* const sample_data, const uint interval, const uint n) {
    clock_t start, finish;
    start = clock();

    uint sample_idx=0;
    for (uint src_idx=0; src_idx<n; src_idx+=interval) {
        sample_data[sample_idx] = src_data[src_idx];
        sample_idx++;
    }

    finish = clock();
    return (double)(finish-start) / CLOCKS_PER_SEC * 1000;
}

GPU代码

线程数量=数字总数，那么变量是？

GRID_DIM!

// 选择样本
__global__ void select_samples_gpu(const ELEM_TYPE* const src_data, ELEM_TYPE* const sample_data, const uint interval) {
    // 线程数量=要选的数字数量，所以就不来额外的限制了
    uint tid = (blockIdx.x*blockDim.x)+threadIdx.x;
    sample_data[tid] = src_data[tid*interval];
}

2.2 样本排序

CPU代码

使用了qsort

// 降序
int compare_down(const void* a, const void* b) {
    return (*(ELEM_TYPE*)b - *(ELEM_TYPE*)a);
}
// 升序
int compare_up(const void* a, const void* b) {
    return (*(ELEM_TYPE*)a - *(ELEM_TYPE*)b);
}
__host__ double sort_sample_cpu(ELEM_TYPE* const sample_data, const uint n) {
    clock_t start, finish;
    start = clock();

    qsort(sample_data, n, sizeof(uint), &compare_up);

    finish = clock();
    return (double)(finish-start) / CLOCKS_PER_SEC * 1000;
}

GPU代码

// 基数排序使用thrust函数库，你瞅啥，基数排序我写过了

2.3 装桶

CPU代码

包括桶子计数，桶子大小前缀和，装桶。通用的二分查找放最后

// 样本桶计