Foundry: Template-Based CUDA Graph Context Materialization for Fast LLM Serving Cold Start Xueshen Liu1,∗ Yongji Wu2,∗,† Yuncheng Yao3 Danyang Zhuo3 Ion Stoica2 1 University of Michigan 2 UC Berkeley 3 Duke
arXiv:2604.06664v1 [cs.DC] 8 Apr 2026
Abstract
often leads to low utilization [36]. Moreover, the optimal parallelism configuration depends on workload characteristics, with different strategies trading off latency, throughput, memory footprint, and communication overhead [9]. To improve resource utilization, modern LLM serving stacks increasingly rely on elasticity. Autoscaling mechanisms scale the number of serving instances in response to real-time loads [13, 19, 26, 27, 48]. Complementarily, parallelism hot switching techniques dynamically reconfigure the parallelism strategy to better match the current workload characteristics [9, 25, 46]. Together, these approaches aim to sustain high utilization while meeting latency/throughput SLOs under rapidly changing traffic patterns. However, both scenarios are bottlenecked by cold start latency [26, 51]. Launching a new serving instance—or reconfiguring an existing one—requires (i) loading (resharding) model weights and (ii) (re)capturing GPU execution graphs, e.g., CUDA graphs on NVIDIA GPUs or HIP graphs on AMD GPUs. Recent systems reduce weight loading to 1–2 seconds via RDMA-based transfer from peer instances, even for trillion-parameter models such as Kimi-K2 [3, 33]. In contrast, graph capture can take several minutes and often dominates end-to-end cold start time [51]. CUDA graphs have become the de facto solution employed by inference frameworks to reduce the CPU kernel launching overhead [43]. With CUDA graphs, a series of independent kernel launches are grouped into a single launching unit. They significantly improve the performance of LLM decoding inference. However, CUDA graphs require a timeconsuming capturing forward to construct. Concretely, the inference framework executes the CPU-side model forward logic, while the CUDA runtime intercepts kernel launches and records their functions, parameters, and execution dependencies to construct the graph. Because graph capture still performs extensive CPU operations as in a normal forward pass, it incurs significant latency overhead [51]. At first glance, this problem seems amenable to simple serialization: the system could dump the topology of captured CUDA graphs offline and reload them during online serving. However, a CUDA graph is not merely a topology description; it is tightly coupled to the execution context in which it was captured. In particular, graph nodes may reference device-side resources, including memory pointers embedded in kernel arguments and kernel function handles resolved by the CUDA runtime. These context-dependent
Modern LLM service providers increasingly rely on autoscaling and parallelism reconfiguration to respond to rapidly changing workloads, but cold-start latency remains a major bottleneck. While recent systems have reduced model weight loading to seconds, CUDA graph capture still takes tens of seconds to minutes and often dominates startup. Unfortunately, CUDA graphs cannot be naively serialized: beyond graph topology, they are tightly coupled to execution context, including device addresses embedded in kernel arguments and kernel code lazily loaded during warmup. Existing approaches either rely on brittle kernel-specific patching or heavyweight process-level checkpoint/restore that are inflexible to dynamic parallelism switching. We present Foundry, a template-based CUDA graph context materialization system that persists both graph topology and execution context during an offline processing stage, and reconstructs executable graphs online with negligible overhead. Foundry enforces deterministic memory layouts, automatically extracts and reloads kernel binaries required by captured graphs, and reduces online reconstruction costs through topology-based templating. For distributed serving, Foundry further enables a single-GPU offline capture to generate templates for multi-GPU deployments by patching only rank-dependent communication state. Across dense and MoE models up to 235B parameters, Foundry reduces cold-start latency by up to 99%, cutting the initialization time of Qwen3-235B-A22B from 10 minutes to 3.9 seconds while preserving the throughput gains of CUDA graphs. Foundry is open-sourced at https://github.com/foundry-org/foundry.
1
Z. Morley Mao1
Introduction
The rapid progress of large language models (LLMs) is reshaping a wide range of applications, from chatbots and coding assistants to browser agents. At the same time, serving LLMs at scale incurs substantial infrastructure cost—recent reports estimate that OpenAI alone has spent roughly $12B on inference compute since 2024 [56]. This cost is exacerbated by the highly dynamic and unpredictable nature of LLM serving workloads: request rates and sequence lengths can shift dramatically over minutes [36, 44, 49]. As a result, static provisioning (e.g., allocating a fixed number of GPUs per model) ∗ Xueshen Liu and Yongji Wu contributed equally. † Corresponding author: Yongji Wu <[email protected]>
1
references make CUDA graphs inherently non-portable and prevent straightforward serialization. Existing approaches explore two different designs to mitigate this problem. One line of work, exemplified by Medusa [51], adopts a patch-based graph restoration mechanism. It materializes the graph topology, but without any device-side resources. At online serving, it applies a hand-crafted rule to manually trigger the loading of relevant device resources and input parameters for each used kernel. Reliance on perkernel patching rules makes Medusa difficult to support rapidly evolving hardware platforms and model architectures, which frequently introduce new custom kernels. At the other end of the spectrum, process-level checkpointing supported by modern GPU drivers, e.g., NVIDIA’s cudacheckpoint [10], can be leveraged to snapshot a LLM serving instance after it is initialized. However, process checkpointing produces substantially larger checkpoint images as it blindly bundles all GPU and CPU states across all worker processes. It also fails at dynamic parallelism switching, as the states of in-flight requests are lost when restoring the full worker processes. In addition, model providers typically serve many models under a wide array of parallelism strategies [9, 25, 42], while they frequently integrate new kernel and engine optimizations to improve serving performance. Hence, offline processing itself must be economical. It is infeasible to frequently spin up the full target number of GPUs for each model under each parallelism configuration, just to create checkpoint images or graph topology files. To address these limitations, we present Foundry, a system that persists CUDA graph states through template-based context materialization. Foundry materializes both the topology and execution context of the captured graphs, making graph restoration kernel-agnostic and eliminating the need for hand-crafted patching rules. To minimize the cost of offline preprocessing, Foundry enables a single GPU to materialize graph templates for distributed inference under different parallelism configurations. During online serving, Foundry instantiates these templates into executable graphs for each GPU by configuring the rank-dependent state of communication libraries (e.g., NCCL [30] or NVSHMEM [31]) referenced by the graphs. However, several challenges remain in realizing Foundry. First, how can Foundry materialize execution context in a general manner that supports arbitrary kernels? Second, how can Foundry transparently enable a single-GPU worker process to record graph templates for multi-GPU inference? Third, how can Foundry minimize online restoration overhead, given that inference engines typically capture many graphs spanning a wide range of batch sizes to optimize performance? To address the first challenge, Foundry intercepts CUDA memory allocations and redirects them to virtual memory 2
NCCL init Model loading
2.7
0
6.7
7.6
10
torch.compile Memory profiling
3.4
20
Graph capture Sampler warmup (0.05s)
Total: 49.9s
29.4
30
Time (s)
40
50
Figure 1. Breakdown of vLLM worker initialization of serving Qwen3-14B on 2xH200s.
management (VMM) APIs to enforce a deterministic memory layout; in addition, Foundry automatically serializes and restores the in-memory binaries of the kernels used in the captured graphs. To address the second challenge, Foundry introduces a stub layer over communication libraries to emulate distributed communication during single-GPU offline processing. Finally, Foundry further reduces restoration overhead by allowing graphs with the same topology to share a common template and by applying in-place updates to graph-node parameters. We implement Foundry in PyTorch and prototype its integration with vLLM [23] to persist CUDA graphs. We evaluate Foundry using both dense and MoE models under data and expert parallelism. The results show that Foundry reduces vLLM cold-start latency by up to 99%, lowering the initialization latency of Qwen3-235B-A22B [47] from 10 minutes to 3.9 seconds, while preserving the inference performance gains of CUDA graphs. In summary, we make the following contributions: • We propose context materialization to mitigate CUDA graph capture overhead in LLM serving cold start. By persisting both graph topology and execution context, Foundry enables graph restoration for arbitrary kernels without kernel-specific patching rules. • We develop a template-based graph restoration mechanism that enables a single GPU to record graph templates for multi-GPU inference. In addition, graphs captured at different batch sizes can share a template when they have the same topology, reducing restoration overhead. • We comprehensively evaluate Foundry, demonstrating its effectiveness in reducing LLM serving cold-start latency across a range of deployment scenarios.
2
Background and Motivation
2.1
LLM Serving Cold Start
LLM service providers such as OpenAI [32], as well as serverless LLM platforms such as Amazon Bedrock [2], typically rely on autoscaling to adjust the number of serving instances for each model in response to changing online demand. When demand spikes, the system launches new serving instances, whose cold-start latency directly increases the time-to-first-token (TTFT) of affected requests. As a result,
Mean TPOT (ms)
40
Materialized CUDA Graphs
35
Materialized Execution Context Memory layout 0x6000 metadata
vLLM with graph vLLM w/o graph
15
0x6800
Deterministic memory layout A: 0x6600
...
Preload binaries Packed kernel binaries
GEMM_kern
(a) Foundry: Context Materialization
10 16
32
64
128
#Prompts
256
A: 0x6600
512
Kernel node GEMM_kern
Figure 2. Time per output token (TPOT) under different batch sizes for Qwen3-30B-A3B on 2xH200s with EP2, using vLLM with and without CUDA graphs.
(b) Medusa: Online Patching
0x8000
0x8800
A: 0x6600
0x8600
....
GEMM_kern
GEMM handle
Layer0.forward()
New memory layout Trigger-based kernel loading
Figure 3. Medusa materializes only the graph topology, while Foundry also materializes the execution context. 2.2
CUDA Graphs are Performance-Critical
CUDA graphs (and HIP graphs on AMD GPUs) provide a graph-based work submission model that represents a GPU workload as a directed graph, where nodes correspond to GPU operations and edges encode their execution dependencies. This abstraction allows a sequence of operations to be defined once and then launched repeatedly as a single execution unit. CUDA graphs have become a de facto mechanism for accelerating LLM inference because LLM decoding consists of many short-lived GPU kernels, especially as GPU architectures continue to improve. If these kernels are launched one by one from the CPU, host-side launch overhead can become a significant fraction of end-to-end latency. CUDA graphs mitigate this overhead by launching an entire graph of kernels at once, rather than issuing each kernel individually. As shown in Figure 2, disabling CUDA graphs noticeably degrades decoding performance. CUDA graphs, however, come with a non-trivial initialization overhead. In modern inference engines, they are usually built through stream capture, in which a regular model forward pass is executed while the CUDA stream is placed in capture mode. During this procedure, the CUDA runtime records all kernel launches and other CUDA operations, such as memcpy, along with the dependencies between them, into a graph. In practice, inference engines also perform several warmup forward passes before capture so that one-time initialization steps, including cuBLAS initialization and Triton autotuning, are completed in advance [43].
cold-start latency often becomes the critical-path bottleneck of autoscaling in LLM serving. The cold-start process of an LLM serving instance consists of two parts. The first is environment initialization, which is largely model-agnostic and includes spawning Python worker processes, importing PyTorch/vLLM and their dependent libraries, and setting up generic host-side runtime state. Prior work shows that this cost is modest (typically under 10s) and can largely be removed from the critical path using pre-warmed workers that have already completed process creation and library import, but do not yet occupy modelspecific GPU resources [19, 51]. The second is worker initialization, which constructs model- and deployment-specific serving state on the GPU, including communicator initialization (e.g., NCCL), model weight loading, compilation cache loading (e.g., torch.compile), KV-cache profiling, and CUDA graph capture for efficient execution. In Figure 1, we breakdown the worker initialization phase of vLLM. Most of the components have already been substantially optimized by prior work. Weight loading can be accelerated via RDMA-based transfer [3, 13, 33, 52, 55], which enables a 1T parameter model to be loaded in under 2s [33]. KV cache profiling can be performed offline to eliminate its runtime overhead [51]. After these optimizations, CUDA graph capture remains the dominant bottleneck. As shown in Figure 1, vLLM already loads warmed torch.compile JIT caches rather than recompiling from scratch, yet it still takes 8s. In contrast, Foundry bypasses even this overhead by directly materializing CUDA graphs together with the compiled kernel binaries embedded in the execution context, eliminating the need for torch.compile’s Python bytecode transformation. Besides cold starts, recent work has explored dynamic parallelism reconfiguration [9, 18, 25, 42] to further improve resource efficiency in LLM serving in response to workload dynamics. These approaches also suffer from the high cost of recapturing CUDA graphs after each parallelism switch.
2.3
Existing Approaches and Their Pitfalls
Recently, Medusa [51] proposed mitigating CUDA graph capture overhead by materializing graph topology during an offline capture phase and reconstructing graphs using CUDA’s explicit graph construction APIs [29]. As shown in Figure 10, graph construction via these APIs is 2–3× faster than stream capture. However, a CUDA graph is defined not only by its topology, but also by the execution context in which it was captured, whereas Medusa materializes only the former, as shown in Figure 3. In particular, each kernel node contains a handle to the corresponding kernel function, while 3
SAVE Phase vLLM Engine Init
init_model()
init_kv_cache()
model.forward() ... torch.matmul(x,w) ...
Opaque struct w_ptr: 0x6200
0x6000 ....
...1a10000 000000000 6200110b1..
0x6400 Model
Materialized CUDA Graphs
Kernel Node: - Argument buffer Node attributes Launch dims - Kernel function
KV cache
0x6800 ....
LOAD Phase Kernel Node: Argument buffer Node attributes Launch dims Kernel function
0x6000 ....
Materialized Execution Context Base addr: 0x6000 Final offset: 0x6800
Binaries are preloaded.
cubin Packed modules
GEMM module
Extract used modules (§4.1.2)
0x6800
Bump
Model
...
Memory preallocation (§4.1.1)
Deterministic memory allocation (§4.1.1)
_MxN_gemm_ kernel handle
Identical memory layout! No argument patching.
Kernel Catalog: module hash -> kernel list module hash -> kernel list
... _MxN_gemm_ _AxB_gemm_ _CxD_gemm_
Reload kernel modules (§4.1.2)
Figure 4. Overview of Foundry’s workflow. its argument buffer may embed pointers to GPU-resident objects such as model weights, the KV cache, and activations. To restore a graph during online serving, Medusa applies hand-crafted, kernel-specific patching rules to rewrite these pointer arguments so that they refer to memory allocated in the current run. To recover kernel handles, Medusa still executes and captures the first layer of the model for each batch size, thereby triggering the loading of the relevant kernel modules. This reliance on hand-crafted patching rules makes Medusa difficult to generalize to new models and hardware platforms, where new kernels are introduced frequently. Moreover, its design becomes increasingly fragile as models and kernel libraries evolve. For example, newer versions of cuBLAS pack pointer arguments into opaque structures, making them difficult to identify and patch (see Appendix B). In addition, many recent models use a dense architecture in the initial layers and MoE architecture thereafter, invalidating Medusa’s assumption that all layers are structurally identical and that the first layer suffices to trigger the loading of all relevant kernels. Meanwhile, process-level checkpoint/restoring (C/R) is recently supported by CUDA driver [10, 39], which can be used in conjunction with CRIU [8] to checkpoint and restore the worker process along with its CUDA states. However, this approach tightly couples CPU and GPU state at the process level, making it ill-suited to dynamic parallelism switching, where request and KV-cache state must be preserved while the parallelism configuration changes. In addition, the CUDA driver’s C/R functionality suffers from slow restoration and does not yet support the IPC memory required for multi-GPU inference. Beyond driver-integrated C/R, many interceptionbased C/R techniques [45, 50] have also been proposed. These systems intercept CUDA driver API calls, record resource
handles, and restore them via API replay. However, none of them supports CUDA graphs. Furthermore, as a single serving instance has scaled to hundreds of GPUs with the popularity of MoE models [11, 53], it is infeasible to frequently allocate a large number of GPUs solely for offline processing. Even if Medusa and C/Rbased approaches apply to multi-GPU scenarios, they still suffer from high processing cost, as they need to prepare a materialized graph or checkpoint for all worker processes, even when those workers share the same computation logic.
3
Overview
We propose Foundry, a system that persists captured CUDA graphs together with the execution context they depend on during a one-time offline capture run, and restores them in a fresh serving instance at negligible cost. A captured CUDA graph is not self-contained: its kernel nodes may embed device addresses that are valid only in the capturing run, and may reference kernels that are lazily loaded during warmup. These context-dependent references make captured graphs non-portable to fresh processes. Foundry addresses this portability problem by materializing the execution context required for graph replay. At a high level, it re-establishes the two conditions that make captured graphs replayable in a fresh process: the memory layout expected by captured kernel arguments, and the kernel functions referenced by graph nodes. Concretely, Foundry enforces a deterministic memory layout by interposing on the CUDA driver’s virtual memory management (VMM) APIs (§4.1.1), and restores the required kernel functions by automatically extracting and reloading the kernel binaries used by the captured graphs (§4.1.2). As a result, Foundry remains library-agnostic and does not require inspecting or patching kernel-specific argument layouts. 4
As shown in Figure 4, Foundry operates in two phases: SAVE and LOAD. During SAVE, Foundry runs the engine’s normal warmup and graph capture once, while intercepting CUDA driver calls to record the graph-dependent state required for replay. The output is a portable archive that packages serialized graph metadata together with the executioncontext information needed for reconstruction, including memory-layout metadata and kernel binaries. During LOAD, each serving process consumes this archive to restore the execution context and reconstruct executable graphs, without re-executing warmup or graph capture. Foundry reduces LOAD cost by exploiting a key structural regularity of CUDA graphs: graphs captured for different batch sizes often share the same topology and differ only in per-node parameters. Concretely, the topology includes node types, dependency edges, and other attributes like cluster dimensions, while the per-node parameters include kernel arguments and launch dimensions, i.e., gridDim and blockDim. Since the CUDA driver supports updating per-node parameters on a constructed graph without re-instantiation, Foundry builds only one template graph for each unique topology instead of reconstructing every captured graph from scratch. Inference engines may capture hundreds of graphs, but these typically collapse to only a small number of unique topologies (e.g., 22 for Qwen3-14B on a H200). Foundry then specializes the template on demand at serving time by applying the target per-node parameters. (§4.2.1). Foundry further extends the templating insight to distributed LLM serving. In SPMD-style parallelism, such as DP, TP, and EP, all ranks follow the same computation flow and thus share the same graph topology, differing only in specific model shards and rank-dependent communication state. Foundry exploits this invariance by performing SAVE on a single GPU with dummy communication, and then reconstructing rank-specific graphs during LOAD by patching in the actual communication handles and rank identifiers. As a result, a single offline capture can be reused across all ranks in a distributed deployment, eliminating redundant warmup runs and reducing archive storage proportionally, e.g., by 64× for a 64-GPU cluster. (§4.2.2). In summary, Foundry performs SAVE once on a single GPU to materialize graph templates and the execution context required for graph replay, and produce a portable archive. During serving, each rank runs LOAD to restore executable graphs from the templates. This design allows offline-captured templates to be reused across batch sizes and distributed ranks, reducing both cold-start latency and offline processing cost.
4
Design
4.1
Execution Context Materialization
Application
Hook Layer
cudaMalloc
VMM APIs
Kernel's first launch
cuModuleLoad series of APIs
Execution Context Assign contiguous address Memory
Dump to disk
cubin
Figure 5. Foundry works as an interposition layer by implementing a CUDA driver hook.
the capturing run, and its kernel nodes may reference functions loaded during the pre-capture warmup runs for each batch size. In a normal engine startup, both dependencies are satisfied implicitly as a side effect of initialization: model and KV-cache initialization allocate the required device state, and warmup execution triggers the loading of the required kernel modules. Foundry makes these dependencies explicit by materializing the execution context required for graph replay, so that LOAD can directly reconstruct executable graphs in a fresh process without re-executing warmup. 4.1.1 Deterministic Memory Layout. Nodes in a captured CUDA graph embed memory pointers in their function arguments, and some CUDA libraries (e.g. cuBLAS) further hide pointers inside opaque structs passed as flat byte arrays (Appendix B). On a second run, CUDA driver may return different memory addresses even for the same program, and thus all of these memory pointers are stale [51]. As discussed in section 2, patching embedded pointers to the new addresses requires per-library knowledge of private parameter layouts and does not generalize. Foundry instead employs deterministic memory allocation mechanism to ensure data required by the graphs is at the same location on each run. Foundry resolves the problem by interposing the CUDA driver’s virtual-memory management (VMM) interface. All CUDA allocation requests are redirected into a reserved virtual-address region starting at a fixed base address chosen big enough to avoid conflicts. For LLM inference engines, these pointers in the graph usually refer to memory for model weights, the KV cache pool, and the input/output cache. Those are large, long-lived objects that are allocated once and never freed throughout the serving lifetime [23]. Besides, framework-level caching allocators retain reserved memory rather than returning it directly to the driver [35]. Foundry therefore exploits this monotonic nature, by placing the address of each allocation contiguously after the previous one during SAVE. Given the same base and the same allocation sequence, this mechanism deterministically produces an identical and continuous address layout. No kernel argument buffer needs to be inspected or rewritten. Because the actual device memory free call during initialization is rare, the fragmentation waste is negligible.
A captured CUDA graph is not self-contained: its kernel arguments may contain device addresses that are valid only in 5
During LOAD, because Foundry skips the graph capture to save time, the transient intermediate buffer allocations during the graph capture window will not occur, while their addresses are also embedded in the captured graph. Therefore, the allocation sequences between SAVE and LOAD are not perfectly identical. Because Foundry tracks memory event, it handles this problem by recording all capture-window allocations during SAVE and replaying them during LOAD, ensuring that the full address space expected by the graph is present (§5). This is also why simply disabling address-space layout randomization (ASLR) [22] would not suffice, as the allocation sequences themselves differ between SAVE and LOAD. The monotonic allocation strategy also enables a performance optimization via preallocation, which is critical to efficient graph reconstruction. Each virtual-memory allocation requires mapping an address range and setting access permissions, and repeated calls to allocate fine-grained buffers introduce significant overhead [34]. Because initialization under Foundry will produce a continuous memory layout, the final offset recorded during SAVE fully describes its extent. LOAD can therefore map the entire range up to that offset in a single allocation before any individual allocation occurs. All subsequent allocation requests, including model weights and KV cache initialization, and the capture-window buffer replay, simply verify that the requested range falls within the preallocated region and advance the pointer offset, reducing each allocation to a nanosecond pointer bump. This also ensures that every device address referenced by graph nodes is live before graph construction begins, which also allows the template-based reconstruction §4.2.1 runs asynchronously from other initialization tasks.
KernelNode KernelNode KernelNode
CUDA Graph
MemsetNode
Topology Key
KernelNode Param1: Arg: 0x006600ff ... Func: AxB_gemm KernelNode Param2: Arg: 0x006811aa ... Func: CxD_gemm
Parameter Sets
Figure 6. CUDA graph topology and node parameter sets.
binary to its contained kernel entry points, keyed by the binary’s content hash and each kernel’s mangled function name. During LOAD, these binaries are restored into the driver using the recorded load paths. When graph reconstruction encounters a kernel node, it resolves the node’s function reference by looking up the same (hash, name) key in the catalog, obtaining a valid kernel function handle without depending on a warmup execution to lazily reload the kernel. Compilation (e.g., torch.compile) and autotuning work paid during cold start are also avoided. Some binaries require additional preparation before they can be loaded. Certain CUDA libraries produce device code as multiple relocatable segments that must be linked into a single binary before they can be loaded as a module. Repeating this linking at LOAD time would add unnecessary latency. Foundry instead pre-links the segments into a ready-to-load cubin during SAVE. Additionally, some modules require postload initialization of device-side runtime state (e.g., DeepEP’s communication kernels depend on NVSHMEM, which must initialize device-side state pointers inside each loaded module). Foundry detects such requirements during SAVE and records a flag in the archive to avoid probing every loaded module during LOAD.
4.1.2 Binary Extraction and Reload. Each kernel node in a CUDA graph holds a reference to a kernel function that performs the actual GPU computation. Kernel functions are organized into a CUmodule (or CUlibrary introduced in CUDA 122 ), including proprietary libraries such as cuBLAS. These modules enter the CUDA context via cuModuleLoad (or cuLibraryLoad) series of APIs. By default, these modules are lazily loaded when a kernel is first launched or explicitly requested. A process that skips warmup never triggers these loads, so directly replaying the graph fails due to the absence of the required kernel functions. Foundry resolves this issue by intercepting those load APIs during SAVE to track which modules/libraries are loaded and restore them at the beginning of LOAD. During SAVE, Foundry extracts the code objects of used modules/libraries from the in-memory shared libraries (e.g., libcublas.so) and records: (1) the binary payload of a module/library, (2) the load API and options used to load it (so LOAD can replay the same driver call), and (3) a catalog that maps each
4.2
Efficient Graph Reconstruction
With deterministic memory layout and preloaded binaries, Foundry can reconstruct replayable graphs from serialized metadata. To scale graph reconstruction in realistic distributed LLM serving, Foundry must address two bottlenecks. Within a single rank, inference engines typically capture one CUDA graph per supported batch size to minimize latency, and rebuilding hundreds of graphs through the CUDA driver is expensive. Across ranks, naïvely producing the archive for an entire cluster requires running SAVE on every rank, consuming the full hardware allocation and storing redundant copies of graphs and binaries. Foundry addresses both through templating. Because captured graphs share identical topology and differ only in per-node parameters, a small number of templates can represent the entire graph set. 4.2.1 Topology-based Intra-Rank Graph Grouping. Reconstructing a CUDA graph incurs many CUDA driver calls, including node addition, attribute setting, and graph instantiation, and the total number of such calls scales with
2 NVIDIA CUDA Driver API, “Library Management,” https: //docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__LIBRARY.html.
6
model size and graph complexity. Although NVIDIA claims that graph objects are not internally synchronized3 , we observe significant contention on driver calls even when building independent graphs on separate threads. The per-call latency increases with thread count and total wall-clock time barely improves (Appendix A). Simply parallelizing graph construction is therefore neither safe nor effective. After analyzing the CUDA graphs for multiple LLM architectures, we realized that graphs captured for nearby batch sizes share the same topology. They share identical node types in the same order with the same dependency structure, and differ only in per-node parameters (Figure 6). For example, among CUDA graphs for Qwen3-14B [37] captured for batch size 1–512 on an H100, there are only 22 unique topologies. On the other hand, we found out CUDA driver provides APIs to update per-node parameters of an instantiated graph executable [29], and it runs up to 20x faster than assemble and instantiate a new graph (§6.4.2). Therefore, rather than building every graph during LOAD, Foundry builds only one template graph per distinct topology and prepares parameter sets of all other graphs in the same group that can be applied to the template on demand. Foundry identifies unique topologies during SAVE to reduce the reconstruction workload of LOAD. After capturing all graphs, Foundry computes a topology key for each graph that encodes its structural and topological properties in a compact form. Graphs with identical keys are grouped together, and the resulting grouping information is recorded in a manifest for later reconstruction. During LOAD, Foundry only constructs template graphs via CUDA driver APIs. Non-template graphs require only node parameter preparation, which involves no driver calls and its speed scales linearly with CPU threads. Because Foundry directly determines the grouping via the manifest, template construction and on-demand parameter preparation run concurrently: multiple worker threads prepare nontemplate graphs in parallel while the main thread builds templates sequentially to avoid driver contention. The details of topology grouping and on-demand update mechanism are described in §5. Furthermore, because the memory are preallocated at the beginning of LOAD (§4.1.1), addresses embedded in the graph are already valid, and thus the whole graph reconstruction can run asynchronously with the foreground initialization (e.g. model weights, kv cache and capture window buffers). At replay time, if the target graph’s batch size matches that of template graph, Foundry replays it directly. Otherwise, Foundry applies the node parameters of the target graph to update the template executable before replay. This is a one-time effort, as repeated calls to the same graph skip
the update entirely. We demonstrate that such on-demand specialization is lightweight in §6.4.2. 4.2.2 Inter-Rank Graph Sharing. In SPMD-style parallel deployment, such as data parallelism (DP), tensor parallelism (TP), and expert parallelism (EP), the computation is invariant across ranks, while the communication kernels differ in rank-specific arguments (e.g., the rank identifier passed to collective kernels). The shared execution flow across participating GPUs ensures that they load identical kernel binaries, while Foundry’s deterministic allocation mechanism enforces a consistent memory layout. Together, these properties allow Foundry to capture a graph on a single GPU, then share it across all participating GPUs. This reduces both the hardware cost of SAVE and the archive footprint from cluster-scale to single-device. In practice, operators often maintain archives for multiple parallelism configurations to meet varying latency and throughput targets, incurring even higher cost. Because a single GPU cannot replay the full collective communication operation without peers, we cannot run the original communication kernels during capture. However, our insight is that the communication kernels rely on common building blocks, such as NCCL and NVSHMEM. Therefore, we can build a stub over these common building blocks and perform a dummy communication during capture. During LOAD, Foundry substitutes the handle with the real kernel function and updates the rank-specific arguments. Inter-rank graph sharing applies whenever all ranks follow the same execution flow and thus produce equivalent graph structure during capture. Pipeline parallelism (PP) is not included, since different ranks execute different model stages and therefore require stage-specific graph capture. Moreover, PP introduces inter-stage bubbles that increase per-request latency during decoding [1], making it less suitable for latency-sensitive serving workloads. Foundry can also support hybrid strategies (e.g., EP combined with TP) as long as the resulting combination preserves topology invariance across ranks and does not introduce rank-specific execution divergence.
5
Implementation
Foundry is implemented as two components: (i) a driver-hook library, a C/C++ interposition layer injected via LD_PRELOAD that enforces deterministic memory allocation and captures kernel binaries, and (ii) a graph extension to PyTorch, a C++ library with Python bindings that serializes, reconstructs, and replays CUDA graphs with template-based sharing. A thin Python integration layer wires these components into the vLLM [23] to demonstrate LLM serving. The driver-hook library and graph extension together comprise approximately 10K lines of C++ and 2K lines of Python.
3 Graph object thread safety, https://docs.nvidia.com/cuda/archive/12.8.1/
cuda-driver-api/graphs-thread-safety.html. 7
5.1
Memory Allocation Redirection
parsing time of 512 graphs from a few seconds to less than 100 milliseconds.
The deterministic allocation mechanism (§4.1.1) is implemented by redirecting CUDA memory allocation to a virtual memory management (VMM) backend. VMM can map a physical memory buffer onto a specified virtual address, as long as it is not occupied. Foundry reserves a large virtual address space at the beginning. For each following allocation, Foundry leverages VMM to assign the virtual address right after the last one, and results in a deterministic and contiguous memory region. Because Foundry tracks all incoming allocation calls, it can record and replay a selected range of memory events as needed.
5.4
Because CUDA graphs are mainly designed for LLM decoding to eliminate the bubbles caused by host-side delay, our optimization target is the decoder of prefill-decode disaggregated serving [54]. To demonstrate multi-GPU capability, Foundry is evaluated on expert-parallel (EP) serving, which is the state-of-the-art LLM deployment strategy for frontier mixture-of-expert (MoE) models [11]. Besides the specialized communication kernels in DeepEP, Foundry seamlessly supports FP8 expert computation (DeepGEMM), thanks to its library-agnostic design. We integrate Foundry into vLLM to accelerate the engine initialization, by implementing a thin layer under its compilation model wrapper to apply our graph extension, and preload our hook layer when each worker process is created. The kv cache size of vLLM can change because it is set via the utilization percentage of currently available memory. We specify the kv cache size before SAVE or LOAD so that the memory layout is consistent across different runs. Because memory is preallocated at the beginning of LOAD (§4.1.1), the graph construction process works asynchronously from the foreground tasks such as kv cache initialization.
5.2 CUDA Module Load Interception The extraction of kernel binaries (§4.1.2) is implemented by intercepting module load driver APIs. CUmodules, which contain kernel handles used by CUDA graphs, are loaded via cuModuleLoad and its series of APIs. During SAVE, Foundry works as a wrapper to track the data pointer passed into the API and extracts its raw binary payload from the memory and computes a hash. Foundry also enumerates through all the kernel handle entrypoints of loaded modules and records a catalog along with the binary payload. Therefore, during LOAD, it can directly resolve the kernel handle from the kernel function name, instead of searching through all the loaded modules. Some modules such as DeepEP use NVSHMEM that requires device-side state initialization after loading. Foundry detects that during SAVE and sets a flag, and calls nvshmemx_cumodule_init 4 on it during LOAD. 5.3
LLM Serving Engine Integration
6
Evaluation
We evaluate Foundry along four dimensions: (1) reduction in cold-start latency compared to vanilla vLLM and CUDAcheckpoint; (2) preservation of serving throughput; (3) effectiveness of template-based graph reconstruction; and (4) storage cost saving.
Graph Grouping and Serialization
Foundry groups hundreds of graphs into a few templates to save time on reconstruction. The on-demand graph update is implemented via cuGraphExecUpdate, which can update node parameters, such as kernel function, argument buffer, and launch dimensions, without re-instantiating the graph executable object. It requires the new graph to have the same topology as the template, including their node types, order, and dependencies between nodes. According to our test, kernel node attributes (e.g., cluster dimensions) set via cuGraphKernelNodeSetAttribute, also remain stale, and thus it is treated as a part of graph topology. Foundry computes a key based on topology and organizes graphs with the same key into a group. Across different model architectures and sizes, graph grouping effectively reduces the number of graphs to build during LOAD (§6.4). We initially saved each cuda graph structure into a JSON file for better readability. However, as model size increases, the parsing incurs significant delay. As a result, Foundry also serializes a binary version, which effectively reduces the
6.1
Experimental Setup
Testbed Configurations. All experiments run on NVIDIA DGX nodes. The primary platform uses 8×H200 GPUs with Intel Xeon Platinum 8480C CPUs and 2 TB host memory. We also evaluate on a second node with 8×B200 GPUs and Intel Xeon Platinum 8570 CPUs to demonstrate Foundry performance on the latest architecture. All GPUs in a single instance are fully connected with NVLink. Both nodes run NVIDIA driver 590.48 with CUDA 13.1. We use vLLM v0.11.2, PyTorch 2.9 and NVSHMEM 3.3.24 in our experiment environment. For all experiments, we load the warmed cache of torch.compile and archive of Foundry from local NVMe storage. Models and parallelism. We evaluate across three model families covering a wide range of architectures and sizes: Qwen3-14B [47], Qwen3-32B, Llama3-8B [17] and Gemma312B [21] are dense models, while Qwen3-30B-A3B and Qwen3235B-A22B are mixture-of-experts (MoE) models. Dense models are deployed with data parallelism (DP1–DP8); MoE models with expert parallelism (EP2–EP8), in both BF16 and FP8
4 NVSHMEM
APIs, https://docs.nvidia.com/nvshmem/api/gen/api/ setup.html#nvshmemx-cumodule-init 8
vLLM w/o graph
Time (s)
103 102
95%
101 1.5
95%
95%
1.7 2.4
5.0
3.3
1.7
1.8
96% 1.7
2.8
96%
96%
2.8 3.5
2.8
4.5
Foundry
154
121
112
83
74
72
48
43
41
36
vLLM with graph
96%
97%
97%12
98%
2.9
2.7
2.9
2.8
9.3
7.8
62
97%
5.5
650
146
136
3.5
6.9
97%
99% 28
3.5
3.9
DP1
DP2
DP4
DP8
EP2-FP8 EP4-FP8 EP8-FP8
Qwen3-14B
EP2
Qwen3-30B-A3B
EP4
EP8
95% 1.2
100 EP4-FP8 EP8-FP8
EP8
Qwen3-235B-A22B
45 11
1.3
95% 2.0
Llama3 Gemma3 8B 12B
Figure 7. [Cold-Start Latency Reduction] Serving engine initialization latency across models and parallelism configurations on H200. Percentages indicate Foundry’s reduction relative to vLLM with CUDA graphs. torch.compile Memory Profile Graph Capture vLLM CUDA-ckpt Foundry
Binary Load Graph Load Toggle Back
8.5
36.1s Qwen3-14B
2.6
25
5.7s
5.2
full performance, we let vLLM releases model weights and the KV cache before checkpointing to avoid data loading bottlenecks. We do not include Medusa [51] because its per-kernel patching rules are implemented for a specific set of kernels and are based on an outdated CUDA driver version (550.54.14). Porting them to newer hardware and kernel libraries would require reimplementing the system. On Hopper and later GPUs, cuBLAS kernels use large argument buffers (Appendix B) with opaque layouts, making such patching impractical.
Wake Up Sampler Warmup
1.0 1.7s
vLLM CUDA-ckpt Foundry
13 1.6
vLLM CUDA-ckpt Foundry
42 59.4s Qwen3-32B
3.9
6.1s
5.4
2.3s 11
28 44.4s Gemma3-12B
5.4
5.9s
5.4 1.1 2.0s
vLLM CUDA-ckpt Foundry
5.5
2.2
5.3
5.7s
20
27.9s Llama3-8B
Metrics. We use cold start time to evaluate Foundry’s ability to reduce initialization latency and use time per output token (TPOT) to measure the serving throughput, as we target the decoding phase of PD-disaggregated serving.
1.3s
0
5
10
Time (s)
15
25 50
6.2 Figure 8. [Cold-Start Latency Reduction] End-to-end engine initialization breakdown comparing vLLM, CUDAcheckpoint, and Foundry across four model configurations. Note: torch.compile only does bytecode transformation and cache loading.
Cold Start Latency Reduction
6.2.1 Comparison with vLLM. Figure 7 compares engine initialization latency across 15 configurations spanning three model families on H200, including both dense and MoE models. For each configuration, we report the cold start latency of three modes: vLLM without CUDA graphs, vLLM with CUDA graphs, and Foundry. Following Medusa [51], for all compared methods, we assume a pool of warm execution environments that eliminate the environment initialization overhead. We also omit weight loading time, as recent systems [3, 48, 52] have reduced it to 1~2s.
precision. All configurations capture 512 CUDA graphs covering batch sizes 1–512 to simulate real world deployment. Baselines. We compare against two baselines: (1) vLLM with CUDA graphs, the default production implementation that performs full warmup and stream capture during initialization; and (2) CUDA-checkpoint [10], which snapshots the entire CUDA context via NVIDIA’s process-level checkpoint/restore API, used in conjunction with CRIU [8]. We also report vLLM without CUDA graphs (eager mode) as a reference for the minimum achievable startup time at the cost of degraded serving performance (§2). CUDA-checkpoint does not support the IPC memory required by communication kernels (e.g., DeepEP) [10], and its restore latency grows disproportionately for multi-GPU data-parallel engines, making it less efficient than launching multiple independent single-GPU instances. We therefore compare with it only on single-GPU settings. To assess its
Dense models. For Qwen3-14B, vLLM with graphs takes 36–48 s depending on the parallelism degree (DP1–DP8). Comparing with vLLM without graphs, graph capture clearly contributes the majority of the cost. Foundry reduces this to 1.7–1.8 s, a 95% reduction that is consistent across all DP configurations. Notably, Foundry’s initialization time remains nearly constant as the number of ranks scales from 1 to 8. This is because both execution context restoration and CUDA graph reconstruction run fully independently, and there is no computation on GPUs. The reduction is generalizable to different model architectures, for example, on Llama3-8B and Gemma3-12B, the reduction is 95% (28 s→1.3 s) and 95% (45 s→2.0 s), respectively. 9
Mean TPOT (ms)
Qwen3-14B (H200) vLLM Foundry
30 20
Qwen3-30B-A3B (H200) 100
DP1
DP2
200
EP2
50 20
40
EP2-FP8
16 32 64 128 256 512
#Prompts
#Prompts
Mixed Models (B200) 200 100 50
235B-EP4 30B-EP4
25 EP4-FP8
30
10 16 32 64 128 256 512
EP8
100 50
DP4
10
Qwen3-235B-A22B (H200)
16 32 64 128 256 512
#Prompts
10.0 7.5
14B-DP4
16 32 64 128 256 512
#Prompts
Avg time per graph (ms)
Figure 9. [Serving Throughput Preservation] Mean time per output token (TPOT) for vLLM and Foundry across batch sizes 16–512. Graph Capture
Template Build
latency of Foundry also increases compared with small models as each graph contains more nodes, but it remains fast as it only builds a few templates and graph reconstruction is much faster than stream capture (Figure 10).
On-Demand Update
198.6
200
2.9x 69.5
100 59.7 1.9x 31.1
50
32x 0.98
0
MoE models with FP8 quantization. By applying FP8 quantization, the vLLM initialization time of both Qwen330B-A3B and Qwen3-235B-A22B decreases, which take 72– 83 s across EP2–EP8 and 136–146 s across EP4–8, respectively. This is because the DeepGEMM kernels greatly accelerate FP8 expert computation, compared with vLLM default Triton implementation on BF16. The LOAD time of Foundry barely changes compared with BF16, as the graph topology is mostly the same, switching kernel functions. Yet, Foundry still consistently achieves 96–97% latency reduction, and completes in 2.8–2.9s on Qwen3-30B-A3B and 3.5s on Qwen3-235BA22B.
24x
DP4
2.89 EP4
Qwen3-14B
Qwen3-30B-A3B
Figure 10. [Effectiveness of Templating] Average pergraph cost of stream capture, template construction via driver APIs, and on-demand parameter update. All experiments capture 512 graphs. Templates
#Graphs
512 384
96%
97%
96%
On-demand Updated 95%
98%
97%
96%
97%
Comparison with initialization without graphs. An alternative to accelerate cold start is to skip graph capture entirely and run in eager mode. The serving engine still needs to load kernels from torch.compile cache, but it is much faster than capturing hundreds of graphs. Figure 7 shows that vLLM without graphs achieves lower initialization latency (1.2–62 s) but, as shown in §2, it significantly degrades decoding performance, while Foundry delivers the full performance of CUDA graphs (Figure 9). Foundry achieves startup times comparable to or faster than eager mode, because all required binaries are already packed into a condensed archive with clear mappings to kernel functions. For example, Qwen3-14B DP8 starts in 5.0 s without graphs versus 1.7 s with Foundry; Qwen3-235B-A22B EP8 starts in 62 s without graphs versus 3.9 s with Foundry.
256 128 0
25 22 19 19 17 15 15 12 Qwen Qwen3 Qwen3 Qwen Qwen3 Qwen3 Llama3Gemma3 14B 30B-FP8 30B 32B 235B-FP8 235B 8B 12B
Figure 11. [Effectiveness of Templating] Number of unique templates versus on-demand updated graphs out of 512 total. Percentages indicate the fraction of graphs served via on-demand update. MoE models with default BF16. Foundry exhibits significant latency reduction on large scale MoE models. Qwen330B-A3B (EP2–EP8) takes 112–154 s to initialize on vLLM with graphs, reduced to 2.7–2.8 s by Foundry (97–98%). The most dramatic case is Qwen3-235B-A22B EP8, where vLLM requires 650 s (about 10 minutes) for initialization, of which graph capture alone accounts for the majority. Foundry reduces this to 3.9 s, a 99% reduction. This extreme speedup arises because graph capture requires model forward on every supporting batch size, and a larger model tends to incur longer forward duration due to computation complexity. The
6.2.2 Comparison with CUDA Checkpoint. Figure 8 provides a detailed phase-level breakdown for four configurations, comparing vLLM, CUDA-checkpoint, and Foundry. CUDA-checkpoint eliminates warmup and graph capture by toggling back snapshotted GPU states. Using vLLM’s sleep and wakeup mechanism, the vLLM model weights and KV cache are released before checkpointing to avoid storing the whole GPU memory (141GB per H200). 10
6.4
Across different models, CUDA-checkpoint consistently restore serving engine in 5.7–6.1 s, achieving 4.9–7.9x faster initialization compared with vLLM. However, Foundry achieves 1.3–2.3 s on these configurations, consistently outperforming CUDA-checkpoint by 2.6–4.4×. The breakdown reveals that the toggle back is slower than the total time of loading binaries and graph in Foundry. This is because Foundry only saves the necessary CUDA states that are most time-consuming to capture, while creating others at runtime, whereas CUDAcheckpoint saves the entire CUDA state, incurring not only higher latency but also larger image size (Table 1).
6.3
Effectiveness of Templating
6.4.1 Topology Sharing Across Models. Figure 11 reports the number of unique topologies (templates) identified by Foundry for each model across 512 captured graphs. The number of templates ranges from 12 (Qwen3-235B-A22B FP8) to 25 (Qwen3-32B), yielding 95–98% of graphs served via on-demand parameter update rather than full construction. The fraction of templates is consistently low across model families we tested (Qwen3, Llama3 and Gemma3), indicating that the topological regularity exploited by Foundry is a fundamental property of LLM inference, not specific to any architecture variant. The low template count means the graphs captured for different batch sizes execute the same sequence of kernel types with the same dependency structure, yet they can execute different kernels tuned for different matrix sizes, and with varied launch dimensions and arguments. Notably, the template count does not grow with model size or complexity: the largest model, Qwen3-235B-A22B (FP8), requires only 12 templates, while the smaller Qwen332B requires 25. Besides, the template count is independent of parallelism degree for SPMD-style multi-GPU serving, as different ranks share the same execution flow as each works individually. The changes to matrix/vector size and rank identifier are covered by using different kernel handles, launch dimensions or arguments.
Serving Throughput Preservation
Fast initialization is only useful if the reconstructed CUDA graphs produce the same serving performance as natively captured ones. As Foundry targets the decoding phase of PD-disaggregated serving, we report TPOT to assess the throughput of the LLM serving. This value also reflects graph replay duration, which would be the same if the loaded graph is equivalent to the captured one. Figure 9 compares mean TPOT between vLLM and Foundry for representative configurations on both H200 and B200. We use the serving benchmark provided by vLLM to simulate incoming requests and test across 16–512 random prompts 10 times, with each prompt generating 128 tokens.
6.4.2 Per-Graph Construction Cost. Figure 10 compares the average per-graph cost of three construction methods: native stream capture, template build via CUDA driver APIs, and on-demand parameter update of an existing template. For Qwen3-14B DP4 and Qwen3-30B-A3B EP4, we observe that graph construction (31.1–69.5 ms) is 1.9–2.9× faster than stream graph capture (59.7–198.6 ms). However, because parallel graph construction results in driver contention (Appendix A), sequentially building 512 graphs for Qwen3-30B-A3B still incurs 512 × 69.5 ≈ 35.6 s. This is not acceptable for elastic serving and it worsens for higher batch sizes. In-place graph update (0.98-2.89 ms) further reduces the latency by 24–32× compared with graph construction, making it possible to execute on demand during serving. The on-demand update is fast because its required parameters have been organized during template construction (§4.2.1), loading them and calling cuGraphExecUpdate to update node parameters in-place is much faster than reinstantiating a graph executable. As a result, the engine only needs to build very small number of graphs on initialization, effectively compresses initialization latency to a few seconds.
No observable throughput degradation. Across all settings, the TPOT curves for Foundry and vLLM with natively captured graphs overlap almost perfectly. The first subplot tests across DP1–4 on Qwen3-14B, confirming that Foundry preserves serving performance for different data parallelism sizes. The second subplot further verifies that the performance holds when applying optimized custom kernels, such as DeepGEMM. Comparing the second and third subplots, we observe that for a large model that requires sharding to multiple GPUs, Foundry still preserves same throughput as vLLM. And the fourth subplot confirms that this preservation generalizes to GPUs across different architectures. This result is expected by Foundry’s library-agnostic design: its execution context materialization ensures that the restored engine contains the exact same kernel binaries, memory layout, regardless of what model is being served. This also confirms that the template-based reconstruction produces graph executables that are semantically equivalent to those produced by stream capture. The on-demand update to the template is very fast (Figure 10) and is shared by multiple engine steps due to continuous batching, resulting in no observable overhead. We also compare the tokens generated by Foundry with vLLM and find that they are identical, confirming the correctness of CUDA graph and execution context restoration.
6.4.3 Storage Cost Saving. Table 1 compares the storage cost between CUDA-checkpoint and Foundry. For CUDAcheckpoint, model weights and KV cache are released before checkpointing. 11
Table 1. Storage cost comparison between CUDAcheckpoint and Foundry. L3, G3 and Q3 indicate Llama3, Gemma3 and Qwen3 respectively. Model Parallel Image size Archive size
L3-8B G3-12B DP1 DP1 3.9 GB 6.6 GB 976 MB 1.3 GB
Q3-14B DP1 3.7 GB 1.1 GB
to different requests; Gyges [9] and Flying Serving [15] perform parallelism transformation to switch running instances across parallelism strategies as request context lengths vary. Elastic expert parallelism (EP) [42] and Expert-as-a-Service (EaaS) [25] have recently been proposed to enable fine-grained scaling and improve fault tolerance for serving MoE models. A common challenge across these systems is that changing the parallelism strategy also changes the computation logic, which requires recapturing CUDA graphs and thus incurs substantial reconfiguration overhead.
Q3-235B EP8 – 2.2 GB
Foundry’s archive is 4–5× smaller than CUDA-checkpoint’s image. For example, on Qwen3-14B, the archive is 1.1 GB versus 3.7 GB for CUDA-checkpoint (3.4× smaller). This is because CUDA-checkpoint captures the entire process state, whereas Foundry stores only graph metadata and kernel binaries, which are compact and rank-independent. Because Foundry allows different ranks to share the same set of kernel binaries and graph metadata, even for Qwen3-235B-A22B EP8, the archive size is just 2.2 GB. While CUDA-checkpoint doesn’t even support expert parallelism, Foundry creates its archive with one GPU.
7.3
GPU checkpoint/restore (C/R) approaches have been extensively studied. Existing system-level GPU C/R techniques can be broadly divided into driver-integrated and interceptionbased approaches. Driver-integrated C/R is vendor-specific; NVIDIA has recently added such support to its proprietary CUDA driver [10, 39], but the current implementation still lacks support for IPC memory and therefore cannot be directly applied to multi-GPU distributed inference. On the other hand, interception-based systems [7, 12, 16, 20, 28, 41] record CUDA resource state during normal execution and reconstruct it during restore via CUDA driver API replay. Recent systems have substantially improved the efficiency of this line of work: PhOS [45] uses validated speculation to enable concurrent GPU C/R, while GCR [50] uses control/data separation to reduce C/R latency and runtime overhead. However, because these systems restore CUDA process state through API replay, they do not directly support restoring captured CUDA graphs and are therefore complementary to Foundry.
Composition of archive. The archive consists of two parts: serialized graph metadata (topology, node parameters, and grouping manifest) and extracted kernel binaries. For Qwen3-235B-A22B EP8, the largest model evaluated, the total archive is 2.2 GB, of which kernel binaries account for 1.4 GB. The binary graph format (§5.3) keeps the graph metadata compact, and Foundry parses 512 serialized graphs in under 100 ms.
7
Related Work
7.1
LLM Cold Start Optimization
8
A growing body of work has sought to reduce cold-start latency in elastic and serverless LLM serving. Many of these efforts focus on model loading. ServerlessLLM [13] caches model weights on local SSDs, while BlitzServe [52], 𝜆Scale [48], and Tensor R-Fork [3] leverage RDMA to fetch weights from peer instances. HydraServe [26] further overlaps model fetching with other stages of cold start, whereas WarmServe [27] and InstaInfer [40] proactively reduce startup latency by predicting future workloads and preparing resources in advance. To mitigate environment initialization overhead, DeepServe [19] uses pre-warmed pods, following a broader line of work on container pre-warming [4– 6, 14, 24, 38]. Medusa, in contrast, targets CUDA graph capture overhead. However, unlike Foundry, it materializes only graph topology and does not preserve the execution context required for general graph restoration. 7.2
GPU Checkpoint/Restore
Conclusion
Modern LLM service providers increasingly rely on autoscaling to improve GPU resource efficiency, yet CUDA graph capture keeps LLM serving cold starts in the tens of seconds to minutes. We presented Foundry, a template-based CUDA graph context materialization system that removes this bottleneck by persisting not only graph topology but also the execution context required for replay. By combining deterministic memory layouts, kernel binary materialization, topology-based templating, and single-GPU offline capture for multi-GPU distributed inference, Foundry enables efficient, kernel-agnostic CUDA graph restoration without hand-crafted patching rules or heavyweight process-level checkpointing. Integrated with vLLM, Foundry reduces coldstart latency by up to 99%.
Parallelism Hot Switching in LLM Serving
Recent work has begun to explore dynamic parallelism reconfiguration in LLM serving. LoongServe [46] introduces elastic sequence parallelism to adapt the degree of parallelism 12
References
models. arXiv preprint arXiv:2407.21783 (2024). [18] Naibin Gu, Zhenyu Zhang, Yuchen Feng, Yilong Chen, Peng Fu, Zheng Lin, Shuohuan Wang, Yu Sun, Hua Wu, Weiping Wang, et al. 2025. Elastic MoE: Unlocking the Inference-Time Scalability of Mixture-ofExperts. arXiv preprint arXiv:2509.21892 (2025). [19] Junhao Hu, Jiang Xu, Zhixia Liu, Yulong He, Yuetao Chen, Hao Xu, Jiang Liu, Jie Meng, Baoquan Zhang, Shining Wan, Gengyuan Dan, Zhiyu Dong, Zhihao Ren, Changhong Liu, Tao Xie, Dayun Lin, Qin Zhang, Yue Yu, Hao Feng, Xusheng Chen, and Yizhou Shan. 2025. DEEPSERVE: Serverless Large Language Model Serving at Scale. In 2025 USENIX Annual Technical Conference (USENIX ATC ’25). https: //www.usenix.org/conference/atc25/presentation/hu-junhao [20] Twinkle Jain and Gene Cooperman. 2020. Crac: Checkpoint-restart architecture for cuda with streams and uvm. In SC20: International Conference for High Performance Computing, Networking, Storage and Analysis. IEEE, 1–15. [21] Aishwarya Kamath, Johan Ferret, Shreya Pathak, Nino Vieillard, Ramona Merhej, Sarah Perrin, Tatiana Matejovicova, Alexandre Ramé, Morgane Rivière, Louis Rouillard, et al. 2025. Gemma 3 technical report. arXiv preprint arXiv:2503.19786 4 (2025). [22] Michael Kerrisk. [n. d.]. proc_sys_kernel(5) — Linux manual page. man7.org. https://man7.org/linux/ man-pages/man5/proc_sys_kernel.5.html Section: /proc/sys/kernel/randomize_va_space, accessed 2026-04-01. [23] Woosuk Kwon, Zhuohan Li, Siyuan Zhuang, Ying Sheng, Lianmin Zheng, Cody Hao Yu, Joseph E. Gonzalez, Hao Zhang, and Ion Stoica. 2023. Efficient Memory Management for Large Language Model Serving with PagedAttention. In Proceedings of the 29th Symposium on Operating Systems Principles (SOSP ’23). doi:10.1145/3600006.3613165 [24] Zhen Lin, Kao-Feng Hsieh, Yu Sun, Seunghee Shin, and Hui Lu. 2021. Flashcube: Fast provisioning of serverless functions with streamlined container runtimes. In Proceedings of the 11th Workshop on Programming Languages and Operating Systems. 38–45. [25] Ziming Liu, Boyu Tian, Guoteng Wang, Zhen Jiang, Peng Sun, Zhenhua Han, Tian Tang, Xiaohe Hu, Yanmin Jia, Yan Zhang, et al. 2025. Expert-as-a-service: Towards efficient, scalable, and robust large-scale moe serving. arXiv preprint arXiv:2509.17863 (2025). [26] Chiheng Lou, Sheng Qi, Chao Jin, Dapeng Nie, Haoran Yang, Yu Ding, Xuanzhe Liu, and Xin Jin. 2026. HydraServe: Minimizing Cold Start Latency for Serverless LLM Serving in Public Clouds. In Proceedings of the 23rd USENIX Symposium on Networked Systems Design and Implementation (NSDI ’26). https://www.usenix.org/conference/nsdi26/ presentation/lou [27] Chiheng Lou, Sheng Qi, Rui Kang, Yong Zhang, Chen Sun, Pengcheng Wang, Bingyang Liu, Xuanzhe Liu, and Xin Jin. 2025. WarmServe: Enabling One-for-Many GPU Prewarming for Multi-LLM Serving. arXiv preprint arXiv:2512.09472 (2025). [28] Akira Nukada, Taichiro Suzuki, and Satoshi Matsuoka. 2023. Efficient checkpoint/Restart of CUDA applications. Parallel Comput. 116 (2023), 103018. [29] NVIDIA 2026. CUDA Driver API: Graph Management. NVIDIA. https://docs.nvidia.com/cuda/cuda-driverapi/group__CUDA__GRAPH.html Accessed March 5, 2026. [30] NVIDIA. 2026. NCCL: Optimized primitives for inter-GPU communication. https://github.com/NVIDIA/nccl [31] NVIDIA. 2026. NVSHMEM: The NVIDIA SHMEM library for GPU clusters. https://github.com/NVIDIA/nvshmem [32] OpenAI. [n. d.]. Models | OpenAI API. https://developers.openai.com/ api/docs/models [33] Perplexity. 2025. Weight Transfer for RL Post-Training in under 2 seconds. https://research.perplexity.ai/articles/weight-transfer-forrl-post-training-in-under-2-seconds
[1] Amey Agrawal, Nitin Kedia, Ashish Panwar, Jayashree Mohan, Nipun Kwatra, Bhargav Gulavani, Alexey Tumanov, and Ramachandran Ramjee. 2024. Taming Throughput-Latency Tradeoff in LLM Inference with Sarathi-Serve. In Proceedings of the 18th USENIX Symposium on Operating Systems Design and Implementation (OSDI ’24). 117–134. https://www.usenix.org/conference/osdi24/presentation/agrawal [2] Amazon Web Services. [n. d.]. Amazon Bedrock. https:// aws.amazon.com/bedrock/ [3] Ant Group DeepXPU Team and SGLang Team. 2025. Let Tensors Fly — Accelerating Large Model Weight Loading with R-Fork. https: //lmsys.org/blog/2025-12-10-rfork/ [4] Vivek M Bhasi, Jashwant Raj Gunasekaran, Prashanth Thinakaran, Cyan Subhra Mishra, Mahmut Taylan Kandemir, and Chita Das. 2021. Kraken: Adaptive container provisioning for deploying dynamic dags in serverless platforms. In Proceedings of the ACM Symposium on Cloud Computing. 153–167. [5] Marc Brooker, Mike Danilov, Chris Greenwood, and Phil Piwonka. 2023. On-demand container loading in {AWS} lambda. In 2023 USENIX Annual Technical Conference (USENIX ATC 23). 315–328. [6] James Cadden, Thomas Unger, Yara Awad, Han Dong, Orran Krieger, and Jonathan Appavoo. 2020. SEUSS: skip redundant paths to make serverless fast. In Proceedings of the Fifteenth European Conference on Computer Systems. 1–15. [7] Shubham Chaudhary, Ramachandran Ramjee, Muthian Sivathanu, Nipun Kwatra, and Srinidhi Viswanatha. 2020. Balancing efficiency and fairness in heterogeneous GPU clusters for deep learning. In Proceedings of the Fifteenth European Conference on Computer Systems. 1–16. [8] checkpoint-restore. [n. d.]. CRIU: Checkpoint/Restore tool. https: //github.com/checkpoint-restore/criu. GitHub repository, accessed April 1, 2026. [9] Haoyu Chen, Xue Li, Kun Qian, Yu Guan, Jin Zhao, and Xin Wang. 2025. Gyges: Dynamic Cross-Instance Parallelism Transformation for Efficient LLM Inference. arXiv preprint arXiv:2509.19729 (2025). [10] NVIDIA Corporation. 2025. cuda-checkpoint: CUDA Checkpoint and Restore Utility. https://github.com/NVIDIA/cuda-checkpoint. Accessed: 2026-03-22. [11] DeepSeek-AI. 2025. Open Infra Index. GitHub repository. https: //github.com/deepseek-ai/open-infra-index Day 6: One More Thing – DeepSeek-V3/R1 Inference System Overview. [12] Niklas Eiling, Jonas Baude, Stefan Lankes, and Antonello Monti. 2022. Cricket: A virtualization layer for distributed execution of CUDA applications with checkpoint/restart support. Concurrency and Computation: Practice and Experience 34, 14 (2022), e6474. [13] Yao Fu, Leyang Xue, Yeqi Huang, Andrei-Octavian Brabete, Dmitrii Ustiugov, Yuvraj Patel, and Luo Mai. 2024. {ServerlessLLM}:{LowLatency} serverless inference for large language models. In 18th USENIX Symposium on Operating Systems Design and Implementation (OSDI 24). 135–153. [14] Alexander Fuerst and Prateek Sharma. 2021. Faascache: keeping serverless computing alive with greedy-dual caching. In Proceedings of the 26th ACM international conference on architectural support for programming languages and operating systems. 386–400. [15] Shouwei Gao, Junqi Yin, Feiyi Wang, and Wenqian Dong. 2026. FLYING SERVING: On-the-Fly Parallelism Switching for Large Language Model Serving. arXiv preprint arXiv:2602.22593 (2026). [16] Rohan Garg, Apoorve Mohan, Michael Sullivan, and Gene Cooperman. 2018. CRUM: Checkpoint-restart support for CUDA’s unified memory. In 2018 IEEE International Conference on Cluster Computing (CLUSTER). IEEE, 302–313. [17] Aaron Grattafiori, Abhimanyu Dubey, Abhinav Jauhri, Abhinav Pandey, Abhishek Kadian, Ahmad Al-Dahle, Aiesha Letman, Akhil Mathur, Alan Schelten, Alex Vaughan, et al. 2024. The llama 3 herd of 13
[34] Ramya Prabhu, Ajay Nayak, Jayashree Mohan, Ramachandran Ramjee, and Ashish Panwar. 2025. vAttention: Dynamic Memory Management for Serving LLMs without PagedAttention. In Proceedings of the 30th ACM International Conference on Architectural Support for Programming Languages and Operating Systems (ASPLOS ’25). 1133–1150. doi:10.1145/3669940.3707256 [35] PyTorch Contributors. [n. d.]. CUDA semantics. https:// docs.pytorch.org/docs/stable/notes/cuda.html. Accessed: 2026-04-01. [36] Yifan Qiao, Shu Anzai, Shan Yu, Haoran Ma, Shuo Yang, Yang Wang, Miryung Kim, Yongji Wu, Yang Zhou, Jiarong Xing, et al. 2024. ConServe: Fine-Grained GPU Harvesting for LLM Online and Offline Co-Serving. arXiv preprint arXiv:2410.01228 (2024). [37] Qwen Team. 2025. Qwen3-14B. https://huggingface.co/Qwen/Qwen314B. Hugging Face model card, accessed 2026-04-01. [38] Rohan Basu Roy, Tirthak Patel, and Devesh Tiwari. 2022. Icebreaker: Warming serverless functions better with heterogeneity. In Proceedings of the 27th ACM International Conference on Architectural Support for Programming Languages and Operating Systems. 753–767. [39] Radostin Stoyanov, Viktória Spišaková, Jesus Ramos, Steven Gurfinkel, Andrei Vagin, Adrian Reber, Wesley Armour, and Rodrigo Bruno. 2025. Criugpu: Transparent checkpointing of gpu-accelerated workloads. arXiv preprint arXiv:2502.16631 (2025). [40] Yifan Sui, Hanfei Yu, Yitao Hu, Jianxun Li, and Hao Wang. 2024. Prewarming is not enough: Accelerating serverless inference with opportunistic pre-loading. In Proceedings of the 2024 ACM Symposium on Cloud Computing. 178–195. [41] Hiroyuki Takizawa, Katsuto Sato, Kazuhiko Komatsu, and Hiroaki Kobayashi. 2009. CheCUDA: A checkpoint/restart tool for CUDA applications. In 2009 International Conference on Parallel and Distributed Computing, Applications and Technologies. IEEE, 408–413. Elastic EP in [42] The Mooncake Team, Volcano Engine. 2026. SGLang: Achieving Partial Failure Tolerance for DeepSeek MoE Deployments. https://www.lmsys.org/blog/2026-03-25-eep-partial-failuretolerance/ [43] vLLM Project. [n. d.]. CUDA Graphs - vLLM. https://docs.vllm.ai/en/ latest/design/cuda_graphs/ [44] Yuxin Wang, Yuhan Chen, Zeyu Li, Xueze Kang, Yuchu Fang, Yeju Zhou, Yang Zheng, Zhenheng Tang, Xin He, Rui Guo, et al. 2025. Burstgpt: A real-world workload dataset to optimize llm serving systems. In Proceedings of the 31st ACM SIGKDD Conference on Knowledge Discovery and Data Mining V. 2. 5831–5841. [45] Xingda Wei, Zhuobin Huang, Tianle Sun, Yingyi Hao, Rong Chen, Mingcong Han, Jinyu Gu, and Haibo Chen. 2025. PhoenixOS: Concurrent OS-level GPU Checkpoint and Restore with Validated Speculation. In Proceedings of the 31st Symposium on Operating Systems Principles (SOSP ’25). doi:10.1145/3731569.3764813 [46] Bingyang Wu, Shengyu Liu, Yinmin Zhong, Peng Sun, Xuanzhe Liu, and Xin Jin. 2024. Loongserve: Efficiently serving long-context large language models with elastic sequence parallelism. In Proceedings of the ACM SIGOPS 30th Symposium on Operating Systems Principles. 640–654. [47] An Yang, Anfeng Li, Baosong Yang, Beichen Zhang, Binyuan Hui, Bo Zheng, Bowen Yu, Chang Gao, Chengen Huang, Chenxu Lv, et al. 2025. Qwen3 technical report. arXiv preprint arXiv:2505.09388 (2025). [48] Minchen Yu, Rui Yang, Chaobo Jia, Zhaoyuan Su, Sheng Yao, Tingfeng Lan, Yuchen Yang, Yue Cheng, Wei Wang, Ao Wang, et al. 2025. {\lambda} Scale: Enabling Fast Scaling for Serverless Large Language Model Inference. arXiv preprint arXiv:2502.09922 (2025). [49] Shan Yu, Jiarong Xing, Yifan Qiao, Mingyuan Ma, Yangmin Li, Yang Wang, Shuo Yang, Zhiqiang Xie, Shiyi Cao, Ke Bao, et al. 2025. Prism: Unleashing gpu sharing for cost-efficient multi-llm serving. arXiv preprint arXiv:2505.04021 (2025). [50] Shaoxun Zeng, Tingxu Ren, Jiwu Shu, and Youyou Lu. 2026. {GPU} {Checkpoint/Restore} Made Fast and Lightweight. In 24th
USENIX Conference on File and Storage Technologies (FAST 26). 239– 254. [51] Shaoxun Zeng, Minhui Xie, Shiwei Gao, Youmin Chen, and Youyou Lu. 2025. Medusa: Accelerating Serverless LLM Inference with Materialization. In Proceedings of the 30th ACM International Conference on Architectural Support for Programming Languages and Operating Systems (ASPLOS ’25). 653–668. doi:10.1145/3669940.3707285 [52] Dingyan Zhang, Haotian Wang, Yang Liu, Xingda Wei, Yizhou Shan, Rong Chen, and Haibo Chen. 2025. {BlitzScale}: Fast and Live Large Model Autoscaling with O (1) Host Caching. In 19th USENIX Symposium on Operating Systems Design and Implementation (OSDI 25). 275–293. [53] Chenggang Zhao, Chengqi Deng, Chong Ruan, Damai Dai, Huazuo Gao, Jiashi Li, Liyue Zhang, Panpan Huang, Shangyan Zhou, Shirong Ma, et al. 2025. Insights into deepseek-v3: Scaling challenges and reflections on hardware for ai architectures. In Proceedings of the 52nd Annual International Symposium on Computer Architecture. 1731–1745. [54] Yinmin Zhong, Shengyu Liu, Junda Chen, Jianbo Hu, Yibo Zhu, Xuanzhe Liu, Xin Jin, and Hao Zhang. 2024. DistServe: disaggregating prefill and decoding for goodput-optimized large language model serving. In Proceedings of the 18th USENIX Conference on Operating Systems Design and Implementation (Santa Clara, CA, USA) (OSDI’24). USENIX Association, USA, Article 11, 18 pages. [55] Wenbin Zhu, Zhaoyan Shen, Zili Shao, Hongjun Dai, and Feng Chen. 2025. Tangram: Accelerating Serverless LLM Loading through GPU Memory Reuse and Affinity. arXiv preprint arXiv:2512.01357 (2025). [56] Edward Zitron. 2025. Exclusive: Here’s How Much OpenAI Spends On Inference and Its Revenue Share With Microsoft. https:// www.wheresyoured.at/oai_docs/
14
A
Parallel CUDA Graph Construction
We build a standalone test that allocates different numbers of threads to build 80 graphs, each with 500 dummy nodes, in parallel. Table 2 shows that wall time is almost consistent when increasing number of threads, and per API call time increases, confirming the contention between CUDA Graph driver APIs. Table 2. Durations of building 80 graphs each with 500 dummy nodes using different number of threads. #Threads 1 2 4 8
Wall Time 28.52 ms 29.32 ms 36.13 ms 37.54 ms
AddNode 0.70 𝜇s 1.44 𝜇s 3.52 𝜇s 7.31 𝜇s
Instantiate 0.8 𝜇s 1.6 𝜇s 3.8 𝜇s 6.6 𝜇s
15
B
Example Kernel Node Structure
"id": 7, "type": "KernelNode", "params": { "blockDimX": 384, "blockDimY": 1, "blockDimZ": 1, "gridDimX": 2, "gridDimY": 62, "gridDimZ": 1, "sharedMemBytes": 206044, "kernel_node_attrs": { "attrQueryAvailable": true, "clusterDimX": 2, "clusterDimY": 1, "clusterDimZ": 1, "clusterSchedulingPolicyPreference": 1, "memSyncDomainMapDefault": 0, "memSyncDomainMapRemote": 1 }, "kernelParams": [ { "index": 0, "offset": 0, "size": 1720 } ], "extra": [ "CU_LAUNCH_PARAM_BUFFER_SIZE", 1720, "CU_LAUNCH_PARAM_BUFFER_POINTER", "null", "CU_LAUNCH_PARAM_END" ], "extra_argBuffer_hex": "3e0000003e00000003000000eeffffffe4388ee304000000f7ffffffe 4388ee303000000fcffffff0000008001000000000102030408101820283037050911192129313806 0a121a222a3239070b131b232b333a0c141c242c343b0d151d252d353c0e161e262e363d0f171f272 f000000000000000000000000000000000000000.....1b600000 (3440 digits in total)", "function_name": "nvjet_tst_168x128_64x5_1x2_h_bz_TNN", "kernel_source_binary_hash": 6788486540864509700, "func_attrs": { "max_dynamic_shared_size_bytes": 206044, "preferred_shared_memory_carveout": -1, "cluster_scheduling_policy_preference": 0, "required_cluster_width": 0, "required_cluster_height": 0, "required_cluster_depth": 0 } }
16