cuda streams asynchronous gpu

**CUDA Streams and Asynchronous Execution** enable **concurrent kernel launches, memory transfers, and host-device synchronization, hiding latencies and improving GPU utilization through fine-grained task scheduling and pipelining.** **Stream Concept and Execution Model** - **CUDA Stream**: Ordered queue of GPU operations (kernels, memory transfers, callbacks). Operations within stream execute sequentially; operations in different streams may execute concurrently. - **Default Stream (Stream 0)**: All operations enqueued to default stream if not specified. Default stream synchronizes with all other streams (implicit barrier). Avoid for concurrent execution. - **Non-Default Streams**: Streams 1, 2, 3,... execute independent of default and each other. Multiple kernels in different streams execute in parallel (if hardware permits). - **Hardware Concurrency Limits**: Modern GPUs (Volta+) support 32-128 concurrent streams. Depends on SM count, kernel complexity, resource contention. **CUDA Events and Synchronization Primitives** - **cudaEvent Creation**: cudaCreateEvent() allocates event. Event timestamps GPU wall-clock time for precise measurements. - **Event Recording**: cudaEventRecord(event, stream) inserts event into stream. GPU records timestamp when reaching that point in stream. - **Event Query**: cudaEventQuery(event) checks if event reached (non-blocking). cudaEventSynchronize(event) blocks host until event reached. - **Elapsed Time Measurement**: (event_end_time - event_start_time) gives kernel duration. More accurate than host timing due to GPU pipeline effects. **Multi-Stream Concurrency and Concurrency Limitations** - **Maximum Concurrent Kernels**: SM (Streaming Multiprocessor) can execute multiple kernels simultaneously if occupancy permits. Typically 8-32 concurrent kernels depending on register/shared memory usage. - **Occupancy Trade-off**: Multiple lightweight kernels co-execute, each using fewer registers. Heavy kernels (high register count) limit concurrent kernel count. - **Load Balancing**: Kernel scheduling across SMs dynamic. GPU scheduler assigns kernels to free SMs, migrates work as kernels complete. - **Resource Contention**: L2 cache, register file, shared memory shared across all kernels. High-memory-footprint kernels increase latency for light concurrent kernels. **Asynchronous Memory Copy with Compute Overlap** - **cudaMemcpyAsync**: Non-blocking memory transfer returns immediately after DMA queued. Compute kernels in other streams progress concurrently with memcpy. - **DMA Engine Limitations**: Most GPUs support 1 host-to-device and 1 device-to-host DMA concurrently (bidirectional, but only 1 H2D and 1 D2H). Cannot overlap multiple H2D transfers. - **Pinned vs Pageable Memory**: Pinned (DMA-able) host memory required for asynchronous transfers. Pageable memory requires intermediate staging → performance loss. - **PCIe Bandwidth**: 16x PCIe 3.0 = 16 GB/s. Theoretical bidirectional = 8 GB/s each direction (practical: 12-14 GB/s unidirectional due to protocol overhead). **Overlap Efficiency** - **Perfect Overlap Conditions**: Kernel computation (10ms) + simultaneous memcpy (10ms) = 10ms total (no additional delay). Requires computation and memcpy duration matched. - **Bottleneck Analysis**: If memcpy faster than kernel (e.g., 5ms memcpy, 10ms kernel), GPU idles 5ms before next memcpy. Padding with extra work hides idle. - **Pipelining**: Stage K computes while Stage K-1 copies output to host. Multiple overlapping stages maintain GPU saturation. - **Profiler Visualization**: Nsight Systems shows timeline of kernel, memcpy, host activities. Overlapping activities side-by-side visualize concurrency. **Stream Priority and Quality of Service** - **Stream Priority**: cudaStreamCreateWithPriority() assigns priority (1-32, higher = higher priority). GPU scheduler prefers high-priority streams. - **Priority Effectiveness**: Only works if GPU has spare resources (not fully saturated). Under full load, priority irrelevant. - **Use Cases**: Critical kernels (audio processing) prioritized over background kernels (profiling). Real-time applications leverage priorities. **Best Practices for Hiding PCIe Latency** - **Batch Transfers**: Multiple cudaMemcpyAsync() calls in rapid succession (all within same stream) amortize PCIe latency. - **Unpinned Memory Workaround**: If pinned memory unavailable, use temporary pinned buffer, copy into unpinned (slower but functional). - **Bidirectional Pipeline**: Overlap H2D (input data), compute (processing), D2H (output data) in 3-stage pipeline. Maintains GPU utilization across phases. - **Persistent Kernels**: Long-running kernels with internal loops reduce kernel launch overhead. Single kernel overlaps internal stages vs multiple kernel launches.