Streams, Events, and Concurrent Kernel Execution

A kernel launch is asynchronous. The host-side call enqueues work and returns; the device executes it later. Everything about ordering is therefore expressed in the queue rather than in the program counter, and the entity that owns a queue is a stream. Two operations in the same stream are ordered, so the second begins only after the first has completed, while operations in different streams have no ordering between them at all unless something explicit creates one. That single rule explains almost every behaviour an application sees when it moves from one stream to several.

The ordering contract inside a stream

Streams give a total order on the operations placed in them, and the order includes copies as well as kernels. A host-to-device copy, a kernel, and a device-to-host copy placed in one stream in that order will execute in that order, and the host-to-device copy’s data will be visible to the kernel. This is what makes the simple program correct without any explicit synchronisation: the default stream, stream zero, provides a serial order that matches the sequence of calls.

The default stream is special. Operations in it are ordered with respect to the blocking streams as well, so a kernel in the default stream waits for everything already enqueued in other blocking streams, and subsequent work in those streams waits for the default-stream kernel. This makes it safe to mix legacy code that never mentions streams with new code that does, and it also serialises them. The per-thread default stream, which makes stream zero private to each host thread, removes that serialisation and is the reason some code changes behaviour when it is enabled: two host threads that previously raced through a shared default stream now run concurrently.

Completion of a stream is observed by synchronising on it — a stream synchronisation call that blocks the host until every operation in the stream has finished — or by querying it. Synchronising is the coarse tool; it waits for everything, including work that the host does not yet need.

Why an enqueued kernel might not overlap

Two kernels in different streams do not overlap by inheritance; they overlap only if the hardware can start the second while the first is still occupying resources. That condition fails in several ordinary situations.

The first is resource exhaustion. A kernel whose blocks use all of a multiprocessor’s registers or shared memory leaves no room for a block of another kernel on the same multiprocessor. If the first kernel’s grid already fills the device’s block slots, the second kernel’s blocks sit in the queue until slots free. The hardware schedules at block granularity, so overlap requires that the first kernel leave a gap that a block of the second kernel fits into.

The second is the concurrency limit on the number of simultaneous kernels. Devices have a bounded number of kernel dispatches they can track at once; beyond it, launches are queued. The limit is a device property, not a small constant, but it is finite and a program that launches hundreds of tiny kernels across streams will hit it.

The third is dependencies that are not visible in the stream structure. A memory copy in stream A and a kernel in stream B that reads the same buffer have no ordering, which is a race rather than a missed overlap. A copy in stream A from a buffer that stream B’s kernel is still writing is equally a race, and it is the more common bug because the copy is issued later in program order and looks safe.

The fourth is implicit synchronisation from the allocator or from a legacy API. A synchronous copy blocks the host and, on the default stream, serialises everything. A pageable-memory copy is staged through a host-side buffer and cannot overlap with a kernel the way an asynchronous copy from pinned memory can. An allocation that must grow the memory pool may synchronise the device. Any of these can silently turn a four-stream pipeline into a single serial one.

Events as the nodes of a graph

An event is a marker placed in a stream; it records a point in that stream’s order. Two uses follow. Waiting on an event makes a second stream’s subsequent operations wait for the first stream to reach the marker, which is how a fork-join pattern is expressed without synchronising the host. Elapsed time between two events, measured with the timing API, gives a device-side interval, which is the only honest way to time a kernel from outside it.

The pattern that matters is the dependency that is finer than a whole stream. A kernel in stream B needs the output of a copy in stream A but nothing else from stream A, so recording an event after the copy and having B wait on it lets the rest of A proceed independently. This is the same idea as a dependency graph with edges derived from data flow, expressed with stream wait calls; the hardware uses the same dependency information to decide what is runnable.

cudaStream_t a, b;
cudaEvent_t copied, done;

cudaStreamCreate(&a);
cudaStreamCreate(&b);
cudaEventCreateWithFlags(&copied, cudaEventDisableTiming);
cudaEventCreate(&done);

// Upload on stream a.
cudaMemcpyAsync(d_in, h_in, bytes, cudaMemcpyHostToDevice, a);
cudaEventRecord(copied, a);

// Stream b's kernel waits only for the upload, not for the rest of a.
cudaStreamWaitEvent(b, copied, 0);
compute<<<grid, block, 0, b>>>(d_in, d_out);
cudaEventRecord(done, b);

// The host waits only at the point where it needs the result.
cudaEventSynchronize(done);
cudaMemcpyAsync(h_out, d_out, bytes, cudaMemcpyDeviceToHost, a);

The event is created with timing disabled because nothing measures it; that flag removes the timestamp machinery from the marker and makes it cheaper. The final copy goes back on stream a, which is safe here only because the host synchronised on done before submitting it. Move that submission before the wait, or make it asynchronous without the wait, and the copy is ordered against nothing in stream b: the event chain runs from a to b, not back again.

Concurrency between copies and computation

A copy can overlap with kernel execution when it is asynchronous and the host memory is pinned. Unpinned memory is not addressable by the device’s direct memory access engine in the same way, so the runtime stages it through a pinned buffer, which imposes a host-side copy and an implicit synchronisation. Pinning moves the page-locked cost to allocation time and makes the transfer a single asynchronous operation.

The number of copy engines is a device property, and the direction matters: many parts have one engine in each direction, some have two. A host-to-device and a device-to-host copy can proceed simultaneously if there are two engines. Two device-to-host copies in different streams generally serialize on one engine, even though the kernels around them overlap. This is why pipelining a workload usually needs at least two streams and why the achieved concurrency plateaus at the engine count rather than the stream count.

Multi-device and peer access

Streams are per-device. A stream created in the context of device 0 orders work submitted to device 0, and passing that stream to a launch intended for device 1 is an error rather than a cross-device dependency. Ordering between devices is expressed with events and with peer-access copies, and only where the platform supports direct peer transfers; otherwise the transfer routes through host memory and the two copies are each a separate operation with its own synchronisation. A pipelined multi-device workload therefore looks much like a single kernel’s staged dependencies, and the same reasoning about event nodes applies.

A copy from device 0 to device 1 enqueued in device 1’s stream, with a wait on an event recorded in device 0’s stream, is the portable form. The dependency is on the event, not on the stream, because the streams belong to different devices and the queue structures cannot see each other. Work distributed across devices that way also has to be balanced explicitly, since the hardware will not move blocks between devices the way it dispatches blocks of one grid; the launch geometry has to be chosen per device, which is the same sizing decision described in grid-stride loops and launch size independence.

Choosing a stream count

The useful number of streams is set by what can overlap, not by what could in principle. A pipeline of $k$ stages that must run in sequence gains nothing from $k$ streams; it gains from two, so that stage $i$ of buffer $j$ overlaps with stage $i-1$ of buffer $j+1$. A batch of independent kernels gains from as many streams as the device can hold concurrently, subject to the dispatch limit and the copy-engine count.

Beyond that point, more streams add launch overhead and a longer host-side queue without adding parallelism, and they make the event graph harder to reason about. The pathological version is one stream per small kernel, where the host enqueues faster than the device retires and the queue becomes the bottleneck; the fix is a fixed pool of streams and a scheduled assignment of work to them, not more streams.

The interaction with occupancy is worth stating plainly: two kernels share a device the same way two blocks of one kernel do. The occupancy analysis that decides how many blocks fit on a multiprocessor — the register and shared-memory budget described in block barriers and divergent synchronization — also decides how the device’s capacity splits between concurrent kernels, and a kernel that demands every register is a kernel that will not overlap with anything. The concurrency controls in the launch configuration — block size, shared memory per block, and any per-block limit on resident blocks — are therefore the same controls that govern multi-kernel scheduling.