The CUDA Handbook ▸ Chapter 6. Streams, Events, and Graphs

6.9 CUDA Graphs

The streams described so far overlap CPU and GPU work, but the application still issues every operation itself, one API call at a time. Each cudaLaunchKernel() or cudaMemcpyAsync() costs CPU time – the driver validates the arguments and writes commands to the GPU – and for a workload made of many small operations, that per-call cost may rival the GPU work it dispatches. CUDA graphs, added in CUDA 10.0, separate the description of a sequence of operations from its execution: the sequence is recorded once, instantiated (“compiled”) once, and then launched repeatedly, each launch costing less CPU time than reissuing the operations one at a time.

The size of the API reference for CUDA Graphs is, frankly daunting; the reason is that graphs must be able to encapsulate every other kind of work that can be submitted to a stream. A node can be any of the following:

• kernel launch,
• memory copy,
• memset,
• host-side callback,
• memory allocation or free,
• event record or wait,
• external-semaphore signal or wait, or
• nested child graph.

Edges between nodes express the dependencies among them. For each family of CUDA operations that may be dispatched from a graph, a corresponding set of graph APIs exist to create nodes, add them to the graph and, when possible, optionally parameterize the node’s operation after graph instantiation. The richness of the API set enables entire heterogeneous operation sequences to be captured and replayed as a unit.

6.9.1 Building, Instantiating, and Launching

A graph can be populated two ways. Stream capture records the operations issued to a stream between cudaStreamBeginCapture() and cudaStreamEndCapture() without executing them, turning working streaming code into a graph with little change. Explicit construction builds the graph node by node – cudaGraphAddKernelNode(), cudaGraphAddMemcpyNode(), and one such function per node type – and wires the dependencies by hand. Explicit construction is precise, but tedious for a large graph, which has motivated the development of higher-level systems that derive the graph automatically from the data dependencies among tasks.⁷

Once created, the populated graph is a template. Before it can run, the graph must be instantiated with cudaGraphInstantiate(), which resolves the dependencies, validates the nodes, and produces an executable graph (cudaGraphExec_t). Instantiation is the expensive step – in effect a compile – and it is done once. The application then launches the executable graph with cudaGraphLaunch() as many times as needed:

cudaGraph_t     graph;
cudaGraphExec_t exec;

// Capture a sequence of operations into a graph.
cudaStreamBeginCapture( stream, cudaStreamCaptureModeGlobal );
    // ... kernel launches, memcpys, memsets, etc., all issued to `stream` ...
cudaStreamEndCapture( stream, &graph );

// Compile it once.
cudaGraphInstantiate( &exec, graph, 0 );

// Launch it many times.
for ( int i = 0; i < nIterations; i++ ) {
    cudaGraphLaunch( exec, stream );
}

When the same work repeats with only its parameters changing – new pointers or scalar arguments each iteration – the graph need not be rebuilt. cudaGraphExecUpdate() applies the changed node parameters to an existing executable graph more efficiently than a fresh instantiation, and a single node can be updated directly with a call such as cudaGraphExecKernelNodeSetParams().

6.9.2 When Graphs Help

The benefit is measured on the CPU: a graph performs the same operations the stream would have, but can be dispatched with less CPU overhead. nullKernelAsyncGraph.cu measures the overhead by capturing a run of empty kernels into a graph and launching the graph in their place. On a GeForce RTX 3060, it reports about 0.54 μs per launch, against the 1.2 μs of the asynchronous stream launch measured in Section 6.1 – roughly half the CPU cost per operation, and one cudaGraphLaunch() in place of the whole run.

The gap widens for kernels with heavier launch configurations and for pipelines that mix in memory copies, whose per-call overhead a graph absorbs the same way; it narrows – or reverses – for a few large kernels, where the GPU work already dwarfs the dispatch cost, and for sequences that change every iteration, where the cost of capture and instantiation cannot be amortized.

Graphs also extend onto the device, which turns them from a launch-overhead optimization into a way to keep whole control-flow structures off the CPU. A graph instantiated with the device-launch flag can be launched from inside a kernel by calling cudaGraphLaunch() on the device, so one grid can fire off an entire prebuilt graph with no round trip to the host.

Conditional nodes let that graph branch and loop on values the GPU computes. A conditional node owns a handle created with cudaGraphConditionalHandleCreate(); a kernel sets the handle’s value with cudaGraphSetConditional(), and an IF node then runs its body only when the value is nonzero, while a WHILE node re-runs its body for as long as the value stays nonzero, re-evaluating between iterations. A loop whose trip count the GPU decides – an iterative solver testing for convergence, a search that stops on the first match – therefore runs entirely on the device, where before it would have returned to the CPU after each step to decide whether to launch again. Conditional nodes overlap in purpose with dynamic parallelism (Section 7.4) and the cooperative grid-wide barrier (Section 7.6); graphs fit the case where the shape of the control flow is fixed ahead of time and only its data-dependent branches and trip counts are left to the GPU.