如果你一直在做数据量比较大的排序任务,多半也遇到过这类尴尬:单线程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, up=True): 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为什么要按段号奇偶交替?因为在更大的合并阶段,相邻的段必须保证方向相反,才能让它们拼接成一个更大的双调序列。你可以把size=2阶段想象成给后续每个size=4段分别准备“升序左半 + 降序右半”,方向自然是一正一反。
整个双调排序的比较总次数是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的位模式是0xBF800000,1.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.0003s | 0.0006s | 0.0002s | 0.0005s |
| 2^16 (65536) | 0.008s | 0.002s | 0.010s | 0.001s |
| 2^20 (1048576) | 0.150s | 0.030s | 0.210s | 0.009s |
| 2^24 (16777216) | 3.200s | 0.420s | 2.800s | 0.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这种“看起来无害”的小值当哨兵,它们会混进真实数据区间,排序结果直接错乱。浮点排序尤其要小心 NaN,NaN 参与比较时min/max的行为不可预测,建议在进入排序前先清理。
第二个坑是 GPU 方向判断错误。静态比较对生成时,升序和降序方向由阶段决定。常见错误是把方向简单写成(i < j),这在局部小规模测试里可能侥幸通过(因为小规模段恰好方向对),一旦规模超过某个阈值就全盘错乱。我后来干脆把方向和索引对一起预生成到常量表里,运行时不判断任何方向,彻底消灭这类 bug。
第三个坑是 bank conflict。GPU 共享内存访问时,如果同一个 warp 里的线程访问的地址相差正好是 32 个 int 的整数倍,就会触发严重的共享内存 bank 冲突,双调排序这种按固定跨度访问的模式极易中招。一个简单的规避办法是把数组的 stride 从 N 改成 N+1,用 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 上小规模局部排序优先用双调排序,两者之间的过渡用基数分桶完成。这套组合在好几个项目里都没有让我失望过。你也可以先跑一跑本文里的递归版和迭代版代码,感受一下双调排序那种“所有比较在同一层同时发生”的节奏,以及基数排序那种“一个字节一个字节搬数据”的踏实感,然后根据你的数据规模、键类型和硬件平台,再决定把重心押在哪一边。