CUDA编程练习（五） 排序

看到大三发的博客，holy，那叫一个羞耻。那时候对某game可喜欢了
时间线是保研前，那时候我还不清楚接下来会发生什么，就像现在一样。。。

某传奇新游电竞选手在回复粉丝问题时曾言：我猛练！

关于这个系列我还没进行一个具体定义，它是：纯手撸（极个别简单到懒得动手的，GPT+个人改改）按照个人记忆与思路、想到哪里写到哪里（临时想到什么优化的点子就会试试）、充满漏洞（以练手为核心，毕竟共识在书里已经整理了，不寻求验证书本）

1 想法

排序算法种类可就多了，也是一两天弄不完的东西。

本篇我打算以这样的方式进行整理：

1. 基数排序
2. 合并排序
3. 奇偶排序
4. 样本排序

全篇对int进行排序（不是哥们，float怎么办呀？）深度网络训练我记得是有个h与d的约定来着，以什么个缩放比例，放进device就很纯粹了，没有指数占位，可以用基数。

有两本书，每本书上排序介绍流程都不一致，我按照一个先写着，后面再补。

2 代码

2.1 基数排序

2.1.1 Dummy

这个版本有较多问题，包括：为什么在全局存储反复读写？为什么要用两份额外空间？

但是吧，这份代码其实已经很优化了，因为考虑了访存合并、bank conflict，里面的交叉眼花缭乱

__device__ void radix_sort(unsigned int * const ps_vector, int vec_length, int leap, int tid,
                                    unsigned int * const sort_tmp_0,
                                    unsigned int * const sort_tmp_1) {  // 呃，貌似可以跨block排序
    // vec_length为整体长度，leap为线程个数，sort_tmp_0暂存为0的值，sort_tmp_1暂存位置为1的值
    for (unsigned int bit=0; bit<32; bit++) {  // 移动位数0-31
        u_int32_t base_excur_0=0;
        u_int32_t base_excur_1=0;  // 基础偏移
        u_int32_t bit_mask = (1<<bit);

        // 分配到两个列表
        for (int i=tid; i<vec_length; i+=leap) {  // 原书里面，i从0开始，这？如果数组长度不是leap整数倍怎么办？最后一轮不是越界了吗
            u_int32_t elem=ps_vector[i];  // 取数
            if (elem & bit_mask) {  // 不为0就是畜
                sort_tmp_1[base_excur_1+tid] = elem; // 记录到1的列表
                base_excur_1+=leap;
            }
            else{
                sort_tmp_0[base_excur_0+tid] = elem; // 记录到1的列表
                base_excur_0+=leap;
            }
        }

        // 移动到原来的数组
        for (int i=tid; i<base_excur_0; i+=leap) {
            ps_vector[i] = sort_tmp_0[i];
        }
        for (int i=tid; i<base_excur_1; i+=leap) {
            ps_vector[base_excur_0+i] = sort_tmp_1[i];
        }
    }
}

2.1.2 共存版

书中解决了上一节说的两个问题，后一个问题还行，前一个问题。。。解决的我一头雾水，您是只用一个block吗？GPU编程指南P125

用了共享内存（不一定会带来提升）少用了一片暂存区（一定提升）

__device__ void radix_sort_v2(int vec_length, int block_leap, int tid,