<<<CUDA C++ Coursegrid · block · warp · lane
EN中文
Stage 3 · Optimization in practice10 / 15about 24 mincuda/08_streams.cu

Streams and async: overlap copies with compute

Stream semantics, event timing, a three-stage pipeline, and CUDA Graph

After this lesson you can
  • Understand in-stream ordering and cross-stream concurrency
  • Build an H2D / compute / D2H overlapping pipeline
  • Time the GPU with `cudaEvent` (not the CPU clock)
  • Know the default-stream trap, and what problem CUDA Graph solves

A stream is an ordered command queue: operations in the same stream execute strictly in submission order; different streams have no ordering constraint and can run concurrently. Modern GPUs have a separate copy engine and compute engine, so one H2D copy, one kernel, and one D2H copy can physically run at the same time, if they belong to different streams.

Serial:
  H2D0 ─ K0 ─ D2H0 ─ H2D1 ─ K1 ─ D2H1 ─ H2D2 ─ K2 ─ D2H2 ─ ...
  |────────────────────── total 12 units ──────────────────────|

Three-stream overlap:
  stream0:  H2D0  K0   D2H0
  stream1:        H2D1  K1   D2H1
  stream2:              H2D2  K2   D2H2
  stream3:                    H2D3  K3   D2H3
            |──────── total 6 units ────────|

  Copy engine and compute engine work in parallel; ideally close to 2x
Serial vs three-stage pipeline (4 chunks)

Create streams and submit work

08_streams.cu: chunked pipeline
1const int nStreams = 4;2cudaStream_t streams[nStreams];3for (int i = 0; i < nStreams; ++i)4    CUDA_CHECK(cudaStreamCreate(&streams[i]));5 6// async copies need pinned host memory, otherwise they silently fall back to synchronous7float *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// wait for every stream to finish26for (int i = 0; i < nStreams; ++i)27    CUDA_CHECK(cudaStreamSynchronize(streams[i]));

Time the GPU with events

Do not time kernels with `std::chrono`. The CPU clock measures how long it took to submit the command. Because launch is asynchronous, you often get a few microseconds, which is meaningless. The correct tool is cudaEvent, which stamps the GPU timeline.

Correct kernel timing
1cudaEvent_t start, stop;2CUDA_CHECK(cudaEventCreate(&start));3CUDA_CHECK(cudaEventCreate(&stop));4 5// warmup first: the first launch includes JIT / context init, so the number is not usable6myKernel<<<blocks, threads>>>(d_in, d_out, n);7CUDA_CHECK(cudaDeviceSynchronize());8 9CUDA_CHECK(cudaEventRecord(start));10for (int i = 0; i < 100; ++i)                     // run several times and average11    myKernel<<<blocks, threads>>>(d_in, d_out, n);12CUDA_CHECK(cudaEventRecord(stop));13CUDA_CHECK(cudaEventSynchronize(stop));           // wait until the stop event actually happens14 15float ms = 0.0f;16CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));17ms /= 100.0f;18 19// compute effective bandwidth; compare it to peak to see how much room is left20double gb = 3.0 * n * sizeof(float) / 1e9;        // read a, read b, write c21printf("%.3f ms, effective bandwidth %.1f GB/s\n", ms, gb / (ms / 1000.0));

Cross-stream dependencies: use events, not synchronize

Make stream1 wait on a point in stream0
1cudaEvent_t done;2CUDA_CHECK(cudaEventCreateWithFlags(&done, cudaEventDisableTiming));  // cheaper when you do not need timing3 4kernelA<<<g, b, 0, stream0>>>(d_x);5CUDA_CHECK(cudaEventRecord(done, stream0));6 7// stream1 waits here; the CPU does not block8CUDA_CHECK(cudaStreamWaitEvent(stream1, done, 0));9kernelB<<<g, b, 0, stream1>>>(d_x);      // guaranteed to start only after kernelA

CUDA Graph: kill launch overhead

Every kernel launch costs a few microseconds on the CPU. When the kernel itself runs for 10 microseconds and you repeat it thousands of times in a loop (the decode phase of LLM inference is the typical case), launch overhead becomes the real bottleneck. CUDA Graph records a whole chain of operations as a graph, then one submit replays every node, collapsing N launches into 1.

Capture and replay
1cudaGraph_t graph;2cudaGraphExec_t instance;3 4// capture mode: operations in this region enter the graph and do not actually run5CUDA_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// each later iteration is a single submit; 40 kernel launches collapse to one15for (int iter = 0; iter < 1000; ++iter) {16    CUDA_CHECK(cudaGraphLaunch(instance, stream));17}18CUDA_CHECK(cudaStreamSynchronize(stream));

Self-check

Answer in your head first, then open the solution. If you cannot, reread that section.