GPU Runtime Execution Lab

In modern deep learning inference, generating a single token doesn't launch one monolithic GPU kernel. During autoregressive decoding, each transformer layer issues a flurry of micro-operations: RMSNorm, RoPE rotary embeddings, QKV projections, attention decode kernels, SwiGLU activations, linear projections, and residual additions. Across 80 layers in a large language model, producing one token requires enqueuing between 300 and 500 individual CUDA operations.

On contemporary accelerators like NVIDIA Hopper or Blackwell, each tiny kernel executes in 2 to 5 microseconds because high-bandwidth memory (HBM) and tensor cores stream through small vectors almost instantly. But the host CPU driver overhead to validate, marshal, and dispatch a single kernel launch takes 3 to 10 microseconds.

When host launch time exceeds device execution time ($T_{\text{launch}} > T_{\text{kernel}}$), the GPU starves. The hardware command processor drains its queue faster than the CPU can push new work, leaving the Streaming Multiprocessors (SMs) sitting completely idle in execution bubbles. Your server might show low GPU throughput, but it's neither memory-bound nor compute-bound. It's host-submission bound.

Why can a GPU appear underutilized even when running an inference service at 100% CPU load?

Answer

When host launch overhead exceeds kernel execution time, the GPU drains its work queue faster than the CPU can push commands. The GPU spends most of its time idling in submission bubbles between micro-kernels.

To master this boundary, we follow a four-operation pipeline from a CPU simulation into a PyTorch CUDA benchmark:

1. Copy input values into a working buffer.
2. Affine transform: compute $2x + 1$.
3. Square: compute each value's square $(2x + 1)^2$.
4. Reduce: sum the elements down to a single scalar.

On toy integer inputs [1, 2, 3, 4], the affine step yields [3, 5, 7, 9], the square step yields [9, 25, 49, 81], and the final reduction sums to 164.

The host can enqueue all four operations into the driver queue and return to user code while the GPU is still working on the initial copy. That isn't a bug. CUDA runtimes are deliberately asynchronous. Bugs appear when engineers mistake host enqueue return for kernel completion, when concurrent streams read shared memory without explicit synchronization, or when graph replay reuses stale pointer addresses.

Two clocks describe one execution

Every CUDA program runs against two independent clocks: the host CPU clock and the device GPU clock. When you call an asynchronous kernel or launch a memory transfer, control returns to your CPU thread almost immediately. Each CUDA stream functions as a first-in, first-out (FIFO) queue managed by the GPU hardware command processor.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

Under the hood, launching a kernel isn't a direct hardware jump. The CPU runtime writes command packets into a user-space driver push-buffer ring in pinned host memory. It then issues a memory-mapped I/O (MMIO) write across PCIe or NVLink to a hardware doorbell register on the GPU Work Launch Engine (WLE). The GPU's front-end command processor polls this doorbell, pulls packets over DMA, and assigns thread blocks to available SMs.

Because enqueue and execution run independently, standard host wall-clock measurements prove very little about device progress:

To visualize this asynchronous separation, consider a logical issue-order simulation. Each host submission costs one tick: the CPU submits operations at ticks 0, 1, 2, and 3, and returns to user code at tick 4. Device operations run on their own schedule, waiting for host submission, preceding stream tasks, and cross-stream dependencies. You can download the complete runtime submission simulation script.

The host returns at tick 4, but the final reduction kernel doesn't finish on the device until tick 10. If an engineer removes the cross-stream dependency between copy and affine, affine launches at tick 1, reading uninitialized buffer memory while the copy engine is still active. Square and reduce require no extra synchronization events because they sit sequentially in the same compute stream.

Stream semantics and synchronization

Commands submitted to a single CUDA stream execute strictly in issue order. Commands in different streams can execute concurrently or in any interleaved order unless an explicit dependency binds them. Having distinct streams grants permission to overlap; it doesn't guarantee that the hardware has the capacity to execute them concurrently.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

The legacy default stream trap

Default stream behavior catches engineers off-guard because CUDA provides three distinct stream operational models:

1. The Legacy Default Stream (NULL stream): By default, CUDA operations submitted without an explicit stream target stream 0 (the legacy NULL stream). The legacy default stream is implicitly synchronizing. Any command enqueued into the legacy NULL stream can't begin until all previously submitted commands across all blocking streams on the device have completed. Even more punishing, no subsequent command in any blocking stream can begin until that legacy NULL stream operation finishes. A single unintentional default stream call in a multi-tenant inference server collapses all concurrent worker streams into a serialized bottleneck.

2. Non-Blocking Streams (cudaStreamNonBlocking): Created via cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking). Non-blocking streams completely opt out of the legacy default stream's implicit barrier. Operations in non-blocking streams run concurrently regardless of pending work in the legacy default stream.

3. Per-Thread Default Streams: When compiled with the --default-stream per-thread compiler flag (or when using cudaStreamPerThread), CUDA replaces the global synchronizing NULL stream with an independent, non-blocking stream unique to each host CPU thread. This eliminates cross-thread synchronization traps in multi-threaded serving frameworks.

Why can an unadorned cudaMemcpy call in one thread stall independent worker streams running in another thread?

Answer

Without non-blocking flags or per-thread default stream compilation, unadorned calls target the legacy NULL stream. That stream acts as a global device barrier, waiting for all active blocking streams to finish and preventing subsequent stream work from starting until it completes.

Cross-stream coordination via events

To coordinate producer and consumer streams without blocking the host CPU, use CUDA events:

The synchronization sequence relies on two runtime calls:

1. cudaEventRecord(ready, copy_stream) appends an event marker into the copy stream's FIFO queue immediately after the copy command.
2. cudaStreamWaitEvent(compute_stream, ready, 0) enqueues a dependency barrier into the compute stream.

cudaStreamWaitEvent returns immediately on the host CPU (taking under 0.5 microseconds). It doesn't block the host thread. Instead, it places a dependency token into the GPU command processor's hardware queue. The device Work Launch Engine pauses execution on the compute stream until the copy stream signals the event flag, freeing the CPU to prepare subsequent pipeline stages.

Remember to record an event before asking a stream to wait on it. Asking a stream to wait on an event that has never been recorded is undefined behavior and won't block future work. Keep source and destination memory buffers pinned and alive until all consuming kernels finish.^{[2]Reference 2CUDA semanticshttps://docs.pytorch.org/docs/stable/notes/cuda.html}

Overlap requirements

Our four-operation demonstration workload is strictly sequential because each kernel consumes the output of its predecessor. However, real-world inference pipelines overlap memory copies for request $n + 1$ with compute kernels for request $n$. That overlap requires meeting strict architectural conditions:

CUDA reports device hardware capabilities like asyncEngineCount (indicating separate copy and compute engines). But hardware support alone doesn't guarantee overlap. Host dispatch delays, resource exhaustion, and implicit stream barriers can serialize operations that look concurrent on paper.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

Measure enqueue, device work, and completion

Measuring GPU execution requires tracking three non-overlapping spans:

1. Host submission time: CPU wall-clock duration around the submission loop, measured without any synchronizations inside the timed region.
2. Device elapsed time: The duration between two CUDA events recorded before and after the workload in the execution stream, evaluated only after the stop event completes.
3. End-to-end time: Host wall-clock duration from the start of submission through the final scoped device synchronization.

Never add host submission time and device elapsed time together. Host and device clocks run concurrently. CUDA event elapsed time reflects the device wall-clock span between two hardware timestamps, not the sum of individual kernel runtimes. If the host CPU takes 8 microseconds to enqueue the next kernel and the current kernel finishes in 3 microseconds, the 5-microsecond starvation bubble is included directly in the measured event duration.^{[3]Reference 3CUDA Runtime API: Event Managementhttps://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__EVENT.html}

A high host submission time paired with tiny kernels points to framework dispatch, Python overhead, or driver contention. Long kernel times paired with a host that stays far ahead indicate GPU compute or memory saturation. Nsight Systems exposes CPU runtime calls alongside GPU streams, memory engines, and kernels. Use it to verify whether host dispatch gaps are starving your SMs before attempting low-level kernel optimizations.^{[4]Reference 4Nsight Systems User Guidehttps://docs.nvidia.com/nsight-systems/UserGuide/[5]Reference 5Nsight Compute Profiling Guidehttps://docs.nvidia.com/nsight-compute/ProfilingGuide/}

Isolating steady-state performance

A disciplined benchmark isolates initialization from steady-state execution:

• Initialize the CUDA context and runtime libraries before starting timers.
• Pre-allocate all input, workspace, and output tensors.
• Warm up the execution pipeline to trigger just-in-time (JIT) compilation, autotuning, and driver memory bindings.
• Prime CUDA timing events: PyTorch initializes its underlying CUDA event handles lazily on their first record call, which can add hundreds of microseconds of artifact noise if done inside the timed loop.^{[6]Reference 6torch.cuda.Eventhttps://docs.pytorch.org/docs/stable/generated/torch.cuda.Event.html}
• Enforce strict numerical tolerance checks against an FP64 CPU reference before trusting any timing speedup.

CUDA Graphs eliminate submission overhead

A CUDA Graph represents a static directed acyclic graph (DAG) of execution nodes. Rather than repeatedly paying CPU driver and kernel launch costs on every step, a CUDA Graph lets you capture the topology once, bake it into hardware-ready descriptors, and replay it with minimal CPU intervention.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

The CUDA Graph lifecycle follows three distinct phases:

1. Capture (cudaGraph_t): Wrapping normal stream commands between cudaStreamBeginCapture and cudaStreamEndCapture intercepts kernel dispatches and constructs an in-memory DAG. No work runs on the GPU during this phase.

2. Instantiation (cudaGraphExec_t): Calling cudaGraphInstantiate compiles the logical DAG into an executable graph. The driver validates node dependencies, assigns hardware resources, binds kernel arguments, and packs the entire pipeline into a pre-built push-buffer command sequence. Instantiation is computationally expensive (often taking 10 to 50 milliseconds), but it runs only once during startup.

3. Launch (cudaGraphLaunch): Replaying the graph requires only a single MMIO doorbell write to the GPU Work Launch Engine. Launching a graph with 200 kernels takes under 2 microseconds of host CPU time, completely eliminating per-kernel launch overhead and preventing execution bubbles between nodes.

Suppose eager execution takes 18 microseconds per iteration (including host launch overhead and kernel execution), while graph replay drops the time to 12 microseconds, but incurs 12 milliseconds of upfront capture and instantiation cost. Each replay saves 6 microseconds:

$\text{Break-even Iterations} = \frac{\text{Setup Cost}}{\text{Per-iteration Savings}} = \frac{12{,}000\,\mu\text{s}}{6\,\mu\text{s}} = 2{,}000$

At 2,000 iterations, both approaches take exactly 36 milliseconds. Replay iteration 2,001 is the first iteration that delivers a net performance improvement:

Capture rules and constraints

Stream capture intercepts stream activity and converts it into graph nodes. It enforces strict runtime constraints:

• Never begin capture on the legacy default stream. Capture must occur on a user-created non-blocking stream or a per-thread default stream.
• If a capture region forks into multiple streams using events, all joined streams must merge back into the origin stream before cudaStreamEndCapture is called.
• Never call host-device synchronizations (cudaStreamSynchronize, cudaDeviceSynchronize, or cudaEventSynchronize) during capture. Doing so invalidates the capture and raises a runtime error.
• Standard host dynamic memory allocations (malloc, cudaMalloc) are strictly prohibited during stream capture. Dynamic allocations alter virtual memory mappings that the static graph can't track. To allocate memory inside a graph, use stream-ordered memory allocators (cudaMallocAsync or cudaGraphAddMemAllocNode) backed by a pre-allocated memory pool.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

PyTorch simplifies this workflow through torch.cuda.graph(), which manages side streams and coordinates with PyTorch's internal caching allocator to keep memory pools stable.^{[2]Reference 2CUDA semanticshttps://docs.pytorch.org/docs/stable/notes/cuda.html}

Virtual addresses are baked into the graph

When a graph is instantiated, the physical 64-bit virtual memory addresses of all input, intermediate, and output buffers are baked directly into kernel argument tables.

Reassigning a Python variable name doesn't update the graph. If you allocate a new tensor with x = torch.empty(...), its virtual memory address changes. Replaying the graph will either read stale values from the old address or corrupt unrelated memory if that address was reused by another allocator.

To use fresh data with a captured graph, you must allocate static input and output buffers once and hold strong references to them for the graph's lifetime. Write fresh data into place using input_buffer.copy_(new_data) before calling graph.replay(). Similarly, read or clone results from static output buffers before launching the next replay.

Why does reassigning an input variable with a freshly allocated tensor break CUDA Graph replay?

Answer

CUDA Graph instantiation bakes 64-bit virtual memory addresses directly into device launch descriptors. Reallocating a tensor changes its memory address, causing graph replay to process the old memory location.

Graph parameter updates

CUDA supports two mechanisms for updating an executable graph without paying for a full reinstantiation:

1. Node Parameter Updates (cudaGraphExecKernelNodeSetParams): Updates the launch configuration (grid dimensions, block dimensions, or scalar parameters) of a specific node in an existing executable graph.
2. Whole-Graph Updates (cudaGraphExecUpdate): Compares a newly captured logical cudaGraph_t against an existing cudaGraphExec_t and updates execution parameters in place.

Updates can adjust kernel arguments, memory pointers, and grid dimensions without rebuilding the underlying command buffer. However, they can't modify graph topology. You can't add nodes, remove nodes, or alter dependency edges. If a workload's execution structure changes, you must instantiate a new graph executable.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

Bucketed graph pools in inference runtimes

Production inference engines like vLLM and SGLang can't rely on a single static CUDA Graph. In continuous batching, the number of active requests changes on every decoding step (e.g., batch sizes 1, 3, 7, 12, or 32). Capturing a new graph on the fly is unacceptable because a 20-millisecond instantiation pause would introduce severe latency spikes.^{[7]Reference 7CUDA Graphshttps://docs.vllm.ai/en/latest/design/cuda_graphs/[8]Reference 8SGLang: Efficient Execution of Structured Language Model Programshttps://arxiv.org/abs/2312.07104}

Inference engines resolve this tension using bucketed graph pools:

1. Pre-Capture Discrete Buckets: During service initialization, the engine captures and instantiates graphs for a fixed set of batch sizes (typically powers of two or tuned intervals: [1, 2, 4, 8, 16, 32, 64, 128, 256]).
2. Padding with Attention Masking: When a batch of 6 requests arrives, the runtime rounds up to the batch-8 graph bucket. The two unused slots are populated with dummy tokens. Attention masks, loss computations, and KV cache updates are configured to ignore these padded slots, ensuring numerical precision remains exact.
3. Eager Fallback for Outliers: When an incoming batch exceeds the largest pre-captured bucket (e.g., a massive prompt prefill or an unusual batch size), the engine falls back to standard eager execution. Eager fallback is an essential correctness mechanism that guarantees reliability for arbitrary input shapes.

Persistent kernels and device-side queues

CUDA Graphs reduce host launch overhead to under 2 microseconds. But for ultra-low latency inference, speculative decoding verification, or tree-search decoding, even 2 microseconds of CPU interaction can become a bottleneck.

Persistent kernels eliminate the CPU from the execution loop entirely. Instead of repeatedly launching kernels from the host, the application launches a single, long-running grid that stays resident on the GPU for the lifetime of the process.

A persistent kernel uses NVIDIA Cooperative Groups (cooperative_groups::grid_group grid = cooperative_groups::this_grid();) launched via cudaLaunchCooperativeKernel. The grid is sized to fit within the physical SM capacity of the GPU (typically 1 or 2 blocks per SM).

Each SM worker block executes a persistent while (running) loop:

1. Threads poll a task ring buffer located in GPU device memory or host-mapped memory using atomic operations (atomicAdd) and memory fences (__threadfence_system()).
2. When a task descriptor arrives, the SM workers execute the math immediately without waiting for an MMIO doorbell write from the CPU.
3. Once finished, workers mark the task complete and poll for the next work unit.

Architectural trade-offs of persistent kernels

While persistent kernels achieve sub-microsecond dispatch latencies, they introduce severe operational trade-offs:

Because persistent thread blocks never yield their SM slots, other CUDA streams and kernels can't schedule work on those multiprocessors. If a persistent worker depends on a secondary kernel that can't launch due to resource starvation, the GPU enters an unrecoverable deadlock. Persistent kernels are powerful, specialized tools that require rigorous lifecycle management, bounded polling timeouts, and dedicated hardware instances.

Run the GPU runtime lab

Download the complete gpu_runtime_lab.py script. It benchmarks our four-operation pipeline (copy -> affine -> square -> reduce) across three submission modes:

• eager: Submits each operation sequentially from the current stream.
• stream: Submits the same chain to an explicit non-blocking stream.
• graph: Warms up the pipeline, captures static buffers, and executes via graph replay.

Run the lab on a system with a supported NVIDIA GPU:

The script evaluates a 4,096-element floating-point vector initialized with linspace(0, 1, elements). The initial step performs a device-to-device copy, ensuring data transfer mechanics are exercised alongside arithmetic compute.

Before running timed iterations, the benchmark performs an automated preflight correctness check. It fills the input buffer with zeros and ones, validating that the outputs produce exact expected sums (elements and 9 * elements) within numerical tolerances. It then restores the original input, warms up the pipeline, primes CUDA timing events, and evaluates final results against an FP64 CPU reference.

Notice how the numbers separate the bottlenecks:

• In eager mode, the host spends ~19.4 microseconds submitting the four operations. Device execution takes ~12.1 microseconds. Host dispatch overhead is the primary bottleneck.
• In graph mode, host submission plunges from 19.4 microseconds down to 2.1 microseconds, reflecting the single doorbell launch.
• On top of that, device execution time drops from 12.1 microseconds to 8.4 microseconds. Why? Because eliminating the CPU dispatch bubbles allows the GPU Work Launch Engine to schedule the four kernels back-to-back without the execution gaps present in eager mode.

Diagnostic guide for runtime bottlenecks

Quiz