---
title: 'Model2Kernel: GPU Memory Safety in LLM Inference'
url: https://www.emergentmind.com/topics/model2kernel
type: topic
---

# Model2Kernel: GPU Memory Safety in LLM Inference

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 large language model 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 [2603.24595]. 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 [1409.5178][1001.2709][2601.05371].

## 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 [2603.24595].

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 [2603.24595].

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 [2603.24595].

## 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 [2603.24595].

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 [2603.24595].

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 [2603.24595].

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 [2603.24595].

## 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\) [2603.24595].

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 \(d_t^0, d_t^1, \ldots\), and element size \(S_t\). 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 [2603.24595].

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 \(t_2 = t_1.\mathrm{sum}(0)\), cuKLEE introduces constraints such as
\[
N_{t2} = N_{t1}/d^0_{t1}, \qquad D_{t2} = D_{t1}-1, \qquad S_{t2} = S_{t1}.
\]
This lets the engine reason about dynamic tensor layouts without entering the full tensor-library implementation [2603.24595].

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 [2603.24595].

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 [2603.24595].

## 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** [2603.24595].

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** [2603.24595].

| 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 [2603.24595].

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` [2603.24595].

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 [2603.24595].

## 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 \(p_M(y \mid x)\) is inserted directly into kernel Bayesian inference by computing the conditional kernel mean
\[
m_{\mathcal Y \mid x} = \int k_{\mathcal Y}(\cdot,y)\, p_M(y \mid x)\,dy,
\]
and the model-based kernel sum rule estimator
\[
\hat m_{Q_{\mathcal Y}} = \sum_{i=1}^{\ell}\gamma_i\, m_{\mathcal Y\mid \tilde X_i}.
\]
The paper’s central claim is that this lets mechanistic models or simulators be used directly inside RKHS inference rather than relearned nonparametrically [1409.5178].

In **“Kernel Topic Models”**, the Dirichlet prior over document-topic proportions is replaced by latent Gaussian document-topic variables \(y_d\), with
\[
\pi_d = \sigma(y_d),
\]
and topic-specific Gaussian-process priors \(h_k(\cdot)\) over document features. In the paper’s formulation, kernels enter through the GP covariance \(\eta_k(\phi_d,\phi_{d'})\), so the model-to-kernel move is to place document-topic logits under kernelized GP priors rather than to kernelize the multinomial likelihood directly [1110.4713].

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,
\[
\prod_{j=1}^{W} p(y_*^{(j)} \mid y, k_\theta),
\]
and an empirical covariance estimator
\[
YY^\top / M - \bar y \bar y^\top.
\]
This is a direct reverse-engineering program in which behavior is mapped to kernel structure [1510.07389].

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
\[
\|K(u,u)-K(u,v)-K(v,u)+K(v,v)\|_{B(Y)}^{1/2} \le \|u-v\|_U,
\]
and uses RKHS regularized least squares to identify nonlinear input-output operators that satisfy incremental integral quadratic constraints [2110.11735].

## 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 \(f=f_2\circ f_1\) induces an equivalent learned kernel machine with
\[
K(x_1,x_2)=K^2(f_1(x_1),f_1(x_2)).
\]
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 [1001.2709].

In **“Multiple Gaussian Process Models”**, additive latent GP components with scales \(\gamma_p^{-1}\) induce the effective combined covariance
\[
K_{\mathrm{comb}} = \sum_{p=1}^P \gamma_p^{-1} K_p.
\]
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 [1110.5238].

In **“Multiple kernel concept factorization algorithm based on global fusion”**, kernel construction is embedded directly inside unsupervised factorization through the fused kernel
\[
K_w=\sum_{i=1}^m \omega_i K^{(i)}, \qquad \omega_i\ge 0,\qquad \sum_{i=1}^m \omega_i=1,
\]
with the practical weight update
\[
\omega_i=\frac{1/e_i}{\sum_{j=1}^m 1/e_j}.
\]
The paper emphasizes that this is a global linear fusion mechanism in which representation learning and kernel construction are coupled [2410.20383].

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 \(D\) converted to
\[
B = -\frac12 J D J,
\]
and coordinates
\[
Z = V_p \Lambda_p^{1/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 [2601.05371].

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 [2603.24595]. 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 [1409.5178][1001.2709][2601.05371].

Source: https://www.emergentmind.com/topics/model2kernel