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

Graph updates and conditional nodes

Here is the convergence loop, as this course has written it since day 34:

do {
    cudaGraphLaunch(exec, stream);                       // day 56's fix
    cudaMemcpy(&r, d_residual, 4, cudaMemcpyDeviceToHost);
} while (r > tol);

Day 56 already grouped the launches into one graph, but the loop still has the wait shown on day 41: that four-byte copy cannot return until everything queued before it finishes. The host must wait once per iteration, deciding something the device already knows. By the end of this page the whole loop is one launch, the decision runs on the device, and you can change a parameter frozen inside the graph without rebuilding it.

A loop the device runs

Since CUDA 12.4, a graph can contain a conditional node: a node that owns a body graph and decides at run time, on the device, whether to run it. The WHILE flavour re-evaluates after every pass, so a body that holds one solver iteration becomes a loop with no host in it. CUDA 12.8 added IF-ELSE and SWITCH flavours; everything on this page needs only 12.4 (https://developer.nvidia.com/blog/dynamic-control-flow-in-cuda-graphs-with-conditional-nodes/ , checked 2026-09-01).

The API has three parts. A handle, created against the parent graph with cudaGraphConditionalHandleCreate, is the variable the condition reads. The node is added with cudaGraphAddNode and cudaGraphNodeTypeConditional; creating it returns its empty body graph in phGraph_out[0], and capturing into that body with cudaStreamBeginCaptureToGraph adds the kernels.

A kernel in the body calls the device function cudaGraphSetConditional(handle, value) to set the next evaluation: nonzero runs the body again, zero stops.

One default makes the shape. Created with defaultLaunchValue = 1 and cudaGraphCondAssignDefault, the handle resets to 1 on every launch of the parent graph, so the body always runs at least once and the node is a do-while. That matches a solver, which cannot know its residual before the first sweep, and it matches the host loop this page replaces, so the two can be required to agree.

Hardware. Conditional nodes need CUDA 12.4+ and, on Linux, driver 550.54.14 or newer (https://docs.nvidia.com/cuda/cuda-toolkit-release-notes/index.html#cuda-driver , checked 2026-09-01); any card this course supports runs them, including the sm_75 floor. On an older toolkit the fallback is this page's strategy 2, the host-checked loop with one graph launch per iteration, and the program compiles the conditional path out and says so rather than failing to build.

Diagram: three ways to stop a solver, on one timeline. Three horizontal bands, each a CUDA API row above a GPU kernel row, one shared time axis, iterations marked by ticks. Band 1, host loop: per iteration, four short launch bars, then a cudaMemcpy bar spanning the kernels queued before it, then a gap on the kernel row. Caption "4 launches, 1 round trip, 1 gap, every iteration." Band 2, graph per iteration: one cudaGraphLaunch bar, the same cudaMemcpy bar and the same gap. Caption "1 launch, still 1 round trip." Band 3, conditional WHILE: a single cudaGraphLaunch bar at the left, then an unbroken kernel row to the end, no copies. Caption "1 launch, 0 round trips, the iteration count decided on the device." Alt text: "Three timelines of one Jacobi solve. A host-checked loop pays one blocking four-byte copy per iteration and the kernel row shows a gap each time. A conditional WHILE graph launches once and the kernel row runs unbroken."

What a graph freezes, and what it does not

A graph stores some values but follows pointers to memory. This distinction depends on whether a kernel argument is passed by value or through a pointer.

A kernel argument passed by value is copied into the node's parameters at capture, so the instantiated executable keeps the old value however often you edit the variable that supplied it. Memory is the opposite: a graph records pointers, not contents, so anything a kernel reads through a pointer can change between launches with no update at all. The residual, iteration counter, and solution buffers in this program change through pointers.

A scalar that must change on each launch can also live behind a pointer, which avoids a graph update.

When a stored value must change, cudaGraphExecUpdate patches the instantiated executable in place. You hand it a second graph with identical topology and new parameters, and it checks compatibility and applies the change, skipping instantiation: "Whole graph update allows the user to supply a topologically identical cudaGraph_t object whose nodes contain updated parameters" (https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/cuda-graphs.html , section 4.2.3, checked 2026-09-01). Prediction 4 tests whether skipping instantiation saves time on this small graph.

The result is stored in a cudaGraphExecUpdateResultInfo that you must read, because a failed update returns an error and leaves the old parameters running.

Deeper. Before conditional nodes, device-driven control flow meant dynamic parallelism, kernels launching kernels. It still exists (https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/dynamic-parallelism.html , checked 2026-09-01), and in 2026 it is rarely the right answer: a convergence loop wants a conditional graph, a pipeline wants streams or a graph, and a truly data-dependent grid usually wants a second launch. This course teaches it as one paragraph on purpose.

One solver, three deciders

Full program in code/day57-graph-update/graph_update.cu. The comparison is fair for three reasons.

Only the decision mechanism moves. All three strategies run the same two Jacobi sweeps and the same residual reduction per iteration, in the same order, from the same zeroed start. The program requires the three iteration counts to match exactly and the three final states to be bit-identical, and it checks the true |A x - b| residual in double on the host, so three strategies agreeing on a wrong answer still fail.

Every gate is a branch, and the loop has a cap. The device-side decision kernel enforces a 500-iteration ceiling as well as the tolerance, because a WHILE body with a bug in its condition is a loop with no ctrl-C.

Events time it, the timeline explains it. Whole solves are timed with CUDA events, mean of ten after three warm-ups, and every phase carries an NVTX range so Nsight Systems can put widths on instantiate, recapture and exec-update.

The conditional graph is built in one function: handle, node, then a capture into the node's body graph on a created stream:

CUDA 13 added dependencyData to cudaGraphAddNode; CUDA 12.6 still uses the earlier five-argument form. The small CUDART_VERSION gate below keeps one source compatible with both signatures. The CUDA 13 run and a fresh CUDA 12.6 compatibility run both pass every gate.

    cudaGraphConditionalHandle handle;
    CUDA_CHECK(cudaGraphConditionalHandleCreate(&handle, graph, 1,
                                                cudaGraphCondAssignDefault));

    cudaGraphNodeParams params = {};
    params.type = cudaGraphNodeTypeConditional;
    params.conditional.handle = handle;
    params.conditional.type = cudaGraphCondTypeWhile;
    params.conditional.size = 1;
    cudaGraphNode_t node = nullptr;
#if CUDART_VERSION >= 13000
    CUDA_CHECK(cudaGraphAddNode(&node, graph, nullptr, nullptr, 0, &params));
#else
    CUDA_CHECK(cudaGraphAddNode(&node, graph, nullptr, 0, &params));
#endif

    // The body graph is owned by the conditional node; capturing into it is
    // how kernels get inside. It is never instantiated or destroyed here.
    cudaGraph_t body = params.conditional.phGraph_out[0];
    cudaStream_t stream;
    CUDA_CHECK(cudaStreamCreate(&stream));
    CUDA_CHECK(cudaStreamBeginCaptureToGraph(stream, body, nullptr, nullptr, 0,
                                             cudaStreamCaptureModeRelaxed));
    enqueueSweeps(d_x, d_y, d_b, d_partials, omega, stream);
    decideAndCount<<<1, kReduceBlocks, 0, stream>>>(d_partials, d_residual,
                                                    d_iterations, kTolerance,
                                                    kMaxIterations, handle);
    CUDA_CHECK(cudaStreamEndCapture(stream, nullptr));
    CUDA_CHECK(cudaStreamDestroy(stream));

The decision itself is eight lines at the end of the body's last kernel, after a reduction that is line for line the one the host-checked strategies run, which is what entitles the program to demand identical iteration counts:

    if (tid == 0) {
        residual[0] = tile[0];
        const int done = *iterations + 1;
        *iterations = done;
        const unsigned int keepGoing =
            (tile[0] > tol && done < maxIterations) ? 1u : 0u;
        cudaGraphSetConditional(handle, keepGoing);
    }

And the update pass changes the relaxation weight from 1.0 to 0.7 in the already-instantiated per-iteration graph, by capturing the same four launches again with the new value and handing the result to cudaGraphExecUpdate:

    nvtxRangePushA("recapture");
    cudaGraph_t updatedGraph = captureIterationGraph(d_x, d_y, d_b, d_partials,
                                                     d_residual, kOmegaSecond);
    nvtxRangePop();

    nvtxRangePushA("exec-update");
    cudaGraphExecUpdateResultInfo info = {};
    CUDA_CHECK(cudaGraphExecUpdate(iterationExec, updatedGraph, &info));
    nvtxRangePop();
    CUDA_CHECK(cudaGraphDestroy(updatedGraph));
    if (info.result != cudaGraphExecUpdateSuccess) {
        std::fprintf(stderr, "cudaGraphExecUpdate result %d, not success\n",
                     static_cast<int>(info.result));
        ok = false;
    }

At omega 0.7 the solve converges more slowly on purpose. A changed iteration count is the one observable that proves the update reached the kernel's frozen argument, and the program fails if the count does not move.

Results

Re-verified 2026-09-02 on the same Tesla T4 with driver 580.173.02 and CUDA 13.0 (V13.0.88); the program printed conditional nodes: available (CUDART 13000). A same-day CUDA 12.6 compatibility run printed CUDART 12060 and passed after the source gained the version gate for CUDA 13's new cudaGraphAddNode dependency-data parameter.

The original and both fresh transcripts, plus the CUDA 13 report, are listed in front matter. The published CUDA 12.6 capture and CSV exports are at content/profiles/graph-updates/: day57-graph.nsys-rep with day57-graph_nvtx_sum.csv, day57-graph_cuda_api_sum.csv, day57-graph_cuda_api_trace.csv and day57-graph_cuda_gpu_kern_sum.csv.

Strategy Iterations Total (ms) Per iteration (ms)
host-loop-sync 26 14.409 0.554
graph-per-iteration 26 14.024 0.539
conditional-graph 26 13.881 0.534

Claims the transcript will grade

  1. All three strategies report the same iteration count, between 15 and 40, and the three final states are bit-identical. The count is set by the Jacobi contraction factor, 0.8 per sweep at omega 1.0, and by nothing the decider touches. After the update to omega 0.7 the count grows by a factor between 1.2 and 2.
  2. In the cuda_api_trace export, the conditional strategy's solves make zero cudaMemcpy calls: between each cudaGraphLaunch of the conditional graph and the synchronize that follows it, the trace shows no copy at all. The host-loop solves show one cudaMemcpy per iteration, each spanning the kernels queued before it, exactly day 41's four-byte container bar.
  3. The ordering is host-loop-sync slowest, conditional-graph fastest, with the whole spread inside 15 percent, because one iteration here is roughly half a millisecond of kernel work and a round trip costs tens of microseconds. If the conditional graph wins by more than 30 percent, my cost model of the round trip is wrong, which would be the more interesting result.
  4. The exec-update NVTX range is narrower than instantiate. That is the claim in the API's name. If it is not, updating buys convenience and not time at this graph size, and the page will say so.

What the run said

  1. Held. All three strategies stopped at 26 iterations with residual 6.914e-06, inside the 15 to 40 band, and the bit-identity gate passed. After cudaGraphExecUpdate moved omega to 0.7, the count grew to 36, a factor of 1.38, inside the predicted 1.2 to 2.
  2. Held, and countable in the published CUDA 12.6 trace. In cuda_api_trace there are exactly 14 launch-to-synchronize windows containing zero cudaMemcpy calls, and they are the conditional graph's 14 whole-solve launches (three warm-ups, ten timed, one correctness). Every other launch-to-sync window in the capture, all 364 of them, carries at least one copy: the four-byte residual crossing per iteration.
  3. Held. host-loop-sync slowest at 14.409 ms, conditional-graph fastest at 13.881, a spread of 3.8 percent, inside the 15 percent bound. At half a millisecond of kernel work per iteration, the round trips this page deletes were never the bill.
  4. Held in the published CUDA 12.6 capture. The exec-update NVTX range spans 12,290 ns against instantiate's 110,150: updating the frozen graph cost a ninth of building it.

The timeline matters more than the measured milliseconds. The conditional kernel row has no gaps or API calls between launch and synchronize. This removes a fixed cost that matters more as kernels get shorter and launch overhead stays fixed.

Run it yourself

Use a CUDA 12.4 or newer toolkit for the conditional path. From the README:

nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo -o graph_update graph_update.cu
nsys profile -t cuda,nvtx -o profile/day57-graph ./graph_update

On CUDA 12.0 to 12.3 the same file builds with strategy 3 compiled out and the update pass intact. If you cannot run nsys, read /setup/learn-cuda-without-a-gpu and work from the shipped reports.

Exercise

Checking the residual costs two reduction kernels per decision, and the device loop makes it cheap to check less often. Change the decision cadence to four double-sweeps per residual check, in all three strategies, so the gates still have three like-for-like solves to compare. Report both timing tables, before and after.

Time: 30 to 45 minutes. Submit: your modified graph_update.cu, the two conditional-graph rows, and one sentence on when checking less often would be wrong.

Check: the program's own gates grade you, and they must all hold: your version prints all gates passed and exits 0, or it names the gate that tripped on stderr, with the iteration counts or the measured |A x - b|_inf next to its printed bound. On top of that, cuda_gpu_kern_sum from your own capture must show about one residualPartial launch per four relaxJacobi pairs; if it shows one per pair, you changed a constant somewhere without changing the graphs.

Hint 1

The body graph is a graph like any other, and nothing says it must hold one iteration. What exactly does the WHILE node re-run, and which lines in this program decide what one pass of each strategy contains?

Hint 2

The change is one loop of four around the sweep pair, in three places: the host loop, the captured iteration graph and the conditional body. Then look at decideAndCount: its counter now has to advance by four per pass, or the 500 ceiling stops binding where the host strategies stop.

Solution

Wrap enqueueSweeps in a for loop of four inside solveHostLoop, inside captureIterationGraph's capture and inside buildConditionalGraph's capture, and make decideAndCount add 4 to the counter, with the host loops counting the same way. All three strategies now overshoot the tolerance by up to three double-sweeps, but they overshoot identically, so the counts still match, the states are still bit-identical, and the |A x - b| gate passes with more room than before, because the extra sweeps tightened the answer.

A device-side loop lets you choose how often to check convergence. Measure the cost of the check kernels against the sweeps between them. The best interval depends on the GPU and workload.

Pitfalls

Your WHILE body never runs. The handle was created with the default value 0 and nothing upstream set it, so the condition is false on entry and the node is skipped, silently and successfully. Create the handle with defaultLaunchValue = 1 and cudaGraphCondAssignDefault for a do-while, or add a kernel before the node that calls cudaGraphSetConditional first.

The graph launch never returns. A device-decided loop with an unreachable tolerance spins forever, and there is no host loop to break out of. Add an iteration cap to the decision kernel next to the tolerance, the way decideAndCount does, before you ever need it.

You changed the variable and the graph kept the old value. Kernel arguments pass by value and are frozen into the executable at instantiation. Re-capture and call cudaGraphExecUpdate, or move the scalar behind a pointer so no update is needed at all.

cudaGraphExecUpdate returned an error and you did not look at why. The updating graph must be topologically identical to the one that was instantiated; an added kernel, a dropped one or a changed dependency fails the whole update, and cudaGraphExecUpdateResultInfo::result names the reason while the executable keeps running the old parameters. Gate on cudaGraphExecUpdateSuccess, as the program does.

The build fails on the conditional types with CUDA 12.0 to 12.3. The conditional API does not exist there, so the compiler rejects the identifiers. The page's program guards that path with #if CUDART_VERSION >= 12040; the fallback is the host-checked loop with one graph launch per iteration, which is strategy 2 unchanged.

Go deeper

Next

Day 58 moves the overlap question inside a single stream: programmatic dependent launch lets a second kernel start while its predecessor is still running. The device-side loop you built today returns in day 60's frame pipeline, where the thing the host must never do per frame is what this page removed per iteration. Day 48's fusion remains the other answer when the launches themselves are the cost.