求适用于CUDA单线程的低Warp发散O(1)空间排序算法

针对单线程独立排序小数组时warp发散严重的问题，以下几种原地排序算法（空间复杂度O(1)）能有效减少分支发散，适配CUDA单线程场景：

1. 固定长度优化的选择排序 #

选择排序本身分支较少，结合N=49的固定特性，可通过循环展开和无分支操作彻底消除动态分支：

• 固定内外层循环次数，让编译器完全展开循环，消除循环终止条件的分支
• 去掉交换前的if (i != min_idx)判断（即使索引相同，交换操作不影响数组，避免分支）
• 用三元表达式实现无分支的最小值索引更新

__device__ void sort_49(double arr[]) {
    #pragma unroll
    for (int i = 0; i < 48; ++i) {
        int min_idx = i;
        #pragma unroll
        for (int j = i + 1; j < 49; ++j) {
            min_idx = (arr[j] < arr[min_idx]) ? j : min_idx;
        }
        // 无分支交换，省掉索引判断
        double temp = arr[i];
        arr[i] = arr[min_idx];
        arr[min_idx] = temp;
    }
}

2. 循环展开的插入排序 #

插入排序对小数组效率优异，通过固定循环次数和循环展开，可大幅减少分支：

• 固定外层循环次数（从1到48），消除动态终止分支
• 用#pragma unroll让编译器展开内层移动循环（或手动展开，进一步减少分支）
• 可选：用CUDA的__selp intrinsic实现无分支的元素移动，替代while循环中的if判断

__device__ void sort_49(double arr[]) {
    #pragma unroll
    for (int i = 1; i < 49; ++i) {
        double key = arr[i];
        int j = i - 1;
        #pragma unroll 5
        while (j >= 0 && arr[j] > key) {
            arr[j + 1] = arr[j];
            j--;
        }
        arr[j + 1] = key;
    }
}

3. 固定步长的梳排序 #

梳排序是冒泡排序的改进，针对N=49可预先计算固定步长序列（如40,13,4,1），全程无动态分支：

• 步长序列固定，循环次数固定，消除步长计算和终止条件的分支
• 无分支的元素交换判断，避免if-else导致的发散

__device__ void sort_49(double arr[]) {
    int gaps[] = {40, 13, 4, 1};
    #pragma unroll
    for (int g = 0; g < 4; ++g) {
        int gap = gaps[g];
        #pragma unroll
        for (int i = gap; i < 49; ++i) {
            double temp = arr[i];
            int j = i;
            while (j >= gap && arr[j - gap] > temp) {
                arr[j] = arr[j - gap];
                j -= gap;
            }
            arr[j] = temp;
        }
    }
}

为什么这些算法比堆排序更适合？ #

堆排序的分支发散源于堆调整过程中依赖数据的条件判断（如左右子节点比较、是否需要下沉），每个线程的数组独立，导致warp内线程分支走向完全不一致，inactive threads占比极高。而上述算法通过固定循环逻辑、无分支操作、循环展开，将动态分支转化为固定逻辑或无分支指令，大幅降低warp发散程度。

内容的提问来源于stack exchange，提问作者Bob