Streams, events, and overlap
Draw a copy-and-kernel dependency graph and implement only its required cross-stream event waits
Before you begin
Follow a linked prerequisite when you cannot explain it from memory. Background skills without links are stated as abilities rather than hidden terminology.
- Chapter 2 explanation of asynchronous GPU timing
- Ability to draw which kernel or copy produces each later input
- Basic host-to-device and device-to-host copy calls
A CUDA program submits work from the host. The submitted operations include memory copies, kernel launches, event records, and waits. A stream gives those operations an order: later work in the same stream does not pass earlier work. The stream does not reserve a physical compute lane or copy engine.
Two streams make overlap possible only when the operations are independent and the device has compatible resources available. A copy and a kernel may overlap when the transfer uses an asynchronous path and the device can execute a copy engine beside compute. Two kernels may overlap when their combined register, shared-memory, pipeline, and bandwidth demands leave room. Separate streams do not create additional hardware capacity.
The main design task is dependency expression. If kernel B reads data produced by kernel A, B must wait for A. An event can represent that specific edge while unrelated work proceeds. A device-wide synchronization also enforces the edge, but it waits for every preceding operation and can introduce dependencies that the algorithm does not require.
Represent the dependency graph
Write each host-to-device copy, kernel, collective, and device-to-host copy as a node. Draw an edge when one operation consumes another operation's output, reuses its storage, or must occur later for another correctness reason. The resulting graph states the algorithm independently of CUDA streams.
Assigning nodes to a stream adds an ordering relation between those nodes. Recording an event after a producer and making another stream wait on that event adds one cross-stream edge. Event-based ordering is more precise than synchronizing the entire device because other streams retain freedom to run operations that do not depend on the producer.
The actual graph can contain edges that are not obvious in the source. Default-stream semantics, pageable-memory staging, synchronous allocation, a scalar read returned to the host, or a library call can block submission or execution. Use a system timeline to compare the intended graph with the observed order.
Determine whether operations can overlap
Independence is necessary but not sufficient. A host-to-device copy normally needs pinned host memory for a fully asynchronous DMA path. The accelerator needs a suitable copy engine and an execution engine that can operate concurrently. Two kernels need resources that can coexist on the device. If either kernel consumes nearly all SM resources, the second may wait even without a data dependency.
Overlap can also reduce performance. Two memory-bound kernels may execute at the same time but divide the same device-memory bandwidth. Their individual durations increase, and the combined completion time may not improve. Two compute kernels that use different pipelines can be more compatible. Judge concurrency by reduction in the application's critical path, not by the presence of overlapping rectangles in a trace.
Build and calculate a staged pipeline
A large input can be divided into chunks that pass through copy-in, compute, and copy-out stages. After startup, copy-in for chunk n+1 can run while compute processes chunk n and copy-out returns chunk n-1. If the stage times are 3 ms, 5 ms, and 1 ms, ideal steady-state throughput approaches one completed chunk every 5 ms because compute is the slowest stage.
The first chunk still requires its stages in dependency order, so pipelining does not reduce its latency to 5 ms. Startup and drain add exposed work at the boundaries. Too few chunks leave stages idle. Too many chunks increase launch overhead, event bookkeeping, and possibly working-set pressure. Calculate both steady-state throughput and per-item latency.
The timeline should show event edges, stage overlap, and remaining gaps. If a copy unexpectedly blocks, verify that host memory is pinned and that the API call is asynchronous for the selected memory type. If kernels do not overlap, inspect their resource use before adding more streams.
Worked example: Building a copy–compute pipeline
A dataset must move from pinned host memory to the GPU, pass through a kernel, and return compact results. The serialized path takes 3 ms copy in, 5 ms compute, and 1 ms copy out per batch.
- The serialized lower bound is 9 ms. With distinct engines and a steady pipeline, ideal throughput approaches the slowest stage, 5 ms per batch, after startup.
- For N batches, the idealized serialized time is 9N ms. The idealized pipeline time is 9 + (N - 1) × 5 ms because the first batch needs all three stages and each later batch completes one compute-stage interval later. For ten batches, this is 90 ms versus 54 ms, or about 1.67× faster, not 9 / 5 = 1.8×, because startup and drain remain visible.
- Split several batches or chunks across streams. Record an event after each input copy, make its kernel wait on that event, and make the output copy wait on the kernel.
- Capture a timeline. Verify that copy and compute overlap and look for hidden gaps from pageable memory, allocation, default-stream behavior, or host submission.
- Measure end-to-end throughput and single-batch latency. Pipelining improves steady throughput but does not necessarily reduce the latency of the first batch.
Conclusion: After startup, the 3 ms input copy and 1 ms output copy can execute during the 5 ms compute stage when the device provides separate resources. The pipeline can approach one completed batch every 5 ms, but the first batch still passes through all required stages.
Common errors
- Assuming one stream maps to one hardware engine. Streams describe order; the scheduler maps operations to available resources.
- Using pageable host memory and expecting truly asynchronous transfer. The runtime may stage or block in ways that erase overlap.
- Adding global synchronization to make timing easy. The measurement can destroy the pipeline it intends to describe.
Reference summary
A CUDA stream orders submitted operations. A stream does not reserve a hardware lane. Operations in different streams may overlap only when no dependency prevents overlap and the device has compatible copy, compute, memory, and residency resources available.
An event represents one point in a stream. Another stream can wait on that event and preserve unrelated concurrency. Device-wide synchronization introduces a much broader dependency and should be used only when the algorithm requires it.
cudaEventRecord(ready, producer_stream);
cudaStreamWaitEvent(consumer_stream, ready);
consumer<<<grid, block, 0, consumer_stream>>>(...);Knowledge check
Why might two independent kernels in separate streams still run sequentially?
Show the answer guide
- Separate streams provide possible concurrency, not dedicated hardware lanes; the device may lack enough compatible compute, copy, register, shared-memory, or memory-bandwidth resources for both kernels.
- A hidden dependency can also serialize the work through events, the legacy default stream, memory reuse, or an earlier producer that both kernels require.
- An Nsight Systems timeline plus resource and dependency checks must distinguish dependency serialization from resource non-overlap before changing the stream structure.
Test stream overlap
Place two independent kernels in separate non-default streams, then vary their block resource use or input size. Compare same-stream, separate-stream, and explicit-event variants while preserving correctness.
Evidence to keep: CUDA source, event dependency graph, Nsight Systems timelines, kernel durations alone and together, complete interval timing, and a classification of dependency versus resource serialization.
After the execution dependencies are correct, optimize data movement at each memory level.
Sources and further reading
The lesson explanations synthesize primary papers and official references. Follow the originals for proofs, exact APIs, architecture details, and version-sensitive behavior.