Model2Kernel: GPU Memory Safety in LLM Inference

• Model2Kernel is a system that automatically verifies CUDA kernel memory safety in LLM inference by integrating model-aware dynamic analysis with CUDA-specialized symbolic execution.
• It employs two components—HFProbe for dynamic tracing and cuKLEE for symbolic execution—to distinguish fixed architectural parameters from user-controlled inputs.
• Empirical evaluation on 120 models and 173 kernels shows that Model2Kernel uncovers 353 memory safety bugs with only 9 false positives, improving efficiency over prior methods.

Searching arXiv for papers on “Model2Kernel” and closely related kernel-learning formulations. Model2Kernel most directly denotes a system for automatically verifying the memory safety of CUDA kernels used in LLM inference. In that specific sense, it combines model-aware dynamic analysis with CUDA-specialized symbolic execution, using knowledge of how a model invokes kernels to distinguish arguments fixed by the model architecture from arguments controlled by model users (He et al., 6 Mar 2026). A broader reading suggested by adjacent kernel literature is that “Model2Kernel” also names a recurring methodological move: translating model structure, probabilistic assumptions, or learned representations into kernel objects, kernel combinations, or kernel search spaces (Nishiyama et al., 2014, Dinuzzo, 2010, Islam et al., 8 Jan 2026).

1. Specific meaning in LLM inference systems

In its explicit arXiv usage, Model2Kernel is a system for CUDA-kernel safety rather than a covariance-kernel learning method. The targeted setting is GPU-accelerated LLM inference, where production systems rely on CUDA kernels to implement core transformer operations. The motivating claim is that these kernels are highly susceptible to memory-safety bugs because of model-dependent tensor layouts, intricate memory indexing, and massive thread-level parallelism, and that such bugs can corrupt model weights, crash inference services, or even enable adversarial attacks (He et al., 6 Mar 2026).

The system is positioned against existing dynamic and static approaches. The paper states that prior techniques either depend on unavailable hardware, incur high overhead, or fail to handle kernel inputs with variable lengths, and that none can effectively detect CUDA memory bugs in LLM inference systems. Model2Kernel is presented as the first practical system for automatically verifying the memory safety of CUDA kernels used in LLM inference, targeting buffer overflows, integer overflows, data races, and NULL pointer dereferences (He et al., 6 Mar 2026).

A central source of difficulty is that kernel safety is model-conditional. The paper gives a representative overflow pattern in which one block is launched per token and the expression blockIdx.x × vec_hidden_size can exceed the 32-bit integer limit when hidden size is 8192, width = 8, vec_hidden_size = 1024, and runtime settings allow batch size up to 64 and sequence length up to 128,000. In the paper’s framing, this is precisely why kernel analysis cannot be decoupled from model architecture and serving-time inputs (He et al., 6 Mar 2026).

2. Model-aware workflow

Model2Kernel consists of two components: HFProbe and cuKLEE. HFProbe performs model-aware dynamic analysis, while cuKLEE performs CUDA-specialized symbolic execution. The overall workflow begins from the model’s forward() method, computes a static call graph over Python model code, identifies Python bindings that connect framework code to CUDA kernels, and determines which kernels a model may invoke (He et al., 6 Mar 2026).

HFProbe then executes models under modified versions of vLLM or Transformers without requiring a GPU. CUDA bindings are replaced with fake functions that record call stacks, parameter types, tensor shapes, tensor dtypes, primitive argument values, and, for integer tensors, min/max values if nonzero. The framework is also modified so that code paths expecting GPU modules still execute, while weight-loading functions create tensors of the correct shape without loading real weights (He et al., 6 Mar 2026).

From these traces, HFProbe derives an analysis context for each kernel usage. The paper states that it infers four classes of constraints: strict linear relationships with batch size, sequence length, or token count; invariant values fixed by model architecture; equality relationships between values; and residual parameters bounded by the maximum observed value. This supports the key classification of kernel arguments into architecture-fixed quantities, such as hidden size or invariant tensor dimensions, and user-controlled quantities, such as batch size, sequence length, or token count (He et al., 6 Mar 2026).

If static analysis finds kernels not exercised dynamically, HFProbe mutates config.json model parameters and framework-level configurations to trigger more kernels. The paper states that GPT-5 in agent mode is used to generate candidate configuration mutations and that failing mutations are discarded (He et al., 6 Mar 2026).

3. CUDA-specialized symbolic execution

cuKLEE starts from wrapper functions rather than isolated kernel bodies. This is important because wrappers compute launch configurations, hidden sizes, token counts, shared-memory sizes, pointer conversions, and other derived arguments. HFProbe-supplied constraints are injected as assumptions over wrapper inputs, including the initial symbolic variables $BS$ for batch size, $SL$ for sequence length, and $TC$ for token count, with

$TC = BS \times SL.$

The default bounds used by the system are $BS \le 1000$ and $SL \le 1{,}000{,}000$ (He et al., 6 Mar 2026).

A key abstraction is the dynamic tensor memory model. For each tensor $t$, cuKLEE creates symbolic metadata for base address $B_t$, number of elements $N_t$, number of dimensions $D_t$, dimensions $SL$0, and element size $SL$1. Tensor regions are modeled as discrete and widely separated memory regions, which avoids the standard symbolic-execution failure mode in which symbolic allocation sizes make later allocation addresses symbolic as well (He et al., 6 Mar 2026).

Tensor methods are summarized rather than inlined. The paper groups 140 Tensor and TensorIterator methods into four categories: property-returning methods such as numel() and data_ptr(), methods that produce new tensors and induce relational constraints, tensor property checks such as check_dim_size, and irrelevant methods that are ignored. For example, if $SL$2, cuKLEE introduces constraints such as

$SL$3

This lets the engine reason about dynamic tensor layouts without entering the full tensor-library implementation (He et al., 6 Mar 2026).

The second major abstraction is symbolic CUDA thread modeling. cuKLEE introduces symbolic variables for blockIdx.x/y/z, threadIdx.x/y/z, gridDim.x/y/z, and blockDim.x/y/z, together with CUDA launch constraints such as blockIdx.x < gridDim.x. This allows one symbolic execution to summarize all threads. For race detection, cuKLEE introduces a second symbolic thread under the same launch constraints and checks whether two different threads can access the same address with at least one write (He et al., 6 Mar 2026).

Bug predicates are then discharged with Z3. Out-of-bounds accesses are checked by asking whether a pointer falls outside a valid tensor region; integer overflows are checked on 32-bit additions and multiplications; NULL dereferences are checked by tracing unknown base pointers; and data races are checked by equating effective addresses across two symbolic threads (He et al., 6 Mar 2026).

4. Empirical results and limitations

The evaluation spans 120 models and 173 kernels, covering 62 vLLM text-generation models, 54 Hugging Face models with custom CUDA kernels, and 4 research-artifact models. The code volume is reported as 170,875 lines of Python and 139,951 lines of CUDA (He et al., 6 Mar 2026).

The main result is that Model2Kernel discovers 353 previously unknown bugs while producing only 9 false positives. The true positives consist of 328 integer overflows and 25 out-of-bounds accesses; the paper reports 0 data races and 0 null-pointer dereferences (He et al., 6 Mar 2026).

Source	True positives	False positives
vLLM	39	9
Hugging Face	290	0
Papers	24	0

The paper also reports 985.95 hours of aggregate analysis time, of which 902.19 hours are attributed to cuKLEE and 103.66 hours to HFProbe, and a GPT-agent monetary cost of \$442.16. HFProbe generates 1,274 distinct model configurations, merged into 8,562 unique contexts for cuKLEE (He et al., 6 Mar 2026).

An ablation result is especially important for the “model-aware” claim. Without HFProbe, cuKLEE produces 5,527 false positives, average analysis time per kernel increases by 2.5×, and 23 kernels time out within one hour, 16 more than with HFProbe. The paper attributes these failures to unrealistic tensor shapes and symbolic values such as hidden sizes near 2,147,483,640 (He et al., 6 Mar 2026).

The stated limitations are equally specific. The current system focuses on Hugging Face models and CUDA kernels with tensor inputs; it does not yet support Triton-based kernels or TensorFlow models. Each kernel is analyzed for at most one hour, large kernels can time out, cuKLEE inherits KLEE path-exploration strategies that were not designed for CUDA, one known bug is missed because std::map is unsupported, and one missed bug arises because tensor shape metadata disagree with underlying memory layout (He et al., 6 Mar 2026).

5. Broader “model-to-kernel” reading in machine learning

A broader interpretation of Model2Kernel is suggested by several papers that explicitly translate models, priors, or architectural structure into kernel objects. In “Model-based Kernel Sum Rule: Kernel Bayesian Inference with Probabilistic Models”, a known conditional probabilistic model $SL$4 is inserted directly into kernel Bayesian inference by computing the conditional kernel mean

$SL$5

and the model-based kernel sum rule estimator

$SL$6

The paper’s central claim is that this lets mechanistic models or simulators be used directly inside RKHS inference rather than relearned nonparametrically (Nishiyama et al., 2014).

In “Kernel Topic Models”, the Dirichlet prior over document-topic proportions is replaced by latent Gaussian document-topic variables $SL$7, with

$SL$8

and topic-specific Gaussian-process priors $SL$9 over document features. In the paper’s formulation, kernels enter through the GP covariance $TC$0, so the model-to-kernel move is to place document-topic logits under kernelized GP priors rather than to kernelize the multinomial likelihood directly (Hennig et al., 2011).

In “The Human Kernel”, human extrapolations are treated as posterior samples and used to infer a covariance structure that reproduces human predictive behavior. The paper gives both a predictive-conditional objective,

$TC$1

and an empirical covariance estimator

$TC$2

This is a direct reverse-engineering program in which behavior is mapped to kernel structure (Wilson et al., 2015).

A related systems-and-control example appears in “Kernel-based models for system analysis”, where desired input-output properties are converted into kernel conditions. The paper defines a nonexpansive operator-valued kernel through

$TC$3

and uses RKHS regularized least squares to identify nonlinear input-output operators that satisfy incremental integral quadratic constraints (Waarde et al., 2021).

6. Kernel composition, fusion, and search

A second cluster of Model2Kernel-style work learns kernels from architectural decomposition, from kernel libraries, or from optimization geometry. In “Kernel machines with two layers and multiple kernel learning”, a two-layer predictor $TC$4 induces an equivalent learned kernel machine with

$TC$5

The paper then shows that multiple kernel learning is a special case in which the second layer is linear and the resulting kernel is learned from the data (Dinuzzo, 2010).

In “Multiple Gaussian Process Models”, additive latent GP components with scales $TC$6 induce the effective combined covariance

$TC$7

Here the model decomposition itself determines the kernel decomposition, while generalized inverse Gaussian priors over the scales implement Bayesian weighting and sparsity control over kernel components (Archambeau et al., 2011).

In “Multiple kernel concept factorization algorithm based on global fusion”, kernel construction is embedded directly inside unsupervised factorization through the fused kernel

$TC$8

with the practical weight update

$TC$9

The paper emphasizes that this is a global linear fusion mechanism in which representation learning and kernel construction are coupled (Li et al., 2024).

In “The Kernel Manifold: A Geometric Approach to Gaussian Process Model Selection”, the emphasis shifts from constructing a kernel to organizing a discrete kernel library geometrically. Pairwise distances between GP priors are embedded by classical multidimensional scaling, with squared distance matrix $TC = BS \times SL.$0 converted to

$TC = BS \times SL.$1

and coordinates

$TC = BS \times SL.$2

Bayesian optimization is then run over the resulting kernel manifold, with the GP log marginal likelihood as the objective and the MDS coordinates as the featurization (Islam et al., 8 Jan 2026).

Taken together, these lines of work suggest that “Model2Kernel” has two distinct but compatible senses in current arXiv usage. In the narrow sense, it is a concrete verification system for CUDA kernels in LLM inference (He et al., 6 Mar 2026). In the broader methodological sense, it describes a family of constructions in which model structure, model behavior, or domain geometry is translated into kernel representations, kernel combinations, or kernel search spaces (Nishiyama et al., 2014, Dinuzzo, 2010, Islam et al., 8 Jan 2026).