Inter-Operator Parallelism and Stream Scheduling

The topological order produced by the first scheduling pass yields one linear sequence. That sequence is a valid schedule and it is the one most runtimes implement, because a single queue of commands is the simplest thing to submit and the easiest to reason about. It is also, for many graphs, needlessly serial. Two branches of a residual block have no dependency between them; neither has a dependency on the branch it is not part of. Running them one after the other leaves whatever resources the first branch did not saturate idle during its execution, and on a graph with many independent small operators the idle fraction can be most of the timeline.

Exploiting that requires two things the graph already provides and one thing it does not. The graph provides the dependency relation: an edge means the consumer must wait, and the absence of an edge means it need not. It provides the independent groups, computable as the connected components of the graph once dependencies are respected, or more generally as any partition of the nodes into levels of the longest-path layering. What it does not provide is a notion of resource conflict: two nodes may be independent in the dataflow sense and still compete for the same bandwidth, the same multiprocessors, or the same cache.

Queues, events, and what an edge costs

A queue is a sequence of commands executed in order. Two queues have no ordering between them, so work submitted to different queues may execute concurrently, interleave, or run in either order. To express a dependency across queues, the runtime needs an event: a marker recorded on one queue that another queue can wait on. The pattern is: record an event after the producer on queue A, make queue B wait for that event, then submit the consumer on queue B.

The cost of that pattern is the reason graphs are usually submitted to one queue. Each event record and each wait is a host-side operation, and each cross-queue wait forces the device front end to track a dependency that a single queue would have expressed for free by ordering. For a graph with $E$ edges crossing between queues, the runtime issues $O(E)$ events. That is more host work than the launches being reordered, unless the graph is coarse enough that each queue carries a long chain.

This sets the granularity threshold. Splitting a graph into two queues pays off when the branches are long enough that the overlap gained exceeds the event overhead, and when the branches are large enough that they can actually run at the same time. Two kernels of five microseconds each will not overlap usefully on a device whose launch granularity is comparable to their duration; two kernels of a millisecond each will.

def assign_queues(order, nodes, consumers, producers):
    """Greedy level assignment: a node joins a queue if nothing on that queue
    produces any of its inputs, else it opens the next queue."""
    queue_of = {}
    n_queues = 0
    for i in order:
        placed = False
        for q in range(n_queues):
            if all(queue_of.get(p) != q for p in producers[i]):
                queue_of[i] = q              # no dependency on this queue
                placed = True
                break
        if not placed:
            queue_of[i] = n_queues
            n_queues += 1
    return queue_of, n_queues

Levels, not streams, are the compiler’s abstraction

A graph compiler should not think in queues. Queues are a runtime concept whose number is a property of the device and the driver, not of the model. The compiler’s abstraction is the level: the longest-path depth of a node from the graph inputs, computed by a single pass over the topological order. Nodes at the same level have no dependencies among them, so any of them may run concurrently with any other at that level.

Level assignment is cheap and it exposes the parallelism the graph contains. It also exposes how little there often is. A chain of operators, which is what a transformer block is once attention is treated as a unit, has one node per level and no exploitable parallelism at all. A graph with a wide fan-out, such as a mixture-of-experts layer dispatching to many experts, has levels with many nodes each and substantial parallelism. The number that matters is the ratio of total work to critical-path work, which bounds the speedup any concurrent schedule can achieve: if the critical path is half the total work, no assignment of queues can produce more than a factor of two.

The compiler emits the level assignment as a scheduling hint and lets the runtime map levels to queues. A runtime with two queues can only exploit two-way parallelism, and a graph with eight-way levels gains nothing from the extra six. That mapping is a deployment decision, which is why the level structure is a better output for the compiler than a fixed queue assignment.

Overlap helps only when resources differ

Two kernels running concurrently on one device share the device. If both are bound by memory bandwidth, running them together does not double the throughput; it divides the available bandwidth between them, and the total time is the sum of the bytes moved divided by the bandwidth. The only gain is the removal of the gap between the two kernels, which is real but small.

The gain is large when the two kernels need different resources. A kernel that saturates the matrix units while issuing few memory requests can run alongside a kernel that moves data and performs little arithmetic, and the two together use more of the device than either alone. This complementarity is the criterion a scheduler should use, and it is why the useful pairings in practice are specific: a compute-heavy convolution with a bandwidth-heavy normalisation, or a device-side kernel with a host-to-device copy of the next layer’s weights.

The complementarity argument extends to communication. In a model sharded across devices, the all-reduce that follows a layer can overlap with the computation of the next layer if the dependency structure permits: the next layer’s input depends on the reduction, but the parts of the next layer that operate on already-reduced shards do not. Overlapping communication with computation is the standard technique in distributed inference and training alike, and its limit is the same: the compute must be independent of the data being transferred, and the transfer must be large enough that its latency is amortised.

/* Two independent branches on two streams, joined by one event. */
cudaEvent_t mid;
cudaEventCreateWithFlags(&mid, cudaEventDisableTiming);
branch_a<<<grid_a, block_a, 0, s1>>>(a_in, a_out);
branch_b<<<grid_b, block_b, 0, s2>>>(b_in, b_out);
cudaEventRecord(mid, s1);                 /* marks A's completion */
cudaStreamWaitEvent(s2, mid, 0);          /* later s2 work waits for A */
join<<<grid_j, block_j, 0, s2>>>(a_out, b_out, out);

Concurrency changes the memory plan

The memory plan from the liveness pass assumes a linear schedule. Two tensors with disjoint intervals may share an offset, because the schedule guarantees the second is written after the first is read. Run the two producers concurrently and the guarantee is gone: the second write can happen while the first tensor is still being read, and the shared offset becomes a data race.

The fix is to define liveness over the partial order rather than the linear order. Two tensors conflict if there exists any node that reads one and any node that writes the other with no ordering constraint between them. Under a level-based schedule, that condition is easier to check: tensors live at the same level, or at levels connected by an unordered path, must not share storage. The result is a larger arena, because the planner can no longer reuse a buffer across two branches that run at the same time. Concurrency and memory are in direct opposition here, and the trade is one buffer per concurrent branch against the time saved by running them together.

A runtime that assigns queues dynamically, without the planner’s knowledge, can therefore break a plan that was correct for the schedule the compiler assumed. This is the reason compiled runtimes that use multiple queues either plan for the concurrent case or forbid concurrent execution of nodes that share an arena.

The limits that make single-queue the default

Three limits bound the benefit. The first is that concurrent kernel execution requires resources: two blocks from different kernels can be resident on one multiprocessor only if their combined register and shared-memory use fits. A kernel already using most of the register file leaves no room for a second, so the concurrency the scheduler intends does not materialise, and the second kernel waits without the runtime being able to see why.

The second is that the overlap window is narrow, which is also what limits the gain from replaying a recorded launch sequence. Kernels are enqueued ahead of execution, so the device can overlap the tail of one with the head of the next even on a single queue, provided the driver does not serialise them. The extra gain from explicit queues is the overlap of the bodies, not the tails, and that is only available when the bodies are long.

The third is that debugging and determinism are easier with one queue. A concurrent schedule that produces a wrong result gives no indication of which interleaving was at fault, and a schedule whose order varies between runs makes a numerical difference between two runs impossible to attribute. Systems that must reproduce a result exactly therefore serialise, and accept the loss.

The reasonable default is one queue and a schedule chosen for memory, with concurrency applied only where the graph’s level structure shows enough independent work to pay for the events, and only where the branches are known to want different resources. For most inference graphs the answer is that they are chains, the parallelism is small, and the single queue is right.