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.