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.