Model2Kernel: Model-Aware Symbolic Execution For Safe CUDA Kernels

Model2Kernel is the first practical system that automatically verifies the memory safety of CUDA kernels in LLM inference by combining model-aware dynamic analysis with specialized symbolic execution to detect bugs caused by variable tensor layouts and user-controlled inputs, successfully identifying 353 previously unknown vulnerabilities with minimal false positives.

The Gist

Imagine you have built a massive, high-speed factory (a Large Language Model) that produces answers to your questions. This factory runs on a super-fast assembly line made of specialized machines called CUDA kernels. These machines are the "muscle" that does the heavy lifting, moving data around at lightning speed on your computer's graphics card (GPU).

However, because these machines are so fast and complex, they are prone to a specific kind of glitch: Memory Safety Bugs.

Think of a memory bug like a forklift driver in a warehouse who gets confused about where the boxes are.

• Buffer Overflow: The driver tries to put a box in a shelf that doesn't exist, knocking over everything nearby.
• Integer Overflow: The driver's clipboard runs out of numbers, so instead of writing "1,000,000," it loops back to "0," causing them to drive into a wall.
• Data Race: Two drivers try to grab the same box at the exact same time, dropping it or smashing it.

In the world of AI, these glitches don't just crash the factory; they can corrupt the AI's "brain" (its weights), shut down the service, or even let hackers take control of the machine.

The Problem: Why is this hard to fix?

Traditionally, checking these machines for bugs has been like trying to inspect a moving train while it's speeding by.

1. Dynamic Shapes: The factory changes its layout depending on how many questions you ask. Sometimes the warehouse is small; sometimes it's huge. Old tools can't handle this flexibility.
2. Thousands of Workers: These machines use thousands of workers (threads) simultaneously. Checking every single worker's path is impossible for humans and too slow for computers.
3. The "Black Box" of Models: The factory manager (the AI model) tells the machines what to do, but the rules change based on the specific model being used. Old tools didn't understand the manager's instructions.

The Solution: Model2Kernel

The authors of this paper built a new system called Model2Kernel. Think of it as a super-intelligent safety inspector that doesn't just watch the factory run; it simulates every possible scenario the factory could ever face, without actually needing the real factory to run.

It has two main parts, working like a perfect team:

1. HFProbe: The "Model Detective"

Imagine you want to test a new car, but you don't know how the driver will behave. HFProbe is the detective that watches the driver (the AI model) in a simulation.

• It runs the model on a "fake" GPU (no real hardware needed).
• It figures out exactly which buttons the driver presses and which parts of the factory they visit.
• Crucially, it tells the inspector: "Hey, the driver always keeps the warehouse width at 7,168 boxes, but the number of people entering changes every time."
• This helps the inspector know what is fixed (safe) and what is variable (risky).

2. cuKLEE: The "Parallel Time-Traveler"

Once the detective gives the clues, cuKLEE takes over. This is the symbolic execution engine.

• Instead of testing the factory with just one set of numbers (e.g., 100 people), cuKLEE uses magic variables. It asks, "What if there are 100 people? What if there are 10,000? What if there are 1 million?"
• It simulates all these scenarios at once.
• It treats the thousands of workers (threads) as a single group, checking if any of them would crash the system under any condition.
• If it finds a path where a forklift hits a wall, it stops and says, "Here is the exact combination of inputs that will break the machine."

The Results: A Safety Revolution

The team tested Model2Kernel on real-world AI factories (like vLLM and Hugging Face models).

• The Findings: They discovered 353 previously unknown bugs. Most were "integer overflows" (the clipboard running out of numbers) and "out-of-bounds" errors (reaching for a non-existent shelf).
• The Accuracy: They only had 9 "false alarms" (mistakenly thinking a safe spot was dangerous).
• The Comparison: When they tried to use older tools to find these same bugs, the older tools found almost nothing. It was like trying to find a needle in a haystack with a magnet, while Model2Kernel used a metal detector.

Why Does This Matter?

As AI becomes part of our daily lives (chatbots, search engines, self-driving cars), the "factories" running them are getting bigger and more complex.

• Safety: Model2Kernel ensures these factories won't crash when you ask them a tricky question.
• Security: It stops hackers from exploiting these memory glitches to steal data or hijack the AI.
• Efficiency: It does this automatically, without needing expensive hardware or slowing down the development process.

In short, Model2Kernel is the ultimate quality control system for the engines powering the AI revolution, ensuring that as these models get smarter, they also get safer.

Technical Summary

1. Problem Statement

The rapid adoption of Large Language Models (LLMs) has made GPU-accelerated inference critical. These systems rely on CUDA kernels to execute core transformer operations. However, these kernels are highly susceptible to memory-safety bugs (e.g., buffer overflows, integer overflows, data races, and null-pointer dereferences) due to:

• Model-Dependent Tensor Layouts: Tensor shapes and memory layouts are often determined at runtime based on the specific model architecture and user input (e.g., sequence length).
• Complex Indexing: Intricate calculations involving thousands of parallel threads and large tensors lead to out-of-bounds accesses and integer overflows.
• System Complexity: Mismatches between model assumptions and kernel implementations can cause memory errors.

Limitations of Existing Techniques:

• Dynamic Analysis (e.g., Compute Sanitizer): High runtime overhead or requires unavailable hardware.
• Static Analysis: Struggles with buffers sized at runtime and variable kernel launch configurations.
• General Symbolic Execution: Fails to handle dynamic arrays efficiently and cannot scale to the thousands of parallel threads in CUDA without excessive path explosion.

Consequently, no existing tool effectively detects memory bugs in LLM inference systems, posing risks of service crashes, weight corruption, and adversarial attacks.

2. Methodology: Model2Kernel

Model2Kernel is a fully automated framework that integrates dynamic model profiling with CUDA-specialized symbolic execution. It consists of two primary components: HFProbe and cuKLEE.

A. HFProbe (Dynamic Model Profiler)

HFProbe analyzes how specific LLM models (e.g., from Hugging Face) invoke CUDA kernels without requiring actual GPU hardware.

• Static Call Graph Construction: Analyzes Python model code (AST) to identify all potential CUDA kernels a model might call, handling dynamic dispatch and __call__ patterns common in PyTorch.
• Mocked Execution: Runs the model using modified inference frameworks (vLLM, Transformers) where actual GPU computations are disabled. Instead, "fake" functions record:
  • Call stacks and kernel arguments.
  • Tensor shapes, data types, and value ranges.
• Configuration Mutation: Uses an LLM agent to generate modified config.json files and framework parameters to trigger kernels that were not activated during initial profiling, ensuring comprehensive coverage.
• Context Generation: Infers constraints on kernel inputs, classifying them as fixed (determined by model architecture, e.g., hidden size) or variable (controlled by users, e.g., batch size, sequence length).

B. cuKLEE (CUDA-Specialized Symbolic Execution Engine)

Built on KLEE and operating on LLVM IR, cuKLEE performs symbolic execution on CUDA kernels using the context provided by HFProbe.

• Discrete Tensor Memory Model: Instead of a contiguous memory model (which fails with symbolic sizes), cuKLEE models each tensor as a distinct, well-separated memory region. It associates symbolic variables with tensor properties (base address, element count, dimensions). This simplifies pointer arithmetic and bounds checking.
• Tensor Method Summarization: Handles 140+ PyTorch Tensor methods by summarizing their effects on symbolic constraints (e.g., sum(0) reduces dimensionality) rather than executing internal logic, improving efficiency.
• Thread-Aware Execution:
  • Treats thread identifiers (threadIdx.x, blockIdx.x) as symbolic variables rather than enumerating values.
  • Executes the kernel logic once to capture the behavior of all threads simultaneously.
  • Detects data races by introducing a second set of symbolic thread IDs to check for conflicting writes to the same memory address.
• Bug Detection: Checks for integer overflows, out-of-bounds accesses, and null-pointer dereferences by generating constraints and solving them with the Z3 solver.

3. Key Contributions

1. cuKLEE: A novel symbolic execution engine designed specifically for CUDA kernels, featuring a discrete memory model for dynamic tensors and a unified approach to handling thousands of parallel threads.
2. HFProbe: A dynamic profiler that extracts model-architecture constraints and generates diverse execution contexts, bridging the gap between high-level model definitions and low-level kernel inputs.
3. Model2Kernel System: The first practical system to automatically verify memory safety in LLM inference systems, capable of handling real-world models from vLLM, Hugging Face, and research papers.

4. Experimental Results

The authors evaluated Model2Kernel on 120 models and 173 CUDA kernels from vLLM, Hugging Face, and recent research papers.

• Bug Detection:
  • Discovered 353 previously unknown memory bugs.
  • Breakdown: 328 integer overflows and 25 out-of-bounds accesses.
  • False Positives: Only 9 (False Discovery Rate of ~2.5%).
• Comparison with Baselines:
  • Compared against Honeycomb, GKLEE, and ESBMC-GPU on 20 known vLLM bugs.
  • Model2Kernel detected 15 of the 20 bugs.
  • Baselines detected 0–5 bugs due to inability to handle dynamic sizes, loop unrolling limits, or lack of thread-awareness.
• Ablation Study:
  • Removing HFProbe caused a massive spike in false positives (5,527) and analysis time due to the lack of architectural constraints.
  • Removing Configuration Mutation missed 32 bugs by failing to trigger specific kernel variants.
  • Running only on real GPUs (without symbolic execution) found zero bugs with small test inputs, proving the necessity of symbolic exploration.

5. Significance

• Security & Reliability: The paper highlights that memory bugs in LLM kernels can lead to service crashes, model weight corruption, and even arbitrary code execution on GPUs. Model2Kernel provides a critical defense mechanism.
• Scalability: It successfully scales symbolic execution to the complex, dynamic environment of modern LLM inference, overcoming the "state explosion" problem of traditional tools.
• Practical Impact: By identifying hundreds of bugs in widely used frameworks (vLLM) and models, it demonstrates that current production systems are significantly more vulnerable than previously realized.
• Future Directions: The approach of combining dynamic profiling with specialized symbolic execution offers a blueprint for securing other hardware accelerators (e.g., NPUs) and complex software stacks.