Kernel Launch Capture and Replay

A kernel launch is cheap per unit of work only when the work is large. The host must validate arguments, place a command in a queue, and return; the driver must eventually translate that command into something the device front end can read, and the device must fetch it. The host-side portion costs on the order of a few microseconds. That is invisible next to a convolution that runs for a millisecond and ruinous next to a fused elementwise kernel that runs for three microseconds. A graph of four hundred small kernels, each launched individually, spends more host time than device time, and the device idles between kernels because the queue starves.

The remedy is to record the sequence of launches once, as a data structure, and replay the whole structure with one submission. This is what CUDA Graphs provide, and what Vulkan command buffers provide in a different form: the commands are captured into a buffer at build time, and execution walks the buffer without host involvement per command.

Two ways to build the record

The first is explicit construction. The application creates a graph object, adds nodes to it programmatically, declaring each node’s dependencies and parameters, and instantiates the result. This is verbose and precise: every node and edge is stated, so the topology is exactly what the programmer intended, and no capture machinery is involved.

The second is stream capture. The application puts a stream into capture mode, then executes its ordinary launch sequence. Every launch, memory copy, and event recorded on that stream is added to a graph rather than executed. Ending capture yields the graph. This is attractive because it requires no restructuring of the launch code: the same functions that issue launches normally issue them into a capture.

Capture has rules that are easy to violate. Any call that would block on the device, such as a synchronous copy or an event synchronise, is illegal during capture, because the point of capture is that nothing executes. Any call that queries the device for a result it does not yet have is likewise illegal. A launch into a stream that is not being captured, from within a capturing region, either joins the capture as a dependency or breaks it, depending on how the streams are related, and getting that wrong produces a graph whose dependencies do not match the program’s intent. Capture also fails when the topology is data-dependent, because the recorded sequence is the one path that happened to execute.

/* Capture: the launches below are recorded into a graph, not executed. */
cudaStreamBeginCapture(stream, cudaStreamCaptureModeThreadLocal);
for (int k = 0; k < n_kernels; k++)
    issue(stream, kernels[k]);        /* same calls as the eager path */
cudaGraph_t graph;
cudaStreamEndCapture(stream, &graph);

Instantiation is the expensive step

Captured graph and executable graph are different objects. Instantiation compiles the topology into a form the device can execute: it resolves node parameters, plans the scheduling of the nodes across the device, and allocates whatever internal state the driver needs. That step is expensive relative to a launch, and it is why an application that captures and instantiates per request is slower than one that launches eagerly. The benefit exists only when one executable graph is reused across many submissions.

Reuse has a precondition: the recorded parameters must still be correct. A captured kernel node stores its arguments, including pointers to buffers. If the buffer addresses change between invocations, the recorded node points at the old memory. The mechanisms for handling this are node parameter updates, which rewrite a specific node’s arguments in an already-instantiated graph, and a whole-graph update operation that takes a newly captured graph and applies its differences to an existing executable. The update path is cheaper than a fresh instantiation but not free, and it fails outright when the topology changed rather than the parameters.

This is the crux of integrating replay with a compiled inference graph. The memory plan must be allocated once at a fixed address, and every replay must use the same arena. That is a strong constraint, and it is satisfiable precisely because the plan is static: the arena is reserved at load time, the offsets are fixed, and the recorded pointers remain valid for the life of the executable graph. A runtime that allocates per request cannot use replay without updating every node’s parameters on every call, which removes most of the benefit.

What replay changes about the schedule

A recorded graph is a dependency structure, not a linear sequence. The driver is free to schedule its nodes according to their declared dependencies, which means the recorded order constrains execution only through the edges. Two nodes with no path between them may run concurrently or in either order. This is a genuine advantage: the topology is available to the scheduler without the host having to express it, and the device can overlap work that the host would have serialised.

It is also a hazard. If two nodes write the same buffer, the graph has no edge to order them unless one was recorded. A capture that relies on stream ordering to serialise two independent nodes loses that ordering when the graph is executed, because stream order is a property of the stream and not of the graph. The captured graph preserves dependencies that were implied by the stream at capture time, which is the saving grace, but only for operations that were on the same captured stream. Operations recorded on different streams are ordered only by explicit events.

The second change is that the launch cost is amortised but not eliminated. Submitting an executable graph is one host operation whose cost grows with the number of nodes, though far more slowly than the same nodes launched individually. A graph with a handful of nodes gains little; a graph with hundreds gains a great deal. The crossover is the only number that matters for deciding whether to use the mechanism, and it is a property of the driver and the device rather than of the graph compiler.

Device-side and conditional nodes

The straightforward replay model assumes the host submits the graph and the device runs it to completion. Two extensions break that assumption in useful ways.

Conditional nodes let the graph contain a branch whose condition is evaluated on the device. The node holds two sub-graphs and a predicate, and the device selects one without returning to the host. This recovers a limited form of data-dependent control flow inside a replayable structure, which matters for models with a mixture-of-experts routing decision or an early-exit path. The limitation is that both branches are instantiated, so the memory for both must be reserved even though only one executes.

Device-side graph launch lets a running kernel launch a graph. The use is a loop: a kernel decides whether to continue and launches the next iteration’s graph directly, avoiding a host round trip per iteration. This is how a replay structure handles a fixed-length decode loop without the host issuing each step. The constraint is that the launching kernel must be resident when it launches, which requires the device to support nested execution and reserves resources for it.

/* Without device-side launch the host still submits once per iteration. */
for (int i = 0; i < N; i++) {
    cudaGraphLaunch(exec, stream);   /* enqueue, do not synchronise */
}
cudaStreamSynchronize(stream);       /* one wait for the whole sequence */

Where replay does not apply

Replay is incompatible with anything that must observe an intermediate result on the host. A model whose output length depends on a generated token needs the token’s value to decide whether to stop, and that value lives on the device until it is copied back. The copy is a synchronisation, and a synchronisation inside a captured region is illegal. The resolution is to capture the body of one step and replay it per step from the host, which recovers the per-launch saving within a step but not across steps.

Replay is also incompatible with dynamic allocation inside the captured region, since an allocation is a host-side operation and capture does not execute anything. It is incompatible with graphs whose topology depends on input values, for the same reason. And it is incompatible with debugging tools that need to intercept individual launches, which is why a runtime usually keeps an eager path available alongside the replay path.

The final limitation is that the mechanism sits below the graph compiler and knows nothing about it. Capture records whatever the runtime issues, including redundant launches that a better fusion pass would have removed. Capturing a badly compiled graph produces a fast replay of too much work. The two optimisations are complementary and neither substitutes for the other: the compiler reduces the node count, and replay reduces the cost per node.