Command Queues, Events, and Cross-Stream Ordering

A kernel launch returns before the kernel runs. That is the entire point of the interface: the host enqueues work into a buffer the device consumes, and the host thread continues. Everything difficult about GPU programming at scale follows from the question of what the host is allowed to assume about that buffer: what has happened, what has not, and what may happen concurrently.

In-order is a statement about completion, not about execution

An in-order queue is the default in CUDA streams and in OpenCL queues created without CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE. The guarantee is that operations complete in enqueue order: operation n finishes before operation n+1 begins. It is not a statement that they run one at a time on the whole device, because the device may overlap the tail of one operation with the head of the next where doing so cannot be observed. It is a statement about the observable ordering, and everything the program can detect is ordered.

The practical value is that a sequence of dependent kernels needs no explicit synchronisation. A kernel that reads what the previous kernel wrote is safe by construction, because the queue serialises them. A copy enqueued after a kernel sees that kernel’s output for the same reason. The default stream, also called the null stream, behaves as if it were such a queue and, under the legacy semantics, additionally serialises with all other blocking streams, a property that makes single-stream code trivially correct and that quietly prevents any concurrency until the program is restructured.

The cost is that in-order means exactly that. A small kernel that cannot start until a large transfer finishes will wait, even though it touches disjoint memory. Two independent matrix tiles that could be computed simultaneously are computed in sequence because they were enqueued in sequence. The queue has no way to know that the dependency is absent, so the program must express the absence by using more than one queue.

Events are the edges of the dependency graph

An event is a handle to a point in a queue. cudaEventRecord or clEnqueueMarkerWithWaitList inserts into the queue a node that completes when everything before it has completed, and other operations can wait on that node instead of waiting on the whole queue.

CUDA exposes this as cudaStreamWaitEvent(stream, event): the operations subsequently enqueued to stream are held until event has completed. A CUDA stream callback or a host-side cudaEventSynchronize gives the host the same ability. OpenCL spells it as a wait list: clEnqueueNDRangeKernel takes a count and an array of events to wait for before starting, which is a strictly more expressive form, since one command can depend on several events at once.

The resulting structure is a directed acyclic graph whose nodes are commands and whose edges are event dependencies. The queue is a convenience that adds edges automatically between consecutive commands; events let a program add edges where they exist and, by using separate queues, decline to add edges where they do not.

// Two independent halves on two streams, joined before the reduction.
cudaMemcpyAsync(d_a, h_a, N * sizeof(float), cudaMemcpyHostToDevice, s1);
cudaMemcpyAsync(d_b, h_b, N * sizeof(float), cudaMemcpyHostToDevice, s2);
kernel_a<<<gs, bs, 0, s1>>>(d_a, d_out);
kernel_b<<<gs, bs, 0, s2>>>(d_b, d_out + N);
cudaEventRecord(ev1, s1);          // completes when kernel_a has finished
cudaEventRecord(ev2, s2);          // completes when kernel_b has finished
cudaStreamWaitEvent(s1, ev2);      // s1 also waits for s2's kernel
reduce<<<gs, bs, 0, s1>>>(d_out, d_result, 2 * N);

The join is what makes the graph a graph rather than a fan-out. Without it, reduce would be ordered only against kernel_a and could read kernel_b’s half of the buffer before it was written.

Vulkan inverts the default

Vulkan compute makes no promise about memory visibility between commands. Commands in a queue begin in submission order, but that says nothing about when one command’s writes become visible to the next, and nothing prevents them from overlapping; the program must insert an explicit barrier or an event dependency, and the specification requires it to declare not merely ordering but the specific memory hazards the barrier covers, naming the source and destination access masks and pipeline stages.

The consequence is that the same program, translated from a CUDA stream sequence, is racy. A dispatch reading a buffer written by a previous dispatch has no memory dependency unless a compute-to-compute barrier with the right access masks sits between them. The barrier is heavier than an event: it blocks, it states which accesses it orders, and an underspecified barrier is a specification violation rather than a conservative estimate.

The design suits a graphics API whose driver has always been handed an explicit command buffer, and where the barrier is the mechanism that lets a tile-based renderer reorder passes. It is hostile to the assumption, built up from CPU threading, that enqueue order implies visibility.

Synchronisation points and what they cost

Four host-side mechanisms appear across these APIs, and they differ in what they wait for and how much they cost.

A queue flush submits enqueued commands to the device without waiting. A queue finish or cudaStreamSynchronize waits for everything enqueued so far in that queue to complete. A wait on a single event waits for one point. A blocking transfer is a finish of one command together with the copy.

The cost of a finish is not the cost of the work; it is the loss of overlap. Calling cudaDeviceSynchronize after every launch converts the device into a synchronous coprocessor and removes its entire advantage. The expensive case is not one such call but the pattern of using one as a debugging aid and leaving it in.

A subtler cost is host-side latency: the round trip through the driver to observe completion is on the order of microseconds even when the work has long finished, so a loop that launches one small kernel and waits for it pays that latency per iteration. Batching many launches behind a single wait amortises it.

Hazard patterns the queue does not cover

Even an in-order queue and a correct event graph leave three classes of bug.

The first is host-side mutation of device-mapped memory while a kernel reads it. The host is not in the queue, so nothing orders it. The second is reuse of a buffer after enqueue but before completion: freeing or recycling memory as soon as the launch returns corrupts the kernel’s input, and the failure appears at an unpredictable later point. The third is the interaction between concurrent streams and a shared buffer: two streams writing overlapping regions have no relative order at all, and unless the writes are disjoint by construction the result is undefined regardless of what the events say.

A fourth belongs to the boundary with the runtime itself: an out-of-order OpenCL queue with no wait list gives no ordering whatsoever between commands, which is occasionally useful and almost never what was intended. The performance-portability layers discussed in Portable Single-Source Layers and Where They Leak generally present a synchronous-looking interface precisely because most programs are written against one, and they reconstruct the queue semantics in their runtime.

Multiple devices in one queue graph

A context or a process may own several devices, and the queue graph extends across them without any additional ordering guarantee. Work on device A’s queue and work on device B’s queue are unordered relative to each other, and a buffer that exists in both address spaces does not exist in both memories automatically: on a system without peer access, the only way to move data from A to B is through the host, which is two transfers and two synchronisation points on two queues.

The portable formulation is per-device queues plus host-mediated transfers, with the host thread acting as the join. Where peer-to-peer access is available the transfers can be issued directly between the devices’ queues, but the destination device’s queue still needs an event wait before it reads what the source wrote.

Cross-device ordering bugs are the easiest to miss because each device’s own stream is internally consistent. A kernel on B that reads stale data will read whatever was in that buffer from a previous iteration, which is usually a plausible value rather than an obviously wrong one.

Reading a queue configuration

A queue specification is a claim about what the program may assume, and the useful discipline is to read it as such: the number of queues, their ordering mode, and the event edges that cross them. Two queues with event joins that always run back to back are one queue with extra bookkeeping. A single queue with a finish after every launch is a synchronous coprocessor. The configuration that actually overlaps work has at least two queues, independent work placed on them, and joins only where a real dependency exists. It also needs the transfers to be asynchronous too, which is a different mechanism with its own constraints, treated in Host-Device Transfers and the Cost of Crossing the Bus.