<<<CUDA C++ 学习路线grid · block · warp · lane
EN中文
阶段 3 · 优化实战10 / 15约 24 分钟cuda/08_streams.cu

Stream 与异步：让拷贝和计算同时进行

流的语义、事件计时、三段流水线，以及 CUDA Graph

学完这一课你会
  • 理解 stream 的顺序语义和跨流并发
  • 搭出 H2D / compute / D2H 三路重叠的流水线
  • 用 cudaEvent 正确地给 GPU 计时（而不是用 CPU 时钟）
  • 知道默认流的坑，以及 CUDA Graph 解决的是什么问题

一个 stream 是一条有序的命令队列：同一个流里的操作严格按提交顺序执行；不同流之间则没有任何顺序约束，可以并发。现代 GPU 有独立的拷贝引擎（copy engine）和计算引擎，这意味着一次 H2D 拷贝、一次 kernel 执行、一次 D2H 拷贝可以在物理上同时进行——前提是它们分属不同的流。

串行执行：
  H2D0 ─ K0 ─ D2H0 ─ H2D1 ─ K1 ─ D2H1 ─ H2D2 ─ K2 ─ D2H2 ─ ...
  |────────────────────── 总时长 12 单位 ──────────────────────|

三流重叠：
  stream0:  H2D0  K0   D2H0
  stream1:        H2D1  K1   D2H1
  stream2:              H2D2  K2   D2H2
  stream3:                    H2D3  K3   D2H3
            |──────── 总时长 6 单位 ────────|

  拷贝引擎和计算引擎并行工作，理想情况下接近 2 倍加速
串行 vs 三段流水线（4 个数据块）

创建流并提交工作

08_streams.cu — 分块流水线
1const int nStreams = 4;2cudaStream_t streams[nStreams];3for (int i = 0; i < nStreams; ++i)4    CUDA_CHECK(cudaStreamCreate(&streams[i]));5 6// 异步拷贝要求 host 侧是 pinned 内存，否则会静默退化成同步7float *h_in, *h_out;8CUDA_CHECK(cudaHostAlloc(&h_in,  bytes, cudaHostAllocDefault));9CUDA_CHECK(cudaHostAlloc(&h_out, bytes, cudaHostAllocDefault));10 11int chunk = n / nStreams;12for (int i = 0; i < nStreams; ++i) {13    int off = i * chunk;14    size_t cb = chunk * sizeof(float);15 16    CUDA_CHECK(cudaMemcpyAsync(d_in + off, h_in + off, cb,17                               cudaMemcpyHostToDevice, streams[i]));18 19    myKernel<<<chunk / 256, 256, 0, streams[i]>>>(d_in + off, d_out + off, chunk);20 21    CUDA_CHECK(cudaMemcpyAsync(h_out + off, d_out + off, cb,22                               cudaMemcpyDeviceToHost, streams[i]));23}24 25// 等所有流跑完26for (int i = 0; i < nStreams; ++i)27    CUDA_CHECK(cudaStreamSynchronize(streams[i]));

用 event 给 GPU 计时

不要用 `std::chrono` 给 kernel 计时。CPU 时钟测的是「提交命令花了多久」，由于 launch 是异步的，你量到的往往是几微秒——毫无意义。正确做法是用 cudaEvent，它在 GPU 的时间线上打时间戳。

正确的 kernel 计时
1cudaEvent_t start, stop;2CUDA_CHECK(cudaEventCreate(&start));3CUDA_CHECK(cudaEventCreate(&stop));4 5// 先热身：第一次 launch 包含 JIT / 上下文初始化，数据没有参考价值6myKernel<<<blocks, threads>>>(d_in, d_out, n);7CUDA_CHECK(cudaDeviceSynchronize());8 9CUDA_CHECK(cudaEventRecord(start));10for (int i = 0; i < 100; ++i)                     // 多跑几次取平均11    myKernel<<<blocks, threads>>>(d_in, d_out, n);12CUDA_CHECK(cudaEventRecord(stop));13CUDA_CHECK(cudaEventSynchronize(stop));           // 等 stop 事件真正发生14 15float ms = 0.0f;16CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));17ms /= 100.0f;18 19// 算一下有效带宽，和硬件峰值对比才知道优化空间还有多少20double gb = 3.0 * n * sizeof(float) / 1e9;        // 读 a、读 b、写 c21printf("%.3f ms, 有效带宽 %.1f GB/s\n", ms, gb / (ms / 1000.0));

跨流依赖：用 event 而不是 synchronize

让 stream1 等 stream0 的某个节点
1cudaEvent_t done;2CUDA_CHECK(cudaEventCreateWithFlags(&done, cudaEventDisableTiming));  // 不计时更轻量3 4kernelA<<<g, b, 0, stream0>>>(d_x);5CUDA_CHECK(cudaEventRecord(done, stream0));6 7// stream1 在此处插入一个等待点，但 CPU 不会阻塞8CUDA_CHECK(cudaStreamWaitEvent(stream1, done, 0));9kernelB<<<g, b, 0, stream1>>>(d_x);      // 保证在 kernelA 之后才开始

CUDA Graph：消灭 launch 开销

每次 kernel launch 都有几微秒的 CPU 侧开销。当你的 kernel 本身只跑 10 微秒、而且要在循环里重复上千次时（LLM 推理的解码阶段就是典型场景），launch 开销会成为真正的瓶颈。CUDA Graph 把一整串操作录制成一个图，之后一次提交就能重放全部节点，把 N 次 launch 的开销压缩成 1 次。

录制并重放
1cudaGraph_t graph;2cudaGraphExec_t instance;3 4// 录制模式：这段里的所有操作只入图，不真正执行5CUDA_CHECK(cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal));6for (int i = 0; i < 20; ++i) {7    kernelA<<<g, b, 0, stream>>>(d_x);8    kernelB<<<g, b, 0, stream>>>(d_x);9}10CUDA_CHECK(cudaStreamEndCapture(stream, &graph));11 12CUDA_CHECK(cudaGraphInstantiate(&instance, graph, nullptr, nullptr, 0));13 14// 之后每一轮只需一次提交，40 个 kernel 的 launch 开销降为一次15for (int iter = 0; iter < 1000; ++iter) {16    CUDA_CHECK(cudaGraphLaunch(instance, stream));17}18CUDA_CHECK(cudaStreamSynchronize(stream));

自测

先自己在心里回答一遍，再展开对照。答不上来的说明这一段值得重读。