CUDA Streams: Asynchronous Execution and Concurrency

Summary: A CUDA stream is an in-order queue of GPU ops. Overlap H→D, kernels, and D→H across streams — and avoid the default-stream trap that serializes everything.

Kernels look fine. The timeline doesn’t.

Nsight shows H→D, then a kernel, then D→H, then the same three bars again. SMs sit idle while the copy engine works. The copy engine sits idle while SMs work. Nothing is “wrong” with the kernel — the schedule never stacked independent engines.

A CUDA stream is an in-order queue of GPU work: kernels, memcpys, events. Same stream → serial. Different streams → the hardware may run phases together if they need different engines (or spare SM capacity).

Two numbers

1. Serial · wall ≈ 3N for N balanced chunks
One stream (or everything on the default stream): each chunk pays H→D + kernel + D→H end-to-end. Engines take turns.

2. Pipelined · wall ≈ N + 2 in steady state
With enough non-default streams, chunk k’s kernel overlaps chunk k+1’s H→D and chunk k−1’s D→H. Three engines busy; wall time drops toward the longest phase, not the sum.

Flip the tapes until that feels mechanical — not magical.

Pipeline: three engines, N streams

Same four chunks. Change stream count and watch H→D · kernel · D→H fill in.

code

// Async API + explicit stream — the building blocks
cudaMemcpyAsync(dst, src, bytes, cudaMemcpyHostToDevice, stream);
myKernel<<<grid, block, 0, stream>>>(args);
cudaMemcpyAsync(host, dst, bytes, cudaMemcpyDeviceToHost, stream);

Kernel launch is already async w.r.t. the host. Overlap lives or dies on which stream you pass and whether the host buffer is pinned.

The default stream is a trap

Every process has a pre-created default stream (stream 0). Bare cudaMemcpy and myKernel<<<g,b>>> land there. Legacy default semantics: that stream implicitly fences every other stream in the process.

You can wire four beautiful non-default streams and still open a process-wide gap with one forgotten sync memcpy.

Synchronization without the sledgehammer

Eventually the host needs “is it done?” Use the lightest tool that fits:

Call	Blocks
`cudaDeviceSynchronize()`	Everything on the device
`cudaStreamSynchronize(s)`	One stream’s queue
`cudaStreamQuery(s)`	Nothing (poll)
Events + `cudaStreamWaitEvent`	One stream waits on another — host free

code

cudaEvent_t done;
cudaEventCreate(&done);

myKernel<<<grid, block, 0, streamA>>>(d_data);
cudaEventRecord(done, streamA);

// streamB waits on GPU work, not on the host thread
cudaStreamWaitEvent(streamB, done, 0);
consumerKernel<<<grid, block, 0, streamB>>>(d_data);

Events also measure GPU time: record before/after, then cudaEventElapsedTime — no CPU-clock skew.

Pitfalls that invent serialization

Same streams idea. Four ways people accidentally kill concurrency.

Stream priorities

code

int lo, hi;
cudaDeviceGetStreamPriorityRange(&lo, &hi);
cudaStream_t hot;
cudaStreamCreateWithPriority(&hot, cudaStreamNonBlocking, hi);

High priority is not preemption mid-kernel — it wins at the next launch boundary. Useful for a latency-sensitive inference path beside a bulk training stream.

PyTorch

code

import torch

s1 = torch.cuda.Stream()
s2 = torch.cuda.Stream()

with torch.cuda.stream(s1):
    a = torch.randn(1024, 1024, device='cuda')
    b = a @ a.T

s2.wait_stream(s1)
with torch.cuda.stream(s2):
    c = b.softmax(dim=-1)

torch.cuda.synchronize()

The cache allocator tracks per-stream ownership and inserts waits so .to() across streams usually “just works.” Custom CUDA ops must still record on the current PyTorch stream or downstream ops will race.

What to do

Never leave critical path work on the legacy default stream — explicit streams (or per-thread default) everywhere you care about overlap.
Pin host memory for cudaMemcpyAsync — pageable “async” is often a silent sync tax.
Pipeline H→D / kernel / D→H with 2–4 streams first; measure before inventing sixteen.
Prefer events and stream waits over cudaDeviceSynchronize in production loops.
Keep allocs out of the hot path — pool or cudaMallocAsync; setup-time malloc only.

code

cudaStream_t streams[4];
for (int i = 0; i < 4; i++)
  cudaStreamCreateWithFlags(&streams[i], cudaStreamNonBlocking);

for (int i = 0; i < nChunks; i++) {
  auto s = streams[i % 4];
  cudaMemcpyAsync(d[i], h[i], bytes, cudaMemcpyHostToDevice, s);
  kernel<<<grid, block, 0, s>>>(d[i]);
  cudaMemcpyAsync(out[i], d[i], bytes, cudaMemcpyDeviceToHost, s);
}
for (int i = 0; i < 4; i++) cudaStreamSynchronize(streams[i]);

When streams are the wrong tool

Multi-process GPU sharing → CUDA MPS, not more streams in one process.
Cross-device work → NCCL / peer memcpy; streams do not cross devices.
One long kernel, nothing to overlap → fix the kernel; streams only hide independent latency.

Streams keep the GPU fed. Get the compute right first, then stop wasting engines between phases.