Multi-Stream Execution, in depth: stream semantics, events, copy-compute overlap, PyTorch side streams, allocator hazards and proving overlap in a trace

What a stream guarantees

Three rules define a stream. First, operations issued to the same stream execute in issue order, and each one starts only after the previous one finishes. Second, operations in different streams have no ordering relationship at all unless you create one. Third, every API call that takes a stream argument returns to the host as soon as the work is queued, so the CPU races ahead of the GPU. Concurrency is permitted, not promised: the device runs work from two streams together only if it has free resources of the right kind.

Those resources come in three types. Copy engines are DMA units that move data across PCIe or NVLink independently of the SMs; most data-centre GPUs have at least two, so a host-to-device and a device-to-host copy can proceed at once. You can read the count from cudaDeviceProp.asyncEngineCount. The SMs themselves run kernels, and two kernels overlap only when the first leaves SMs, registers or shared memory unused. Finally the host driver feeds the device through a limited number of hardware queues; the environment variable CUDA_DEVICE_MAX_CONNECTIONS sets how many the driver uses (the default is 8, the maximum 32). With more streams than connections, streams share queues and can pick up false dependencies on each other.

There is also the default stream. With the legacy default stream, any work issued to stream 0 waits for all other blocking streams and blocks them in turn, which is the most common reason a multi-stream program is secretly serial. Create side streams with cudaStreamNonBlocking, or compile with --default-stream per-thread so each host thread gets its own ordinary default stream. PyTorch's default stream on each device is the legacy one, which is why side-stream code in PyTorch must be explicit about ordering.

Ordering across streams with events

The tool for ordering across streams is the event. cudaEventRecord(e, s) drops a marker into stream s; cudaStreamWaitEvent(t, e, 0) makes everything later issued to stream t wait until that marker is reached. The wait happens on the GPU, so the host does not block. If the event has never been recorded, the wait completes immediately, which is convenient for the first iteration of a double-buffered loop and dangerous if you expected it to block. The classic pattern is a two-buffer pipeline: copy chunk i into buffer i mod 2 on a copy stream while the compute stream processes the other buffer.

cudaStream_t copy_s, comp_s;
cudaStreamCreateWithFlags(&copy_s, cudaStreamNonBlocking);
cudaStreamCreateWithFlags(&comp_s, cudaStreamNonBlocking);

cudaEvent_t ready[2], done[2];
for (int b = 0; b < 2; ++b) {
    cudaEventCreateWithFlags(&ready[b], cudaEventDisableTiming);
    cudaEventCreateWithFlags(&done[b],  cudaEventDisableTiming);
}

// h_in must be pinned (cudaMallocHost) or the copies are not asynchronous.
for (int i = 0; i < n_chunks; ++i) {
    int b = i % 2;
    cudaStreamWaitEvent(copy_s, done[b], 0);          // buffer b no longer read by compute
    cudaMemcpyAsync(d_in[b], h_in + (size_t)i * chunk, chunk * sizeof(float),
                    cudaMemcpyHostToDevice, copy_s);
    cudaEventRecord(ready[b], copy_s);

    cudaStreamWaitEvent(comp_s, ready[b], 0);         // data for chunk i has landed
    process<<<grid, block, 0, comp_s>>>(d_in[b], d_out + (size_t)i * chunk, chunk);
    cudaEventRecord(done[b], comp_s);
}
cudaStreamSynchronize(comp_s);

Two events per buffer encode both directions of the hazard. ready[b] stops the kernel from reading a buffer before the copy finishes. done[b] stops the next copy from overwriting a buffer the kernel is still reading. Dropping either produces a program that passes small tests and fails under load, because the race only shows when timing shifts.

Side streams in PyTorch

PyTorch wraps the same ideas. torch.cuda.Stream() creates a stream, the context manager torch.cuda.stream(s) routes every operation inside it to s, a.wait_stream(b) is record-plus-wait in one call, and torch.cuda.Event exposes record, wait, synchronize and elapsed_time. The most useful application is a prefetcher that copies batch n+1 on a side stream while the model trains on batch n.

import torch

class CudaPrefetcher:
    # Overlap the H2D copy of the next batch with compute on the current one.
    # The DataLoader must use pin_memory=True or non_blocking copies are synchronous.

    def __init__(self, loader, device="cuda"):
        self.it = iter(loader)
        self.device = device
        self.stream = torch.cuda.Stream(device=device)
        self.next_batch = None
        self._preload()

    def _preload(self):
        try:
            batch = next(self.it)
        except StopIteration:
            self.next_batch = None
            return
        with torch.cuda.stream(self.stream):
            self.next_batch = [t.to(self.device, non_blocking=True) for t in batch]

    def __iter__(self):
        return self

    def __next__(self):
        if self.next_batch is None:
            raise StopIteration
        main = torch.cuda.current_stream(self.device)
        main.wait_stream(self.stream)          # GPU-side: compute waits for the copy
        batch = self.next_batch
        for t in batch:
            t.record_stream(main)              # tell the allocator who else uses this memory
        self._preload()                        # start copying the following batch
        return batch

The record_stream line is the one people omit. PyTorch's caching allocator tracks memory per stream: a block allocated on the side stream is, by default, considered free for reuse on that stream as soon as the Python tensor dies. If the main stream is still running kernels that read it, a later copy on the side stream can overwrite the block mid-read. record_stream(main) tells the allocator to wait until the main stream's work queued at the time of the free has completed. The cost is that such memory is returned later, which raises peak usage slightly; the alternative, allocating on the main stream and only copying on the side stream, avoids the call at the price of more bookkeeping. The allocator's behaviour is covered further in CUDA memory pools.

Hidden synchronisation

Most failed attempts at overlap are not races; they are hidden synchronisation points that make the host wait, so it never gets far enough ahead to enqueue the overlapping work. The usual suspects:

• Copies from pageable host memory. The driver stages them through a pinned bounce buffer and the call does not return until the staging is done. Pin the source; see page-locked host memory.
• cudaMalloc and cudaFree, which can synchronise the device. Use a stream-ordered allocator or PyTorch's caching allocator and keep the steady state allocation-free.
• Anything that reads a GPU value on the host: .item(), .cpu(), print(tensor), Python if on a tensor, and ops with data-dependent output shapes such as torch.nonzero or boolean-mask indexing.
• Any work on the legacy default stream while side streams use blocking semantics.
• torch.cuda.synchronize() left in from a timing experiment.

PyTorch can flag these for you: torch.cuda.set_sync_debug_mode("warn") emits a warning whenever an operation synchronises with the device, which is the fastest way to find the one stray .item() in a training loop.

Kernel concurrency, communication and priorities

Copy-compute overlap is the reliable win. Kernel-kernel overlap is rarer than diagrams suggest. A large matrix multiply launches enough thread blocks to occupy every SM several times over, so a kernel on a second stream simply waits for blocks to drain; it may slip into the tail of the first kernel, but the total time hardly moves. Kernel concurrency pays off when individual kernels are too small to fill the device: many tiny per-expert matmuls, independent branches of a multi-tower model, small inference requests from different users, or a communication kernel that uses only a few SMs alongside compute.

That last case is where multi-stream execution matters most in distributed training. NCCL collectives run on their own stream, so the all-reduce for one bucket of gradients can proceed while backward computes the next bucket. Because NCCL kernels also occupy SMs, frameworks sometimes limit their footprint or carefully order launches; Megatron-LM, for example, documents setting CUDA_DEVICE_MAX_CONNECTIONS=1 for some overlap configurations so the hardware sees kernels in launch order. The full treatment is in collective communication overlap.

Stream priorities help when one stream is latency-sensitive. cudaDeviceGetStreamPriorityRange returns the allowed range, where a numerically lower value is a higher priority, and cudaStreamCreateWithPriority or torch.cuda.Stream(priority=-1) creates a high-priority stream. Priority decides which waiting thread blocks the scheduler dispatches next; it does not stop blocks that are already running, so a long low-priority kernel still delays a high-priority one by up to one wave of blocks.

Worked example: hiding the input copy

Take a vision training job. Each step copies a 64 MB batch of uint8 images to the GPU, and forward plus backward take 12 ms. Over a PCIe Gen4 x16 link the copy from pinned memory runs at roughly 20 to 25 GB/s in practice, so it takes about 2.6 to 3.2 ms. On one stream a step costs about 12 + 3 = 15 ms. With the prefetcher, the copy of batch n+1 runs on a copy engine during the 12 ms of compute, and the step time falls to about 12 ms: a 20 percent throughput gain for some fifty lines of code.

Now suppose the batch is pageable. The copy call blocks the host for the whole transfer, the host cannot enqueue the next forward pass during it, and the trace shows the side stream's copy and the main stream's kernels laid end to end, exactly as on one stream. Switching pin_memory=True on in the DataLoader restores the overlap. The lesson generalises: measure before and after, because multi-stream code that does not overlap looks identical in source to code that does.

Proving overlap in a trace

Overlap is a claim about time, so verify it in a timeline, not with a stopwatch around the loop. Nsight Systems shows one row per stream plus rows for memory copies:

nsys profile --trace=cuda,nvtx,osrt -o step_trace python train.py --max-steps 30

In the trace, look for three things. Copy bars on the side stream should sit directly beneath kernel bars on the compute stream. Host-side, the CUDA API row should show launches running ahead of the GPU, not long cudaStreamSynchronize or cudaMemcpy calls. And the GPU utilisation row should have no gaps between steps. The PyTorch profiler shows the same streams in its Chrome trace; LLM trace analysis walks through reading one.

Failure modes

The last row deserves a note. CUDA graphs can capture multi-stream work, but only if every side stream forks from and joins back to the capturing stream through events; otherwise the capture is invalid. CUDA streams and graphs covers capture in detail.

Trade-offs

Multi-stream code buys throughput with complexity. Each additional stream adds a dependency to reason about, a lifetime rule for every tensor that crosses it, and a new class of bug that unit tests rarely catch. Keep the design small: one compute stream, one copy stream for inputs, perhaps one for outputs or communication, each with an owner and a comment that states what waits on what. Prefer library mechanisms that already encapsulate the pattern, such as DataLoader pinning, DDP's bucketed all-reduce, or an inference server's per-request streams, over hand-written schemes. And remember that if the GPU is already compute-bound at high utilisation with no gaps in the trace, more streams cannot help; the win comes only from filling idle time that you can see.

What to do next

1. Profile one training or inference step with Nsight Systems and note every gap on the compute stream and its cause.
2. Turn on torch.cuda.set_sync_debug_mode("warn") for a few steps and remove the synchronising calls it reports from the hot loop.
3. Confirm the input path uses pinned memory and non_blocking=True.
4. Add the prefetcher above, with wait_stream and record_stream, and re-profile to check the copies sit under compute.
5. Compare outputs bit for bit against the single-stream version on a fixed seed before you keep the change.
6. Write down each stream's role and dependencies next to the code so the next person does not remove a wait that looks redundant.