先记住一句话

同一 stream 内严格按序；不同 streams 只是“允许并发”，真正 overlap 还需要无依赖、足够硬件资源、正确 memory 与没有隐式同步。

画 dependency DAG→分配 streams→用 events 连边→pinned memory + async copies→double buffer/graph capture→timeline 验证 overlap
Stream 0H2D₀ → K₀ → D2H₀
Stream 1H2D₁ → K₁ → D2H₁
events enforce only true edges
CUDA Graphreplay fixed DAG
Streams 表达并发队列，events 表达依赖，Graphs 降低重复提交固定 DAG 的 CPU overhead。

1. Stream 是 device work queue

kernel launch、async memcpy、event record 按 enqueue 顺序进入 stream。同一 stream 后项等前项；host 通常不等 device。多个 streams 可让 runtime 从就绪 work 中调度，但 priority 只是 hint。

2. Event 是 dependency marker

producer kernel in stream A
cudaEventRecord(ready, A)
cudaStreamWaitEvent(B, ready)
consumer kernel in stream B

这样只建立必要 dependency，不用粗暴 cudaDeviceSynchronize()。events 也可在同 stream 前后记录并给 device elapsed time。

3. Copy/compute overlap 条件

  • 硬件有 copy engines 与 concurrent capability；
  • 使用 async copy 和非阻塞 streams；
  • host buffer 是 pinned/page-locked；
  • chunks 足够大，copy 与 compute 时间可重叠；
  • 没有默认 stream/allocator/host API 引入全局同步；
  • device buffer 生命周期覆盖所有异步 work。

4. Double buffering

chunk 0: H2D ─ compute ─ D2H
chunk 1:      H2D ─ compute ─ D2H
chunk 2:           H2D ─ compute ─ D2H

用两/多组 buffers 和 streams 做 pipeline。chunk 太小则 launch/copy setup 主导，太大又不能充分 overlap 且增加 latency/memory。

5. Default stream 陷阱

legacy NULL stream 与 blocking streams 有特殊同步语义，插入一次无意的 default-stream operation 可能串行化本可并发的工作。framework 通常使用 per-device/current streams；custom op 必须在调用者 current stream 上 launch，不能私自用默认 stream。

6. 隐式同步来源

  • pageable host memory copy；
  • 某些 allocation/free 或 device-wide API；
  • 读取 device scalar 回 CPU、.item()；
  • debug/logging 触发同步；
  • 错误的 default stream；
  • profiling mode 或 deterministic setting。

Nsight Systems timeline 最容易看到 CPU gaps 与 streams 是否真正 overlap。

7. CUDA Graphs

Graph 把 kernels、copies 等 nodes 和 dependencies 与执行分离：先 capture/construct、instantiate，然后反复 launch executable graph。收益来自减少每个 tiny kernel 的 host driver setup，并让 runtime 看到完整 DAG。

8. Graph 适合什么

适合困难
重复、shape/memory 稳定的训练 step/inference动态 allocation、data-dependent control/shape
大量短 kernels、enqueue-boundhost callbacks/不支持 capture 的 API
固定 buffer addresses 可复用每次地址/拓扑大幅变化

graph 不能加速单 kernel 内部 arithmetic；若已经 compute-bound，收益小。

9. 测量与安全

分别记录 enqueue time、GPU compute time、end-to-end wall time。graph capture 前 warmup/JIT，capture 期间避免非法 API；每个 captured graph/context 的并发和 buffer ownership 要明确。

10. 四个异步执行手算

例 1：串行 pipeline

H2D=3 ms、kernel=7 ms、D2H=2 ms，单 batch latency=12 ms。处理 10 batches 串行需 120 ms。

例 2：理想 double buffering

稳态中 copy 与 compute 可 overlap，stage bottleneck=max(3,7,2)=7 ms。10 batches 理想总时约首尾开销加 9×7，接近 75 ms，而不是 120 ms。

例 3：event elapsed time

GPU start/end events 记录为 1.250 ms 与 3.875 ms，elapsed=2.625 ms。Host wall clock 若未 sync 可能只测到 enqueue 的几十 μs。

例 4：Graph launch overhead

一个迭代含 50 tiny kernels，每次 host launch 5 μs，提交开销约 250 μs。Graph replay 若整体 launch 15 μs，单迭代省约 235 μs；kernel compute 本身不因此加速。

常见误解：

创建多个 streams 不会自动让 kernels 并发。如果单 kernel 占满 SM、存在 event dependency、copy engine 忙或 default stream 同步，timeline 仍完全串行。

自测

1. event 比 device synchronize 好在哪里？

它只表达必要的 producer-consumer dependency，不阻塞无关 streams/host work。

2. pageable host memory 为什么妨碍 async copy？

DMA 需要稳定物理页，runtime 常要先 staging 到 pinned buffer，导致额外 copy/同步。

3. CUDA Graph 最主要减少什么？

稳定 operation DAG 每次逐 kernel 提交的 CPU/driver launch overhead。

官方资料

CUDA Programming Guide 的 Asynchronous Execution定义 streams/events/default-stream；CUDA Graphs覆盖 capture、instantiate、更新与限制。