CUDA Agent: Large-Scale Agentic RL for High-Performance CUDA Kernel Generation

🎯 Pitch

A language model trained with reinforcement learning in an interactive CUDA development environment generates kernels that dramatically beat torch.compile, where even the strongest proprietary models fail. By giving the model profiling tools and execution feedback, it learns to discover algebraic simplifications and novel tiling strategies that static compilers miss, more than doubling the geometric mean speedup over compile on the hardest operator fusion tasks.

1. Executive Summary

This paper introduces CUDA Agent, a large-scale agentic reinforcement learning system that trains an LLM to generate high-performance CUDA kernels by jointly scaling data synthesis, a skill-augmented development environment with automated verification and profiling, and stability-oriented RL techniques. Using Seed1.6 as the base model and training on 6,000 synthesized operator tasks, CUDA Agent operates within an interactive agent loop that interleaves coding, compilation feedback, and performance profiling across up to 200 turns, optimized via PPO with a multi-stage warm-up strategy — rejection fine-tuning for the actor and value pretraining for the critic — to prevent training collapse from the severe distribution mismatch between pretraining data and CUDA code. On KernelBench, CUDA Agent achieves 100%, 100%, and 92% faster rate over torch.compile on Level-1, Level-2, and Level-3 splits respectively, substantially outperforming the strongest proprietary models Claude Opus 4.5 and Gemini 3 Pro by approximately 40% on the hardest Level-3 setting, establishing that learned optimization policies can consistently surpass static compiler heuristics particularly on complex operator fusion tasks while leaving fundamentally out-of-capability problems — where no method makes meaningful progress — as an unsolved boundary.

2. Context and Motivation

The Core Problem: LLMs Cannot Write Competitive GPU Kernels

The fundamental problem this paper addresses is deceptively simple to state but extraordinarily difficult to solve: large language models, despite remarkable proficiency in general-purpose programming, cannot produce GPU kernels that outperform standard compiler optimizations. Specifically, on the KernelBench benchmark [17], even the strongest proprietary models—Claude Opus 4.5 and Gemini 3 Pro—produce kernels where only 66–69% actually run faster than torch.compile, a default PyTorch just-in-time compiler that applies static, rule-based optimizations. Even when those models do produce faster kernels, the geometric mean speedup is modest: roughly 1.46–1.42× over torch.compile across all difficulty levels (Table 1).

This gap is not marginal. torch.compile is a general-purpose compiler backend—it applies predefined patterns like kernel fusion, memory layout optimization, and operator substitution without any understanding of the specific computation's mathematical structure. It has no awareness of algebraic simplifications (e.g., recognizing that a diagonal matrix multiplication is equivalent to row-wise scaling), no ability to discover novel tiling strategies tailored to a particular operator composition, and no iterative refinement loop driven by profiling feedback. In principle, a human CUDA expert armed with profiling tools can dramatically outperform torch.compile by exploiting these domain-specific insights. Yet LLMs—which have ingested vast quantities of human-written code and documentation—fail to approach this level of performance. The paper's central question is therefore: can we train an LLM to close this gap, transforming it from a naive code generator into a genuine systems optimizer for GPU computing?

Why This Problem Matters: Three Levels of Significance

Practical impact on deep learning infrastructure. GPU kernels are the computational backbone of modern AI. Every matrix multiplication, convolution, attention mechanism, and normalization operation in a neural network ultimately executes as one or more CUDA kernels on NVIDIA hardware. The performance of these kernels directly determines training throughput, inference latency, and energy consumption at datacenter scale. Even a modest improvement—say, 20% faster execution for a commonly used operator—can translate to millions of dollars in saved compute costs and reduced carbon footprint when deployed across thousands of GPUs running continuously.

The status quo for kernel optimization relies on a scarce resource: expert human engineers who deeply understand both the mathematical structure of neural network operations and the microarchitectural details of specific GPU generations (shared memory banks, warp scheduling, tensor core throughput, memory bandwidth hierarchies). These experts are expensive, their throughput is low, and their optimized kernels often need to be rewritten for each new GPU architecture. If an LLM-based system could automate this expertise—producing kernels that are consistently faster than compiler baselines across diverse operator types—it would fundamentally change the economics of deep learning deployment.

The paper's focus on operator fusion (Level 2 of KernelBench) is particularly significant here. Modern neural networks increasingly rely on custom fused operators—combinations of primitive operations executed as a single kernel to avoid the overhead of multiple kernel launches and intermediate tensor materialization in global memory. FlashAttention is the canonical example: by fusing the entire attention computation into a single carefully tiled kernel, it achieves order-of-magnitude speedups over naive implementations. But FlashAttention required months of expert engineering. The vision behind CUDA Agent is that an LLM, trained with execution-driven RL, could automatically discover similar fusion opportunities for arbitrary operator compositions—effectively democratizing the kind of kernel optimization that currently requires elite human expertise.

Theoretical significance: bridging code generation and performance optimization. Prior LLM code generation research has largely focused on functional correctness—does the generated code pass unit tests? For CUDA kernels, correctness is table stakes. The real challenge is performance optimization, which requires reasoning about a fundamentally different set of constraints: memory bandwidth, cache hierarchies, warp-level parallelism, and hardware-specific throughput characteristics. A kernel that computes the correct mathematical result but uses uncoalesced memory accesses, excessive global memory round-trips, or suboptimal block sizes will be functionally correct but practically useless.

This distinction matters because it exposes a capability gap that extends beyond CUDA. The same pattern—LLMs can produce correct solutions but cannot optimize them—likely applies to database query optimization, distributed systems configuration, FPGA synthesis, and other domains where performance depends on deep understanding of the underlying hardware or system architecture. By tackling CUDA kernel optimization as a concrete, measurable instance of this broader challenge, the paper develops techniques (execution-driven RL, skill-augmented environments, robust reward design) that may generalize to other performance-critical code generation tasks.

The training-inference gap for low-resource programming domains. The paper identifies a more subtle but equally important problem: the severe distribution mismatch between pretraining data and CUDA code. The authors estimate that CUDA kernel code constitutes less than 0.01% of the pretraining data for models like Seed1.6 (Section 3.3, citing analyses of The Stack [10] and StarCoder [13]). This means that even a model with 230 billion total parameters has seen vanishingly few examples of CUDA kernel implementations during pretraining. The token probabilities the model assigns to CUDA-specific syntax, idioms, and optimization patterns are near the numerical precision floor (approximately $10^{-9}$ ), creating a fragile foundation for any fine-tuning approach.

This is not unique to CUDA. Many specialized programming domains—shader languages, Verilog/VHDL for hardware description, scientific computing with specialized libraries—are similarly underrepresented in general code pretraining corpora. The instability analysis in Section 3.3 (training collapses after 17 steps without proper initialization) provides concrete evidence that standard RL fine-tuning approaches break down when the target domain is this far from the pretraining distribution. The solutions developed—rejection fine-tuning for the actor, value pretraining for the critic, and multi-stage warm-up—constitute a general recipe for adapting LLMs to low-resource specialized programming domains.

Where Existing Approaches Fall Short

The paper identifies two broad paradigms of prior work, both of which fail to fundamentally solve the problem for different reasons.

Several systems attempt to improve CUDA kernel generation without modifying the underlying model weights, instead relying on hand-designed workflows that orchestrate the model's interactions with execution feedback.

STARK [5] constructs a multi-agent system with specialized roles—planning, coding, and debugging—that explore a tree-structured search space of possible kernel implementations. Each agent has a fixed function: the planner proposes optimization strategies, the coder implements them, and the debugger fixes compilation or correctness errors. ReGraphT [6] takes a retrieval-augmented approach, distilling optimization trajectories into a reasoning graph that can be searched via Monte Carlo Graph Search to guide kernel generation. EvoEngineer [8] formulates kernel optimization as evolutionary code editing, where an LLM iteratively mutates kernel implementations and selection pressure favors faster, correct versions. CudaForge [26] employs a two-agent system where a Judge agent analyzes Nsight Compute profiling data to provide targeted feedback to a Coder agent.

The critical limitation of all these approaches—which the paper states directly in Section 1—is that they "do not remedy the fundamental lack of CUDA-coding abilities in the base models." The base model's probability distribution over CUDA tokens remains essentially unchanged; the training-free orchestration merely samples from it more intelligently. The performance gains are therefore "significantly capped by the model's intrinsic capabilities." If the base model has never learned that diagonal matrix multiplication can be optimized to row-wise scaling (Case D.2), no amount of tree search or evolutionary pressure will discover that transformation—the model simply assigns negligible probability to the relevant code patterns.

This is empirically visible even in the strongest proprietary models. Claude Opus 4.5 and Gemini 3 Pro achieve respectable pass rates (91–95%) on KernelBench, meaning they can write correct kernels. But their faster rates against torch.compile remain at 66–69%—a third of their correct kernels are no faster than what a static compiler produces automatically. The training-free refinement approaches can improve on this baseline, but their ceiling is the base model's optimization knowledge, which—as the pass-rate-vs-faster-rate gap reveals—is limited.

Fine-Tuning with Execution Feedback

A second line of work attempts to improve the model itself through training, but the paper identifies fundamental limitations in how existing methods construct the training process.

Kevin [4] introduces multi-turn RL for CUDA kernel generation, modeling the iterative developer workflow where the model sees compilation errors and refines its implementation. However, the paper notes a critical issue: Kevin trains on a subset of KernelBench itself, meaning "the reported gains may partially reflect benchmark-specific adaptation rather than generalizable kernel generation capability" (Appendix C). This is a form of data leakage—the model is evaluated on problems it was trained on—which inflates apparent performance.

CUDA-L1 [14] suffers from an even more severe version of this problem, directly constructing supervised fine-tuning data from KernelBench reference implementations. The paper states this renders the results "not directly comparable to approaches, such as ours, that strictly avoid using KernelBench for training."

ConCuR [11] synthesizes CUDA kernels with reasoning traces using the Kevin-32B model, then fine-tunes a separate model (KernelCoder) on this data. But since Kevin-32B was itself trained on KernelBench, the contamination propagates downstream—"the performance of KernelCoder do not reflect training from independently curated or real-world kernel optimization data."

Beyond the data contamination issue, the paper identifies a deeper problem with existing fine-tuning approaches: they use fixed multi-turn refinement loops. The training protocol prescribes a specific sequence of actions—generate a kernel, see compilation errors, fix them, profile, optimize—rather than allowing the model to learn autonomously what sequence of actions is most effective. The difference is analogous to supervised learning with a fixed curriculum versus reinforcement learning where the agent discovers its own exploration strategy. A fixed loop cannot teach the model, for example, that it should run the profiler before attempting optimization to identify bottlenecks, or that certain types of errors should trigger a complete restart rather than incremental debugging. The model learns to follow the prescribed script rather than developing genuine CUDA debugging and optimization expertise.

Finally, all prior fine-tuning approaches operate at limited scale. The paper notes they are "fundamentally constrained by the scarcity of high-quality training data, limited training scale, and hand-designed optimization loops, which collectively cap their performance improvements" (Section 2.2). CUDA kernel datasets are small (KernelBench has only 250 problems), and manually writing reference implementations is prohibitively expensive for RL training, which typically requires thousands to tens of thousands of training tasks.

How This Paper Positions Itself

CUDA Agent's positioning relative to these prior approaches is defined by three strategic choices that directly address the identified failures.

From fixed loops to autonomous agents. Rather than prescribing a specific refinement protocol, CUDA Agent trains the model to operate as an autonomous agent within a realistic development environment (Section 3.2). The model is given standard software engineering tools (bash shell, file editing, search, compilation) and CUDA-specific utilities (verification scripts, profiling scripts), but it is not told what sequence of actions to take. Through RL, the model learns its own optimization strategy—it discovers when to profile, when to fuse kernels, when to use cuDNN libraries, and when to restart from scratch. The paper frames this as a move toward "active systems optimizers" rather than "passive code generators" (Section 5).

This is a fundamentally different problem formulation. Training-free methods apply external intelligence (hand-designed heuristics, multi-agent orchestration) to a static model. Fixed-loop fine-tuning imposes external structure on the learning process. CUDA Agent instead treats the development workflow itself as the action space and lets the model learn an internal policy over that space. The 200-turn interaction horizon and 128k token context window enable complex, multi-stage optimization trajectories that no hand-designed protocol would anticipate.

From benchmark contamination to independent data synthesis. The paper explicitly avoids using KernelBench for training, instead constructing a synthetic dataset of 6,000 operator tasks through a three-stage pipeline: crawling seed operators from PyTorch and Transformers libraries, using an LLM to compose them into fused multi-operator tasks, and filtering for executability, determinism, and appropriate difficulty (Section 3.1). This synthetic data is verified to have low AST-level similarity to KernelBench test cases (the maximum similarity across all training samples is below 0.9 after filtering; Appendix A, Figure 7).

The key insight driving this data synthesis is that operator fusion creates genuinely new optimization problems. When multiple operators are composed—say, a matrix multiplication followed by a reduction followed by a scaling—the optimal fused implementation is not equivalent to optimizing each operator individually and chaining them. Fusion reshapes the optimization landscape because it eliminates intermediate global-memory materialization, couples stages through shared register and shared memory constraints, and requires a unified parallel mapping and data layout that may favor downstream consumption. These are exactly the kinds of optimization challenges that torch.compile struggles with (as evidenced by CUDA Agent's massive 2.80× speedup on Level 2, where operator fusion is the primary challenge), making the synthetic data highly relevant to the target capability.

From training instability to multi-stage warm-up. The paper's third strategic contribution is addressing the fundamental instability that arises when fine-tuning a model on a domain far from its pretraining distribution. The observation that initial RL attempts collapsed after just 17 steps (Section 3.3) is not incidental—it reveals that the CUDA domain's extreme underrepresentation in pretraining data creates numerical instabilities in the PPO importance sampling ratios. When token probabilities are near the precision floor ( $\pi_\theta(a_t | s_t) \approx 10^{-9}$ ), small floating-point errors from mixed-precision training (BF16 vs. FP16) cause the importance ratio $\rho_t(\theta) = \frac{\pi_\theta(a_t|s_t)}{\pi_{\theta_{\text{old}}}(a_t|s_t)}$ to fluctuate wildly or explode.

The proposed solution—single-turn RL warm-up followed by RFT-based actor initialization and value-pretraining-based critic initialization (Figure 3)—is not ad hoc. Each stage addresses a specific failure mode identified through ablation analysis:

Single-turn RL provides a first pass at adapting the base model to produce CUDA code, reducing the probability gap before the model enters the more complex agent loop.
RFT (Rejection Fine-Tuning) provides a behavioral prior that constrains entropy growth during PPO. Without it, the policy distribution becomes "increasingly diffuse, producing incoherent and poorly structured outputs" (Figure 4b), leading to catastrophic reward collapse.
Value Pretraining ensures the critic can immediately provide meaningful advantage estimates. Without it, the uninitialized critic fails to penalize fruitless search paths, leading to an "explosion in trajectory length" (Figure 5b) as the agent explores endlessly without guidance.

This multi-stage approach enables stable training for 150 steps—nearly 9× longer than the 17-step collapse point—with consistent reward growth, making large-scale agentic RL on a highly specialized domain feasible for the first time.

The Broader Framing: LLMs as Systems Optimizers

Underlying all of this is a broader thesis that the paper articulates in its conclusion: "equipping foundation models with structured environments and reliable execution-based rewards can transform them from passive code generators into active systems optimizers." This framing positions CUDA Agent not merely as a solution to a specific benchmark but as a proof of concept for a new class of AI systems—models that learn to optimize performance-critical software through interaction with execution environments, rather than through imitation of human-written examples.

The implications extend beyond CUDA. If this approach generalizes, the same methodology—synthetic task generation, skill-augmented environments, execution-based reward, stability-oriented RL—could be applied to optimizing SQL queries, configuring distributed systems, designing hardware layouts, or tuning compiler passes. Each domain has its own execution environment (query planners, simulators, synthesis tools) that can provide reliable reward signals, and each has a vast space of possible optimizations that static heuristics cannot fully explore. CUDA Agent demonstrates that the combination of LLM reasoning capabilities with RL-driven exploration can navigate these spaces more effectively than either hand-designed rules or human experts alone.

The paper explicitly acknowledges that it does not yet realize this full vision—it doesn't combine PRM tree-search with revisions, doesn't handle problems where the base model's capability is fundamentally insufficient (difficulty bin 5 in KernelBench, where no method achieves meaningful performance), and doesn't compare against more sophisticated compiler frameworks like TVM (Appendix E). But by establishing that learned optimization policies can consistently outperform static compiler heuristics—100% faster rate on Level 2, 2.80× geometric mean speedup—it provides the strongest evidence to date that this direction is viable and worth pursuing at scale.

3. Technical Approach

3.1 Reader Orientation

CUDA Agent is a training system — not a single model — that uses large-scale reinforcement learning to turn a general-purpose LLM into an autonomous CUDA kernel optimization agent: the system learns, through trial and error in a realistic development environment, to generate CUDA kernels that are both correct and consistently faster than torch.compile, a standard compiler-based optimization baseline. The problem it solves is the severe under-representation of CUDA code in pretraining data (less than 0.01% of pretraining tokens), which prevents LLMs from developing GPU-aware optimization strategies through standard supervised fine-tuning or inference-time refinement alone; the solution is a three-component architecture — (1) a synthetic data pipeline that generates 6,000 diverse operator tasks for RL training, (2) a skill-augmented agent loop with rigorous reward anti-hacking protections that provides reliable correctness and performance feedback, and (3) a multi-stage RL warm-up strategy that prevents training collapse by bridging the distribution gap between general-purpose code and specialized CUDA kernel programming.

3.2 Big-Picture Architecture (Diagram in Words)

The CUDA Agent training system has five major components through which data and models flow:

Scalable Data Synthesis Pipeline — crawls seed PyTorch/Transformers operators, uses an LLM to compositionally synthesize fused multi-operator tasks, and filters for executability, determinism, and appropriate difficulty. Produces the CUDA-Agent-Ops-6K dataset of 6,000 training problems.
Skill-Integrated Agent Environment — provides a sandboxed development workspace where the LLM, acting through standard software engineering tools (bash, file editing, search, compilation) and CUDA-specific skill scripts (verification and profiling), generates kernel implementations, receives compilation errors, runtime correctness feedback, and nanosecond-level latency measurements. This environment enforces rigorous anti-hacking constraints: file permission controls prevent modification of evaluation scripts, context managers block fallback to torch.nn.functional, and profiling uses warm-up iterations with device synchronization.
Multi-Stage RL Training Pipeline — the core learning loop with three sequential stages: (a) single-turn RL warm-up on the base model using PPO to adapt to CUDA code distribution, (b) rejection fine-tuning (RFT) on agent trajectories to initialize the actor model with a strong behavioral prior, and (c) value pretraining on those same trajectories to initialize the critic model with meaningful advantage estimates. The main PPO training then operates stably for 150 steps with 128K context windows and up to 150 training turns (200 at evaluation).
GPU Sandbox Pool — a CPU–GPU decoupled architecture where a Docker-based terminal sandbox handles CPU tasks (compilation, file operations) while a dedicated pool of 128 NVIDIA H20 GPUs provides process-isolated verification and profiling, eliminating inter-process interference for stable nanosecond-precision latency measurements.
PPO Optimization Loop — samples agent trajectories using the current policy, computes token-level advantages via Generalized Advantage Estimation (GAE), and updates both the actor (via clipped surrogate objective with asymmetric clipping ϵ_lower = 0.2, ϵ_higher = 0.28) and the critic (via mean squared error against GAE-computed targets).

Information flow: synthetic operators enter the agent environment → the LLM (guided by SKILL.md and tool specifications) generates multi-turn optimization trajectories — analyzing bottlenecks, writing CUDA kernels, compiling, verifying correctness, profiling latency, and iteratively refining → each trajectory yields an outcome reward based on the robust discrete schedule (Equation 1) → GAE computes per-token advantages along the trajectory → PPO updates the actor and critic to increase the probability of actions that led to high-performance, correct kernels.

3.3 Roadmap for the Deep Dive

First, the scalable data synthesis pipeline (Section 3.1) — the three-stage process of crawling, compositional synthesis, and rubric-based filtering that produces 6,000 training problems. Understanding this first is critical because the entire RL system depends on having diverse, executable operator tasks with reliable reward computation.
Second, the skill-integrated agent loop (Section 3.2) — the development environment, tool specifications, CUDA coding skill (SKILL.md), reward computation with the robust discrete schedule, and anti-hacking protections. This defines the action space and reward signal that the RL algorithm optimizes.
Third, the robust reward scheduling equation (Equation 1, Section 3.2) — the discrete {-1, 1, 2, 3} reward scheme that jointly evaluates correctness, speedup over eager execution, and speedup over torch.compile. This is the objective function that shapes all learning, and its design (discrete milestones vs. continuous speedup ratios) has a decisive effect on optimization outcomes.
Fourth, the stability analysis and multi-stage warm-up strategy (Section 3.3) — the root cause of training instability (probability values near numerical precision floor causing PPO importance ratio explosion), and the three-stage solution (single-turn RL → RFT → Value Pretraining) with each stage's specific objective function (Equations 2–5).
Fifth, the core PPO algorithm with GAE advantage estimation (Section 3.3) — the clipped surrogate objective, the asymmetric clipping bounds, the GAE formulation (Equation 3), and the critic loss (Equation 4). This is the engine that drives policy improvement across the 150 training steps.

3.4 Detailed, Sentence-Based Technical Breakdown

This is primarily a systems paper whose core idea is that large-scale agentic RL — combining synthetic training data, execution-driven reward, and multi-stage initialization — can bridge the severe distribution gap between general-purpose code pretraining and specialized CUDA kernel optimization, enabling an LLM to learn autonomous optimization strategies that consistently outperform static compiler heuristics.

Scalable Training Data Synthesis Pipeline

Problem: the existing data landscape is insufficient for RL. Supervised fine-tuning requires reference implementations — expert-written optimized CUDA kernels — which are prohibitively expensive to produce at scale. Reinforcement learning requires only a correct PyTorch reference implementation and an execution environment that can measure the generated kernel's correctness and speed, which is far cheaper. However, the existing public datasets for CUDA kernel benchmarking are small (KernelBench has only 250 problems) and cannot be used for training without data leakage. The paper therefore designs a three-stage pipeline to synthesize a large, diverse training set from existing well-maintained libraries.

Stage 1: Seed Problem Crawling. The system crawls individual operator implementations from the torch and transformers Python libraries. Each operator is represented as a torch.nn.Module subclass with an __init__ method for parameter initialization and a forward method specifying the computation. The paper explicitly restricts crawling to these two libraries because "these operator classes are widely used and well-maintained," and excludes "individually maintained repositories that lack sufficient code quality" (Section 3.1). Each crawled operator becomes a seed problem — a reference PyTorch implementation that the agent must accelerate — accompanied by two auxiliary functions: get_init_inputs() which constructs the inputs needed to instantiate the operator, and get_inputs() which generates runtime inputs for the forward method, making each sample a self-contained executable training task.

Stage 2: Combinatorial Problem Construction. To expand the dataset and introduce higher complexity — particularly the operator fusion challenges central to CUDA optimization — the system uses an LLM to synthesize composite operators from the crawled seeds. The prompt instructs the LLM to sample no more than 5 operator classes from the torch library and compose them sequentially by stacking their computations into a single forward method. The paper makes a deliberate exclusion: "We do not sample operator classes from the transformers library, as these operators are typically higher-level modules that already encapsulate multiple primitive operations" (Section 3.1). The key insight driving this composition is that fused multi-operator tasks are not equivalent to optimizing each operator individually and concatenating the results. Fusion avoids intermediate global-memory materialization, couples computation stages through shared register and shared memory constraints, and requires a unified parallel mapping that may favor downstream consumption. A kernel that fuses a matrix multiply with a subsequent reduction and scaling faces an entirely different optimization landscape than a naive chain of three separate kernels — exactly the kind of challenge that torch.compile's static pattern-matching struggles with and that the RL agent must learn to navigate.

Stage 3: Rubric-Based Problem Filtering. The synthesized operators are validated against four execution-based criteria:

Executability: The operator must run successfully in both PyTorch Eager and torch.compile modes, ensuring the reference implementation is correct and the baseline measurements can be computed.
Determinism: Operators with inherent stochasticity are excluded, since non-deterministic outputs would corrupt the reward signal — the correctness verification (comparing generated kernel output against reference output) requires bit-level or near-bit-level reproducibility.
Anti-hacking validation: The system verifies that outputs for different random inputs are neither constant values nor numerically indistinguishable. Without this check, the agent could learn to produce a kernel that returns a constant (e.g., always outputting zeros for a reduction operation), which might pass correctness checks for certain input distributions.
Computational load filtering: The execution time in Eager mode is constrained to the range 1 ms to 100 ms. Problems below 1 ms have too little room for optimization (kernel launch overhead dominates), while problems above 100 ms are impractically slow for RL training, which requires thousands of rollouts.

Additionally, an explicit decontamination step is applied using an off-the-shelf AST-based code similarity tool (PythonASTSimilarity). The system extracts the Model class from each training sample, computes pairwise structural similarity against all KernelBench evaluation samples, and removes any training sample whose maximum similarity exceeds a threshold of 0.9. After filtering, the distribution of maximum similarities (Appendix A, Figure 7) shows that no training sample exceeds the threshold, with the majority having low similarity scores — confirming that the training data is structurally distinct from the evaluation benchmark.

Final dataset composition. The filtered dataset, named CUDA-Agent-Ops-6K, contains 6,000 samples with the following composition (Table 3 in Appendix A): single torch operators (3.40%), two-operator compositions (83.77%), three-operator compositions (7.62%), four-operator compositions (2.80%), five-operator compositions (1.23%), and standalone transformers operators (1.18%). The dominance of two-operator compositions reflects that moderate fusion complexity provides the richest training signal — simple enough to be solvable by the agent, complex enough to require non-trivial optimization beyond what torch.compile achieves automatically.

Skill-Integrated Agent Loop

Design philosophy: the agent learns its own workflow rather than following a fixed script. The paper frames the agent loop design as a deliberate departure from prior work that uses "fixed multi-turn refinement loops driven by code execution feedback" (Section 1). In those approaches, the training protocol prescribes a specific sequence — generate kernel, compile, fix errors, profile, optimize — which teaches the model to follow the prescribed steps rather than developing genuine CUDA debugging and optimization expertise. CUDA Agent instead provides the model with a standard development environment (tools, verification scripts, profiling scripts, a structured workflow description) and uses RL to let the model discover its own optimization strategy over up to 200 interaction turns.

Environment architecture. The environment follows the ReAct-style paradigm (Yao et al., 2022), where each interaction turn consists of the model producing a reasoning trace (analyzing the current situation), selecting an action (e.g., running a shell command, editing a file, invoking the profiler), receiving an observation (the command output, compilation errors, profiling results), and integrating that observation into its next reasoning step. The LLM is provided with a standard suite of software engineering tools designed for the OpenHands framework (Wang et al., 2024):

BashTool: Executes shell commands in a persistent session with safety constraints (command quoting rules, directory validation). Used for compilation, dependency management, and running verification/profiling scripts.
Read / Write: Provide read-only and write access to local files, with write operations guarded by a read-before-write policy to prevent blind overwriting.
Edit / MultiEdit: Support deterministic string-level code modifications. Edit performs single replacements; MultiEdit enables multiple atomic edits within a single file, ensuring consistency across dependent code changes (e.g., updating both a kernel launch configuration and its corresponding binding code simultaneously).
Glob: Performs file discovery using glob patterns (e.g., **/*.py), enabling the agent to navigate the workspace structure efficiently.
Grep: A structured code search interface based on ripgrep, supporting regex search, file-type filtering, and contextual line retrieval for debugging and code inspection.
NotebookEdit: Enables structured modification of Jupyter notebook cells, supporting mixed code–analysis workflows.
BashOutput / KillBash: Stream incremental outputs from background shell processes (for monitoring long-running compilation or profiling jobs) and terminate sessions when jobs hang.

CUDA Coding Skill (SKILL.md). The paper adopts the Agent Skills paradigm (Anthropic, 2025), deliberately placing CUDA-specific instructions, tools, and workflows into a structured specification document (SKILL.md) rather than hard-coding them into the environment or training protocol. This separation allows the model to learn when and how to use CUDA-specific knowledge rather than having it imposed externally. The SKILL.md document (reproduced in full in Appendix B.2) defines:

Critical restrictions: A list of forbidden actions — no torch::* operations in C++ code (to prevent the agent from simply calling PyTorch's built-in kernels), no modifications to infrastructure files (binding.cpp, binding_registry.h, utils/), no third-party libraries except cuBLAS (for GEMM operations) and cuDNN (mandatory for convolutions), and focus on implementing kernels in the kernels/ directory.
Workspace structure: A fixed file layout with designated directories for agent-generated code (kernels/ for CUDA .cu and binding .cpp files, model_new.py for the optimized PyTorch model) and protected directories for infrastructure (utils/ for compilation, verification, and profiling tools).
Four-stage optimization workflow: (1) Analyze baseline performance using the provided profile.py script to identify bottlenecks, (2) implement custom CUDA operators by writing kernel source files, binding code, and an optimized model, (3) compile and evaluate in the sandbox environment, and (4) iterate until the implementation achieves at least 5% speedup over torch.compile while passing correctness checks.
Optimization strategy priority ordering: Algorithmic optimizations (kernel fusion, shared memory tiling, memory coalescing) with >50% expected impact, hardware utilization optimizations (vectorized loads, warp-level primitives, occupancy tuning) with 20–50% impact, and fine-tuning optimizations (instruction-level parallelism, mixed precision, prefetching) with <20% impact.
Code templates: Concrete CUDA kernel structure with grid-stride loops, dynamic configuration selection via switch statements, and extern "C" launcher functions that avoid PyTorch dependency in .cu files.
Compilation, verification, and profiling commands: Explicit instructions for using TORCH_CUDA_ARCH_LIST=9.0 bash utils/compile.sh for compilation targeting Hopper GPUs, and sudo python3 -m utils.verification / sudo python3 -m utils.profiling for sandbox execution.
Debugging guidance: Common compilation errors (undefined symbols, missing kernel images), correctness failure patterns (wrong outputs, NaN/Inf results, shape mismatches), and performance issues (poor occupancy, uncoalesced access, high kernel launch count) mapped to likely causes and solutions.

Crucially, the SKILL.md document provides the specification of what a CUDA optimization workflow looks like, but the RL training allows the model to learn how to execute that workflow effectively — which transformations to apply, in what order, and with what degree of aggressiveness — through trial and error guided by reward signals. The model is not constrained to follow the SKILL.md steps rigidly; it can discover non-obvious optimization sequences, skip steps when appropriate, or apply transformations not explicitly mentioned in the document.

Robust Reward Scheduling. The reward signal that drives all RL training jointly evaluates two orthogonal dimensions: functional correctness and execution latency. The paper identifies a critical problem with the naive approach of using raw speedup ratios as continuous rewards: "operators vary substantially in optimization difficulty, making raw speedup an unreliable proxy for code quality" (Section 3.2). A kernel that achieves 10× speedup on a trivially optimizable operator is not "better" code than one that achieves 1.5× speedup on a heavily optimized operator — but a continuous speedup ratio reward would treat them as such, biasing the policy toward easy operators and creating high variance in the reward signal.

The paper proposes a normalized, discrete reward scheme:

$r = \begin{cases} -1 & \text{if correctness check fails} \\ 3 & \text{if } b(t, t_{\text{eager}}) \land b(t, t_{\text{compile}}) \\ 2 & \text{if } b(t, t_{\text{eager}}) \\ 1 & \text{otherwise} \end{cases}$

where $t$ is the generated kernel's measured runtime, $t_{\text{eager}}$ is the runtime of the PyTorch eager implementation, $t_{\text{compile}}$ is the runtime of torch.compile, and $b(t, t_0) = \mathbb{I}[(t_0 - t) / t_0 > 5\%]$ is a binary indicator of whether the generated kernel achieves at least 5% speedup over baseline $t_0$ .

What it computes: The reward function maps each training trajectory to one of four discrete levels. Level -1 is assigned when the generated kernel fails the correctness verification — the output does not match the reference PyTorch implementation within numerical tolerances — making this a hard failure with negative reward regardless of any potential speed gains. Level 3 is the highest reward, assigned when the kernel is both correct and achieves at least 5% speedup over both the Eager baseline AND the torch.compile baseline — this is the target outcome, representing a kernel that outperforms the compiler. Level 2 is an intermediate achievement: correct and faster than Eager execution, but not faster than torch.compile — the kernel shows optimization skill but hasn't surpassed the compiler baseline. Level 1 is the minimum positive reward: correct kernel that is not significantly faster than either baseline — a functionally correct but unoptimized implementation.

Why this form: The discrete milestone structure serves three purposes. First, it normalizes away the operator difficulty variance: achieving level 3 on a hard operator and achieving level 3 on an easy operator both receive the same reward magnitude, preventing the bias toward easy optimization targets that a continuous speedup-ratio reward would create. Second, the 5% threshold for $b(t, t_0)$ prevents reward noise from nanosecond-level measurement fluctuations — small timing variations that don't represent genuine optimization improvements are thresholded out. Third, the multi-level structure provides a curriculum: the model can receive positive reinforcement for achieving correctness and beating the Eager baseline (levels 1–2) even before it learns to surpass torch.compile (level 3), preventing reward sparsity early in training. The paper validates this design empirically (Section 4.3.2): replacing the discrete schedule with a continuous speedup reward ( $r_s = t_{\text{compile}} / t_{\text{gen}}$ for correct solutions, -1 for incorrect) yields "comparable functional correctness, but substantially weaker optimization performance," indicating that the milestone-based structure more reliably guides the policy toward genuine speed improvements.

Anti-hacking protections. The paper implements a comprehensive set of safeguards to prevent the agent from obtaining high rewards through trivial or dishonest behaviors rather than genuine kernel optimization, addressing known failure modes from prior work (the paper explicitly cites the "150× speedup that was actually 3× slower" case from Lange et al., 2025, as a cautionary example):

File permission controls: The evaluation Python scripts for correctness verification and performance profiling are protected via system-level file permissions, preventing the agent from modifying the evaluation logic to report inflated speedups or bypass correctness checks.
Context manager enforcement: Execution-time constraints using Python context managers explicitly block the agent from invoking fallback implementations via torch.nn.functional. Without this constraint, the agent could produce a model_new.py that simply calls torch.nn.functional.conv2d or torch.matmul — which would appear to be a "generated kernel" but actually delegates to PyTorch's built-in implementations, achieving no actual optimization.
Multi-input correctness validation: For each problem, the verification script tests kernel outputs against five randomly sampled inputs, strictly following the KernelBench evaluation protocol. This prevents the agent from overfitting to a single input configuration or exploiting edge cases.
Measurement rigor: The profiling pipeline includes proper CUDA device synchronization (torch.cuda.synchronize()), warm-up iterations to eliminate cold-start effects (first-kernel-launch overhead, cache warming), and repeated measurements with averaging to reduce run-to-run variance. Without these, the agent could exploit measurement artifacts — for example, timing a kernel before GPU warm-up and then comparing against a properly warmed-up baseline.
No external information retrieval: The agent is not provided with web search or external information retrieval tools, ensuring all solutions are derived purely from the local execution environment through trial-and-error experimentation. This prevents the agent from simply searching for and copying existing optimized kernel implementations.

Agent loop execution flow. For each training problem, the interaction proceeds as follows: The model receives the problem specification (the PyTorch operator class in model.py and the workspace structure) and the SKILL.md content. It begins by analyzing the baseline implementation — typically running sudo python3 -m utils.profiling to measure Eager and torch.compile runtimes and identify bottlenecks. It then generates CUDA kernel code (writing .cu files with __global__ kernel functions and _binding.cpp files with PyTorch tensor handling and registration), updates model_new.py to use the custom operators, compiles with TORCH_CUDA_ARCH_LIST=9.0 bash utils/compile.sh, verifies correctness with sudo python3 -m utils.verification, and profiles performance. If the verification fails, the model interprets compilation errors or numerical mismatches and debugs the kernel. If performance is insufficient, the model applies optimization transformations — kernel fusion, shared memory tiling, vectorized loads, occupancy tuning, etc. — guided by the SKILL.md optimization hierarchy. This cycle repeats until the model determines it has achieved the best possible performance (or exhausts its turn budget), at which point it cleans up intermediate files and finalizes the implementation.

Multi-Stage Warm-Up Strategy

The root cause of training instability: severe distribution mismatch. The paper reports that initial RL attempts — directly applying PPO to the base model in the agent loop — collapsed after only 17 training steps, with performance rapidly degrading rather than improving (Section 3.3). The diagnostic analysis traces this collapse to a numerical precision issue arising from the extreme rarity of CUDA code in pretraining data. The pretraining corpora analyzed by Kocetkov et al. (2022) and Li et al. (2023) indicate that CUDA kernel code constitutes less than 0.01% of training tokens. Consequently, the base model assigns extremely low probabilities to CUDA-specific tokens and syntactic patterns — values on the order of $\pi_\theta(a_t | s_t) \approx 10^{-9}$ , near the representational floor of 16-bit floating-point formats.

The PPO algorithm computes importance sampling ratios to correct for the mismatch between the current policy and the policy that generated the training trajectory:

$\rho_t(\theta) = \frac{\pi_\theta(a_t | s_t)}{\pi_{\theta_{\text{old}}}(a_t | s_t)}$

where $\pi_\theta$ is the current policy (being optimized), $\pi_{\theta_{\text{old}}}$ is the behavior policy that generated the trajectory, and $a_t$ is the action (token) at position $t$ .

What it computes: the ratio of how much more (or less) likely the current policy is to produce token $a_t$ in state $s_t$ compared to the policy that originally generated it. This ratio is multiplied by the advantage estimate $\hat{A}_t$ in the PPO objective to increase the probability of tokens that led to good outcomes and decrease the probability of tokens that led to bad outcomes.

Why this form and why it fails here: When token probabilities are near the precision floor ( $10^{-9}$ ), small numerical errors — for example, from the training engine using BF16 while the inference engine uses FP16, a common setup for efficiency — cause the numerator and denominator to diverge in ways that don't reflect actual policy change. A token whose true probability under both policies is $10^{-9}$ might be computed as $1.1 \times 10^{-9}$ in BF16 and $0.9 \times 10^{-9}$ in FP16, producing an importance ratio of $\sim 1.22$ — a 22% apparent change when the actual policy hasn't changed at all. For tokens near the precision floor, these spurious fluctuations accumulate, causing the PPO clipped surrogate objective's gradient estimates to become dominated by noise. The paper connects this to the analysis in Liu et al. (2025), which formalizes how training-inference numerical precision mismatches cause PPO collapse when operating near representational boundaries.

The consequence is not just noisier updates but catastrophic unlearning: the policy's entropy sharply increases (Figure 4b), the model begins producing increasingly diffuse and incoherent outputs, and rewards collapse as the policy loses the ability to generate syntactically valid CUDA code.

Stage 1: Single-Turn RL Warm-up. The first stage addresses the probability gap head-on by performing PPO training in a simpler setting before introducing the full agent loop. The base model is trained on single-turn CUDA generation — the model receives a problem specification and produces the complete optimized kernel and binding code in a single response, without any interaction with compilation, verification, or profiling feedback. The optimization uses the same PPO algorithm and hyperparameters as the main training (global batch size 1024, actor learning rate $3 \times 10^{-6}$ , critic learning rate $6 \times 10^{-6}$ , context window 32,768 tokens). This stage serves to raise the model's probability mass on CUDA-appropriate tokens from the $10^{-9}$ range to a regime where the PPO importance ratios become numerically stable — it doesn't teach sophisticated optimization strategies (those require the agent loop's execution feedback), but it provides a safe "warm start" for the next stage.

Stage 2: Actor Initialization via Rejection Fine-Tuning (RFT). After single-turn RL, the warmed-up model is used to collect full agent trajectories by running it in the complete skill-integrated agent loop (Section 3.2). The model generates multi-turn optimization trajectories — sequences of reasoning, action, and observation that may include multiple rounds of profiling, kernel writing, compilation, debugging, and performance tuning. From these collected trajectories, only high-quality rollouts are retained according to a two-stage rejection criterion:

Outcome filtering: trajectories must achieve positive reward ( $R > 0$ ), meaning the final kernel must be at least correct (reward level ≥ 1). This filters out trajectories where the model failed to produce a working kernel.
Pattern filtering: trajectories exhibiting inefficient or invalid behaviors are discarded, specifically those with "redundant multi-turn loops or hallucinations that violate the predefined tool-call schema." This prevents the model from learning degenerate exploration patterns — endlessly cycling through the same failed optimization attempts or producing tool calls that don't conform to the expected interface.

The filtered trajectories are used to fine-tune the actor model via standard supervised learning:

$\mathcal{L}_{\text{RFT}}(\theta) = -\mathbb{E}_{\tau \sim \mathcal{D}'}\left[\sum_{t=1}^T \log \pi_\theta(a_t \mid s_t, a_{<t})\right]$

where $\tau = (s_0, s_1, \ldots, s_{T-1})$ is a filtered agent trajectory from dataset $\mathcal{D}'$ , $\pi_\theta$ is the policy parameterized by $\theta$ , $a_t$ is the action token at position $t$ , and $s_t$ is the state (the entire interaction history up to that point).

What it computes: the standard next-token prediction cross-entropy loss, summed over all tokens in the trajectory. For each position, it maximizes the log-probability that the model assigns to the action token $a_t$ that was actually taken in the successful trajectory, given the preceding state.

Why this form: this is maximum-likelihood estimation over the set of high-reward trajectories. By training only on trajectories that achieved positive rewards and exhibited efficient behavior, the objective pushes the model's distribution toward patterns that led to success — correct kernel generation, effective use of profiling, strategic optimization choices — without requiring explicit reward modeling at each step. This provides a strong behavioral prior that anchors the policy in a region of the action space known to produce valid, reward-achieving behavior.

The empirical effect is visible in Figure 4: without RFT, the policy's entropy spikes sharply as PPO training progresses, and training rewards collapse within approximately 10 steps. With RFT, the entropy remains controlled — the policy stays within a well-structured output distribution, and rewards remain stable for the full 150-step training run. The paper diagnoses the collapse as the policy becoming "increasingly diffuse, producing incoherent and poorly structured outputs" when not anchored by RFT's behavioral prior.

Stage 3: Critic Initialization via Value Pretraining. While RFT stabilizes the actor, an uninitialized critic creates a different failure mode. The critic network $V_\phi$ is responsible for estimating the expected future return from each state, which the PPO algorithm uses to compute advantages — how much better or worse a particular action was compared to the average expected outcome. Without meaningful value estimates, the advantage computation becomes noisy and unreliable.

The paper performs Value Pretraining on the same filtered agent trajectories used for RFT. For each trajectory $\tau = (s_0, s_1, \ldots, s_{T-1})$ , the outcome reward $r$ is assigned only at the final token (i.e., $r_t = 0$ for $t < T-1$ and $r_{T-1} = r$ , where $r$ is the discrete reward from the robust reward schedule). Target values are computed using Generalized Advantage Estimation (GAE) (Schulman et al., 2018):

$V_t^{\text{targ}} = V_\phi(s_t) + \hat{A}_t$

where $\hat{A}_t$ is the GAE advantage estimate at position $t$ , $V_\phi(s_t)$ is the critic's current value estimate for state $s_t$ , and $V_t^{\text{targ}}$ is the regression target.

The GAE advantage is computed as:

$\hat{A}_t = \sum_{l=0}^{T-1-t} (\gamma \lambda)^l \delta_{t+l}$

where $\gamma = 1$ is the discount factor (set to 1 since the task has a finite horizon and the reward occurs only at the end), $\lambda = 0.95$ is the GAE trace-decay parameter, and:

$\delta_t = r_t + \gamma V_\phi(s_{t+1}) - V_\phi(s_t)$

is the temporal difference (TD) error — the difference between the actual reward plus the estimated future value and the current value estimate. The termination value is defined as $V_\phi(s_T) = 0$ (no reward after the trajectory ends).

What GAE computes: a weighted sum of future TD errors, with exponentially decaying weights controlled by $\lambda$ . Each TD error $\delta_{t+l}$ measures whether the outcome at step $t+l$ was better or worse than the critic had predicted. The $(\gamma \lambda)^l$ weighting gives more influence to near-term TD errors while still incorporating distant outcomes. When $\lambda = 0$ , GAE reduces to the one-step TD error $\delta_t$ ; when $\lambda = 1$ , it approaches the Monte Carlo return (sum of all future rewards). The intermediate value $\lambda = 0.95$ balances bias (from imperfect critic estimates) and variance (from stochastic trajectory outcomes).

What $V_t^{\text{targ}}$ computes: for each state $s_t$ , it produces a target value that equals the critic's current estimate $V_\phi(s_t)$ plus the GAE advantage estimate. In expectation, this converges to the true expected return from $s_t$ as the critic becomes more accurate (since GAE's advantage estimates become zero-mean when the value function is correct). The critic is then trained to minimize the mean squared error against these targets:

$\mathcal{L}_{\text{VP}}(\phi) = \frac{1}{2} \mathbb{E}_{\tau \sim \mathcal{D}}\left[ \frac{1}{T} \sum_{t=0}^{T-1} \left(V_\phi(s_t) - V_t^{\text{targ}}\right)^2 \right]$

where $\mathcal{D}$ is the collection of agent trajectories, $T$ is the trajectory length, and $V_\phi(s_t)$ is the critic's predicted value for state $s_t$ .

What it computes: the average squared difference between the critic's predicted value and the GAE-computed target value, averaged over all positions in all trajectories. The $1/2$ factor is conventional for gradient convenience (it cancels the factor of 2 from the squared term's derivative). This is a standard regression objective: the critic learns to predict, for any given state in a CUDA optimization trajectory, what final reward the agent is likely to achieve if it continues following the current policy.

Why this form: Value Pretraining provides the critic with an informed initialization — it learns to distinguish promising states (where the generated kernel is on track to be correct and fast) from unpromising ones (where the agent is stuck in a debugging loop or producing inefficient code) before the main PPO training begins. Without it, the paper reports two critical failures (Figure 5): (1) the explained variance of the value function (a measure of how well the critic predicts actual returns) remains near zero throughout training, indicating the critic has learned essentially nothing about the value landscape; (2) the response length clipped ratio — the fraction of trajectories that hit the maximum turn limit — explodes, because the uninitialized critic "fails to penalize fruitless or redundant search paths," causing the agent to explore endlessly without any signal about which directions are productive and which are dead ends. Value Pretraining ensures that from the first PPO update, the critic provides meaningful advantage estimates that guide the agent toward efficient optimization paths.

The three-stage warm-up is the paper's key enabler: single-turn RL raises token probabilities out of the precision-floor danger zone, RFT anchors the actor in a stable output distribution, and Value Pretraining equips the critic with initial value estimates that prevent infinite-loop pathologies. Together, they enable stable PPO training for 150 steps — nearly 9× the 17-step collapse point — with consistent reward growth throughout.

PPO Training with Generalized Advantage Estimation

With the actor and critic initialized, the main RL training phase applies PPO (Schulman et al., 2017) to optimize the agent's CUDA optimization policy through interaction with the skill-integrated environment. The training operates with context windows of 131,072 tokens, a maximum of 150 agent turns during training rollouts (relaxed to 200 turns during evaluation), a global batch size of 1,024 matching the mini-batch size for online PPO updates, actor learning rate $3 \times 10^{-6}$ , and critic learning rate $6 \times 10^{-6}$ .

The clipped surrogate objective. PPO optimizes the policy $\pi_\theta$ by maximizing a conservative surrogate objective that prevents destructively large policy updates:

$\mathcal{L}_{\text{CLIP}}(\theta) = \mathbb{E}_{\tau \sim \mathcal{D}}\left[ \frac{1}{T} \sum_{t=0}^{T-1} \min\left( \rho_t(\theta) \hat{A}_t, \ \text{clip}\left(\rho_t(\theta), 1 - \epsilon_{\text{lower}}, 1 + \epsilon_{\text{higher}}\right) \hat{A}_t \right) \right]$

where $\rho_t(\theta) = \frac{\pi_\theta(a_t|s_t)}{\pi_{\theta_{\text{old}}}(a_t|s_t)}$ is the importance sampling ratio between the current and old policies (same as defined above), $\hat{A}_t$ is the GAE advantage estimate at position $t$ , $\epsilon_{\text{lower}} = 0.2$ and $\epsilon_{\text{higher}} = 0.28$ are the asymmetric clipping bounds, and the $\min$ operation is taken element-wise over the two terms.

What it computes: The objective takes the expected value over all trajectories and all token positions of the minimum between two quantities: (1) the unclipped importance-weighted advantage $\rho_t(\theta) \hat{A}_t$ , which is the standard policy gradient update — increase probability of tokens with positive advantage, decrease for negative advantage, scaled by how much the current policy differs from the old policy; and (2) the clipped version $\text{clip}(\rho_t(\theta), 1 - \epsilon_{\text{lower}}, 1 + \epsilon_{\text{higher}}) \hat{A}_t$ , which caps the importance ratio to the range $[0.8, 1.28]$ — preventing any single token's probability from changing by more than a factor of 0.8 or 1.28 relative to the old policy.

Why this form and the asymmetric clipping: The standard PPO algorithm uses symmetric clipping ( $\epsilon = 0.2$ ), meaning probability increases and decreases are equally constrained. This paper adopts asymmetric clipping ( $\epsilon_{\text{lower}} = 0.2$ , $\epsilon_{\text{higher}} = 0.28$ ) following the DAPO approach (Yu et al., 2025). The asymmetry allows the policy to increase probabilities of good actions more aggressively (up to 1.28×) than it can decrease probabilities of bad actions (down to 0.8×). The intuition is that in multi-turn agent trajectories, many token-level actions are "neutral" — they represent routine file editing or shell commands that don't directly cause the final outcome — and aggressively down-weighting them would destabilize the policy's basic interaction capabilities. The looser upper bound encourages faster adoption of clearly beneficial patterns (successful optimization strategies, effective debugging sequences) without destabilizing the base interaction behavior.

The $\min$ operation ensures the objective is a pessimistic bound: when the advantage is positive and the policy would like to increase $\pi_\theta(a_t|s_t)$ substantially, the clipping caps the effective objective at the clip boundary, preventing over-optimistic updates. When the advantage is negative and the policy would like to decrease $\pi_\theta(a_t|s_t)$ , the clipping similarly caps the penalty. This conservatism is what makes PPO robust to the noisy advantage estimates that arise from the long, multi-turn trajectories — without it, a single high-advantage action in an otherwise mediocre trajectory could cause a destructively large policy update.

Critic training during PPO. The critic continues to be trained throughout the PPO phase using the same GAE-based target computation and MSE loss as in Value Pretraining (Equation 4), but now on fresh trajectories sampled by the current policy rather than the frozen pre-RFT trajectories. The critic's value estimates are used to compute the GAE advantages $\hat{A}_t$ that drive the actor's updates, creating a co-adaptation loop where improved value estimates lead to better advantage estimates, which lead to better policy updates, which produce higher-reward trajectories, which provide better training data for the critic.

Training dynamics. Over the 150 training steps, the policy learns to navigate the optimization space autonomously. The paper doesn't provide detailed training curves (beyond the ablation-related Figures 4a and 5a), but the ablation results in Table 2 demonstrate the necessity of each component: removing the agent loop drops the faster rate from 96.8% to 14.1% (the model loses the ability to use execution feedback for iterative refinement), removing the robust reward drops optimization performance from 2.60× to 1.70× geometric mean speedup (the continuous speedup reward misguides the policy), removing RFT drops to 1.56× speedup before training collapses, and removing Value Pretraining drops to 1.49× speedup with trajectory length explosion. The full system achieves 2.60× geometric mean speedup over Eager and 2.11× over torch.compile, with a 98.8% pass rate and 96.8% faster rate over Eager — indicating that the combined architecture successfully produces an agent that not only learns to write correct CUDA kernels but learns to optimize them to consistently outperform static compiler heuristics.

Case study evidence of learned strategies. The three case studies in Appendix D illustrate the types of optimization strategies the trained agent discovers. In Case D.2 (Level 1: diagonal matrix multiplication), the agent recognizes the algebraic equivalence between a diagonal matrix construction followed by GEMM and simple row-wise scaling, reducing $O(N^2 M)$ to $O(NM)$ and achieving 73.31× speedup over torch.compile. In Case D.3 (Level 2: matrix multiply, division, summation, scaling), the agent algebraically rearranges the computation to pre-compute column-wise weight sums, then fuses the remaining operations into a single kernel with vectorized float4 loads and shared-memory tree reduction, achieving 24.04× speedup. In Case D.4 (Level 3: ResNet BasicBlock), the agent folds batch normalization into convolution weights (eliminating a kernel launch), uses cudnnConvolutionBiasActivationForward to fuse convolution, bias, and ReLU into a single cuDNN call, enables TF32 tensor core computation, and writes a custom fused add-ReLU kernel for the residual connection, achieving 3.59× speedup. These examples demonstrate that the agent has learned to combine high-level algebraic reasoning, library-aware optimization, kernel fusion, and hardware-specific configuration — a breadth of strategies that no hand-designed refinement loop would prescribe.

4. Key Insights and Innovations

Innovation 1: Reframing LLM Code Generation as a Systems Optimization Problem Rather Than a Pattern-Matching Task

The paper's most fundamental conceptual contribution is not any single technical method but the reframing of what it means for an LLM to generate code in a performance-critical domain. Prior work on CUDA kernel generation—both training-free orchestration systems (STARK, EvoEngineer, CudaForge) and fine-tuning approaches (Kevin, CUDA-L1, ConCuR)—implicitly treated the problem as one of retrieving and adapting known optimization patterns: given a PyTorch operator, find or synthesize a CUDA kernel that implements it more efficiently. The base model's role was to serve as a repository of code patterns, and the training or search mechanism's role was to navigate toward the right pattern.

CUDA Agent fundamentally reframes this. The problem is not pattern retrieval but autonomous systems optimization: an agent operating in an execution environment, making sequential decisions about what to measure, what to modify, and what to try next, guided solely by feedback from the hardware itself. This reframing has several consequences that distinguish it from the prior paradigm:

From static pattern knowledge to learned exploration strategy. In the pattern-retrieval view, the model's capability is bounded by what optimization patterns it has seen during training. A model that has never encountered the algebraic equivalence between diagonal matrix multiplication and row-wise scaling cannot apply that optimization—regardless of how many search iterations or evolutionary rounds you wrap around it. CUDA Agent demonstrates that through RL, the model can discover optimizations it has never explicitly seen, by experimenting with transformations and observing the performance consequences. The 73.31× speedup in the diagonal matrix multiplication case (Appendix D.2) is not a pattern retrieved from memory—it emerges from the agent's learned strategy of analyzing the mathematical structure of the computation and algebraically simplifying before writing kernels.

From demonstration-driven learning to outcome-driven learning. Prior fine-tuning approaches (Kevin, ConCuR) rely on expert demonstrations—reference implementations that show the model what good kernels look like. This is expensive to scale and limits the model to optimizations that humans have already invented. CUDA Agent's RL approach requires only a correct PyTorch reference implementation and a way to measure the generated kernel's speed—no expert CUDA code is needed for training. The model learns what works rather than what humans have done. This is a more general learning paradigm: it applies to any domain where correctness can be verified and performance can be measured, regardless of whether expert implementations exist.

From fixed workflows to emergent optimization strategies. Prior training-free systems encode a specific optimization workflow—STARK's plan-code-debug tree search, CudaForge's profile-feedback-edit loop—that the model must follow. These workflows represent the system designer's hypothesis about how optimization should proceed. CUDA Agent provides tools and a skill document describing possible actions, but lets RL discover the actual sequence that works best. The agent learns autonomously when to profile (before optimization, not after), when to use cuDNN library functions versus writing custom kernels, when to apply algebraic simplification versus hardware-level tuning, and when to abandon a failed approach and restart. The case studies in Appendix D reveal sophisticated multi-stage strategies—folding batch normalization into convolution weights, then enabling TF32, then fusing the residual addition with ReLU—that no hand-designed workflow explicitly prescribes.

This reframing is significant beyond CUDA. It establishes a template for using RL to train LLMs as performance engineers in any domain with executable specifications and measurable performance: database query optimization, distributed systems configuration, compiler pass ordering, hardware design. The key ingredients—synthetic task generation, execution-based reward, skill-augmented environments—are domain-agnostic. The paper thus contributes not just a better CUDA kernel generator but a new class of AI capability: foundation models that learn to optimize systems through interaction rather than through imitation.

Innovation 2: Diagnosing and Solving the "Precision-Floor Collapse" as the Central Barrier to RL in Low-Resource Programming Domains

A second distinctive contribution is the paper's identification and systematic resolution of a specific failure mode that arises when applying RL to domains far from the pretraining distribution. This is more than an engineering fix—it is a diagnostic concept with implications for any effort to adapt LLMs to specialized, low-resource programming tasks.

The failure mode: catastrophic entropy explosion from numerical precision artifacts. The paper reports that direct PPO training on the base model in the agent loop collapsed after just 17 steps (Section 3.3). The diagnostic analysis reveals a mechanism that is subtle and non-obvious: when the model assigns probabilities to domain-specific tokens near the representational floor of 16-bit floating-point formats ( $\pi_\theta(a_t | s_t) \approx 10^{-9}$ ), the PPO importance sampling ratio $\rho_t(\theta) = \pi_\theta(a_t|s_t) / \pi_{\theta_{\text{old}}}(a_t|s_t)$ becomes numerically unstable. Small differences between training precision (BF16) and inference precision (FP16)—a standard setup for throughput—cause these near-floor probabilities to diverge in ways that don't reflect actual policy change. The resulting noise in importance ratios causes the policy gradient to become dominated by artifacts, the policy entropy spikes as the model begins producing increasingly random outputs, and performance collapses.

Why this is a conceptual contribution, not just an engineering detail. Prior work on RL for language models (the RLHF literature, the DAPO paper by Yu et al., 2025) has focused on reward hacking, distributional shift between the policy and reference model, and KL-constrained optimization as the primary failure modes. The precision-floor collapse identified here is a qualitatively different failure mode that arises specifically from domain rarity—it occurs not because the reward signal is corrupted or the optimization is too aggressive, but because the model's own probability estimates for the target domain are numerically unreliable. This has not been systematically characterized in prior work on domain-specific LLM adaptation.

The implication is significant: any domain where the target output distribution is severely underrepresented in pretraining—shader languages, hardware description languages, scientific computing DSLs, legacy system configuration formats—will face this same barrier. The paper thus contributes a general diagnostic: when adapting an LLM to a specialized programming domain via RL, check whether token probabilities are near the precision floor, and if so, expect PPO to be unstable until this is addressed.

The solution is principled, not ad hoc. The three-stage warm-up strategy (single-turn RL → RFT → Value Pretraining) addresses distinct aspects of the problem. Single-turn RL raises token probabilities out of the danger zone by providing initial gradient signal on the target domain distribution—this is not just a "warm start" but a numerical conditioning step. RFT provides a behavioral prior that constrains entropy growth (Figure 4b documents the catastrophic entropy spike without it). Value Pretraining prevents the exploration pathologies—excessively long trajectories, fruitless search—that arise from an uninitialized critic (Figure 5b). Each stage targets a specific symptom of the underlying precision-floor collapse, and the ablation results in Table 2 demonstrate that removing any one stage causes either training collapse or substantial performance degradation.

A negative result with positive implications. The paper's candid reporting of the 17-step collapse is itself a contribution—it establishes a boundary condition for when standard RL fine-tuning is expected to work and when specialized initialization is required. This is the kind of negative result that saves other researchers from replicating the failure, and it provides a concrete benchmark (17 steps before collapse, 150 steps with the full warm-up) against which alternative solutions can be measured.

Innovation 3: Replacing Continuous Speedup Rewards with a Discrete Milestone-Based Schedule as a General Principle for Performance Optimization RL

The paper's robust reward schedule (Equation 1) appears at first glance to be a minor implementation choice—discretize the continuous speedup ratio into four levels. But the empirical and conceptual analysis reveals it as a fundamentally different approach to shaping optimization objectives, with implications that extend well beyond CUDA.

The problem with continuous speedup rewards. The natural objective in performance optimization RL is to maximize speedup: $r = t_{\text{baseline}} / t_{\text{generated}}$ for correct solutions, a penalty for incorrect ones. This seems intuitive—faster kernels get higher rewards, creating a smooth gradient that should guide the policy toward better implementations. But the paper identifies two failure modes:

Operator difficulty bias: Speedup ratios vary dramatically across operators based on how well-optimized the baseline is. A naive PyTorch implementation of a trivial operation might yield 50× speedup from basic fusion; a heavily optimized operation might max out at 1.2×. The continuous reward makes the former worth 25× more than the latter, even though both represent the best achievable optimization. This biases the policy toward easy targets and creates high variance in the reward signal.
Measurement noise amplification: Nanosecond-level timing fluctuations—from GPU clock variations, memory controller contention, thermal throttling—add noise to the speedup ratio, which the continuous reward transmits directly to the policy gradient. The result is that the model receives conflicting signals: the same kernel might yield $r = 2.1$ on one run and $r = 1.9$ on another, with no change in code quality.

Why discretization works. The discrete schedule $\{-1, 1, 2, 3\}$ solves both problems simultaneously. By mapping speedup to binary thresholds ( $>5\%$ improvement over each baseline), it normalizes away operator difficulty—achieving level 3 is equally rewarded regardless of whether the absolute speedup is 1.1× or 73×. By thresholding at 5%, it eliminates the impact of sub-threshold measurement noise—a timing fluctuation from 1.04× to 1.06× speedup translates to the same reward (level 2 or 3, depending on the compile baseline) rather than a 2% reward difference.

The empirical evidence in Table 2 validates this design choice decisively: replacing the discrete schedule with a continuous speedup reward ("w/o Robust Reward") causes the geometric mean speedup to drop from 2.11× to 1.25× over torch.compile, with the faster rate falling from 96.8% to 60.4%. This is not a marginal difference—it represents a qualitative change in the policy's optimization behavior, from consistently outperforming the compiler to mostly failing to do so.

Why this is a conceptual contribution, not just a hyperparameter choice. The paper's analysis reveals that the discrete schedule is not merely thresholding a continuous signal—it is changing the optimization landscape from regression (predict the speedup ratio) to classification (achieve discrete performance milestones). This shift has several properties that make it particularly well-suited to performance optimization:

Curriculum structure: The four levels (incorrect → correct-but-slow → faster-than-eager → faster-than-compile) form a natural curriculum. Early in training, the model receives positive reinforcement for achieving level 1 (correctness), which prevents reward sparsity. As training progresses, the policy is pushed toward levels 2 and 3. A continuous reward would penalize correct-but-slow kernels almost as much as incorrect ones (since their speedup ratio is near 1.0), providing no intermediate learning signal.
Robustness to degenerate optima: A continuous speedup reward incentivizes the policy to find any transformation that increases the ratio, including ones that exploit measurement artifacts (e.g., timing before GPU warm-up) or that produce illegitimately "fast" kernels (e.g., returning constant outputs). The discrete thresholds require the kernel to actually pass correctness verification AND achieve a meaningful speed improvement, closing these reward-hacking pathways.
Alignment with practical goals: In deployment, what matters is whether a kernel is faster than the compiler baseline—not whether it's 1.7× or 1.9× faster. The discrete schedule directly optimizes for this binary outcome (level 3 = beats compile), making the training objective aligned with the evaluation metric.

This finding generalizes: any RL task where the natural continuous reward (speedup, cost reduction, latency improvement) exhibits high variance due to instance difficulty or measurement noise can benefit from discretizing into meaningful milestones. The paper's ablation provides strong evidence that this is not a minor tuning detail but a first-order determinant of optimization success.

Innovation 4: Demonstrating That Learned Optimization Policies Can Consistently Outperform Static Compiler Heuristics Through Autonomous Strategy Discovery

While the three previous innovations focus on training methodology and conceptual reframing, this fourth innovation concerns the empirical finding that the resulting system achieves something qualitatively different from prior approaches: it does not just match or slightly exceed compiler performance, but consistently and substantially outperforms static compiler heuristics across diverse operator types and difficulty levels, through strategies that the model discovered autonomously rather than through human-designed optimization recipes.

The magnitude and consistency of the result. Table 1 reports that CUDA Agent achieves a 100% faster rate—meaning every single correct kernel is faster than torch.compile—on KernelBench Level 2 (operator sequences), with a 2.80× geometric mean speedup. This is not a marginal improvement over the strongest proprietary models (Claude Opus 4.5 at 69% faster rate, 1.60× speedup) but a qualitative leap: the model has learned to optimize operator fusion problems so reliably that static compilation is never competitive. Even on Level 3—the hardest setting, comprising realistic neural network building blocks like ResNet blocks—CUDA Agent achieves 90% faster rate and 1.52× speedup, compared to Claude Opus 4.5 at 50% and Gemini 3 Pro at 52%. The gap widens as difficulty increases, suggesting the learned optimization strategies scale better with problem complexity than static compiler patterns.

What makes this more than just "model got better." The significance is not the absolute numbers but what they reveal about the nature of the learned capability. Static compilers like torch.compile apply predefined, rule-based optimization patterns—operator fusion templates, memory layout transformations, launch configuration heuristics—that were designed by compiler engineers. These patterns capture common optimization opportunities but are inherently limited: they cannot perform the algebraic reasoning that turns diagonal matrix multiplication into row-wise scaling (Case D.2), they cannot recognize that a specific sequence of matrix multiply → divide → sum → scale can be rearranged to pre-compute column-wise sums (Case D.3), and they cannot integrate high-level graph transformations (batch norm folding) with low-level hardware configuration (TF32 enablement) and custom kernel fusion (Case D.4).

CUDA Agent's case studies in Appendix D show that the trained agent discovers these cross-abstraction optimizations autonomously. The agent does not merely apply known patterns faster or more consistently—it synthesizes strategies that span multiple levels of the optimization stack, combining algebraic insight, library selection (cuBLAS vs. custom kernel vs. cuDNN fused APIs), hardware-specific configuration, and low-level kernel design. This is the kind of multi-level optimization reasoning that human experts do, and that static compilers fundamentally cannot do because they operate on a fixed intermediate representation without semantic understanding of the computation.

The significance for the compiler vs. learned optimization debate. There is an ongoing debate in the systems community about whether learned optimization policies can surpass decades of compiler engineering. The dominant assumption—implicit in the design of systems like torch.compile, TVM, and XLA—is that compiler heuristics, refined over years by domain experts, represent a strong ceiling that learned approaches will struggle to match. CUDA Agent provides the strongest evidence to date against this assumption in the GPU kernel domain: a learned policy, trained purely through RL with execution feedback and no compiler-derived supervision, can achieve a 100% faster rate on operator fusion tasks where static compilers are specifically designed to excel.

This does not mean compilers are obsolete—torch.compile remains the practical baseline for most users, and the paper acknowledges it does not compare against more sophisticated frameworks like TVM (Appendix E). But it challenges the framing that compilers represent an upper bound on automated optimization. The paper's results suggest instead that compilers and learned optimizers have complementary strengths: compilers excel at consistent, low-latency application of known patterns, while learned optimizers can discover patterns that compiler engineers haven't encoded. The 2.80× speedup on Level 2—where operator fusion is the primary challenge—is precisely where this complementarity is most pronounced.

A boundary condition: capability floors. The paper is careful to acknowledge that its approach has a fundamental limitation: it cannot create capability from nothing. On problems where the base model's understanding is simply insufficient—presumably the hardest Level 3 tasks where the geometric mean speedup, while still positive at 1.52×, is lower than Level 2's 2.80×—the RL training cannot manufacture new mathematical or algorithmic insights. This establishes a clear boundary: learned optimization amplifies existing capability but does not create it from nothing, much like the finding in the original example paper that test-time compute cannot help on problems where the base model's pass@1 is near zero. Recognizing this boundary is as important as the positive result itself, because it tells researchers where to invest effort—improving base model reasoning capabilities for the hardest mathematical transformations, rather than scaling RL training further.

5. Experimental Analysis

Evaluation Methodology

Dataset. All experiments use KernelBench (Ouyang et al., 2025), specifically the Level 1, Level 2, and Level 3 subsets comprising a total of 250 distinct operator tasks — 100 for Level 1 (single operators), 100 for Level 2 (operator sequences), and 50 for Level 3 (full model building blocks like ResNet blocks). The paper explicitly avoids using KernelBench for training, verifying separation through AST-based decontamination (no training sample exceeds 0.9 similarity to any test sample; Appendix A, Figure 7).
Base model. The base model is Seed1.6 (Bytedance Seed), a Mixture-of-Experts model with 23B active and 230B total parameters. The choice is motivated as "representative of the capabilities of many contemporary LLMs" (Section 4.1) with sufficient scale to benefit from RL — not so small that it lacks fundamental CUDA reasoning capability, not so large that training is prohibitively expensive.
Metrics. The paper reports three metrics, computed per-task and aggregated across subsets: (1) Pass Rate — the percentage of tasks where the agent generates a kernel that successfully compiles and passes functional correctness checks; (2) Faster Rate — the percentage of tasks where the generated kernel is both correct AND achieves faster execution than the specified baseline (Eager or torch.compile); and (3) Speed-up — the geometric mean of the execution speedup ratio relative to the baselines, computed exclusively for correct solutions (incorrect kernels are excluded from the geometric mean). For each task, the best-performing solution along the trajectory — the one achieving maximum speedup over torch.compile — is extracted for final metric computation. Overall metrics across all three levels are weighted by the number of problems in each level (Level 1: 100, Level 2: 100, Level 3: 50).
Baselines. Four models are used as baselines (Table 1): Claude Opus 4.5 (Anthropic, 2025), Gemini 3 Pro (Google DeepMind, 2025), GLM 4.6 (Zhipu AI, 2025), and Kimi K2 (Kimi Team, 2025). The first two represent the strongest proprietary coding models available; the latter two represent leading open-source coding models. All baselines are evaluated under the same agent loop as CUDA Agent — same tools, same SKILL.md, same sandbox environment — ensuring that performance differences reflect the model's learned CUDA optimization capability rather than the evaluation protocol. The paper notes that ChatGPT-5 series models were excluded because they "consistently declined to respond to CUDA-related prompts" (Section 1 footnote).
Generation budget / compute accounting. The primary unit of test-time compute is agent turns — each turn consists of the model producing a reasoning trace, selecting an action, and receiving an observation. During training, rollouts are capped at 150 turns; during evaluation, the cap is relaxed to 200 turns. The context window is 131,072 tokens during agentic RL (32,768 during single-turn RL warm-up). All models are evaluated under identical turn budgets and environment constraints, making the comparison fair. The paper does not report FLOP counts or wall-clock time — the turn count and context length serve as the budget proxy.
Cross-validation / statistical protocol. The paper does not employ cross-validation for the main results. The training set (CUDA-Agent-Ops-6K) and test set (KernelBench) are fixed with explicit decontamination. For ablation studies (Table 2), each leave-one-out variant is trained separately and evaluated on the full KernelBench test set. The paper notes that for variants without RFT or Value Pretraining — which exhibit training collapse — results are reported from "the final validation step before training collapse." No confidence intervals or statistical significance tests are reported for any metrics, which is a limitation.
Hardware specification. Training uses a pool of 128 NVIDIA H20 GPUs for sandbox evaluation. The CPU–GPU decoupled architecture places Docker-based terminal sandboxes on CPU for compilation and file operations, while verification and profiling jobs are dispatched to dedicated GPUs with process-level isolation. This architecture is described as necessary for "stable latency measurements and guaranteed HBM capacity" (Section 4.1). Evaluation hardware is unspecified, but the paper uses TORCH_CUDA_ARCH_LIST=9.0 throughout, indicating Hopper-architecture GPUs.

Main Quantitative Results

The paper reports results across three axes: (1) main benchmark comparison against baselines on KernelBench Levels 1–3, (2) ablation analysis of each architectural component, and (3) case studies of learned optimization strategies. There is no separate "search vs. revisions" axis (as in the original example paper) since the system is a unified agent; the ablation study in Table 2 serves as the primary decomposition.

Benchmark Results: CUDA Agent vs. Proprietary and Open-Source Baselines

Headline result (Table 1, overall metrics): CUDA Agent achieves a 98.8% Pass Rate, 96.8% Faster Rate over torch.compile, and 2.11× geometric mean speedup over torch.compile across all 250 KernelBench problems. The closest proprietary competitor, Claude Opus 4.5, achieves 95.2% Pass Rate, 66.4% Faster Rate, and 1.46× speedup over torch.compile. The gap in Faster Rate — 96.8% vs. 66.4% — means that approximately one-third of Claude's correct kernels are no faster than static compilation, while CUDA Agent's correct kernels are almost universally faster. The gap in geometric mean speedup — 2.11× vs. 1.46× — indicates that even when Claude produces faster kernels, the magnitude of improvement is substantially smaller.

Level-by-level breakdown (Table 1):

Level 1 (single operators, 100 problems): CUDA Agent achieves 100% Pass Rate, 97% Faster Rate over torch.compile, and 1.87× geometric mean speedup. Claude Opus 4.5 achieves 96% Pass Rate, 72% Faster Rate, and 1.54× speedup. The 25-percentage-point gap in Faster Rate (97% vs. 72%) indicates that even on the simplest benchmark problems — single operators with relatively well-understood optimization strategies — Claude leaves substantial performance on the table, failing to beat torch.compile on 28% of correct kernels. CUDA Agent produces correct kernels and beats the compiler in nearly all cases, with the remaining 3% gap attributed to "trivially simple" operators where no optimization is possible.
Level 2 (operator sequences, 100 problems): This is where CUDA Agent's advantage is most dramatic. CUDA Agent achieves 100% Pass Rate, 100% Faster Rate over torch.compile, and 2.80× geometric mean speedup over torch.compile — a perfect score on faster rate, meaning every single correct kernel outperforms the compiler. Claude Opus 4.5 achieves 98% Pass Rate but only 69% Faster Rate and 1.60× speedup. The 31-percentage-point gap in Faster Rate (100% vs. 69%) and the 1.75× ratio in geometric mean speedup (2.80 vs. 1.60) demonstrate that CUDA Agent has learned to handle operator fusion — the defining challenge of Level 2 — far more effectively than any baseline model. The perfect 100% faster rate implies the agent has internalized a general fusion strategy rather than succeeding only on specific operator combinations.
Level 3 (full model building blocks, 50 problems): The hardest subset. CUDA Agent achieves 94% Pass Rate, 90% Faster Rate over torch.compile, and 1.52× geometric mean speedup. Claude Opus 4.5 achieves 88% Pass Rate, 50% Faster Rate, and 1.10× speedup. The 40-percentage-point gap in Faster Rate (90% vs. 50%) is the largest in absolute terms across all levels, and the speedup gap (1.52× vs. 1.10× ratio of 1.38) shows that on the most challenging problems — realistic neural network components like ResNet blocks — CUDA Agent's learned optimization strategies are dramatically more effective than what general-purpose LLMs can achieve through the agent loop alone. However, the geometric mean speedup of 1.52× (vs. 2.80× on Level 2) also reveals the difficulty ceiling: complex models with multiple interacting operators leave less room for dramatic speedups, even for the trained agent.

Performance relative to Eager mode (Table 1): The Eager baseline comparisons show an even larger gap. CUDA Agent achieves a 2.60× geometric mean speedup over Eager overall, compared to Claude Opus 4.5 at 1.99×. On Level 2 specifically, CUDA Agent achieves 3.27× over Eager (vs. 2.24× for Claude) — indicating that the agent's fusion optimizations provide a 3.27× average speedup over naive PyTorch execution for operator sequences.

The gap between Pass Rate and Faster Rate reveals optimization capability, not just coding capability. Across all baselines, the Pass Rate is high (75–95%), but the Faster Rate is substantially lower (19–69%). This gap represents kernels that compile and run correctly but fail to achieve meaningful speed improvements — they are "naïve" implementations that essentially replicate what torch.compile already does automatically. The baselines' high Pass Rates confirm that general-purpose LLMs can write correct CUDA code, but their low Faster Rates confirm they lack optimization expertise. CUDA Agent nearly closes this gap: 98.8% Pass Rate vs. 96.8% Faster Rate over torch.compile, meaning only 2% of correct kernels fail to beat the compiler. This near-unity ratio is the strongest evidence that the RL training has specifically improved optimization capability rather than coding capability in general.

Speedup variability across difficulty levels. A notable pattern: CUDA Agent's geometric mean speedup over torch.compile decreases monotonically from Level 2 (2.80×) to Level 1 (1.87×) to Level 3 (1.52×). This is intuitively sensible — Level 2 problems (operator sequences) have the most room for fusion-based optimization, Level 1 problems (single operators) are already relatively well-optimized by torch.compile, and Level 3 problems (full model blocks) have complex data dependencies that constrain the speedup achievable through single-kernel fusion. The baseline models show a similar but less pronounced pattern, with their speedup advantages consistently lower at every level.

What the Baseline Comparison Tests (and Doesn't Test)

The experimental design compares CUDA Agent against strong proprietary models evaluated in an identical agent loop — same tools, same skill specification, same turn budget, same sandbox environment. This controls for the effect of the agent scaffold itself and isolates the difference in the models' learned CUDA optimization capabilities. The baselines therefore test: does a general-purpose LLM, given the same interactive development environment that CUDA Agent was trained in, achieve comparable kernel optimization performance? The answer is clearly no.

What this comparison does NOT isolate is whether CUDA Agent's advantage comes from (a) the base model's pretraining quality, (b) the single-turn RL warm-up, (c) the agentic RL training, or (d) the specific training data (CUDA-Agent-Ops-6K). Since the baselines are entirely different model families trained on different data, the comparison demonstrates overall system superiority but cannot attribute it to specific components. The ablation study in Table 2 partially addresses this by isolating the effect of removing individual training components (agent loop, reward design, RFT, Value Pretraining) from the full CUDA Agent system, but it does not test, for example, whether Claude Opus 4.5 fine-tuned on the same RL pipeline would outperform CUDA Agent — which would test whether the advantage is primarily architectural or primarily the RL training methodology.

Ablation Studies and Robustness Checks

Table 2 presents the central ablation results, comparing the full CUDA Agent against four leave-one-out variants under the same agent loop evaluation. Each ablation removes a single major component and tests the resulting system on the full KernelBench benchmark.

Impact of the Skill-Integrated Agent Loop (Table 2, "w/o Agent Loop"): This variant replaces the multi-turn agent interaction with single-turn code generation — the model produces the complete kernel and binding code in one response without execution feedback — but uses the same training data and RL algorithm as CUDA Agent. The results are stark: the Faster Rate over torch.compile collapses from 96.8% to 14.1%, and the geometric mean speedup over torch.compile drops from 2.11× to 0.69× — meaning the generated kernels are on average slower than what the static compiler produces automatically. The Pass Rate also drops from 98.8% to 77.1%, indicating that without iterative compilation feedback, the model cannot reliably produce even syntactically valid CUDA code. This result confirms that the interactive agent loop is not merely a helpful addition — it is essential for learning optimization. Without exposure to compilation errors, profiler output, and iterative refinement during training, the model learns only pattern-matching from problem to kernel, which fails comprehensively on a domain where correctness requires precise hardware specification.

Impact of Robust Reward Design (Table 2, "w/o Robust Reward"): This variant replaces the discrete milestone-based reward schedule (Equation 1) with a continuous speedup reward: $r_s = t_{\text{compile}} / t_{\text{gen}}$ for correct solutions and -1 for incorrect ones. The Faster Rate over torch.compile drops from 96.8% to 60.4%, and the geometric mean speedup over torch.compile drops from 2.11× to 1.25×. The Pass Rate remains comparable (96.8% vs. 96.8%), meaning the continuous reward produces kernels that are just as likely to be correct but substantially less likely to be fast. This is a non-obvious finding: the continuous reward, despite providing finer-grained optimization signal, worsens optimization outcomes. The paper attributes this to operator difficulty bias (easy-to-optimize operators dominate the reward signal) and measurement noise amplification (nanosecond-level timing variations add noise to the continuous reward that the thresholds filter out). The magnitude of the effect — a 36-percentage-point drop in Faster Rate — confirms that reward design is a first-order concern in performance optimization RL, not a minor hyperparameter detail.

Impact of Rejection Fine-Tuning (Table 2, "w/o RFT"): This variant skips the RFT-based actor initialization and proceeds directly from single-turn RL to agentic PPO. The reported results are from "the final validation step before training collapse." The geometric mean speedup over torch.compile drops to 1.05× (vs. 2.11× for the full system), and the Faster Rate drops to 49.8% (vs. 96.8%). These numbers understate the actual failure mode: Figure 4a shows the training reward collapsing catastrophically within approximately 10 steps, and Figure 4b shows the corresponding entropy spike as the policy becomes "increasingly diffuse, producing incoherent and poorly structured outputs." The ablation demonstrates that RFT is not optional — it is necessary for stable training. Without the behavioral prior provided by supervised fine-tuning on high-quality trajectories (outcome-filtered and pattern-filtered), PPO's exploration in the vast action space of 200-turn CUDA development trajectories leads to rapid entropy explosion and policy collapse.

Impact of Value Pretraining (Table 2, "w/o Value Pretraining"): This variant skips the critic initialization stage and starts PPO with a randomly initialized value network. Like the RFT ablation, results are reported from the final validation step before training collapse. The geometric mean speedup drops to 1.00× over torch.compile — essentially no improvement over the compiler — and the Faster Rate drops to 50.9%. Figure 5a shows that the explained variance of the value function (a standard diagnostic for critic quality) remains near zero throughout training, indicating the critic has learned essentially nothing about the value landscape of different optimization states. Figure 5b shows the consequence: the response length clipped ratio — the fraction of trajectories hitting the 150-turn limit — explodes, because the uninitialized critic "fails to penalize fruitless or redundant search paths" (Section 4.3.3). The agent explores endlessly, producing excessively long trajectories that are computationally expensive but yield no performance improvement. This ablation demonstrates that value pretraining is not merely helpful — without it, PPO's advantage estimates are too noisy to provide useful learning signal, and the agent enters a pathological exploration loop.

Interaction between ablations: The paper does not report combinatorial ablations (e.g., removing both RFT and Value Pretraining simultaneously, or removing the agent loop AND the robust reward). Given that each individual ablation either causes training collapse or severe performance degradation, it is reasonable to infer that combinations would be worse, but the absence of these experiments means we cannot assess whether the components provide partially redundant benefits or are strictly complementary. The paper also does not report ablations of the warm-up stage ordering (e.g., what if Value Pretraining precedes RFT?) or of the specific hyperparameter choices (the 5% threshold in the reward, the asymmetric clipping bounds, the GAE λ value), which are standard ablation dimensions in RL systems.

Decontamination validation (Appendix A, Figure 7): The paper verifies that training data does not overlap with evaluation data using AST-based similarity. The maximum similarity between any training sample and any evaluation sample is below 0.9 after filtering, with the distribution showing most training samples at very low similarity scores. This is a necessary robustness check given the history of data leakage in prior CUDA fine-tuning work (Kevin, CUDA-L1, ConCuR) and confirms that CUDA Agent's performance reflects generalization rather than memorization of KernelBench problems.

Case studies as qualitative evidence of learned strategies (Appendix D): Three detailed case studies analyze the agent's optimization trajectories on representative problems from each difficulty level. Case D.2 (Level 1: diagonal matrix multiplication) shows the agent discovering algebraic simplification from $O(N^2 M)$ matrix multiplication to $O(NM)$ row-wise scaling, achieving 73.31× speedup over torch.compile. Case D.3 (Level 2: matrix multiply, division, summation, scaling) shows the agent rearranging the computation algebraically to pre-compute column-wise sums, then fusing remaining operations into a single kernel with vectorized float4 loads and shared-memory tree reduction, achieving 24.04× speedup. Case D.4 (Level 3: ResNet BasicBlock) shows the agent combining batch norm folding, cuDNN API selection with bias+activation fusion, TF32 enablement, and a custom fused add-ReLU kernel, achieving 3.59× speedup. These cases provide concrete evidence that the learned strategies span multiple abstraction levels — algebraic, algorithmic, library, and hardware — which no single hand-designed optimization recipe or compiler pattern would prescribe. However, they are cherry-picked examples (the paper selects "representative" cases), and the distribution of optimization strategies across all 250 test problems is not characterized.

The base model's standalone performance (Table 1, "Seed1.6 (base model)"): The base model without any RL training achieves only 74.0% Pass Rate, 27.2% Faster Rate over torch.compile, and 0.69× geometric mean speedup over torch.compile — meaning on average, it produces kernels slower than the compiler. This provides the baseline against which all training gains are measured. The jump from 0.69× to 2.11× geometric mean speedup over torch.compile represents the total effect of the full training pipeline (single-turn RL + RFT + Value Pretraining + agentic PPO). The paper does not report intermediate checkpoints (e.g., after single-turn RL only, after RFT only) in the main results, so the incremental contribution of each training stage to the final benchmark performance cannot be precisely quantified from Table 1 alone; the ablation study in Table 2 addresses this indirectly by measuring what happens when each component is removed from the full system.

Critical Assessment

The experiments broadly support the paper's central claims, but with specific boundary conditions, methodological limitations, and missing comparisons that qualify the strength of the conclusions.

Claim 1: "CUDA Agent achieves state-of-the-art results on KernelBench, delivering 100%, 100%, and 92% faster rate over torch.compile on Level-1, Level-2, and Level-3." The data directly supports the numerical claims (Table 1). The claim is substantively stronger than a typical "state-of-the-art" claim because the evaluation protocol is fair — all baselines are evaluated under the identical agent loop with identical tools and turn budgets — and the margin is substantial (96.8% vs. 66.4% overall Faster Rate, 2.11× vs. 1.46× geometric mean speedup). The claim would be further strengthened by: (a) confidence intervals on the metrics (the test set is only 250 problems, and variance across problems is not characterized; Level 3 has only 50 problems, making the 90% Faster Rate estimate relatively noisy); (b) evaluation on multiple GPU architectures (all results are on Hopper GPUs with TORCH_CUDA_ARCH_LIST=9.0 — it's unclear whether the learned optimizations transfer to other architectures or are Hopper-specific); and (c) comparison against additional compiler baselines (not just torch.compile; the paper acknowledges in Appendix E that TVM comparison is missing, which is a genuine gap — TVM's auto-tuning capabilities could provide a stronger compiler baseline than torch.compile's static patterns, though the paper's justification about "substantial tuning overhead and complex deployment requirements" is reasonable for the RL training setting).

Claim 2: "Outperforming the strongest proprietary models Claude Opus 4.5 and Gemini 3 Pro by about 40% on the hardest Level-3 setting." This claim is numerically accurate (90% vs. 50% Faster Rate on Level 3 for Claude, 90% vs. 52% for Gemini — differences of 40 and 38 percentage points respectively; Table 1). However, "about 40%" uses percentage-point differences rather than relative improvement, which can be misleading. The relative improvement in Faster Rate is 80% (from 50% to 90%) for Claude and 73% (from 52% to 90%) for Gemini — deploying the "about 40%" framing without qualification could be read as understating the effect. More importantly, the claim compares a specifically trained system against general-purpose models — this is a comparison of system + training rather than model architecture. A fairer comparison would test whether Claude or Gemini, if fine-tuned on the same CUDA-Agent-Ops-6K dataset with the same RL pipeline, would match or exceed CUDA Agent. The paper does not attempt this (likely due to API access limitations for proprietary models), so we cannot distinguish whether CUDA Agent's advantage comes primarily from the training methodology or from the base model's specific pretraining characteristics.

Claim 3: "Establishing that learned optimization policies can consistently outperform static compiler heuristics, particularly on complex operator fusion tasks." This is the paper's most significant empirical claim, and it is well-supported by the Level 2 results: 100% Faster Rate with 2.80× geometric mean speedup over torch.compile on operator sequences (Table 1). The "consistently" qualifier is justified — a 100% faster rate means every single correct kernel beats the compiler, which is about as consistent as possible given the finite test set. However, three caveats qualify the generality of this claim:

First, the "consistently" only applies to the specific operator types represented in KernelBench Levels 1–3. KernelBench covers single operators, operator sequences, and model building blocks — a reasonable but not exhaustive sample of real-world kernel optimization challenges. Missing categories include: operators with dynamic shapes, sparse operations, attention mechanisms (which have received extensive hand-optimization in the literature), and custom operators for novel neural network architectures. It's possible that torch.compile performs better on operator types outside KernelBench's coverage, or that CUDA Agent's learned strategies don't transfer to these unseen categories.

Second, the paper does not compare against any learned optimization baseline OTHER than the general-purpose LLMs. A critical missing comparison is against a model fine-tuned with supervised learning on expert-written CUDA kernels (e.g., from NVIDIA's cuBLAS/cuDNN libraries or open-source optimized kernel collections). Would 6,000 synthetic RL training tasks outperform fine-tuning on, say, 500 expert-written kernel implementations? The paper's argument that expert CUDA code "is prohibitively expensive to produce at scale" (Section 3.1) is valid for the RL training paradigm, but it doesn't test whether that paradigm is actually more effective per unit of training data than supervised approaches.

Third, the claim's boundary condition is hinted at but not systematically analyzed. The speedup on Level 3 (1.52×) is substantially lower than on Level 2 (2.80×), and the paper's own difficulty analysis doesn't exist for the agentic RL setting — it's unclear whether the 6% of Level 3 problems that CUDA Agent fails on (10% of 50 = 5 problems where pass rate is 94%, plus some correct-but-slower kernels to account for the 90% faster rate) are consistently the hardest problems in a difficulty sense, or whether the failure pattern is random. A difficulty-bin analysis (analogous to the one in the original example paper) would have been informative for understanding when learned optimization policies reach their ceiling.

Claim 4 (implicit): "The multi-stage warm-up strategy (RFT + Value Pretraining) is necessary for stable training, and each component contributes independently." The ablation results (Table 2, Figures 4–5) strongly support the claim that removing either RFT or Value Pretraining causes training collapse. However, the claim of "independent contribution" is harder to assess without combinatorial ablations — we don't know if RFT alone (without Value Pretraining) could achieve stable training if other hyperparameters were adjusted, or whether Value Pretraining's benefit is primarily from providing better advantage estimates (which RFT might partially compensate for by producing better trajectories). The specific diagnostic mechanisms proposed — entropy explosion for RFT (Figure 4b) and trajectory length explosion for Value Pretraining (Figure 5b) — are consistent with the data but represent a single mechanistic interpretation; alternative explanations (e.g., RFT might primarily help by reducing the KL divergence from the base model rather than specifically controlling entropy) are not tested.

Methodological limitations that qualify all claims:

Single model family, single GPU architecture. All results are with Seed1.6 on Hopper GPUs. The claim that the approach "systematically improves the base model's CUDA coding and optimization abilities" (Section 1) would be stronger with at least one replication on a different base model (e.g., applying the same pipeline to a LLaMA or Qwen variant). The GPU architecture constraint is particularly important — CUDA optimization is highly architecture-specific (register file sizes, shared memory banks, tensor core throughput vary across generations), and a policy trained exclusively on Hopper GPUs may not transfer to Ampere, Ada, or future architectures.
No characterization of variance. The test set has 250 problems (only 50 for Level 3), and no confidence intervals, standard deviations, or bootstrap estimates are reported. The reported metrics are point estimates whose reliability is unknown — a 90% Faster Rate on 50 problems could be 45/50 or could represent a different underlying rate with wide confidence bounds. Given the paper's emphasis on the magnitude of improvement over baselines, statistical characterization is needed to assess whether observed differences are reliable.
Training data contamination control is based on AST similarity, not semantic similarity. Two operators could have low AST similarity but represent nearly identical computational patterns (e.g., a matrix multiplication followed by a ReLU vs. a matrix multiplication followed by a sigmoid — structurally different at the AST level but presenting nearly identical optimization challenges). AST-based decontamination catches literal code copying but may miss semantically similar problems, which could partially inflate CUDA Agent's performance on KernelBench if the training distribution overlaps with the test distribution in optimization-pattern space.
The ablation evaluation may be confounded by training collapse. For the "w/o RFT" and "w/o Value Pretraining" ablations, results are reported from "the final validation step before training collapse" — which means these are at different training steps with different degrees of optimization. The reported metrics may reflect an intermediate point that happens to perform best before collapse, rather than a stable trained model. This makes the ablation comparison less clean than a stable-training comparison would be — we're comparing a stable 150-step model against models at (presumably) 10–17 steps of training, which may understate what these variants could achieve if the instability were addressed through other means (e.g., different learning rates, different clipping parameters).
Missing infrastructure baselines. The paper's primary compiler baseline is torch.compile, which is a JIT compiler using static patterns. More advanced frameworks — TVM with auto-tuning, Triton kernels, or even hand-optimized kernels from libraries like FlashAttention — are not compared. The paper acknowledges the TVM gap (Appendix E) with a practical justification ("difficult to integrate into large-scale RL training"), but for the evaluation, where only inference occurs, integrating TVM as a per-problem baseline would have been feasible and would have provided a stronger test of the "outperforms compilers" claim.

Experiments that would have strengthened the paper:

Difficulty-bin analysis. Binning KernelBench problems by the base model's success rate or by torch.compile's relative performance (analogous to the original example paper's difficulty analysis) would reveal whether CUDA Agent's gains are concentrated in specific difficulty regimes or are uniform.
Scaling analysis. How does performance vary with the number of training problems (e.g., 1K vs. 3K vs. 6K)? How does it vary with the number of RL training steps (e.g., 50 vs. 100 vs. 150)? Understanding the scaling curves would inform whether further investment in data or training steps would yield further improvement.
Transfer to unseen GPU architectures. Evaluating the Hopper-trained policy on Ampere or Ada GPUs would test whether the learned optimizations are architecture-specific or general CUDA optimization strategies.
Supervised fine-tuning baseline. Training on a comparable amount of expert-written CUDA kernel data (if available) or on the CUDA-Agent-Ops-6K dataset with supervised learning (using expert-generated kernel implementations as targets, if they could be produced) would contextualize the RL approach's benefit relative to simpler training paradigms.
Cross-model replication. Applying the same training pipeline to a different base model (e.g., Qwen, DeepSeek-Coder) would test whether the findings are Seed1.6-specific or generalize to other model architectures and pretraining distributions.
Combinatorial ablations. Testing RFT + Value Pretraining interactions (both removed, each individually removed vs. both present) would clarify whether the components provide independent or partially redundant benefits.

6. Limitations and Trade-offs

Capability Floor: No Improvement on Problems Outside the Base Model's Reach

The assumption or constraint. CUDA Agent's RL training amplifies the base model's existing CUDA optimization capability but cannot create it from nothing. The paper implicitly assumes the base model possesses sufficient mathematical reasoning and CUDA programming knowledge for the RL training to build upon. This is explicit in the case studies (Appendix D), where the agent discovers algebraic simplifications, kernel fusion opportunities, and hardware-specific configurations — but only because the base model already understands these concepts at some level. The paper does not claim that RL can teach fundamentally new mathematical or algorithmic reasoning from scratch.

The consequence. The most direct evidence of this capability floor is the Level 3 performance: despite achieving 90% Faster Rate over torch.compile, the geometric mean speedup is only 1.52× (Table 1), substantially lower than Level 2's 2.80×. The remaining 10% of Level 3 problems where the agent cannot beat torch.compile — and the 6% where it cannot even produce a correct kernel (94% Pass Rate) — represent a hard ceiling. The paper does not characterize these failure cases, but the pattern is consistent with problems requiring optimization insights that the base Seed1.6 model simply lacks. Furthermore, the gap between the base model's standalone performance (27.2% Faster Rate over torch.compile, 0.69× speedup; Table 1) and CUDA Agent's final performance represents the total gain achievable through RL with this base model — a larger gap than training can bridge does not exist. For domains or operators where Seed1.6's initial capability is even lower than on KernelBench Level 3, the RL approach would likely fail entirely, since the model would never produce correct kernels to generate positive reward signals.

What evidence exists in the paper. The paper provides this evidence indirectly through the Level 3 metrics (Table 1) and the base model baseline, but does not conduct a difficulty-bin analysis that would directly measure the relationship between base model capability and RL improvement magnitude. Without such an analysis, we cannot estimate the capability threshold below which RL training becomes ineffective — we only know it exists somewhere below the Level 3 difficulty level. In the original example paper on test-time compute scaling, this limitation is explicitly characterized: "On the hardest questions (bin 5), no method makes meaningful progress — the base model simply lacks the capability to produce correct solutions regardless of how the budget is allocated." CUDA Agent lacks an equivalent characterization, making its capability boundary undocumented.

Mitigation status. The paper implicitly acknowledges this limitation through the Level 3 results but does not analyze it systematically. Section 5 ("Conclusion") frames RL as enabling models to "move beyond syntactic code generation toward hardware-aware performance optimization," but does not discuss what happens when the base model lacks the prerequisite hardware-awareness. The limitation is not addressed in the proposed future work.

Difficulty Estimation Cost: The Warm-Up Stages Require Massive Upfront Computation That Is Not Accounted for in Headline Results

The assumption or constraint. The multi-stage warm-up strategy — single-turn RL, RFT, and Value Pretraining — requires collecting full agent trajectories from the partially trained model, filtering them, and performing supervised fine-tuning before the main PPO training can even begin. This is an enormous computational investment that is entirely separate from the 150-step PPO training reported in the headline results. The paper does not quantify the cost of these warm-up stages in GPU-hours, FLOPs, or wall-clock time, nor does it amortize this cost into the reported training efficiency. A practitioner attempting to replicate this approach needs to budget for: (1) single-turn RL warm-up (training steps and GPU hours not reported), (2) trajectory collection by running the partially trained agent on all 6,000 training problems at up to 150 turns each, (3) RFT supervised fine-tuning on filtered trajectories, and (4) Value Pretraining on those same trajectories. Each stage requires significant GPU resources beyond what the 150-step PPO phase consumes.

The consequence. The headline numbers — 150 training steps, 128 NVIDIA H20 GPUs — understate the true computational cost by an unknown but likely substantial factor. For a practitioner deciding whether to adopt this method, the total cost matters: if the warm-up stages consume, say, 3× the compute of the main PPO phase, then the effective cost-per-unit-performance is correspondingly higher. More importantly, the warm-up stages create a minimum viable scale for this approach — you need enough GPU resources to complete all stages before seeing whether the approach works for your domain. This is a barrier to adoption for smaller research groups or for applications of the methodology to domains with different characteristics.

What evidence exists in the paper. The paper provides almost no quantitative evidence about warm-up costs. Section 3.3 describes the stages conceptually and provides ablation results showing they are necessary (Table 2, Figures 4-5), but the number of training steps, the trajectory collection cost, and the filter pass rate for RFT are not reported. The paper states that the base model's initial RL attempts "could only train stably for 17 steps before the model's performance collapsed" (Section 3.3), which gives a lower bound on the problem severity, but the cost of the solution is unquantified.

Mitigation status. Not addressed. The paper does not discuss warm-up cost, does not propose cheaper alternatives, and does not flag this as an area for future work. The acknowledgment in Appendix E that the training pipeline "relies on a large GPU pool with process-level isolation, which incurs considerable computational and engineering cost" is about the sandbox architecture, not the warm-up overhead specifically.

Single Model Family, Single GPU Architecture: Results May Not Transfer to Different Base Models or Hardware Generations

The assumption or constraint. All experiments use a single base model (Seed1.6, 23B active / 230B total parameters; Section 4.1) on a single GPU architecture (Hopper, via TORCH_CUDA_ARCH_LIST=9.0; Appendix B.2). The training data (CUDA-Agent-Ops-6K), the agent environment, and the RL algorithm are all designed around and tested with this specific combination. The paper does not evaluate on alternative base models (e.g., LLaMA, Qwen, DeepSeek-Coder) or alternative GPU architectures (Ampere, Ada, or consumer-grade GPUs).

The consequence. There are two distinct generalization risks. Model generalization: Seed1.6's pretraining data mixture, architecture (Mixture-of-Experts), and CUDA-specific knowledge may differ substantially from other model families. The paper's own diagnosis of the stability problem — that CUDA code constitutes less than 0.01% of pretraining data (Section 3.3) — applies universally, but the degree of underrepresentation and the model's resulting baseline CUDA capability likely vary. A model with even less CUDA exposure than Seed1.6 might fail to benefit from the same warm-up strategy; a model with more exposure might not need all three stages. Without cross-model validation, we cannot distinguish which aspects of CUDA Agent's success are specific to Seed1.6 versus general properties of the training methodology. Hardware generalization: CUDA optimization is notoriously architecture-specific. Register file sizes, shared memory bank configurations, tensor core throughput and precision (FP16 vs. TF32 vs. FP8), and memory bandwidth vary across GPU generations. A policy trained exclusively on Hopper GPUs has learned Hopper-specific optimization strategies — the case study in Appendix D.4 explicitly mentions enabling TF32 for Tensor Core utilization, which is Hopper-specific behavior. Whether the learned strategies transfer to Ampere GPUs (no TF32, different tensor core instructions) or consumer GPUs (smaller register files, fewer SMs) is unknown.

What evidence exists in the paper. None. The paper does not include any cross-model or cross-architecture experiments, and does not discuss these generalization dimensions in the limitations section (Appendix E discusses TVM as a missing compiler baseline and computational cost, but not model or hardware generalization).

Mitigation status. Not addressed. The paper's claims are implicitly bounded by the specific model and hardware used, but this boundary is not stated explicitly. A practitioner deploying this on LLaMA or on Ampere GPUs has no guidance on expected performance or required adaptations.

AST-Based Decontamination May Miss Semantic Overlap Between Training and Test Data

The assumption or constraint. The paper verifies that training data does not leak KernelBench evaluation problems using AST-based code similarity, removing any training sample with maximum similarity above 0.9 to any test sample (Appendix A, Figure 7). The implicit assumption is that AST similarity is a sufficient decontamination criterion — that problems with low AST similarity present genuinely different optimization challenges.

The consequence. AST structure can differ substantially while the underlying computational pattern is essentially identical. Consider two KernelBench Level 2 test problems: one that composes matmul → relu → layernorm and another that composes matmul → gelu → batchnorm. These would have distinct AST structures (different operator classes with different parameterizations) and would likely pass the 0.9 similarity threshold, but the optimization challenge — fusing a matrix multiplication with an element-wise activation and a normalization — is structurally identical. If the training set (CUDA-Agent-Ops-6K) contains compositions like matmul → silu → instancenorm, the learned fusion strategies would transfer almost perfectly to the test problems despite low AST similarity. This is not "data leakage" in the classical sense (the model isn't memorizing specific test cases), but it means the training and test distributions overlap in optimization-pattern space in ways that AST similarity cannot detect. The consequence is that CUDA Agent's performance on KernelBench may partially reflect the similarity between the synthetic training data's optimization patterns and the test set's patterns, rather than genuine generalization to novel optimization challenges. The 2.80× speedup on Level 2 — where operator fusion is the primary challenge and where the training data consists of 83.77% two-operator compositions (Table 3) — is particularly susceptible to this concern.

What evidence exists in the paper. The paper provides the AST similarity distribution (Figure 7) but no semantic similarity analysis. The training data composition (Table 3) shows a deliberate focus on 2-operator compositions (83.77% of the dataset), which aligns with KernelBench Level 2's operator sequence structure. The paper does not analyze whether training problems with the same "fusion pattern" as test problems disproportionately contribute to performance, and does not characterize the diversity of optimization patterns in CUDA-Agent-Ops-6K compared to KernelBench.

Mitigation status. The paper does not address this limitation. The AST decontamination is presented as sufficient (Appendix A: "we confirm that no training sample exhibits high AST-level similarity with the evaluation set"), and semantic overlap is not discussed. This is a gap because the paper's central claim — that CUDA Agent learns generalizable optimization strategies — would be weakened if a significant fraction of its advantage comes from pattern overlap between the synthetic training data and the specific KernelBench test set.

The 5% Speedup Threshold and Discrete Reward Schedule Are Not Ablated or Calibrated Against Alternative Designs

The assumption or constraint. The robust reward schedule (Equation 1, Section 3.2) uses a 5% threshold in the binary indicator $b(t, t_0) = \mathbb{I}[(t_0 - t) / t_0 > 5\%]$ to determine whether a speedup is "significant." This threshold is applied uniformly across all operators, all difficulty levels, and all baseline comparisons (Eager and Compile). The paper treats this 5% value as a fixed design choice and does not experiment with alternative thresholds (e.g., 1%, 10%, or an operator-specific threshold based on measurement variance).

The consequence. The 5% threshold has two effects that interact with the experimental results in unknown ways. First, it determines the granularity of the reward signal: kernels with 4.9% speedup receive reward level 1 or 2 instead of 2 or 3, even though the difference between 4.9% and 5.1% speedup may be within measurement noise. This introduces a cliff in the reward landscape that could either help (by thresholding out noise) or harm (by denying positive reinforcement for near-threshold improvements). Second, it creates a baseline dependence: torch.compile itself achieves varying speedups over Eager mode for different operator types. For an operator where torch.compile already achieves 20× speedup over Eager, achieving an additional 5% over torch.compile (reward level 3) is much harder than for an operator where torch.compile achieves only 1.1× speedup. The uniform 5% threshold does not account for this heterogeneity, potentially making reward level 3 much harder to achieve on well-compiled operators.

What evidence exists in the paper. The ablation in Table 2 ("w/o Robust Reward") demonstrates that the discrete schedule outperforms continuous speedup reward, but this ablation changes the entire reward structure (discrete vs. continuous) rather than the specific threshold value. No experiment tests whether the 5% threshold is near-optimal or whether performance is sensitive to this choice. The paper provides no measurement noise characterization — no data on run-to-run variance in the profiling pipeline, which would be necessary to assess whether 5% is above or below the noise floor.

Mitigation status. Not addressed. The reward design is presented as a finished choice, and threshold sensitivity analysis is absent from both the main experiments and the limitations section. This matters because the 5% threshold could be a consequential hyperparameter — too low, and measurement noise causes spurious reward assignments; too high, and the model receives insufficient positive reinforcement early in training. A practitioner replicating this method on a different domain or hardware platform would need to recalibrate this threshold without guidance from the paper.

No Comparison Against Supervised Fine-Tuning Baselines or Alternative RL Formulations

The assumption or constraint. CUDA Agent is evaluated exclusively against general-purpose LLMs (Claude Opus 4.5, Gemini 3 Pro, GLM 4.6, Kimi K2) in an identical agent loop, and against ablated versions of itself (Table 2). The paper does not compare against the most natural alternative training approaches: (1) supervised fine-tuning on expert-written CUDA kernels, (2) supervised fine-tuning on the CUDA-Agent-Ops-6K dataset using synthetically generated "expert" solutions (e.g., from a stronger model or from iterative refinement), or (3) alternative RL formulations such as best-of-N rejection sampling without the full PPO machinery. The paper's argument that "manual implementation of expert-level reference code is prohibitively expensive" (Section 3.1) justifies the choice of RL, but does not eliminate the need to compare against alternative training paradigms to establish that RL specifically — rather than simply having 6,000 training problems and a multi-turn agent loop — is responsible for the performance gains.

The consequence. We cannot determine how much of CUDA Agent's advantage over the baseline models comes from the RL training algorithm specifically, versus from (a) having access to 6,000 training problems tailored to the target difficulty distribution, (b) the agent loop scaffold itself (which the baselines also use, but which they were not trained to exploit), or (c) the warm-up stages independent of the PPO optimization. A supervised fine-tuning baseline — training on the CUDA-Agent-Ops-6K problems with expert-generated or strong-model-generated kernel solutions — would isolate the contribution of execution-driven RL versus demonstration-driven learning. If SFT on the same data achieved comparable performance, the paper's central claim about RL's unique value would be weakened. Conversely, if SFT substantially underperformed, the claim would be strengthened. Neither result exists.

What evidence exists in the paper. The ablation study (Table 2) compares the full system against degraded versions of itself, which tests the contribution of system components to the specific RL training pipeline but does not test whether RL is necessary versus alternative training paradigms. The base model's baseline performance (27.2% Faster Rate over torch.compile; Table 1) provides a lower bound, but the jump from 27.2% to 96.8% could theoretically be achieved through a combination of supervised fine-tuning on the synthetic data (using a strong teacher model's outputs) and the agent loop's iterative refinement at test time, without RL. The paper provides no evidence against this hypothesis.

Mitigation status. Not addressed. The paper does not discuss alternative training paradigms, does not include SFT baselines, and does not flag this as a limitation. This is a significant gap because it leaves the paper's central methodological claim — that agentic RL is the right approach for this problem — supported by elimination of prior approaches rather than by positive comparison against the most natural alternatives.

7. Implications and Future Directions

How This Work Changes the Landscape

This work establishes a new paradigm for adapting LLMs to specialized, performance-critical programming domains. Before CUDA Agent, the dominant approaches to AI-assisted CUDA kernel optimization fell into two camps: training-free orchestration systems that wrap a static LLM in hand-designed refinement workflows (STARK, CudaForge, EvoEngineer), and fine-tuning approaches that train on expert demonstrations or benchmark data (Kevin, CUDA-L1, ConCuR). Both implicitly treat the LLM as a pattern-retrieval engine—the model's role is to recall and adapt known optimization patterns from its training data. The fundamental limitation of both camps is the same: the model cannot exceed the optimization knowledge encoded in its training data or prompted by the workflow designer.

CUDA Agent breaks this ceiling by reframing the problem. Rather than treating the LLM as a pattern store, it treats the LLM as an autonomous systems optimizer that learns optimization strategies through trial-and-error interaction with a realistic execution environment. The critical shift is from demonstration-driven learning (the model imitates what humans have done) to outcome-driven learning (the model discovers what works through reward). This is not an incremental improvement—it is a qualitative change in what kind of capability the resulting system possesses. The case studies in Appendix D provide concrete evidence: the trained agent discovers algebraic simplifications (diagonal matrix multiplication → row-wise scaling, a 73.31× speedup), multi-kernel fusion with shared-memory tree reduction (24.04× speedup on a fused operator sequence), and cross-abstraction optimization combining batch norm folding, cuDNN API selection, TF32 enablement, and custom kernel fusion (3.59× on ResNet blocks). None of these strategies were explicitly taught—the agent discovered them because RL over execution feedback rewards outcomes, not adherence to prescribed patterns.

This reframing has three specific landscape-changing consequences:

First, it establishes that learned optimization policies can consistently outperform static compiler heuristics—and the gap is large, not marginal. On KernelBench Level 2 (operator sequences), CUDA Agent achieves a 100% faster rate over torch.compile with 2.80× geometric mean speedup, compared to Claude Opus 4.5 at 69% and 1.60×. The conventional wisdom in the systems community has been that compiler heuristics, refined over years by domain experts, represent a strong ceiling that learned approaches will struggle to approach. CUDA Agent provides the strongest counter-evidence to date: a learned policy, trained purely through interaction with hardware feedback and without compiler-derived supervision, can beat the compiler on every single problem in the operator fusion category—exactly where compilers are supposed to be strongest. This does not make compilers obsolete, but it reframes them from a ceiling to a baseline: the question is no longer whether learned optimization can match compilers, but on which problem classes and by what margin it can exceed them.

Second, it provides a general diagnostic and solution for the "precision-floor collapse" that arises when adapting LLMs to low-resource programming domains. The paper's identification of the root cause of training instability—token probabilities near the 16-bit floating-point representational floor creating numerically unstable PPO importance ratios—is a specific failure mode that generalizes beyond CUDA. Any domain where the target output distribution constitutes less than ~0.01% of pretraining tokens (shader languages, Verilog/VHDL, legacy system configuration formats, scientific computing DSLs) will face this barrier. The three-stage warm-up strategy (single-turn RL → RFT → Value Pretraining) is not merely an engineering fix but a principled prescription: raise token probabilities above the danger zone, anchor the actor with a behavioral prior from high-quality trajectories, and equip the critic with initial value estimates to prevent pathological exploration. Prior work on RL for LLMs had focused on reward hacking, distributional shift, and KL-constrained optimization as the primary failure modes; CUDA Agent adds a new failure mode to this taxonomy and provides both diagnostic evidence (the 17-step collapse, the entropy explosion in Figure 4b, the trajectory length explosion in Figure 5b) and a validated solution.

Third, it reconciles a seeming contradiction in prior results. Training-free systems like STARK and CudaForge achieved notable speedups over base models on some KernelBench problems, while pure fine-tuning approaches like Kevin reported gains but at limited scale. The contradiction was: if training-free refinement can get substantial gains, why invest in model training? And if fine-tuning helps, why do the reported gains cap out well below what a human expert could achieve? CUDA Agent's results resolve this: both approaches are limited by the base model's intrinsic optimization capability. Training-free refinement can surface latent knowledge more effectively than naive prompting, but cannot create knowledge the model doesn't have. Limited fine-tuning on benchmark data or expert demonstrations can teach specific patterns, but doesn't build the general optimization skill of autonomously exploring the design space through profiling and iterative refinement. CUDA Agent's agentic RL approach builds this skill directly, which is why the performance ceiling is dramatically higher—the model learns not just specific optimizations but a meta-strategy for discovering optimizations.

Research directions that become more attractive:

Learned optimization across other performance-critical domains. The ingredients—synthetic task generation, execution-based reward, multi-turn agent environments—are domain-agnostic. Database query optimization (reward: query latency), distributed systems configuration (reward: throughput or tail latency), compiler pass ordering (reward: generated code speed), and FPGA synthesis (reward: LUT/FF utilization and clock frequency) all share the structure that CUDA Agent exploits. The successful scaling to 6,000 training problems and 200-turn trajectories demonstrates that the approach is feasible at operationally relevant scale.
Combining learned optimization with compiler infrastructure. CUDA Agent currently treats torch.compile as a black-box baseline. A natural next step is to use RL to augment compiler optimization passes rather than replace them—the learned policy could propose transformations (fusions, tiling strategies, data layout choices) that the compiler then verifies, code-generates, and profiles, combining the exploration power of RL with the correctness guarantees and deployment infrastructure of mature compilers.
Improving verifier robustness and reward design for performance optimization. The paper's finding that a discrete milestone-based reward substantially outperforms continuous speedup (Table 2: 2.11× vs. 1.25× speedup) reveals that reward design is a first-order concern in performance optimization RL. Further work on adaptive thresholds, operator-specific reward shaping, and combining compiler-estimated speedup with measured speedup could yield substantial improvements.

Research directions that become less attractive:

Pure training-free refinement as a primary research direction for CUDA kernel generation. If a 150-step RL pipeline can produce models that achieve 100% faster rate on operator fusion (2.80× speedup) compared to training-free methods that rely on base model capabilities (Claude Opus 4.5 achieves 69% faster rate at 1.60× in the same agent loop), the marginal value of designing increasingly complex multi-agent orchestration systems for static models is substantially reduced. Training-free methods remain useful when model training is infeasible, but as a research direction toward the goal of matching or exceeding human expert performance, they are capped by the base model's knowledge. CUDA Agent shows that breaking this cap requires training.
Fine-tuning on benchmark-derived data as a publishable contribution. The paper's analysis in Appendix C explicitly documents data leakage issues in Kevin (trains on KernelBench subset), CUDA-L1 (constructs SFT data from KernelBench reference implementations), and ConCuR (trains on Kevin-32B outputs, propagating contamination). CUDA Agent's rigorous decontamination (AST similarity < 0.9, independent data synthesis) establishes a standard that makes prior approaches with known benchmark contamination difficult to justify as evaluations of generalizable capability. The field should expect future work to either use independent training data or clearly label results as benchmark-specific.

Follow-Up Research This Work Enables

Difficulty-adaptive allocation of agentic RL compute. CUDA Agent trains a single policy for all problem difficulties, but the performance gap between Level 2 (2.80× speedup over torch.compile) and Level 3 (1.52×) suggests that training and inference compute could be allocated adaptively. The original example paper on test-time compute scaling demonstrated 4× efficiency gains by routing easy problems to one strategy and hard problems to another. The analogous experiment here would be: train a difficulty estimator on CUDA-Agent-Ops-6K (e.g., using the base model's pass rate or profiling data as a difficulty signal), then train separate policies or allocate different turn budgets per difficulty bin during inference. A strong follow-up would measure whether difficulty-conditioned turn allocation—spending 50 turns on easy operators and 200 turns on hard ones—recovers the Level 3 performance gap while reducing total inference cost.

Cross-architecture transfer of learned optimization strategies. CUDA Agent is trained and evaluated exclusively on Hopper GPUs (TORCH_CUDA_ARCH_LIST=9.0), and the optimizations it discovers—TF32 tensor core usage, specific shared memory configurations, warp-level primitives—may be Hopper-specific. A critical stress-test would evaluate the trained policy on Ampere (SM 8.0/8.6) or Ada (SM 8.9) GPUs without retraining. Does the policy transfer? If it does, the learned strategies are architectural abstractions (coalescing, fusion, tiling) rather than hardware-specific tuning. If it doesn't, the approach requires per-architecture training, which changes the cost calculus substantially. A strong follow-up would run the exact KernelBench evaluation on an A100 (Ampere) or RTX 4090 (Ada) and measure the Pass Rate and Faster Rate. Interleaved training on multiple architectures during RL—where each rollout randomly samples a GPU target—would test whether the policy can learn to be architecture-aware (querying device properties and adapting its strategy) rather than architecture-specific (implicitly hard-coding Hopper assumptions). The case study in Appendix D.4 explicitly shows the agent enabling TF32, which is Hopper-specific—an Ampere evaluation would reveal whether the agent can suppress Hopper-specific optimizations when they're unavailable.

Scaling laws for synthetic training data: how many operator compositions are enough? The paper trains on 6,000 synthesized problems with a specific composition distribution: 83.77% two-operator compositions, and roughly 12% compositions of three or more operators (Table 3). Is this distribution near-optimal? A follow-up could train CUDA Agent variants on dataset sizes of 1K, 3K, 6K, and 12K problems, each holding the composition distribution constant, and measure the scaling curve of KernelBench performance versus dataset size. If the curve has not saturated at 6K, further data synthesis would yield predictable gains. Orthogonally, varying the composition depth distribution—e.g., training primarily on 3-operator compositions versus the current 2-operator-dominated mix—would test whether training on harder problems (more operators per composition) improves transfer to the hardest KernelBench Level 3 problems, which are the current performance ceiling. The paper's finding that Level 3 speedup (1.52×) lags behind Level 2 (2.80×) makes this a high-priority follow-up. A strong experiment would train three variants with composition distributions skewed toward 1-op, 2-op, and 3-op respectively, and measure Level 3 Faster Rate for each.

Supervised fine-tuning baseline with synthetic expert demonstrations. The paper argues that RL is necessary because expert CUDA code is scarce, but doesn't test whether a strong teacher model (e.g., Claude Opus 4.5 in the agent loop, generating high-quality trajectories that are then verified and filtered) could produce training data for SFT that rivals RL performance. A controlled experiment would: (1) run Claude Opus 4.5 in the CUDA Agent environment on the 6,000 training problems with 200 turns and relaxed filtering, (2) extract successful trajectories (reward level 3, fewer than 50 turns, no tool-call violations), (3) SFT-fine-tune Seed1.6 on these trajectories using the same RFT objective (Equation 2), and (4) evaluate the resulting model on KernelBench under the identical agent loop. This would isolate the contribution of the RL optimization algorithm versus the quality of the training trajectories. If SFT on Claude trajectories matches or approaches CUDA Agent's performance, it suggests that the primary value of the system is in trajectory quality (exploration guided by execution feedback) rather than in the PPO optimization specifically—a finding that would substantially simplify deployment for practitioners who want to avoid RL's complexity.

Dynamic reward thresholding based on measurement noise characterization. The robust reward schedule (Equation 1) uses a fixed 5% threshold to determine "significant" speedup. This threshold is uncalibrated against measurement noise—if the profiling pipeline has ±3% run-to-run variance on certain operator types, the 5% threshold may be near the noise floor, causing spurious reward assignments. A methodological follow-up would: (1) characterize the profiling pipeline's noise distribution per operator type by running each KernelBench operator through the profiling script 100 times and computing the coefficient of variation, (2) set per-operator thresholds at 2–3 standard deviations above the noise floor, and (3) compare this adaptive thresholding against the fixed 5% threshold in a full RL training run. If adaptive thresholding improves the Level 3 Faster Rate (currently 90%, with some correct-but-not-fast-enough kernels), it would demonstrate that reward design for performance optimization should be measurement-noise-aware. This is a relatively low-cost experiment (no new training data or model architecture needed) with potentially high impact on the final performance ceiling.

Combining learned optimization with classical compiler auto-tuning. CUDA Agent currently operates independently of torch.compile—it either beats the compiler or doesn't, with no intermediate relationship. A more sophisticated integration would have the learned policy propose optimization transformations (fusions, tiling configurations, data layouts) and use the compiler to validate and code-generate them, combining the exploration capability of RL with the correctness guarantees and code generation quality of a mature compiler backend. A concrete experiment: on KernelBench Level 3, after CUDA Agent proposes a set of transformations (as documented in the Appendix D.4 case study), compare the agent's hand-written CUDA kernel against a TVM-generated kernel implementing the same transformations (same fusion, same tiling, same data layout). If TVM's auto-tuning can further improve the agent's proposed strategy, this establishes a hybrid workflow where the agent proposes the high-level optimization plan and the compiler handles low-level code generation and auto-tuning. This could close the 1.52× speedup gap on Level 3 by leveraging compiler strengths (auto-tuning, architecture-specific code generation) that the agent currently bypasses by writing raw CUDA.

Practical Applications and Downstream Use Cases

Automated kernel optimization for in-house deep learning deployments. Organizations deploying custom neural network architectures at scale—for recommendation systems, domain-specific vision models, or scientific computing—routinely encounter operator combinations (custom attention variants, fused activation functions, specialized normalization layers) that torch.compile cannot optimize beyond basic fusion and layout transformations. The current workflow requires scarce CUDA experts to hand-write and profile custom kernels, with turnaround times of days to weeks per operator. CUDA Agent—once trained on the organization's target GPU architecture—could automate this pipeline, processing 6,000+ operator compositions with 200-turn optimization trajectories, achieving the 100% faster rate on operator fusion demonstrated on Level 2 (Table 1). At the 2.80× geometric mean speedup observed, a recommendation model training pipeline consuming 10,000 GPU-hours could be reduced to approximately 3,570 GPU-hours, saving thousands of GPU-hours per training run and reducing hardware provisioning costs. The paper's anti-hacking protections (file permission controls, context manager enforcement, multi-input validation) make this suitable for production environments where kernel correctness is non-negotiable—the 98.8% Pass Rate (Table 1) means fewer than 2% of generated kernels fail correctness checks, and those can be routed to human review.

Data generation for self-improving kernel optimization pipelines. The paper's training methodology itself enables a bootstrapping loop: deploy CUDA Agent to generate optimized kernels for a large corpus of operator compositions, filter for kernels with reward level 3 (correctness + speedup over both Eager and Compile), and use these trajectories as training data for a next-generation model. This is directly analogous to the self-improvement pipelines the original example paper envisions for test-time compute, but applied to performance optimization rather than reasoning. The key enabler is that CUDA Agent produces not just correct kernels but optimization trajectories—sequences of profiling, code generation, compilation, debugging, and refinement—that encode the strategy, not just the outcome. A next-generation model could be trained on these trajectories using the same RFT objective (Equation 2), potentially learning more sophisticated strategies than the current model discovered. The 6,000-problem training set could be expanded to 50,000+ compositions through combinatorial synthesis from additional libraries (e.g., torchvision, timm, diffusers), creating a large-scale dataset of optimization trajectories that bootstraps progressively more capable kernel agents.

On-premise deployment of optimized LLM inference with custom operators. Organizations running LLM inference on-premise frequently need to optimize model architectures for their specific hardware configuration—consumer GPUs, edge devices, or older datacenter GPUs—where vendor-optimized libraries (cuBLAS, cuDNN) may not provide optimal implementations for custom attention mechanisms, quantization schemes, or activation functions. CUDA Agent's methodology could be adapted to train an optimization agent specifically for the target hardware, using the same synthetic data pipeline (crawling seed operators, combinatorial synthesis, rubric-based filtering) but with profiling and verification scripts calibrated to the target GPU's specifications. The paper's architecture-specific training on Hopper GPUs (SM 9.0) sets a precedent: retraining on the target architecture with appropriate TORCH_CUDA_ARCH_LIST settings. For an organization deploying inference on Ampere GPUs (A100, A6000), a CUDA Agent trained with TORCH_CUDA_ARCH_LIST=8.0 could optimize custom operators for that specific hardware, potentially achieving speedups comparable to the 2.80× observed on Hopper Level 2 (Table 1) but for the target architecture's memory bandwidth and tensor core specifications.