假设你管理着一座拥有上万名分拣员的智能物流分拣中心。过去安排工作,你要亲自盯着每一个岗位:包裹放哪层货架、按什么顺序排列、每个区域的包裹量怎么汇总、不同加急等级的包裹怎么分流……这些底层细节极其琐碎,稍有差错整条流水线就堵塞。现在,总部下发了一本标准化作业手册:你只需在指令单上写一句“把所有包裹按编号排序”,上千名分拣员立刻并行执行完毕;再写一句“汇总各区域包裹总量”,各小组自动两两合并、逐级上报。这本手册就是 Thrust——CUDA 官方为 GPU 量身定制的高级算法库。它把 sort 排序、reduce 归约、scan 前缀和、partition 分区、merge 合并等复杂的并行算法,压缩成日常 C++ 程序员也看得懂的几行代码,让 GPU 编程的门槛从“刀耕火种”降为“按图索骥”。
学会 Thrust 高级 API,等于拿到了 GPU 并行算法的“快捷指令集”,从此告别手写每一个并行核函数的苦役,几分钟就能完成过去几天的算法落地工作。
目录
可视化演示
展示Thrust的sort排序与scan前缀和过程、zip迭代器多元数据组合及GPU算法性能对比,理解GPU并行编程核心。
本节配有交互动效(进入视口后自动加载,约 3KB)。
sort排序与scan前缀和 sort_by_key按键排序同时重排关联值数组 scan计算前缀和,支持inclusive和exclusive模式 sort(A) → scan(A) → prefix_sum[i] Thrust排序后计算前缀和,GPU并行加速数据处理
算法超市初览:sort排序
新入职的分拣员第一天首先学会“上架排架”:把一批随机到达的包裹按编号从小到大摆进货架。主管不用教他怎么两两交换、怎么递归拆分,只需在指令单上写下"sort",整套排架流程自动开始。这就是 Thrust 的 thrust::sort——你把一个 device_vector(设备端数组,货架区)的起止位置交给它,它马上调动成千上万个线程并行整理,全程不需要你碰底层归并树或比特排的细节。
Thrust 的排序内部会依据数据规模与特征自动选择并行归并排序或基数排序。降序排列时只需传入比较器 thrust::greater<int>();如果包裹有“编号+目的地”两列数据,thrust::sort_by_key 会以第一列为键排序,同时让第二列的值跟着一起挪动,就像一个人搬箱子时顺手把标签也带走。它的复杂度可用归并递归式刻画:
对比手写排序核函数,sort 的价值正如主管不必培训员工“如何交换、如何分区”——手册已把最优并行策略固化好了。
- 问题规模 \(n\)待排包裹总数量,它决定排序成本的上限。
- 子问题拆分 \(2T(n/2)\)把 \(n\) 份包裹平均切成左右两半,两条并行流水线各自递归排序,耗时翻倍但规模减半。
- 线性合并 \(O(n)\)两半排好后再按顺序归并成完整货架,这一趟扫描与 \(n\) 成正比。
- 递归基只剩一个包裹时无需排序,直接结束递归。
thrust::device_vector<int> box_ids = {42, 7, 19, 3, 88};thrust::sort(box_ids.begin(), box_ids.end());// box_ids: 3, 7, 19, 42, 88
reduce归约:汇总数据的引擎
仓库主管每天要统计各区域包裹总量。他不必派一个人跑遍所有货架逐一相加,而是喊一句"reduce":每个小组内部先两两相加,再把结果上报到上一级,层层合并,最后汇总成一个总数。这就是归约(reduction)——把一段数据通过某个二元操作折叠成一个结果。thrust::reduce 用法极其简洁:
第三个参数是初值(identity)。它不仅是数学上的“零起点”,更决定了返回类型:传入 0 返回 int,传入 0.0f 返回 float。如果忘写初值,接口默认以元素类型的默认构造值开始,对内置数值类型通常是 0,但请注意:对 float 求和时,初值 0 会被隐式转成 int 参与运算吗?不会,因为 Thrust 用模板推导初值类型,所以显式写 0.0f 是好习惯。归约不限于加法,任何满足结合律的操作都可以:
使用自定义操作时,传入函数对象即可,例如 thrust::plus<int>()、thrust::maximum<int>(),甚至可以写自己的仿函数(functor)。这就像主管自定义“汇总口径”——今天按件数、明天按重量。
- 数据元素 \(x_i\)第 \(i\) 个待汇总的值,好比小组里每人的包裹数。
- 二元操作 \(\oplus\)一种满足结合律的运算,如加法、乘法、取最小值,对应主管规定的汇总口径。
- 求和展开把 \(n\) 个值连成一个长串,串的顺序不影响结果,因为结合律允许任意两两分组并行计算。
- 并行执行GPU 把原串看成叶子节点,每轮相邻两个合成一个内部节点,经过 \(\log_2 n\) 轮得到根节点。
int total = thrust::reduce(vec.begin(), vec.end(), 0);scan前缀和:逐站累计的流水账
一条传送带旁有 16 个扫码站,每个站想知道“我到达时,前面已经扫过多少个包裹”。朴素做法是第 1 站等第 0 站传数字,第 2 站等第 1 站……这串行等待非常浪费。scan(前缀和)算法让所有站点同时拿到答案:每个站记录“前面所有包裹之和”作为自己的前缀值。thrust::inclusive_scan 与 thrust::exclusive_scan 的区别是:inclusive 把当前位置的值也算进前缀中,exclusive 则只计算之前所有元素的和(首元素恒为初值)。
前缀和的核心公式如下,exclusive 版本:
scan 最经典的用途是任务分配:每个 GPU 线程处理长度不等的碎片数据,用 exclusive_scan 算出每个碎片的起始偏移量,线程便知道从哪个位置开始干活。就像扫码站知道了自己前面的包裹总数,才可以安排后续货架的落位。
- 索引 \(i\)第 \(i\) 个站点,\(i\) 从 0 开始计数。
- 前 \(i\) 项求和exclusive 表示只累加排在当前站之前的元素,不含自己。
- 首元素初值第 0 站前面什么都没有,所以前缀等于传入的 identity(如 0)。
- 相邻递推每站前缀只需在前一站基础上多并入一个元素,这正是并行扫描的“波前传播”思想。
thrust::inclusive_scan(input.begin(), input.end(), output.begin());// input: 1 2 3 4// output: 1 3 6 10 thrust::exclusive_scan(input.begin(), input.end(), output.begin(), 0);// output: 0 1 3 6
partition分区:按规则分流的传送带
分拣中心经常要按“是否加急”把包裹分成两条传送带,一条送往飞机,一条送往货车。partition 就是干这个的:接受一个谓词(判断规则),把满足条件的元素挪到前面,不满足的挪到后面。thrust::partition_copy 则更彻底——它不原地打乱,而是把两类分别写入两个不同的输出容器,整个过程并行且不改变各自内部相对顺序(配合 thrust::stable_partition 可保证稳定的相对顺序):
分区操作的行为可以用集合论描述:把所有元素按谓词 \(f\) 分成两个子集,前一个子集是使 \(f\) 为真的元素,后一个是使 \(f\) 为假的元素:
在图形学中,partition 常用来把发射的光线分为“击中物体”与“未击中物体”两组,分别走不同着色管线;在数据库算子中,它负责按哈希桶把记录分流到不同分区。一条指令,两条传送带,GPU 千名分拣员同时开工。
- 谓词 \(f(x)\)判断规则,输入一个元素输出布尔值,例如“是否加急”。
- 真子集满足规则的元素组成第一部分,它们从原传送带被搬到加急通道。
- 假子集不满足规则的元素组成第二部分,进入普通通道。
- 两部分拼接两个子集内部各自保持采样顺序(stable 版本),整体结果是原数组的一个重排。
thrust::partition_copy(items.begin(), items.end(), urgent.begin(), normal.begin(), [](const Box& b) { return b.priority == URGENT; });
merge合并:有序汇流的双通道
两条传送带各自已经拍好序——A 带按编号从小到大,B 带也如此。现在要把它们汇合成一条依然有序的带子。如果直接倒进同一个池子再重新排序,那太浪费了;正确做法是用 thrust::merge 同时从两条带子头部取值,谁小谁先出,一路并行推进:
这个操作的前提是输入必须已经有序,因为合并是“归并排序”的收尾环节,它利用了两个序列各自的局部有序性,每步只需比较两个头部。thrust::merge_by_key 则额外携带键值对:键合并时对应的值也一并移动,类似 sort_by_key 的配套做法。合并结果长度等于两序列长度之和,因此输出容器需要预先分配足够的空间。用符号表示,合并操作把两个有序序列 \(A\) 与 \(B\) 交织成一个新的有序序列:
merge 在工业界是“多路归并”的基石:搜索引擎把各分片返回的有序结果合并成最终排名,日志系统把多台服务器的有序日志流汇成全局时间线。两条有序的传送带,不需要打乱重排,只要一个巧妙的汇流阀。
- 有序前提\(A\) 内部与 \(B\) 内部都已满足 \(a_1 \le a_2 \le \cdots\)、\(b_1 \le b_2 \le \cdots\)。
- 结果长度合并后元素总数是两序列元素数之和。
- 全局有序结果序列中任意相邻两元素满足 \(z_k \le z_{k+1}\),即整体保持递增。
- 逐位推进GPU 并行比较 A 与 B 中对应区间的头部,选出较小者填入结果,再将指针后移。
thrust::merge(a.begin(), a.end(), b.begin(), b.end(), result.begin());定制规则与多元数据:比较器、算子与zip
分拣中心每天都有新玩法:今天按包裹重量排序,明天按客户等级排序,后天要求“重量相同则按优先码再比”。Thrust 允许你自定义比较器——一个返回布尔值的函数对象,作为 sort 的第三个参数:
同理,reduce 与 transform 的“操作算子”也可自定义:写一个仿函数把两个元素合并成一个。此外,现实中数据往往是多元的——收件人、地址、电话各放一个数组。这时可以用 thrust::make_zip_iterator 把多个迭代器“捆”成一个元组迭代器,每次解引用拿到一个 thrust::tuple,再通过 thrust::get<0>、thrust::get<1> 取出对应字段:
用数学语言描述 zip 操作:两个等长序列按位置配对生成一个新的“复合序列”:
这就像分拣员把“包裹面单的三联”订在一起处理——面单、标签、回执虽然分属三个盒子,但 zip 让它们永远同步移动,不会出现“名字和地址错位”的尴尬事故。
- 两个输入序列\(A\) 存放收件人,\(B\) 存放地址,长度都是 \(n\)。
- 按位配对相同下标的元素组成一个二元组,相当于把两张表单钉在同一行。
- 结果集合zip 的新序列长度仍为 \(n\),每个元素是一个复合数据,供后续排序、归约或 transform 统一处理。
struct by_weight_then_priority { __host__ __device__ bool operator()(const Box& a, const Box& b) const { if (a.weight != b.weight) return a.weight < b.weight; return a.priority > b.priority; }};thrust::sort(boxes.begin(), boxes.end(), by_weight_then_priority());
auto zip_begin = thrust::make_zip_iterator( thrust::make_tuple(names.begin(), addresses.begin(), phones.begin()));auto t = *zip_begin; // thrust::tuple 实例std::string n = thrust::get<0>(t);
混合与性能:跨越流水线的协奏曲
再好的标准流水线也有覆盖不到的特殊包裹——比如需要三角运算、查表或者复杂分支逻辑。此时标准库函数不够用,你需要自己写 CUDA kernel。Thrust 提供了“混合编程”的接口:用 thrust::raw_pointer_cast 把 device_vector 的内部设备指针取出来,传给自定义 kernel;kernel 算完后再把指针包装回 device_vector 继续使用 Thrust 的算法。也可以反向操作:在自定义 kernel 内部直接调用 thrust::sort(thrust::device, ptr, ptr + n),让库函数替你完成局部排序。
混合编程并非免费,每次调用 Thrust 算法都有内核启动与同步开销。总耗时可分解为几个阶段的和:
因此,当数据规模很小时,分配与传输的开销可能超过计算本身,这时 CPU 的 std::sort 反而更快——就像一个只有三个包裹的货架,你没必要动用上千人分拣。经验法则是:数据量达到百万级,Thrust 的优势才会充分显现;频繁小规模调用时应尽量合并计算,避免“每调一次库就搬运一次数据”。
- 内存分配 \(T_{\text{alloc}}\)创建设备端的货架(device_vector)需要向 GPU 申请内存。
- 数据搬运 \(T_{\text{transfer}}\)把输入数据从 CPU 内存迁到 GPU 显存,通常走 PCIe 总线,是大头开销。
- 内核执行 \(T_{\text{kernel}}\)GPU 上千线程并行计算的实际耗时,通常远小于单核 CPU。
- 同步等待 \(T_{\text{sync}}\)主机等待设备完成、检查错误的隐式同步耗时。
int* raw_ptr = thrust::raw_pointer_cast(vec.data());my_kernel<<<grid, block>>>(raw_ptr, n); // 自写核函数处理thrust::sort(thrust::device, raw_ptr, raw_ptr + n); // 再交给库
AI 应用场景
大规模特征排序与 Top-K 筛选
阶段:数据预处理 / 特征工程
解决的问题:推荐系统每天要对上亿用户产生候选特征向量,按预估点击率排序截取 Top-K。传统做法是在 CPU 上跑单线程排序,耗时数秒无法满足在线延迟要求;Thrust 的 thrust::sort 配合 sort_by_key 直接在 GPU 上完成全量排序,耗时降到毫秒级。
具体案例:某电商推荐服务,模型对 1 亿条商品特征打分后,把分数与商品 ID 放入两个 device_vector,调用 thrust::sort_by_key 按分数降序排列,再取前 1000 条返回,端到端耗时从 2.1 秒降到 38 毫秒。
物理模拟中的并行力汇总
阶段:模型训练前的仿真数据生成
解决的问题:分子动力学模拟中,每个粒子受周围数千邻居的合力,汇总所有相互作用力是典型的大规模归约(reduce)任务。手写归约 kernel 容易出错且难调优,thrust::reduce / transform_reduce 提供现成的并行归约方案。
具体案例:某蛋白质折叠仿真,将 320 万对粒子间的相互作用力分量拆分到 GPU 各线程局部求和,再用 thrust::reduce 做树形合并,每帧受力更新从 CPU 的 900 毫秒降到 12 毫秒,迭代速度提升 70 倍。
多路日志流有序归并
阶段:大数据流水线后处理
解决的问题:分布式系统中各节点产生带时间戳的日志流,需要按全局时间合并成一条有序记录。多路 merge 是典型应用场景,thrust::merge 可同时合并两条已各自有序的日志流,merge_by_key 还能保证“时间戳+日志内容”同步移动。
具体案例:某实时风控平台把 16 台采集服务器的日志先各归各排序,再用两级 thrust::merge 树合并成全局时间序列,融合耗时从 5 秒降至 200 毫秒,支撑每秒 120 万条日志的实时归并分析。
Python 演示
import random def parallel_reduce(values, op=lambda a, b: a + b): """模拟 GPU 并行归约:每轮相邻两两合并。""" arr = values[:] round_no = 0 print("初始:", arr) while len(arr) > 1: next_arr = [op(arr[i], arr[i + 1]) for i in range(0, len(arr) - 1, 2)] if len(arr) % 2 == 1: next_arr.append(arr[-1]) arr = next_arr round_no += 1 print("第 %d 轮:" % round_no, arr) return arr[0] def exclusive_scan(values, identity=0): """模拟前缀和:每个位置保存之前所有元素之和。""" result, acc = [], identity for v in values: result.append(acc) acc += v return result if __name__ == "__main__": data = [random.randint(1, 20) for _ in range(16)] print("总和:", parallel_reduce(data)) print("前缀和:", exclusive_scan(data))
while len(arr) > 1:只要还有两个以上元素就继续下一轮合并,对应 GPU 归约的递归树形结构。[op(arr[i], arr[i + 1]) for i in range(0, len(arr) - 1, 2)]相邻两两配对并应用二元操作,模拟并行线程同时处理多个对。if len(arr) % 2 == 1: next_arr.append(arr[-1])奇数长度时把最后一个元素直接透传到下一轮,对应 GPU 归约中的“落单处理”。acc += v前缀和的核心递推:后一个位置的结果由前一个位置的累计值加上当前元素得到。result.append(acc)exclusive_scan 先记录当前累计值再累加,确保第 i 个结果不包含第 i 个元素本身。
常见误区
误区“Thrust 不就是给 STL 套了一层壳吗?sort 内部肯定还是调用 std::sort,把数据传到 CPU 排完再送回 GPU。”
正解Thrust 的算法完全在 GPU 设备端执行,sort 内部根据数据规模动态选择并行归并排序或基数排序,由 GPU 上千个线程协作完成,不存在把整批数据拷回 CPU 再排序的过程。只有当你传入 host_vector 的迭代器时,它才会走主机路径;只要传入 device_vector 的迭代器,全程都在显存内运行。
为何常见:Thrust 的调用习惯(begin/end 迭代器)与 C++ STL 高度相似,初学者容易把两者底层实现混为一谈。
误区“thrust::reduce(v.begin(), v.end()) 不写第三个参数也默认是 0,整型浮点型都能用,不用管。”
正解初值不仅参与运算,还决定返回类型。对 float 数据不写初值时,模板从元素类型推导出 identity,通常是 T() 即 0.0f;语义上没问题,但一旦自定义归约操作为 thrust::minimum<int>(),初值应为 INT_MAX 而非 0,否则结果恒为 0。建议始终显式传入与业务语义匹配的初值。
为何常见:加法归约中初值 0 恰好正确,让很多人误以为初值只是“可填可不填”的装饰。
误区“make_zip_iterator 打包后,*iter 返回的就是一个结构体,直接 .field 访问成员就行。”
正解zip 迭代器解引用返回的是 thrust::tuple,必须用 thrust::get<0>(*iter)、thrust::get<1>(*iter) 提取对应字段,不能使用 .field 语法。另外,zip 要求各序列长度相同,超出最短序列的尾部元素不会被访问。
为何常见:Python 的 zip 返回可解包的元组,C++ 用户顺理成章地以为也能自动解包,忽略了 Thrust tuple 是显式模板类型。
误区“既然 Thrust 是官方库,算法肯定都是最优的,我只要无脑调用就行,肯定比我手写的快。”
正解Thrust 胜在通用与开发效率,但极致性能需要结合具体场景。数据量很小时,kernel 启动开销(微秒级)与内存分配反而主导耗时,此时简单 CPU 排序更快;而对高度定制化的细粒度计算,手写 kernel 可以针对性优化访存模式,超过通用实现。
为何常见:教程中 Thrust 动辄展示“百万数据几毫秒”的效果,容易让人忽略小数据规模与传输开销的边界条件。
面试题
Q1. thrust::sort(vec.begin(), vec.end()) 对 device_vector<int> 操作时,数据在什么位置完成排序?
- A.CPU 内存中完成后再拷回 GPU
- B.GPU 显存中并行完成← 正确答案
- C.CPU 与 GPU 各排一半合并
- D.由 CUDA 驱动决定位置
答案:B。传入 device_vector 的迭代器时,Thrust 全程在 GPU 设备端执行,数据不会拷回 CPU。选项 B 正确。A 是典型误解(Thrust 不等于 std::sort),C、D 没有这样的机制。
Q2. thrust::reduce(v.begin(), v.end(), 0) 的返回值类型是什么?
- A.迭代器类型
- B.与 v 元素相同的类型
- C.由第三个参数初值决定← 正确答案
- D.固定为 int 类型
答案:C。归约结果类型由初值参数推导。传 0 返回 int,传 0.0f 返回 float,传自定义对象返回该对象类型。因此 C 正确。A、B 都与 Thrust 模板推导机制不符,D 过于绝对。
Q3. exclusive_scan 与 inclusive_scan 的核心区别是?
- A.exclusive 不包含当前元素← 正确答案
- B.inclusive 不包含当前元素
- C.两者完全相同
- D.exclusive 只支持加法
答案:A。exclusive_scan 计算结果前缀时排除当前位置的元素,首元素为初值(如 0);inclusive_scan 则把当前位置元素算入结果。A 正确,B 描述反了,C、D 均错误。
Q4. thrust::sort_by_key(keys.begin(), keys.end(), vals.begin()) 的作用是?
- A.只对 keys 排序,vals 不动
- B.对 keys 排序,vals 中与 keys 绑定的值同步移动← 正确答案
- C.对 vals 排序,keys 同步移动
- D.对 keys 和 vals 分别独立排序
答案:B。sort_by_key 以第一段范围为键排序,第二段范围作为“值”跟随键同步移动,保证键值对应关系不变。B 正确。A 是 sort 的语义,C 说反了主次,D “分别独立排序”会破坏对应关系。
Q5. 以下哪个是 thrust::partition_copy 的必要条件?
- A.输入必须有序
- B.需要两个输出容器的迭代器← 正确答案
- C.谓词必须返回整数
- D.只能用于 device_vector
答案:B。partition_copy 按谓词把元素分写到两个独立输出范围,分别承接 true 与 false 结果,因此至少需要两个输出迭代器。B 正确。A 是 merge 的前提,C 谓词返回 bool 即可,D 也可用于 host_vector。
Q6. 关于 thrust::merge,下列说法正确的是?
- A.两个输入序列可以无序
- B.输出容器长度可小于两序列长度之和
- C.两个输入序列必须各自有序← 正确答案
- D.merge 等价于 sort
答案:C。merge 利用输入序列各自有序的性质做线性归并,若输入无序则结果不保证有序,因此 C 正确。A 错误,B 会越界写,D 混淆了排序与合并。
Q7. 数据规模很小时,Thrust 排序可能比 CPU 的 std::sort 更慢,最主要原因是?
- A.Thrust 算法复杂度更高
- B.CPU 计算能力比 GPU 强
- C.GPU 并行度在小规模下无法发挥
- D.内核启动与内存传输开销主导← 正确答案
答案:D。小规模数据时 GPU 无法吃饱并行度,同时每次调用要承担 kernel 启动、设备内存分配与潜在的 PCIe 传输开销,这些固定成本超过排序本身的收益。D 正确。A 复杂度相同或更低,B、C 是表面现象而非主因。
Q8. 自定义比较器传给 thrust::sort 时,仿函数的 operator() 应标注哪些限定符才能在设备端被调用?
- A.
__host__ - B.
__device__ - C.
__host__ __device__← 正确答案 - D.不需要任何限定符
答案:C。Thrust 算法可能在 GPU kernel 内部调用比较器,因此仿函数 operator() 必须同时标注 __host__ __device__,从而既可在主机编译也可在设备编译。C 正确。只标一个或都不标会出现编译错误。
Q9. 关于 thrust::make_zip_iterator,下列说法正确的有?
- A.可把多个迭代器绑定成一个复合迭代器← 正确答案
- B.解引用时返回 thrust::tuple← 正确答案
- C.各输入序列要求等长← 正确答案
- D.只能用于整型数据
答案:ABC。zip 迭代器将多个迭代器按位置绑定,解引用得到 tuple,且要求序列等长否则超出部分不访问,因此 A、B、C 都正确。D 错误,zip 支持任意 POD 类型的序列,与元素类型无关。
Q10. 下列哪些关于 Thrust 混合编程的说法是正确的?
- A.raw_pointer_cast 可取出 device_vector 内部指针← 正确答案
- B.自定义 kernel 与 Thrust 可互相调用← 正确答案
- C.每次 Thrust 调用都有内核启动开销← 正确答案
- D.混合编程没有任何额外成本
答案:ABC。raw_pointer_cast 确实能拿到设备端裸指针供自定义 kernel 使用,kernel 内也可调用 Thrust 算法,且每次调用自带启动与同步开销,因此 A、B、C 正确。D 错误,混合编程有 kernel 启动、同步与数据搬运的隐形成本,规模过小时可能得不偿失。