如果你一直在做数据量比较大的排序任务多半也遇到过这类尴尬单线程std::sort跑在几十万甚至上百万条记录上性能肉眼可见地拉胯换成并行快排又担心负载不均衡、pivot 选不好导致某些线程空转。我最初是在给某图像处理 Demo 做特征向量排序时被这个问题卡住才开始系统研究并行排序算法最后把注意力锁定在双调排序和基数排序这两个看起来方向完全不同的家伙身上。这篇文章不是什么教科书翻译而是我把两种并行排序从原理、实现到实测选型完整走了一遍之后的实践记录适合正在为大规模数据排序发愁、打算在 CPU 多核或 GPU 上做并行排序的开发者参考。先说结论双调排序是典型的“比较网络型”并行算法比较顺序与数据无关天生适合 SIMD/GPU 这类无分支、有固定计算模式的硬件基数排序则是“数据驱动型”算法靠按位分桶绕过比较吞吐量极高是处理超大规模整数/浮点键排序时的首选。两者不是替代关系在真实工程里甚至可以组合使用。下面我按“算法路线 → 双调排序 → 基数排序 → 实测对比 → 踩坑与混合方案”的顺序展开尽量把每个关键的为什么都说清楚。1. 两种算法路线比较网络与数据驱动排序算法到了并行的语境里最重要的分类标准不再是“快排 vs 归并”而是“是否能在排序开始前就确定所有比较操作”。这个维度把主流方向分成两派。传统比较排序的思路是先比较再决定下一步比较什么。快排的 partition 结果依赖数据本身pivot 好坏直接影响子任务大小所以在多线程里很容易出现“一个线程忙死、其他线程闲死”的负载不均衡。归并排序稍微好一些但归并树的每一层都需要同步而且合并操作本身就是串行瓶颈。这些问题的根源都在于比较操作之间有一条动态的数据依赖链我们只能等它跑完一层才知道下一层该怎么做。双调排序走的是另一条路。它的全部“比较-交换”操作都能在算法开始前静态列出来不依赖任何键值内容。你可以把整个排序过程想象成一张固定电路图数据从左边进去经过若干层比较器从右边出来就是有序的。这种性质叫“比较网络”实现起来特别直白——每个比较器都是独立可并行的层与层之间只需要一次同步。基数排序则根本不用“比较”这个词。它把键拆成数字位按照桶的编号把元素搬来搬去。单趟扫描是 O(n)总复杂度取决于键的位数和每次处理的 digit 宽度数据是否接近有序对它几乎没影响。它的难点不在比较而在并行写桶时的竞争以及访存局部性。这两个方向的一个核心差异值得记住比较网络路线追求“关键路径尽量短”哪怕总比较数多一点也无所谓数据驱动路线追求“单趟吞吐尽量大”但要把桶分布和内存访问模式照顾好。理解了这一点后面看实测数据就不会觉得奇怪——双调排序总操作数明明更多某些规模下却比快排更快因为它每一层的所有比较都在同一时刻发生。2. 双调排序从双调序列到固定比较网络双调排序这个名字听起来有点冷门但它的数学基础非常漂亮一个序列被称为双调的是指它先单调递增、再单调递减或者可以通过循环移位变成这种形状。比如[3, 5, 8, 7, 4]就是典型的双调序列先升到 8 再降[5, 8, 7, 4, 3, 6]也是双调的因为它把[5, 8, 7, 4, 3, 6]从 6 处切开再拼接就等价于先增后减的序列。2.1 双调序列的定义与双调合并定理双调排序依赖一个非常反直觉的定理。假设a[0..n-1]是双调序列其中n是 2 的幂。把下标i和i n/2的元素两两比较小的放在前半部分大的放在后半部分。做完这一轮“跨点半比较”之后有两个神奇的结果前半部分的所有元素都不大于后半部分的所有元素前半部分和后半部分各自仍然是双调序列。这个性质意味着我们可以递归地继续处理两个子段最终得到一个完全有序的序列。整个过程不需要知道双调序列的峰在哪里不需要记录任何动态状态只需要反复做同一件事把距离为n/2的两个元素按大小交换。这正是它适合并行实现的数学根基。2.2 递归实现与工程实现的分野要给任意序列做双调排序先得把序列改造成双调序列。一个简单办法是分治对左半段做升序双调排序对右半段做降序双调排序拼接起来就是双调序列然后对整个序列做一次双调合并。下面这个递归版本的 Python 代码我建议先跑一遍找感觉def bitonic_merge(a, low, n, up): if n 1: return k n // 2 for i in range(low, low k): if (a[i] a[i k]) up: a[i], a[i k] a[i k], a[i] bitonic_merge(a, low, k, up) bitonic_merge(a, low k, k, up) def bitonic_sort(a, low, n, up): if n 1: return k n // 2 bitonic_sort(a, low, k, True) # 前半升序 bitonic_sort(a, low k, k, False) # 后半降序 bitonic_merge(a, low, n, up) # 合并成整体有序 arr [4, 2, 7, 1, 8, 3, 6, 5] bitonic_sort(arr, 0, len(arr), True) print(arr) # [1, 2, 3, 4, 5, 6, 7, 8]注意两个细节。第一up参数代表这个子段最终是升序还是降序它是贯穿全过程的“全局方向”第二递归版本里后半段排序方向是False这是为了构造双调序列和最终方向无关。递归实现最容易理解但真正跑性能时不会有人写递归因为最后一层比较粒度太小递归调用栈开销占比太高。工程上要把递归展开成三层循环外层枚举段大小size中层枚举合并时的半步长k内层遍历段内比较对。段的方向由段号奇偶决定奇数段和偶数段方向相反。写一个可运行的迭代版本供参考def bitonic_sort_iter(a, n, upTrue): size 2 while size n: for base in range(0, n, size): direction up if (base // size) % 2 0 else (not up) k size // 2 while k 0: for i in range(base, base k): j i k if (a[i] a[j]) direction: a[i], a[j] a[j], a[i] k // 2 size * 2这个版本的逻辑和递归完全等价。direction为什么要按段号奇偶交替因为在更大的合并阶段相邻的段必须保证方向相反才能让它们拼接成一个更大的双调序列。你可以把size2阶段想象成给后续每个size4段分别准备“升序左半 降序右半”方向自然是一正一反。整个双调排序的比较总次数是n/2 * log(n) * (log(n) 1)属于O(n log^2 n)比最优比较排序的O(n log n)多一截。但它换来了两个在并行世界极其珍贵的东西比较次序完全静态层与层之间只有O(log^2 n)次同步点。它给人感觉就像一条有固定深度的流水线数据进去、排序完成期间没有任何分支判断。2.3 在 GPU/SIMD 上预生成比较对既然比较对是静态的最直接的工程技巧是“预生成索引表”。CPU SIMD 下可以把比较-交换操作映射成两条向量指令一条取两向量的逐元素min一条取max然后用掩码写回正确位置。GPU 上更简单每个线程负责一对元素的比较线程编号按预生成的索引表来取数据。一个常用的 GPU 内核片段是这样// 伪代码比较对由异或直接生成 __global__ void bitonic_pass(int* data, int step, int direction) { int i blockIdx.x * blockDim.x threadIdx.x; int j i ^ step; // 关键技巧异或生成配对下标 int low min(data[i], data[j]); int high max(data[i], data[j]); // 根据位置决定写回顺序 data[i] (i step) 0 ? low : high; data[j] (i step) 0 ? high : low; }i ^ step这个技巧特别值得记。当step是 2 的幂时每个下标i和i ^ step正好是一对配对下标而且不会出现重复配对。用位运算代替数组索引查找在 GPU 上能省掉一整段显存访问。如果你在 CUDA 里跑双调排序比较对和方向完全可以在编译期展开内核函数里没有if、没有循环、没有动态寻址这对 GPU 的吞吐率非常友好。也正因为比较对是静态的双调排序可以配合共享内存做小块排序。比图每个 block 处理 256 个元素块内用完整展开的比较网络不需要任何原子操作也几乎不会有内存加载不均。这类“小块内排序”正是双调排序在 GPU 上最舒服的场景。2.4 双调排序的适用边界双调排序不是万能的。第一个限制是长度必须是 2 的幂不是 2 的幂时需要用哨兵元素补齐第二个限制是O(n log^2 n)比较总次数在数据规模变大之后会被基数排序拉开差距。我自己的实测体会是双调排序最适合的元素规模大致在2^10到2^14之间也就是作为 GPU 上一个 block 能装下的小块排序器。一旦数据量大到需要跨 block、跨核显存通信双调排序的优势就会被同步开销慢慢吃掉。另外纯键值比较的双调排序不是稳定排序。如果需求是“先按优先级再按原始序”你得把比较键扩展成(key, original_index)的复合比较或者干脆选基数排序。这个点我在第 5 章还会提到。3. 基数排序按位分桶的并行化之路基数排序的策略和比较网络完全不同。它不碰键与键之间的比较而是把每个键拆成多个 digit每轮按其中一个 digit 分桶桶的编号天然决定顺序。经典实现大多走 LSD从低位到高位路线原因只有一个LSD 能保证排序稳定性上一轮的桶内顺序可以在下一轮被保留。3.1 LSD 基线流程与三个并行难点先看一个最朴素的串行 LSD 基数排序这里以无符号 32 位整数为例每轮处理 8 位def lsb_radix_sort(data): n len(data) out [0] * n for shift in (0, 8, 16, 24): count [0] * 256 for x in data: count[(x shift) 0xff] 1 # 前缀和 total 0 for i in range(256): count[i], total total, total count[i] for x in data: digit (x shift) 0xff out[count[digit]] x count[digit] 1 data, out out, data return data这段代码逻辑简单但并行化之后会遇到三个问题计数阶段的竞争。每个线程都需要对自己的元素做digit统计如果多个线程同时更新同一个全局计数桶就要用原子操作。散射阶段的写冲突。统计完前缀和后每个元素要写回out的某个位置这个位置是所有线程共享的必须用“桶游标”来逐元素分配。访存局部性。如果键值分布均匀散射时写回的目标地址几乎是随机的缓存命中率会很差。这三点是并行基数排序优化的主线。3.2 block 级计数与散射偏移计算“偷懒”的直接方案是全局atomicAdd计数然后再做一次全局前缀和。这种写法在数据量小、分布均匀时能跑但数据稍微倾斜就会陷入严重的原子竞争。正确做法是分两级先让每个 block 维护自己的一组局部计数然后只把局部计数的汇总结果用原子操作合并到全局。具体流程是假设每个 block 处理一段连续数据先把这段数据里的元素按 digit 分到 256 个本地桶里桶的计数存放在共享内存数组local_count[256]中。然后对每个桶号做一次跨 block 的归约得到全局前缀偏移global_offset[bucket]。散射阶段每段再算出自己在全局中的起始地址直接写入输出数组全程不需要全局原子操作竞争只发生在 block 内部压力小得多。这里有个工程细节要注意如果每个线程都用共享内存计数器做原子操作256 个线程同时抢 256 个计数器的场景仍然会有明显的原子延迟。更进一步的做法是每个线程先在寄存器里统计自己负责的那批元素再合并进共享内存这样共享内存原子操作的次数能减少到原来的 1/elements_per_thread。散射偏移计算也要分两步每个 block 先算出自己每个桶的本地前缀和再叠加 block 的全局基地址。这样写目标地址完全确定不会有两个线程写同一个位置散射阶段可以无锁并发。3.3 浮点数与有符号整数的位模式映射这是基数排序最容易栽跟头的地方。浮点数的内存位模式直接按无符号整数比较顺序是错的所有负数在无符号整数下会排在正数之后因为负数的符号位是 1。比如-1.0的位模式是0xBF8000001.0是0x3F800000按无符号比较会错误地认为「1.0 小于 -1.0」。解决办法是把浮点位模式映射成一个保持浮点数数值顺序的无符号整数。经典做法是uint32_t float_to_sorted_bits(float f) { uint32_t b; memcpy(b, f, sizeof(b)); // 取出 IEEE 754 位模式 uint32_t mask (b 0x80000000u) ? 0xffffffffu : 0x80000000u; return b ^ mask; }原理不复杂正数符号位 0翻转符号位把它们统统移到无符号数域的高半段并且保持它们内部的大小顺序负数符号位 1翻转全部位把原本“越大越靠后”的负数位模式反转成“数值越小位模式越小”的顺序。做完这个映射后浮点数的排序就退化成无符号整数排序排序完之后再用反向映射还原。有符号整数也可以用同样的思想非负数翻转符号位进高半段负数取反进低半段。NaN 是个例外它的位模式没有全序性建议在排序前就过滤掉或统一映射到极值。这个位模式技巧看起来不起眼但它在实际项目里非常值钱。因为很多排序任务最终要排的数据不是纯整数而是浮点分数、坐标、距离忽略这一步会让基数排序直接给出错误结果。3.4 自适应跳过空 digit 的小技巧32 位键在 8 位 digit 下需要 4 趟但真实数据往往没有用满 32 位。某个模拟项目的测试数据里我们拿到的是一批 16 位分数直接跑 4 趟等于白做两趟。一个很实用的优化是排序前先采样一批元素统计最高非零 digit只从那个 shift 开始跑。如果键空间很小甚至可以只跑一趟。另一种情况是高位 digit 完全恒定。比如按用户 ID 排序但 ID 高字节永远是0x00那么这一趟的 256 个桶里只有少数几个桶有数据其余全是空桶。与其在空桶上浪费前缀和不如把分桶宽度从 8 位调成 11 位甚至 16 位减少趟数。你把位宽调大后桶数组变大了但缓存能装下 2048 个 int 计数数组约 8KB效果依然不错。4. 实测数据谁的墙钟时间更短光看复杂度推导容易产生误解特别是双调排序理论复杂度明显劣于基数排序但实测不一定输。我在这台测试机上8 核 16 线程 CPU数据为随机生成的 uint32以下时间为多次运行取中值记录了一组对比数据规模std::sort 单线程OpenMP 并行快排双调排序(SIMD, 多线程)LSD 基数排序(8bit, 多线程)2^12 (4096)0.0003s0.0006s0.0002s0.0005s2^16 (65536)0.008s0.002s0.010s0.001s2^20 (1048576)0.150s0.030s0.210s0.009s2^24 (16777216)3.200s0.420s2.800s0.092s这里要说明一点双调排序在多核 CPU 上并不总是赢因为每层比较之间都要同步一次但在 SIMD 向量化和固定比较网络的加持下它赢在局部性和无分支。它的墙钟时间和数据内容没关系波动很小而快排在数据分布不好时可能从 0.2s 跳到 1s 以上。基数排序在这个表里几乎是一骑绝尘。原因很简单它只需要固定 4 趟扫描每趟都是线性读取和线性写回数据吞吐非常高。当数据规模来到百万级基数排序和比较排序的差距已经不是一个量级而是接近一个数量级。但注意它的优势建立在一个前提上键是整数或者能映射成整数的类型。如果排的是变长字符串基数排序的优势就没有了。4.1 为什么双调排序输在总比较数、赢在关键路径为什么双调排序比较次数明明更多在中等规模却不比快排慢太多把乘法模型换成深度模型就清楚了。串行快排的关键路径是O(n log n)数量级的比较一个接一个执行双调排序的关键路径是O(log^2 n)同一层的比较全部同时发生。在 16 核机器上你比的是深度而不是总工作量的一半——这就是并行排序和串行排序最容易被忽视的差异。GPU 上这个差异更明显。某个 GPU 内核里我们把 512 个元素的小块排序从原来的串行插入排序改成双调排序后内核总耗时减少了大约 40%。因为双调排序消除了所有动态分支一个 block 内 512 个线程全部做同样的事延迟被完全隐藏。4.2 选型规则把实测结果浓缩成一条简单的经验法则场景推荐算法理由元素数量 2^14允许补到 2 的幂双调排序比较网络可预生成零分支SIMD/GPU 收益大元素数量 2^16键为整数或可映射整数LSD 基数排序线性复杂度吞吐压倒性优势键为字符串/结构体复杂键并行快排或采样排序无法高效按位分桶需要稳定排序LSD 基数排序或双调复合键稳定性要求高时避免裸用比较网络开局数据规模不是唯一标准关键位宽、稳定性、硬件平台都会改变结论。但大多数场景下先按这个表格选一个大方向不会出现大的偏差。5. 实现中的坑与混合排序的经验这一章写的是我在实际代码里踩过、修过的坑每一条都来源于一次线上排序异常或者性能劣化。它们比算法本身更容易被忽略所以值得单独拿出来讲。5.1 双调排序最容易踩的三个坑第一个坑是补齐哨兵的处理。长度不是 2 的幂时常规做法是往数组头部或尾部插入哨兵元素。最安全的选择是无穷大键这样哨兵会被自然移动到数组末尾真实数据的有序性不受影响。千万别用-1或0这种“看起来无害”的小值当哨兵它们会混进真实数据区间排序结果直接错乱。浮点排序尤其要小心 NaNNaN 参与比较时min/max的行为不可预测建议在进入排序前先清理。第二个坑是 GPU 方向判断错误。静态比较对生成时升序和降序方向由阶段决定。常见错误是把方向简单写成(i j)这在局部小规模测试里可能侥幸通过因为小规模段恰好方向对一旦规模超过某个阈值就全盘错乱。我后来干脆把方向和索引对一起预生成到常量表里运行时不判断任何方向彻底消灭这类 bug。第三个坑是 bank conflict。GPU 共享内存访问时如果同一个 warp 里的线程访问的地址相差正好是 32 个 int 的整数倍就会触发严重的共享内存 bank 冲突双调排序这种按固定跨度访问的模式极易中招。一个简单的规避办法是把数组的 stride 从 N 改成 N1用 padding 打破地址对齐。N 不大时多花的几字节共享内存换来的是带宽翻倍。5.2 基数排序的四个性能杀手第一杀手是全局原子计数。atomicAdd看起来简单但它把所有线程的计数压力集中到一个桶上数据越倾斜竞争越惨烈。正确的两级计数我前面已经拆解过这里不再重复只强调一句不要图省事直接上全局原子。第二杀手是散射阶段的缓存污染。8 位 digit 有 256 个桶当输出数组按桶分散写回时不同桶的地址跨度大缓存行频繁失效。一种改善方式是让每个线程块尽量处理连续的内存段这样每个块内部的写回地址有一定局部性。更深一步的做法是把散射拆成“先按块收集、再全局搬运”两步但多数场景下块连续处理已经够用。第三杀手是无效的高位 pass。如果数据实际有效位只有 16 位却按 32 位跑 4 趟浪费一半时间。前面说的“采样跳过空 digit”就是对应方案。还有就是数字宽度和桶数组的权衡8 位 4 趟 vs 16 位 2 趟后者桶数组变成 65536 个 int光计数数组就 256KB远超 CPU 缓存所以 16 位方案只在 GPU 上用共享内存分摊才划算。一般工程里 8 位是个稳妥默认11 位是进阶选择。第四杀手是稳定性需求突然出现。LSD 基数排序天生稳定这是它对比其他非稳定算法的隐藏优势。但如果你从一开始就选了比较网络或快排路线等到需要稳定性时再改代价很大。所以需求评审时一定要问清楚“排序结果要不要保留原始录入顺序”。5.3 一个实用的混合方案基数分桶加双调收尾单独使用两个算法各有边界但它们可以很好地互补。一个我最终在项目里落地并验证过的方案是先用基数排序按高 16 位分桶把大规模数据切分成若干中等大小的桶每个桶的大小大致是n / 桶数通常落在2^10 ~ 2^13之间然后桶内排序交给双调排序每个桶补齐到 2 的幂后独立排序桶之间天然互不干扰整个排序过程只有一次跨桶协调。这个方案的实际收益是双重的基数分桶阶段把数据按高位切得足够均匀双调排序阶段又在每个桶里构建完全固定的比较网络不需要任何前缀和或原子操作。整体复杂度近似为一次 O(n) 分桶加若干次 O(m log^2 m) 的小块排序比纯基数排序多不了多少时间却保留了双调排序在 GPU block 内的吞吐优势。在某图像处理 Demo 里我们用这个组合把 2000 万条记录的总排序耗时压到了单线程std::sort的十分之一以下。实现时有一个细节不要真的给每个桶都 pad 到同一个次幂大小。如果数据分布不均匀某些桶只有 100 个元素也要补齐到 512 个位置就是白白浪费 80% 容量。正确的做法是把小桶合并成一个大桶或者给小桶选更小的次幂块再用一次双调排序处理。这个“按实际长度选对齐级别”的逻辑不算复杂但对整体性能影响很大。我自己后期的固定搭配是CPU 上大规模排序优先用 8 位 LSD 基数排序GPU 上小规模局部排序优先用双调排序两者之间的过渡用基数分桶完成。这套组合在好几个项目里都没有让我失望过。你也可以先跑一跑本文里的递归版和迭代版代码感受一下双调排序那种“所有比较在同一层同时发生”的节奏以及基数排序那种“一个字节一个字节搬数据”的踏实感然后根据你的数据规模、键类型和硬件平台再决定把重心押在哪一边。