GPU Stream and Event Synchronization is the CUDA programming model for managing concurrent operations on the GPU by organizing kernels and memory transfers into streams (ordered sequences of operations) and using events to synchronize between them — effective stream usage enables overlapping computation with data transfer, concurrent kernel execution, and precise timing measurements.
CUDA Stream Fundamentals:
- Stream Definition: a stream is a sequence of GPU operations that execute in order — operations in different streams may execute concurrently if the GPU has available resources
- Default Stream: operations without an explicit stream use stream 0 (the default/legacy stream) — the default stream implicitly synchronizes with all other streams unless compiled with --default-stream per-thread
- Stream Creation: cudaStreamCreate(&stream) creates a non-blocking stream — cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking) creates a stream that doesn't synchronize with the default stream
- Stream Destruction: cudaStreamDestroy(stream) releases stream resources — any pending operations in the stream complete before destruction
Overlapping Computation and Transfer:
- Concurrent Copy and Execute: launch memory copies (cudaMemcpyAsync) on one stream and kernels on another — the GPU's copy engines (DMA) and compute engines (SMs) operate independently
- Double Buffering Pattern: alternate between two buffers — while the GPU computes on buffer A, transfer results from the previous iteration and load input for the next iteration using buffer B
- Triple Buffering: three buffers provide maximum overlap — one being computed, one being uploaded, one being downloaded — hides both upload and download latency simultaneously
- Pinned Memory Requirement: cudaMemcpyAsync requires pinned (page-locked) host memory allocated with cudaMallocHost — unpinned memory forces synchronous copies regardless of stream assignment
Multi-Stream Concurrency:
- Concurrent Kernels: multiple small kernels on different streams can execute simultaneously if the GPU has enough SMs — Ampere A100 supports up to 128 concurrent kernels
- Resource Partitioning: concurrent kernels share SM resources — each kernel gets a portion of SMs proportional to its grid size, enabling fine-grained GPU sharing
- Stream Priority: cudaStreamCreateWithPriority(&stream, flags, priority) assigns a scheduling priority — high-priority streams preempt low-priority ones at thread block boundaries
- Breadth-First Launch: launch all streams' operations in interleaved order (kernel1_stream1, kernel1_stream2, ...) rather than depth-first (all of stream1, then all of stream2) — ensures the GPU sees concurrent work from all streams
CUDA Events:
- Event Creation: cudaEventCreate(&event) creates an event marker — events are recorded into streams and can be queried or waited on from the host or other streams
- Recording: cudaEventRecord(event, stream) inserts the event into the stream's operation queue — the event is "completed" when all preceding operations in that stream finish
- Host Synchronization: cudaEventSynchronize(event) blocks the host thread until the event completes — more fine-grained than cudaStreamSynchronize which waits for all stream operations
- Query: cudaEventQuery(event) returns cudaSuccess if the event is complete or cudaErrorNotReady if still pending — enables non-blocking polling from the host
Inter-Stream Synchronization:
- cudaStreamWaitEvent: cudaStreamWaitEvent(stream, event, 0) makes all subsequent operations in stream wait until event is complete — establishes a dependency between two streams without blocking the host
- Use Case: launch data preparation on stream 1, record an event, then make stream 2 (which runs the compute kernel) wait on that event — ensures data is ready before computation begins
- Fork-Join Pattern: launch multiple independent kernels on separate streams, create events for each, and have a final stream wait on all events before launching the reduction kernel — maximizes concurrency while maintaining correctness
- Graph Capture: cudaStreamBeginCapture captures a sequence of stream operations into a CUDA Graph — replaying the graph eliminates per-launch overhead (5-10 µs per kernel launch saved)
Timing with Events:
- Elapsed Time: cudaEventElapsedTime(&ms, start, stop) returns the time in milliseconds between two recorded events — GPU-side timing with ~0.5 µs resolution, unaffected by host-side scheduling
- Pattern: cudaEventRecord(start, stream) → launch kernel → cudaEventRecord(stop, stream) → cudaEventSynchronize(stop) → cudaEventElapsedTime(&ms, start, stop) — the standard GPU timing pattern
- Profiling: events provide accurate per-kernel timing without external profiler overhead — essential for performance optimization and regression detection in production code
Common Pitfalls:
- Default Stream Serialization: accidentally using the default stream for one operation serializes the entire GPU — always use explicit streams for concurrent work
- Insufficient Concurrency: launching many small kernels sequentially on one stream wastes GPU resources — distribute across multiple streams to utilize idle SMs
- Missing Synchronization: reading GPU results on the host without synchronization leads to undefined behavior — always synchronize (event, stream, or device) before accessing results
- Pinned Memory Exhaustion: allocating too much pinned memory with cudaMallocHost can degrade system performance — pinned memory can't be swapped, limiting available system memory
Stream and event management is the programmer's primary tool for maximizing GPU utilization — well-structured multi-stream applications achieve 90-95% GPU utilization by overlapping transfers with computation, hiding latency behind concurrent operations, and minimizing synchronization barriers.