极客工场

Appearance

共 4,808 字预计阅读约 16 分钟

CUDA stream 异步：用多流流水线把 PCIe 拷贝藏到计算背后 ​

更新时间：2026-09-25。在第一个 CUDA 程序里，代码通常长这样：cudaMemcpy 把数据搬进去，启动 kernel，cudaDeviceSynchronize 等它算完，再 cudaMemcpy 把结果搬回来。能跑，但 GPU 的时间线上有大段空闲：搬数据时 SM 在发呆，计算时拷贝引擎在发呆。这一篇讲 CUDA 用什么机制把两件事叠起来——stream（流）。它和统一内存解决的是同一组痛点的两个侧面：统一内存让你少写拷贝代码，stream 让你把不可避免的拷贝时间藏掉。

本文要回答的问题 ​

  • 什么是 stream？为什么说即使代码里从没出现过 stream，你的 kernel 也一直在某个流上执行？
  • cudaMemcpyAsync 「异步」到底是什么意思，为什么拿普通 malloc 的内存做异步拷贝会被驱动偷偷同步？
  • 多流流水线长什么样？H2D 搬第 N+1 批数据的同时 GPU 算第 N 批，时间线如何重叠，加速比怎么算？
  • event（事件）除了给 kernel 计时，还能怎么表达跨流依赖？cudaStreamWaitEvent 用在什么场景？
  • 流开得越多越快吗？GPU 上的拷贝引擎有几个，流数量和硬件引擎数量是什么关系？

一、默认流：没创建流，也已经在流上 ​

先给定义：stream 是一条按顺序执行的操作队列。你往流里丢入的每一个操作——一次 H2D 拷贝、一次 kernel 启动、一次 D2H 拷贝、一次 memset——都进入队尾，GPU 按入队顺序逐个执行。同一个流内，前一个操作没完成，后一个不会开始，这是硬性顺序保证。

大多数入门代码从没显式创建过流，因为 CUDA 提供了一条默认流（default stream，编号 stream 0）。不指定流参数的 cudaMemcpy、<<<>>> 启动、cudaMemset 全部进入默认流。所以你写的每一行 GPU 代码其实都在流上，只是这条流由运行时替你管着。

默认流带来一个隐蔽的串行陷阱。看这个时间线，H2D 拷贝由 PCIe 上的拷贝引擎（copy engine，本质是一个独立 DMA 控制器）完成，kernel 由 SM 完成，D2H 再由拷贝引擎完成——三者在硬件上本来是不同的执行单元，但因为排在同一条 FIFO 队列里，谁也不许抢先：

chip-design/intermediate/cuda-streams-async-copy.md D2 示意图

两批数据各 50 ms（H2D 10 + 计算 30 + D2H 10），串行跑完正好 100 ms。注意图里大片的「空闲」：H2D 引擎一个批次只忙 10 ms，利用率 20%；SM 只忙 30 ms，利用率 60%。硬件单元之间互不相干，软件队列却把它们绑成了一列——这正是 stream 要解开的结。

异步还带来一个必须养成的习惯：错误检查要「延后」。kernel 启动是异步的，启动那一行即使越界、参数非法，返回值也往往是 cudaSuccess——故障要等后面某个同步点（cudaStreamSynchronize、下一次拷贝甚至另一个 kernel 启动）才以错误码的形式浮出来，报错位置和真正出错的位置可能隔着十几行。调试阶段可以在每个 kernel 之后插一句 cudaGetLastError() 配合流同步，把错误「钉」在案发地；发布前再撤掉。反过来，流内顺序保证也是程序员手里最重要的工具：同一条流里 H2D 完成后 kernel 才会启动，kernel 完成后 D2H 才会开始，这层保证不靠任何显式同步，入队顺序即依赖关系。

补一个历史细节：老版编译模式（CUDA 7 之前的默认行为）下，默认流还会和设备上所有其他流做隐式同步，像一道全局栅栏，谁都越不过它。现在用 --default-stream per-thread 编译后默认流是每线程独立的普通流，不再阻塞别的流，但同流内串行这一条永远不变。流还有优先级属性，cudaStreamCreateWithPriority 可以创建高低优先级流，硬件就绪队列里高优先级流的 block 会被优先调度，但它只影响排队等待时的顺序，不能抢占已经在 SM 上执行的 kernel。

二、cudaMemcpyAsync 与 pinned memory：异步的入场券 ​

显式创建流只有三行：

cpp
cudaStream_t stream;
cudaStreamCreate(&stream);
cudaMemcpyAsync(d_a, h_a, bytes, cudaMemcpyHostToDevice, stream);
kernel<<<grid, block, 0, stream>>>(d_a);          // 第 4 个参数是 shared memory 字节数
cudaMemcpyAsync(h_a, d_a, bytes, cudaMemcpyDeviceToHost, stream);
cudaStreamDestroy(stream);

Async 的含义是：调用立刻返回，操作只是被排进了队列，CPU 线程不等它完成。这带来两个直接后果：其一，CPU 在 GPU 干活期间可以去准备下一批数据、做日志、跑别的线程，而不是干等；其二，也是关键——多个流里排队的操作，GPU 硬件在资源允许时可以让它们并发执行。

但异步有一个常被忽略的前提：主机侧的内存必须是锁页内存（pinned / page-locked memory），用 cudaMallocHost 分配：

cpp
float *h_a;
cudaMallocHost(&h_a, bytes);   // 分配后物理页被锁定，操作系统不会换出到磁盘
// ... 使用 h_a ...
cudaFreeHost(h_a);

为什么普通 malloc 的内存不行？操作系统管理的可分页内存（pageable memory）随时可能被换页、移动，物理地址不固定。而 GPU 的 DMA 引擎一旦启动传输，它拿到的是一个物理地址，传输期间没人能替它翻译或修正地址——DMA 直接在 GPU 和物理内存之间搬数据，不经过 CPU。于是驱动面对可分页内存时只能走一条弯路：先在内部一块锁页暂存区（staging buffer）上做一次 CPU 参与的拷贝，再从暂存区发起 DMA。更糟的是，对异步调用驱动无法保证暂存拷贝何时完成，很多情况下它会在排队前直接做一次隐式同步——你以为写的是 cudaMemcpyAsync，实际效果退化成了 cudaMemcpy，异步收益全部消失，还多搬了一次数据。

锁页内存消除了这个中间环节：物理页常驻、地址固定，DMA 可以直接全速读写用户缓冲区，操作才能真正排队后立刻返回。代价也要清楚：锁页内存不能被换出，占多了整机可用内存减少、系统响应变慢；它的分配和释放也比 malloc 慢。实践中只把需要高带宽传输的缓冲区锁页，容量控制在物理内存的几分之一，绝不把几个 GB 的堆全钉住。

cudaMallocHost 还有一个同族接口 cudaHostAlloc，接受几个有用的标志。cudaHostAllocMapped 分配的锁页内存同时被映射进 GPU 地址空间，可以用 cudaHostGetDevicePointer 拿到设备侧指针，kernel 直接跨过 PCIe 读写主机内存，称为 zero-copy 访问。它省掉了显式拷贝，但每次访问都付 PCIe 延迟，只适合传输量极小、访问稀疏或只想让 GPU 顺手写几个状态字的场景；大块数据走它会慢到可以当反面教材。cudaHostAllocPortable 则让这块锁页内存在多线程、多设备上下文里都被当作锁页对待，多 GPU 程序里偶尔需要。日常记住主结论即可：想要异步拷贝跑出满速，源和目的主机缓冲区必须是锁页的，这条没有替代品。

一个验证小技巧：cudaMemcpyAsync 返回后立刻在主机侧改写源缓冲区是合法的（操作已入队），但要确保真正的传输完成前数据不能被踩——这正是需要 event 或分批次独立缓冲区的原因。工程上通常为每个批次准备各自独立的 pinned 缓冲区，CPU 在第 N+2 批的缓冲里填数据，GPU 同时在第 N+1 批上做 DMA、在第 N 批上计算，三块数据互不干扰，谁也不用等谁。

三、多流流水线：搬第 N+1 批，算第 N 批 ​

现在把数据切成多个批次，每一批用独立的流串起「H2D → 计算 → D2H」，让三条硬件单元像流水线一样转起来。以三批为例，每批仍然是 H2D 10 ms、计算 30 ms、D2H 10 ms：

chip-design/intermediate/cuda-streams-async-copy.md D2 示意图

读这张时间线：第 10 ms 时 H2D #1 刚完、H2D #2 开始，同一时刻 SM 开始算第 1 批；第 40 ms 计算 #1 结束、结果经 D2H 回传，而 H2D 早已把三批数据全部送完，SM 无缝衔接到计算 #2。拷贝和计算在时间轴上重叠了。

实现上并不需要给每批新建和销毁流，常见写法是预先建好少量流、为每个批次准备独立的缓冲，在一个循环里把全部操作一次排完：

cpp
const int NB = 8;
for (int b = 0; b < NB; ++b) {
  cudaMemcpyAsync(d_in[b],  h_in[b],  sz, cudaMemcpyHostToDevice,  stream[b % NSTREAMS]);
  kernel<<<grid, block, 0, stream[b % NSTREAMS]>>>(d_in[b], d_out[b]);
  cudaMemcpyAsync(h_out[b], d_out[b], sz, cudaMemcpyDeviceToHost,  stream[b % NSTREAMS]);
}
cudaDeviceSynchronize();   // 生产代码里更精细的做法是用 event 逐批回收缓冲

这段循环在 CPU 上几十微秒就跑完了——它排的是未来几十毫秒的全部工作计划，之后 CPU 与三块硬件单元各自忙各自的。深度学习训练里的 prefetch、推理服务里的连续 batch，底下都是同一个结构：CPU 线程不断往队列里填批，GPU 稳态地以计算时间为节拍吐结果。需要留意批次到流的映射：相邻批次如果共用同一条流，它们之间会退回串行，所以批次数应是流数的整数倍并按取模均匀铺开。

算一笔账。三批串行需要 3 × 50 = 150 ms；流水线总耗时约 110 ms——启动阶段有 10 ms 无法重叠的填充（prefill），收尾阶段计算 #3 加最后一次 D2H 共 40 ms 的排空（drain），中间的稳态期每 30 ms 吐出一批。批次数越多，首尾开销摊得越薄：

$$T_{pipeline} = (N-1)\times T_{compute} + T_{H2D} + T_{compute} + T_{D2H}$$

代入 N 批、每段 10/30/10 ms：$T = 30N + 20$ ms，而串行是 $50N$ ms。10 批时为 320 ms 对 500 ms，吞吐加速比约 1.56 倍；批次趋于无穷时稳态上限是 50/30 ≈ 1.67 倍。数字本身不性感，但注意它的含义：每批 20 ms 的往返 PCIe 时间被完全藏掉了，没换硬件、没改 kernel，只改了任务的排队方式。如果计算更短、拷贝占比更高（比如小 batch 推理、数据预处理在 CPU 侧），重叠收益反而更大；反过来若 kernel 一次算几百毫秒，拷贝那十几 ms 藏不藏都无所谓——先看时间线再决定值不值得做。

三种写法摆在一起：

对比维度同步模式（cudaMemcpy + DeviceSynchronize）单流异步多流流水线
入队后 CPU 线程阻塞，等每一步完成不阻塞，继续往下跑不阻塞，连续排完所有批次
同流内拷贝与计算严格串行严格串行各流内部仍严格串行
跨操作并发无仅与 CPU 侧工作重叠H2D / 计算 / D2H 三类硬件单元并发
PCIe 时间能否藏掉不能不能能，藏进 kernel 执行时间
主机内存要求pageable 也能跑（驱动代搬）pageable 会退化成同步必须 pinned 才能真正重叠
同步方式每次调用即同步cudaStreamSynchronize 等流event 或逐流等待
典型适用场景调试、一次性小传输CPU 与 GPU 同时干活批量推理、训练数据管线、流式处理

四、event：计时与跨流依赖 ​

异步之后，clock() 这类 CPU 侧计时全部失效——调用返回时活还没干。CUDA 提供 event（事件），它是可以插进流里的时间点标记：

cpp
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, stream);
kernel<<<grid, block, 0, stream>>>(d_a);
cudaEventRecord(stop, stream);
cudaEventSynchronize(stop);
float ms = 0;
cudaEventElapsedTime(&ms, start, stop);   // GPU 自己打的时间戳，精度到微秒级

两个 event 都排在同一条流里，计时的是它们之间 GPU 实际执行的区间，不含 CPU 入队和排队等待，这是测量 kernel 真实耗时的标准做法。

event 的第二个用途更重要：表达跨流依赖。不同流默认互不等待，但有时数据关系要求 B 流必须等 A 流走到某个点。比如 A 流负责把一张查找表 H2D 搬入并做初始化，B、C 两个计算流都得在表就绪后才能启动 kernel。你不能在 CPU 上 cudaStreamSynchronize(A) 再排队——那会阻塞 CPU，也浪费入队时机。正确写法：

cpp
cudaEvent_t tableReady;
cudaEventCreate(&tableReady);
cudaEventRecord(tableReady, streamA);          // A 流走到此处时事件触发
cudaStreamWaitEvent(streamB, tableReady, 0);   // B 流排队一个等待：事件触发后才继续
cudaStreamWaitEvent(streamC, tableReady, 0);   // 同一个 event 可以被多条流等
// 立刻把 B、C 的后续操作排进去，CPU 不阻塞

cudaStreamWaitEvent 排入的是一个流内的等待节点：B 流执行到该节点处暂停，直到事件在 A 流中被触发，再继续执行后面排队的操作。这就用软件画出了一条跨流的边，把多条独立队列在指定点缝合起来，又不破坏各自其余部分的并发。它在多 GPU 场景还有一个变体：cudaEventRecord 配合 cudaEventDisableTiming 标志并启用 IPC，可以让一张卡的流等待另一张卡上的事件，配合 NVLink/PCIe P2P 使用。

值得花一句话区分三种「等待」的粒度：cudaDeviceSynchronize 等的是整张卡上所有流的全部操作，范围最大，常在程序收尾或确定要回收所有缓冲时使用；cudaStreamSynchronize(stream) 只等一条流排到队尾，其他流不受影响，批量循环里按流回收结果用它；cudaEventSynchronize(event) 只等一个时间点，粒度最细，计时场景够用就不必升级成等整条流。等待范围越大，被白白拖住的并发越多，能细就别粗。

写 event 代码有两个常见坑：一是忘记 cudaEventSynchronize 或等错了流，读取 elapsed time 时数据还是旧值（函数会返回错误码但不报错退出）；二是把纯同步用的 event 也走完整时间戳硬件路径——不需要计时的 event 加上 cudaEventDisableTiming 创建，开销更低，同步场景一律建议带上。还有一个概念上的辨析：event 触发只保证在同一条流里排在它前面的操作全部完成，它不是一块「数据已就绪可读」的证书；CPU 是否能安全读回某个缓冲区，仍取决于对应的 D2H 拷贝排在 event 之前且属于同一流，别把流内顺序保证张冠李戴到缓冲区可见性上。

五、硬件边界：拷贝引擎只有 1~2 个，流不是越多越好 ​

流是软件队列，并发度最终由硬件单元数决定，这里必须回到GPU 的硬件结构上看：

  • 拷贝引擎数量有限。消费级卡通常只有 1 个拷贝引擎，意味着 H2D 和 D2H 不能同时进行，两个方向的传输彼此串行；数据中心卡（如 A100/H100 一档）一般有 2 个拷贝引擎，入方向和出方向各一个，可以一边搬入一边搬回。设备属性里的 asyncEngineCount 直接给出这个数字。
  • kernel 之间的并发有条件。不同流的 kernel 在 SM 资源（block 槽位、shared memory、寄存器）有空余时才能同时驻留；一个 kernel 若占满全卡，其他流的 kernel 只能排队。MPS（Multi-Process Service）或 MIG 会改变这些规则，但「流数不等于并发 kernel 数」始终成立。
  • 传输与计算的并发最稳。拷贝引擎和 SM 是物理上独立的单元，一条流做 DMA、另一条流占着 SM，这种重叠在任何支持 async engine 的卡上都成立，也是多流优化收益最确定的部分。

因此流的数量按「能制造多少条独立工作的硬件流水线」来定，通常 2~4 条流即可覆盖「一条 H2D + 一条计算 + 一条 D2H」的全部重叠机会；超过引擎数和 kernel 并发能力的流只会增加调度开销和内存占用（每条流的批次缓冲往往是独立的），不会再提速。少数场景例外：流上挂的操作延迟差异大、需要细粒度交错时，可以适当多开，但要用 profiler（Nsight Systems 的时间线视图）验证重叠是否真实发生——凭感觉开几十条流是典型的负优化。

最后给一份落地顺序：先用 profiler 或 event 量出 H2D、kernel、D2H 各占多少时间，确认拷贝占比值得藏；把传输缓冲区换成 cudaMallocHost 分配的 pinned memory；数据切批，每批一条流或在少量流上循环排多批；用 cudaStreamWaitEvent 缝合必须保证的跨流先后关系；再跑一次 profiler，看到拷贝块和计算块在时间线上真正咬合，才算完成。每批的字节数也要留意：太小的拷贝被固定启动开销主导，带宽跑不满，批大小以单批传输能跑到峰值带宽六七成以上为宜。

这一篇在知识体系里的位置 ​

这一篇属于 GPU 编程的「执行模型与数据移动」专题。向基础看，kernel 启动和线程组织在第一个 CUDA 程序里；向硬件看，SM、warp 调度与 HBM 带宽决定了操作为什么能、以及能并发到什么程度，见 GPU 架构深入：SIMT、warp 调度；向另一条数据管理路线看，不想手写 cudaMemcpyAsync 时可以读 CUDA 统一内存，理解页面迁移与显式流水线各自的取舍。stream 是上层推理框架数据 loader、训练框架 prefetch pipeline 底下共同的原语，看懂时间线之后再看那些框架的配置项，会发现它们只是在替你排流。

总结 ​

CUDA stream 是 GPU 操作的 FIFO 队列，默认流让所有操作串行；配上 cudaMallocHost 的锁页内存和 cudaMemcpyAsync，再把数据分批投到多条流上，就能让 H2D 拷贝、kernel 计算、D2H 拷贝在不同硬件单元上并发，把 PCIe 传输时间整段藏进计算背后——典型得到约 1.5~1.7 倍吞吐提升，而 cudaEventRecord 与 cudaStreamWaitEvent 负责计时和在必要处缝合跨流依赖，流的数量则受限于 GPU 仅 1~2 个拷贝引擎和 SM 的实际空余资源。