CUDA Streams
CUDA Streams are queues of GPU operations that execute in order within the stream but potentially overlap with operations in other streams — enabling concurrent execution of kernels, memory transfers, and other GPU operations to maximize GPU utilization.
Default Stream Behavior
- Without streams: All operations go to default stream → execute sequentially.
- GPU underutilized: Memory transfer blocks kernel execution.
- With streams: Overlap transfers and kernels → higher utilization.
Creating and Using Streams
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// Launch kernel on stream1
my_kernel<<<grid, block, 0, stream1>>>(...);
// Transfer data on stream2 concurrently
cudaMemcpyAsync(dst, src, size, cudaMemcpyHostToDevice, stream2);
// Wait for both streams
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);Overlap Patterns
Transfer-Compute Overlap:
- Stream 1: Copy batch N to device.
- Stream 2: Compute on batch N-1 (already transferred).
- Requires pinned (page-locked) host memory for async transfer.
Kernel-Kernel Overlap:
- Multiple independent kernels on different streams execute concurrently.
- Limited by SM occupancy — if one kernel uses all SMs, no overlap possible.
CUDA Events for Timing and Synchronization
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventRecord(start, stream1); // Mark start in stream1
// ... kernel ...
cudaEventRecord(stop, stream1);
cudaEventSynchronize(stop);
float ms;
cudaEventElapsedTime(&ms, start, stop);Stream Dependencies
- `cudaStreamWaitEvent(stream2, event, 0)`: Stream2 waits for event from stream1.
- Enables fine-grained inter-stream dependencies.
CUDA Graphs (2019+)
- Record stream operations into a graph → replay graph with minimal overhead.
- Eliminates CPU-GPU synchronization overhead for repeated workloads.
- 5–30% speedup for inference workloads with many small kernels.
CUDA streams are the key to achieving high GPU utilization in production inference pipelines — overlapping data transfer with compute through multi-stream design can recover 20–40% performance that single-stream sequential execution leaves unused.