Stream Ordering: What a Queue Guarantees and What It Does Not

A stream is an in-order queue of commands executed by one device. That sentence is the entire contract, and most misuse of the abstraction comes from reading more into it than it contains. The queue executes its commands in submission order, one after another. That order binds only the commands in this queue. The queue belongs to one device. Nothing about it says anything about other queues, other devices, or the host.

Holding those four statements apart is the work. The first, in-order execution, is what makes the programming model usable, because a kernel submitted after a copy can assume the copy happened. The second, that the order stops at the queue’s edge, is what makes a second stream necessary whenever two pieces of work must proceed at the same time. The third and fourth define where the guarantees stop, and every explicit synchronization operation in the API exists to cross one of those boundaries.

The queue and the front end

Commands enter a stream through a host-side submission path. The runtime writes a command packet into a ring buffer in host memory, advances a write pointer, and rings a doorbell register on the device. The device’s front end reads the packet, decodes it, and dispatches it. The host call returns as soon as the packet is in the ring and the doorbell has been rung; it does not wait for the device to read it.

That structure explains a cost that surprises people measuring launch overhead: a kernel launch is not a single instruction. It is a packet of tens of bytes written to memory, a memory barrier to make it visible, a doorbell write over the bus, and on some paths a check that the ring is not full. The dominant term for small kernels is often the host-side cost of building the packet rather than the device-side cost of scheduling it.

The ring has a finite capacity. When it fills, the submission call blocks until the device has consumed enough entries. This is the mechanism by which a host thread that submits faster than the device can execute ends up throttled, and it is why a program that submits millions of tiny operations shows a host-bound profile even though the device is idle between them.

Commands are consumed in order and dispatched to different engines. A copy packet goes to a copy engine, a kernel packet goes to the compute front end, and each engine carries out the work its packet names. Separate engines do not mean overlapping commands within the queue: the next command is not started until the current one has completed, which is exactly what makes a kernel submitted after a copy safe to write. The concurrency a device offers is obtained by issuing work on more than one stream, and the concurrency a single kernel contains comes from its own threads, not from the queue.

#include <cuda_runtime.h>

__global__ void scale(float *x, float k, int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) x[i] *= k;
}

__global__ void add(const float *x, const float *y, float *z, int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) z[i] = x[i] + y[i];
}

void pipeline(float *a, float *b, float *c, int n, cudaStream_t s0, cudaStream_t s1)
{
    const int t = 256, g = (n + t - 1) / t;

    /* b is unrelated to a, so this launch is free to overlap the s0
       work: two queues, and no ordering between them. */
    scale<<<g, t, 0, s1>>>(b, 0.5f, n);

    /* Same stream: the second launch sees the first one's writes. */
    scale<<<g, t, 0, s0>>>(a, 2.0f, n);
    scale<<<g, t, 0, s0>>>(a, 3.0f, n);

    /* The result depends on both a and b, so it needs an explicit edge.
       cudaEventRecord on s0 and cudaStreamWaitEvent on s1 create one
       without blocking the host; without it this launch would race the
       kernels above. */
    cudaEvent_t done;
    cudaEventCreateWithFlags(&done, cudaEventDisableTiming);
    cudaEventRecord(done, s0);
    cudaStreamWaitEvent(s1, done, 0);

    /* s1 is already ordered after its own scale, so both inputs are
       ready when this runs. */
    add<<<g, t, 0, s1>>>(a, b, c, n);
}

Where the host sits relative to the queue

The host is outside the stream. A kernel launch returns when the packet is queued, so host code that immediately reads a value the kernel was supposed to compute reads the old value. This is not a subtle race; it is the ordinary behaviour of an asynchronous interface, and the reason every runtime provides a synchronization call that blocks the host until the stream drains.

The distinction that matters in practice is between synchronizing the host and synchronizing the device. Blocking the host on a stream is a full stop: the calling thread does nothing until the device has finished. Making one stream wait on another, through an event, keeps the host running and lets the device maintain its own ordering. A pipeline that synchronizes the host after every stage gets none of the benefit of having stages.

Events are the primitive that crosses stream boundaries. Recording an event captures a point in a stream’s execution; waiting on that event in another stream creates an edge in the dependency graph without involving the host. The device-side effect is that the waiting stream’s subsequent commands are held until the recorded point completes. The host-side effect is that the wait call returns immediately, having queued the dependency rather than blocked on it.

There is a cost to each event. Recording one is a command in the stream, and on some architectures the completion of a recorded event requires a write to host-visible memory that the device must flush. A pipeline with an event between every pair of small kernels spends a meaningful fraction of its time on those flushes. Coarser dependencies, where one event covers a batch of work, are cheaper and usually sufficient.

Two streams, no ordering

The most common bug in stream-based code is the assumption that stream 1’s work begins after stream 0’s work finishes. It does not. Two streams are two independent queues, and the device may interleave their commands arbitrarily. A kernel in stream 1 that reads a buffer written by a kernel in stream 0 races with it, and the outcome depends on scheduling.

The failure is intermittent in the worst way: on an idle device the two kernels often do run in submission order, so the bug is invisible until another stream is added or the machine is loaded. Nothing in the API detects it. There is no dataflow analysis that inspects the buffers a kernel touches and infers a dependency; that information is not available to the runtime, because a kernel’s accesses are computed at run time from arguments.

Memory allocated on the device does not carry a stream affinity. A buffer created by an ordinary allocation can be touched by any stream, and the API neither records which stream wrote it nor enforces anything. Some runtimes offer pooled allocations that are scoped to a stream for the purposes of reuse ordering, but the scope affects when the memory may be recycled, not when writes become visible.

Two specific hazards deserve naming. The first is a free while work is in flight: releasing device memory is legal once the host has released it, regardless of whether a queued kernel still references it, and the allocator may hand the same memory to the next allocation while the old kernel is still running. The second is a copy out of a buffer that a kernel in another stream is writing, which produces a truncated or interleaved result that looks like a data bug rather than an ordering bug.

When ordering guarantees are stronger than the stream

Certain operations carry device-wide or system-wide effects that the stream model does not describe. A device-wide synchronization call waits for all streams on the device, not just the one it was issued on. Operations that change device state, such as setting a cache configuration or resetting the device, have similar scope.

The other direction is a weakening. Some operations are documented as completing asynchronously with respect to the stream in which they were enqueued, notably certain allocations and frees, which may be deferred until a later synchronization point. An application that assumes a free has completed because the stream drained can be wrong on a platform where the free was deferred, and the symptom is memory that is not returned to the driver for longer than expected rather than a correctness failure.

Across devices the stream model says nothing at all. A stream belongs to a device, and work in a stream on device 0 has no ordering relationship with a stream on device 1 unless an event is recorded on one and waited on by the other, which the runtime supports for peer-capable pairs. Without that, a kernel on device 1 reading a buffer that device 0’s kernel is filling is a plain race, and the topology matters for whether the read is even coherent; see peer-to-peer transfers and link topology.

Synchronization scope and the memory model

Stream ordering is a statement about execution order, and execution order is not the same as visibility. On a device with caches that are not coherent with the host, a write performed by a kernel may sit in a cache that the host cannot see even after the kernel has completed and the stream has drained. The synchronization calls that drain a stream are documented to make prior writes visible to the host, and that flush is part of what the call costs.

The same distinction recurs inside kernels. Two threads in different blocks have no ordering between them through the stream model, because the stream model does not describe thread execution at all. Establishing order between them requires the memory-model machinery of system-scope atomics and memory ordering across the bus, where the scope of an atomic operation determines whether a write becomes visible to a host thread, to another device, or only within the current one.

A useful way to hold the whole model is as three nested scopes. Within a stream, every command completes before the next begins. Across streams on one device, order exists only where an event creates it. Across devices and the host, order exists only where an explicitly scoped synchronization operation creates it, and the strength of that operation depends on the memory system’s coherence properties rather than on the stream abstraction.

The buffer lifetime questions that streams create are covered in pinned host memory and asynchronous transfers, and the way the address space makes cross-device pointers meaningful at all is described in unified virtual addressing and pointer provenance.