关卡阅读材料 // 02
任务并发
用多个 Stream 表达数据块之间的独立性，让复制、计算和回传获得重叠机会，同时保留每个数据块内部的正确顺序。
建议边读边写下答案；时长包含代码推演与练习。
进入互动关卡→- 01导读与目标05 MIN
- 02心智模型08 MIN
- 03概念深挖15 MIN
- 04代码推演12 MIN
- 05历史生态08 MIN
- 06练习复盘12 MIN
把 GPU 应用视为完整工作流，使用时间线找到 CPU 提交、数据搬运与 GPU 计算之间的空隙。
关卡用 Stream 时间线讲清依赖；实验把它转成真实 Nsight Systems Trace，要求区分“API 异步”和“设备真的重叠”。
- 课程 00–01
- 理解 H2D / D2H
- 知道同一 Stream 内保持顺序
- Nsight Systems
- CUDA Streams / Events
- Page-locked Host Memory
导读与目标
先记住这一句
并发来自独立工作与正确依赖；异步 API 只是把这种可能性告诉运行时。
完成后你应该能够
- 01
画出单 Stream 与多 Stream 的依赖时间线。
- 02
判断哪些操作可以重叠，哪些必须由事件或 Stream 顺序保护。
- 03
解释页锁定 Host Memory、Copy Engine 与默认 Stream 对重叠的影响。
- 04
把同步从“每一步都等”移动到真正的数据消费边界。
建立心智模型
- 01队列
同一 Stream 内的命令按提交顺序执行，天然形成依赖链。
- 02拆分
每个独立数据块使用自己的 Stream，可让不同块的 H2D、Kernel 与 D2H 交错。
- 03汇合
先提交所有独立工作，最后只在真正消费结果前同步。
概念深挖
Stream 是顺序语义
它描述一条有序命令流。不同 Stream 之间可以并发，但不承诺一定同时执行。
依赖应跟着数据走
按数据块拆分时，每块的复制、计算、回传保持同一 Stream，避免丢失局部依赖。
重叠需要条件
异步传输通常还要求合适的设备能力、页锁定 Host Memory、足够资源以及没有隐式同步。
Stream 先定义顺序，再提供并发机会
同一 Stream 中的命令按提交顺序观察执行关系，因此 H2D → Kernel → D2H 可以自然表达一个 Chunk 的流水。把三个 Chunk 放进同一 Stream 会保持正确性，却也让互相独立的 Chunk 排队。按 Chunk 拆成多个 Stream 后，每条局部链仍正确，链与链之间才获得并发机会。
不同 Stream 的工作可能并发，不代表必然并发。资源不足、隐式同步、默认 Stream 语义或数据依赖都会让它们串行。Stream 是程序对独立性的声明，硬件与运行时决定在当下能实现多少重叠。
停下来想一想为什么按 H2D、Kernel、D2H 三种操作各建一条 Stream 可能破坏正确性？＋
参考答案同一 Chunk 的 Kernel 可能在它的 H2D 完成前启动，D2H 也可能过早发生；需要额外事件重建按数据的依赖。
传输重叠依赖 Host 与 Device 两侧
cudaMemcpyAsync 的 Host 行为与真实 DMA 路径取决于内存类型和传输方向。页锁定 Host Memory 让设备能够稳定地直接 DMA，通常是 H2D/D2H 与计算重叠的重要条件。普通 Pageable Memory 可能触发 staging 或同步行为。
设备还需要可用的 Copy Engine，Kernel 与传输也不能争用到无法并行的资源。应通过时间线工具验证，而不是从 API 名称推断。小 Chunk 过多时，启动与管理开销也可能吞掉流水收益。
停下来想一想API 返回得很快，能否证明传输已经与 Kernel 重叠？＋
参考答案不能。快速返回只说明 Host 没有一直等待；需要设备时间线确认 DMA 与 Kernel 的实际重叠。
Event 表达局部依赖，避免全局停机
cudaDeviceSynchronize 会等待整个 Device 上此前工作完成，边界非常宽。若 Stream B 只依赖 Stream A 中某个操作，可以在 A 记录 Event，再让 B 等待该 Event。这样既保证正确性，又不会阻塞无关 Stream 或 Host。
当相同依赖图反复执行时，CUDA Graph 可以减少重复提交开销，但它不会创造新的独立性。先把依赖设计正确，再考虑用 Event 或 Graph 更精确、更低开销地表达。
停下来想一想Stream B 只依赖 Stream A 的一次 H2D，为什么 Event 通常比 cudaDeviceSynchronize 更合适？＋
参考答案Event 只建立需要的跨 Stream 边，其他无关工作仍可继续；设备级同步会等待所有此前工作。
把概念放回代码
01for (int s = 0; s < 3; ++s) {02 cudaMemcpyAsync(d[s], h[s], bytes, H2D, stream[s]);03 kernel<<<grid, block, 0, stream[s]>>>(d[s]);04 cudaMemcpyAsync(out[s], d[s], bytes, D2H, stream[s]);05}06cudaDeviceSynchronize(); // 消费结果之前代码逐段推演
- 01LOOP每次迭代拥有独立 Chunk
s 同时选择 Host Buffer、Device Buffer 与 Stream，避免不同 Chunk 误用同一份存储。
- 02H2D异步提交输入搬运
若 Host Buffer 可被立即复用，程序还需等待对应传输完成；异步不等于生命周期自动安全。
- 03KERNEL同 Stream 保留局部顺序
Kernel 会在本 Chunk 的 H2D 之后执行，不需要 Host 在中间同步。
- 04JOIN在消费所有结果前统一等待
把同步移到提交循环之后，为不同 Chunk 的复制、计算与回传保留最大重叠窗口。
知识图谱
补充参考 · 不要求一次读完，也不计入核心 60 分钟
核心术语
- Stream
- 一条有序的 Device 命令流；不同 Stream 表达潜在独立性，但不保证物理并发。
- Event
- 记录 Stream 中某个完成点，并为计时或跨 Stream 依赖提供局部同步。
- Pinned Memory
- 不会被操作系统换出的 Host Memory，使 DMA 路径更直接，也是异步传输的重要条件。
- Copy Engine
- 执行数据传输的硬件引擎；数量和方向能力影响复制能否与计算或另一方向复制重叠。
- CUDA Graph
- 预先描述并重复执行一组操作与依赖，以降低反复提交的 CPU 开销。
带数字推演
- 01串行基线
3 × (2 + 5 + 2) = 27 ms。每个 Chunk 都等前一个完整结束。
- 02找节拍
最慢阶段是 5 ms 的 Kernel，因此稳态每 5 ms 推进一个 Chunk。
- 03加首尾
首个 Chunk 走完需 9 ms，后两个各增加 5 ms：9 + 2 × 5 = 19 ms。
诊断手册
用了多个 Stream 仍是一条直线
- 先检查
- 默认 Stream 语义、Host Memory 类型、同步位置、设备并发能力
- 需要的证据
- Nsight Systems 中 Host API、Memcpy 和 Kernel 的同轴时间线
Stream 越多越慢
- 先检查
- Chunk 是否太小、启动开销、Context/资源竞争与内存带宽饱和
- 需要的证据
- Stream 数量扫描、CPU 提交时间和 GPU 利用率
偶发读到旧数据
- 先检查
- 跨 Stream 生产者—消费者边、Buffer 生命周期与 Host 复用时机
- 需要的证据
- Event 依赖图、每个 Buffer 的所有者和同步点
从硬件到生态
- 01异步 Kernel + Stream
Host 可以提交工作后继续执行。
程序能表达 CPU、复制和 GPU 计算的重叠。 - 02Kepler Hyper-Q
更多硬件工作连接减少独立队列被错误串行化。
软件已表达的并发更容易抵达硬件。 - 03Event + CUDA Graph
依赖从宽泛等待变成可复用的局部执行图。
优化重点从“更多队列”转向“更精确的依赖与更低提交开销”。
硬件与生态坐标
CUDA Streams 早于 Kepler。Hyper-Q 的历史意义在于提供更多硬件工作连接，减少多个独立 Stream 被错误串行化的情况，让软件已经表达出的并发更有机会抵达硬件。
练习与复盘
为什么不应该在每次 cudaMemcpyAsync 后立刻 cudaDeviceSynchronize？
- A. 它会让拷贝结果不正确
- B. 它会阻塞 Host，提前消除后续工作与当前工作的重叠窗口
- C. 它会把显存清空
查看答案＋
同步位置就是并发边界。过早等待会把原本独立的命令重新排成一条串行链。
三个 Chunk 共用默认 Stream，各执行 H2D→Kernel→D2H。最多会形成多少条有序链？
提示＋
同一 Stream 的命令按提交顺序。
参考答案＋
一条包含九个操作的长链。Chunk 之间即使独立，也没有从程序中暴露出来。
提交每个 Chunk 后立刻 cudaStreamSynchronize(stream[s])。正确但不快，为什么？
提示＋
Host 何时能继续提交下一个 Chunk？
参考答案＋
Host 会等当前 Chunk 完整结束才提交下一个，跨 Chunk 的流水窗口被关闭。应先提交所有 Chunk，再在消费对应结果前等待。
Stream B 的 Kernel 必须读取 Stream A 产生的数据。如何避免设备级同步？
提示＋
在生产者完成点记录局部信号。
参考答案＋
在 Stream A 记录 Event，让 Stream B 通过 cudaStreamWaitEvent 等待，再提交消费者 Kernel。
代码使用多个 Stream 和 cudaMemcpyAsync，时间线仍无重叠。列出三个检查方向。
提示＋
从 Host 内存、硬件能力和同步三层考虑。
参考答案＋
检查 Host Buffer 是否页锁定、设备 Copy Engine/并发能力、是否存在默认 Stream 或显式/隐式同步；还应检查任务是否太小或资源已饱和。
可选实践实验
不计入核心 60 分钟 · 需要相应 CUDA / GPU 环境
实验目标
把三块数据变成可验证的复制—计算流水线
从一个默认 Stream 的九段串行任务出发，构建每 Chunk 一条 Stream 的版本，并用时间线验证重叠与依赖。建议步骤
- 01
先用 Pageable Host Memory 和默认 Stream 运行三个 Chunk，保存基线时间线。
- 02
改用 Pinned Host Memory、三个非阻塞 Stream；每条 Stream 内保持 H2D→Kernel→D2H。
- 03
把同步移动到所有任务提交之后，再用 Event 表达一个刻意加入的跨 Stream 依赖。
- 04
在 Nsight Systems 中标注 Host API、Copy Engine 与 Kernel 区间，解释重叠或未重叠原因。
提交优化前后两张时间线、同步位置说明，以及端到端时间。必须说明硬件能力与任务大小限制。
把稳定的提交序列捕获为 CUDA Graph，并比较 CPU 提交开销，而不是只比较 GPU Kernel 时间。