CUDA Graph Replay Pitfalls

The one rule

What a graph records

Inside with torch.cuda.graph(g): PyTorch puts a side stream into capture mode. Your Python runs normally, but every kernel launch, memcpy and memset sent to that stream is not executed. It is appended to the graph as a node instead. A kernel node stores the function, its launch shape and a copy of its argument bytes. Tensors allocated during capture come from a memory pool owned by the graph, so their addresses stay valid between replays. g.replay() then enqueues the whole node list with one call.

Your Python during capture (idx = 3)

def func(idx, x, y):

if idx % 2 == 0:CPU decides once: odd

z = x + ynever launched

else:

z = x - y→ node 1

return z * x→ node 2

What the graph keeps

node 1 · elementwise sub

reads: &x, &y
writes: &z (graph pool)
shape: 1 element, grid fixed

node 2 · elementwise mul

reads: &z, &x
writes: &out (graph pool)
shape: 1 element, grid fixed

Simplified. idx appears nowhere: it was only ever a CPU value.

Replay re-issues

every captured kernel, with its saved argument bytes (pointers and scalars)
captured memcpy and memset nodes, with their saved addresses
the same launch shapes

Replay never re-runs

Python: branches, loops, counters, prints, hooks
allocations and frees
host syncs such as .item() (illegal inside capture)
work sent to any stream that is not part of the capture

And capture itself computes nothing

Values produced inside the with block exist only after the first replay. That is why every example below replays once right after capturing.

Every failure on the page is one of four things that the recording freezes or leaves out:

frozen choice

Decisions made on the CPU

Branches, Python scalars, shapes and Python side effects only happen at capture.

frozen pointers

Addresses

The graph reads and writes fixed pointers, whatever lives there now.

missing nodes

Which work was recorded

Only the capture stream is recorded. Other work ran once, then disappears.

async replay

Replay is asynchronous

Replay only enqueues. The CPU and other streams race the queued work.

How to read the measured strips below

warmup (eager) capture, then one replay replay only observed value differs from correct

Each cell shows the observed value large and the correct (eager) value as exp beneath it. All values come from the repo's run log 02-cuda-graph-numerical-errors/results/20261019-145927.txt.

Frozen choice

Decisions made on the CPU

Anything Python decides while capturing becomes a constant of the graph: which branch, which scalar, how many elements, what the host-side state was. The fix is always one of two moves. Either turn the decision into data the GPU reads at replay time, or keep it outside the graph.

Dynamic control flow

A Python if on a runtime value picks one branch at capture. Every replay repeats that branch.

Bug page code

def func(idx, x, y):
    if idx % 2 == 0:  # ❌ Dynamic control flow!
        z = x + y
    else:
        z = x - y
    return z * x

for i in range(6):
    if i < 3:      z = func(i, x, y)      # warmup
    elif i == 3:              # capture: idx=3 is odd
        with torch.cuda.graph(graph):
            z = func(i, x, y)
        graph.replay()
    else:          graph.replay()

Fix torch.where, from the page

def func(idx_tensor, x, y):
    condition = (idx_tensor % 2 == 0)   # GPU tensor
    add_result = x + y
    sub_result = x - y
    z = torch.where(condition, add_result, sub_result)  # ✅
    return z * x

idx_tensor = torch.zeros(1, device="cuda", dtype=torch.int32)
for i in range(6):
    idx_tensor.fill_(i)  # update the GPU value every iteration
    ...

Bug, measured. Same as the page.

All four fixes (torch.where, a branching CUDA kernel, two graphs chosen on the CPU, partial capture) gave the same correct strip.

Underneath

idx % 2 is evaluated by Python, once, at capture. The graph holds a subtract node and a multiply node, and idx is not an argument of either.
torch.where works because the condition becomes a tensor. The select kernel reads it through a pointer at replay time, so fill_(i) before each replay changes the result. The cost is that both branches always run.
The custom kernel reads idx[0] from device memory and branches per thread, so only one path does work. It must be launched on getCurrentCUDAStream(), otherwise it is not captured (see library streams). It compiled with load_inline for sm_121 and ran correctly.
Two graphs (one per branch) and partial capture (branch in eager code, then a graph for the static rest) keep the decision on the CPU, outside the recording.
CUDA 12.4 added conditional graph nodes, but the page notes that PyTorch's graph API does not expose them yet.

cases: control-flow/bug · fix-where · fix-kernel · fix-multi-graph · fix-partial

Capture-aware code

Code that asks is_current_stream_capturing() and computes something different gets that different computation baked into every replay.

Bug page code

def func(x, y):
    z = x + y
    if torch.cuda.is_current_stream_capturing():
        z *= -1  # ❌ Different logic during capture!
    return z

Fix change how, never what

def func(x, y):
    z = x + y
    if torch.cuda.is_current_stream_capturing():
        torch.cuda._sleep(1000)  # ✅ graph-safe sleep alternative
    else:
        time.sleep(0.001)
    return z

Bug, measured (y.fill_(i) each iteration). Same as the page.

Underneath

The check returns True during exactly one execution of the Python: the capture. The negation kernel is recorded, and at replay time the check does not run at all, so it cannot "switch back". Changing how the same result is computed is fine. Both fixes printed 0..5.

cases: capture-aware/bug · fix · fix-sleep

Python scalars become constants

A Python number passed to an op becomes a kernel argument. Its capture-time value is copied into the graph node.

Bug page code

def func(tensor, exponent: int):  # exponent is a Python int
    return torch.float_power(tensor, exponent)

x = torch.tensor(2, device="cuda", dtype=torch.int32)
# warmup i=0..2, capture at i=3, replay i=4,5
y = func(x, i)  # ❌ i is baked in at capture

Fix page code

exponent_tensor = torch.zeros(1, device="cuda", dtype=torch.int32)

with torch.cuda.graph(graph):
    y = func(x, exponent_tensor)  # ✅ captures the address

for i in range(3, 6):
    exponent_tensor.fill_(i)  # update the value in place
    graph.replay()

Bug, measured (2ⁱ). Same as the page.

Underneath

A kernel node stores a copy of the launch's argument bytes. A Python scalar is passed by value into those bytes: 3 is in the node itself.
A tensor is passed as a pointer. The kernel loads whatever is at that address when it runs, so fill_ between replays changes the result.
Constants that never change (a dropout rate, a hidden size) are safe as Python values. Anything computed per step is not.

cases: dynamic-scalar/bug · fix

Hidden scalars: learning rate, step counter, GradScaler, RNG

The page warns about scalars that hide inside training loops. Each fragment was made runnable.

Bug learning-rate schedule

LRS = [1.0, 0.5, 0.25, 0.125]
lr = LRS[0]
with torch.cuda.graph(graph):
    w.sub_(grad * lr)  # ❌ lr is a Python float

for step in range(4):
    lr = LRS[step]     # the graph keeps LRS[0]
    graph.replay()

Fix a device tensor

lr_tensor = torch.zeros(1, device="cuda")
with torch.cuda.graph(graph):
    w.sub_(grad * lr_tensor)  # ✅

for step in range(4):
    lr_tensor.fill_(LRS[step])
    graph.replay()

Learning-rate bug, measured: w after each step, with grad = 1.

Underneath

Step counter. Same mechanism: model(x, step=global_step) froze the step at its capture value. A step tensor updated with fill_ fixed it.
GradScaler. The loss scale is already a device tensor (scaler._scale). scale(loss) multiplies by it through a pointer, and update() rewrites it in place. So capturing forward plus scaler.scale(loss).backward() and calling step() and update() in eager code is safe. With the scale doubling every step, the weights and scale matched an eager run exactly. The page's ❌ variant puts step() inside the graph. That does not fail silently: step() checks found_inf on the host, a cuStreamSynchronize, and capture aborts with operation not permitted when stream is capturing.
RNG. A random kernel in a graph reads its seed and offset through device pointers. PyTorch's generator header describes how each replay first advances that offset by the graph's total consumption. So 5 replays of a captured dropout gave 5 different masks. On torch 2.14 this also held for a custom torch.Generator that was never passed to register_generator_state. The generator creates per-capture state on first use.

cases: hidden-scalar/lr-bug · lr-fix · step-counter-bug · step-counter-fix · grad-scaler-bug · grad-scaler-fix · rng/default-auto · default-registered · custom-unregistered · custom-registered

Dynamic shapes

Launch shapes are fixed at capture. A graph captured for 4 elements processes 4 elements, always.

Bug page code

x = torch.ones(4, device="cuda", dtype=torch.int32)
with torch.cuda.graph(graph):
    y = x * 2

x_large = torch.ones(8, device="cuda", dtype=torch.int32) * 5
x.copy_(x_large[:4])  # ❌ only 4 elements fit the captured input
graph.replay()

Fix pad to a maximum, or bucket

static_input = torch.zeros(1024, device="cuda")
with torch.cuda.graph(g):
    static_output = model(static_input)

for data in dataloader:
    n = data.shape[0]
    static_input[:n].copy_(data)
    static_input[n:].zero_()             # ✅ zero padding
    g.replay()
    result = static_output[:n].clone()   # trim

Bug, measured: the 8 requested outputs. Same as the page.

Underneath

The grid size and the element count are part of the recorded launch, and the input and output buffers were allocated for 4 elements. There is nowhere to put element 5. Padding to one maximum wastes compute on zeros. Bucketing captures one graph per size class: sizes 3, 8 and 5 ran through buckets 4, 8 and 8 and came out correct. vLLM does the same per padded batch size. In topic 01, Step C, that came to 2995 graph instantiations.

cases: dynamic-shape/bug · fix-pad · fix-bucket

Host-side state mutated inside the graph

A counter incremented inside the captured function increments during capture only. Eager code after the graph then reads a frozen value.

Bug page code

def graphed_func(state, x):
    y = x * 1
    state[0] = state[0] + 1  # ❌ skipped during replay
    return y

def eager_func(state, x):
    if state[0] % 2 == 0:
        x *= -1
    return x

Fix a GPU tensor for the state

def graphed_func(state_tensor, x):
    y = x * 1
    state_tensor += 1  # ✅ captured and replayed
    return y

# eager code reads state_tensor.item() after replay
# (or: do state[0] += 1 outside the graph)

Bug, measured: y after eager_func. The state reads 1, 2, 3, 4, 4, 4 instead of 1..6. Same as the page.

Underneath

state[0] += 1 is a CPU instruction. It never becomes a node. The same applies to list.append, dictionary updates and self.step = i. If code outside the graph depends on the value, either mutate it outside, or make it a tensor that a captured kernel updates.

cases: host-state/bug · fix-outside · fix-gpu-tensor

Frozen pointers

Addresses

A graph doesn't know about Python variables or tensor objects, only raw addresses. Anything that changes what lives at an address breaks it silently: rebinding a name, freeing a tensor, reusing a buffer across replays, two graphs handed the same block.

Reassigned inputs

Rebinding a Python name to a new tensor moves the name. The graph keeps reading the old address.

Bug page code

for i in range(2):                  # warmup
    lr = base_lr + i  # ❌ Creates NEW tensor each iteration
    out = 1 * lr

with torch.cuda.graph(graph):       # captures lr's address (holds 1)
    out = 1 * lr

for i in range(2, 6):
    lr = base_lr + i    # new tensor, new address
    graph.replay()

Fix page code

lr = torch.zeros(1, device="cuda", dtype=torch.int32)  # once

for i in range(2):
    lr.copy_(base_lr + i)  # ✅ Update values, same address
    out = 1 * lr

with torch.cuda.graph(graph):
    out = 1 * lr

for i in range(2, 6):
    lr.copy_(base_lr + i)
    graph.replay()

Bug, measured. Same as the page.

What the name and the graph point at, at i = 4

block A

lr from i=1 · value 1 · freed at i=2← the graph's mul node reads here

block B

lr from i=4 · value 4← the Python name lr now points here

Underneath

Block A is freed at i = 2, the moment the name moves on. This run still read 1 only because the caching allocator happened to hand other free blocks to the later tensors. Once A is recycled, the graph reads someone else's data (next example). Watch for hidden inputs that cross the graph boundary: globals, preprocessing outputs, module attributes and caches.

cases: dynamic-input/bug · fix

The autocast weight cache

Under autocast, the FP16 copy of a weight is cached. A graph that uses the cached copy never sees the weight change.

Bug page code

w = torch.ones(1, 1, device="cuda", requires_grad=True)
with torch.amp.autocast("cuda"):
    _ = x @ w                 # w cast to FP16 and cached
    with torch.cuda.graph(g):
        y = x @ w             # ❌ reads the cached FP16 copy

g.replay()                    # y = 1
with torch.no_grad():
    w.fill_(2)
g.replay()                    # y is still 1

Fix page code, plus a warmup

# Option 1: no cache, so the cast kernel is captured
with torch.amp.autocast("cuda", cache_enabled=False):
    with torch.cuda.graph(g):
        y = x @ w             # ✅

# Option 2: autocast inside the capture
with torch.amp.autocast("cuda"):
    _ = x @ w                 # warmup (cuBLAS init), see below
with torch.cuda.graph(g):
    with torch.amp.autocast("cuda"):
        y = x @ w             # ✅

Bug, measured. Same as the page.

Underneath

Autocast caches the casted copy of a leaf tensor that requires grad, for as long as the autocast context lives. In the bug, the warmup filled the cache. The captured matmul therefore reads the cached FP16 buffer, and no cast kernel is recorded. Writing new values into w never reaches that buffer. After the context exits, the cache entry is dropped, so its memory can also be recycled under the graph. With the cache disabled, or with autocast entered inside the capture, the cast is a node, and replay re-casts w each time.

On torch 2.14

Option 2 as printed on the page has no warmup. In a fresh process its matmul is the first one, so cuBLAS creates its handle inside the capture, and capture fails with CUBLAS_STATUS_NOT_INITIALIZED. One eager matmul first, the usual "warm up before capture" rule, fixes it.

cases: autocast-cache/bug · fix-no-cache · fix-inside

Freed inputs

When the last Python reference to a graph input goes away, the caching allocator gives its block to the next request of that size.

Bug page code

def capture_graph():
    static_input = torch.ones(4, device="cuda")  # ❌ local
    g = torch.cuda.CUDAGraph()
    with torch.cuda.graph(g):
        output = static_input * 2
    return g, output        # static_input freed here

g, output = capture_graph()
other_tensor = torch.randn(4, device="cuda")
g.replay()                  # reads other_tensor's values

Fix page code

def capture_graph():
    static_input = torch.ones(4, device="cuda")
    g = torch.cuda.CUDAGraph()
    with torch.cuda.graph(g):
        output = static_input * 2
    return g, static_input, output  # ✅ keep it alive

g, static_input, output = capture_graph()
static_input.fill_(3)
g.replay()                  # [6., 6., 6., 6.]

Bug, measured: the 4 output elements.

Addresses from the run

0x32ee00000

static_input (freed on return)other_tensor = randn(4)← the graph's mul node still reads here

The output is 2 × randn: the replay doubled whatever other_tensor holds.

On torch 2.14: the CPU-tensor variant

The page's second example frees a CPU tensor whose H2D copy was captured (torch.tensor(i).to("cuda", non_blocking=True)). On torch 2.14 the copy is refused at capture: Cannot copy between CPU and CUDA tensors during CUDA graph capture unless the CPU tensor is pinned. Both of the page's fixes fail the same way, including torch.tensor(i, device="cuda"), which is itself a pageable copy. A graph-safe replacement is torch.full((), i, device="cuda"), a fill kernel.

We rebuilt the bug with pinned memory, and it still did not reproduce. After capture, new pinned tensors landed at 0x32ee00200, never at the captured block 0x32ee00000. In eager mode the same allocator does reuse a freed pinned block once its copy has finished. So PyTorch now withholds host blocks used by a capture.

cases: freed-cuda-input/bug · fix · freed-cpu-input/bug · fix-keep-alive · fix-gpu-tensor · bug-pinned · fix-pinned-keep-alive · fix-full

Output tensor reuse

A graph's output is one buffer that every replay overwrites. Appending it to a list appends the same buffer three times.

Bug page code

with torch.cuda.graph(graph):
    y = x ** 2 + 1

ys = []
for i in range(3):
    x.fill_(i)
    graph.replay()
    ys.append(y)  # ❌ Appending reference to same tensor!
out = torch.sum(torch.cat(ys))

Fix page code

ys = []
for i in range(3):
    x.fill_(i)
    graph.replay()
    ys.append(y.clone())  # ✅ Clone to preserve each result
out = torch.sum(torch.cat(ys))

Bug, measured. Same as the page.

Underneath

In eager mode, each forward pass allocates a fresh output, so keeping references is safe. In a graph, the output was allocated once, in the graph's pool, and its address is part of the recording. This matters for accumulating predictions, metrics, gradient accumulation and pipeline stages. It also matters across graphs that share a pool: another graph's replay can overwrite your output too (next example).

cases: output-reuse/bug · fix

Graphs sharing a memory pool

With pool=graph1.pool(), graph 2 is allocated as if it runs after graph 1. Replaying in another order lets graph 1's scratch memory land on graph 2's output.

Bug page code

def func1(x1):
    t1 = x1 * 3       # intermediate
    return t1 + 5     # y1 = 8
def func2(x2):
    return x2 ** 2    # y2 = 4

with torch.cuda.graph(graph1):
    y1 = func1(x1)
with torch.cuda.graph(graph2, pool=graph1.pool()):
    y2 = func2(x2)

graph2.replay()       # y2 = 4
graph1.replay()       # ❌ t1 = 3 written over y2

Fix any one of three

# 1. read each output before another graph replays
graph2.replay(); result2 = y2.clone()
graph1.replay(); result1 = y1.clone()

# 2. don't share the pool
with torch.cuda.graph(graph2):   # no pool= argument
    y2 = func2(x2)

# 3. replay in capture order  ✅
graph1.replay()
graph2.replay()

Bug, measured. Same as the page, and all three fixes gave 8 and 4.

The shared pool, with addresses printed during capture

0x327e00000

graph 1: t1 (scratch, freed at end of capture)graph 2: y2 (output)

0x327e00200

graph 1: y1 (output)

graph2.replay() writes y2 = 4 into 0x327e00000.
graph1.replay() writes t1 = 1 × 3 = 3 into 0x327e00000, then y1 = 8.
Reading y2 now gives 3.

Underneath

During capture the pool behaves like the normal caching allocator. When graph 1's capture ends, t1 has no Python reference, so its block goes back to the pool's free list, and graph 2's capture is handed that block for y2. That assignment is correct only if graph 1 always finishes before graph 2 runs. The page's rule of thumb: when sharing a pool, replay in capture order, and never concurrently.

A bug this harness made itself

The first version of the "read immediately" fix printed 17 instead of 8. The helper that built the graphs created x1 locally and didn't return it, so x1 (at 0x32ee00000) was freed. The y2.clone() was handed that block and wrote 4 into it. graph1 then computed 4 × 3 + 5 = 17. That is exactly the freed-input bug, made by accident in code written to show a different one. The other shared-pool cases allocated nothing in between, so they looked fine.

cases: shared-pool/bug · fix-read-immediately · fix-separate-pools · fix-capture-order

Replay order with `make_graphed_callables`

Graphing several callables at once makes six graphs (a forward and a backward each) in one pool. Interleaving them the way pipeline parallelism does corrupts gradients.

Bug page code

func = lambda x: x ** 2 + 2 * x + 1      # d/dx = 2x + 2
graphed = torch.cuda.make_graphed_callables(
    (func, func, func), ((x0,), (x1,), (x2,)))

y0 = graphed[0](x0)    # ❌ interleaved
y1 = graphed[1](x1)
y0.backward()
y2 = graphed[2](x2)
y1.backward()
y2.backward()

Fix the order the source documents

graphed = torch.cuda.make_graphed_callables(
    (func, func, func), ((x0,), (x1,), (x2,)))

ys = [g(x) for g, x in zip(graphed, (x0, x1, x2))]
for y in reversed(ys):  # ✅ fwd 0,1,2 then bwd 2,1,0
    y.backward()

# also correct: one make_graphed_callables call per function
# (one pool each), or capturing manually in replay order

Bug, measured gradients for x = 1, 2, 3. Same as the page.

Capture order versus replay order (F = forward graph, B = backward graph)

Captured, per torch/cuda/graphs.py

F0F1F2B2B1B0

Captured, as the page says

F0B0F1B1F2B2

Bug replay

F0F1B0F2B1B2

6, 3, 8

Page's "fix" replay

F0B0F1B1F2B2

6, 3, 8

Documented order

F0F1F2B2B1B0

4, 6, 8

Highlighted steps are out of place relative to the real capture order.

Underneath

The pool's addresses were assigned by simulating execution in capture order: a buffer that is dead after one graph is handed to the next. Saved activations that a backward graph needs are only guaranteed intact if the graphs in between are the ones the capture assumed. The source comment in make_graphed_callables says it captures "fwd 1, fwd 2, ... fwd N, then bwd N, bwd N-1, ... bwd 1". Its docstring says the tuple order must be the order the callables run in the live workload.

On torch 2.14

The page's fix 1 (y0.backward() right after f0, and so on) assumes alternating capture. It is itself a mismatch and produced the same wrong gradients as the bug. Its other two fixes (separate pools, or capturing manually in replay order) work.

cases: replay-order/bug · fix-capture-order (page, wrong) · fix-documented-order · fix-separate-pools · fix-capture-in-replay-order

Missing nodes

Which work was recorded

Capture is per stream. Work on the capturing stream, or on a stream joined to it through events, becomes nodes. Work sent anywhere else executes once, immediately, and is gone from every replay. Python that dispatches the work, like a submodule's forward and its hooks, is not part of the graph either.

A library that launches on the default stream

A library that hard-codes the default stream works in eager mode. Under capture, its kernel is never recorded.

Bug page code

def library_call(x):
    stream = torch.cuda.default_stream()  # ❌ assumes default stream
    with torch.cuda.stream(stream):
        return x * 2

with torch.cuda.graph(graph):
    y = x + 1              # captured on the capture stream
    y = library_call(y)    # runs on the default stream

Fix page code

# A: use the stream you are given (the capture stream)
def library_call_a(x, stream):
    with torch.cuda.stream(stream):
        return x * 2

# B: fork an internal stream from it and join back
def library_call_b(x, stream):
    internal = torch.cuda.Stream()
    internal.wait_stream(stream)        # ✅ fork
    with torch.cuda.stream(internal):
        result = x * 2
    stream.wait_stream(internal)        # ✅ join
    return result

Bug, measured: y = (i + 1) × 2. Same as the page.

Underneath

torch.cuda.graph captures on a side stream. x + 1 goes there and becomes a node writing into the graph pool. It has not executed.
x * 2 goes to the default stream and executes right away. It reads the pool buffer that x + 1 will only fill at replay time. In this run that buffer held 0, so the result was 0, written into an ordinary tensor. The Python name y now points at that tensor, and no replay ever writes it.
No error was raised in this run, so the only symptom is the wrong value.
Option B works because wait_stream records an event on the capturing stream and makes the internal stream wait on it. That pulls the internal stream into the same capture. The join at the end is required before capture can end. This is how multi-stream work, such as the overlap patterns from topic 01, ends up inside one graph.
The page's diagnosis tool is a profiler timeline that shows which stream each kernel ran on. That is what 01-cuda-streams/nsys_streams.py reports.

cases: library-stream/bug · fix-current-stream · fix-pass-stream · fix-internal-stream

Module hooks that never fire

After make_graphed_callables, only the top module's hooks run. The submodules' forward is never called.

Bug page code

model = torch.nn.Sequential(Embedding(4096, 4096),
                            Linear(4096, 4096), Linear(4096, 4096))
model = torch.cuda.make_graphed_callables(model, (inp,))

for module in model.modules():          # 4 modules
    module.register_forward_hook(hook)  # registered AFTER graphing
model(inp)                               # ❌ only 1 hook fires

Workaround one top-level hook

def hook(module, input, output):
    # ✅ recurse=True reaches every parameter in the tree
    for name, p in module.named_parameters(recurse=True):
        ...

model.register_forward_hook(hook)
# measured: the hook saw all 5 parameters

Bug, measured. Same as the page.

Underneath

Hooks are dispatched by Module.__call__. make_graphed_callables replaces the top module's forward with a function that replays the graph, so the top module's __call__ still runs (and its hook fires), but nothing ever calls the three submodules. Logging, activation caching or gradient tweaks done in per-layer hooks silently stop. Either capture at a finer granularity with CUDAGraph, or move that logic outside the graphed region.

cases: module-hooks/bug · workaround-top-level

Async replay

Replay is asynchronous

graph.replay() returns as soon as the work is queued. Anything the CPU, or another stream, does to the graph's memory afterwards races the queued work. These two examples belong to the CUDA streams topic as much as to graphs.

Writing pinned memory the GPU has not read yet

The graph's H2D copy reads a pinned host buffer. The CPU loop rewrites that buffer long before the queued copy gets to it.

Bug page code

def func(x):
    torch.cuda._sleep(int(50 * cycles_per_ms))  # 50 ms of GPU work
    x_cuda = x.to(device="cuda", non_blocking=True)
    return x_cuda ** 2

x = torch.empty(1, device="cpu", pin_memory=True)
with torch.cuda.graph(graph):
    y = func(x)
for i in range(4):
    x.fill_(i)        # ❌ CPU writes while the GPU may still read
    graph.replay()    # asynchronous: returns immediately
    outputs.append(y.clone())

Fix page code

event = torch.cuda.Event()
for i in range(4):
    event.synchronize()  # ✅ wait for the previous replay only
    x.fill_(i)
    graph.replay()
    outputs.append(y.clone())
    event.record()

# or: torch.cuda.synchronize() before each write,
# or: copy into a static CUDA tensor with x_cuda.copy_(x_cpu)

Graph bug, measured (x²). Same as the page.

Eager version on a side stream, measured. Same as the page: not a graph-only bug.

Graph bug, drawn to scale (0 to 200 ms)

CPU

← four writes, x = 0, 1, 2, 3, all before the first sleep ends

GPU

sleep 50 ms sleep 50 ms sleep 50 ms sleep 50 ms

050100150200 ms

CPU writes the pinned bufferqueued H2D copy reads it: always 3

Underneath

A non_blocking copy from pinned memory is a DMA that reads the host buffer when the copy engine gets to it, not when you call it. The graph recorded that host address.
replay() and clone() only enqueue, so the CPU finishes its four iterations while the GPU is still in the first 50 ms sleep. Every copy then reads 3, and every output is 9.
An event recorded after the replay and synchronized before the next write lets the CPU wait for exactly the work that reads the buffer, without draining the whole device. All three fixes gave 0, 1, 4, 9.
x_cuda.copy_(x_cpu) is safe for a different reason. A blocking copy from pinned memory makes the host wait until the copy is done, so it is ordered behind the previous replay. It also gives up the overlap.
Topic 01 showed the other half of this trade-off: pageable copies block the host (safe but slow), and pinned ones don't (fast, but now you own the synchronization).

cases: pinned/bug-eager · bug-graph · fix-synchronize · fix-event · fix-cuda-input

Replaying pool-sharing graphs on several streams

Three graphs captured into one pool, replayed concurrently on three streams, write each other's scratch buffers.

Bug page code

f1 = lambda x: x ** 2 + 2 * x + 1     # x1 = 1 -> 4
f2 = lambda x: x ** 2 + 2 * x + 2     # x2 = 2 -> 10
f3 = lambda x: x ** 2 + 2 * x + 3     # x3 = 3 -> 18
with torch.cuda.graph(g1): y1 = f1(x1)
with torch.cuda.graph(g2, pool=g1.pool()): y2 = f2(x2)
with torch.cuda.graph(g3, pool=g1.pool()): y3 = f3(x3)

for _ in range(100):
    with torch.cuda.stream(s1): g1.replay()  # ❌ concurrent
    with torch.cuda.stream(s2): g2.replay()
    with torch.cuda.stream(s3): g3.replay()

Fix page code

# 1. separate pools: safe to run concurrently  ✅
with torch.cuda.graph(g2): y2 = f2(x2)   # no pool=

# 2. or chain the streams so replays never overlap  ✅
s1.wait_stream(torch.cuda.current_stream())
with torch.cuda.stream(s1): g1.replay()
s2.wait_stream(s1)
with torch.cuda.stream(s2): g2.replay()
s3.wait_stream(s2)
with torch.cuda.stream(s3): g3.replay()

Bug, measured with 268,435,456-element (1 GiB) tensors. Every element of y2 and y3 was wrong. The page shows 13, 12, 19; the values depend on timing.

Underneath

Each f needs scratch buffers (x**2, 2*x). Captured one after another into one pool, the three graphs were given the same scratch blocks, which is correct only if they run one at a time.
On three streams they overlap (each kernel is large enough to run tens of milliseconds), and all three write those blocks at once.
The measured values fit that picture: 7 = 1 + 4 + 2 (another graph's 1² next to its own 2·2) and 11 = 4 + 4 + 3 (graph 2's temporaries in graph 3's sum). The addresses weren't logged for this case, so this is consistent arithmetic, not proof.
Both fixes left 0 wrong elements out of 268,435,456 per output. Topic 01's rule applies unchanged: streams give permission to overlap, and overlap is only correct when the two sides share no memory.

cases: parallel-replay/bug · fix-separate-pools · fix-serialize

Diagnosis

Finding the first bad layer

The page's recipe: make execution deterministic, seed identically, clone every layer's output in eager mode and inside the capture, replay, then compare layer by layer and stop at the first difference. Its snippet leaves the model out. Here the model is four Linear + tanh blocks, and block 2 contains the page's "different precision during capture" anti-pattern.

Planted bug block 2

def forward(self, x):
    x = torch.tanh(self.lin(x))
    if self.buggy and torch.cuda.is_current_stream_capturing():
        x = x.half().float()  # ❌ different precision during capture
    return x

Recipe page code, condensed

torch.use_deterministic_algorithms(True)
torch.manual_seed(42); torch.cuda.manual_seed(42)
# eager: append layer(x).clone() for each layer
# graph: same loop inside torch.cuda.graph(g), then g.replay()
for i, (e, g) in enumerate(zip(eager_out, graph_out)):
    diff = (e - g).abs().max()
    if diff > 1e-5:
        print(f"Issue first appears at layer {i}")  # ✅
        break

Layer	Max \|eager − graph\|, planted bug	Without the bug
0	0.00e+00	0.00e+00
1	0.00e+00	0.00e+00
2	2.31e-04 · stops here	0.00e+00
3	not compared	0.00e+00

Underneath

Without the bug every layer matches exactly. Replay launches the same kernels on the same data as eager mode, so any non-zero difference is a real behavioral change, not noise. 2.3e-4 is about the size of an FP16 rounding error for values just below 1, which is where tanh outputs live. Note that the clones taken inside the capture are themselves captured kernels: they only hold values after g.replay().

7 of 66 cases

Where torch 2.14 differs from the page

Every difference was traced to a cause. Most of them are recent PyTorch turning a silent failure into a loud one, or into a non-issue.

The replay-order fix is wrong

make_graphed_callables captures all forwards, then the backwards in reverse. The page's "fixed" alternating order still gives 6, 3, 8. Details

Pageable host copies are refused during capture

The freed-CPU-tensor bug and both of its fixes fail at capture instead of returning garbage. torch.full is a graph-safe way to make the tensor on the GPU. Details

Captured pinned blocks are not recycled

Rebuilt with pinned memory, the freed-host-tensor bug did not reproduce: the captured block was withheld from later allocations. Details

Unregistered custom generators still advance

Per-capture generator state is created on first use, so explicit register_generator_state is optional on 2.14. Details

GradScaler step inside the graph fails capture

Its host-side found_inf check is a stream sync, so the page's ❌ variant raises instead of using a stale scale. Details

The autocast-inside fix needs a warmup

Without one, cuBLAS initializes inside the capture and the capture fails with CUBLAS_STATUS_NOT_INITIALIZED. Details

The first five items account for the seven differing cases; the pageable-copy item covers three of them. The sixth is a gap in the page's snippet: with the warmup added, that case matches the page.

Practice

Before you capture

Audit everything that crosses the boundary of the captured region. Most bugs on this page are correct at the capture iteration and wrong only afterwards, so compare against eager mode on several iterations.

Every tensor read from outside

Allocated once, kept alive for the graph's lifetime, updated with copy_ or fill_. Never rebound.

Every Python value

A branch, scalar or shape is a constant of the graph. Make it a device tensor, or keep it outside and capture one graph per case.

Every side effect

Counters, appends, hooks and attribute writes happen once. Move them out, or into a GPU tensor.

Every library launch

It must use torch.cuda.current_stream(), or fork and join with wait_stream. Check in a profiler.

Every output you keep

The next replay overwrites it, and so may any graph that shares the pool. clone() what you keep.

Every host buffer and stream

Wait on an event before rewriting pinned inputs. Never replay pool-sharing graphs concurrently or out of capture order.

Run log

All 66 cases

A bug case is reproduced when its output differs from the correct one. A fix case is correct when it matches. Each case ran in its own process.

Show the full verdict table

Case	Kind	Verdict

What a graph records

Your Python during capture (idx = 3)

What the graph keeps

Replay re-issues

Replay never re-runs

And capture itself computes nothing

Decisions made on the CPU

Addresses

Which work was recorded

Replay is asynchronous

Decisions made on the CPU

Dynamic control flow

Capture-aware code

Python scalars become constants

Hidden scalars: learning rate, step counter, GradScaler, RNG

Dynamic shapes

Host-side state mutated inside the graph

Addresses

Reassigned inputs

The autocast weight cache

Freed inputs

Output tensor reuse

Graphs sharing a memory pool

Replay order with make_graphed_callables

Which work was recorded

A library that launches on the default stream

Module hooks that never fire

Replay is asynchronous

Writing pinned memory the GPU has not read yet

Replaying pool-sharing graphs on several streams

Finding the first bad layer

Where torch 2.14 differs from the page

The replay-order fix is wrong

Pageable host copies are refused during capture

Captured pinned blocks are not recycled

Unregistered custom generators still advance

GradScaler step inside the graph fails capture

The autocast-inside fix needs a warmup

Before you capture

Every tensor read from outside

Every Python value

Every side effect

Every library launch

Every output you keep

Every host buffer and stream

All 66 cases

Replay order with `make_graphed_callables`