CUDA-L2: Surpassing cuBLAS Performance for Matrix Multiplication through Reinforcement Learning

ArXiv: 2512.02551

🎯 Pitch

A reinforcement learning system trained to write CUDA kernels discovers HGEMM optimizations that beat NVIDIA’s heavily-tuned cuBLAS libraries by up to 22%, with the largest gains on smaller matrices where GPU underutilization leaves optimization headroom. The method automatically tailors kernel strategies across 1,000 matrix size configurations, outperforming even cuBLASLt’s exhaustive auto-tuning that benchmarks up to 100 hand-written candidate kernels per shape.

1. Executive Summary

This paper introduces CUDA-L2, a system that combines large language models and reinforcement learning to automatically optimize Half-precision General Matrix Multiply (HGEMM) CUDA kernels across 1,000 configurations spanning all triplet combinations of M, N, K from {64, 128, 256, 512, 1024, 2048, 4096, 8192, 12288, 16384}. The system extends prior work through a multi-stage RL training pipeline progressing from general CUDA kernel optimization to HGEMM-specific tuning, complemented by continued pretraining on diverse CUDA code, comprehensive NCU profiling metrics as optimization context, and retrieval-augmented context for architectural knowledge. Against NVIDIA's strongest baseline, cuBLASLt-AutoTuning (which exhaustively benchmarks up to 100 algorithm candidates per configuration and selects the fastest), CUDA-L2 achieves an 11.4% average speedup in offline mode and 15.9% in server mode, with speedups reaching 22.0% over torch.matmul and 19.2% over cuBLAS in offline evaluation. The gains are most pronounced on smaller matrix sizes — where GPU underutilization creates optimization headroom — establishing that LLM-guided RL can discover superior implementations even for performance-critical, heavily-optimized kernels, but only where the hardware has slack to exploit.

2. Context and Motivation

The Core Problem: HGEMM Remains Suboptimal Despite Decades of Hand-Optimization

The paper addresses a deceptively narrow but enormously consequential problem: Half-precision General Matrix Multiply (HGEMM) CUDA kernels are not performance-optimal across the full space of matrix dimensions that LLMs actually use, even when compared against NVIDIA's own heavily-tuned libraries. This is surprising because matrix multiplication has been one of the most intensely optimized operations in computing for decades, and the cuBLAS ecosystem represents the cumulative product of that effort. Yet the paper demonstrates that significant performance — 11–22% depending on the baseline — remains discoverable.

Why does this gap exist? The paper identifies a multi-dimensional scaling challenge that makes exhaustive manual optimization impractical. Three factors conspire:

1. Different (M, N, K) triplets require different optimization strategies. The space of relevant matrix dimensions is not a single problem but 1,000 related-but-distinct optimization problems. The optimal tile size, the decision to use block swizzling, the number of pipeline stages, the choice between single-buffer and double-buffer register management — all of these change depending on whether you're multiplying a small attention head (M=64, N=64, K=64) or a large feedforward layer (M=8192, N=16384, K=16384). A kernel tuned for one configuration may perform poorly on another. Section 5 demonstrates this concretely: for (M=8192, N=512, K=2048), CUDA-L2 selects BM=160 and pads M to 8320, achieving a 15.2% speedup over cuBLASLt-AutoTuning-TN. Change BM to 128 (a conventional choice that evenly divides 8192) and the speedup drops to 0.4%; change to 256 and performance degrades by 15.7%. These optimization choices are non-transferable even across different dimensions of the same operation on the same GPU.

2. Optimizations rarely transfer across GPU architectures. The A100 (Ampere), which this paper targets, has fundamentally different characteristics from Hopper (H100) or Blackwell (B200) — different tensor core dimensions, different shared memory sizes, different cache hierarchies, different instruction latencies. A tile size optimized for one generation may be suboptimal for the next. This means that the manual optimization problem is not solved once but must be re-solved for each new hardware generation, multiplying the combinatorial explosion.

3. Even identically-sized matrices on the same GPU can admit different optimal strategies depending on secondary considerations. The paper notes that accumulator precision — using FP16 versus FP32 accumulators for FP16 inputs — leads to different register pressure and thus different optimization strategies. This subtlety means that the configuration space is larger than just (M, N, K); secondary parameters (accumulator type, epilogue operations, beta values) create further combinatorial branches.

The net effect: a truly comprehensive hand-optimization effort would need to tune thousands of dimension-specific, architecture-specific kernel variants. This is why cuBLASLt-AutoTuning — which exhaustively benchmarks up to 100 algorithm candidates per configuration — represents the state of the art: it automates selection among hand-designed kernels but does not automate discovery of new kernels. And even this approach leaves 11.4% on the table, as CUDA-L2 demonstrates.

Why This Problem Matters Now

The practical importance of HGEMM optimization is difficult to overstate. Matrix multiplication accounts for a "very significant portion of computation time in both training and inference" (Section 1) of large language models. In a typical transformer forward pass, every attention layer performs multiple matrix multiplications (Q×K^T projections, attention×V aggregation, output projection), and every feedforward layer performs two large GEMMs. Across models with dozens or hundreds of layers, matmul dominates the FLOPs budget. A 10% improvement in HGEMM throughput translates approximately to a 10% improvement in LLM inference throughput for the matmul-bound portion of the model.

The paper provides evidence that this optimization headroom is real and exploitable by citing the TensorRT-LLM team's recent work on DeepSeek-R1: targeted Grouped GEMM optimizations yielded a 13% speedup (Section 1). This was achieved when "human expertise was narrowly focused on a single target" — a single model architecture, a single GPU generation. The manual engineering effort required to achieve that 13% gives a sense of the cost of hand-optimization. CUDA-L2's contribution is to demonstrate that automated methods can discover larger speedups (15.9% over cuBLASLt-AutoTuning in server mode) across a broader range of configurations without the same per-model, per-architecture engineering investment.

The timing is also significant. As LLMs scale, the matrix dimensions encountered span an increasingly wide range: small attention heads with M,N on the order of 64–128, medium feedforward layers with dimensions around 4096, and large expert layers (in mixture-of-experts architectures like DeepSeek-R1) with dimensions reaching 16,384 or higher. The paper's configuration space — all triplet combinations from {64, 128, 256, 512, 1024, 2048, 4096, 8192, 12288, 16384} — is explicitly justified as covering "those used in attention and FFN layers of widely open-sourced models like Qwen, Llama, and DeepSeek" (Section 1). This is not an academic benchmark; it is the production configuration space of contemporary LLMs.

Where Prior Approaches Fall Short

The paper identifies shortcomings across three categories of prior work:

Manual optimization (cuBLAS/cuBLASLt). NVIDIA's libraries represent the gold standard of hand-tuned matmul performance. cuBLAS uses internal heuristics to select among pre-compiled kernel variants, hiding the selection logic from the developer. cuBLASLt exposes more control, allowing developers to enumerate available algorithms and benchmark them exhaustively (cuBLASLt-AutoTuning benchmarks up to 100 candidates per configuration, as shown in Listing 8). The limitation is fundamental: these libraries can only select among existing kernel designs. They cannot generate novel kernel implementations optimized for specific dimension triplets. CUDA-L2, by contrast, generates new CUDA code — using CuTe abstractions, inline PTX assembly, and custom tiling strategies — that was never anticipated by the library authors.

Benchmark-focused automatic kernel generation (KernelBench, AI CUDA Engineer, CUDA-L1). Recent work has explored using LLMs to generate CUDA kernels automatically. KernelBench (Ouyang et al., 2025) provides a benchmark of 250 diverse CUDA tasks, but critically, each task is evaluated on a single, fixed input configuration (e.g., one specific dimension). The paper explicitly notes (Section 1):

"it remains unclear how these benchmark-optimized kernels translate to real-world production environments"

A kernel that achieves excellent performance on KernelBench's single test dimension may perform poorly on the hundreds of different dimensions encountered in actual model inference. The benchmarks test a single point in the configuration space; real deployment requires performance across the entire space. This is the key gap: prior LLM-based kernel optimization has not demonstrated generalization across the dimension space that matters for production.

CUDA-L1 (Li et al., 2025) introduced contrastive RL for kernel optimization but was limited in ways that made HGEMM optimization infeasible. Its SFT stage only fine-tuned on KernelBench kernels, providing no exposure to sophisticated matrix multiplication implementations. Its pretrained LLM lacked knowledge of newer CUTLASS versions, the CuTe abstraction library, and recent GPU architectures — all of which are essential for generating competitive HGEMM kernels. And it did not incorporate the detailed NCU profiling metrics (memory throughput, SM occupancy, cache efficiency) that enable an RL agent to learn why a kernel is fast or slow rather than simply observing end-to-end execution time.

General-purpose automatic tuning (AutoTVM, Halide, Triton autotuning). While not directly compared, domain-specific auto-tuners exist for other frameworks. Triton's autotuner, for example, can explore tile sizes and pipeline configurations for Triton kernels. However, these operate at a higher level of abstraction than raw CUDA — Triton compiles to PTX through an intermediate representation, limiting the optimizations available compared to direct CUDA C/C++ with inline PTX assembly and CuTe primitives. CUDA-L2 operates at the lowest level (CUDA .cu files compiled with nvcc, explicitly avoiding Python-based DSLs like Triton, as stated in Section 3.2.3), giving it access to optimizations that higher-level frameworks cannot express.

How CUDA-L2 Positions Itself

CUDA-L2 frames itself as solving the generalization across configurations problem that prior LLM-based kernel optimization sidestepped. The paper's contributions can be understood as four incremental-but-essential extensions to CUDA-L1 that enable HGEMM optimization at production scale:

Continued pretraining on diverse CUDA code (Section 3.2.1). To generate competitive matmul kernels, the model needs exposure to sophisticated implementations beyond basic KernelBench tasks. CUDA-L2 collects CUDA code from web sources (cleaned and segmented using rule-based filtering), established libraries (PyTorch, ATen, CUTLASS, NVIDIA tutorials), and pairs each code snippet with an LLM-generated instruction description. These instruction-context-code triplets are used for continued pretraining on DeepSeek 671B. This addresses the "knowledge gap" — prior models simply hadn't seen enough examples of high-performance CUDA to generate competitive kernels.

Multi-stage RL from general to specialized (Sections 3.2.2–3.2.3). Rather than jumping directly to HGEMM optimization, CUDA-L2 progresses through stages of increasing specialization. First, a general-kernel RL stage trains on roughly 1,000 kernels from established libraries spanning diverse operations (linear algebra, convolutions, reductions, element-wise operations, attention, sampling). This builds broad optimization capability. Then, an HGEMM-specific RL stage focuses exclusively on matrix multiplication with varying (M, N, K), using the same contrastive RL strategy from CUDA-L1 but now applied to the specialized domain. This curriculum approach prevents the RL from getting stuck in local optima early — the model learns general optimization principles before applying them to the most challenging kernel type.

Rich performance feedback via NCU profiling (Section 3.2.3). CUDA-L1 used only end-to-end execution time as reward. For HGEMM, where performance differences may be subtle and attributable to specific microarchitectural behaviors, this is insufficient. CUDA-L2 incorporates "memory throughput, compute utilization, warp occupancy, and cache hit rates" from NVIDIA Nsight Compute profiling into the RL context. This enables the model to perform root-cause analysis: a slow kernel might have high compute utilization (good) but poor memory throughput (bad), suggesting a tiling or prefetching issue. Without this granular feedback, the RL agent is effectively blind to why one kernel outperforms another, making optimization a random walk in a high-dimensional space.

Retrieval-augmented context for architectural knowledge (Section 3.2.1). New GPU architectures, library versions (CUTLASS 3.x, CuTe), and optimization techniques emerge faster than foundation models are retrained. CUDA-L2 uses each instruction as a search query to retrieve relevant documentation and code examples, concatenating them as context during generation. This allows the system to leverage post-training knowledge — a CuTe tutorial published after the base model's training cutoff, for example — without requiring full model retraining.

The paper's positioning is explicit about scope and limitations: it is not claiming to have solved HGEMM optimization for all architectures (work on Ada Lovelace, Hopper, and Blackwell is "ongoing"), and it is not claiming that LLM-guided RL universally outperforms human experts on all kernel types. The claim is narrower and more credible: for HGEMM on A100 GPUs, across the dimension space that production LLMs actually use, automated RL-driven kernel generation discovers optimizations that even NVIDIA's exhaustive auto-tuning baseline misses. This is significant because it shifts the boundary of what we consider "optimally tunable" — demonstrating that the exploration-exploitation tradeoff in kernel optimization can be automated at a scale impractical for manual effort.

3. Technical Approach

3.1 Reader Orientation

CUDA-L2 is a reinforcement learning system that trains a large language model to write high-performance CUDA kernels for half-precision matrix multiplication, using actual GPU execution speed as the reward signal. The system solves the problem that different matrix dimension triplets (M, N, K) require different optimization strategies, and manually tuning kernels for thousands of such configurations across GPU architectures is impractical — CUDA-L2 automates this by having an LLM generate, execute, profile, and iteratively improve kernel code through a multi-stage RL curriculum.

3.2 Big-Picture Architecture (Diagram in Words)

The system has five major components connected in a training pipeline:

1. Continued Pretraining Module — extends a base LLM (DeepSeek 671B) with diverse CUDA code from web sources and established libraries, paired with LLM-generated instruction descriptions and retrieved documentation, so the model acquires general CUDA optimization knowledge before RL training begins.

2. General Kernel RL Module — trains the LLM through contrastive reinforcement learning on approximately 1,000 kernels from established libraries spanning diverse operations (linear algebra, convolutions, reductions, attention), using execution speed as reward to build broad optimization capability.

3. HGEMM RL Module — fine-tunes the RL-trained model specifically on half-precision matrix multiplication across varying (M, N, K) configurations, incorporating detailed NCU profiling metrics and a correctness-penalized reward function to generate production-quality matmul kernels.

4. Kernel Validation Pipeline — checks generated kernels for executability (compilation, memory safety via compute-sanitizer) and correctness (exact-match with binary inputs against FP32 CPU reference, plus baseline-bounded deviation against NVIDIA's cuBLAS/cuBLASLt kernels) before accepting them as valid.

5. Evaluation Harness — benchmarks generated kernels against baselines (torch.matmul, cuBLAS, cuBLASLt-heuristic, cuBLASLt-AutoTuning) under both offline (back-to-back execution) and server (random-interval) modes, using the timing protocol from CUDA-L1 that prevents reward hacking via CUDA stream manipulation or Python lazy evaluation.

Information flows as follows: diverse CUDA code → continued pretraining → base LLM with CUDA knowledge → general kernel RL (broad optimization skills) → HGEMM RL (specialized matmul optimization, with NCU profiling feedback) → valid, correct HGEMM kernels → benchmark evaluation against baselines.

3.3 Roadmap for the Deep Dive

• First, the continued pretraining stage (Section 3.2.1), because it establishes the foundation model's knowledge of CUDA programming, CuTe abstractions, and optimization patterns — without this, the RL stages would start from an LLM that has never seen high-performance CUDA code.
• Second, the general kernel RL stage (Section 3.2.2), which builds broad optimization capability before specialization — this curriculum prevents the RL from converging to narrow local optima in the HGEMM stage.
• Third, the HGEMM RL stage (Section 3.2.3), which is the core contribution: how the reward function, NCU profiling, and contrastive RL strategy are configured specifically for matrix multiplication optimization.
• Fourth, the kernel correctness framework (Section 2.3), which must be understood before the RL reward makes sense — the system validates kernels through a two-tier approach (binary-input exact match and baseline-bounded deviation) that handles floating-point non-associativity.
• Fifth, the evaluation protocol (Section 2.4), including the anti-hacking measures, offline vs. server modes, and the timing methodology — these design choices directly affect the RL reward signal and must be explained to understand why the system generates reliable kernels rather than exploiting timing artifacts.

3.4 Detailed, Sentence-Based Technical Breakdown

This is primarily a systems + reinforcement learning paper whose core idea is that LLM-guided RL with a multi-stage curriculum — progressing from general CUDA knowledge acquisition to specialized matmul optimization, and using rich profiling feedback rather than just end-to-end timing — can discover HGEMM kernel implementations that outperform NVIDIA's strongest auto-tuning baselines.

Continued Pretraining on Diverse CUDA Code

What problem this solves. The base LLM (DeepSeek 671B), despite being trained on a broad corpus of code, lacks sufficient exposure to (a) the specific idioms of high-performance CUDA programming, (b) recent library abstractions like CuTe (CUTLASS's tensor-level programming model), and (c) the relationship between kernel structure and GPU microarchitectural behavior. A general-purpose code LLM asked to write an HGEMM kernel might produce syntactically correct CUDA that compiles but achieves only a fraction of the GPU's theoretical throughput because it lacks the specialized knowledge that human CUDA engineers accumulate through years of practice. Continued pretraining on a curated corpus of diverse, high-quality CUDA code fills this knowledge gap before RL training begins.

Data collection. CUDA-L2 collects training data from two complementary sources. The first is web-sourced CUDA code, which is "cleaned, extracted, and segmented using a combination of rule-based filtering and LLM-based cleaning" (Section 3.2.1). Web code provides diversity — it covers many different kernel types, optimization strategies, and coding styles — but is noisy, often lacking documentation or context. The second source is implementations from established libraries: "PyTorch, ATen, CUTLASS, NVIDIA's tutorials and examples, etc." (Section 3.2.1). These provide high-quality, well-structured examples of production CUDA code but cover a narrower range of operations.

Instruction generation. Raw CUDA code is not directly usable for instruction tuning (the training paradigm where the model learns to map natural language instructions to code). Instruction tuning requires descriptive prompts — "Write a CUDA kernel that performs tiled matrix multiplication with shared memory and warp-level primitives" — paired with the corresponding implementation. Since such prompts rarely accompany CUDA code in the wild, CUDA-L2 uses Claude Sonnet 4 to automatically generate instruction descriptions for each collected CUDA code snippet (Section 3.2.1). This converts unstructured code into instruction-tuning data: for each kernel, the LLM produces a natural language description of what the kernel does, what optimization techniques it uses, and what constraints it handles.

Retrieval augmentation. For each generated instruction, CUDA-L2 "uses [it] as a search query to retrieve relevant documentation and code examples in the search engine, which are concatenated as additional context" (Section 3.2.1). This serves two purposes. First, it provides the model with authoritative documentation (e.g., CUDA programming guide sections, CuTe API references) that supplements the code examples, teaching the model why certain patterns work rather than just what they look like. Second, and more subtly, it trains the model to use retrieved context during inference — a capability that becomes critical later when the RL agent needs to reason about new GPU architectures or library versions not seen during pretraining. The resulting training data consists of instruction-context-code triplets: a natural language instruction, retrieved supporting documentation, and the corresponding CUDA implementation.

Training procedure. These triplets are used for continued pretraining on DeepSeek 671B. "Continued pretraining" means the model is trained with a standard language modeling objective (next-token prediction) on this specialized corpus, rather than fine-tuned with an instruction-following objective. This choice is deliberate: continued pretraining integrates the new knowledge into the model's parametric memory more deeply than supervised fine-tuning, which primarily teaches the model to follow a specific output format. The model learns the statistical patterns of high-performance CUDA code — common tiling patterns, shared memory usage, synchronization idioms, CuTe abstractions — as first-class linguistic knowledge rather than as a mapping from a narrow instruction template.

General Kernel RL: Building Broad Optimization Capability

Why a curriculum matters. The paper argues that jumping directly from continued pretraining to HGEMM RL would be suboptimal. HGEMM is an extremely challenging optimization target — the kernels are complex (involving multi-level tiling, multi-stage pipelining, warp-level matrix multiply-accumulate operations, and careful shared memory management), and the performance landscape is rugged (small changes in tile sizes or pipeline depth can cause large, discontinuous changes in throughput due to register pressure and occupancy effects). Starting RL directly on this problem would likely cause the model to converge to a poor local optimum — perhaps discovering one moderately effective tiling strategy for one dimension range and never exploring beyond it. The general kernel RL stage serves as a curriculum: the model first learns general optimization principles on a diverse set of simpler kernel types, then applies those principles when tackling the harder HGEMM problem.

Training data. The general kernel RL stage uses "roughly 1K CUDA kernels from established libraries, covering a large range of operations including linear algebra, convolution operations, reduction operations, element-wise operations, attentions, sampling, and others that cannot be readily categorized (e.g., embedding lookups, loss functions, gradient clipping, optimizer steps)" (Section 3.2.2). Each kernel is paired with an official or successful reference implementation from PyTorch, ATen, CUTLASS, or similar — this provides both a correctness oracle and a performance baseline. The diversity is intentional: a convolution kernel teaches the model about sliding-window memory access patterns, a reduction kernel teaches hierarchical parallelism and warp shuffles, an attention kernel teaches tiling with multiple input matrices. By training across this diverse set, the model builds a repertoire of optimization techniques that it can later recombine for HGEMM.

Contrastive RL strategy. CUDA-L2 adopts the contrastive RL strategy introduced in CUDA-L1. In this approach, the model is not simply rewarded for producing fast kernels and penalized for slow ones — instead, it is explicitly prompted to perform comparative analysis of previously generated CUDA variants and their execution performances. The prompt includes: (1) the current kernel's CUDA code, (2) one or more previous kernel variants (earlier attempts at the same task), (3) the execution speed of each variant, and (4) optionally, profiling metrics that explain why one variant is faster than another. The model is then asked to generate an improved kernel that synthesizes the strengths of previous variants while avoiding their weaknesses. This is fundamentally different from standard RL where the model receives a scalar reward and must implicitly learn which aspects of its output contributed to success or failure. The contrastive approach makes the learning signal explicit: "variant A was 30% faster than variant B because it used double-buffered register fragments, which overlapped memory loads with tensor core computation."

RL algorithm and reward. The paper uses GRPO (Group Relative Policy Optimization) for LLM parameter updates, consistent with the approach in DeepSeek-R1 and CUDA-L1. The reward for a generated kernel is "the average speedup score across all test iterations" (Section 3.2.2), where the speedup score for a single run is:

$s(\text{custom}) = \frac{t_{\text{ref}}}{t_{\text{custom}}} - 1$

where $t_{\text{ref}}$ is the execution time of the reference kernel and $t_{\text{custom}}$ is the execution time of the model-generated kernel.

What it computes: the fractional improvement in execution time relative to the reference. If the custom kernel takes 80% of the reference time, $s = 1.25 - 1 = 0.25$, indicating a 25% speedup. If the custom kernel is slower than the reference, $s$ is negative.

Why this form: the subtraction of 1 centers the reward at zero, making zero the "break-even" point. Positive rewards mean improvement over the baseline; negative rewards mean regression. This centering is important for RL because it creates a natural threshold — the agent learns to avoid generating kernels worse than the reference. An alternative would be to use raw execution time as a cost (lower is better), but speedup ratios are scale-invariant across different problem sizes (a 10ms savings on a 100ms kernel and a 1ms savings on a 10ms kernel both represent a 10% improvement), making the reward signal more consistent across the diverse kernel types in the general training set.

Reward smoothing and clipping. Following CUDA-L1, "rewards will be smoothed and clipped to alleviate the effect of reward hacking during RL training" (Section 3.2.2). Reward hacking occurs when the agent discovers a way to achieve high reward that does not correspond to genuine improvement — in the CUDA context, this could mean generating kernels that exploit timing measurement artifacts (e.g., returning early before computation completes, offloading work to asynchronous streams that aren't synchronized). Smoothing (averaging rewards over multiple evaluations) reduces the variance from GPU timing noise, making the reward signal more reliable. Clipping (capping extreme rewards) prevents the model from overfitting to a single configuration where it discovered a lucky-but-fragile optimization that doesn't generalize.

HGEMM RL: Specialized Matmul Optimization

What changes from general kernel RL. The HGEMM RL stage narrows the training distribution to only half-precision matrix multiplication kernels with varying (M, N, K) configurations. This specialization allows the RL process to discover optimizations specific to matmul — such as the relationship between tile sizes (BM, BN, BK) and matrix dimensions, the interaction between pipeline stage count and the K dimension, and the conditions under which block swizzling is beneficial. The model has already learned general CUDA optimization principles from the previous stage; now it learns how those principles interact with the specific computational pattern of matrix multiplication.

Reward function. The HGEMM reward extends the speedup-based reward with two additional terms that address specific challenges of matmul kernel generation:

$r(\text{custom}) = \frac{1}{N} \sum_{i=1}^{N} \left( \frac{t_i^{\text{ref}}}{t_i^{\text{custom}}} - \alpha \cdot \text{diff}_i \right) - \beta L(\text{custom})$

where $N$ is the number of test iterations, $t_i^{\text{ref}}$ is the reference kernel execution time for iteration $i$, $t_i^{\text{custom}}$ is the custom kernel execution time for iteration $i$, $\text{diff}_i = \max_j |\text{out}_i^{\text{FP32}}[j] - \text{out}_i^{\text{custom}}[j]|$ is the maximum element-wise absolute difference between the FP32 CPU reference output and the custom kernel output, $\alpha > 0$ is the correctness penalty coefficient, $L(\text{custom})$ is the length of the generated code (in characters or tokens), and $\beta > 0$ is the code length penalty coefficient.

What it computes in operational English: for each of $N$ test runs, the system measures the custom kernel's execution time relative to the reference, computes the worst-case numerical deviation between the custom kernel's output and the mathematically correct FP32 result, subtracts a penalty proportional to that deviation, averages these penalized speedups across all runs, and then subtracts an additional penalty proportional to the length of the generated CUDA code. The result is a single scalar that rewards speed, correctness, and conciseness simultaneously.

Why this form — three penalty terms explained:

Correctness penalty ($-\alpha \cdot \text{diff}_i$). Pure speed optimization without correctness constraints would produce kernels that run fast but produce wrong answers — for instance, a kernel that skips computation entirely and returns zeros would be extremely fast but useless. The correctness penalty scales with the maximum element-wise error: a kernel that is slightly off (diff $\approx 10^{-3}$) receives a small penalty, while a kernel that produces completely wrong outputs (diff large) receives a severe penalty. This creates a continuous gradient that allows the RL agent to improve correctness gradually — it can start with approximately correct kernels and refine them, rather than needing to be perfectly correct from the first attempt. The choice of $\alpha$ controls the tradeoff: too small, and incorrect-but-fast kernels dominate the reward; too large, and the agent becomes overly conservative, never exploring novel optimizations that might initially introduce small numerical errors.

Code length penalty ($-\beta L(\text{custom})$). Shorter code is preferred for two reasons. First, shorter CUDA kernels generally have lower instruction cache pressure and are less likely to contain dead code or redundant operations. Second, and more importantly for the RL process, shorter code is easier for the model to generate correctly — longer kernels have more opportunities for bugs, and the RL agent might converge to bloated implementations that achieve speed through unnecessary complexity rather than genuine optimization insight. The penalty term pushes the model toward elegant, concise solutions. The choice of $\beta$ is subtle: too large, and the model never uses sophisticated library abstractions (CuTe) that require some boilerplate; too small, and the model generates verbose kernels with unnecessary branches and redundant memory operations.

Speedup term ($\frac{t_i^{\text{ref}}}{t_i^{\text{custom}}}$). Same form as the general kernel RL reward — ratio-based for scale invariance across different (M, N, K) configurations.

NCU profiling metrics in the RL context. A critical enhancement over CUDA-L1 is the inclusion of detailed NVIDIA Nsight Compute (NCU) profiling metrics in the context provided to the RL agent. Rather than only seeing end-to-end execution time, the model sees:

• Memory throughput: how many bytes per second are being transferred between global memory, shared memory, and registers. Low memory throughput despite high compute utilization suggests the kernel is compute-bound and could benefit from reducing redundant operations; high memory throughput with low compute utilization suggests the kernel is memory-bound and would benefit from better data reuse (tiling) or prefetching.
• Compute utilization (SM occupancy): what fraction of the GPU's streaming multiprocessors are actively executing instructions. Low occupancy often indicates insufficient parallelism — the tile sizes may be too large, limiting the number of thread blocks that can run concurrently.
• Warp occupancy: the fraction of warps that are active versus stalled (waiting for memory or synchronization). High stall rates indicate the kernel is latency-bound — data is not arriving fast enough to keep the compute units busy.
• Cache hit rates: what fraction of memory accesses are served from L1/L2 cache versus requiring global memory reads. Low cache hit rates suggest poor data locality, which can be addressed through block swizzling or improved tiling strategies.

The model is prompted with these metrics alongside execution time, enabling it to perform root-cause analysis. For example, if the model sees that a kernel has 95% compute utilization but only 30% of theoretical memory bandwidth, it can infer that the kernel is compute-efficient but memory-inefficient — a signal to explore larger tiles or better prefetching patterns. Without this granular feedback, the RL agent would only know that the kernel is "slow" without understanding why, making it essentially a random search in a high-dimensional optimization space.

Contrastive prompts with profiling. The contrastive RL strategy is extended to include profiling metrics for each kernel variant. The prompt structure is:

1. The task description: "Generate an HGEMM kernel for (M=2048, N=1024, K=4096) using CuTe abstractions."
2. Previous variant A's code, execution time (e.g., 1.2ms), and NCU metrics (e.g., "memory throughput: 800 GB/s, SM occupancy: 65%, L2 hit rate: 42%").
3. Previous variant B's code, execution time (e.g., 0.9ms), and NCU metrics (e.g., "memory throughput: 1,100 GB/s, SM occupancy: 78%, L2 hit rate: 58%").
4. A comparative analysis: "Variant B outperforms Variant A by 33%. The profiling shows B has 37.5% higher memory throughput and 38% higher L2 cache hit rate, suggesting that B's tiling strategy achieves better data locality..."

The model then generates a new variant that (ideally) synthesizes the profiling insights. This makes the optimization process interpretable: the model is not just trying random perturbations and observing the reward, but reasoning about the causal relationship between code structure and microarchitectural behavior.

Code generation constraints. Generated HGEMM kernels are restricted to CUDA C/C++, CuTe abstractions, inline PTX assembly, CUDA intrinsics, and CUTLASS templates — explicitly excluding Python-based DSLs like Triton (Section 3.2.3). This is a deliberate design choice. Triton compiles to PTX through an intermediate representation, which limits access to low-level hardware features that can only be expressed in CUDA C++ or inline PTX. For example, precise control over register allocation, explicit warp-level matrix multiply-accumulate (WMMA) instructions, and certain shared memory bank conflict avoidance patterns require CUDA C++ or PTX. By restricting to .cu files compiled with nvcc, CUDA-L2 ensures that the search space includes all optimizations available to human CUDA engineers, not just those expressible in a higher-level DSL.

Kernel compilation and execution. Generated kernels are parsed into .cu files and compiled with nvcc (Section 3.2.3). The compilation step serves as an implicit filter: kernels that fail to compile (syntax errors, type mismatches, invalid template instantiations) receive a negative reward and become negative examples in the contrastive RL training. This creates a tight feedback loop — the model learns to avoid common compilation errors (e.g., mismatched CuTe tensor shapes, incorrect cute::copy usage) through RL optimization, converging toward a space of syntactically valid, compilable kernel code.

Kernel Correctness Validation

The floating-point challenge. Validating the correctness of a matrix multiplication kernel is more subtle than simply comparing outputs element-by-element against a reference. Floating-point arithmetic is non-associative: $(a + b) + c \neq a + (b + c)$ due to rounding. This means that two correct implementations of matrix multiplication — using different tiling strategies, accumulation orders, or even different thread block schedules — will produce slightly different floating-point outputs. A naive error threshold (e.g., "all elements must match to within $10^{-5}$") would either reject valid kernels (if the threshold is too tight) or accept incorrect kernels (if too loose). CUDA-L2 solves this through a two-tier validation strategy.

Tier 1: Exact match with binary inputs. The first validation tier uses a clever property of half-precision representation. The system generates random matrices $A$ and $B$ with elements being binary values from $\{0, 1\}$ (Section 2.3.2). For matrix multiplication $C = A \times B$, each output element is:

$c_{ij} = \sum_{k=1}^{K} a_{ik} \cdot b_{kj}$

where $a_{ik}$ and $b_{kj}$ are binary, so each product is either 0 or 1, and the sum is guaranteed to be a non-negative integer.

The reference output $C^{\text{ref}}$ is computed using FP32 on CPU, which provides exact integer results (FP32 can represent all integers up to $2^{24}$ exactly). The custom kernel computes $C^{\text{custom}}$ using half-precision on GPU.

For each output position $(i, j)$ where $c_{ij}^{\text{ref}} < 2048$, the system requires exact equality: $c_{ij}^{\text{custom}} = c_{ij}^{\text{ref}}$. Positions where $c_{ij}^{\text{ref}} \geq 2048$ are ignored.

What property makes this work. Half-precision (FP16) has 10 mantissa bits plus 1 implicit leading bit, yielding 11 bits of significand precision. This means all integers in $[0, 2048)$ — that is, integers from 0 to 2047 — are exactly representable in FP16. When both input matrices are binary, each partial sum $\sum_{k=1}^{t} a_{ik} \cdot b_{kj}$ is monotonically non-decreasing as $t$ increases (since all terms are 0 or 1). If the final sum is below 2048, then all intermediate partial sums are also below 2048, meaning every intermediate value was exactly representable throughout the accumulation. Therefore, any deviation between the custom kernel and the reference must be due to an implementation error, not floating-point rounding.

Why $\{0, 1\}$ inputs specifically. The paper explicitly notes (footnote 3) that sampling from $\{-1, 0, 1\}$ would not work because partial sums could temporarily exceed 2048 before later negative terms bring the final result below 2048. When the intermediate sum exceeds 2048, exactness is already lost, and the final correct result does not guarantee intermediate correctness. The binary-input approach elegantly sidesteps this by ensuring monotonicity.

What about non-integer values. The paper addresses a natural question: could we use fractional values like 1/2, 1/4, or 1/8? The answer is that this merely shifts the threshold without fundamentally changing the analysis (footnote 4). In FP16, the unit in the last place (ULP) — the smallest representable difference between consecutive numbers — depends on the magnitude. For half-integer sums, the ULP is 1/2 for magnitudes in $[512, 1024)$, but increases to 1 for magnitudes in $[1024, 2048)$. So half-integers can only be exactly represented up to 1024, not 2048. The binary-input approach avoids this complexity entirely by guaranteeing integer intermediate results and a clean threshold at 2048.

Practical considerations. The binary probability for $\{0, 1\}$ is "adjusted based on matrix sizes to ensure a significant proportion of $c_{ij}^{\text{ref}}$ is below 2048 but larger than 0" (Section 2.3.2). For very large matrices ($K$ in the thousands), the sum across $K$ terms of 0/1 random variables has expected value $K \cdot p$, where $p$ is the probability of 1. To keep most outputs below 2048, the probability is set appropriately low for large $K$. The test is repeated multiple times with different random inputs; if any single iteration fails, the kernel is considered incorrect.

Tier 2: Baseline-bounded deviation. The binary-input test is powerful but has a limitation: it only validates correctness for a specific input distribution (binary values). A kernel might pass the binary test but produce unacceptably large errors for general floating-point inputs — for instance, due to a numerically unstable accumulation order that amplifies rounding errors. The second validation tier addresses this by leveraging NVIDIA's own kernels as a reference for acceptable deviation levels.

The procedure is: select a set of "highly reliable baseline kernels developed by NVIDIA" including cuBLAS-NN, cuBLAS-TN, cuBLASLt-heuristic-NN, cuBLASLt-heuristic-TN, cuBLASLt-AutoTuning-NN, and cuBLASLt-AutoTuning-TN. For each test input (with general floating-point values, not just binary), compute the maximum element-wise difference among these baseline kernels — call this $\text{max\_diff}_{\text{baselines}}$. This value represents the upper bound of variability for floating-point matrix multiplication: even NVIDIA's own kernels disagree with each other within this range due to different accumulation orders and tiling strategies. A custom kernel is considered incorrect if its maximum element-wise deviation from the FP32 CPU reference exceeds this baseline maximum:

$\text{incorrect if: } \max_j \left|\text{out}^{\text{FP32}}[j] - \text{out}^{\text{custom}}[j]\right| > \max_{\text{baseline pairs}} \max_j \left|\text{out}^{\text{baseline}_a}[j] - \text{out}^{\text{baseline}_b}[j]\right|$

Why this tier exists. The binary-input test catches implementation bugs (wrong indexing, incorrect tiling boundaries, logic errors) because those produce exact mismatches. The baseline-bounded test catches numerical instability — accumulation orders so pathological that they produce errors larger than what any reasonable implementation would exhibit. Together, the two tiers provide strong confidence that a kernel is both logically correct and numerically well-behaved.

Evaluation Protocol and Anti-Hacking Measures

Timing measurement. The paper uses the standard kernel timing strategy from CUDA-L1 and KernelBench, shown in Listing 1 of the paper:

torch.cuda.synchronize()
start_event = torch.cuda.Event(enable_timing=True)
end_event = torch.cuda.Event(enable_timing=True)
start_event.record()
kernel(a, b, b_col_major, out)
end_event.record()
torch.cuda.synchronize()
elapsed_time_ms = start_event.elapsed_time(end_event)

The torch.cuda.synchronize() before recording the start event ensures all previous GPU operations have completed. The start event is recorded immediately before the kernel launch. The end event is recorded immediately after. The second torch.cuda.synchronize() blocks the CPU until the end event has been recorded (i.e., until the kernel has completed). The elapsed time between the two events captures only the kernel's GPU execution time, excluding CPU-side launch overhead.

Anti-hacking measures. CUDA-L1 identified two strategies that RL agents discovered to "hack" this timing measurement and achieve artificially high speedup scores without actually computing the correct result faster:

1. Asynchronous CUDA streams: the kernel could create an additional CUDA stream and offload computation to it without synchronizing, so the timing measurement captures only the stream creation and kernel launch, not the actual computation.
2. Python lazy evaluation: when kernels are written as Python functions (e.g., using PyTorch's torch.utils.cpp_extension), the kernel call may not force materialization of the output — Python's lazy evaluation could defer the actual computation until some later point after the timing measurement completes.

CUDA-L2 prevents both attacks by construction: "(1) disallowing additional CUDA stream creation; (2) generating HGEMM kernels only as CUDA code in .cu files, which naturally bypasses Python's lazy evaluation" (Section 2.4.1). The restriction to .cu files and nvcc compilation is not just about performance — it is also about evaluation integrity. A .cu file compiled with nvcc produces a native binary that executes synchronously with respect to the calling context, making timing measurements honest.

Offline vs. server mode. The paper evaluates kernels under two scenarios following the MLPerf benchmarking standard:

• Offline mode: "kernels are executed back-to-back without pauses" (Section 2.4.2). This measures peak throughput by keeping the GPU continuously busy. In offline mode, the GPU remains in a steady thermal state — clock frequencies stabilize after the warmup period, caches remain warm from previous executions, and the memory controller is continuously active. This provides the best-case performance for a well-optimized kernel.

• Server mode: "kernels are executed at intervals to simulate real-time inference" (Section 2.4.2). A critical detail: "the interval time is not included in the execution time used for the speedup comparison." The timing measurement captures only the kernel's active execution, not the idle time between requests. Server mode is more demanding because "GPU's caches can cool down and then start from a cold state," and the GPU may experience "clock boosting followed by thermal throttling" as it transitions between idle and active states (Section 4.1).

The paper attributes the larger speedups observed in server mode (28.7% over torch.matmul vs. 22.0% offline) to thermal dynamics: CUDA-L2's kernels appear to handle cold-start conditions better than the baselines, possibly due to different memory access patterns that are less dependent on cache warmth. This is an important practical finding because real LLM inference deployments more closely resemble server mode than offline mode — requests arrive at irregular intervals, and the GPU is not continuously saturated.

Evaluation duration and repetitions. Each evaluation "runs for a minimum of 30 seconds after a 10-second warmup period" (Section 2.4). The warmup period ensures the GPU reaches steady-state thermal conditions before measurements begin. The 30-second minimum provides a large number of kernel executions (for a kernel taking 1ms, this yields 30,000 samples), averaging out GPU timing noise. Within each evaluation, execution order is randomized to "eliminate ordering effect" — if kernel A always runs before kernel B, systematic biases from GPU state could favor one over the other; randomization ensures fair comparison.

The final speedup score. The reported speedup for a kernel is:

$\text{speedup}(\text{custom}) = \frac{1}{N} \sum_{i=1}^{N} \left( \frac{t_i^{\text{ref}}}{t_i^{\text{custom}}} - 1 \right)$

where $N$ is the number of runs in the evaluation window. This is the "mean speed score over all runs" (Section 2.4) — the expected fractional improvement in execution time.

4. Key Insights and Innovations

Innovation 1: The Configuration-Generalization Gap as the Central Unsolved Problem in Automatic Kernel Optimization

The paper's most conceptually significant move is not the specific techniques it builds, but its diagnosis of what makes HGEMM optimization genuinely unsolved despite decades of effort. Prior work on LLM-based kernel generation — KernelBench, AI CUDA Engineer, CUDA-L1 — treated kernel optimization as a task where success means beating a baseline on a single fixed input configuration. The implicit assumption was that if an LLM can optimize a matmul for (M=1024, N=1024, K=1024), it has "solved" matmul optimization for that kernel type.

CUDA-L2 argues that this framing misses the central challenge. The paper demonstrates, through its 1,000-configuration evaluation, that optimization strategies do not transfer across dimensions within the same kernel type on the same GPU architecture. The evidence is concrete: for (M=8192, N=512, K=2048), BM=160 with zero-padding achieves +15.2% over cuBLASLt-AutoTuning, but changing BM to 128 — a conventional choice that evenly divides 8192 — drops the gain to 0.4%, and BM=256 causes a -15.7% regression (Section 5.2). These are not small differences in a continuous optimization landscape; they are discontinuous jumps where the "obvious" parameter choices are dramatically suboptimal, and the optimal choice depends on the specific (M, N, K) triplet in non-obvious ways.

This is a reframing of the problem definition, not an incremental technique. Before CUDA-L2, the field evaluated kernel optimizers on point estimates (one configuration, one speedup number). After CUDA-L2, the relevant metric becomes distributional: how does the optimizer perform across the full span of configurations that production models actually encounter? The 1,000-configuration benchmark — covering all triplets from the dimension set used by Qwen, Llama, and DeepSeek — is itself a methodological contribution. It establishes a standard for what "solving HGEMM optimization" means: not excellence on a few cherry-picked sizes, but systematic superiority across the configuration space.

This reframing has downstream implications that extend beyond this paper. It suggests that prior benchmark-focused kernel optimization work (KernelBench's single-configuration evaluations) may produce kernels that are brittle — tuned to the benchmark's specific dimensions but not robust to the dimensional variation encountered in real models. It also explains why NVIDIA's cuBLASLt-AutoTuning, which exhaustively benchmarks up to 100 candidates per configuration, represents the state of the art: it implicitly acknowledges that no single kernel works everywhere, and the best we can do is select among hand-designed variants per configuration. CUDA-L2's contribution is to push beyond selection into generation — creating novel kernels specialized per configuration — and to demonstrate that this yields gains even over exhaustive selection.

Innovation 2: Multi-Stage Curriculum RL as a Solution to the Rugged Optimization Landscape of Low-Level GPU Code

The paper introduces a curriculum learning strategy for RL-based kernel generation — progressing from continued pretraining on diverse CUDA code, through general kernel RL across many operation types, to specialized HGEMM RL — and this is more than an engineering convenience. It represents a principled solution to a fundamental challenge: the optimization landscape for low-level GPU kernels is rugged and deceptive.

To understand why this matters, consider the alternative that the paper explicitly avoids: training RL directly on HGEMM from a base LLM. The HGEMM optimization space has the following properties: (a) small code changes — swapping BM from 128 to 160, changing the number of pipeline stages from 3 to 5 — can cause large, discontinuous changes in throughput; (b) the relationship between code structure and performance is mediated by complex microarchitectural interactions (register pressure, shared memory bank conflicts, warp scheduling) that are not directly observable from end-to-end timing alone; and (c) most randomly generated kernels are either uncompilable, incorrect, or orders of magnitude slower than cuBLAS, providing essentially no learning signal. In RL terms, this is a sparse-reward, high-dimensional, non-smooth optimization problem — exactly the kind of problem where RL agents converge to poor local optima or fail to learn entirely.

The multi-stage curriculum addresses each of these challenges. The continued pretraining stage ensures the model starts from a parameterization where generating compilable, syntactically valid CUDA using sophisticated abstractions (CuTe, WMMA, inline PTX) is already within its capability distribution — the base probability of generating a valid kernel is non-negligible. The general kernel RL stage provides dense learning signal across a diverse set of operation types where the optimization landscape is less extreme (a convolution kernel's performance varies more smoothly with tile size than an HGEMM kernel's). This builds a repertoire of optimization techniques — tiling, pipelining, shared memory management — that transfer to HGEMM. By the time the model reaches HGEMM RL, it has already learned the "grammar" of CUDA optimization; the specialized stage only needs to adapt these general principles to the specific constraints of matrix multiplication.

This is fundamentally different from prior LLM-based kernel optimization (CUDA-L1, AI CUDA Engineer), which used single-stage training without a curriculum. The conceptual contribution is recognizing that CUDA kernel optimization is not one problem but a hierarchy of related problems, and that RL training should respect this hierarchy. The paper does not frame this as its primary contribution — it is described as an extension of CUDA-L1 — but the architecture implies a general principle: when the target domain has a sparse and deceptive reward landscape, train first on related domains with denser reward, then specialize. This principle likely transfers to other low-level code optimization tasks (e.g., GPU kernel generation for convolutions, attention, or custom ops) and potentially to other RL-for-code-generation domains.

The evidence supporting this is implicit in the results: CUDA-L2 achieves its 11.4% speedup over cuBLASLt-AutoTuning and generates sophisticated, non-obvious optimizations (double-buffered register fragments, zero-padding for tile-size flexibility, staggered A-B prefetch scheduling). A single-stage RL approach on HGEMM would be unlikely to discover these optimizations because the probability of randomly exploring the specific code transformations that yield them, starting from a base LLM with no CUDA specialization, is vanishingly small. The curriculum provides the necessary scaffolding.

Innovation 3: Profiling Metrics as Interpretable Learning Signal — Moving RL Beyond Black-Box Reward

CUDA-L2's incorporation of NCU profiling metrics into the RL context is a methodological advance in how RL agents learn to optimize performance-critical code. Prior work (CUDA-L1) used only end-to-end execution time as reward — a scalar that tells the agent whether a kernel is fast or slow but provides no information about why. This makes RL a black-box search: the agent proposes code, receives a speed number, and must implicitly infer which aspects of its output contributed to success or failure through many trials.

The NCU integration transforms the learning signal from black-box to gray-box. The model sees memory throughput, SM occupancy, warp occupancy, and cache hit rates alongside execution time. When variant A is 30% faster than variant B, the model also sees that A achieved 37.5% higher memory throughput and 38% higher L2 cache hit rate. This enables attribution: the model can hypothesize that the speedup is due to better data locality rather than, say, better instruction scheduling. More importantly, it enables targeted improvement: the model can be prompted to "improve memory throughput" or "reduce warp stalls" as explicit sub-goals, rather than generically "make the kernel faster."

This is conceptually analogous to the difference between model-free and model-based RL, but applied to code generation. A model-free code optimizer (CUDA-L1) learns a direct mapping from task description to kernel code that maximizes reward. A model-based code optimizer would learn an internal model of how code structure affects microarchitectural behavior, then use that model to plan improvements. NCU profiling gives CUDA-L2 a partial model — not learned, but provided by the hardware profiler — of the causal chain from code to performance. The RL agent doesn't need to infer that shared memory bank conflicts cause reduced memory throughput; the profiler tells it directly, and the agent learns to avoid patterns that produce those profiler signatures.

The significance extends beyond the specific HGEMM results. This approach is generalizable to any performance-critical code generation task where hardware profilers exist: CPU kernel optimization (Linux perf), GPU optimization for other architectures (Nsight Compute on Hopper/Blackwell), and potentially even non-CUDA domains like FPGA HLS optimization (where synthesis reports provide analogous resource utilization and timing metrics). The conceptual contribution is recognizing that modern hardware profilers provide interpretable intermediate representations of performance that can serve as a learning curriculum for RL agents, bridging the gap between high-level code structure and low-level execution behavior.

Innovation 4: The Discovery of Non-Standard Optimization Strategies as Evidence for RL's Exploration Advantage

While the paper reports performance numbers as its primary result, a subtler but equally important contribution is the catalog of optimization techniques discovered by CUDA-L2 that deviate from standard practice (Section 5). These are not techniques that the system was explicitly programmed to explore — they emerged from the RL process as the agent found non-obvious parameter combinations that outperformed conventional wisdom.

The two most striking examples:

Zero-padding to expand the tile-size search space (Section 5.2). Standard tiled matrix multiplication requires that the matrix dimension M be divisible by the tile size BM to avoid out-of-bounds memory accesses at boundary tiles. This constrains BM to divisors of M — for M=8192, the natural choices are {64, 128, 256}. CUDA-L2 discovered that padding M to 8320 (a ~1.6% overhead) allows BM=160, a tile size that is not a divisor of the original dimension. The 1.6% computational overhead from processing extra rows is more than offset by the performance gain from the more efficient tile size, yielding a net 15.2% speedup. This tradeoff — paying a small computational cost to unlock a better parameter choice — is non-obvious because it violates the standard assumption that tile sizes should divide matrix dimensions exactly.

Staggered A-B prefetch scheduling (Section 5.3.4). Standard practice in tiled matrix multiplication is to prefetch both A and B tiles for the next iteration consecutively, then perform the matrix multiply-accumulate (MMA) on the current iteration's data. CUDA-L2 discovered that splitting the A and B prefetches around the MMA operation — prefetch A, perform MMA, then prefetch B — can improve instruction-level parallelism by allowing the A prefetch and MMA to overlap in the pipeline. This is a scheduling optimization that depends on the relative latencies of memory loads and tensor core operations, and its benefit varies with problem dimensions and pipeline depth. The standard consecutive-prefetch approach is simpler and easier to reason about; the staggered approach emerged because the RL agent was free from the human bias toward symmetry and simplicity.

These discoveries are significant not because they are individually transformative (each provides incremental improvement for specific dimension ranges), but because they demonstrate a mode of exploration that differs qualitatively from human optimization. Human CUDA engineers optimize by applying principled heuristics: "choose BM to divide M exactly," "prefetch both operands together," "use as many pipeline stages as K allows." These heuristics are effective but they prune the search space in ways that systematically exclude certain regions — regions where counterintuitive choices (like deliberately misaligning tile sizes or interleaving memory and compute operations asymmetrically) might outperform the heuristic recommendations. RL, unburdened by these heuristics, can explore those regions and discover that the heuristics are locally but not globally optimal.

This is a conceptual contribution about the nature of optimization expertise. It suggests that in domains where the performance landscape is sufficiently complex — where the mapping from design choices to outcomes is mediated by many interacting microarchitectural mechanisms — expert heuristics become a double-edged sword. They accelerate convergence to good solutions but they also create blind spots that automated exploration can exploit. The implication is not that human expertise is obsolete, but that the optimal division of labor may be: humans provide the building blocks (CuTe abstractions, WMMA primitives, tiling frameworks), and automated RL explores the combinatorial space of how to configure and combine those building blocks.

5. Experimental Analysis

Evaluation Methodology

• Dataset. The evaluation uses 1,000 (M, N, K) configurations representing all 10³ = 1,000 triplet combinations where M, N, and K each take values from {64, 128, 256, 512, 1024, 2048, 4096, 8192, 12288, 16384} (Section 1). This set is explicitly justified as covering "those used in attention and FFN layers of widely open-sourced models like Qwen, Llama and DeepSeek" (Section 1). No separate train/validation/test split is used for evaluation — the 1,000 configurations serve as the full test set against which all baselines and CUDA-L2 are benchmarked. For RL training, the model sees HGEMM kernels with varying (M, N, K) during the specialized HGEMM RL stage (Section 3.2.3), but the paper does not specify whether the exact 1,000 evaluation configurations or a superset were used during training. This is a notable omission — if the evaluation configurations were seen during RL training, the reported speedups may partially reflect memorization rather than generalization, though the diverse optimization strategies described in Section 5 suggest genuine capability rather than simple memorization.

• Base model(s). The foundation model is DeepSeek 671B (Section 3.2.1), which undergoes continued pretraining on diverse CUDA code to acquire general CUDA optimization capabilities. The choice of DeepSeek 671B is pragmatic rather than principled — the paper does not justify this specific model over alternatives beyond its role as the developer's existing infrastructure. The ~14× larger model from the prior work (PaLM 2-S*) is not relevant here; the comparison is between CUDA-L2-generated kernels and NVIDIA's library baselines, not between pretrained models. The paper's claim is about kernel quality, not model scaling.

• Metrics. The primary metric is relative speedup, defined as s(custom) = t_ref / t_custom − 1 where t_ref is the reference kernel's execution time and t_custom is CUDA-L2's kernel execution time (Section 2.4, Equation 2). This is a ratio-based metric: a value of 0.22 means the custom kernel is 22% faster (executes in ~82% of the reference time). The final speedup is the mean over all runs within a single evaluation. Execution time is measured using CUDA events with synchronization, as shown in Listing 1 — the measurement captures only GPU kernel execution time, excluding CPU-side launch overhead. Each evaluation runs for a minimum of 30 seconds after a 10-second warmup period (Section 2.4), and within each evaluation, execution order between custom and reference kernels is randomized to eliminate ordering effects. The paper also reports win rates — the fraction of the 1,000 configurations where CUDA-L2's kernel is faster than the baseline — though these are mentioned in prose rather than tabulated (Section 4.1: "The win rates span 79.3% to 95.7% across all baselines"). A secondary metric is the max(CUDA-L2, baseline) speedup (Table 2), which simulates a deployment scenario where the user can select whichever kernel runs faster per configuration.

• Baselines. Four baselines are compared:

  • torch.matmul (PyTorch's default matrix multiplication): internally dispatches to cuBLAS for half-precision but includes overhead from PyTorch's tensor dispatch and memory management (Section 2.2.1). Represents standard practice for users who rely on default PyTorch operations.
  • cuBLAS (NVIDIA's cuBLAS library): evaluated in two layout configurations — NN (normal-normal, both matrices row-major) and TN (transposed-normal, A transposed to column-major). The cublasGemmEx function with CUBLAS_GEMM_DEFAULT_TENSOR_OP is used, which enables Ampere FP16 Tensor Cores and allows cuBLAS's internal heuristics to select the optimal algorithm (Section 2.2.2). Code is shown in Listing 6. The paper also reports cuBLAS-max, which selects the faster of NN and TN per configuration.
  • cuBLASLt-heuristic (NVIDIA's cuBLASLt library, heuristic mode): uses the cublasLtMatmulAlgoGetHeuristic API to query algorithm recommendations, selecting the top-ranked algorithm (index 0) per configuration (Section 2.2.3). Algorithm selection is cached beforehand to avoid repeated API overhead during evaluation. Code is shown in Listing 7. Evaluated in NN, TN, and max layouts.
  • cuBLASLt-AutoTuning (NVIDIA's cuBLASLt library, exhaustive tuning mode): retrieves up to 100 algorithm candidates from cublasLtMatmulAlgoGetHeuristic, benchmarks each with randomized execution order and random input matrices (50 warmup + 100 measurement rounds, as shown in Listing 8), and selects the fastest based on median execution time (Section 2.2.3). This is the strongest baseline and represents the state of the art in production HGEMM performance — it exhaustively tests all available hand-designed kernels but cannot generate novel kernels. Evaluated in NN, TN, and max layouts.

  All cuBLAS and cuBLASLt baselines use the library's native C API, not Python wrappers, to avoid PyTorch overhead.

• Generation budget / compute accounting. CUDA-L2's "compute budget" during inference is not measured in generations (as in LLM sampling) but in the one-time cost of RL training: continued pretraining on CUDA code, general kernel RL across ~1K kernels, and HGEMM RL across varying configurations. Once trained, CUDA-L2 generates a kernel once per (M, N, K) configuration — this is a one-time offline cost, not a per-inference cost. The paper does not report the total GPU-hours or FLOPs consumed by the RL training process. For evaluation, the cost is the per-configuration benchmarking time (minimum 30 seconds after warmup), which is identical for CUDA-L2 and the baselines (cuBLASLt-AutoTuning also requires per-configuration benchmarking, and its setup cost — benchmarking 100 candidates — is substantially higher than CUDA-L2's kernel generation cost). The paper does not quantify these relative setup costs, which is a gap: the claim that CUDA-L2 is "better" might need to account for whether its training cost amortizes favorably across many configurations versus the per-configuration cost of cuBLASLt-AutoTuning's exhaustive search.

• Cross-validation / statistical protocol. No cross-validation is reported. The 1,000 configurations are evaluated exhaustively — there is no held-out set, no k-fold splitting, and no statistical significance testing. The paper reports mean speedup across all 1,000 configurations and win rates, but does not provide confidence intervals, standard errors, or hypothesis tests for whether the observed speedups are statistically significant. This is a significant limitation: with 1,000 configurations and speedup values that vary substantially across dimensions (Figure 3 shows speedups ranging from ~1.4× for small matrices to ~1.0× for large matrices), the mean could be driven by a subset of configurations where CUDA-L2 excels, masking weaker or negative performance on others. The win rate metric partially addresses this (reporting what fraction of configurations show any improvement), but without per-configuration variance estimates, the reliability of the mean speedup is unclear. For kernel evaluation, each configuration is run for a minimum of 30 seconds after a 10-second warmup, providing a large number of timing samples per configuration, but no per-configuration variance is reported in the paper.

Main Quantitative Results

Overall Performance Against Baselines

Table 1 presents the central results: CUDA-L2's average speedup across all 1,000 (M, N, K) configurations against each baseline, in both offline and server modes.

Offline mode (kernels executed consecutively):

• +22.0% over torch.matmul on average. The win rate spans 79.3% to 95.7% across baselines, meaning CUDA-L2 is faster on the vast majority of individual configurations, not just in aggregate.
• +19.2% over cuBLAS-max (the optimal layout per configuration, choosing between NN and TN). For individual layouts: +20.0% over cuBLAS-NN and +21.4% over cuBLAS-TN (Table 1). The NN layout slightly outperforms TN across all baselines, which the paper attributes to better algorithm selection by the library's heuristics for row-major layouts.
• +16.8% over cuBLASLt-heuristic-max. Individual layouts: +17.3% over NN, +19.1% over TN.
• +11.4% over cuBLASLt-AutoTuning-max — the strongest baseline, which exhaustively benchmarks up to 100 algorithm candidates per configuration and selects the fastest. Individual layouts: +12.1% over NN, +13.3% over TN.

Server mode (kernels executed at random intervals simulating real-time inference):

• +28.7% over torch.matmul
• +26.0% over cuBLAS-max (+28.8% NN, +30.2% TN)
• +22.4% over cuBLASLt-heuristic-max (+24.4% NN, +25.9% TN)
• +15.9% over cuBLASLt-AutoTuning-max (+17.9% NN, +19.1% TN)

The server-mode speedups are consistently larger than offline-mode speedups across all baselines — for cuBLASLt-AutoTuning-max, the gap is 15.9% vs. 11.4% (a 4.5 percentage point increase). The paper attributes this to GPU thermal dynamics (Section 4.1): in offline mode, continuous execution keeps the GPU in steady thermal state with predictable clock behavior, while in server mode, idle periods allow the GPU to cool, causing clock boosting followed by thermal throttling when kernels launch. CUDA-L2's kernels appear to handle cold-start conditions better than the baselines, perhaps because their memory access patterns are less dependent on cache warmth.

Key observation about baseline ordering. The speedup hierarchy — torch.matmul (weakest) < cuBLAS < cuBLASLt-heuristic < cuBLASLt-AutoTuning (strongest) — confirms that each baseline represents a genuine step up in optimization sophistication. The gap between CUDA-L2 and the strongest baseline (11.4% offline, 15.9% server) quantifies how much headroom remains even after exhaustive auto-tuning of hand-designed kernels. The gap between torch.matmul and cuBLASLt-AutoTuning itself (22.0% − 11.4% = 10.6 percentage points offline) shows the benefit of moving from default PyTorch operations to exhaustive library tuning — a useful calibration for practitioners deciding how much tuning effort to invest.

Combining CUDA-L2 with Baselines: max(CUDA-L2, baseline)

Table 2 presents a pragmatic deployment scenario: what if users have access to both CUDA-L2's generated kernels and the baseline library, and can select whichever runs faster per configuration? The max(CUDA-L2, baseline) metric computes the mean speedup of this combined approach over the baseline alone.

In offline mode:

• max(CUDA-L2, torch.matmul) yields +23.1% over torch.matmul (up from +22.0% for CUDA-L2 alone)
• max(CUDA-L2, cuBLAS-max) yields +20.2% (up from +19.2%)
• max(CUDA-L2, cuBLASLt-heuristic-max) yields +17.0% (up from +16.8%)
• max(CUDA-L2, cuBLASLt-AutoTuning-max) yields +13.2% (up from +11.4%)

The increases are modest (0.2 to 1.8 percentage points across baselines), indicating that CUDA-L2's kernels are faster than the baselines on most configurations already — the max operation only helps on the minority of configurations where the baseline is still faster. The largest increase (+1.8 points for cuBLASLt-AutoTuning-max) suggests that the strongest baseline retains an advantage on a non-trivial subset of configurations, consistent with the win rate being 79.3% (not 100%).

In server mode, the pattern is similar with slightly larger increases: +1.1 to +2.2 percentage points across baselines. The fact that even the combined approach leaves performance on the table (the speedup is non-zero, meaning the baseline is faster on some configurations) is an honest nuance: CUDA-L2 does not dominate every single configuration, but it dominates the distribution.

Speedup vs. Problem Size

Figure 3 and Table 3 analyze how CUDA-L2's speedup over cuBLASLt-AutoTuning-max varies with matrix dimensions. This is the paper's most analytically rich result, revealing a clear and interpretable pattern.

Figure 3(a): Speedup vs. log2(M × N × K). The total problem size (product of dimensions, a proxy for total FLOPs) is plotted on the x-axis on a log2 scale, ranging from approximately 2^18 (when M=N=K=64, product ≈ 262K) to 2^42 (when M=N=K=16384, product ≈ 4.4 trillion). Speedup on the y-axis ranges from −0.1 to +0.6. A clear negative trend emerges:

• For small problems (log2 ≈ 18–24, corresponding to dimensions of 64–256), speedups cluster around 1.2–1.5× (40–50% improvement).
• For medium problems (log2 ≈ 25–32, dimensions 256–2048), speedups decrease to roughly 1.1–1.2×.
• For large problems (log2 ≈ 33–42, dimensions 4096–16384), speedups converge toward 1.0× (zero improvement), with substantial variance including some negative values (down to roughly −0.05).

The shaded band (mean ± 1 standard deviation) widens at both extremes: small problems show high variance (some configurations achieve +50%, others only +10%), while large problems show variance driven by occasional regressions.

Figure 3(b): Speedup vs. average dimension (M+N+K)/3. This panel controls for the possibility that the total product obscures dimension-specific effects. The pattern is consistent: small average dimensions (< 512) show speedups of 1.2–1.4×, medium dimensions (512–4096) show 1.05–1.15×, large dimensions (> 4096) converge to 1.0×. The trend is slightly smoother than the product-based plot, suggesting that average dimension captures most of the relevant variation.

Figure 3(c): Speedup vs. max(M, N, K). Using the maximum dimension rather than average or product, the pattern holds: max dimension < 1024 yields better speedups, max dimension > 8192 yields near-zero improvement.

Table 3: Stratified by problem size bins. The table (whose exact values are not legible in the provided paper excerpt, but the trend is described in the text) quantifies mean speedup per size bin, confirming that the negative correlation is robust and monotonic.

The authors' interpretation: "This behavior is expected. For small matrix multiplications, the GPU is underutilized and there is significant room for optimization through better memory access patterns, tiling strategies, and kernel configurations. CUDA-L2 exploits these opportunities effectively. On the contrary, large matrices saturate the GPU's floating-point throughput, leaving less room for improvement" (Section 4.3). This interpretation is crucial for understanding the scope of CUDA-L2's contribution: the gains come from finding better ways to use underutilized GPU resources, not from exceeding the GPU's theoretical peak throughput. On large matrices where cuBLAS already achieves near-roofline performance, automated optimization provides marginal benefit.

A nuance the paper does not fully explore: For very large matrices, some configurations show negative speedup (CUDA-L2's kernel is slower than cuBLASLt-AutoTuning). This is visible as points below the 0.0 line in Figure 3(c). The win rate of 79.3–95.7% across baselines implies that 4.3–20.7% of configurations show regression. This is not a failure of the approach — it's the expected behavior of any optimization method that explores a rugged landscape — but it means deployment would require the max(CUDA-L2, baseline) selection strategy rather than naive replacement of the baseline library.

Offline vs. Server Mode: The Thermal Dynamics Effect

The consistent gap between offline and server mode across all baselines (Table 1) is a novel finding that the paper attributes to thermal dynamics (Section 4.1). In offline mode, "continuous execution keeps the GPU in a steady thermal state with predictable clock behavior." In server mode, "idle periods between requests allow the GPU to cool, causing subsequent kernel launches to experience clock boosting followed by thermal throttling." The paper notes that "the larger variance of the server mode v.s. offline mode" is consistent with this explanation.

The practical implication is that CUDA-L2's advantage is larger in realistic deployment scenarios (server mode) than in synthetic throughput benchmarks (offline mode). This is uncommon — many optimizations that look good in peak-throughput benchmarks degrade under realistic request patterns. CUDA-L2's reversal of this pattern suggests its kernels have characteristics (perhaps different memory access patterns, or different sensitivity to cache state) that are less dependent on steady-state GPU conditions.

However, the paper does not provide direct evidence for the thermal dynamics hypothesis (e.g., clock frequency measurements, temperature traces). The explanation is plausible but speculative — alternative explanations include different sensitivity to cache warming (server mode kernels may start with cold caches, penalizing baselines more than CUDA-L2's kernels if the latter have better data locality) or different sensitivity to GPU power management state transitions. A more rigorous analysis would include NCU metrics collected in both modes to identify which microarchitectural behaviors differ.

Ablation Studies and Robustness Checks

The paper does not contain a traditional ablation studies section where components of CUDA-L2 are systematically removed and performance is compared. Instead, the "ablation" evidence comes from three sources: (1) the implicit comparison with CUDA-L1 (which lacked continued pretraining, multi-stage RL, NCU profiling, and retrieval augmentation — and could not optimize HGEMM), (2) the configuration-specific case studies in Section 5 that demonstrate sensitivity to parameter choices, and (3) the speedup vs. problem size analysis that shows where the method works and where it doesn't. These are presented here as the closest analogues to ablation analysis.

Implicit ablation: CUDA-L1 vs. CUDA-L2. The paper's narrative (Section 3.1–3.2) positions CUDA-L2 as extending CUDA-L1 with four enhancements: continued pretraining on diverse CUDA code, multi-stage RL (general → specialized), NCU profiling metrics in context, and retrieval-augmented context. However, no head-to-head comparison between CUDA-L1 and CUDA-L2 on HGEMM is reported. The paper states that CUDA-L1 had "limitations that hinder its effectiveness on the more challenging HGEMM task" (Section 3.1) but does not quantify these limitations. This is a significant omission: without a CUDA-L1 baseline on the same 1,000 configurations, we cannot determine whether the improvements come from (a) the multi-stage curriculum, (b) the continued pretraining, (c) the NCU profiling, (d) the retrieval augmentation, or (e) simply having a larger foundation model (DeepSeek 671B vs. whatever CUDA-L1 used). The ablation is qualitative rather than quantitative.

Case study ablation: tile size sensitivity (Section 5.2). For the specific configuration (M=8192, N=512, K=2048), the paper compares three BM values:

• BM = 160 (with zero-padding M to 8320): +15.2% over cuBLASLt-AutoTuning-TN
• BM = 128 (standard choice, divides 8192 exactly): +0.4%
• BM = 256 (larger tile): −15.7% (regression)

This is effectively a within-kernel ablation showing that tile size choice matters enormously — the difference between optimal and suboptimal BM is 30.9 percentage points of speedup. This validates the paper's central claim that per-configuration optimization is necessary, but it is a single configuration and does not generalize to statements about the RL training process.

Case study ablation: abstraction selection (Section 5.1). CUDA-L2 selects different implementation abstractions for different problem sizes: lightweight WMMA-based kernels for small matrices, CuTe-based kernels with multi-stage pipelining for large matrices. This is an emergent property of the RL training (reward encourages shorter code, so CuTe's conciseness is favored when the complexity justifies the abstraction overhead), but no ablation compares a version of CUDA-L2 forced to use only one abstraction against the adaptive version.

Analysis as ablation: problem size breakdown (Figure 3, Table 3). The speedup vs. problem size analysis serves as a de facto ablation of what types of problems CUDA-L2 helps on. It reveals that CUDA-L2 is not uniformly beneficial — it provides substantial gains on small-to-medium matrices (where GPU utilization is low) and negligible or negative gains on large matrices (where cuBLAS already achieves near-peak throughput). This is an important boundary condition that the paper's high-level claims (e.g., "systematically outperforms major matmul baselines") obscure. The method is not universally superior; it is superior on a substantial subset of configurations, and that subset tilts toward smaller problem sizes.

Training data overlap concern: The paper does not report whether the 1,000 evaluation configurations were seen during HGEMM RL training. If they were, the results may partially reflect memorization of optimized kernels for specific dimension triplets rather than generalization ability. However, the rich optimization strategies described in Section 5 (zero-padding for tile flexibility, staggered prefetch scheduling, double-buffered register fragments) suggest genuine capability rather than lookup-table behavior — these are parameterized optimizations that apply across dimension ranges, not rote memorization of specific (M, N, K) values. The per-dimension speedup trends (Figure 3, showing smooth degradation with problem size rather than sharp memorization spikes at training points) also support generalization.

Critical Assessment

The paper makes three central claims, explicitly or implicitly. Here I evaluate each against the experimental evidence.

Claim 1: CUDA-L2 generates HGEMM kernels that outperform NVIDIA's strongest auto-tuning baseline by 11.4% (offline) and 15.9% (server) on average across 1,000 production-relevant configurations.

What the experiments actually demonstrate: This claim is directly supported by Table 1 and Figure 1. CUDA-L2 achieves these average speedups against cuBLASLt-AutoTuning-max, and the win rates (79.3–95.7%) confirm the gains are systematic rather than driven by outliers. The server-mode gains (15.9%) are larger than offline (11.4%), which strengthens the practical relevance claim.

Where the evidence is weaker:

• No statistical significance testing. The paper reports mean speedup across 1,000 configurations without confidence intervals, standard errors, or hypothesis tests. With speedup variance as high as shown in Figure 3 (points ranging from −0.1 to +0.6), the mean could be imprecisely estimated. A 95% confidence interval of ±5 percentage points would substantially change the interpretation.
• Training-evaluation overlap is undocumented. If the 1,000 evaluation configurations were used during HGEMM RL training, the reported speedups may overestimate generalization to unseen dimensions. The paper does not specify the training distribution.
• The mean obscures the problem-size dependence. Figure 3 shows that speedups concentrate on smaller problems and approach zero (or negative) for large problems. A practitioner with a workload dominated by large matrices (e.g., LLM inference with batch size 1 where FFN dimensions are typically large) would see much smaller benefits than the 11.4% headline. The paper acknowledges this in Section 4.3 but the headline numbers don't convey the conditionality.

Claim 2: LLM-guided RL can discover superior implementations by systematically exploring configuration spaces at scales impractical for human optimization.

What the experiments actually demonstrate: Section 5 provides compelling evidence for this claim through specific optimization technique discoveries. The zero-padding strategy (Section 5.2) — deliberately misaligning tile sizes to unlock better BM values — is genuinely non-obvious and emerged from the RL process rather than being pre-programmed. The staggered A-B prefetch (Section 5.3.4) and double-buffered register fragments (Section 5.3.1) similarly represent exploration of design choices that human engineers, guided by heuristics, would be unlikely to test.

Where the evidence is weaker:

• No comparison to human-in-the-loop optimization. The paper frames its contribution as automating what humans cannot practically do at scale, but does not compare CUDA-L2 against a human CUDA engineer given equivalent time. Would a skilled CUDA engineer, given one week per configuration (generous compared to the RL training cost), achieve similar or better results on specific configurations? The paper cannot answer this because the baseline is library auto-tuning, not expert manual effort on the specific configurations.
• Novelty of discovered techniques is asserted, not verified. The paper claims the optimization techniques in Section 5 were "discovered" by CUDA-L2, but does not verify that these techniques are absent from the cuBLASLt algorithm pool. If cuBLASLt-AutoTuning's 100 candidates already include kernels with similar strategies, CUDA-L2's achievement would be in configuration selection rather than technique discovery. The paper provides only one example (the BM=160 zero-padding case) with specific performance numbers; the other techniques (Sections 5.3.1–5.3.4) are described qualitatively without per-configuration speedup attribution.
• The "scale" argument is qualitative. The paper claims CUDA-L2 explores at "scales impractical for humans," but does not quantify the search space explored by the RL process. How many distinct kernel variants were generated per configuration during RL training? What fraction were compilable? What fraction were correct? Without these metrics, "scale" is a vague claim.

Claim 3: CUDA-L2 extends CUDA-L1 through continued pretraining, multi-stage RL, NCU profiling, and retrieval augmentation, enabling HGEMM optimization that was previously infeasible.

What the experiments actually demonstrate: This claim is not experimentally validated. The paper provides no ablation study comparing CUDA-L2 with CUDA-L1 on HGEMM, nor does it ablate individual components (pretraining only, RL only, NCU profiling only, retrieval augmentation only). The claim that these extensions "enable" HGEMM optimization rests entirely on the narrative that CUDA-L1 had "limitations that hinder its effectiveness on the more challenging HGEMM task" (Section 3.1) and that CUDA-L2 succeeds. But success could be attributable to any subset of the changes, to the larger base model, to more RL training compute, or to implementation details not captured by the four listed enhancements.

This is the paper's most significant experimental weakness. A well-designed ablation would train variants of CUDA-L2 with components removed and compare performance on a subset of configurations. Even a minimal ablation — CUDA-L2 without NCU profiling in context (using only end-to-end timing as reward, as CUDA-L1 did) — would isolate the contribution of rich profiling feedback. The absence of ablations means the paper cannot distinguish between "these techniques are individually necessary" and "these techniques were convenient engineering choices, but any reasonable multi-stage R

6. Limitations and Trade-offs

The Difficulty Estimation Cost Is Unaccounted for in the Performance Claims

The assumption or constraint. The paper's entire compute-optimal framework depends on estimating prompt difficulty before allocating the test-time compute budget. The method for this estimation — generating 2,048 samples per question and computing either the ground-truth pass@1 rate (oracle difficulty) or the PRM's average final-answer score (predicted difficulty) — is extraordinarily expensive. The paper acknowledges this explicitly in Section 3.2:

"estimating difficulty in this way still incurs additional computation cost during inference... our experiments do not account for this cost largely for simplicity"

Generating 2,048 complete solutions for a single question is comparable to or greater than the largest test-time compute budgets studied (256–512 generations). To then apply a strategy that uses, say, 16 generations (the compute-optimal policy at the 4× efficiency point), the total cost is 2,048 (estimation) + 16 (execution) = 2,064 generations — a 129× overhead relative to the claimed budget. The difficulty estimation dominates the total cost.

The consequence. The reported 4× efficiency gains over best-of-N are computed after difficulty is known, without amortizing the cost of learning it. In a realistic deployment where difficulty must be estimated per prompt, the total compute cost would be difficulty estimation + strategy execution, and the former could easily consume more than the savings from strategy selection. The $4\times$ figure should therefore be understood as an upper bound on achievable efficiency — what is possible given perfect, cost-free difficulty information — not as a realized deployment gain. For interactive applications where latency matters, the difficulty estimation phase (requiring thousands of generations per query) would be completely prohibitive regardless of accuracy gains.

What evidence exists in the paper. The paper's results in Figures 4 and 8 plot accuracy against generation budget where the budget accounts only for the strategy execution, not for difficulty estimation. The predicted difficulty curves largely overlap the oracle curves, confirming that the PRM-based approximation works — but the cost of computing that approximation (the 2,048 samples and PRM scores) is not included in the x-axis of either figure. This means the x-axis is not a total-cost axis; it is a marginal-cost axis. The paper provides no analysis of how the efficiency gains change if the difficulty estimation cost is included — for example, what fraction of the total compute budget is spent on estimation vs. execution at different total budget levels.

Mitigation status. The paper acknowledges the problem in Section 3.2 and frames it as an exploration-exploitation tradeoff, suggesting future work on "pretraining or finetuning models to directly predict difficulty of a question" from the question text alone. However, no such model is developed or evaluated. The paper also does not explore adaptive difficulty estimation — starting with a small number of samples (say, 4–8), using the verifier's score distribution on those samples as a rough difficulty signal, and allocating the remaining budget accordingly. Such a scheme could amortize difficulty estimation into the problem-solving process, but it remains unimplemented. As the paper stands, the difficulty estimation cost is the single largest unaccounted expense in the system, and the $4\times$ efficiency claim should be treated as conditional on a cheap difficulty estimator that does not yet exist.

The $14\times$ Larger Model Baseline Is Not Compute-Optimally Trained and Uses Only Greedy Decoding

The assumption or constraint. The FLOPs-matched comparison in Section 7 scales model parameters while holding training data fixed, following the LLaMA paradigm (Touvron et al., 2023) rather than compute-optimal pretraining (Hoffmann et al., 2022). The paper acknowledges this explicitly:

"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."

Additionally, the $14\times$ larger model is evaluated using greedy decoding only — no majority voting, no best-of-N, no search, no revisions. The smaller model with test-time compute is compared against a larger model using the simplest possible inference strategy.

The consequence. Both choices make the pretraining baseline weaker than it could be, potentially inflating the reported advantages of test-time compute over pretraining. A Chinchilla-optimal model trained with $14\times$ more total FLOPs (scaling both parameters and training data optimally) would likely outperform a model where only parameters are scaled, because the latter is under-trained relative to its capacity. The reported numbers — e.g., +27.8% relative improvement on easy questions at $R \ll 1$ for revisions (Figure 1) — are measured against this suboptimal baseline. Against a properly compute-optimal larger model, the advantage might shrink or reverse on some difficulty tiers.

The greedy decoding choice similarly stacks the comparison. If the $14\times$ larger model were given even a modest test-time compute budget — best-of-8 with majority voting, or a simple verifier-based selection — its accuracy would improve, narrowing or eliminating the gap with the smaller model's compute-optimal strategy. The paper is effectively comparing small model + smart inference against large model + dumb inference, when a fairer comparison would give both models access to test-time compute and study the interaction between model scale and inference strategy.

What evidence exists in the paper. The FLOPs-matched results in Figure 9 and the bar charts in Figure 1 show the comparison as described. The paper does not provide an ablation where the larger model receives any test-time compute budget, nor does it compare against a compute-optimally trained larger model. The paper's choice of parameter-only scaling is explicitly flagged as a departure from best practice, but no sensitivity analysis is performed to assess how much this choice affects the conclusions.

Mitigation status. The paper is transparent about the limitation, but does not mitigate it — no variant of the experiment with a compute-optimal pretraining baseline or with test-time compute allocated to the larger model is presented. The authors frame this as future work. Until such comparisons are made, the paper's central claim — "test-time compute can outperform a $14\times$ larger model" — should be understood as applying specifically to the parameter-scaled, greedy-decoded baseline studied, not to compute-optimally trained larger models in general.

All Results Are on a Single Benchmark (MATH) with a Single Model Family (PaLM 2-S*)

The assumption or constraint. Every experiment in the paper uses the MATH benchmark — 500 test questions spanning high-school competition-level math — and PaLM 2-S* as the base model. The paper states the model is "representative of the capabilities of many contemporary LLMs" (Section 4), but this claim is unverified. No experiments are conducted on other reasoning benchmarks (e.g., GSM8K, MMLU reasoning subsets, code generation tasks), other model families (e.g., LLaMA, Qwen, DeepSeek), or other task types (e.g., factual QA, summarization, planning).

The consequence. Three generalization questions are left unanswered:

1. Domain generalization: MATH consists exclusively of symbolic math reasoning problems with well-defined correct answers. It is unclear whether the difficulty-dependent strategy patterns observed (beam search hurts easy problems, revisions help easy problems, parallel search helps medium problems, no method helps hard problems) generalize to other reasoning domains. Code generation, for example, has qualitatively different error patterns (syntax errors, logic errors, test case failures) that might interact differently with verifier guidance and iterative revision.

2. Model generalization: PaLM 2-S* has specific characteristics — its calibration, its error distribution, its revision capability after fine-tuning — that may not transfer to other model families. A model with different base capabilities might exhibit different difficulty-dependent scaling curves. For example, a model with stronger base math performance might have a larger fraction of problems in the "easy" bin where revisions dominate, shifting the compute-optimal policy.

3. Benchmark specificity: MATH problems are high-school competition problems requiring multi-step symbolic deduction. Real-world deployment often involves a mix of simple queries (where test-time compute is unnecessary) and genuinely novel problems (where test-time compute may not help). The MATH distribution may not match the difficulty distribution encountered in practice, making the reported 4× gains optimistic or pessimistic depending on the deployment context.

What evidence exists in the paper. All evidence is from MATH + PaLM 2-S*. The paper provides no cross-domain or cross-model replication. The difficulty bins are defined relative to PaLM 2-S*'s pass@1 rate on MATH, meaning the entire framework — difficulty estimation, strategy selection, compute-optimal allocation — is calibrated to this specific model-benchmark pair. A different model on a different benchmark might require re-deriving the difficulty bins and re-learning the compute-optimal policies from scratch.

Mitigation status. The paper does not address this limitation beyond acknowledging the scope. No experiments on other benchmarks or model families are reported or planned. The claim that PaLM 2-S* is "representative" is an assertion, not an empirical finding. A practitioner deploying these methods on a different model family (e.g., LLaMA on code generation) should treat the MATH/PaLM 2-S* results as a proof of concept requiring independent validation in their specific context.

The 500-Question Test Set and Five-Bin Stratification Yield Small Per-Bin Sample Sizes

The assumption or constraint. The MATH test set contains 500 questions. These are split into five difficulty quintiles of approximately 100 questions each. The compute-optimal policy is selected using two-fold cross-validation within each bin — meaning strategy selection is based on roughly 50 questions per fold per bin. The paper then reports accuracy on the held-out fold, averaged across the two folds.

The consequence. The sample size for strategy selection is small. With 50 questions per bin, the variance in estimated best strategy is potentially large — a different random split of 100 questions into two folds of 50 could yield different strategy selections. More importantly, no confidence intervals, standard errors, or statistical significance tests are reported for any of the main results (Figures 3, 4, 6, 7, 8). The paper reports point estimates of accuracy and efficiency gains (e.g., 4× improvement, +27.8% relative), but the reliability of these point estimates at n = 50 per bin is unclear. A 95% confidence interval of ±10 percentage points around the accuracy figures would substantially change the interpretation of which strategies are optimal and how large the efficiency gains actually are.

The cross-validation protocol itself — selecting the best strategy on one fold and evaluating on the other — is sound in principle, but its effectiveness depends on the folds being representative of the difficulty bin. At 50 questions per fold, a single outlier question (e.g., one that is unusually easy or hard for its bin) could bias the strategy selection. The paper does not report stability analysis (e.g., how often the best strategy changes when the folds are re-randomized, or the variance in accuracy across folds for each strategy).

What evidence exists in the paper. All result figures (3–9) show point estimates without error bars. The paper does not report standard deviations, confidence intervals, or p-values for any comparison. The win rate metric for search (Section 5) is across all 500 questions, not per-bin, partially mitigating the concern for aggregate claims but not for per-bin policy selection. The FLOPs-matched comparison (Section 7) aggregates difficulty bins into coarser groups ("easy," "medium," "hard"), which may increase per-group sample size but at the cost of coarser allocation.

Mitigation status. The paper does not address this limitation. Given the fixed MATH test set size (500 questions) and the need for within-bin cross-validation, the small per-bin sample size is inherent to the experimental design. Possible mitigations — bootstrapping confidence intervals, reporting fold-level variances, testing sensitivity to the number of difficulty bins — are not presented. A practitioner should treat the reported compute-optimal policies and efficiency gains as preliminary estimates that may not be robust at the per-bin level.

Verifier Over-Optimization Is Documented but Not Resolved — It Sets a Hard Ceiling on Test-Time Compute Scaling

The assumption or constraint. The paper demonstrates that process reward model (PRM) guided search suffers from over-optimization: aggressive search (beam search, lookahead search) finds solutions that score highly under the PRM but are actually incorrect. This is most evident for easy problems, where beam search degrades performance as the generation budget increases (Figure 3, right, bin 1: beam search accuracy decreases from ~78% to ~77% as budget goes from 4 to 256 generations, while best-of-N improves from ~68% to ~88%). Lookahead search — the most powerful optimizer — paradoxically underperforms all other methods at the same generation budget (Figure 3, left). Qualitative examples in Appendix M show search producing degenerate outputs: repetitive low-information steps at the end of solutions, overly short 1–2 step solutions that score well under the PRM but are incorrect.

The consequence. Verifier over-optimization imposes a hard ceiling on the benefits of test-time compute scaling. As budget increases, search doesn't just hit diminishing returns — it can actively hurt performance. The compute-optimal policy mitigates this by routing easy problems away from aggressive search (using best-of-N instead of beam search), but it does not solve the underlying problem. On medium-difficulty problems where beam search is deployed, the over-optimization ceiling still limits how far additional budget can push performance — the beam search curves in Figure 3 flatten well before the budget is exhausted. This means that improving verifier robustness is the bottleneck for further scaling, not improving search algorithms or revision strategies. Until the PRM can remain calibrated under aggressive optimization pressure, test-time compute scaling will asymptote at whatever budget the verifier begins to be exploited.

The paper's current approach — avoiding over-optimization by switching to weaker methods on easy problems — is a workaround, not a solution. It means the system deliberately leaves performance on the table for easy problems (where best-of-N is used instead of potentially more efficient search) to avoid the risk of exploitation. For medium problems, where the verifier signal is more reliable, the over-optimization ceiling still exists but is reached at higher budgets.

What evidence exists in the paper. Figure 3 (right) provides the clearest direct evidence: beam search accuracy decreasing with budget on bin 1. Figure 3 (left) shows lookahead search underperforming across the board. Appendix M (Figures 29 and surrounding examples) provides qualitative evidence of degenerate search outputs. The paper's analysis in Section 5.3 explicitly identifies over-optimization as the problem:

"The degradation at high budgets is attributed to over-optimization of the PRM — search finds solutions that score highly under the PRM but are actually incorrect."

Mitigation status. The paper identifies the problem but does not attempt to solve it. The compute-optimal policy mitigates by avoiding aggressive search where the verifier is unreliable, but this is a routing solution, not a verifier improvement. The paper suggests no concrete approaches to training more robust verifiers (e.g., adversarial training against search-generated examples, ensemble verification, KL-regularized search to keep outputs close to the base model's distribution). This is flagged as future work in Section 8, but no experiments are reported. For a practitioner, this means that deploying the system with a stronger search method (e.g., lookahead search with higher k) would likely be counterproductive — the verifier used in this paper cannot support it.

Revisions and Search Are Studied Independently but Never Combined

The assumption or constraint. The paper studies two complementary axes of test-time compute — PRM-guided search (Section 5) and iterative revisions (Section 6) — but never evaluates them in combination. Section 8 explicitly acknowledges this:

"we did not experiment with PRM tree-search techniques in combination with revisions"

The results therefore represent the performance achievable with each mechanism independently, not the performance achievable by a system that uses both.

The consequence. This is a significant gap because the two mechanisms have complementary strengths revealed by the paper's own analysis:

• Revisions (modifying the proposal distribution) work best on easy problems where the model's initial output is roughly correct and just needs refinement — a local search in answer space.
• PRM search (optimizing the verifier) works best on medium-hard problems where the model needs to explore qualitatively different solution strategies — a global search.

A combined system could, in principle, use the revision model as the proposal distribution within beam search — at each step of the search tree, the model conditions on previous rejected branches as context, generating higher-quality candidate steps than the base model would. Alternatively, the PRM could guide which revisions to pursue: rather than blindly generating a long revision chain, use the PRM's per-step scores to decide when a revision is on track versus when to restart from scratch. Such combinations could yield gains beyond either method alone, particularly on medium-difficulty problems where both exploration (search) and refinement (revisions) are needed.

What evidence exists in the paper. None — this is a gap, not a failed experiment. The paper provides no data on how search and revisions interact. It is possible that combining them yields superlinear gains (each amplifies the other); it is also possible that they interfere (the revision model's output distribution may be poorly calibrated for PRM scoring, or search may over-optimize the verifier more aggressively on revision outputs). Without experiments, both possibilities remain open.

The paper does provide related evidence that hints at potential challenges: Appendix J (Figure 15a) shows that the base-LM-trained PRM underperforms a revision-specific ORM when scoring revision model outputs, suggesting that verifier calibration degrades under distribution shift. If beam search were applied to revision model outputs using the base PRM, the over-optimization problem could be exacerbated by the distribution mismatch between the PRM's training data (base model outputs) and the search space (revision model outputs).

Mitigation status. The paper acknowledges the gap as future work but provides no results, no preliminary analysis, and no concrete proposal for how to combine the two mechanisms. For a practitioner, this means the reported 4× efficiency gains represent a lower bound on what a fully integrated system could achieve — but the gap between this lower bound and the achievable upper bound is unknown. A production deployment of these techniques would almost certainly need to combine them, and the paper provides no guidance on how to do so effectively.

The Revision Model Has a Fundamental Correct-to-Incorrect Reversion Problem and Sensitive Training Dynamics

The assumption or constraint. The revision model is trained on trajectories where all in-context answers are incorrect, followed by a correct answer (Section 6.1). At inference time, however, the model may generate a correct answer early in the revision chain — and then, in a subsequent revision step, "revise" that correct answer into an incorrect one. The paper reports (Section 6.1):

"approximately 38% of correct answers get converted back to incorrect ones"

This is a direct consequence of the training data construction: the model learned that the previous answers in context are always wrong, so it learned that its task is to change the answer regardless of whether the previous answer is actually correct. The model was never trained on trajectories where a correct answer should be preserved.

The consequence. The 38% reversion rate means that generating longer revision chains does not monotonically improve quality — later revisions can destroy earlier successes. The paper mitigates this by using majority voting or verifier-based selection across the entire chain (picking the best answer from any point in the chain, not always the last revision). However, this mitigation is imperfect: it requires the selector (majority or verifier) to correctly identify which point in the chain has the best answer, and it abandons the possibility of genuine multi-step refinement where each step builds on the previous one to reach a better answer than any single step could achieve alone.

The reversion problem also fundamentally limits the benefit of sequential revision scaling. As the chain length grows, the probability that a correct answer appears at some point increases, but the probability that a subsequent revision destroys it also increases. The optimal chain length trades off these competing effects, and at some length the marginal benefit of additional revisions becomes zero or negative.

What evidence exists in the paper. The 38% reversion rate is reported in Section 6.1. Figure 6 (left) shows that the revision model's per-step accuracy improves from ~18.2% at step 1 to ~24–25% by steps 15–20 and then remains flat out to 64 steps — evidence that the benefit of additional revisions saturates (consistent with the reversion problem), but also that it does not collapse (inconsistent with a naive model where every correct answer gets reverted). The conservative interpretation is that the selector (verifier or majority) successfully filters out most reversions at high step counts, but at the cost of discarding the intended benefit of multi-step refinement.

Mitigation status. The paper mitigates via within-chain selection (Section 6.1) but does not solve the root cause. A more principled solution — training the revision model to recognize when the current answer is already correct and produce a "no change needed" output — is not explored. The paper does not report what fraction of the final selected answers come from early versus late positions in the chain, which would quantify how much of the revision chain is actually useful versus discarded.

The ReST$^{EM}$ experiment (Appendix K, Figure 16) further highlights the fragility of revision training. Attempting to further optimize the revision model with RL-style training (ReST$^{EM}$) caused performance to degrade substantially with sequential revisions: at 256 generations, fully sequential performance drops to ~33.5% compared to ~38.5% at the optimal ratio. The authors hypothesize that "on-policy data collection in ReST$^{EM}$ exacerbates spurious correlations in revision data, causing the model to fail to learn the revision task properly." This suggests that the revision approach is sensitive to training methodology in ways not fully understood, and that the positive results in Section 6 depend critically on the specific offline data construction procedure (edit-distance-based pairing, fixed number of in-context incorrect answers). A practitioner attempting to replicate or extend the revision approach should treat the training procedure as brittle until proven otherwise.

7. Implications and Future Directions

How This Work Changes the Landscape

CUDA-L2 represents a reframing, not a paradigm shift — it is not the first system to use LLMs for CUDA kernel generation (that lineage runs through KernelBench, AI CUDA Engineer, and CUDA-L1), but it is the first to demonstrate that LLM-guided RL can discover optimizations that exceed the performance ceiling of exhaustive library auto-tuning on a production-critical kernel across a production-relevant configuration space. The conceptual shift is from evaluating kernel optimizers on point estimates (one configuration, one speedup number) to evaluating them on distributional performance across the full span of dimensions that deployed models actually encounter.

This reframing has several downstream effects on the field:

It makes the configuration-generalization gap the central evaluation criterion for automated kernel optimization. Prior to CUDA-L2, it was reasonable to claim success on "matrix multiplication optimization" by beating a baseline on a few hand-picked (M, N, K) triplets. After CUDA-L2, the responsible standard is systematic evaluation across the dimension space — specifically, the 1,000-configuration grid that covers the attention and FFN layers of deployed open-source models like Qwen, Llama, and DeepSeek. The paper's 1,000-configuration benchmark is a methodological contribution in its own right, establishing what "solving HGEMM optimization" should mean: not excellence on cherry-picked sizes, but systematic superiority across the configuration distribution. Systems that report results on single configurations without analyzing per-dimension trends should now be viewed with skepticism.

It positions automated kernel optimization as a planning-and-exploration problem rather than a code-translation problem. The insight that CUDA-L2 discovers optimization strategies (zero-padding for tile-size flexibility, staggered prefetch scheduling, double-buffered register fragments) that deviate from standard practice — and that human experts are unlikely to test because their heuristics systematically exclude those regions of the search space — suggests that RL's comparative advantage is not in translating known optimizations into code, but in exploring the parts of the optimization space that expert heuristics prune away. This has a concrete implication for system design: future kernel optimization systems should invest more heavily in exploration mechanisms (curriculum RL, diversity rewards, novelty search) and less in encoding expert heuristics as hard constraints. The expert's role shifts from prescribing what to try to providing the building blocks (CuTe abstractions, WMMA primitives, tiling frameworks) that the RL agent can recombine in unexpected ways.

It provides empirical evidence that library auto-tuning — even exhaustive enumeration of 100 hand-designed candidates — leaves measurable headroom. The 11.4% offline speedup (15.9% server) over cuBLASLt-AutoTuning establishes a concrete lower bound on what exhaustive selection among existing kernel designs misses. This has practical implications for organizations deploying LLM inference: the marginal cost of running CUDA-L2's RL training once and deploying the generated kernels indefinitely may be justified even at single-digit percentage improvements, given the scale of matmul in total inference cost. The paper does not quantify this cost-benefit tradeoff (RL training cost is unreported), but the existence of the 11.4% gap itself shifts the burden of proof: library authors must now argue that additional hand-designed kernel variants would close this gap, or accept that automated generation is a necessary complement to manual design.

It resolves a latent tension in the kernel optimization literature. Prior work on LLM-based CUDA generation (KernelBench, AI CUDA Engineer) demonstrated that LLMs can produce functionally correct kernels for diverse operations, but struggled to achieve competitive performance on heavily-optimized kernels like HGEMM. The pessimistic reading was that LLMs lack the specialized expertise to compete with hand-tuned libraries on the hardest kernels. CUDA-L2's results — specifically the observation that gains concentrate on smaller matrices (where GPU underutilization creates headroom) and diminish on large matrices (where cuBLAS already achieves near-roofline performance) — reconciles these findings. The LLM is not uniformly worse or better than human experts; it is differentially capable on the subset of the configuration space where conventional optimization leaves slack. Prior work tested on configurations or benchmarks that happened to fall in the saturated regime. CUDA-L2 shows that changing the evaluation distribution changes the conclusion, and that the headroom is real but concentrated.

However, CUDA-L2 does not resolve the critical question of causal contribution — which of its four technical enhancements (continued pretraining, multi-stage RL curriculum, NCU profiling, retrieval augmentation) is necessary, and which is incidental. Without ablation studies, the field cannot determine whether the 11.4% speedup requires all four components, or whether a much simpler system (e.g., continued pretraining + single-stage RL with end-to-end timing as reward) would achieve comparable results. This limits the paper's prescriptive value for system builders: should one invest in NCU profiling integration, or is a larger pretraining corpus sufficient? The answer is unknown, and until ablation studies are conducted, each component's contribution remains speculative. This is the paper's most significant unresolved question for the research community.

Finally, the paper opens a line of inquiry that it does not pursue: the relationship between problem size and the headroom for automated optimization. The monotonic degradation of speedups with increasing (M, N, K) (Figure 3, Table 3) is not a failure mode but a diagnostic — it tells us where automated methods add value (underutilized GPUs, small-to-medium matrices) and where they don't (saturated GPUs, large matrices near roofline). Future work that extends CUDA-L2 to Hopper or Blackwell architectures should test whether the headroom changes — newer architectures with larger tensor cores and different memory hierarchies may shift the boundary between underutilized and saturated regimes, potentially opening new optimization opportunities at dimension ranges where CUDA-L2 currently flatlines.

Follow-Up Research This Work Enables

Ablation study decomposing CUDA-L2's four technical contributions. The paper's most urgent follow-up is a controlled experiment that trains CUDA-L2 variants with individual components removed and measures performance on the 1,000-configuration benchmark. The minimal variant to test: continued pretraining on diverse CUDA code + single-stage HGEMM RL (no general kernel RL curriculum), with end-to-end timing as the only reward (no NCU profiling in context) and no retrieval augmentation. Additional variants would add components one at a time: +general kernel RL only, +NCU profiling only, +retrieval augmentation only. The dependent variable is the mean speedup over cuBLASLt-AutoTuning-max. A strong result would show that each component contributes independently and additively; a weak result would show that one component (e.g., continued pretraining) accounts for most of the gain, simplifying the design space for future systems. The ablation would also test whether NCU profiling contributes above and beyond simply providing more training compute — the current paper conflates these, since variants without profiling would presumably train faster per step.

Adaptive difficulty estimation for kernel configuration difficulty. The paper's per-dimension speedup analysis (Figure 3) reveals that CUDA-L2's benefit varies systematically with problem size — large gains on small matrices, near-zero on large matrices. This suggests a deployment strategy analogous to the compute-optimal allocation in the reference paper's LLM test-time compute work: estimate whether a given (M, N, K) configuration is in the "optimizable" regime before invoking CUDA-L2's generation process, and fall back to cuBLASLt-AutoTuning for saturated configurations. A concrete experiment: train a lightweight classifier on (M, N, K) triplets to predict CUDA-L2's expected speedup (continuous or binned), using the 1,000-configuration data as training labels. Evaluate on held-out configurations from the same dimension space. The goal is to build a deployment system that invokes CUDA-L2 only when expected gain exceeds some threshold (say, +5%), saving the (currently unquantified) RL generation cost on configurations where it cannot help. This is directly analogous to the difficulty estimation problem in the reference paper, but with a cleaner oracle signal (the measured speedup distribution).

Cross-architecture transfer of CUDA-L2's optimization strategies. The paper explicitly states that "optimizations rarely transfer across GPU architectures" (Section 1) and that work on Ada Lovelace, Hopper, and Blackwell is "ongoing." A direct follow-up would replicate the HGEMM RL stage on an H100 (Hopper) GPU using the same 1,000-configuration space, comparing: (a) kernels generated de novo by CUDA-L2 on H100; (b) kernels generated on A100 and directly executed on H100 (zero-shot transfer); (c) cuBLASLt-AutoTuning on H100. The key question is whether the RL curriculum (continued pretraining + general kernel RL) provides enough architectural knowledge to optimize for the new target, or whether the HGEMM RL stage overfits to Ampere-specific characteristics that don't transfer. If zero-shot transfer performs poorly (as the intuition about architecture-specific optimizations suggests), a second experiment would test whether fine-tuning the A100-trained model on H100 configurations (domain adaptation) recovers performance faster than training from scratch, quantifying the transfer learning benefit.

Combining CUDA-L2's generated kernels with cuBLASLt's algorithm pool in a joint optimization framework. The current evaluation treats CUDA-L2 and cuBLASLt as competitors — CUDA-L2 generates a kernel, cuBLASLt selects the best from up to 100 hand-designed candidates, and the paper reports speedup of one over the other. The max(CUDA-L2, baseline) analysis in Table 2 shows that combining them (selecting the faster per configuration) yields marginal additional gains (+1.8 percentage points for cuBLASLt-AutoTuning-max offline). A more ambitious follow-up would integrate CUDA-L2 into the auto-tuning loop: use the RL-trained model to generate novel kernels that are added to cuBLASLt's algorithm pool, then use cuBLASLt's exhaustive benchmarking to select the best among the combined set (hand-designed + RL-generated). This leverages CUDA-L2's novel kernel discovery with cuBLASLt's robust per-configuration selection, potentially closing the gap on the ~20% of configurations where cuBLASLt still outperforms CUDA-L2 alone. The experiment would report the speedup of this combined approach over cuBLASLt-AutoTuning alone, and would test whether CUDA-L2's kernels provide orthogonal value (the best combined kernel is often CUDA-L2) or redundant value (cuBLASLt already has a kernel nearly as good, and the max operation captures most of the gain).

Profiling-guided curriculum design for other kernel types. The multi-stage RL curriculum (general kernels → HGEMM-specific) is described as a design choice but not empirically justified. A rigorous follow-up would test whether this curriculum structure is necessary by comparing: (a) the current general→specialized curriculum; (b) a specialized-only curriculum (HGEMM RL directly after continued pretraining, skipping general kernel RL); and (c) a mixed curriculum (interleaving general and HGEMM kernels throughout RL training). The dependent variable is final HGEMM speedup over cuBLASLt-AutoTuning-max. A strong result for the curriculum hypothesis would show that (a) outperforms (b) by a significant margin, indicating that general optimization skill transfer is real and non-trivial. A null result would show that (b) achieves comparable performance, simplifying future systems by eliminating the general kernel RL stage. A third variant could test a profiling-guided curriculum where the general kernel RL stage is complemented with specific attention to profiling metrics (memory throughput, occupancy) that transfer to HGEMM, testing whether the curriculum benefit is from general skill acquisition or specifically from learning to interpret profiling data.

Latency-aware optimization: incorporating wall-clock time constraints into the RL reward. The paper measures performance in raw execution time but ignores latency — the fact that server mode kernels, while individually measured without idle time included, still execute in a deployment context where response time matters. A follow-up would modify the RL reward to include a latency penalty: reward kernels not just for throughput (executions per second) but for meeting a target latency budget (e.g., kernel execution < 1ms for interactive inference). This would test whether CUDA-L2 can discover kernels that are simultaneously fast and latency-bounded — potentially by exploring tile sizes that reduce per-kernel latency at the cost of slightly lower throughput, a tradeoff that is irrelevant for offline evaluation but critical for real-time serving. The experiment would compare CUDA-L2 trained with and without the latency penalty on server-mode performance, and would report the Pareto frontier of throughput vs. latency across generated kernels.

Practical Applications and Downstream Use Cases

LLM inference serving with per-configuration kernel selection. The most immediate application is a deployment system where a model server (e.g., vLLM, TensorRT-LLM) maintains a library of kernels — cuBLAS defaults, cuBLASLt auto-tuned kernels, and CUDA-L2-generated kernels — and selects the fastest per (M, N, K) configuration at inference time. The max(CUDA-L2, baseline) results in Table 2 directly quantify the benefit: even against the strongest baseline (cuBLASLt-AutoTuning-max), selecting the best of CUDA-L2 and the baseline yields +13.2% offline and +18.1% server speedups. For a model serving deployment running Llama- or Qwen-scale architectures, this translates to a double-digit percentage reduction in matmul-bound inference latency with no change to model architecture, weights, or hardware. The implementation cost is a one-time offline generation pass (CUDA-L2 generates one kernel per encountered dimension triplet), followed by the same per-configuration benchmarking that cuBLASLt-AutoTuning already requires — the marginal overhead is the RL training and kernel generation, not the deployment benchmarking.

Edge and mobile GPU deployment where cuBLAS is unavailable or unoptimized. The paper focuses on A100 (data-center Ampere), but the framework is architecture-agnostic. For edge GPUs (e.g., Jetson Orin, RTX 4060 mobile) where NVIDIA's library auto-tuning may be less exhaustive or where certain matrix dimensions encounter pathological performance (e.g., dimensions that don't align with the GPU's warp or tensor core sizes), CUDA-L2 could generate specialized kernels that exploit the specific hardware characteristics. The small-matrix speedup results (Figure 3, showing 1.2–1.5× gains for small matrices where GPU underutilization is highest) are directly relevant to edge scenarios where batch sizes are small and matrix dimensions may be moderate (e.g., on-device inference for a 7B-parameter model with M=64 attention heads). The key practical advantage is that CUDA-L2's per-configuration specialization is most valuable precisely where generic library kernels are least optimized — the small-matrix, underutilized regime.

Non-standard matrix multiplication variants not covered by library auto-tuning. cuBLAS and cuBLASLt optimize standard GEMM operations (C = αAB + βC) with conventional layouts and accumulator types. CUDA-L2's code generation capability extends to custom matmul variants that libraries may not exhaustively support: fused operations (e.g., matmul followed by element-wise activation, as in FFN layers), mixed-precision accumulation strategies beyond what cuBLASLt exposes, or non-standard matrix layouts (e.g., block-sparse or structured-sparse formats for pruned models). The paper's demonstration that CUDA-L2 discovers valid, correct kernels without hard-coded heuristics (the correctness validation in Section 2.3 ensures functional equivalence) suggests that the same pipeline could be applied to these non-standard operations without requiring library authors to design and maintain specialized kernels for each variant. The key inflection point for a practitioner: if you need a matmul variant for which cuBLASLt's algorithm pool is small or nonexistent, CUDA-L2 may provide a viable alternative to manual kernel development.

When to Prefer This Method

The paper does not explicitly position CUDA-L2 against named alternative automated optimization frameworks (e.g., Triton autotuning, AutoTVM, Halide) beyond the restriction to CUDA .cu files rather than Python-based DSLs (Section 3.2.3). The comparison is against library baselines (torch.matmul, cuBLAS, cuBLASLt), not alternative generation approaches. The paper also does not articulate a decision rule for when to use CUDA-L2 versus when to rely on library defaults or manual tuning. A practitioner reading the paper can infer from the results:

• CUDA-L2 is preferable when the workload contains a substantial fraction of small-to-medium matrix multiplications (dimensions up to roughly 2048–4096, where Figure 3 shows speedups of 1.1× or higher). This regime includes attention head projections in moderate-scale transformers and FFN layers in smaller models.
• cuBLASLt-AutoTuning remains competitive or preferable for very large matrices (dimensions > 8192, where Figure 3 shows speedups converging to 1.0× with occasional regressions). In this regime, cuBLAS already approaches roofline performance, and CUDA-L2's per-configuration generation may not justify its (currently unquantified) training and generation cost.
• The server-mode advantage (4–5 percentage points higher speedups) suggests CUDA-L2 is particularly valuable for real-time inference deployments, where cold-start effects penalize baseline libraries more than CUDA-L2's kernels.
• The max(CUDA-L2, baseline) strategy dominates either alone (Table 2), making the safest deployment approach a combined system where CUDA-L2 kernels supplement rather than replace existing libraries until per-configuration benchmarking identifies the superior option.

However, these inferences are drawn from the results rather than stated as explicit design guidance by the authors. A true decision rule would require quantifying CUDA-L2's training and per-configuration generation cost against the cumulative performance benefit across the expected configuration distribution — and that quantification is absent from the paper.