100 DAYS OF GPU Course preview
Day 56Module 6
in-technical-review

CUDA graphs

CUDA graphs let you record a sequence once and replay it with one call instead of launching each kernel separately. A captured pipeline may still show no speedup.

Day 48 explains why. Its four-kernel chain at 16 Mi elements spends milliseconds on memory traffic and microseconds on four launches. A graph can remove only part of the launch cost.

This page captures that chain and launches it both ways at five sizes. The results show when cudaGraphLaunch saves a large share of run time and when the change is too small to matter.

What a capture records, and what it does not

A CUDA graph is your work as data: a set of nodes (kernels here, but copies and memsets qualify) plus the dependencies between them. You build it once, cudaGraphInstantiate turns it into an executable, and each cudaGraphLaunch submits the whole thing to a stream as one call.

Stream capture records the code you already have. Between cudaStreamBeginCapture and cudaStreamEndCapture, "all operations pushed into the stream will not be executed, but will instead be captured into a graph" (https://docs.nvidia.com/cuda/archive/12.6.0/cuda-runtime-api/group__CUDART__STREAM.html , checked 2026-09-01). Nothing runs during capture.

The four launches become four nodes with three edges. The program checks that count with cudaGraphGetNodes before using the graph.

What the capture freezes matters as much as what it records. Grid dimensions, block shapes, kernel arguments, the pointers themselves: all of it is stored at capture time. Replaying the graph replays the chain at the size it was captured at, so this program builds one graph per problem size, and changing a node without rebuilding is day 57's subject.

Graphs do not reduce this chain's memory traffic. The captured chain still writes three intermediates to global memory and reads them back, 36 bytes per element against the fused kernel's 12. Fusion removes that traffic, while a graph reduces host submissions.

The two changes solve different costs and can be used together.

Diagram: one iteration of the chain, three timelines. Three horizontal bands, each with a host API row above a GPU kernel row on one time axis. Band 1, individual launches at a small size: four cudaLaunchKernel bars on the API row, four short kernel bars below, with visible gaps where the device waits for the next submission. Caption "4 submissions per iteration." Band 2, the graph at the same size: one cudaGraphLaunch bar on the API row, and still four kernel bars below, packed tighter. Caption "1 submission, the same 4 kernels." Band 3, either path at 16 Mi elements: the four kernel bars are milliseconds wide and the API bars are slivers at the left edge. Caption "the kernels are about 1,000 times wider than the submissions." Alt text: "Capturing four kernel launches as one CUDA graph turns four host submissions into one. The four kernel bars remain, and at sixteen million elements they dwarf either submission cost."

Fewer launches help only when launches limit the run

Four host calls becoming one does not cut GPU work by three quarters. Launches are asynchronous: the host queues work and runs ahead, as day 9 showed. While a kernel runs for milliseconds, the host can queue the next three, so their submission cost overlaps with GPU work.

Submission only lands on the critical path when the device drains its queue faster than the host refills it, which is the small-kernel regime: an inference loop at batch size one, a solver iterating on a few thousand cells, day 48's chain at the bottom of its sweep. There the gap between kernel bars is launch overhead, the timeline from day 41 shows it as white space on the kernel row, and a graph exists to close it.

The limit is simple: a graph can save at most the submission cost, which is usually measured in microseconds. Whether that is a rounding error or your whole frame budget depends on how many launches you make and how long they run, not on the API.

One function, timed as launches and replayed as a graph

Full program in code/day56-graphs/graphs.cu. Four things pin the comparison down.

The kernels are day 48's, unchanged. The four stages, their constants and the input generator are copied from code/day48-fusion/fusion.cu with a marker comment at each block, so the chain this page prices is byte for byte the chain day 48 priced.

One launchChain function is both paths. Part 2 times it as four launches, and the capture records it. There is no second copy of the chain that could drift into launching something else.

The graph must replay the chain exactly. Before timing anything, the program runs the chain both ways at 1,049,187 elements, an odd size no launch shape covers exactly, checks both against a double reference, then requires the two outputs to match bit for bit and the node count to be four. Every gate is a real branch returning EXIT_FAILURE.

Events on the stream that owns the work. The timing helper is the course's timeKernel with its event records moved onto the stream, because nothing here runs on the default stream. Each timed run is a batch of 20 whole-chain iterations, so a per-iteration difference of microseconds sits well above the half-microsecond event resolution, and the warm-up inside the helper covers every kernel and the graph itself, which matters under lazy module loading.

Here is the capture. It is the program's own launch code, recorded:

static cudaGraphExec_t captureChain(cudaStream_t stream, const float* d_x,
                                    const float* d_r, float* d_t1, float* d_t2,
                                    float* d_t3, float* d_y, size_t n,
                                    size_t* nodeCount) {
    cudaGraph_t graph = nullptr;
    CUDA_CHECK(cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal));
    launchChain(stream, d_x, d_r, d_t1, d_t2, d_t3, d_y, n);
    CUDA_CHECK(cudaStreamEndCapture(stream, &graph));
    CUDA_CHECK(cudaGraphGetNodes(graph, nullptr, nodeCount));

    cudaGraphExec_t exec = nullptr;
    CUDA_CHECK(cudaGraphInstantiate(&exec, graph, 0));
    // The executable graph is self-contained, so the recording can go.
    CUDA_CHECK(cudaGraphDestroy(graph));
    return exec;
}

And the whole comparison is two loops that differ in one line:

for (int it = 0; it < kItersPerBatch; ++it) {
    launchChain(stream, d_x, d_r, d_t1, d_t2, d_t3, d_y, n);  // 4 calls
}
for (int it = 0; it < kItersPerBatch; ++it) {
    CUDA_CHECK(cudaGraphLaunch(exec, stream));                // 1 call
}

Capture and instantiation sit outside the timed region because the program pays them once before replay. Read their cost from the timeline: the cudaGraphInstantiate bar in the shipped Nsight Systems capture, wrapped in its own NVTX range. A chain you will run once should stay four launches; the exercise below makes you find the break-even.

Results

Re-verified 2026-09-02 on the same Tesla T4 with driver 580.173.02 and CUDA 13.0 (V13.0.88). Correctness and the size-dependent shape reproduced. At 1,024 elements the four-launch path improved from 19.753 to 13.015 us, so the graph saving narrowed from 10.014 to 3.918 us.

At 16,777,216 elements both paths still round to the same 1.00 ratio. Both transcripts and the fresh CUDA 13 report are listed in front matter. The predictions predate both measurements; their verdicts follow the table.

One: the win per iteration is microseconds, never milliseconds. At 1,024 elements, where four kernels finish in less time than the host takes to submit them, the graph saves somewhere between 2 and 15 microseconds per iteration. If the saved column shows a millisecond anywhere, either the measurement or this page's whole model is broken.

Two: the saving becomes negligible at the top of the sweep. At 16,777,216 elements the chain moves about 604 MB per iteration and the two columns stay within 2 percent of each other, because both are the same four memory-bound kernels and submission is hidden behind milliseconds of traffic. Anyone promising a graph speedup here is selling the wrong fix: this chain needs fusion, which cuts its bytes by three, not a graph, which cuts its submissions by four.

Three: the timeline shows one API bar where there were four, and still four kernel bars. In the shipped capture, the graph NVTX range holds cudaGraphLaunch calls on the API row while the kernel row underneath still shows all four kernels per iteration. A graph changes how work is submitted, not what runs.

Four: instantiation costs more than one iteration saves. The cudaGraphInstantiate bar inside the capture range is wider than the per-iteration saving at any size, so a graph launched once is a loss and the break-even is several replays in.

n 4 launches (us/iter) 1 graph (us/iter) saved (us) ratio
1,024 13.015 9.097 3.918 1.43
16,384 13.934 9.946 3.988 1.40
262,144 43.238 39.129 4.109 1.11
4,194,304 609.423 598.689 10.734 1.02
16,777,216 2,410.840 2,406.538 4.303 1.00

One held: the largest saving in the table is 10.734 microseconds and the headline one at n = 1,024 is 3.918, inside the predicted 2 to 15. Two held: the largest-size row's columns differ by 0.2 percent and round to a ratio of 1.00.

Three held in the published CUDA 12.6 capture: 1,301 cudaGraphLaunch calls on the API row where the four-launch passes needed 5,228 cudaLaunchKernel, with all four kernel bars still present per iteration underneath. Four held: cudaGraphInstantiate averages 80.53 microseconds in that cuda_api_sum, eight times that run's 1,024-element saving, so a graph launched once is a loss and break-even sits several replays in.

This page ships its timeline at content/profiles/cuda-graphs/: day56-graphs.nsys-rep plus the nvtx_sum and cuda_api_sum CSV exports, so the exercise is doable without running nsys yourself. Nsight Systems needs no root and no performance counters, so on your own card the capture commands in the README reproduce it as an ordinary user.

A graph can save at most the submission cost. Check that cost on the timeline before treating a graph as an optimization.

Run it yourself

Use a CUDA GPU supported by the program:

nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo -o graphs graphs.cu

The program holds six 64 MiB buffers and runs 20-iteration batches across five sizes and two submission paths. If you cannot run it, read /setup/learn-cuda-without-a-gpu and work from the shipped reports.

Exercise

Change captureChain so one capture records two whole iterations of the chain, eight nodes instead of four, and halve the replay count so the total work is unchanged. Predict what happens to the saved column at the smallest size before you run, then measure and explain the shape.

Time: 25 to 40 minutes. Submit: your measured table at n = 1,024 and 16,777,216, your prediction, and two sentences on which one survived.

Check: a shape rather than an absolute number, because absolute microseconds vary by card. The program's node-count gate must now expect 8, and the bit-identical gate must still pass; if it fails, your second iteration read a buffer the first one no longer feeds. On pass, compare saved microseconds per iteration against the unmodified run: the remaining submission cost is per cudaGraphLaunch, and you halved how often you pay it.

Hint 1

What exactly is left in the graph column's cost? Not the kernels, they run either way. Count the calls the host still makes per iteration of real work, before and after your change.

Hint 2

At the small size the four-launch path pays 4 submissions per iteration and the one-graph path pays 1. Your change makes it half a submission per iteration. Does the saved column grow by the same step it grew from 4 to 1, or a smaller one?

Solution

The saved column grows, but by less than the first step did. Going from four submissions to one deleted three per iteration; going from one to a half deletes half of one. Batching k iterations into a graph leaves 1/k of a submission per iteration, so the returns diminish geometrically while the graph's memory footprint and staleness grow linearly, and somewhere past a handful of iterations the remaining cost stops being submission at all.

A graph amortises a fixed cost, so its value is the cost times the count. Price the submission first, then count how many you can batch into one replay before something else becomes the widest bar.

Pitfalls

You captured your pipeline and nothing got faster. Almost always the kernels are big enough that submission was already hidden. Open the timeline before reaching for the API: if the kernel row has no white space between bars, there is no overhead to reclaim, and your problem is bytes, which is day 48's subject.

Capture fails immediately on the default stream. The runtime is explicit: "Capture may not be initiated if stream is cudaStreamLegacy" (https://docs.nvidia.com/cuda/archive/12.6.0/cuda-runtime-api/group__CUDART__STREAM.html , checked 2026-09-01). Create a stream and capture on that, which this program does anyway because everything since day 51 runs on one.

cudaStreamEndCapture returned a null graph. Something between begin and end broke the capture rules, a synchronous cudaMemcpy or a cudaDeviceSynchronize being the usual suspects: "If capture was invalidated, due to a violation of the rules of stream capture, then a NULL graph will be returned" (same page, checked 2026-09-01). Work during capture is recorded, not run, so a call that must wait for results cannot be inside one.

You re-capture and re-instantiate every iteration. Then the one-off cost is not one-off, and prediction four above says each instantiation costs more than a replay saves, so the graph path loses to plain launches. Build once, replay many; when a parameter changes each iteration, update the instantiated graph instead, which is day 57.

The graph replays stale sizes or stale pointers. Capture stores every argument, so replaying after n changed runs the old grid on the old byte counts and the correctness gate is what catches it. One graph per configuration, or day 57's update path.

Go deeper

Next

Day 57 keeps the graph and moves the parts: cudaGraphExecKernelNodeSetParams to change a node without re-instantiating, and a conditional node to put the convergence loop from day 41's PageRank inside the graph, where checking the residual no longer means a host round trip. Day 58 then attacks the overhead a graph cannot reach, the gap between dependent kernels, with programmatic dependent launch on Hopper.