cuda stream concurrency

**CUDA Stream Concurrency** is the **GPU programming technique that uses multiple independent execution streams to overlap kernel execution, memory transfers, and host computation — enabling the GPU to simultaneously execute a kernel on one stream while transferring data for the next kernel on another stream, hiding memory transfer latency and increasing overall GPU utilization from 40-60% to 85-95%**. **The Problem Streams Solve** Without streams, GPU execution is serialized: copy data H→D, launch kernel, copy results D→H, repeat. The GPU sits idle during transfers, and the PCIe/NVLink bus sits idle during computation. On a typical workload, the GPU may be actively computing only 50% of the time. **How Streams Work** A CUDA stream is an ordered sequence of operations (kernels, memcpy, events) that execute in issue order within the stream but can execute concurrently with operations in other streams. The hardware has independent engines: - **Compute Engine(s)**: Execute kernels (multiple kernels from different streams can run concurrently if they don't fully occupy the GPU) - **Copy Engine(s)**: Execute H→D and D→H transfers (modern GPUs have 2 copy engines, enabling simultaneous upload and download) **Double-Buffering Pattern** ``` Stream 0: [Copy chunk 0 H→D] [Kernel chunk 0] [Copy chunk 0 D→H] Stream 1: [Copy chunk 1 H→D] [Kernel chunk 1] [Copy chunk 1 D→H] Stream 2: [Copy chunk 2 H→D] [Kernel chunk 2] ... ``` While the kernel processes chunk 0, chunk 1's data is being uploaded. While chunk 1's kernel runs, chunk 0's results download and chunk 2 uploads. The GPU's compute and copy engines are always busy. **Synchronization Mechanisms** - **cudaStreamSynchronize(stream)**: Blocks the host until all operations in the specified stream complete. - **cudaEventRecord / cudaStreamWaitEvent**: Records a timestamp in one stream; another stream can wait on that event. Enables fine-grained inter-stream dependencies without blocking the host. - **cudaDeviceSynchronize()**: Wait for all streams — the nuclear option, use sparingly. **Concurrency Limits** - **Kernel Concurrency**: Multiple small kernels from different streams can execute simultaneously if they collectively don't exceed the GPU's SM count. A single kernel that occupies all SMs blocks other kernels regardless of streams. - **False Dependencies**: Operations issued on the default stream (stream 0) synchronize with all other streams, destroying concurrency. Always use explicit non-default streams. - **Hardware Queue Depth**: The GPU has finite hardware queues for scheduling. Excessive streams (>16-32) provide no additional benefit and add scheduling overhead. CUDA Stream Concurrency is **the technique that transforms the GPU from a batch processor into a pipelined system** — keeping every hardware engine continuously fed with work by interleaving independent operations across multiple streams.