CUDA Stream Concurrency

CUDA Streams and Asynchronous GPU Execution are the programming mechanisms that enable overlapping of multiple GPU operations — kernel execution, memory transfers, and host computation — through independent queues (streams) of work, extracting maximum hardware utilization by keeping all GPU subsystems (compute, copy engines, host CPU) busy simultaneously rather than executing operations one-after-another.

The Default (Synchronous) Problem

Without streams, GPU operations execute sequentially: copy H→D, kernel launch, copy D→H, repeat. The GPU compute units are idle during transfers, and the copy engines are idle during kernel execution. For a kernel taking 5 ms and transfers taking 2 ms each, utilization is only 5/(5+2+2) = 56%.

Stream Semantics

• A stream is an in-order queue of GPU operations. Operations within the same stream execute sequentially.
• Operations in different streams may execute concurrently if hardware resources are available.
• The default stream (stream 0) synchronizes with all other streams (legacy behavior) or is per-thread independent (per-thread default stream).

Overlap Patterns

Double Buffering (Ping-Pong):

Stream A: [Copy H→D chunk 0] [Kernel chunk 0] [Copy D→H chunk 0]
Stream B:                      [Copy H→D chunk 1] [Kernel chunk 1] [Copy D→H chunk 1]

While the kernel processes chunk 0, the copy engine simultaneously transfers chunk 1 to the device. Total time approaches max(compute, transfer) instead of compute + transfer.

Multi-Stream Pipeline: With N streams and N data chunks, all pipeline stages run simultaneously. The steady-state throughput is limited by the slowest stage.

Hardware Requirements

• Separate Copy Engines: Modern GPUs have 2-3 independent copy engines (H→D, D→H, D→D / P2P). Bidirectional transfers overlap with each other AND with kernel execution.
• Concurrent Kernels: GPUs can execute multiple small kernels simultaneously on different SMs when a single kernel doesn't saturate all SMs. Up to 16-128 concurrent kernels depending on architecture.
• Hyper-Q (MPS): Multiple CPU processes or threads can submit to independent hardware queues, enabling fine-grained concurrent kernel execution without false dependencies.

Synchronization Primitives

• cudaStreamSynchronize(stream): Host blocks until all operations in the stream complete.
• cudaEventRecord / cudaStreamWaitEvent: Fine-grained inter-stream dependencies. Stream B can wait for a specific event recorded in Stream A without synchronizing with the host.
• cudaEventElapsedTime: Measure GPU-side elapsed time between two events for accurate kernel timing.

Common Pitfalls

• Implicit Synchronization: Some CUDA API calls (cudaMalloc, cudaMemcpy without Async) implicitly synchronize the device, serializing all streams. Use Async variants exclusively in multi-stream code.
• Page-Locked Memory: Async transfers require pinned (page-locked) host memory. Pageable memory forces synchronous copies.

CUDA Streams are the concurrency abstraction that converts GPU programming from sequential batch processing into pipelined, overlapped execution — extracting the last 20-40% of hardware utilization that separates a well-optimized GPU application from a naive one.