Thrust库：CUDA版STL

可视化演示

Thrust初识

你刚接管一间大后厨，听老师傅说，以前做菜得自己从磨刀学起——对应着 CUDA 编程里手写 kernel、管理线程块和网格的繁琐步骤。但你真正想做的是“用一万份食材做出同样味道的招牌菜”。这时老厨师告诉你：别慌，我们有一位厨师长叫 Thrust。你只管把菜单递给他，他负责让一千个帮厨在同一秒钟同时开工。

Thrust 是 NVIDIA 在 CUDA 之上封装的高层模板库，头文件放在 <thrust/...> 里。它的接口刻意模仿 C++ 标准库：正如 std::vector 之于 CPU，thrust::device_vector 之于 GPU；正如 std::sort 之于数组，thrust::sort 之于设备内存。你不需要手动分配显存、指定网格尺寸、写 __global__ 函数，只需要描述“要做什么”。

一次 thrust::sort(d_vec.begin(), d_vec.end()) 调用，库内部会自动选择合适的 kernel 并调度所有线程。底层依赖 C++ 模板与 functor/lambda 机制，所以你的代码风格从“指挥每一步”变成了“描述意图”——这正是厨师长存在的意义：菜单上写什么，他就做什么，怎么做是他的事。

$$\text{加速比}_{\text{理想}} = \frac{T_{\text{单核}}}{T_{\text{并行}}} \approx N_{\text{线程}}$$

1. T_单核（串行耗时）：一个厨师按照菜谱从头做到尾的总时间。

$$T_{\text{单核}}$$

2. T_并行（并行耗时）：厨师长把菜谱拆给 N 个帮厨后，大家同时开工完成的总时间。

$$T_{\text{并行}}$$

3. N_线程（帮厨人数）：GPU 上同时干活的线程数量，通常成千上万。

$$N_{\text{线程}}$$

4. 加速比（快了多少倍）：串行时间除以并行时间，理想情况下近似等于线程数 N。

$$\frac{T_{\text{单核}}}{T_{\text{并行}}} \approx N_{\text{线程}}$$

当然，理想的加速比不会真正拿到——帮厨之间要传话、锅要轮流用，这就是后文会遇到的传输开销与调度开销。但至少，你把“指挥一千双手”这件事，完全交给了厨师长。

设备向量

开一家餐厅，第一件事是找地方存食材。CPU 侧的 std::vector 是“一楼前台冰箱”（主机内存）；GPU 侧的 thrust::device_vector 是“后厨冷藏柜”（显存）。冷藏柜紧挨着炒锅（处理器），帮厨取料最快——但你别隔三差五把食材从一楼搬到二楼，来回跑楼梯可要命了。

device_vector 的用法几乎和 std::vector 一模一样：thrust::device_vector<int> v(1024); 创建一个能装 1024 个整数的冷藏柜；v.size() 看容量；v[0] 访问单个元素；thrust::fill(v.begin(), v.end(), 7) 把所有格子都填上 7。乍一看毫无 GPU 痕迹，这就是 STL 风格接口的魅力。

真正需要注意的是数据搬运。thrust::host_vector 对应主机内存，thrust::device_vector 对应显存，两者之间用 thrust::copy 同步。每次 CPU↔GPU 拷贝都相当于把食材从一楼搬上二楼再搬下来，PCIe 带宽有限，来回跑非常费时间。所以最佳实践是：数据一次搬进 GPU，连续完成多步算法，最后再搬回一次。

$$\text{搬运占比} = \frac{T_{\text{H2D}} + T_{\text{D2H}}}{T_{\text{计算}} + T_{\text{H2D}} + T_{\text{D2H}}}$$

1. T_H2D（主机到设备拷贝）：把食材从一楼前台搬进后厨冷藏柜的时间。

$$T_{\text{H2D}}$$

2. T_D2H（设备回传主机）：把成品从后厨端回一楼前台的时间。

$$T_{\text{D2H}}$$

3. T_计算（真正处理的时间）：帮厨们实际切菜炒菜所用的时间。

$$T_{\text{计算}}$$

4. 搬运占比（开销比例）：如果这个比值接近 1，说明你一直在搬货而不是做菜——这也是 GPU 菜鸟最常见的性能浪费。

$$\frac{T_{\text{H2D}} + T_{\text{D2H}}}{T_{\text{计算}} + T_{\text{H2D}} + T_{\text{D2H}}}$$

记住：设备向量是后厨的冷藏柜。食材住得越久、来回搬的次数越少，菜（程序）就上得越快。让数据“住”在设备端，是玩转 Thrust 的第一条心法。

核心算法

厨师长最拿手的三道主厨招牌菜：变换、归约、排序。你只需要报菜名，他就指挥全体帮厨按分工行动。这三道菜分别对应 thrust::transform、thrust::reduce、thrust::sort，是绝大多数 GPU 任务的“三板斧”。

transform 变换

transform 对每个食材（元素）独立执行同一个操作——就像把一万根胡萝卜同时放进切丝机，每根出来都是丝。它适合做逐元素运算：取绝对值、缩放、归一化、套激活函数。因为每个输出只依赖对应的一个输入，天然并行，GPU 上轻松跑满。

$$\displaystyle y_i = f(x_i), \qquad i = 0, 1, \dots, n-1$$

1. x_i（第 i 个输入）：第 i 份原始食材，来自输入容器的当前位置。

$$x_i$$

2. f（处理函数）：切丝机，由用户自定义，比如取平方、加偏置、sigmoid 激活。

$$f(\cdot)$$

3. y_i（第 i 个输出）：处理之后的食材，与输入一一对应，输出容器同位存放。

$$y_i$$

4. i 从 0 到 n-1（遍历全部）：所有元素同时处理，顺序无关紧要，所以可以彻底并行。

$$i = 0, 1, \dots, n-1$$

reduce 归约

reduce 把许多食材“炖”成一锅汤：求和、求最大值、求均值都用它。GPU 上的秘诀是分层的两两相加——先让 N 个帮厨两两配对，把结果合并，只剩一半，再继续配对合并，直到炖成一锅。这样并行深度只有约 log₂(n) 轮，而不是串行的 n 次加法。

$$\displaystyle \text{sum} = \sum_{i=0}^{n-1} x_i = ((x_0 + x_1) + (x_2 + x_3)) + \cdots$$

1. x_i（第 i 个元素）：第 i 份食材，是归约的输入。

$$x_i$$

2. 求和符号（Σ i=0 到 n-1）：把所有食材都倒进同一口锅里。

$$\sum_{i=0}^{n-1}$$

3. 两两相加的展开式：并行归约的秘诀——两两同时进行加法，而不是一个一个累加。

$$((x_0 + x_1) + (x_2 + x_3)) + \cdots$$

4. 并行深度（约 log₂ n）：串行要 n-1 次加法，并行只需要约 log₂(n) 轮合并。

$$O(\log n)$$

sort 排序

sort 按大小把食材码放整齐。Thrust 底层封装了高度调优的排序 kernel（CUDA 后端常基于 CUB 的基数排序与归并排序组合），对大规模数组，性能接近手写最优实现。调用只有一行，复杂度为经典的 O(n log n)。

$$\displaystyle T(n) = O(n \log n)$$

1. T(n)（耗时随规模增长）：把 n 份食材排整齐所需的操作次数。

$$T(n)$$

2. O(n log n)（复杂度级别）：规模翻倍时，耗时略超两倍；这是通用比较排序的理论下界附近。

$$O(n \log n)$$

3. 托底保障：Thrust 不要求你关心内部用基数排序还是归并排序，只保证这个复杂度级别。

$$n \log n$$

三道菜可以自由组合：饭店流水线常常“先变换、再归约、最后排序取 TopK”。厨师长对每条流程都胸有成竹，你只管点菜。

迭代器体系

帮厨们怎么知道“这份是牛肉、那份是豆腐”？靠的是食材上的标签纸。Thrust 的迭代器就是贴在数据上的标签。device_vector 返回的迭代器内部包装了设备指针（thrust::device_ptr），它说明“这份食材在后厨冷藏柜的哪个架子上”——比裸指针更安全，因为它知道数据类型和所属容器。

除了普通迭代器，Thrust 还有几张“魔法标签”：constant_iterator 让每个位置都读到同一个常数，就像所有食材都贴着“特价 5 元”；counting_iterator 自动生成 0、1、2、3…… 的连续序列，省掉你手动造一个下标数组；zip_iterator 则把多个容器的对应元素打包成元组，像把“菜名、价格、图片”三张标签贴在同一份菜单条目上——一次 transform 就能同时处理三份相关数据。

这些特殊迭代器的妙处在于：它们只描述“数据长什么样”，并不会真的把数据搬来搬去。算法按需取值，延迟生成序列，省掉了大量临时数组和中间内存。

$$\text{zip}(a_i, b_i) \mapsto (a_i, b_i) \quad (i=0,1,\dots,n-1)$$

1. a_i、b_i（两个序列）：两张不同的标签，例如价格序列和库存序列。

$$a_i, \ b_i$$

2. zip（打包）：把同一下标的两个元素合成一个元组，这就是 zip_iterator 的行为。

$$\text{zip}(a_i, b_i)$$

3. 映射结果 (a_i, b_i)：一次遍历里同时拿到两组信息，避免反复访问多个容器。

$$(a_i, b_i)$$

4. i 的范围：要求两个序列长度一致，配对必须一一对应，否则越界。

$$i = 0, 1, \dots, n-1$$

有了标签纸，厨师长不需要问你“这份在哪儿”，只需要说“按标签干活”，帮厨就能精准取料、毫不含糊。

执行策略

同样一道菜，你可以让一个帮厨慢慢做，也可以让一百个帮厨同时做——这就是执行策略。Thrust 提供 thrust::seq、thrust::par、thrust::device 三档“人手配置”，作为算法的第一个参数传入，例如 thrust::sort(thrust::device, first, last)。

thrust::seq 表示串行执行：一个帮厨从头做到尾，适合小数据量或调试时看过程。thrust::par 表示允许并行调度，配合后端决定实际执行方式。thrust::device 则明确指定在 GPU 设备上执行——相当于写清楚“这道菜必须在后厨大灶上做”。注意，策略决定的是“怎么执行”（串行还是并行），迭代器指向的设备内存才决定“在哪里取食材”。

策略选择本质上是一笔经济账：调度本身有固定成本。数据量小时，喊话、分派、汇总结果的开销比干活本身还大，不如一个帮厨直接做完；数据量大时，并行分摊的优势才凸显出来。存在一个“拐点”，越过它并行才划算。

$$T_{\text{总}} = T_{\text{调度}} + \frac{T_{\text{工作}}}{N_{\text{线程}}}$$

1. T_总（实际总耗时）：从下单到上菜的真实时间。

$$T_{\text{总}}$$

2. T_调度（调度开销）：厨师长喊话、分派任务、汇总结果的固定成本，与数据量无关。

$$T_{\text{调度}}$$

3. T_工作 / N_线程（分摊后工作量）：工作总量被 N 个帮厨平分后的理想耗时。

$$\frac{T_{\text{工作}}}{N_{\text{线程}}}$$

4. 拐点判断：当 T_调度 远超分摊后的工作量时，并行反而更慢。

$$T_{\text{调度}} \gg \frac{T_{\text{工作}}}{N_{\text{线程}}}$$

人手多不一定更快——如果任务只有一颗土豆，十个帮厨抢着削反而添乱。理解执行策略，就是学会算清楚这笔账。

自定义算子

厨师长的固定菜谱不够用的时候，你需要给他一张即写即用的配方卡：“每份牛排上撒 5 克盐”。Thrust 里的配方卡就是自定义算子，可以用 functor（仿函数）或 lambda（匿名函数）来写。

functor 是一个重载了 operator() 的 struct / class。比如定义一个“平方器”：struct square { __device__ float operator()(float x) { return x * x; } };，然后传给 thrust::transform。C++11 之后更常用 lambda，一行搞定：thrust::transform(dv.begin(), dv.end(), dv.begin(), [](float x){ return x * x; });。lambda 其实就是编译器帮你生成的匿名 functor，殊途同归。

使用时有几个硬性注意点：在 GPU 上执行的算子必须用 __host__ __device__ 标记；lambda 捕获变量必须按值捕获（设备端访问不了主机内存）；不能用 std::vector 这类主机容器去捕获。想象一下：配方卡是发给千里之外的帮厨的，卡上只能写“盐 5 克”这种可复制的内容，不能写“去一楼冰箱拿”——他根本够不着。

$$\text{lambda}: \quad x \mapsto x^2 + 1$$

1. x（入参）：传入的每个元素，相当于“每份食材”本身。

$$x$$

2. x² + 1（规则）：配方卡上写的具体操作，对应 lambda 的函数体。

$$x^2 + 1$$

3. ↦（映射箭头）：表示从输入到输出的对应关系，是变换算法的核心语义。

$$\mapsto$$

4. lambda 与 functor 的关系：lambda 会被编译器展开成一个匿名的 functor 类，本质相同。

$$\text{lambda} \equiv \text{functor}$$

配方卡要写得清楚、写得短。所有帮厨都会照着它执行，写错一处就是一万份菜全部做错——所以自定义算子要精心设计，先在 CPU 上验证逻辑，再搬上 GPU。

性能对比

有人觉得手写 kernel 像“自己挽袖子进厨房”，更精细、更可控；Thrust 像“交给厨师长”，省事但怕大锅饭不好吃。真实对比到底如何？结论是：看场景。

对 sort、reduce、scan 这类已有成熟并行算法的操作，Thrust 底层调用高度调优的实现（CUDA 后端基于 CUB），性能通常与手写最优接近，甚至更优——毕竟库作者花了多年打磨，还针对不同数据规模做了分派。手写 kernel 的赢面在于：你可以针对特定访存模式、特定数据分布（比如稀疏、重复度高）做极致定制，避开 Thrust 的通用模板开销。

另一个被忽略的成本是“启动开销”：小规模、多次启动的场景下，每次调用都要付出调度成本，手写一个融合的大 kernel 可能更划算。Amdahl 定律为这一切划定了上限：无论你有多少线程，串行部分始终卡住加速比。

$$S = \frac{1}{(1-P) + \frac{P}{N}}$$

1. P（并行比例）：整个流程中能被帮厨们同时做的那部分占比。

$$P$$

2. 1 - P（串行比例）：必须一个人完成的部分，比如拷贝、汇总结果、依赖步骤。

$$1 - P$$

3. N（线程数）：帮厨人数，也就是参与并行的处理器数量。

$$N$$

4. S（加速比）：最终快了多少倍；当 N 趋近无穷时，S 收敛到 1/(1-P)，串行部分成了天花板。

$$S = \frac{1}{(1-P) + \frac{P}{N}}$$

厨师长的价值在于“用 80% 的代码量拿到 90% 的手写性能”。当性能瓶颈清晰、访存模式特殊时，再亲自下场手写不迟——但默认，请相信厨师长。

实战应用

最终，一切要落到一套完整的后厨流程上：下单（CPU 数据）→ 搬进冷藏柜（device_vector）→ 厨师长做菜（算法链）→ 端回前台（回传 CPU）。把这套流程固化下来，就是 Thrust 的典型工作流。

实战中最常见的模式是：先用 thrust::copy 把原始数据从 host_vector 拷进 device_vector；然后链式调用 transform（归一化）→ reduce（求统计量）→ sort（排序）→ 截断取 TopK；最后把结果拷回主机。每个环节都尽量留在设备端，避免中途回传。比如：

常见陷阱也值得记住：频繁启动小数组，调度开销会吞噬并行收益；忘记 copy 就直接用主机指针访问设备内存，程序崩溃或读到垃圾值；lambda 捕获不当导致编译错误。调试技巧也很简单：把设备结果拷回主机分段打印，缩小范围对比；用 cuda-memcheck 检查越界访问。把“最小复现”思想用在 GPU 上，问题往往一目了然。

$$\text{CPU 数据} \xrightarrow{\text{copy}} \text{GPU 显存} \xrightarrow{\text{transform/sort/reduce}} \text{GPU 结果} \xrightarrow{\text{copy}} \text{CPU}$$

1. CPU 数据（前台订单）：最初位于主机内存的原始数据。

$$\text{CPU 数据}$$

2. → copy（搬运）：用 thrust::copy 把数据搬进设备向量，这是唯一一次大搬运。

$$\xrightarrow{\text{copy}}$$

3. GPU 显存（后厨）：搬入之后，所有算法都在设备端连续执行，不回传。

$$\text{GPU 显存}$$

4. → copy 回 CPU（上菜）：把最终结果传回主机，整个过程只搬两次。

$$\xrightarrow{\text{copy}} \text{CPU}$$

把这套流程练熟，你就从“只会自己切菜的初学者”变成了“能指挥整间厨房的掌勺”。以后无论排序、统计还是推荐打分，都可以沿用同一套点菜方式——这就是 Thrust 带给你的职业底气。

thrust::device_vector<float> scores = host_scores;thrust::transform(scores.begin(), scores.end(), scores.begin(), normalize);float mean = thrust::reduce(scores.begin(), scores.end()) / scores.size();thrust::sort(scores.begin(), scores.end(), thrust::greater<float>());

AI 应用场景

Python 演示

# -*- coding: utf-8 -*-"""模拟 Thrust 风格链式调用: transform -> reduce -> sort"""import math  class DeviceVector:    """极简版 thrust::device_vector 模拟: 数据常驻"设备端"(这里就是 list)"""     def __init__(self, data):        self._data = list(data)     def transform(self, fn):        # 对应 thrust::transform, 逐元素应用函数        return DeviceVector([fn(x) for x in self._data])     def reduce(self, fn):        # 对应 thrust::reduce, 两两归约成一个值        acc = self._data[0]        for x in self._data[1:]:            acc = fn(acc, x)        return acc     def sort(self, reverse=False):        # 对应 thrust::sort        return DeviceVector(sorted(self._data, reverse=reverse))     def to_host(self):        # 对应 thrust::copy 回 host        return self._data  def main():    scores = [0.2, 8.1, 5.5, -1.3, 9.9, 3.2]    # 数据搬上 GPU(设备向量)    dv = DeviceVector(scores)    # transform: 逐元素求 exp, 相当于配方卡    probs = dv.transform(lambda x: math.exp(x))    # reduce: 求和后再归一化    total = probs.reduce(lambda a, b: a + b)    norm = probs.transform(lambda p: p / total)    # sort: 按概率从大到小排序    ranked = norm.sort(reverse=True)    # 数据拷回 CPU(主机向量)    top3 = [round(p, 4) for p in ranked.to_host()[:3]]    print("softmax 后概率 Top-3:", top3)    assert abs(sum(ranked.to_host()) - 1.0) < 1e-6  if __name__ == "__main__":    main()

1. class DeviceVector:模拟 thrust::device_vector，数据在“设备端”常驻。
2. def __init__(self, data): self._data = list(data)底层存储就是一份数据列表。
3. def transform(self, fn):对应 thrust::transform，逐元素应用函数。
4. return DeviceVector([fn(x) for x in self._data])生成一个新的设备向量，返回新向量而非原地乱改。
5. def reduce(self, fn):对应 thrust::reduce，把整个序列归约成一个值。
6. probs = dv.transform(lambda x: math.exp(x))lambda 就是正文里的“配方卡”。
7. total = probs.reduce(lambda a, b: a + b)归约求和，对应 GPU 上的并行加法树。
8. norm = probs.transform(lambda p: p / total)归一化，全部留在“设备端”完成。
9. ranked = norm.sort(reverse=True)排序取 Top，对应 thrust::sort。
10. top3 = [round(p, 4) for p in ranked.to_host()[:3]]只有最后这一步才“拷回主机”。

常见误区

面试题