Streams and Asynchronous Execution
Streams let a GPU run kernels, host-device transfers, and CPU work concurrently by queuing independent operations that can overlap.
Ordered queues
A stream is an ordered queue of GPU operations. Work within one stream executes in issue order, but operations in different streams may run concurrently if the hardware has the resources. This is the mechanism for overlapping a data transfer over the PCIe or NVLink bus with kernel execution, and for running several small kernels at once when each alone cannot fill the device.
Overlapping transfer and compute
The canonical use is a pipeline: split the input into chunks, and while chunk k is being copied to the device in one stream, the kernel processes chunk k-1 in another stream, and results from chunk k-2 are copied back in a third. With enough chunks the transfer time is almost entirely hidden behind computation. Overlap requires pinned (page-locked) host memory, since only then can the copy engine run asynchronously with the CPU.
- Operations within a stream are ordered; across streams they may overlap.
- Overlapping copy with compute hides bus transfer latency.
- Pinned host memory is required for asynchronous copies.
- Events allow one stream to wait on another for dependencies.
Events and dependencies
When streams do have a dependency (a kernel in stream B needs output from stream A), events record a point in one stream that another can wait on. Events also serve as lightweight timers for measuring kernel duration. The default stream has special synchronizing behavior, so performance-critical code generally creates explicit non-default streams.
# pseudocode: three-way overlap
# for k in chunks:
# cudaMemcpyAsync(d_in[k], h_in[k], HtoD, stream[k%3])
# kernel<<<..., stream[k%3]>>>(d_in[k], d_out[k])
# cudaMemcpyAsync(h_out[k], d_out[k], DtoH, stream[k%3])In practice
A Hyperion parameter sweep that runs many independent field evaluations can assign them to separate streams so transfers and kernels overlap, keeping the GPU busy rather than idling during each copy. This is a routine way to raise device utilization for throughput-oriented workloads.