GPU Runtime Execution Lab

Four CUDA operations transform [1, 2, 3, 4] into one scalar:

1. copy values to a working buffer;
2. compute $2x + 1$, giving [3, 5, 7, 9];
3. square each value, giving [9, 25, 49, 81];
4. reduce the four values to 164.

The host can enqueue all four operations and return while the GPU is still executing the copy. That isn't a bug. CUDA normally works this way. The bug appears when a benchmark calls host return "latency," when a second stream reads the buffer without a dependency, or when graph replay reuses storage whose address changed.

This lab keeps that same four-operation chain from the first simulation through a real PyTorch CUDA experiment. Stable work makes runtime choices visible. Kernel math doesn't change, so each timing difference has an execution-layer owner.

What two facts must you keep separate before optimizing this chain?

Answer

The host has submitted four operations, and the device has produced 164. A host API return proves only submission unless that API is documented as synchronizing.

Two clocks describe one execution

CUDA launches and asynchronous copies usually place work into a stream, then return control to the host before that work completes. Each stream is an in-order queue, but host progress and device progress are separate.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

The next program is a logical scheduler, not a GPU performance model. Its ticks have no time unit. They expose the two clocks without requiring CUDA hardware. Download the complete runtime submission simulation.

Host submission ends at logical tick 4. The dependent device chain ends at tick 10. Synchronizing immediately after every operation would force those clocks together, but it would also erase the queueing behavior this lab needs to inspect.

Streams carry order

Commands in one CUDA stream execute in issue order. Commands in different streams may execute concurrently or in either order unless a dependency connects them. "Different streams" therefore means permission to overlap, not a promise that hardware can or will overlap them.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

Default-stream behavior needs an explicit contract because CUDA supports two modes:

• The legacy default stream (also called the NULL stream) synchronizes with blocking streams. A NULL-stream operation submitted between otherwise independent operations can prevent their overlap.
• A stream created with nonblocking semantics doesn't synchronize with that legacy stream.
• Per-thread default-stream mode makes each host thread's default stream a regular stream instead of one process-wide synchronization point.
• A library may run on its current stream, a caller-supplied stream, or an internal stream. Confirm its API contract before assuming order.

Relying on accidental default-stream ordering makes code sensitive to build flags, stream creation flags, host thread, and library integration. Record an event after the producer instead. Make the consumer stream wait for that event.

cudaEventRecord(ready, copy_stream) inserts the event after earlier copy-stream work. cudaStreamWaitEvent(compute_stream, ready) prevents later compute-stream work from passing that event. The host doesn't have to wait. CUDA events can also carry timestamps, but dependency-only events can be created with timing disabled when no timestamp is needed.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

Overlap has prerequisites

Our four operations are intentionally dependent, so their critical path stays serial. A pipeline can overlap copy for request $n + 1$ with compute for request $n$ only after it gives the requests separate buffers and correct event edges.

CUDA exposes device properties such as asynchronous engine count and concurrent-kernel support, but a capability bit doesn't prove overlap for a particular trace. Resource pressure, transfer direction, dependencies, and host submission cadence still decide the timeline.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

Two operations appear in different streams but remain serialized. Which four owners should you inspect before changing code?

Answer

Inspect dependencies, pinned-memory eligibility, copy-engine or SM capacity, and host submission gaps. Separate stream names permit overlap; they don't prove workload and hardware conditions allow it.

Measure enqueue, device work, and completion

One timing number can't locate launch overhead. Keep three intervals:

1. Host submission time: CPU wall time around the enqueue loop, without a synchronization inside it.
2. Device elapsed time: CUDA events placed before and after the chain in the execution stream, read only after the stop event completes.
3. End-to-end time: CPU wall time from before submission through the final scoped synchronization.

These intervals answer different questions. A large host interval with narrow device work points toward Python, framework, driver, or launch dispatch. Long kernels with a small host interval point toward device execution. Large end-to-end time outside both intervals points toward setup, synchronization, memory movement, or other process work.

Nsight Systems shows CUDA API calls beside GPU streams, copies, and kernels. Use it to separate API time, queue gaps, and kernel time before reaching for a kernel profiler. Nsight Compute then explains a selected kernel's resource use and instruction or memory behavior.^{[2]Reference 2Nsight Systems User Guidehttps://docs.nvidia.com/nsight-systems/UserGuide/[3]Reference 3Nsight Compute Profiling Guidehttps://docs.nvidia.com/nsight-compute/ProfilingGuide/}

Remove cold work from steady state

A valid steady-state comparison records setup separately:

• create the CUDA context and initialize libraries;
• allocate output and workspace buffers;
• trigger module loading, just-in-time (JIT) compilation, autotuning, and lazy library setup;
• run warmup iterations on a side stream when graph capture will follow;
• synchronize once before starting timed work;
• keep shape, dtype, stream policy, power state, and background load documented.

Warmup isn't a fixed magic count. Repeat until the execution path and timing distribution stabilize, then report the chosen count. A cold-start receipt is also useful, but differs from a steady-state replay experiment.

A CPU timer reports 40 microseconds around ten launches, while a CUDA event reports 120 microseconds through same stream. What did each timer measure?

Answer

CPU timer measured roughly 40 microseconds of host enqueue work. Synchronized CUDA events measured 120 microseconds of device progress through stream. Dividing host number by ten doesn't produce kernel latency.

Graph replay removes repeated submission

A CUDA Graph represents a dependency graph of operations. Runtime cost has distinct phases:

1. Define: build nodes explicitly or capture operations issued to streams.
2. Instantiate: validate the graph and create an executable snapshot with launch resources prepared.
3. Replay: enqueue the executable graph into a stream.

Instantiation isn't replay, and capture isn't a benchmark iteration. Report graph setup outside steady-state replay timing. Graph replay can reduce repeated CPU and driver submission overhead while preserving the same GPU nodes and dependencies.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

Capture is an ordering proof

Stream capture records work issued to an origin stream and any joined streams. Raw CUDA permits any origin except the legacy NULL stream, including a per-thread default stream, and rejects patterns whose ordering can't be represented safely.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

• Don't begin capture on the legacy default stream.
• If capture forks into another stream, record an event from the captured stream, make the joined stream wait on it, and rejoin the origin before ending capture.
• Don't query or synchronize a captured stream or captured event during capture.
• Avoid unrelated uncaptured CUDA work in the same process while PyTorch capture is active.
• Ordinary CPU work executes during capture but isn't replayed. Raw CUDA can represent an explicit CPU-function graph node.
• CPU-GPU synchronization, unsupported stream APIs, or capture-unsafe library work can make a region unsafe to capture.

PyTorch's torch.cuda.graph context uses a side stream. Raw PyTorch capture requires a nondefault stream. Its safe recipe is warmup, allocate long-lived inputs and outputs, capture only safe CUDA work, then copy new values into the same input addresses before replay.^{[4]Reference 4CUDA semantics.https://docs.pytorch.org/docs/2.13/notes/cuda.html}

Addresses are graph inputs too

Raw graph replay uses the same kernel arguments and virtual addresses observed during capture. Holding a Python variable with the same name doesn't preserve its address. Reallocating a tensor, allowing a captured output to be freed, or changing a view's layout can invalidate replay or silently redirect work.

Keep strong references to captured inputs, outputs, parameters, and workspaces for the graph's full lifetime. Copy fresh values into static input buffers. Consume or clone static outputs before another replay overwrites them. When capture includes asynchronous allocation or free nodes, dependency order must still prevent access before allocation and after free.^{[4]Reference 4CUDA semantics.https://docs.pytorch.org/docs/2.13/notes/cuda.html[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

One executable graph (cudaGraphExec_t) can't run concurrently with itself. CUDA orders another launch of the same executable after its previous launch. A runtime that needs overlapping graph executions must use distinct executable instances or another execution strategy, while keeping each instance's static buffers and memory-pool lifetime safe.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

Updates have a structural boundary

CUDA offers two update styles:

• update parameters of an individual node in an executable graph;
• compare a new graph definition with an existing executable using whole-graph update.

Parameter updates can change compatible kernel arguments, memcpy parameters, or addresses without rebuilding every launch resource. They don't permit arbitrary topology changes. Whole-graph update requires matching topology and compatible node types and ordering. A changed node count, dependency shape, or incompatible function contract generally needs a new graph executable. Successful updates apply to later launches, not one already running.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

Dynamic shapes cross that boundary often. A changed size may alter grid dimensions only, or it may change temporary allocation, operator choice, kernel count, and control flow. Ordinary PyTorch graph captures therefore suit static shapes and control flow. Current PyTorch can represent supported GPU-data-dependent branches through torch.cond(), but arbitrary Python branching and CPU decisions still aren't replayed. Use fixed shape buckets, padding with masked semantics, framework-managed graph trees, or eager fallback when the topology or memory plan changes.^{[4]Reference 4CUDA semantics.https://docs.pytorch.org/docs/2.13/notes/cuda.html}

CUDA also supports device-side graph launch for constrained graph node types and launch modes. That is an advanced scheduling mechanism, not a loophole around lifetime rules. Device-launched executables require the appropriate instantiation and upload path, stay on one device, have node-type restrictions, and can't be updated from device code. Don't mix host and device launches of the same executable concurrently.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

A request changes from batch 8 to batch 9. What must remain compatible before you update one captured grid dimension?

Answer

Check allocation sizes, attention path, loop and kernel counts, dependencies, and output layout. If topology or memory plan changes, select another captured bucket, instantiate compatible graph, or run eager.

Run the GPU lab

Download gpu_runtime_lab.py. It runs the same copy, affine, square, reduce chain in three modes:

• eager: issue each operation from the current stream;
• stream: issue the same chain on an explicit nondefault stream;
• graph: warm, capture fixed buffers, then replay the captured graph.

Use an isolated CUDA machine. The script requires a CUDA-enabled PyTorch build compatible with the installed driver. uv creates the script environment; if its default PyTorch wheel lacks support for your CUDA setup, install the official compatible PyTorch build in a project environment and run the same file there.

Start with the defaults. Increase iterations only after checking shared-machine policy and total runtime. Don't run a persistent stress loop on a display GPU or shared production accelerator. Long-running kernels can monopolize resources or distort neighboring workloads. Supported devices can use compute preemption to improve sharing, but context switches add overhead.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

The script prints measured JSON rather than a promised speedup. It records device and driver identity, workload, graph setup, captured buffer addresses, timing totals, per-iteration averages, reference value, observed value, error, tolerance, and pass status. Results depend on GPU, driver, framework build, clocks, process load, and workload size. Preserve the receipt instead of copying somebody else's numbers.

Each timing total covers the full loop. The per-iteration fields divide that total by iterations. End-to-end timing performs one scoped synchronization after the loop, so its per-iteration value measures amortized steady-state execution. It isn't the latency of a request that synchronizes after every chain.

For a useful comparison:

1. save the unedited JSON receipt;
2. repeat runs in a fresh process and inspect distribution, not one sample;
3. confirm every mode passes the reference check;
4. profile one representative run with Nsight Systems;
5. explain gaps using trace evidence before changing stream or graph structure.

The explicit-stream mode is not expected to accelerate this dependent chain. Its purpose is to prove stream ownership and timing scope. Graph mode targets repeated submission overhead. If kernels dominate device time, replay may produce little end-to-end change.

Decide whether replay owns the bottleneck

Graph replay doesn't fuse kernels. Our graph still launches copy, affine, square, and reduce nodes. A compiler such as torch.compile, or a custom Triton or CUDA kernel, may combine compatible pointwise work and reduce intermediate memory traffic.^{[5]Reference 5torch.compilehttps://docs.pytorch.org/docs/main/generated/torch.compile.html[6]Reference 6Triton Programming Guide: Introductionhttps://triton-lang.org/main/programming-guide/chapter-1/introduction.html} Fusion and graphs therefore attack different boundaries:

• fusion changes GPU work by reducing nodes or memory traffic;
• graph replay changes how a stable sequence is submitted;
• streams and events change permission and order among operations;
• persistent kernels move scheduling into a long-lived device program.

Measure in that order of ownership. If fusion turns four tiny kernels into one substantial kernel, graph benefit may shrink because fewer host dispatches remain. If a required library operation forms a fusion boundary, piecewise capture can still reduce submission around it.

Inference runtimes dispatch by execution contract

Inference requests vary in batch size, token phase, sequence length, key-value (KV) cache layout, adapter, collective pattern, and attention backend. A production runtime can't safely send every request through one raw graph.

vLLM's current CUDA Graph design names five configurations: NONE, PIECEWISE, FULL, FULL_DECODE_ONLY, and FULL_AND_PIECEWISE. Its dispatcher selects FULL, PIECEWISE, or eager NONE execution from a runtime mode and a batch descriptor that includes token count, request count, uniformity, and Low-Rank Adaptation (LoRA) presence. Full graphs can cost more startup time and memory; piecewise graphs leave incompatible regions eager.^{[7]Reference 7CUDA Graphshttps://docs.vllm.ai/en/latest/design/cuda_graphs/}

SGLang likewise keeps separate prefill and decode graph policies with captured-size lists and full, breakable, tc_piecewise, or disabled backends. Its runtime code owns static capture buffers, warmup, runner selection, and eager fallbacks.^{[8]Reference 8SGLang Source Repositoryhttps://github.com/sgl-project/sglang} Those structures encode four production rules:

1. Dispatch key: shape isn't enough when backend, phase, adapter, dtype, or collective structure changes execution.
2. Static memory plan: captured addresses and workspaces must outlive replay.
3. Coverage policy: capture hot reusable buckets; don't force rare requests through unsafe padding.
4. Fallback: eager execution is part of correctness, not evidence that graphs failed.

Padding also needs a semantic proof. Padded tokens must be masked from attention, reductions, sampling, cache updates, and service accounting. Selecting a captured size only for larger capacity can be fast and wrong.

Persistent work queues move the boundary

A persistent kernel stays resident, reads work descriptors from a device-visible queue, executes tasks, and loops. It can remove repeated host launches and react to device-side work, but it exchanges runtime simplicity for a custom scheduler.

CUDA doesn't guarantee block scheduling order. A persistent design must size resident blocks using occupancy evidence and avoid assuming that a queued producer block will run before a waiting consumer block.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/} Preemption can prevent one long kernel from monopolizing a GPU, but saving and restoring execution state adds overhead. Shared-accelerator policy remains a deployment constraint, not a detail to discover under traffic.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

Graph replay is usually the lower-risk first move for a stable host-driven sequence. Persistent kernels fit specialized schedulers whose queue protocol, fairness, cancellation, observability, and recovery justify owning more runtime machinery.

Debug symptom by owner

Mastery check

Keep four receipts for one experiment:

• simulation output showing host tick 4, device tick 10, and result 164;
• eager, explicit-stream, and graph JSON with environment and all three intervals;
• Nsight Systems trace locating API gaps, stream order, copies, and kernels;
• correctness record with reference, tolerance, maximum error, and captured dispatch key.

Use this rubric:

• Incomplete: reports one CPU duration or one fast replay result without synchronized device timing and correctness.
• Operational: separates enqueue, device, and end-to-end intervals; identifies stream ownership; preserves static captured storage.
• Production-ready: also records setup, dispatch keys, fallback behavior, shape and lifetime contracts, profiler evidence, and shared-GPU risk.

Common pitfalls:

• synchronizing after every launch, then claiming independent streams never overlap;
• timing graph capture or instantiation as if it were replay;
• comparing cold eager execution with warmed graph execution;
• replacing a captured tensor instead of copying into its static buffer;
• padding shape without masking every downstream semantic effect;
• treating graph replay as kernel fusion;
• adopting a persistent queue without fairness, cancellation, or timeout drills.

If you can explain why the stream mode may match eager, why graph mode may reduce only host time, and why every mode must still produce the same scalar within tolerance, you can distinguish runtime optimization from timing theater.

Quiz