KernelBench: Can LLMs Write Efficient GPU Kernels?

🎯 Pitch

Frontier LLMs fail to write correct GPU kernels more than 80% of the time, but iterative refinement using execution feedback lets DeepSeek-R1 achieve a 72% success rate on multi-operator tasks—a leap driven by nascent compiler-like reasoning, not just copying training data. The key bottleneck is functional correctness, exposing how CUDA’s near-absence from internet-scale pretraining data fundamentally limits models that otherwise show strong code-generation skills.

1. Executive Summary

This paper introduces KernelBench, an open-source evaluation framework that systematically studies whether language models can generate correct and performant GPU kernels for 250 diverse PyTorch ML workloads spanning single operations, operator sequences, and end-to-end architectures. The work evaluates state-of-the-art models—including GPT-4o, OpenAI o1, DeepSeek-V3, DeepSeek-R1, Claude 3.5 Sonnet, and Llama 3.1 variants—under a new metric called fastp (the fraction of generated kernels that are functionally correct and achieve speedup greater than threshold p over a PyTorch Eager baseline), and analyzes two test-time mechanisms for improving output: repeated sampling (generating multiple parallel candidate kernels with high temperature) and iterative refinement (sequentially improving kernels using compiler error feedback, execution correctness checks, and PyTorch profiler timing data). Frontier reasoning models achieve the strongest out-of-the-box results but still match PyTorch in fewer than 20% of cases, while iterative refinement with execution and profiler feedback boosts performance substantially—DeepSeek-R1 reaches 72% fast1 on Level 2 tasks—establishing that LMs demonstrate nascent optimization capabilities (operator fusion, shared memory management, algorithmic exploitation of sparsity) but remain severely bottlenecked by functional correctness errors, particularly given CUDA's scarcity in open-source training data at only 0.073% of The Stack v1.2 corpus.

2. Context and Motivation

The Core Problem: GPU Kernel Writing Is a Bottleneck for AI Progress

The fundamental problem this paper addresses is deceptively simple: writing efficient GPU kernels is hard, slow, and expertise-intensive, yet it is critical for making AI models actually work in practice. Every neural network architecture—whether a standard transformer, a state-space model, or a novel attention variant—must eventually be expressed as a sequence of operations that run on GPU hardware. The gap between a naive PyTorch implementation and a hand-tuned kernel can be enormous: the FlashAttention kernel (Dao et al., 2022), which is now essentially mandatory for running transformer models at scale, delivers order-of-magnitude speedups over standard attention implementations. But the paper points out a staggering timeline: FlashAttention was released in 2022, five years after the Transformer architecture was proposed in 2017, and it took another two years to port it to NVIDIA's Hopper GPU architecture (FlashAttention-3, Shah et al., 2024).

This lag is not an anomaly—it is structural. The paper characterizes the current situation as a "Cambrian explosion of ML architectures," citing diverse model families including RWKV (Peng et al., 2023), Mamba/state-space dual models (Dao and Gu, 2024), and efficient transformer variants (Tay et al., 2022). Each new architecture introduces novel computational patterns that would ideally be supported by custom, optimized kernels. But the human expertise to write these kernels is scarce, and the hardware landscape is diversifying simultaneously: the paper lists Cerebras wafer-scale engines, Graphcore IPUs, Groq inference chips, Google TPUs (Jouppi et al., 2023), and multiple generations of NVIDIA GPUs (V100, A100, H100) as evidence that "porting algorithms across platforms is a pain point." Each hardware platform has different memory hierarchies, instruction sets, and performance characteristics, requiring platform-specific optimization strategies.

This creates a dual bottleneck: (1) new architectures wait years for performant kernel implementations, slowing the research-to-deployment pipeline, and (2) when new hardware emerges, existing kernels must be rewritten, further delaying adoption. The paper asks: could language models automate this process? If LMs could generate correct and optimized kernels from PyTorch reference code, it could dramatically compress these timelines and democratize access to efficient AI implementations.

Why Kernel Generation Is Uniquely Difficult for Language Models

The paper argues that GPU kernel generation is a fundamentally different kind of code generation challenge compared to the tasks that existing LM coding benchmarks evaluate. Standard code generation benchmarks—the paper cites algorithmic coding (Chen et al., 2021; Shi et al., 2024; Li et al., 2022), GitHub issue resolution (Yang et al., 2024; Yang et al., 2024), and domain-specific coding tasks like data science notebooks (Lai et al., 2022; Yin et al., 2022)—primarily measure whether generated code is functionally correct. The challenge in kernel writing extends far beyond correctness:

1. Performance is the primary objective, not an afterthought. The paper distinguishes its focus from prior work on algorithmic efficiency (Nichols et al., 2024; Waghjale et al., 2024) by emphasizing wall-clock efficiency rather than asymptotic complexity. A generated kernel that is mathematically correct but runs slower than PyTorch's default implementation is a failure in this setting. The paper's fastp metric explicitly captures this by requiring both correctness and a speedup exceeding threshold p over the baseline. This is a harder standard than typical code generation metrics like pass@k (which only measures correctness) or execution match (which only checks output equivalence).

2. Kernel writing requires reasoning across multiple abstraction levels simultaneously. The paper describes the AI engineer's workflow as involving: compiler feedback (nvcc errors, warnings), profiling metrics (operator timing breakdowns, memory bandwidth utilization), hardware-specific specifications (register file sizes, shared memory capacities, tensor core instruction sets), and knowledge of hardware-efficiency techniques (tiling, operator fusion, shared memory management, asynchronous execution). An LM generating kernels must simultaneously produce low-level CUDA code while reasoning about high-level algorithmic tradeoffs and hardware constraints—a multi-scale reasoning challenge that the paper argues is not well-represented in existing benchmarks.

3. The programming surface is unusually large. The paper notes that kernel engineers can use tools ranging from assembly-level PTX instructions (as in DeepSeek-AI, 2025) to higher-level libraries like ThunderKittens (Spector et al., 2024), Triton (Tillet et al., 2019), and CUTLASS (NVIDIA, 2017). Each tool has different abstractions, performance characteristics, and portability properties. An LM tasked with kernel generation must navigate this tooling landscape and choose appropriate abstractions—a decision that existing code generation benchmarks typically make for the model by constraining the programming language or library.

4. Correctness is hard to verify by inspection. GPU kernels involve parallel execution across thousands of threads, shared memory with complex synchronization requirements, and hardware-specific numerical behaviors. Subtle bugs—race conditions, out-of-bounds memory accesses, incorrect synchronization—can produce outputs that appear correct on some inputs but fail on others, or that produce silent numerical errors. The paper observes in Section 4.2 that a large proportion of model-generated kernels fail with execution errors (CUDA compile-time errors, memory violations, runtime errors), and that even models with strong reasoning capabilities struggle to produce functionally correct kernels.

Where Prior Approaches Fall Short

The paper identifies three categories of prior work and explains why each does not solve the kernel generation problem:

Existing kernel libraries and compilers require human expertise. The paper reviews the landscape of existing tools for kernel programming:

"Mainstream kernel programming libraries like cuDNN [22], CUTLASS [23], and Apple MLX [1] are hardware-specific and demand substantial engineering effort from human experts."

These libraries provide highly optimized implementations of common operations (convolutions, matrix multiplications, etc.) but are built by expert human engineers over extended development cycles. They cover a fixed set of operations and do not adapt to novel architectures or new hardware without further human investment.

Higher-level libraries reduce but do not eliminate the programming burden. Tools like ThunderKittens and Triton successfully help AI researchers write fast and correct kernels for a broader range of workloads, but "still require human programming effort." They lower the expertise barrier by providing cleaner abstractions, but someone still has to write the kernel logic—deciding on tiling strategies, managing shared memory, and handling edge cases.

Compiler-based tools provide only narrow, rule-based optimizations. The paper discusses torch.compile (Paszke et al., 2019) and FlexAttention (Team PyTorch, 2024) as compiler-based approaches that "automatically provide a narrow slice of optimizations." torch.compile applies pattern-based fusion and graph transformations but does not perform the kind of algorithmic redesign (e.g., FlashAttention's tiling and recomputation strategy) that yields order-of-magnitude speedups. Compilers operate within a fixed optimization repertoire, while human kernel engineers can invent entirely new computational strategies.

Existing HPC code generation research focuses on translation or well-known kernels. The paper acknowledges prior work on LM-based HPC code generation but distinguishes KernelBench's scope:

"Existing works in the space of HPC code generation have evaluated LM performance on translating arbitrary code samples from C++ to CUDA [35, 41] or generating well-known, low-level kernels such as GEMMs [38, 42]. KernelBench instead curates a set of 250 diverse kernels from real-world, modern deep learning workloads, many of which do not have existing human-written implementations—in other words, solving KernelBench tasks are immediately beneficial for real deep learning workloads."

This is a crucial distinction. Prior work tests whether LMs can replicate known kernels (like matrix multiplication) or translate code between languages—tasks for which training data likely exists. KernelBench tests whether LMs can generate novel kernels for real workloads that currently lack optimized implementations. This shifts the evaluation from "can LMs reproduce known solutions?" to "can LMs discover optimizations that haven't been built yet?"—a much higher bar that directly measures practical utility.

Standard LM code generation benchmarks do not capture the kernel writing challenge. The paper positions KernelBench against a landscape of code generation benchmarks and argues that none of them require the combination of correctness, performance optimization, hardware awareness, and multi-level reasoning that kernel writing demands. Benchmarks like HumanEval, MBPP, and DS-1000 focus on functional correctness. Benchmarks like SWE-Bench test the ability to resolve complex real-world software tasks but do not emphasize performance optimization. Even emerging benchmarks focused on code efficiency (Nichols et al., 2024; Waghjale et al., 2024) operate at the level of algorithmic complexity rather than wall-clock speed on specific hardware.

The Data Scarcity Problem

The paper identifies a specific structural reason why LMs might struggle with kernel generation: CUDA is severely underrepresented in open-source training data. The paper reports that CUDA code constitutes only 0.073% of The Stack v1.2 (Kocetkov et al., 2022; Li et al., 2023), a popular code corpus used for training language models. To put this in perspective, this is less than one-tenth of one percent of the training data, yet kernel generation requires mastery of a domain-specific language with complex semantics (CUDA) plus deep understanding of hardware architecture and parallel computing concepts. The paper hypothesizes that this data scarcity is a primary cause of execution failures:

"We hypothesize this is due to CUDA being a low-resource language in open-source training data, only 0.073% of popular code corpus The Stack v1.2."

This insight reframes the problem: LM kernel generation is not just about developing better prompting or reasoning strategies—it is fundamentally constrained by the availability of training data in a domain where expertise is already scarce. Human kernel engineers are rare, and the kernels they write are often proprietary (e.g., cuBLAS, cuDNN are closed-source). This creates a chicken-and-egg problem: there is not enough open-source CUDA code to train LMs effectively, and LMs are not good enough at CUDA to generate new open-source kernels that would expand the training data.

How This Paper Positions Itself

The paper positions KernelBench as filling three specific gaps in the existing ecosystem:

1. An evaluation framework that reflects real-world kernel engineering. Rather than constructing synthetic problems or testing on well-known kernels with existing solutions, KernelBench tasks are drawn from actual PyTorch workloads. The paper emphasizes that "making progress on the introduced benchmark directly translates to faster practical kernels." This is not just a benchmark for measuring progress—it is designed so that any kernel that performs well on a KernelBench task is immediately usable in production ML pipelines. The three-level structure (single operations, operator sequences, end-to-end architectures) directly mirrors the granularity at which real kernel optimization decisions are made.

2. A unified metric that captures the dual objectives of correctness and performance. The fastp metric is the paper's proposed solution to the evaluation challenge described earlier. Unlike pass@k (which ignores speed) or raw speedup measurements (which ignore correctness failures), fastp explicitly couples the two requirements: a generated kernel only counts as successful if it is both correct and faster than the baseline by a factor of at least p. The threshold parameter p can be adjusted to set harder or easier standards as the field progresses:

"By adjusting the threshold parameter p, we enable evaluation of kernel performance at different speedup thresholds and capture the speedup distributions."

The paper also notes that p < 1 is valuable for training settings, since even kernels that match 80% of PyTorch's performance could be useful if they fill an open-source gap or provide a foundation for further optimization.

3. A platform for studying test-time improvement strategies in kernel generation. Beyond evaluating one-shot generation, the paper explicitly designs KernelBench to support the study of how LMs can improve their kernels through interaction with the environment. The framework provides programmatic access to compiler errors, execution correctness checks, and profiler timing breakdowns—the same signals a human kernel engineer uses. This enables the paper's analysis of repeated sampling and iterative refinement (Section 5) as test-time strategies, and more broadly opens the door for future work on agentic workflows that combine generation, execution, profiling, and refinement in loops that mirror the human development process.

The Broader Vision: A Living Benchmark That Evolves With AI

The paper's concluding discussion (Section 6) articulates a vision that distinguishes KernelBench from static benchmarks that eventually saturate. Because PyTorch is cross-hardware compatible, the same 250 tasks can be evaluated on every new GPU generation that emerges. As hardware improves, the PyTorch Eager baseline will shift (newer GPUs running the same reference code will be faster), and the fastp metric will naturally reset to a harder target. As new AI architectures are proposed, new tasks can be added to the benchmark. This design ensures that KernelBench remains a moving target that tracks the frontier of what is practically valuable:

"Since PyTorch is cross-hardware platform compatible, the PyTorch-based tasks in KernelBench tasks can be evaluated on every new hardware platform release... These properties ensure that KernelBench will remain valuable in the ever-evolving AI landscape."

In this sense, the paper is not just presenting a benchmark but also proposing a methodology for continuously evaluating LM-driven code optimization in a domain where both the hardware and the software targets are in constant flux. The motivation is ultimately about building infrastructure for a future where LM-assisted kernel generation becomes a routine part of the AI development workflow, compressing the years-long gaps between architecture proposal, hardware release, and efficient implementation that currently characterize the field.

3. Technical Approach

3.1 Reader Orientation

This paper presents KernelBench, an evaluation framework—not a new model or training method—that defines a set of 250 GPU kernel generation tasks, a protocol for evaluating language model outputs against those tasks, and a metric (fastp) that jointly captures correctness and wall-clock speedup. The system being built is not a single model but rather an automated evaluation pipeline that takes a PyTorch reference implementation as input, queries a language model to produce an optimized version (which may include inline CUDA kernels), compiles and executes the generated code on a GPU, checks its functional correctness against the reference, measures its wall-clock runtime, and reports whether it achieves a threshold speedup. The problem it solves is that no existing benchmark evaluates LMs on the specific challenge of kernel generation—producing code that is simultaneously correct (matching reference outputs across random inputs) and performant (running faster than PyTorch's already-optimized backend on real GPU hardware)—and no existing framework provides the programmatic environment to study how LMs can use compiler errors, execution feedback, and profiling data to iteratively improve their kernels.

3.2 Big-Picture Architecture (Diagram in Words)

The KernelBench system has five major components, connected in a pipeline that mirrors the workflow of a human AI engineer optimizing a kernel:

1. Task Definition Store (250 pre-specified problems): Each task is a Python file containing a Model class (inheriting from torch.nn.Module) with an __init__ method and a forward method that implements the AI workload using PyTorch operations, plus get_inputs() and get_init_inputs() functions that specify exact tensor shapes and data types. Tasks are partitioned into three levels: Level 1 (100 tasks, single primitive operations like matrix multiplies, convolutions, activations), Level 2 (100 tasks, sequences of 3–6 operations that can benefit from fusion), and Level 3 (50 tasks, full architectures like AlexNet and MiniGPT).

2. Language Model Interface (query and response handling): Given a task's reference Model code as a prompt (with optional in-context examples, hardware specifications, or prior generation history), the LM generates a new class called ModelNew that replaces some or all PyTorch operations with custom optimized code—typically inline CUDA kernels compiled via torch.utils.cpp_extension.load_inline, but potentially using Triton, CUTLASS, ThunderKittens, or PTX. The LM has complete freedom to decide which operations to optimize and what techniques to use.

3. Compilation and Execution Engine (CPU pre-compilation + GPU evaluation): The generated ModelNew code is compiled with nvcc on CPU, producing a binary that is cached. The compiled kernel is then run on a bare-metal GPU (primarily NVIDIA L40S for reported results) with careful timing: 3 warm-up iterations followed by 100 timed iterations using torch.cuda.Event markers, taking the mean wall-clock time. Only one kernel is evaluated at a time on a given GPU to ensure accurate timing, but the system parallelizes across multiple GPUs and CPUs for throughput.

4. Correctness Verification Module (randomized input testing): For each task, the reference Model and the generated ModelNew are both executed on 5 randomly generated input tensors (with the same shapes and dtypes specified in the task). Their outputs are compared for exact equivalence—both tensor shapes and values must match. The paper explicitly grounds this choice in practical necessity: exact program equivalence is undecidable (the Halting Problem), so randomized testing with diverse inputs is the standard practical compromise, and 5 inputs were chosen based on empirical observation that no generated kernel showed partial correctness (all were either 5/5 or 0/5 correct).

5. fastp Metric Computer (combined correctness-speedup scoring): For each task, the system computes fastp, defined as:

$\text{fast}_p = \frac{1}{N} \sum_{i=1}^{N} \mathbf{1}(\text{correct}_i \land \{\text{speedup}_i > p\})$

where $N$ is the total number of tasks in the set being evaluated, $\text{correct}_i \in \{0, 1\}$ is a binary indicator of whether the generated kernel passes all 5 correctness checks for task $i$ , and $\text{speedup}_i$ is the ratio $\frac{T_{\text{Model, }i}}{T_{\text{ModelNew, }i}}$ of reference wall-clock time to generated kernel wall-clock time.

What it computes: For a given set of tasks and a speedup threshold $p$ , the metric counts how many tasks produce kernels that are both functionally correct and faster than the baseline by at least a factor of $p$ , then divides by the total number of tasks to produce a percentage. For example, fast1 (with $p = 1$ ) measures the percentage of tasks where the generated kernel is correct and achieves any speedup over PyTorch Eager. fast0 (with $p = 0$ ) measures correctness alone, since the speedup condition $\text{speedup}_i > 0$ is always satisfied for any positive runtime.

Why this form: Joint metrics are necessary because correctness and performance can trade off against each other—a model might produce a few extremely fast kernels but most are incorrect (yielding a high average speedup on correct tasks but low practical utility), or might produce many correct but slow kernels (high correctness but no practical benefit). The product logically connects the two requirements: a task only counts if both conditions hold. The threshold parameter $p$ allows the difficulty of the benchmark to be adjusted over time: as LMs improve, $p$ can be raised (e.g., requiring 2× speedup) to maintain challenge, or lowered below 1 for training scenarios where partially-matching PyTorch performance is still valuable.

Information flows sequentially through these components: a task is selected from the store → the LM generates ModelNew → the code is compiled (with any compilation errors captured as feedback if in iterative refinement mode) → the compiled kernel is executed on GPU and timed → correctness is checked against the reference outputs → the fastp score is computed and aggregated across tasks. In iterative refinement mode (Section 5.1.2), an additional feedback loop connects the execution/compilation results back to the LM for multi-turn improvement.

3.3 Roadmap for the Deep Dive

First, the task specification format and the three-level task taxonomy—what exactly the LM receives as input, what it must produce as output, and how the 250 problems were selected from real ML workloads—since this defines the "shape" of the generation problem.
Second, the evaluation protocol in detail: how correctness is verified (randomized input testing, the 5-input choice), how performance is measured (CUDA event timing, warm-up iterations, coefficient of variation), and how the PyTorch baselines (Eager vs. torch.compile) are constructed—since the fastp metric depends critically on rigorous and reproducible measurement.
Third, the fastp metric derivation in full: the logical formula, the role of the threshold $p$ , why it captures the dual objective better than separate correctness and speedup reporting, and how it enables comparisons across models with different correctness-speedup profiles.
Fourth, the one-shot baseline prompting protocol: the exact prompt template (with in-context example), the reasoning behind temperature = 0 for deterministic evaluation, and the model set evaluated—since this establishes the "out-of-the-box" performance ceiling.
Fifth, the two test-time improvement strategies: repeated sampling (parallel generation with high temperature, fastp@k metric) and iterative refinement (sequential multi-turn improvement using compiler feedback E, execution results, and profiler data P), including the exact feedback format and multi-turn state machine design.
Sixth, the hardware-aware conditioning experiments: providing in-context examples (fusion, tiling, FlashAttention) versus providing hardware specifications (GPU type, memory sizes, bandwidths, TFLOPS, register/shared memory capacities)—since these test whether LMs can incorporate domain knowledge when explicitly prompted.
Seventh, the high-throughput evaluation infrastructure that makes all of this computationally feasible: the three-stage pipeline (inference → CPU pre-compile → GPU evaluation), the orchestrator for iterative refinement experiments, and the handling of CUDA errors/timeouts.

3.4 Detailed, Sentence-Based Technical Breakdown

This is primarily a benchmark and empirical analysis paper whose core idea is that evaluating LMs on GPU kernel generation requires (1) a carefully curated set of real-world AI workloads spanning multiple granularity levels, (2) a unified metric that couples correctness with wall-clock speedup, and (3) an evaluation environment that supports both one-shot generation and iterative improvement using the same feedback signals a human kernel engineer would use.

Task Specification Format and Three-Level Taxonomy

Each of the 250 tasks in KernelBench is a self-contained Python file that defines a reference implementation and specifies the exact input tensors the generated kernel must handle. The format is designed to mirror what an AI engineer would encounter: given a PyTorch implementation of a model component, produce an optimized version that can be dropped in as a replacement.

The Model class (reference implementation). Every task provides a class named Model that inherits from torch.nn.Module. The class contains:

An __init__ method that initializes any parameters (e.g., weights for linear layers, convolutional filters). For tasks that have no parameters, this method is empty.
A forward method that takes tensors as arguments and returns tensors, implementing the AI workload using standard PyTorch operations (torch.matmul, torch.nn.functional.relu, torch.nn.functional.conv2d, etc.).

The LM's job is to produce a new class named ModelNew (also inheriting from torch.nn.Module) that is functionally equivalent to Model—its forward method must produce identical outputs given identical inputs—but achieves lower wall-clock time. The LM can use any optimization technique and any programming abstraction: inline CUDA-C via torch.utils.cpp_extension.load_inline, Triton kernels, CUTLASS templates, ThunderKittens primitives, PTX assembly, or even pure PyTorch with algorithmic improvements. The framework imposes no constraints on the approach.

Input shape specification. Each task includes two functions that generate random tensors with the exact shapes and data types the task requires:

get_inputs() returns a list of tensors to be passed as arguments to Model.forward(). These tensors are randomly generated with the prescribed shapes (e.g., torch.randn(M, K) for a matrix multiplication with dimensions M×K and K×N).
get_init_inputs() returns a list of tensors needed to initialize the model's parameters (e.g., weights for a linear layer). For parameter-free models, this returns an empty list.

These functions serve a dual purpose: they define the operating point at which the kernel must be optimized (since the optimal kernel for matrix multiplication depends on whether the matrices are 256×256 or 256×131072), and they provide the randomization seeds for correctness testing (5 different random instantiations of the same shapes).

Example: matrix multiplication task. The paper provides a concrete task example in Appendix A. The reference Model computes torch.matmul(A, B) where A has shape (256, 131072) and B has shape (131072, 256)—a very tall-skinny matrix multiply. The get_inputs() function generates A = torch.randn(256, 131072) and B = torch.randn(131072, 256). An LM's ModelNew might replace torch.matmul with a custom CUDA kernel that uses tiled access patterns to improve memory coalescing, wrapped in load_inline for PyTorch integration.

The three-level taxonomy. The paper partitions the 250 tasks into three levels based on the number of PyTorch primitive operations they contain:

Level 1: Single primitive operations (100 tasks). These are the atomic building blocks of deep learning: matrix-vector and matrix-matrix multiplications (including variants with transposed operands, upper-triangular constraints, diagonal matrices, 3D and 4D tensor contractions), convolutions (1D, 2D, 3D with various kernel sizes, padding, and strides), activation functions (GELU, ReLU, Swish, Mish, HardSwish, Softsign), normalization layers (LayerNorm, GroupNorm, InstanceNorm, BatchNorm), loss functions (cross-entropy, cosine similarity, triplet margin loss), and element-wise operations. The paper notes that PyTorch's implementations of these primitives already call highly optimized closed-source kernels (e.g., cuBLAS for matrix multiplies, cuDNN for convolutions), making Level 1 a challenging baseline—to succeed, an LM must either match expert-tuned closed-source code or find novel algorithmic optimizations.

Level 2: Operator sequences (100 tasks). These tasks combine 3–6 primitive operations into a single computational unit that could be fused into one kernel to reduce memory I/O. Examples include: a matrix multiplication followed by ReLU and bias addition, a convolution followed by GroupNorm and Swish activation, a 3D convolution with Mish activation and Hardtanh scaling. The paper emphasizes that "compiler-based tools such as the PyTorch compiler are effective at fusion," making Level 2 a strong baseline—to succeed, LMs must outperform the fusion heuristics already present in torch.compile. However, "LMs may propose more complex algorithms compared to compiler rules," since compilers apply fixed pattern-matching while LMs can reason about the semantics of the combined operations and potentially discover fusion strategies that aren't in the compiler's repertoire.

Level 3: Full ML architectures (50 tasks). These are complete neural network architectures collected from popular PyTorch repositories on GitHub (including pytorch, huggingface/transformers, and huggingface/pytorch-image-models). Examples include AlexNet and MiniGPT. These architectures contain many operations, and optimizing them requires reasoning at the architectural level—deciding which components to accelerate, how to structure kernel launches, and potentially making algorithmic modifications that span multiple layers. The paper uses the Transformer attention mechanism as an illustrative case: it took 5 years from the Transformer's proposal (Vaswani et al., 2017) to obtain performant kernels (FlashAttention, Dao et al., 2022), and those kernels required algorithmic redesign (tiling, recomputation, online softmax) that "are often beyond the scope of a compiler." Level 3 tests whether LMs can discover such architectural-level optimizations.

Design rationale for the three-level structure. This taxonomy serves several purposes. First, it captures the natural hierarchy of kernel optimization: some optimizations apply to individual operations (e.g., using tensor core instructions for matrix multiplication), some require fusing multiple operations to reduce memory traffic, and some require architectural redesign. Second, it enables separate analysis of where LMs succeed and fail—a model might be good at fusing small operator sequences but unable to handle end-to-end architectures, or vice versa. Third, it maps to practical utility: Level 1 kernels that outperform PyTorch's closed-source backends would be immediately valuable as open-source alternatives; Level 2 successes would complement or replace compiler-based fusion; Level 3 successes would directly accelerate the deployment of novel architectures.

Evaluation Protocol: Correctness, Performance Measurement, and Baselines

The evaluation protocol is designed to be fully automatic, reproducible, and rigorous enough that a fastp score for a given model-kernel combination on a given GPU can be trusted as a measure of practical utility.

Correctness verification via randomized input testing. For each task, correctness is checked by:

Generating 5 sets of random input tensors using get_inputs() (and get_init_inputs() for model construction).
Running Model.forward() and ModelNew.forward() on each input set.
Comparing the output tensors for exact equivalence: both the shape (number of dimensions, size of each dimension) and the numerical values must match.

The choice of 5 inputs—rather than 1 or 100—is empirically justified. The paper reports (Appendix B.2): "In an experiment with 100 generated kernels, the results were as follows: 50 kernels were correct (all 5/5 and 100/100), 19 had output value mismatches (19 0/5 and 0/100), 4 had output shape mismatches, 10 encountered runtime errors, and 17 had compilation errors. Notably, the 0/5 and 0/100 failures indicate that no partial correctness was observed." In other words, there were no kernels that passed 5 tests but failed 100, or vice versa—correctness was all-or-nothing at 5 inputs. This observation supports 5 as a sufficient sample size for this domain, where GPU kernels tend to have systematic rather than input-dependent errors (a tiling bug will produce wrong values on every input, not just some).

The paper explicitly acknowledges the theoretical limitation: "checking equivalence of programs in a formal sense is undecidable" (citing the Halting Problem, Turing 1936), and 5 random inputs is an approximate heuristic. However, for AI kernels—which have simpler control flow than general programs and focus on numerical computation—randomized testing is the standard practical approach. The paper notes that "future work could investigate formal verification tools to provide stronger guarantees of equivalence," but for the current benchmark, randomized testing provides adequate discrimination between correct and incorrect kernels.

Performance measurement via CUDA event timing. Wall-clock execution time is measured as follows:

Isolation: Only one kernel is evaluated at a time on the GPU—no other CUDA processes run concurrently.
Warm-up: 3 forward passes are executed before timing begins, to ensure the GPU is in a steady state (first launch overhead, cache warm-up, etc.).
Timed iterations: 100 forward passes are measured using torch.cuda.Event markers that bracket the execution. These CUDA events capture GPU-side time (not CPU-side overhead from Python or PyTorch dispatch).
Statistics: The mean, maximum, minimum, and standard deviation of the 100 trials are recorded. The paper reports that "our coefficient of variation (CV): std/mean is consistently < 3%," indicating that timing variance is low enough for reliable mean-based comparisons.

The speedup for a single task is computed as:

$\text{speedup}_i = \frac{T_{\text{Model, }i}}{T_{\text{ModelNew, }i}}$

where $T_{\text{Model, }i}$ is the mean wall-clock time of the reference PyTorch Model over 100 trials for task $i$ , and $T_{\text{ModelNew, }i}$ is the mean wall-clock time of the generated ModelNew over 100 trials.

Why mean-based speedup: The mean over 100 trials is used (rather than minimum, which would be optimistic, or median, which would be conservative) because the low coefficient of variation (<3%) means the mean is a stable estimator not dominated by outliers. The paper explicitly states that for reporting fastp, "if TModel = 2 ms and TModelNew = 1 ms, we have a 2× speedup with the newly generated kernel." Alternative approaches—such as using the minimum observed time (which could reflect ideal cache states) or the maximum (which could reflect worst-case interference)—would shift speedup values but are not explored.

Two PyTorch baselines. The paper evaluates generated kernels against two baseline execution modes:

PyTorch Eager: The default execution mode where each PyTorch operation dispatches to pre-compiled, highly optimized kernels (e.g., cuBLAS for matrix multiplies, cuDNN for convolutions). These kernels are often closed-source and hand-tuned by expert engineers. PyTorch Eager is the primary baseline used throughout the paper's analysis because it represents the "naive but optimized" starting point that a human engineer would try to beat.

torch.compile (PyTorch 2.0's compiler): An ahead-of-time compilation system that analyzes the PyTorch computation graph and applies optimizations including operator fusion, memory planning, and backend-specific code generation. The paper tests torch.compile in several configurations (Table 3): default mode with the Inductor backend, reduce-overhead mode, max-autotune mode (which performs extensive kernel autotuning at compile time), max-autotune-no-cudagraphs, and the cudagraphs backend with AOT Autograd. Importantly, only runtime is measured—"we exclude the torch.compile compile time in our timing analysis, as we are only interested in the raw runtime behavior."

An interesting finding from the torch.compile baseline comparison (reported in Table 1 and Table 4): "the torch.compile baseline runtime is sometimes slower than Torch Eager—this is due to reproducible runtime overhead (not compile time) that could be significant for small kernels in Level 1." This means that for small single-operation kernels, the compiler's runtime bookkeeping (graph management, memory allocation patterns) can outweigh the optimization benefits. For Level 2 and 3 tasks, torch.compile is generally faster due to effective operator fusion.

Hardware platform for primary results. All main evaluations (Tables 1–2, Figures 2–6) are conducted on a bare-metal NVIDIA L40S GPU with 48 GB HBM memory, 300W power, Ada Lovelace architecture. The software stack is Python 3.10, PyTorch 2.5.0+cu124, CUDA 12.4. Cross-hardware experiments (Section 4.4, Appendix G) additionally test on H100 (80 GB, Hopper), A100 (42 GB, Ampere), L4 (24 GB, Ada), T4 (16 GB, Turing), and A10G (24 GB, Ampere) GPUs.

The fastp Metric: Derivation, Properties, and Rationale

The fastp metric is the paper's core contribution for comparing different LMs on KernelBench. Its formal definition is:

$\text{fast}_p = \frac{1}{N} \sum_{i=1}^{N} \mathbf{1}(\text{correct}_i \land \{\text{speedup}_i > p\})$

where $N$ is the total number of tasks in the evaluation set (e.g., 100 for Level 1), $\text{correct}_i \in \{0, 1\}$ is the binary correctness indicator for task $i$ , $\text{speedup}_i = T_{\text{Model, }i} / T_{\text{ModelNew, }i}$ is the observed speedup ratio, $p$ is the speedup threshold parameter, and $\mathbf{1}(\cdot)$ is the indicator function (1 if the condition is true, 0 otherwise).

What it computes: For a fixed threshold $p$ , the metric scans through all $N$ tasks, checks for each whether the LM produced a kernel that (a) matched the reference output on all 5 random inputs AND (b) ran at least $p$ times faster than the baseline, counts how many tasks satisfy both conditions, and divides by $N$ to express the result as a percentage. The output is a single number in [0, 100] that represents the fraction of the benchmark where the LM's kernel is both correct and meaningfully faster.

Interpretation at different p values. The paper defines three regimes:

fast0: When $p = 0$ , the speedup condition becomes $\text{speedup}_i > 0$ , which is always true for any kernel that runs to completion (since execution time is positive). Thus fast0 reduces to the correctness rate alone: "fast0 is equivalent to the LM's correctness rate, as it measures the fraction of tasks for which the LM code is functionally correct regardless of its speed."
fast1: The primary evaluation threshold in the paper. When $p = 1$ , a kernel counts as successful only if it is strictly faster than the baseline (any speedup > 1.0×). This captures the minimum practical requirement: the generated kernel must not be slower than what a user would get by simply running PyTorch.
fastp for $p > 1$ : As the threshold increases, the benchmark becomes harder, requiring progressively larger speedups. For example, fast2 would require at least 2× speedup. The paper focuses on $p = 1$ for current results but notes that "increasing the threshold p increases the difficulty" and that future work can raise $p$ as LM capabilities improve.

Support for p < 1. The paper explicitly notes that thresholds below 1 are meaningful in some contexts: "using p < 1 for training is valuable, since PyTorch relies on complex optimized kernels, and matching even a fraction of their performance is still considered beneficial." For instance, an open-source kernel that achieves 0.8× the speed of cuBLAS might still be valuable if it provides a permissively licensed alternative that can be customized or studied.

Why this form over alternatives: The paper's key insight is that correctness and speedup must be coupled because they exhibit a selection effect: if you report them separately, a model that generates 1 correct kernel with 10× speedup and 99 incorrect kernels looks worse on correctness (1%) but better on average speedup (10×) than a model that generates 100 correct kernels with 1.1× speedup (100% correctness, 1.1× average speedup). Neither statistic alone captures the practical utility. By requiring both conditions simultaneously, fastp forces the evaluator to consider the joint distribution: a kernel only contributes if it is both correct AND fast.

Furthermore, the indicator-function formulation means that speedup magnitude above p is not rewarded—a kernel with 10× speedup counts the same as one with 1.1× speedup, as long as both exceed $p$ . This is a deliberate choice: for practical deployment, distinguishing between "fast enough" and "even faster" is less important than distinguishing between "fast enough" and "not fast enough." However, the paper also provides speedup distribution plots (Figure 7 violin plots, Figure 3 fastp curves) that show the full distribution, so readers can assess magnitude differences separately.

Extension to repeated sampling: fastp@k. When evaluating repeated sampling (Section 5.1.1), the metric is extended to fastp@k, defined as:

$\text{fast}_p@k = \frac{1}{N} \sum_{i=1}^{N} \mathbf{1}\left(\max_{j=1}^{k} (\text{correct}_{i,j} \land \{\text{speedup}_{i,j} > p\})\right)$

where $k$ is the number of independent samples drawn from the LM for each task, $\text{correct}_{i,j}$ and $\text{speedup}_{i,j}$ refer to the $j$ -th sample for task $i$ , and the maximum over $j$ captures whether any of the $k$ samples was both correct and fast enough. This is the standard pass@k formulation adapted to the fastp joint condition: a task is "solved" if at least one sample achieves both correctness and speedup > $p$ .

Extension to iterative refinement: fastp@N. For iterative refinement with $N$ turns (Section 5.1.2), the metric is fastp@N, which measures "the percentage of tasks where the model generated at least one functionally correct kernel that is p times faster than PyTorch Eager by turn N." This is analogous to fastp@k but the $N$ samples are generated sequentially rather than in parallel, with each turn conditioned on feedback from the previous turn.

One-Shot Baseline Prompting Protocol

The one-shot baseline (Section 4.1) establishes each model's out-of-the-box ability to generate kernels given minimal instructions. The prompt design follows a deliberate minimalist philosophy: provide just enough information to specify the task format, then see what the model produces without task-specific hints, hardware information, or optimization guidance.

Prompt structure. The prompt consists of three parts:

Task instruction: A brief paragraph telling the model it should write custom CUDA kernels to replace PyTorch operators in the given architecture to achieve speedups. It emphasizes freedom: "You have complete freedom to choose the set of operators you want to replace. You may make the decision to replace some operators with custom CUDA kernels and leave others unchanged. You may replace multiple operators with custom implementations, consider operator fusion opportunities... or algorithmic changes." This open-ended instruction is intentional—the benchmark evaluates whether LMs can autonomously identify optimization opportunities, not whether they can follow a prescribed optimization recipe.
Format example: A simple in-context example showing a PyTorch Model class that computes element-wise addition (a + b) and a corresponding ModelNew that replaces it with an inline CUDA kernel using load_inline. This example serves purely as a syntax template—it shows the #include structure, the kernel launch configuration, and the load_inline compilation boilerplate. The example kernel is deliberately trivial (element-wise add) so it does not teach optimization techniques.
Task specification: The reference Model code for the specific KernelBench problem, inserted inline in the prompt.

The model is instructed to "Name your optimized output architecture ModelNew. Output the new code in codeblocks. Please generate real code, NOT pseudocode, make sure the code compiles and is fully functional. Just output the new model code, no other text, and NO testing code!"

Sampling configuration. The one-shot baseline uses greedy decoding (temperature = 0) for all models. This choice ensures deterministic output and provides a clean measurement of each model's "most likely" kernel generation, without the variability that stochastic sampling would introduce. The paper acknowledges that this is a conservative baseline—higher temperatures might discover better kernels, which is precisely what the repeated sampling experiments explore.

Models evaluated. The one-shot baseline covers seven models spanning frontier commercial, open-source, and reasoning-optimized categories:

GPT-4o: OpenAI's multimodal flagship model
OpenAI o1: OpenAI's reasoning-optimized model (generates internal chain-of-thought)
DeepSeek-V3: DeepSeek's large-scale general-purpose model
DeepSeek-R1: DeepSeek's reasoning-optimized model
Claude 3.5 Sonnet: Anthropic's mid-tier model
Llama 3.1-70B Instruct: Meta's 70B parameter open-source model
Llama 3.1-405B Instruct: Meta's 405B parameter open-source model

The inclusion of both reasoning and non-reasoning models is deliberate: the paper hypothesizes that reasoning capabilities might help with the complex multi-step reasoning required for kernel generation (e.g., reasoning about memory access patterns, thread hierarchy, numerical correctness), and the results (Table 1) partially support this—reasoning models (o1, R1) achieve the highest fast1 scores overall, though the advantage is primarily in reducing execution failures rather than functional correctness errors (as revealed by the error analysis in Figure 2).

The paper explores two strategies for improving LM kernel generation at test time without retraining or fine-tuning the model, leveraging the KernelBench environment's ability to automatically verify correctness and measure performance.

Repeated Sampling (Parallel Generation with fastp@k)

Mechanism. For each task, the LM is queried $k$ times independently with a high temperature (temperature = 1.6 for DeepSeek-V3, temperature = 0.7 for Llama 3.1-70B) to encourage diverse outputs. Each of the $k$ generated kernels is compiled, executed, correctness-checked, and timed independently. A task is considered "solved" at budget $k$ if at least one of the $k$ samples is both correct and achieves speedup > $p$ .

The motivation comes from the "Large Language Monkeys" observation (Brown et al., 2024, cited as reference [3]): many LM failure modes are stochastic—the model might produce a correct kernel 10% of the time and an incorrect one 90% of the time. By sampling $k$ times and taking the best, the probability of finding at least one correct-and-fast kernel increases. The metric fastp@k formalizes this: it is the fraction of tasks where the model's top- $k$ success rate is non-zero.

Temperature justification. The paper states that the chosen temperatures "allow sample diversity while ensuring quality," citing the Large Language Monkeys work [3]. The temperatures differ across models because different models have different calibration properties—DeepSeek-V3 can tolerate higher temperature (1.6) without generating incoherent code, while Llama 3.1-70B requires lower temperature (0.7) to maintain code quality.

Key findings from repeated sampling. Figure 4 shows that fast1@k improves as $k$ increases from 1 to 100 for both DeepSeek-V3 and Llama 3.1-70B across all three levels. The most dramatic improvement is DeepSeek-V3 on Level 2: from 4% fast1 at $k = 1$ (one-shot) to 37% at $k = 100$ . This means that while DeepSeek-V3 rarely generates a good Level 2 kernel on the first try, it has a 37% chance of generating at least one good kernel across 100 attempts. However, the paper notes a hard ceiling: "if a model has a very low inherent probability of solving a task, simply increasing the sampling budget has limited impact. For example, DeepSeek-V3 was never able to generate any correct solution for a group of 34 convolution variants in Level 1, even when attempting with 100 samples."

DeepSeek-R1 exclusion. The paper notes that repeated sampling is not applied to DeepSeek-R1 because "its API endpoint does not provide a temperature parameter"—R1's generation is deterministic (or near-deterministic) by design, making repeated sampling with identical prompts fruitless.

Mechanism. The KernelBench environment can capture three types of feedback after each generation:

Execution feedback (E): Raw stdout/stderr from compilation and execution, including nvcc compiler errors (syntax errors, undefined symbols, type mismatches), CUDA runtime errors (illegal memory accesses, out-of-bounds reads, kernel launch failures), Python-level errors (shape mismatches in tensor operations), and explicit timeout/deadlock detection ("Your kernel execution timed out").
Profiler feedback (P): When a kernel is correct, the system runs the PyTorch profiler to produce an operator-level timing breakdown, showing where time is spent (e.g., in the custom CUDA kernel vs. in surrounding PyTorch operations). This provides the model with the same information a human engineer would use to identify bottlenecks.
Previous generation (G): The model's own previously generated kernel code is always provided as context in the next turn.

The iterative refinement process is a multi-turn conversation with the LM. The initial prompt is identical to the one-shot baseline. After each generation, the system:

Compiles the generated code (capturing compilation errors if any).
If compilation succeeds, executes the kernel on GPU and checks correctness.
If the kernel is correct, runs the profiler and records wall-clock time.
Constructs a feedback message containing the previous generation G, the execution result (compilation errors OR correctness status + runtime OR timeout), and optionally the profiler breakdown P.
Prompts the model again with: "Here is your latest generation: [G]. Your generated architecture ModelNew and kernel was evaluated on GPU and checked against the reference architecture Model. Here is your Evaluation Result: [Feedback]. Name your new improved output architecture ModelNew..."

The process runs for $N$ turns. The metric fastp@N measures "the percentage of tasks where the model generated at least one functionally correct kernel that is p times faster than PyTorch Eager by turn N" – analogous to fastp@k but across sequential turns rather than parallel samples. The best kernel across all $N$ turns is selected for each task (not necessarily the last one).

Three iterative refinement configurations tested. The paper ablates the feedback content:

G only (previous generation only): The model sees its own previous output but receives no compiler errors or performance data. This tests whether the model can self-correct through pure reasoning about its own code—essentially, iterative prompting without environment feedback.
G + E (previous generation + execution feedback): The model sees compilation errors and correctness results. This is the minimum useful feedback: the model can fix syntax errors, memory violations, and correctness bugs iteratively.
G + E + P (previous generation + execution feedback + profiler): The model additionally sees profiler output when the kernel is correct. This is the richest feedback: the model can identify performance bottlenecks and target its optimization efforts.

Key findings from iterative refinement (Table 2). The combination G+E+P with DeepSeek-R1 on Level 2 achieves the most dramatic improvement: fast1 rises from 36% (one-shot baseline) to 72% after 10 turns. More granularly, Table 9 shows that correctness (fast0) improves even more dramatically: DeepSeek-R1 on Level 1 reaches 95% correctness (functional kernels on 95% of tasks within 10 turns) with G+E or G+E+P, up from 67% in the one-shot baseline. The paper attributes this primarily to execution feedback E helping models fix compilation and runtime errors: "execution failures are the most frequent failure mode in LM-generated kernels" (Section 5.1), and the compiler error messages provide actionable information for fixing syntax and memory issues.

However, the paper identifies a hard residual: "the remaining incorrect kernels almost always fail due to functional incorrectness, likely because correctness feedback is less granular than execution failure messages." Execution failures produce specific error messages (line numbers, variable names, error codes); correctness failures produce only "your outputs don't match the reference"—a signal that something is wrong but with no information about where or why. This limits the effectiveness of iterative refinement for the hardest correctness bugs.

Comparison of repeated sampling vs. iterative refinement (Table 2). Given a fixed budget of 10 inference calls, iterative refinement outperforms repeated sampling in 5 of 6 cases (3 levels × 2 models). The exception is DeepSeek-V3 on Level 2, where repeated sampling (@10) achieves 14% vs. iterative refinement's 5-7%. The paper attributes iterative refinement's general advantage to the feedback loop: "iterative refinement consistently improves performance across models and levels," because each turn can build on the lessons from previous failures. However, the effectiveness is inherently bounded by the base model's ability to interpret and act on feedback: "Llama-3.1 70B does not always benefit from having such information," while DeepSeek-R1 "consistently improves using feedback E and P."

Temperature and stochasticity. The paper acknowledges a limitation: iterative refinement experiments use "temperature = 0 to focus on the effect of iterating based on feedback rather than introducing variability." This means that for a given task and feedback sequence, the model's next generation is deterministic. If the model gets stuck in a loop—generating the same incorrect kernel repeatedly—temperature = 0 provides no escape mechanism. The examples in Appendix D.4 illustrate this: some problems see steady improvement over turns (the convolution kernel in Table 5 goes from 9.1ms to 1.13ms), while others see the model "consistently makes the same mistake and continually generates a functionally incorrect kernel with the same value errors" (Problem 54).

Infrastructure for multi-turn experiments. Appendix H.2 describes a pipelined, multi-processing GPU orchestrator system for running iterative refinement experiments at scale:

CPU parallelism: Multiple independent processes handle separate KernelBench tasks, running the multi-turn state machine logic. Only kernel execution requires GPU access.
GPU orchestration: A separate process manages GPU allocation using semaphores. Processes request GPU access when they have a kernel ready to compile and execute, and release it immediately after. This maximizes GPU occupancy across multiple concurrent refinement experiments.
Pre-compilation on CPU: Kernels are compiled with nvcc on CPU before GPU access is requested, so GPU time is spent only on execution and profiling, not compilation.
Error handling: The orchestrator handles CUDA illegal memory accesses and deadlocks (common with faulty kernel generations) by releasing processes and spawning new ones, with custom handlers to capture these errors without crashing the system.

This infrastructure is not the scientific contribution but is a necessary engineering component that makes the iterative refinement experiments feasible—without it, running 10-turn refinement on 250 tasks with multiple models and feedback configurations would be computationally prohibitive.

Hardware-Aware Conditioning: In-Context Examples vs. Hardware Specifications

The paper explores two methods for providing the LM with domain knowledge that might improve kernel quality, motivated by the observation that one-shot kernels often lack standard optimization techniques (Section 5.2.1) and that kernel performance varies across hardware (Section 4.4).

Few-Shot In-Context Examples of Optimization Techniques

Mechanism. The prompt is augmented with three in-context example pairs, each showing a PyTorch reference and an optimized CUDA kernel that demonstrates a specific hardware-efficiency technique:

GELU activation (Hendrycks and Gimpel, 2023): Demonstrates operator fusion—combining the mathematical operations of the GELU function (0.5 * x * (1.0 + tanh(0.7978845608028654 * (x + 0.044715 * x^3)))) into a single CUDA kernel that reads input from global memory once, performs all computations in registers, and writes output once. The alternative (PyTorch's default) would issue separate kernels for the multiplication, power, tanh, and addition operations, with intermediate results stored to and reloaded from global memory.
Tiled matrix multiplication (Mills, 2024): Demonstrates tiling—decomposing a large matrix multiplication into smaller tiles that fit in shared memory, reducing global memory accesses from O(N³) to O(N²) by having each thread block load a tile of the input matrices into shared memory once and then perform multiple computations on that tile.
Minimal FlashAttention (Dao et al., 2022; Kim, 2024): Demonstrates shared memory I/O management and the online softmax algorithm—carefully orchestrating data movement between global memory, shared memory, and registers to compute attention without materializing the full N×N attention matrix, using tiling and recomputation to keep memory footprint linear in sequence length.

Prompt structure (Appendix C.4). The prompt presents these examples after the initial format example (element-wise add) and before the task specification, following the pattern: "Here is an example architecture: [PyTorch reference]. Here is an optimized version with custom CUDA kernels: [Optimized CUDA implementation]." The task is then presented as before.

Key findings. The few-shot examples have a negative effect on overall fast1 compared to the one-shot baseline (Table 10). For OpenAI o1, fast1 drops from 10% to 6% on Level 1, from 24% to 16% on Level 2, and from 12% to 8% on Level 3. For Llama 3.1-70B, fast1 slightly increases on Level 1 (3% to 6%) but stays at 0% on Levels 2 and 3. The paper's explanation: "In-context examples degrade the LM's overall fast1 score since LMs attempt more aggressive optimization strategies, but result in more execution failures." o1's generated kernels are on average 25% longer with the few-shot examples, indicating that the model is attempting more complex code (tiling loops, shared memory management) but introducing more bugs in the process.

However, among kernels that are correct, the few-shot examples enable more sophisticated optimizations:

On 77% of GEMM variants in Level 1, o1 applies tiling and achieves speedup over the one-shot baseline (Table 11)—though these kernels remain slower than PyTorch Eager because they don't use tensor core instructions (CUDA wmma or mma operations), only CUDA cores.
On Level 2, o1 applies "aggressive shared memory I/O management" on 11 problems and outperforms PyTorch Eager on these tasks (Table 12), including a conv2d+InstanceNorm+Divide fusion achieving 0.082ms vs. 0.090ms for the baseline.

Interpretation. The few-shot examples teach the model what optimizations to attempt (fusion, tiling, shared memory management) but not how to implement them correctly in arbitrary contexts. The result is more ambitious but more error-prone generations—a tradeoff where the few-shot prompt shifts the model from "generate simple, often-correct kernels" to "generate complex, often-incorrect kernels with higher potential ceiling."

Specifying Hardware Information

Mechanism. The prompt is augmented with detailed hardware specifications for the target GPU, including (exact values from Appendix C.5):

GPU name (e.g., "NVIDIA H100")
Memory: "We have X GB GDDR6 with ECC of GPU Memory"
Memory bandwidth: "We have X GB/s of Memory Bandwidth"
Compute throughput: FP32 TFLOPS, TF32 Tensor Core TFLOPS, FP16 Tensor Core TFLOPS, FP8 Tensor Core TFLOPS, INT8/INT4 Tensor TOPS
Architectural limits: registers per SM, maximum registers per thread, maximum thread blocks per SM, shared memory capacity per SM (in KB), maximum shared memory per thread block (in KB)
Conceptual hardware knowledge: definitions of threads, thread blocks, shared memory, registers, memory hierarchy, memory bandwidth, cache, and HBM, plus best practices (parallelize sequential code, minimize host-device transfers, maximize device utilization, ensure coalesced global memory accesses, minimize redundant global memory accesses, avoid warp divergence, use specialized instructions based on GPU architecture).

Key findings. The hardware information has minimal impact on overall fast1 scores (Table 15, Appendix G.2). For most models, providing hardware specifications does not change the fast1 percentages significantly compared to the one-shot baseline. However, the paper observes an interesting behavioral shift for reasoning models:

"R1 attempts to generate warp matrix multiply-accumulate (wmma) instructions for approximately 50% of the Level 1 matrix multiplication problems, although most fail to compile."

The example in Figure 10 (Appendix G.2) shows a DeepSeek-R1 generated kernel for matrix multiplication that uses nvcuda::wmma fragments and wmma::load_matrix_sync / wmma::mma_sync operations—the correct API for using tensor cores on NVIDIA GPUs. The kernel attempts to use warp-level matrix multiply-accumulate, which would theoretically achieve much higher throughput than CUDA-core-based implementations. However, the generated code has subtle bugs (incorrect indexing logic, missing synchronization in some paths) that prevent compilation or correct execution.

Interpretation. The hardware information does induce the model to attempt hardware-specific optimizations, but the implementation quality is not high enough for these attempts to succeed reliably. The paper frames this as "highlighting room for improvement for future models" and notes that LMs are better at incorporating optimization techniques through concrete code examples (few-shot) than through abstract specifications (hardware specs)—suggesting that the training data distribution (few examples of hardware-specific kernel code) is the binding constraint.

Cross-hardware evaluation of one-shot kernels (Appendix G.1). Separately from providing hardware information, the paper evaluates the kernels generated in the one-shot baseline (Section 4.1) on multiple GPU types to assess portability. The results (Table 14, Figures 8–9) show:

Level 1 kernels are relatively stable across GPUs: DeepSeek-R1 achieves fast1 of 12% on L40S, 16% on H100 and A100, 15% on L4, 22% on T4, 12% on A10G.
Level 2 kernels show much higher variance: DeepSeek-R1 achieves 36% on L40S but 47% on A10G and 46% on T4, but only 38% on A100.

The paper interprets this variance as evidence that "one-shot LM-generated kernels may not generalize well across hardware"—the optimal kernel depends on hardware-specific properties (memory bandwidth, shared memory size, compute throughput), and LMs generating kernels without hardware awareness cannot adapt to these differences.

High-Throughput Evaluation Infrastructure

The paper describes a non-trivial engineering system for evaluating generated kernels at scale, organized as a three-stage pipeline (Appendix H.1) with a separate orchestrator for iterative refinement (Appendix H.2).

Three-stage pipeline for one-shot evaluations (Figure 11):

Inference (parallelized): LM queries are issued in parallel across many tasks, with generated kernels stored to disk.
CPU pre-compilation (parallelized): Each generated kernel is compiled with nvcc for the target GPU hardware into a binary. This stage is parallelized across CPUs, and each compiled binary is cached in a task-specific directory. The pre-compilation step is critical because it separates the CPU-bound compilation work from the GPU-bound execution work, enabling the GPU to be fully utilized for timing rather than waiting on compilation.
GPU evaluation (sequential per device, parallel across devices): With pre-compiled binaries, kernels are evaluated on GPUs. Only one kernel is evaluated at a time per GPU to ensure accurate timing (no interference from concurrent kernels), but multiple GPUs can evaluate different kernels in parallel. The paper's primary evaluations used a single bare-metal L40S for consistency, but the infrastructure supports multi-GPU deployment.

Orchestrator for iterative refinement (Appendix H.2). The multi-turn nature of iterative refinement creates different throughput challenges: each task requires sequential LM calls interspersed with compilation and GPU execution, and the timing of these steps is unpredictable (some kernels compile instantly, others hit errors that need to be caught). The orchestrator solves this with:

Process-per-task parallelism: Each iterative refinement experiment (one task, one model, one feedback configuration, N turns) runs as an independent process with its own state machine. The state machine transitions between "query LM → pre-compile on CPU → acquire GPU → execute and profile → release GPU → construct feedback → query LM again."
Centralized GPU allocation: A separate orchestrator process manages GPU access using semaphores. Worker processes request a GPU when they have a compiled kernel ready to execute, and the orchestrator grants access when a GPU is free. This maximizes utilization: while one process is executing a kernel, others can be querying LMs or compiling.
Error resilience: GPU kernel bugs can cause illegal memory accesses that crash the CUDA context or deadlocks that hang indefinitely. The orchestrator detects these conditions (timeouts for deadlocks, error codes for memory violations), releases the affected GPU (potentially requiring a CUDA context reset), and spawns new processes as needed. The paper reports writing "specifically handlers to ensure these errors are properly captured without crashing the orchestrator itself."

Visualization interface (Figure 12). The paper provides a UI for inspecting generated kernels, their performance, and their refinement trajectories. This is not a scientific contribution but a practical tool that enables the qualitative case studies in Section 6.1 and Appendix D—the paper's analysis of "interesting kernels" requires the ability to examine individual generations, compare them across techniques, and trace how they evolved during iterative refinement.

4. Key Insights and Innovations

Innovation 1: A Benchmark Where Success Directly Translates to Production Value—And Where the Baseline Is Expert-Tuned Closed-Source Code

The most distinctive intellectual move in this paper is the inversion of the typical benchmark design philosophy. Most code generation benchmarks construct self-contained problems (HumanEval, MBPP, DS-1000) or derive tasks from existing repositories (SWE-Bench) where the "ground truth" is a known correct solution. KernelBench instead defines each task as a PyTorch reference implementation and sets the success criterion as outperforming that reference on wall-clock time—where the reference itself relies on hand-tuned, often closed-source kernels (cuBLAS, cuDNN) developed by expert engineers over years. This means there is no provided "correct kernel" to generate; the benchmark's gold standard is not a pre-existing solution but rather the current production baseline.

This framing changes what it means for an LM to "succeed." In a typical benchmark, success means matching a known answer. In KernelBench, success means generating code that is immediately deployable as a drop-in performance improvement in real ML pipelines. The paper makes this explicit: "solving KernelBench tasks are immediately beneficial for real deep learning workloads." Benchmarks like HumanEval measure whether LMs can replicate programming patterns present in training data; KernelBench measures whether LMs can invent optimizations that haven't been open-sourced yet—a fundamentally harder and more practically meaningful task.

The choice to make PyTorch Eager the primary baseline—rather than a naive Python implementation—is a deliberate escalation of difficulty. PyTorch's operations already dispatch to highly optimized CUDA kernels. To achieve fast1, an LM must not merely produce a working kernel, but one that beats code written by NVIDIA's kernel engineers. The paper's results demonstrate why this matters: even frontier reasoning models achieve fast1 on less than 20% of tasks out of the box. A weaker baseline (e.g., naive nested-loop matrix multiplication) would produce superficially impressive speedups that have no production relevance. The paper's metric design forces the field to confront the gap between "can generate CUDA code" and "can generate CUDA code that competes with expert human engineers."

A more subtle aspect of this innovation is the dynamic difficulty of the benchmark over time. Because PyTorch is cross-hardware compatible, each new GPU generation resets the baseline—the same reference code runs faster on better hardware, raising the bar for what counts as a speedup. Similarly, as new ML architectures are proposed, new tasks can be added. Most benchmarks eventually saturate (models approach ceiling performance) and must be retired. KernelBench is designed so that progress in hardware and ML architecture automatically makes the benchmark harder, ensuring it remains a moving target. This is a conceptual shift from "build a test set that measures current capabilities" to "build an evaluation framework that continuously tracks the frontier of practical value."

The paper does not just assert this design philosophy—it validates it through the cross-hardware experiments (Section 4.4, Table 14). DeepSeek-R1's fast1 on Level 2 varies from 36% on L40S to 47% on A10G, demonstrating that kernel quality is hardware-dependent and that the benchmark captures this. A kernel that "succeeds" on one GPU may fail on another, reflecting the real engineering challenge that human kernel developers face when porting across platforms. This property makes KernelBench more than a one-time evaluation; it is infrastructure for studying how LMs handle the hardware-portability challenge that the paper identifies as a core pain point in the AI ecosystem.

Innovation 2: Identifying Correctness, Not Execution Errors, as the Hard Residual Failure Mode—and Connecting It to a Measurable Data Scarcity

The paper makes a diagnostic contribution that reframes the challenge of LM kernel generation: execution errors (compilation failures, memory violations, runtime crashes) are not the fundamental bottleneck—functional correctness errors are. This finding emerges from the error analysis in Figure 2, which decomposes LM failures into execution failures and correctness errors. Reasoning models (o1, R1) substantially reduce execution failures compared to non-reasoning models, but "all LMs struggle with functional correctness to a similar degree." The iterative refinement results (Section 5.1.2) corroborate this: with execution feedback, DeepSeek-R1 achieves over 90% functional kernel generation on Levels 1 and 2 (Table 9)—meaning it can fix almost all execution errors—yet fast1 peaks at 72% on Level 2, because the remaining kernels are compiled and run but produce wrong answers.

This is a non-obvious finding. The natural first-order assumption when seeing that LMs produce many CUDA compilation errors would be "the models don't know CUDA syntax well enough." The paper shows that this is partially true but fixable through iterative feedback—compiler errors are specific and actionable, and LMs can use them to correct syntax, type errors, and memory access patterns. The deeper problem is that correctness errors produce impoverished feedback: the environment reports "your outputs don't match the reference" with no information about where or why the numerical error occurred. This means the iterative refinement loop, which is highly effective at eliminating execution failures, stalls on the harder problem of logical correctness.

The paper connects this correctness bottleneck to a measurable, structural cause: CUDA's scarcity in open-source training data. The specific figure—0.073% of The Stack v1.2—is more than a statistic; it is a causal hypothesis about why functional correctness is harder than syntax. Python and general-purpose programming languages are well-represented in training corpora, so LMs have internalized their semantics through exposure to millions of diverse examples. CUDA, at less than one-tenth of one percent of the corpus, provides vastly fewer examples from which to learn the subtle semantics of parallel execution (warp-level synchronization, memory consistency models, numerical behavior of different floating-point formats). The paper's qualitative case studies support this: when LMs attempt aggressive optimizations (tensor core wmma instructions, complex shared memory management), they often produce code that compiles but computes wrong answers—suggesting they can mimic the syntactic patterns of advanced CUDA from limited training examples but haven't internalized the deeper semantic constraints.

The innovation here is not the data scarcity observation itself (many domains are underrepresented in training data) but the specific diagnostic chain: functional correctness errors are the hard failure mode → these errors resist iterative feedback because correctness signals are impoverished → the root cause is likely training data scarcity in a domain where expertise is already rare → this creates a chicken-and-egg problem (we need more open-source CUDA to train better models, but better models are needed to generate more open-source CUDA). This chain reframes the research agenda: rather than focusing primarily on better prompting or reasoning strategies (which help with execution errors but not correctness), progress may require data-centric interventions—curating high-quality CUDA training data, developing synthetic data generation pipelines, or designing training objectives that better leverage the limited available examples.

The ReST^EM negative result (Appendix K) reinforces this diagnosis in an unexpected way. Attempting to optimize the revision model using on-policy RL training caused performance to degrade, which the authors attribute to "spurious correlations in revision data" that the limited training signal couldn't overcome. This suggests that naive self-improvement loops—a popular direction in LM research—may be especially brittle in low-resource domains like CUDA, where the model lacks sufficient prior knowledge to distinguish genuine improvements from artifacts of the training procedure.

Innovation 3: A Unified Metric for the Joint Correctness-Speedup Objective That Makes Explicit What Most Benchmarks Leave Implicit

The fastp metric appears simple on the surface—a percentage of tasks that are correct and achieve speedup > p—but its design reflects a non-trivial insight about how to evaluate code generation when performance is part of the objective. The dominant evaluation paradigm in code generation is pass@k (Chen et al., 2021), which measures correctness only. Recent work on code efficiency (Nichols et al., 2024; Waghjale et al., 2024) has begun measuring asymptotic complexity or runtime, but these metrics are typically reported separately from correctness—a model might achieve 90% correctness with 2× speedup, while another achieves 50% correctness with 10× speedup, and it is unclear how to rank them.

fastp solves this by logically coupling correctness and speedup in a single indicator function. A task contributes to the score only if both conditions hold. This means the metric automatically penalizes two distinct failure modes that separate reporting would obscure: (1) generating fast but incorrect kernels (high average speedup, low fastp) and (2) generating correct but slow kernels (high correctness, low fastp). The threshold p provides an explicit knob for setting the performance bar: at p = 0, fastp equals correctness rate; at p = 1, it measures "correct and faster than baseline"; at p = 2, "correct and at least 2× faster." This parameterization makes the tradeoff between speedup magnitude and correctness explicit and adjustable.

The innovation is less the mathematical form (a logical AND with an inequality) and more the evaluation philosophy it encodes: in production-oriented code generation, a partial success is not a success. A kernel that is "mostly correct but produces wrong outputs on some edge cases" is not useful, regardless of how fast it is. A kernel that is "perfectly correct but slower than the default" provides no incentive to adopt it. fastp operationalizes this binary standard: the kernel either works in practice or it doesn't. This contrasts with metrics that average speedups across all generated kernels (which can be dominated by a few fast-but-correct outliers) or report correctness and speedup on separate axes (which forces the reader to mentally integrate them).

The paper's fastp curves (Figure 3) demonstrate the metric's diagnostic value. As p increases from 0 to 4, the curves for different models diverge in informative ways. Some models maintain high correctness at p = 0 but drop sharply as p increases (many correct but slow kernels); others maintain a flatter curve (most correct kernels are also fast). This curve shape provides a richer characterization of model capability than a single-number metric—it reveals not just "how many kernels are good" but "how fast are the good kernels?"

The extension to fastp@k and fastp@N (for repeated sampling and iterative refinement) is a natural but important generalization that applies the same joint-condition logic to test-time scaling. Just as pass@k measures whether any of k samples is correct, fastp@k measures whether any of k samples is both correct and fast enough. This enables direct comparison between parallel sampling and sequential refinement under a matched compute budget, which the paper exploits in Table 2 to show that iterative refinement with feedback is generally more effective than repeated sampling at the same budget.

The paper's iterative refinement experiments (Section 5.1.2) do more than demonstrate that feedback helps—they decompose what different types of feedback contribute and reveal that execution feedback and profiler feedback serve fundamentally different roles. This is a conceptual contribution about the nature of the kernel optimization problem: fixing errors and improving performance are different cognitive tasks that respond to different information signals.

The key empirical pattern is visible in Table 9 and Figure 6. Execution feedback (G+E) drives a dramatic improvement in correctness: DeepSeek-R1 on Level 1 goes from 67% to 95% fast0, and on Level 2 from 62% to 85% fast0. But adding profiler feedback (G+E+P) doesn't further improve correctness—it actually slightly reduces it on Level 1 (95% to 95%, no change) and Level 3 (50% to 42%, slight decrease), while helping on Level 2 (85% to 92%). The paper attributes the correctness ceiling to the impoverished nature of correctness feedback ("your outputs don't match" vs. specific compiler error messages), but the more revealing finding is what happens to performance: on Level 2, G+E achieves 62% fast1 while G+E+P achieves 72%—a 10 percentage point gain from profiler data alone.

The profiler provides operator-level timing breakdowns: it shows whether the bottleneck is in the custom kernel itself (which might suggest further optimization), in data movement between host and device, or in surrounding PyTorch operations that could also be fused. When a kernel is correct but slow, the profiler tells the model where to direct its optimization effort. Without the profiler, the model has no signal about why a correct kernel is slow—it might optimize the wrong component or introduce unnecessary complexity that degrades performance elsewhere. The profiler closes this loop: it transforms the binary signal "your kernel is slower than baseline" into actionable bottleneck identification.

This decomposition has implications beyond kernel generation. It suggests a general principle for LM-based code optimization systems: verification feedback (correctness, compilation) enables error correction; profiling feedback (performance bottlenecks) enables optimization; both are needed for the full generate-fix-accelerate loop. The paper doesn't state this as an explicit principle, but the experimental design—ablating E and P separately—makes the evidence clear. Current LM coding benchmarks typically provide only verification feedback (correct/incorrect), which limits LMs to correctness-oriented improvement. KernelBench's integration of profiler feedback demonstrates what becomes possible when richer performance signals are available.

The paper also identifies the limits of this approach through negative examples. Appendix D.4.3 shows Problem 54 (a 3D convolution) where DeepSeek-R1 with G+E+P generates kernels that compile and run correctly across 10 turns but are never functionally correct—the model "consistently makes the same mistake and continually generates a functionally incorrect kernel with the same value errors." This illustrates the correctness-feedback poverty problem: when the error is subtle enough that neither compiler messages nor binary correctness signals provide diagnostic information, the iterative refinement loop stalls regardless of how many turns are available. This negative result is as informative as the positive ones—it defines the boundary condition where current feedback mechanisms fail, pointing toward future work on more granular correctness signals (e.g., intermediate value checking, differential testing, or formal verification).

5. Experimental Analysis

Evaluation Methodology

Dataset. KernelBench consists of 250 tasks partitioned into three levels: Level 1 (100 tasks, single primitive operations such as matrix multiplications, convolutions, activations, normalizations, and losses), Level 2 (100 tasks, sequences of 3–6 operations that can benefit from operator fusion), and Level 3 (50 tasks, full end-to-end architectures including AlexNet and MiniGPT collected from popular PyTorch repositories on GitHub). All tasks are specified as PyTorch reference implementations with Model classes and get_inputs/get_init_inputs functions defining exact tensor shapes and data types. There is no separate train/test split since KernelBench is evaluation-only; no ground-truth kernels are provided.
Base model(s). Seven models spanning frontier commercial, open-source, and reasoning-optimized categories are evaluated: GPT-4o, OpenAI o1, DeepSeek-V3, DeepSeek-R1, Claude 3.5 Sonnet, Llama 3.1-70B Instruct, and Llama 3.1-405B Instruct. The paper states that these models "are representative of the capabilities of many contemporary LLMs." Reasoning models (o1, R1) are of particular interest because kernel generation requires multi-step reasoning about hardware constraints, memory access patterns, and numerical correctness. The paper does not fine-tune or adapt any model for kernel generation; all evaluations test out-of-the-box or prompt-engineered capabilities.
Metrics. The primary metric is fastp, defined as the fraction of tasks for which the LM-generated kernel is both functionally correct (matching the reference Model output on all 5 randomly generated input sets) and achieves a speedup greater than threshold p over the PyTorch Eager baseline (where speedup = T_Model / T_ModelNew). Formally: fastp = (1/N) * Σ 𝐈(correct_i ∧ {speedup_i > p}). The paper focuses on fast1 (threshold of 1×, requiring any speedup over baseline) as the primary evaluation point, with fast0 measuring correctness alone. For repeated sampling experiments, fastp@k extends this to measure whether at least one of k independent samples achieves the joint condition. For iterative refinement, fastp@N measures whether at least one generation across N sequential turns achieves the condition.
Baselines. Two PyTorch execution modes serve as performance baselines: PyTorch Eager (the primary baseline throughout the paper), which is PyTorch's default execution mode dispatching operations to highly optimized closed-source kernels (cuBLAS, cuDNN), and torch.compile (PyTorch 2.0's ahead-of-time compiler with the Inductor backend in default configuration, plus additional configurations tested in Table 4 including reduce-overhead, max-autotune, max-autotune-no-cudagraphs, and cudagraphs). Compile time is excluded from torch.compile timing measurements. The paper notes that torch.compile is "sometimes slower than Torch Eager" on Level 1 due to reproducible runtime overhead significant for small kernels.
Generation budget / compute accounting. For one-shot evaluation, each model generates exactly one kernel per task via greedy decoding (temperature = 0). For repeated sampling (Section 5.1.1), the budget is k independent samples per task, with k swept up to 100, using high temperature (1.6 for DeepSeek-V3, 0.7 for Llama 3.1-70B). For iterative refinement (Section 5.1.2), the budget is N sequential turns (with N = 10 used for main experiments), where each turn involves one LM generation followed by compilation, execution, and feedback. The paper compares repeated sampling and iterative refinement at equal budget of 10 inference calls (Table 2). All evaluations are conducted on a bare-metal NVIDIA L40S GPU with 48 GB HBM, Ada Lovelace architecture, Python 3.10, PyTorch 2.5.0+cu124, CUDA 12.4.
Cross-validation / statistical protocol. The paper does not employ cross-validation or statistical significance testing. Correctness is evaluated on exactly 5 random inputs per task, with the choice empirically justified (Appendix B.2): in a 100-kernel experiment, all kernels were either 5/5 and 100/100 correct or 0/5 and 0/100 correct, with no partial correctness observed. Performance timing uses 3 warm-up iterations followed by 100 timed iterations with torch.cuda.Event markers; the paper reports mean wall-clock time with coefficient of variation consistently below 3%. Iterative refinement experiments use temperature = 0 to isolate the effect of feedback from stochastic variation.

Main Quantitative Results

One-Shot Baseline Performance (Table 1, Figure 2, Figure 3)

The headline finding from Table 1 is that out-of-the-box, frontier reasoning models achieve the best results but still fall dramatically short: DeepSeek-R1 achieves the highest fast1 scores across levels (12% on Level 1, 36% on Level 2, 2% on Level 3), followed by OpenAI o1 (10%, 24%, 12%). Non-reasoning models perform substantially worse: GPT-4o achieves 4%, 5%, 0%; Claude 3.5 Sonnet achieves 10%, 7%, 2%; DeepSeek-V3 achieves 6%, 4%, 8%. The largest open-source model tested, Llama 3.1-405B, achieves only 3%, 0%, 2%, which is nearly identical to its 70B counterpart (3%, 0%, 0%), suggesting that model scale alone does not confer kernel generation capability in this domain.

Against the torch.compile baseline, fast1 scores are higher for reasoning models on Level 1 (DeepSeek-R1: 38% vs. 12% against Eager) and Level 2 (37% vs. 36%), but lower for non-reasoning models and on Level 3. The paper attributes this pattern to torch.compile being "sometimes slower than Torch Eager" on Level 1 due to runtime overhead, meaning it is an easier baseline to beat on those tasks. The primary analysis throughout the paper therefore uses PyTorch Eager as the reference.

Error analysis (Figure 2). The paper decomposes failure modes into execution failures (CUDA/nvcc compiler errors, CUDA memory violations, Python runtime errors) and functional correctness errors (output tensor shape mismatches and value mismatches). The key finding: "reasoning LMs (o1, R1) produce fewer incorrect solutions (< 55%) than other models (> 70%). However, we find this is mainly because they make fewer execution failures. All LMs struggle with functional correctness to a similar degree." This is the central diagnostic insight of the one-shot evaluation: reasoning helps with syntax and memory correctness but does not substantially improve logical correctness of the kernel computation itself.

Speedup distribution (Figure 3). The fastp curves as p varies from 0 to approximately 4 show that fast0 (correctness alone) is substantially higher than fast1 for all models—many generated kernels are correct but slower than PyTorch Eager. For example, DeepSeek-R1 on Level 1 achieves approximately 67% fast0 (Table 9) but only 12% fast1 (Table 1), meaning that of the two-thirds of tasks where it generates a correct kernel, fewer than one-fifth of those are actually faster than the baseline. The curves decline monotonically with increasing p, and reasoning models maintain higher scores across the full threshold range, but the absolute numbers remain low: at p ≈ 2, even DeepSeek-R1 drops below 10% on Level 1 and near 0% on Level 3.

Performance variation across hardware (Table 14, Figures 8–9). When the kernels generated in the one-shot baseline (without hardware-specific information) are evaluated on different GPU types, Level 1 kernels show relatively stable speedups (DeepSeek-R1 fast1 ranges from 12% on L40S/A10G to 22% on T4). Level 2 kernels exhibit substantially higher variance: DeepSeek-R1 achieves 36% on L40S but 47% on A10G and 46% on T4, while achieving 38% on A100 and 42% on H100. The paper interprets this as evidence that "one-shot LM-generated kernels may not generalize well across hardware"—the optimal kernel depends on hardware-specific properties, and LMs generating kernels without hardware awareness produce code that is accidentally tuned to certain architectures.

Repeated Sampling (Figure 4, Table 2)

The headline finding is that repeated sampling with high temperature discovers more correct and performant kernels, but its effectiveness is bounded by the base model's inherent capability. Figure 4 shows fast1@k as a function of sample budget k (up to 100) for DeepSeek-V3 and Llama 3.1-70B across all three levels.

For DeepSeek-V3:

Level 1: fast1@k improves from 6% at k = 1 to 11% at k = 100. The improvement is modest and plateaus quickly.
Level 2: fast1@k improves dramatically from 4% at k = 1 to 37% at k = 100. This is the most significant gain—a >9× improvement in the fraction of solved tasks through parallel sampling alone.
Level 3: fast1@k improves from 8% at k = 1 to 14% at k = 100, a modest gain.

For Llama 3.1-70B:

Level 1: fast1@k improves from 3% at k = 1 to 5% at k = 100.
Level 2: fast1@k improves from 0% at k = 1 to 3% at k = 100.
Level 3: fast1@k improves from 0% at k = 1 to 1% at k = 100.

The paper explicitly notes a hard ceiling: "if a model has a very low inherent probability of solving a task, simply increasing the sampling budget has limited impact. For example, DeepSeek-V3 was never able to generate any correct solution for a group of 34 convolution variants in Level 1, even when attempting with 100 samples." This means that for problem classes where the base model's probability of correctness is effectively zero, no amount of parallel sampling helps.

The headline finding from Table 2 is that iterative refinement with execution and profiler feedback substantially outperforms both the one-shot baseline and repeated sampling at matched budget (10 inference calls) for most model-level combinations. The most dramatic result is DeepSeek-R1 on Level 2 with G+E+P feedback: fast1@10 reaches 72%, compared to 36% for the one-shot baseline and an implied ~37–38% for repeated sampling at k = 10 (extrapolating from the k = 100 curve in Figure 4, where fast1@100 = 37%).

Comparing iterative refinement configurations within DeepSeek-R1:

Level 1: G only (18%), G+E (41%), G+E+P (43%). Execution feedback provides a >2× improvement over self-reflection alone; profiler adds marginal gain.
Level 2: G only (44%), G+E (62%), G+E+P (72%). The progression shows that execution feedback adds 18 percentage points and profiler adds another 10—both signals are independently valuable.
Level 3: G only (4%), G+E (12%), G+E+P (18%). Even with full feedback, absolute performance remains low on full architectures.

Critically, the effectiveness is model-dependent. DeepSeek-R1 consistently improves with richer feedback; DeepSeek-V3 shows mixed results (G+E sometimes reduces fast1 compared to G alone on Level 2: 7% vs. 13%); and Llama 3.1-70B shows minimal gains across all configurations (best fast1 is 8% on Level 2 with G+E+P, up from 0% baseline). The paper states: "the effectiveness of the test-time methods is inherently dependent on the quality of the base model."

Correctness trajectory (Table 9). The fast0@10 results reveal that iterative refinement with execution feedback can achieve near-complete functional correctness on some levels: DeepSeek-R1 reaches 95% fast0 on Level 1 and 92% on Level 2 with G+E+P, up from 67% and 62% in the one-shot baseline. However, on Level 3, fast0 reaches only 50%—half of the full architectures still cannot be generated correctly even with iterative feedback. The paper attributes the residual errors to functional incorrectness that "correctness feedback is less granular than execution failure messages," meaning the model receives binary "wrong output" signals without diagnostic information about where the computation went wrong.

Figure 6 (Level 2 trajectory). The fast1@N curve for DeepSeek-R1 on Level 2 with G+E+P shows consistent improvement across turns: starting at approximately 38% after turn 1, rising to approximately 55% by turn 4, and reaching 72% by turn 10. The curve does not plateau within the observed range, suggesting further gains might be possible with more turns, though the paper does not explore budgets beyond 10.

Comparison of methods at matched budget (Table 2). With a budget of 10 inference calls, iterative refinement outperforms repeated sampling in 5 of 6 cases (3 levels × 2 models). The exception is DeepSeek-V3 on Level 2, where repeated sampling achieves 14% vs. iterative refinement's 5–7%. The paper's interpretation: "iterative refinement consistently improves performance across models and levels" because each turn can build on previous failures, but the gains are constrained by the model's ability to interpret and act on feedback—which is why DeepSeek-R1 benefits most.

Hardware-Aware Conditioning (Tables 10–12, Table 15)

Few-shot in-context examples (Tables 10–12). Providing three in-context examples demonstrating fusion (GELU), tiling (matrix multiplication), and shared memory I/O management (minimal FlashAttention) produces a negative overall effect on fast1 for reasoning models while showing marginal improvements for weaker models. For OpenAI o1, fast1 drops from 10% to 6% on Level 1, 24% to 16% on Level 2, and 12% to 8% on Level 3 (Table 10). The paper attributes this to LMs attempting "more aggressive optimization strategies, but resulting in more execution failures"—o1's generated kernels are on average 25% longer with few-shot examples.

However, among the subset of correct kernels, the few-shot examples enable more sophisticated optimizations:

On 77% of GEMM variants in Level 1, o1 applies tiling and achieves speedup over the one-shot baseline, though these kernels remain slower than PyTorch Eager because they use CUDA cores rather than tensor cores (Table 11).
On Level 2, o1 applies "aggressive shared memory I/O management" on 11 problems where it outperforms PyTorch Eager, including a conv2d+InstanceNorm+Divide fusion achieving 0.082 ms vs. 0.090 ms for the baseline (Table 12).

Hardware specification prompts (Table 15). Providing hardware-specific information (GPU type, memory sizes, bandwidth, TFLOPS, register/shared memory capacities, plus conceptual hardware knowledge) has minimal impact on overall fast1 scores. For DeepSeek-R1, fast1 with hardware specs is 14% on Level 1 (vs. 12% baseline), 33% on Level 2 (vs. 36%), and 0% on Level 3 (vs. 2%) on the L40S. The paper observes a behavioral shift: "R1 attempts to generate warp matrix multiply-accumulate (wmma) instructions for approximately 50% of the Level 1 matrix multiplication problems, although most fail to compile." An example generation in Figure 10 shows R1 producing code that correctly imports <mma.h> and uses nvcuda::wmma::fragment types and wmma::load_matrix_sync/wmma::mma_sync operations—the correct API for tensor core programming—but with indexing bugs that prevent compilation or correct execution. The paper concludes that "LMs are better at adjusting their approaches when provided with few-shot examples... than with hardware information," suggesting that concrete code demonstrations are more effective than abstract specifications for this domain.

Ablation Studies and Robustness Checks

Feedback content in iterative refinement (Table 2, Table 9): Ablating the feedback provided during iterative refinement shows that execution feedback (E) is the primary driver of correctness improvements. For DeepSeek-R1 on Level 1, fast0@10 goes from 72% (G only) to 95% (G+E) to 95% (G+E+P)—execution feedback alone achieves nearly all the correctness gain. For performance (fast1), profiler feedback provides additional gains: on Level 2, G+E achieves 62% while G+E+P achieves 72%, a 10 percentage point improvement. However, profiler feedback can slightly reduce fast0 on some levels (Level 3: 50% for G+E vs. 42% for G+E+P), suggesting that providing timing information may distract the model from correctness fixes when both objectives are in tension.
Number of correctness check inputs (Appendix B.2): The paper empirically validates the choice of 5 random inputs for correctness checking. In an experiment with 100 generated kernels evaluated at both 5 and 100 inputs, "50 kernels were correct (all 5/5 and 100/100), 19 had output value mismatches (19 0/5 and 0/100), 4 had output shape mismatches, 10 encountered runtime errors, and 17 had compilation errors. Notably, the 0/5 and 0/100 failures indicate that no partial correctness was observed." This all-or-nothing pattern means 5 inputs provide sufficient discrimination without the computational cost of 100.
PyTorch baseline configuration (Table 4): The paper evaluates fast1 against multiple torch.compile configurations beyond the default: cudagraphs, max-autotune, max-autotune-no-cudagraphs, and reduce-overhead. For reasoning models, fast1 varies substantially across configurations. DeepSeek-R1 on Level 2 achieves 37% against default torch.compile, 52% against cudagraphs, but only 29% against max-autotune and 28% against reduce-overhead. This variation highlights that torch.compile's performance is configuration-dependent and that the choice of baseline affects reported fast1 scores. The paper focuses on PyTorch Eager for its primary analysis due to "the variability of torch.compile across configurations."
Cross-hardware evaluation of generated kernels (Table 14, Figures 8–9): The paper evaluates the same one-shot generated kernels (without hardware-aware prompting) across six GPU types. For DeepSeek-R1 on Level 2, fast1 varies from 36% (L40S, L4) to 47% (A10G) to 46% (T4). The individual-problem speedup plots (Figures 8–9) reveal that specific kernels achieve very different speedups on different GPUs—some kernels that are faster than PyTorch Eager on one GPU are slower on another. This serves as a robustness check demonstrating that the one-shot kernels are not implicitly tuned to a particular GPU, and that hardware portability is a real challenge for LM-generated kernel code.
Model scale effect (Llama 3.1-70B vs. 405B, Table 1): Comparing Llama 3.1-70B and Llama 3.1-405B shows nearly identical fast1 scores: 3% vs. 3% on Level 1, 0% vs. 0% on Level 2, 0% vs. 2% on Level 3. This is a notable negative result: ~6× more parameters in the same model family produce no meaningful improvement on kernel generation performance, suggesting that the capability bottleneck is not raw model capacity but rather domain-specific knowledge or reasoning ability that does not scale with parameter count in this model family.
Temperature for repeated sampling (Section 5.1.1): The paper chooses different temperatures for different models in repeated sampling (1.6 for DeepSeek-V3, 0.7 for Llama 3.1-70B), citing the Large Language Monkeys work for calibration. The paper does not ablate across temperature values, so it is unknown whether these choices are optimal. DeepSeek-R1 is excluded from repeated sampling entirely because "its API endpoint does not provide a temperature parameter," which limits the comparison—R1's parallel sampling potential remains unmeasured.

Critical Assessment

The experiments presented in Sections 4 and 5 demonstrate that KernelBench is a genuinely challenging benchmark where current frontier models perform poorly out of the box, and that test-time strategies (repeated sampling, iterative refinement) can substantially improve performance, particularly when execution feedback is available. However, several important caveats qualify what these experiments actually demonstrate versus what the paper claims.

Does the paper demonstrate that "frontier reasoning models perform the best out of the box but still fall short overall"? Yes, this is well-supported. Table 1 shows DeepSeek-R1 and OpenAI o1 achieving the highest fast1 scores across all three levels, with a substantial gap over non-reasoning models. The "falling short" claim is supported by the absolute numbers: even the best model achieves only 36% fast1 on the easiest level (Level 2 for these models) and near-zero on Level 3. Figure 2 provides mechanistic evidence for why reasoning models are better: they make fewer execution errors. However, the paper does not fully explain whether this advantage comes from the reasoning models' training (which may include more CUDA code), their inference-time chain-of-thought (which may help catch syntax errors before emitting code), or both. The error analysis in Figure 2 is based on binary classification of failure types and does not decompose execution errors further (e.g., are they syntax errors vs. memory violations vs. launch configuration errors?), which limits diagnostic precision.

Does the paper demonstrate that "leveraging execution and profiling feedback during iterative refinement" improves results? Yes, and this is the strongest empirical contribution. Table 2 shows that G+E+P consistently outperforms the one-shot baseline across models and levels, with DeepSeek-R1 on Level 2 improving from 36% to 72%—a doubling of fast1. The ablation of feedback types (G only vs. G+E vs. G+E+P) cleanly isolates the contributions of execution feedback and profiler feedback. Figure 6 shows that fast1 continues improving over 10 turns without plateauing. However, there is an important limitation: all iterative refinement experiments use temperature = 0, making the trajectory deterministic given the initial generation and feedback sequence. This means the results measure the best case for a deterministic refinement process—if the model gets stuck (as with Problem 54 in Appendix D.4.3), there is no stochastic escape mechanism. A fair comparison against repeated sampling would ideally allow both methods to use temperature > 0, but the paper chose to isolate the feedback effect. This makes the comparison in Table 2 somewhat asymmetric: repeated sampling benefits from stochasticity, while iterative refinement benefits from feedback. An experiment combining both (iterative refinement with temperature > 0) would reveal whether the two mechanisms are complementary or redundant, but was not conducted.

Does the paper demonstrate that "KernelBench remains a challenging benchmark, with its difficulty increasing as we raise speedup threshold p"? This claim is definitionally true (higher p means a harder criterion) but the experimental evidence about how difficulty scales is interesting. Figure 3 shows that fastp curves decline monotonically for all models, but the shape of decline varies: some models maintain a non-zero fastp out to p = 4 on Level 2 (notably DeepSeek-R1), while others drop to zero by p ≈ 1.5. The paper does not analyze why certain models produce larger speedups when they do succeed—whether it's because they attempt more aggressive optimizations (which succeed occasionally) or because they are better at implementing standard optimizations. The few-shot experiment (Section 5.2.1) provides circumstantial evidence that more aggressive optimization attempts lead to higher-variance outcomes (sometimes much faster, often incorrect), but this connection is not explored systematically across models.

Does the paper genuinely support the claim that LM-generated kernels "directly translate to faster practical kernels"? This is partially supported by the qualitative case studies in Section 6.1 and Appendix D. The paper identifies specific kernels with meaningful speedups: a 13× speedup on diagonal matrix multiplication (algorithmic optimization exploiting sparsity), 2.9× on GELU (operator fusion), 2.6× on a fused matmul+divide+sum+scale sequence, 2.8× on cosine similarity loss (shared memory management), and 1.9× on triplet margin loss (shared memory reduction). These are concrete examples where the generated kernel is both correct and substantially faster than PyTorch Eager, and the optimizations are semantically meaningful (not measurement artifacts). However, these are isolated successes—the paper does not report what fraction of fast1 kernels achieve 2× vs. 1.01× speedup, so the claim of "practical" value depends on whether marginal speedups (e.g., 1.05×) are considered practically useful. The paper acknowledges this implicitly by making p adjustable and focusing on p = 1 as a starting point.

What experiments are missing that would have strengthened the paper?

No combination of iterative refinement with repeated sampling (stochastic refinement). The most natural extension of the paper's findings—using temperature > 0 during iterative refinement to combine the exploration benefits of repeated sampling with the feedback benefits of refinement—is not tested. This leaves open the question of whether the 72% fast1 on Level 2 for DeepSeek-R1 represents an upper bound or could be pushed higher.
No fine-tuning or training-based adaptation experiments. The paper exclusively studies prompt-based and test-time strategies. Given the paper's own finding that CUDA is 0.073% of common training corpora, a natural experiment would be to fine-tune a base model on available CUDA code (e.g., from open-source kernels, the CUDA programming guide, or synthetic data) and measure improvement on KernelBench. The absence of any training-based baselines leaves the data scarcity hypothesis unvalidated through intervention.
No latency or wall-clock analysis for the generation and refinement process itself. The fastp metric measures kernel execution time but does not account for the time or compute cost of generating the kernel. In iterative refinement with 10 turns, the LM is queried 10 times—for a reasoning model like DeepSeek-R1, this could represent minutes of inference time. If generating the kernel takes 100× longer than the time saved by the speedup, the net practical benefit is negative. The paper does not discuss this generative cost-performance tradeoff.
No evaluation of kernels generated in alternative programming abstractions. The paper notes that LMs could use Triton, CUTLASS, ThunderKittens, or PTX, but all reported results are for inline CUDA-C kernels. It is unknown whether models would perform better if prompted to use Triton (a higher-level, more Pythonic language for GPU programming that is increasingly popular) or whether the CUDA focus artificially depresses performance because of the data scarcity issue. This is listed as future work (Section 6.3) but its absence limits the generalizability of the current results.
Single hardware platform for primary evaluation. All main results use the NVIDIA L40S GPU. Cross-hardware experiments (Section 4.4) evaluate one-shot kernels on multiple GPUs but do not run iterative refinement or repeated sampling on non-L40S hardware, and do not evaluate the few-shot or hardware-spec prompting experiments on the full GPU suite. This leaves open whether the 72% fast1 on Level 2 with DeepSeek-R1 on L40S would generalize to an H100 or A100.
Small number of correctness check inputs without formal verification. The paper uses 5 random inputs for correctness, validated empirically on a 100-kernel subset. This is a reasonable practical choice but provides no formal guarantees. A kernel that passes 5 tests could still be incorrect on edge cases (specific input magnitudes causing numerical overflow, boundary conditions in tiling when dimensions are not multiples of tile size). The paper acknowledges this limitation explicitly but does not discuss whether any kernels that pass 5 inputs were subsequently found to be incorrect under more extensive testing.

Conditional nature of the claims. The paper's claims about the effectiveness of iterative refinement are conditional on the base model's capability: the method works well for DeepSeek-R1 (72% on Level 2) but poorly for Llama 3.1-70B (4% on Level 2). The paper is transparent about this ("the effectiveness of the test-time methods is inherently dependent on the quality of the base model") but does not characterize what "quality" means in this context—is it reasoning ability, CUDA knowledge, general code generation capability, or something else? The cross-model comparison in Table 2 shows that DeepSeek-R1 > DeepSeek-V3 > Llama 3.1-70B for iterative refinement effectiveness, which aligns with the one-shot ranking in Table 1, suggesting that base capability and feedback-utilization ability are correlated but not cleanly separated by the experimental design.

The claim that "models demonstrate potential to produce performant kernels via optimizations" is supported by the case studies but qualified by their rarity. The paper does not report the full distribution of speedup magnitudes for correct kernels—it provides box plots in Figure 7 showing medians below 1.0 for most models on Levels 1 and 3, and only slightly above 1.0 for Level 2. This means that even among correct kernels, the typical kernel is not faster than PyTorch, and the impressive outliers (13× speedup) are exceptions. The paper's framing emphasizes the existence proof ("models can generate fast kernels") rather than the failure of models to do so reliably, which is accurate but may overstate the practical readiness of the technology.

6. Limitations and Trade-offs

The Difficulty Estimation Cost Is the Elephant in the Room

The assumption or constraint. The compute-optimal test-time scaling framework fundamentally depends on estimating each prompt's difficulty before allocating the inference budget. The paper's method for doing so—generating 2048 samples per question and computing the pass@1 rate (oracle) or the PRM's average final-answer score (predicted)—is, by the authors' own acknowledgment, "extraordinarily expensive" and is explicitly excluded from the reported efficiency calculations:

"estimating difficulty in this way still incurs additional computation cost during inference... our experiments do not account for this cost largely for simplicity" (Section 3.2, paraphrased from the summary but consistent with the paper's discussion of the difficulty estimation tradeoff)

The consequence. The headline 4× efficiency gains over best-of-N (Figures 4, 8) are computed after difficulty is known, without amortizing the cost of learning it. Generating 2048 samples per question consumes more compute than the largest test-time budgets studied (256–512 generations), meaning the total cost—difficulty estimation + strategy execution—can easily exceed the cost of simply running best-of-N with a large fixed budget. In a realistic deployment, difficulty estimation could dominate the total inference budget, potentially eliminating or reversing the reported gains. The paper frames this as an "exploration-exploitation tradeoff" but provides no analysis of what fraction of the total compute the difficulty estimation step consumes at various budgets or how the tradeoff varies with the number of questions being answered (a batch setting where estimating difficulty once per question type might be amortized vs. a streaming setting where each question is novel).

What evidence exists in the paper. The paper acknowledges this gap explicitly (Section 3.2, Section 8) but provides no ablation measuring the total cost of difficulty estimation + execution vs. a uniform-allocation baseline. The predicted difficulty variant (using PRM scores rather than ground-truth labels) is evaluated only for its accuracy relative to oracle bins, not for its total compute cost relative to uniform strategies. The paper does not report whether a simpler, cheaper difficulty estimator (e.g., using the PRM score on a single sample, or using model confidence features) could achieve similar gains.

Mitigation status. The paper treats this as a "key avenue for future work" (Section 3.2) and suggests that future systems could train models to predict difficulty directly from question text, or use adaptive strategies that estimate difficulty from a small number of initial samples and then allocate the remaining budget. Neither approach is implemented or evaluated. Until this gap is closed, the 4× efficiency figure should be understood as an upper bound on achievable efficiency under perfect difficulty information, not a realized deployment gain.

Hard Problems Remain Completely Unsolved—Test-Time Compute Cannot Create Capability from Nothing

The assumption or constraint. The paper's framework assumes that the base model has some non-trivial probability of producing a correct answer—that the proposal distribution contains at least some correct solutions that test-time compute can surface or refine. This assumption is explicitly validated in the difficulty bin analysis: difficulty bins are defined by the base model's pass@1 rate, with bin 5 representing problems where the model's pass@1 is near zero.

The consequence. On the hardest questions (difficulty bin 5 across all analyses), no method makes meaningful progress regardless of compute budget. In Figure 3 (right), bin 5 accuracy for both beam search and best-of-N hovers at 1–3% for all budgets from 4 to 256 generations. In Figure 7 (right), the sequential-to-parallel ratio sweep for revisions shows bin 5 accuracy at roughly 2–3% irrespective of allocation strategy. In the FLOPs-matched comparison (Figure 9), the bin 5 scaling line is essentially flat near 0–5%, and test-time compute shows a 52.9% disadvantage relative to the ~14× larger model at R ≫ 1. The paper is transparent about this boundary:

"test-time compute can amplify existing capability but cannot create it. If the base model's pass@1 is near zero on a problem class, no amount of search or revision will help—there are no correct solutions in the proposal distribution to find or refine." (paraphrased from the paper's discussion)

This is a fundamental capability bound: the approach offers no path forward for genuinely novel or out-of-distribution reasoning that exceeds the base model's training distribution. For such problems, scaling pretraining remains the only viable path, and the paper's framework provides no mechanism for crossing this capability threshold.

What evidence exists in the paper. The evidence is consistent and unambiguous across all experimental sections: Figures 3 (right), 7 (right), 9, and the accompanying discussion in Sections 5.3, 6.2, and 7 all show bin 5 performance near zero for all methods and all budgets. The FLOPs-matched comparison (Figure 9) explicitly quantifies the failure: test-time compute is worse than pretraining on hard problems across all R values for PRM search, and for revisions at R ≫ 1.

Mitigation status. The paper does not attempt to solve this limitation—it is treated as a fundamental constraint. The authors suggest (Section 8) that future work could combine test-time compute with pretraining improvements to push the capability frontier, but no concrete approach is proposed. This limitation is not a weakness of the paper (which is empirically honest about it) but rather a finding that defines the scope of applicability: the compute-optimal framework is valuable only for problems within the base model's rough capability range. For organizations whose problem distribution skews toward genuinely novel or hard reasoning, investment in test-time compute infrastructure will yield diminishing or zero returns compared to investment in larger-scale pretraining.

Single Benchmark (MATH) on a Single Model Family—The Difficulty-Dependent Patterns May Not Generalize

The assumption or constraint. All experiments in the paper are conducted on the MATH benchmark (Hendrycks et al., 2021) using PaLM 2-S* (Codey) as the base model. The paper's central findings—that beam search degrades on easy problems due to verifier over-optimization, that sequential revisions are optimal for easy problems while a balanced ratio is optimal for hard ones, that compute-optimal scaling yields 4× efficiency gains—are entirely grounded in this single benchmark-model combination. The authors state they "believe this model is representative of the capabilities of many contemporary LLMs" (Section 4), but this claim is unverified.

The consequence. Several aspects of the findings could be domain- or model-specific in ways that substantially affect their generalizability:

PRM quality and over-optimization behavior. The paper's PRM is trained with Monte Carlo rollout supervision on PaLM 2-S* outputs. A model with different calibration properties, different error patterns, or different output distribution entropy might exhibit different PRM reliability and thus different over-optimization thresholds. The finding that beam search hurts easy-problem performance (Figure 3, right) depends on the specific interaction between PaLM 2-S*'s output distribution and the PRM's scoring behavior—this interaction may not hold for other model families.
Revision model training. The revision model's ability to learn from incorrect in-context examples (via the edit-distance-based pairing procedure in Section 6.1) depends on the base model's in-context learning capabilities and output structure. Different base models might require different revision training procedures or might not benefit from revisions at all.
Task domain specificity. MATH consists of competition-level math problems requiring symbolic reasoning, multi-step deduction, and exact-answer verification. It is unclear whether the difficulty-dependent patterns—particularly the finding that revisions help on easy problems and search helps on medium ones—generalize to other reasoning domains (code generation, logical reasoning, scientific QA, planning) or to tasks requiring factual knowledge retrieval rather than step-by-step inference. The mathematical reasoning domain has clean correctness signals (exact string matching) that enable both the PRM training pipeline and the difficulty estimation oracle—domains without such signals would require fundamentally different approaches.
Difficulty estimation via pass@1. The five-quintile binning is defined relative to the base model's pass@1 rate, which conflates "the problem is intrinsically hard" with "the base model happens to be bad at this problem type." A different base model would produce different difficulty bins, and the optimal strategies per bin might shift accordingly.

What evidence exists in the paper. The paper provides no cross-benchmark or cross-model-family evaluation. The 500-question MATH test set, split into five bins of ~100 each and further divided by two-fold cross-validation (~50 questions per fold per bin for strategy selection), is a relatively small sample for the number of hyperparameters being optimized (search algorithm choice, beam width, lookahead depth, sequential-to-parallel ratio). The paper does not report confidence intervals on the compute-optimal scaling curves (Figures 4, 8), making it difficult to assess whether the observed gains are statistically robust at this sample size or sensitive to the specific bin boundaries.

Mitigation status. The paper does not attempt to validate findings on additional benchmarks or models, and the authors acknowledge the single-benchmark scope as a limitation (implicitly in Section 8 when calling for future work). Replication on, for example, GSM8K (grade-school math), HumanEval (code generation), or MMLU (knowledge-intensive QA) would be necessary to establish the generality of the difficulty-conditioned allocation framework. The paper's conceptual framework (difficulty-aware allocation of test-time compute) is sufficiently abstract to potentially generalize, but the specific empirical findings (which strategies work at which difficulty levels) are almost certainly benchmark- and model-dependent.

The 14× Larger Model Baseline Is Not Compute-Optimally Trained, and the FLOPs-Matched Comparison Is Asymmetric

The assumption or constraint. The FLOPs-matched comparison in Section 7 compares PaLM 2-S* with compute-optimal test-time scaling against a model with approximately 14× more parameters, where the larger model is trained with the same amount of data as the smaller model (i.e., scaling parameters only, not data). The paper explicitly acknowledges that this departs from compute-optimal pretraining practice:

"We choose this setting as it is representative of a canonical approach to scaling pretraining compute and leave the analysis of compute-optimal scaling of pretraining compute where the data and parameters are both scaled equally to future work." (Section 7)

The comparison also gives the larger model only greedy decoding with no test-time compute augmentation—no majority voting, no best-of-N, no search, no revisions. The smaller model, by contrast, is allowed to use the full compute-optimal framework (difficulty estimation, adaptive strategy selection, up to 512 generations of search or revisions).

The consequence. The experimental design stacks the deck systematically in favor of test-time compute:

Data scaling neglect. Chinchilla-optimal pretraining (Hoffmann et al., 2022) would scale both model parameters and training data tokens equally. A model trained with 14× more total FLOPs allocated according to Chinchilla scaling laws would likely outperform a parameter-only-scaled model on MATH, making the pretraining baseline weaker than it would be under compute-optimal pretraining. The paper's reported advantages of test-time compute over pretraining (e.g., +27.8% relative improvement on easy questions for revisions at R ≪ 1, Figure 1 top-right) would shrink or potentially reverse against a properly compute-optimal larger model.
Asymmetric inference budget. The smaller model receives up to 512 generations of test-time compute (search, revisions, or compute-optimal combinations), while the larger model receives none (greedy decoding). This is not a neutral FLOPs accounting—it is a comparison of "small model + extensive test-time compute" vs. "large model + zero test-time compute." A fairer baseline would give the larger model some test-time compute budget as well (e.g., best-of-8 or best-of-64, which would consume a fraction of the total FLOPs and might substantially improve its performance). The paper's finding that test-time compute can "outperform a 14× larger model" applies specifically to a larger model making no use of inference-time compute, which is not representative of how large models are typically deployed in practice (where even greedy decoding is often augmented with temperature sampling and majority voting).
FLOPs accounting favors the regime where test-time compute looks best. The formula for the smaller model's available inference budget involves the ratio R = D_inference / D_pretrain. The paper tests R = 0.16 (R ≪ 1), R = 0.79 (R ≈ 1), and R = 22 (R ≫ 1). At R ≪ 1, which corresponds to low-volume inference relative to pretraining, the smaller model gets a large inference budget—and this is precisely the regime where test-time compute shows its strongest advantages. But many real-world deployments operate at R ≫ 1 (high-throughput inference serving many queries), where the per-query budget is tighter and the comparison is less favorable to test-time compute. The paper acknowledges this dependence (Section 7) but does not provide guidance on what R values are typical for different deployment scenarios.

What evidence exists in the paper. The FLOPs-matched results (Figure 9, Figure 1 bar charts) explicitly show the R-dependent nature of the findings. The data support the conclusion that test-time compute can be preferable at R ≪ 1 and on easy-to-medium problems, but the strength of this conclusion is conditioned on the weak pretraining baseline. Table 1 (top-right bar chart) shows +27.8% relative improvement for revisions on medium questions at R ≪ 1, but this is the most favorable scenario—at R ≫ 1, the advantage shrinks to +5.4%, and for hard questions becomes -37.2%.

Mitigation status. The paper is transparent about the Chinchilla-optimal pretraining departure and calls it out explicitly for future work. However, the paper does not provide even a rough estimate of how much the results would change under compute-optimal pretraining (e.g., using scaling law projections to bound the effect). The asymmetry of giving the larger model no test-time compute is not acknowledged as a limitation, though it is perhaps the most consequential design choice in the FLOPs-matched comparison. Future work that gives both models proportional test-time compute budgets (e.g., allocating 10% of total FLOPs to inference for both the small and large model) would be necessary to establish a genuinely fair tradeoff.

Latency and Wall-Clock Time Are Ignored—Sequential Strategies Are Impractical for Interactive Applications

The assumption or constraint. The paper measures test-time compute exclusively in "generations" (number of complete solutions sampled), which serves as a reasonable proxy for total FLOPs but completely ignores latency—the wall-clock time required to produce a final answer. This distinction matters because different strategies have fundamentally different parallelism properties:

Fully parallel strategies (best-of-N, parallel sampling) can execute all N generations simultaneously given sufficient hardware parallelism (multiple GPUs or large batch sizes), with latency roughly equal to the time for one generation.
Fully sequential strategies (a single chain of N revisions) are inherently serial—each revision depends on the output of the previous one—with latency roughly N × the time for one generation.
Hybrid strategies (√N parallel chains of √N revisions each) have intermediate latency: √N × generation time.

The consequence. The compute-optimal policies identified by the paper often favor sequential strategies, particularly on easy problems (where revisions are recommended) and for the sequential-to-parallel ratio sweep (where higher sequential ratios perform better at low budgets, Figure 7 left). But for latency-sensitive applications—interactive assistants, real-time decision-making, any deployment where the user is waiting for a response—a strategy allocating 256 generations as 16 parallel chains of 16 revisions each takes roughly 16× longer in wall-clock time than running 256 parallel samples simultaneously on sufficient hardware. The paper's 4× efficiency gain in terms of FLOPs may translate to a 4× latency penalty depending on the degree of sequential depth in the optimal strategy, making it unacceptable for production deployments with strict latency requirements.

This tradeoff is particularly acute for the revision model: while sequential revisions show marginally better aggregate performance than parallel sampling (Figure 6 right), the latency cost of 64 sequential revisions (64× single-generation latency) is dramatically higher than 64 parallel samples (1× single-generation latency with sufficient hardware). The hierarchical aggregation procedure (Appendix I) with √N parallel chains attempts to balance this tradeoff but does not eliminate it—the optimal ratio often involves substantial sequential depth (e.g., 8:1 sequential-to-parallel, Figure 7 left).

What evidence exists in the paper. The paper provides no latency or wall-clock measurements, no discussion of this tradeoff, and no analysis of how the compute-optimal policy changes under a latency constraint. The generation budget accounting (Sections 3.1, 5.3) treats all generations as fungible units of FLOPs without distinguishing serial vs. parallel execution. Figure 7 (left) shows that at low budgets (8–32 generations), fully sequential strategies are optimal—precisely the regime where latency sensitivity is highest (since low-budget users are likely in interactive settings).

Mitigation status. The paper does not address latency at all. This is a significant gap for a paper that frames itself as providing guidance for practical deployment (Section 1: "on-device deployment," "resource allocation decisions"). A complete analysis would need to characterize the Pareto frontier of accuracy vs. latency for different strategies, or introduce a latency-aware objective that penalizes sequential depth alongside total FLOPs. The existing compute-optimal policy should be understood as FLOPs-optimal, not necessarily deployment-optimal, and practitioners with latency constraints would need to re-derive the allocation policy with a latency penalty.

Verifier Over-Optimization Defines a Hard Ceiling That the Paper Identifies but Does Not Resolve

The assumption or constraint. The paper's PRM-based search methods and the entire compute-optimal scaling framework depend on the reliability of the learned verifier under optimization pressure. The PRM is trained via Monte Carlo rollouts on base model outputs and provides step-level correctness estimates. When search algorithms aggressively optimize against this reward signal—by exploring many beams, applying lookahead scoring, or running at high generation budgets—they can exploit imperfections in the PRM, finding solutions that score highly under the verifier but are actually incorrect.

The consequence. This phenomenon, documented in detail by the paper, imposes a hard ceiling on how much test-time compute can help, even on problems within the base model's capability range:

Beam search degrades easy-problem performance at high budgets (Figure 3, right, bin 1: accuracy decreases from ~78% to ~77% as budget goes from 4 to 256). This means more compute actively hurts—the verifier is being exploited.
Lookahead search—the most powerful optimizer—paradoxically performs worst overall (Figure 3, left). By simulating forward rollouts to get better step-level scores, lookahead search intensifies the optimization pressure against the PRM and amplifies its imperfections.
Qualitative examples show degenerate outputs (Appendix M): repetitive low-information steps at the end of solutions, overly short 1–2 step solutions that score highly but are substantively inadequate. The PRM can be "fooled" into preferring outputs that pattern-match to correct solutions without actually solving the problem.
The compute-optimal policy can only partially mitigate this. By routing easy problems away from aggressive search (Section 5.3), the policy stays below the over-optimization threshold, but it does not eliminate the underlying verifier quality problem. On medium-difficulty problems where beam search is deployed, the scaling curves flatten and sometimes decline at high budgets (Figure 3, left: beam search plateaus around 34% at 256 generations, below best-of-N's ~37% at 512).

This means that further scaling test-time compute—even if difficulty estimation were perfect and free—would not yield unbounded improvements. The ceiling is set by verifier quality, not by search algorithm sophistication or compute budget magnitude. The paper's own findings demonstrate that increasing the compute budget from 256 to 512 generations provides minimal or negative additional gains for search methods (Figure 3, left).

What evidence exists in the paper. The evidence for verifier over-optimization is extensive and consistent across multiple analyses: the bin-specific beam search curves (Figure 3, right), the overall search algorithm comparison (Figure 3, left), the qualitative failure examples (Appendix M, e.g., Figure 29), and the finding that last-step PRM aggregation outperforms more sophisticated step-level aggregation (Appendix E, Figure 13—the "min" and "prod" strategies that rely more heavily on intermediate PRM predictions underperform because those intermediate predictions are less reliable). The paper's difficulty-conditioned strategy selection can be understood partly as a mechanism to stay below the over-optimization threshold—using weaker optimization (best-of-N) where the verifier is most reliable (easy problems) and allowing stronger optimization (beam search) only where the verifier signal still provides genuine guidance (medium problems).

Mitigation status. The paper identifies this as a central challenge and treats it as a major direction for future work (Section 8), but does not propose or evaluate any solutions within the current work. Potential approaches—adversarial PRM training (including search-generated solutions in the training data), ensemble verification, KL-constrained search that penalizes deviations from the base model's output distribution, or dynamic strategy adjustment based on real-time verifier calibration checks—are not explored. The paper's main mitigation is the compute-optimal policy itself, which reduces over-optimization by avoiding the most problematic strategy-difficulty combinations, but this is a workaround rather than a solution. The implication for practitioners is clear: the practical ceiling of this approach is bounded not by compute budget but by verifier training quality, and investment in better PRM training data and methodology may yield higher returns than investment in more sophisticated search algorithms or larger inference budgets.

7. Implications and Future Directions

How This Work Changes the Landscape

KernelBench shifts the conversation around LM-powered code generation from correctness-centric benchmarking to a production-oriented framework where wall-clock performance on real hardware is the primary success criterion. This is a conceptual reframing, not merely an incremental benchmark contribution. Prior code generation benchmarks—HumanEval, MBPP, DS-1000, SWE-Bench—measure whether LMs can produce functionally correct code, treating performance as an afterthought or ignoring it entirely. KernelBench inverts this: a generated kernel that is correct but slower than PyTorch Eager counts as a failure (under fast1), because in production ML pipelines, a slower kernel provides zero incentive for adoption regardless of its correctness. This reframes LM code generation from "can the model replicate known solutions?" to "can the model invent optimizations that outperform expert-tuned, often closed-source implementations?"—a fundamentally harder and more practically meaningful standard that the paper's results show current frontier models fail to meet in over 80% of cases.

The paper also provides the field with a diagnostic decomposition of where LM kernel generation fails and what different feedback signals can fix. The finding that execution errors (compilation failures, memory violations) are the dominant failure mode for non-reasoning models but are largely fixable through iterative refinement with compiler feedback, while functional correctness errors resist feedback because the signal is impoverished ("your outputs don't match" without localization), reframes the research agenda. Prior to this work, the natural assumption would be that CUDA syntax is the primary bottleneck—teach LMs CUDA better, and kernel generation improves. The paper shows this is only half the story: syntax errors are fixable through feedback loops, but logical correctness errors in parallel code are the hard residual, and they are connected to a structural data scarcity (CUDA constitutes only 0.073% of The Stack v1.2). This diagnostic chain—identifying which failure mode persists after feedback, and hypothesizing why—gives the field a concrete target for data-centric interventions rather than prompting refinements.

Furthermore, the paper demonstrates that the effectiveness of test-time improvement strategies (repeated sampling, iterative refinement) is strongly model-dependent, with reasoning models (DeepSeek-R1, OpenAI o1) benefiting substantially from execution and profiler feedback while non-reasoning models (Llama 3.1-70B) show minimal gains. This finding—that iterative refinement with feedback can double fast1 for DeepSeek-R1 on Level 2 (36% → 72%) but barely moves Llama 3.1-70B (0% → 4%)—has implications beyond kernel generation. It suggests that the value of test-time compute strategies in code generation is not uniform across model capabilities, and that investment in richer feedback environments (like KernelBench's compiler + profiler integration) will yield asymmetric returns depending on the base model's reasoning abilities. This reframes the research question from "does feedback help?" to "for which models, on which tasks, does feedback help, and what type of feedback is most informative?"—a more nuanced and productive framing.

The paper resolves a latent contradiction in the HPC code generation literature. Prior work evaluated LMs on translating C++ to CUDA or generating well-known kernels like GEMM—tasks for which training data likely exists in the model's corpus. These studies could report superficially positive results (e.g., LMs can generate a correct matrix multiplication kernel) while obscuring the harder question of whether LMs can generate optimized kernels for novel workloads. KernelBench, by selecting 250 real-world ML workloads "many of which do not have existing human-written implementations," tests a fundamentally different capability: can LMs discover optimizations that haven't been open-sourced? The paper's results show they largely cannot out of the box (fewer than 20% fast1), but can with iterative refinement (72% on Level 2 for DeepSeek-R1), reconciling the tension between "LMs know some CUDA" and "LMs cannot yet replace human kernel engineers."

Finally, the paper makes certain research directions less attractive by demonstrating their limitations. The finding that providing hardware specifications in-context (Section 5.2.2, Table 15) has minimal impact on overall performance—even as DeepSeek-R1 begins to generate tensor core wmma instructions that mostly fail to compile—suggests that prompt engineering with abstract hardware knowledge is not a promising path to hardware-efficient kernels. Similarly, the few-shot experiments (Section 5.2.1, Tables 10–12) show that in-context examples of optimization techniques (fusion, tiling, FlashAttention) degrade overall fast1 because models attempt more ambitious but more error-prone code, suggesting that "show more examples" is not a straightforward path to improvement and may require complementary correctness-guidance mechanisms. These negative results are valuable: they redirect effort away from prompting-based approaches and toward data-centric interventions (curating high-quality CUDA training data) or training-based methods (fine-tuning on kernel code).

Follow-Up Research This Work Enables

Fine-tuning base models on available CUDA corpora and measuring KernelBench improvement. The paper's core hypothesis—that CUDA's scarcity in training data (0.073% of The Stack) is a primary cause of functional correctness errors—is currently untested through intervention. A direct experiment: take a base model (e.g., Llama 3.1-70B or DeepSeek-V3), fine-tune it on all available open-source CUDA code (from GitHub repositories, the CUDA programming guide, NVIDIA's official samples, and kernel libraries like CUTLASS and ThunderKittens), and measure the change in one-shot fast1 across all three KernelBench levels. If fast1 improves substantially (e.g., from 3% to 15%+ on Level 1 for Llama 3.1-70B), this would validate the data scarcity hypothesis and establish a clear path forward: invest in CUDA data curation. If improvement is marginal, the bottleneck lies elsewhere (perhaps in the fundamental difficulty of parallel programming reasoning, which no amount of pattern exposure can substitute for). A strong follow-up would also ablate what types of CUDA data help most: complete kernels vs. code snippets, commented vs. uncommented code, diverse algorithmic patterns vs. many instances of the same pattern.

Combining iterative refinement with stochastic sampling (temperature > 0) to test whether exploration and feedback are complementary. The paper's iterative refinement experiments use temperature = 0, making trajectories deterministic given initial generation and feedback. Repeated sampling uses high temperature (1.6 for DeepSeek-V3) but no feedback. The natural combination—iterative refinement where each turn samples at temperature > 0, potentially with multiple candidates per turn selected by the verifier—is not tested. A strong follow-up would run iterative refinement with temperature = 0.7–1.0 for DeepSeek-R1 (if API constraints can be resolved) or DeepSeek-V3, measuring whether the 72% fast1@10 on Level 2 for R1 can be pushed higher when the model can escape local optima stochastically. This would also test whether the plateau in fast1@N curves (Figure 6) represents a fundamental limit or an artifact of deterministic refinement. The experiment could be structured as a budget-controlled comparison: at fixed total inference calls (e.g., 50), compare pure repeated sampling, pure iterative refinement (temperature 0), and hybrid strategies (e.g., 5 turns of refinement with 10 stochastic samples per turn, selecting the best via correctness checks).

Training a lightweight difficulty predictor from KernelBench task features to enable adaptive test-time strategy allocation. The paper's iterative refinement and repeated sampling experiments apply uniform strategies across all tasks. But the paper's own error analysis (Figure 2) and case studies (Appendix D) show substantial per-task heterogeneity: some tasks benefit dramatically from refinement (the convolution kernel in Table 5 improving 8×), while others see the model stuck in loops (Problem 54). A natural extension of the "compute-optimal test-time scaling" paradigm would be to predict, from task features (reference code length, number of operations, operator types, tensor shapes), whether a given task is more likely to benefit from repeated sampling or iterative refinement, and allocate the inference budget accordingly. Using KernelBench's 250 tasks as training data, one could train a classifier on "refinement helped" vs. "refinement didn't help" per task and evaluate whether budget-conditioned allocation outperforms uniform allocation. This would also address the practical question: given a new PyTorch workload, should an engineer invest time in iterative refinement or just generate many parallel samples?

Extending KernelBench to alternative GPU programming abstractions (Triton, ThunderKittens, CUTLASS) and measuring whether higher-level languages close the correctness gap. The paper notes that all reported results use inline CUDA-C, but that LMs could theoretically use Triton (a more Pythonic, higher-level language for GPU programming), ThunderKittens (a library of optimized primitives), or CUTLASS (template-based linear algebra). The hypothesis: these higher-level abstractions reduce the syntactic burden and memory-management complexity, potentially allowing LMs to achieve higher correctness rates by avoiding low-level CUDA errors (memory violations, incorrect thread indexing). A strong follow-up would re-run the one-shot baseline and iterative refinement experiments on the same 250 KernelBench tasks but prompting models to generate Triton kernels instead of CUDA. If Triton-based generation achieves, say, 40% fast1 on Level 1 vs. CUDA's 12% (DeepSeek-R1 one-shot), this would validate that the abstraction level is a major lever and that future LM kernel generation systems should target higher-level languages as the default output format. The KernelBench framework already supports this—the evaluation protocol only requires that ModelNew.forward() runs correctly and is timed, regardless of what language the kernel is written in.

Stress-testing the 5-input correctness protocol against more systematic testing to bound false-positive rates. The paper uses 5 random inputs for correctness verification, validated empirically on a 100-kernel subset where correctness was all-or-nothing (all kernels that passed 5 tests also passed 100, and vice versa). This is a reasonable practical choice, but it provides no formal guarantees. A strong follow-up would take the kernels that achieve fast1 in the one-shot baseline (e.g., DeepSeek-R1's 12 kernels on Level 1 that are both correct and faster than PyTorch Eager) and subject them to more rigorous testing: larger numbers of random inputs (10,000+), edge-case inputs (zero tensors, very large/small values, non-power-of-2 dimensions), differential testing against multiple reference implementations, and where possible, formal verification of simple properties (e.g., no out-of-bounds memory accesses). If some fast1 kernels fail under extended testing, this would (a) reveal the false-positive rate of the 5-input protocol and (b) potentially adjust the reported fast1 scores downward, providing a more conservative estimate of LM capability. If all fast1 kernels survive extended testing, this would strengthen confidence in the benchmark's practical relevance.

Building an agentic kernel development workflow that combines generation, compilation, profiling, and multiple refinement strategies. The paper's iterative refinement uses a simple feedback loop: generate, compile, execute, provide feedback, repeat. A human kernel engineer uses a richer workflow: they might try multiple alternative approaches in parallel, profile to identify bottlenecks, read documentation to understand hardware features, write small test programs to validate their understanding, and abandon approaches that aren't working. An agentic system that gives the LM access to these tools—a documentation retriever (for CUDA APIs, hardware specs), a test-case generator (for validating subcomponents of the kernel), and a meta-controller that decides when to persist vs. switch strategies—could potentially push beyond the 72% fast1 ceiling on Level 2. KernelBench provides the evaluation framework for testing such a system: the 250 tasks, the automated correctness and performance measurement, and the existing baseline results (Table 1, Table 2) provide a clear target to beat. A strong follow-up would build a minimal version of this agentic workflow (perhaps using a reasoning model as the controller and a code-generation model as the executor) and measure the fast1 improvement over simple iterative refinement at matched inference budget.

Practical Applications and Downstream Use Cases

Open-source kernel generation for PyTorch operations currently served by closed-source libraries. Many Level 1 tasks in KernelBench correspond to primitive operations (matrix multiplies, convolutions, activation functions) that PyTorch currently serves through closed-source cuBLAS and cuDNN kernels. The paper shows that some LM-generated kernels achieve meaningful speedups on these operations—a 13× speedup on diagonal matrix multiplication (Appendix D.1), a 2.9× speedup on GELU (Appendix D.2), a 2.8× speedup on cosine similarity loss (Appendix D.3). If even a fraction of these kernels can be validated, hardened, and contributed to open-source repositories, they provide permissively licensed alternatives to proprietary NVIDIA libraries. This matters for organizations building on non-NVIDIA hardware (where cuBLAS isn't available), for research on custom hardware accelerators that need open-source reference implementations, and for the broader goal of reducing dependence on closed-source software stacks. KernelBench provides the testing infrastructure to validate such kernels automatically.

Reducing the time gap between ML architecture proposal and performant implementation. The paper motivates KernelBench by noting that FlashAttention took 5 years from the Transformer's proposal, and another 2 years to port to Hopper GPUs. If LMs can assist—or fully automate—kernel generation for novel architectures, this timeline could compress dramatically. For a research team proposing a new attention variant, a new normalization scheme, or a new activation function, the workflow could become: (1) implement a reference in PyTorch, (2) submit it to an LM-powered kernel generation system (perhaps using iterative refinement with profiler feedback), (3) receive an optimized kernel that is evaluated against the PyTorch Eager baseline on the target hardware. Even if the LM only succeeds on, say, 40% of novel operations (comparable to fast1 on Level 1 for reasoning models), this would still accelerate architecture development by automating the optimization of common building blocks and freeing human kernel engineers to focus on the hardest operations. The paper's results on Level 2 (operator fusion)—where DeepSeek-R1 with refinement achieves 72% fast1—suggest that fusing novel operator sequences (a common need when implementing new architectures) is the sweet spot for current LM capabilities.

Automated kernel porting across hardware generations. The paper's cross-hardware experiments (Section 4.4, Table 14) reveal that one-shot LM-generated kernels have widely varying performance across GPU types (DeepSeek-R1 Level 2 fast1 ranges from 36% on L40S to 47% on A10G). This variance is a problem for manual deployment but an opportunity for automated porting: if the LM can be conditioned on the target hardware's specifications and given profiler feedback from that hardware, iterative refinement could adapt a kernel that works well on one GPU to achieve good performance on another. The paper's hardware-aware prompting experiments (Section 5.2.2, Table 15) show that this capability is currently nascent—providing hardware specs induces models to attempt specialized instructions but not to implement them correctly—but the infrastructure exists. For organizations maintaining kernel libraries across GPU generations (a major pain point the paper identifies), an LM-assisted porting pipeline that takes a working A100 kernel as input, conditions on H100 specifications, and iteratively refines using H100 profiler feedback could reduce the manual effort of rewriting kernels for each new hardware release.