CUDA-L1: Improving CUDA Optimization via  Contrastive Reinforcement Learning
Authors: Xiaoya Li
Xiaofei Sun
Albert Wang
Jiwei Li
Chris Shum
Venue: ICLR 2026 Poster (2026)
Keywords: CUDA Optimization
Reinforcement Learning
LLMs
Primary area: applications to computer vision, audio, language, and other modalities

Abstract: The exponential growth in demand for GPU computing resources has created an urgent need for automated CUDA optimization strategies. 
While recent advances in LLMs show promise for code generation, current state-of-the-art models achieve low success rates in improving CUDA speed. In this paper, we introduce CUDA-L1, an automated reinforcement learning (RL) framework for CUDA optimization that employs a novel contrastive RL algorithm. 

CUDA-L1 achieves significant performance improvements on the CUDA optimization task: trained on NVIDIA A100, it delivers an average speedup of {\bf ×3.12} with a median speedup of {\bf ×1.42} against default baselines over across all 250 CUDA kernels of KernelBench, with peak speedups reaching {\bf ×120}. In addition to the default baseline provided by KernelBench, CUDA-L1 demonstrates {\bf ×2.77} over Torch Compile, {\bf ×2.88} over Torch Compile with reduce overhead, and {\bf ×2.81} over CUDA Graph implementations. Furthermore, the model also demonstrates portability across GPU architectures, achieving average speedups of {\bf ×3.85} (median {\bf ×1.32}) on H100, {\bf ×3.13} (median {\bf ×1.31}) on L40, {\bf ×2.51} (median {\bf ×1.18}) on RTX 3090, and {\bf ×2.38} (median {\bf ×1.34}) on H20 despite being optimized specifically for A100. 

Beyond these benchmark results, CUDA-L1 demonstrates several properties: CUDA-L1 1) discovers a variety of CUDA optimization techniques and learns to combine them strategically to achieve optimal performance; 2) uncovers fundamental principles of CUDA optimization, such as the multiplicative nature of optimizations; 3) identifies non-obvious performance bottlenecks and rejects seemingly beneficial optimizations that actually harm performance.
The capabilities demonstrate that, RL can transform an initially poor-performing LLM into an effective CUDA optimizer through speedup-based reward signals alone, without human expertise or domain knowledge. 
In this process, it identifies CUDA optimization patterns, discovers new techniques, synthesizes them to achieve speedups, and more importantly, 
extends the acquired reasoning abilities to new kernels.
This paradigm opens possibilities for automated optimization of CUDA operations, and holds promise to substantially promote GPU efficiency and alleviate the rising pressure on GPU computing resources.

Published as a conference paper at ICLR 2026
CUDA-LI: IMPROVING CUDA OPTIMIZATION VIA CONTRASTIVE REINFORCEMENT LEARNING
Xiaoya Li, Xiaofei Sun, Albert Wang, Jiwei Li and Chris Shum
DeepReinforce Team
Q gittob.co//ccpreeooorre///DDDDD
ABSTRACT
The exponential growth in demand for GPU computing resources has created urgent need for automated CUDA optimization strategies. While recent ad- vances in LLMs show promise for code generation, current state-of-the-ar models achieve low success rates in improving CUDA speed. In this paper, we introduce CUDA-LI, an automated reinforcement learning (RL) framework for CUDA op- timization that employs a novel contrastive RL algorithm.
CUDA-LI achieves significant performance improvements on the CUDA opti- mization task: trained NVIDIA A100, it delivers average speedup of x3.12 with a median speedup of 1.42 against default baselines over across all 250 CUDA kernels of KernelBench, with peak speedups reaching x 120. In addition to the default baseline provided by KernelBench, CUDA-LI demonstrates x x2.77 over Torch Compile, x2.88 over Torch Compile with reduce overhead, and x2.81 over CUDA Graph implementations. Furthermore, the model also demonstrates portability across GPU architectures. CUDA-LI opens possibilities for automated optimization of CUDA operations, and holds promise to substantially promote GPU efficiency and alleviate the rising pressure on GPU computing resources. b
INTRODUCTION
The exponential growth in demand for GPU computing resources, driven primarily by the rapid advancement and deployment Large Language Models (LLMs), has created an urgent need for highly efficient CUDA optimization strategies. Traditionally, CUDA optimization has been highly manual and time-intensive process, where skilled engineers must meticulously analyze memory access patterns, experiment with different thread block configurations, and iteratively profile their code through extensive trial-and-error cycles.
Recent advances in LLMs (Team et al.) 2023 Shengyu et al., 2023 Team al. 2024 Grattafiori et al. 2024; Yang et al. 2025 Hurst et al.
et al. 2024 Li er al. 2024a OLMo et al. 2024 especially those powered with RL (Jacch et al.) 2024. Guo et al. 2025 Wang et al. 2024
Muennighoff et al., (025). have demonstrated remarkable capabilities in code generation and al- gorithm design. Despite the promise, current performance remains limited. State-of-the-ar LLM models such as DeepSeek-RI (Guo et 2025) and OpenAl-ol (Jaech et al.) (024) achieve low suc- cess rates in generating optimized CUDA code (only approximately 15% on KernelBench (Ouyang et al. 2025)) To address these limitations and unlock the potential of LLMs for automated CUDA optimization, this work, we propose CUDA-LI, an LLM framework powered by contrastive rein- forcement learning for CUDA optimization. CUDA-LI is pipelined framework, the core of which a newly-designed contrastive RL framework.
Different from previous RL models (William 1992: Shao et al. 2024, Schulman et al. 2017)
trastive RL performs comparative analysis of previously generated CUDA variants alongside their execution performance, enabling the model to improving through distinguishing between effective and ineffective optimization strategies. Contrastive-RL simultaneously optimizes the foundation
Email: (xiaoya]ii, xiaofei_ sun, albert_ wang, jiwei_li, chris shumm_ @deep-reinnorrce.cmm

---

Published as a conference paper at ICLR 2026
Stage SFT via Data Augmentation
Stage 2: supevvuue Learning
Taah Daveridoe
LLM
Self SASEErAAio widd LIM
Seed Data
LLM Generated
Data Synthesis
LLM Generated
Code Collection
Code Collection
Iterative Fine turing
Stage 3: Constrastive
Gradient Update
LUM Demenated
Srae to cmmematioo
Fonere
0.219
Exemplars Buckets
Constrative RL
Policy Network 71
Database
Sample bucket based on scores
Instruction Prompt
Normalized Reward Scores
LLM Generated Cuda Code
and
Update Exemplar Buckets
Figure 1: Overview of the CUDA-LI training pipeline.
model through gradient-based parameter updates while fulfilling the maximum potential from the current model through contrastive analysis from high-performance CUDA variants, creating evolutionary dynamic that drives superior CUDA optimization performance.
CUDA-LI delivers significant improvements on the CUDA optimization task: trained NVIDIA A 100, achieves an average speedup of x3.12 (median x 1.44) over the default baseline across all 250 KerneBBench CUDA kernels, with maximum speedups reaching 1200 In addition, CUDA- LI demonstrates x x2.77 over Torch Compile, x2.88 over Torch Compile with reduce overhead, x2.1 over CUDA Graph implementations. Furthermore, the CUDA codes optimized specifically for A100 demonstrate strong portability across GPU architectures, with similar optimization patterns observed across different baseline configurations: achieving average speedups of x3.85 (median x 1.32) on H100, x3.13 (median x 1.31) on L40, x2.51 (median x 1.18) on RTX 3090, and x2.38 (median x x1.34) on H20. Similar performance improvements over Torch Compile, Torch Compile with reduce overhead, CUDA Graph are consistently observed across all GPU types.
CUDA-LI reveals a remarkable capability of RL in autonomous learning for CUDA optimization:
1. Even starting with foundation model with poor CUDA optimization ability, by using code speedups as RL rewards and proper contrastive RL training techniques, we can still train RL system capable of generating CUDA optimization codes with significant speedups.
2. Without human prior knowledge, RL systems can independently discover CUDA optimiza- tion techniques, learn to combine them strategically, and more importantly, extend the ac- quired CUDA reasoning abilities to unseen kernels. This capability unlocks the potential for variety of automatic CUDA optimization tasks, e.g., kernel parameter tuning, mem- ory access pattern optimization, and different hardware adaptations, offering substantial promises to enhance GPU utilization.
Another contribution of this work is the enrichment of the KernelBench dataset with CUDA Graph implementations. Please refer to supplementary materials. We release these implementations to the community, providing substantially stronger baselines for performance comparison.
2

---

Published as a conference paper at ICLR 2026
2 CUDA-LI
2.1 OVERVIEW
Existing large language models (Guo et al. 2025 Yang et al. 2025 Grattation et al. 2024) demon- strate significant limitations in generating executable and correct C CUDA code with speedup, ported in prior research (Quyang et al.] 2025). This deficiency likely stems from the insufficient representation of CUDA code in the training datasets of these models. To address this fundamental gap, we introduce three-stage pipelined training strategy for CUDA-LI, i.e., Supervised fine-tuning via data augmentation, Self-supervised learning and Contrastive reinforcement learning, aiming to progressively enhances the model's 's CUDA programming capabilities:
Before we delve into the details of each stage. we provide key definitions adopted throughout the rest of this paper:
1. Executability: CUDA code executable if it successfully compiles, launches, and ex- ecutes to completion within 1000 the runtime of the reference implementation. Code exceeding this runtime threshold is considered unexecutable]
2. Correctness: A CUDA code is correct if it produces equivalent outputs to the reference
implementation across 1000 random test inputs P
3. Success: A CUDA code is successful if it is both executable and correct.
2.2 SFT VIA DATA AUGMENTATION
In the SFT stage, collect dataset by using existing LLMs to generate CUDA code snippets and selecting successful one. This dataset is directly used to fine-tune the model. Throughout this paper, deepseek-v3-677
(Liu et al. 2024a) as the model backbone. Please refer the details of data collection in Appendix A. The collected dataset is used to finetune the foundation model. The instruction the model the same the prompt for dataset generation, where the reference code is included in the instruction and the model is asked to generate an improved version. The model is trained to predict each token in di,j given the instruction.
2.3 SELF-SUPERVISED LEARNING
Now we are presented with the finetuned model after the SFT stage, where the model can potentially generate better CUDA code with higher success rates than the original model without finetuning. We wish further improve the model's s ability to generate successful CUDA code by exposing it to more code snippets generated by itself.
We achieve this iteratively by sampling CUDA code from the model, evaluating for executability and correctness, removing the unsuccessful trials and keeping the successful ones. Successful ones are batched and used to update the model parameters. Using the updated model, we repeat the process: generating code, evaluating it, and retraining the model. It is worth noting that during the self-supervvsed learning stage, we focus exclusively on the exccutability and correctness of the generated code, without considering speed
metric. This design choice reflects our primary objective of establishing reliable code generation before optimizing for performance.
2.4 CONTRASTIVE REINFORCEMENT LEARNING
Now we have a model capable of generating successful CUDA code at reasonable success rate. Next, we aim to optimize for execution speed.
One straightforward approach is to apply existing reinforcement learning algorithms such as RE- INFORCE (Williams) 1992). GRPO (Shao et al. 2024). or PPO (Schulman al.) 2017). In this approach, we would ask the model to first perform cfain-of-thought reasoning (Wei et al. 2022). then generate code, evaluate it, and use the evaluation score to update the model parameters. How- ever, our experiments reveal that these methods perform poorly in this task. The issue is as follows: standard RL algorithms compute a sealar reward for each generated CUDA code sample. During training, this reward undergoes algorithm-specife processing (e.g.. baseline subtraction in REIN- FORCE, advantage normalization in GRPO, importance sampling in PPO). The processed reward then serves loss weighting term for gradient updates, increasing the likelihood of high-reward
1 'This threshold is reasonable since code with 1000 slower performance contradicts our speedup optimiza- tion goals.
Prior work uses only 5 random inputs, which we found insufficient for robust validation.
3

---

Published as a conference paper at ICLR 2026
sequences while decreasing the likelihood of low-reward sequences. Critically, in this paradigm, the reward signal is used exclusively for parameter updates and never provided as input to the LLM. Consequently, the LLM cannot directly reason about performance trade-offs during code generation. To address this limitation, we propose incorporating reward information directly into the reasoning process by embedding performance feedback within the input prompt. Specifically, model with multiple code variants alongside their corresponding speedup scores. Rather than simply generating code, the LLM is trained first conduct comparative analysis of why certain implemen- tations achieve superior performance, then synthesize improved solutions based on these insights. Each generated code sample undergoes evaluation to obtain performance score, which serves dual purposes in our training framework: the primary purpose is to the score acts as reward signal for gradient-based parameter optimization, updating model weights. The score functions as a reward signal for gradient-based parameter optimization, directly updating the model weights; the other is to construct prompt for future training stages. The scored code sample becomes part of the exemplar set for subsequent training iterations, enriching the contrastive learning dataset.
present the
This dual-utilization strategy enables iterative optimization across two complementary dimensions:
Foundation Model Enhancement Parameter updates progressively improve the model's funda- mental understanding and capabilities for CUDA optimization tasks, expanding its representational capacity.
Solution Optimization The contrastive approach seeks extract the maximum potential from the current model's parameters by leveraging comparative analysis of high-quality exemplars.
These two optimization processes operate synergistically: enhanced foundation models enable more accurate contrastive reasoning, while improved reasoning strategies provide higher-quality training signals for parameter updates of foundation models. This co-evolutionary dynamic drives con- vergence toward optimal performance. We term this approach contrastive reinforcement learning (contrastive-RL for short).
2.4.1 PROMPT CONSTRUCTION
Here we describe the construction of prompts provided to the LLM. The prompt provided to the LLM during Contrastive-RL. training comprises the following structured components:
1) Task Descrpition: A detailed description of the computational problem, including in- put/output specifications, performance requirements, and optimization objectives.
III Previous Cuda Codes with Scores: Previously generated CUDA implementations paired with their corresponding performance scores (e.g., execution time, throughput, memory efficiency). providing concrete examples of varying solution quality.
III) Generation Protocol: Explicit instructions defining the required output format and com- ponents.
IV) Requirements and Restrictions: Requirements and restrictions to prevent reward hacking in RL.
The model's response must contain the following three structured components:
Performance Analysis: A comparative analysis identifying which previous kernel imple- mentations achieved superior performance scores and the underlying algorithmic or imple- mentation factors responsible for success.
II) Algorithm Design: A high-level description of the proposed optimization strategy, outlin- ing the key techniques to applied, presented as numbered points in natural language.
III) Code Implementation: The complete CUDA kernel implementation incorporating opti-
mizations.
A detailed demonstration for the prompt in shown in Table
2.4.2 CONTRASTIVE EXEMPLAR SELECTION
The selection of code exemplars for prompt construction is critical, as core of Contrastive-RL is to perform meaningful comparative analysis. The selection strategy needs to addresses the follow- ing two key requirements: first, for achieving competitive performance, the exemplar set should include higher-pertie perfforin implementations to guide the model toward competitive CUDA codes, avoiding local minima that result from analyzing and comparing inferior codes; second, to ensure
4

---

Published as a conference paper at ICLR 2026
performance diversity, the selected codes must exhibit substantial performance differences to enable effective contrastive analysis.
We employ sampling strategy akin that adopted by evolutionary LLM models: Let N denote the number of code exemplars included in each prompt (set to .N = 2 in our experiments). During RL training, we maintain performance-indexed database of all successful code samples generated during RL training. Codes are organized into performance buckets B₃ based on discretized score intervals, where bucket B; contains codes with scores range [sk, 8k + As).
We first sample N distinct buckets according to a temperature-scaled softmax distribution:
exp ((5. - -/w/TT exp ((5} 1*)/T)
P(B.) exp
where denotes the aggregate score of bucket B,, computed as the mean of its constituent code scores, μ4, mean(f5) )=1) represents the global mean of all bucket scores, and is the tempera- ture parameter governing the exploration-expllotttion tradeoff. The sampling strategy in Equation differs from conventional temperature sampling in evolutionary LLM approaches through a mod- ification: the deduction of μı stabilizes the distribution by centering scores around zero, which prevents absolute score magnitudes from dominating the selection.
From each selected bucket B;, we uniformly sample one representative code to construct the final prompt set. This approach satisfies both design criteria: Regarding competitive Performance, score- weighted bucket sampling biases selection toward perffor-nr implementations, ensuring the exemplar set contains competitive solutions; Regarding performance Diversity, enforcing selection from N distinct buckets ensures sufficient performance variance for effective contrastive analysis.
A more sophisticated alternative is to use an island-based approach for exemplar selection. However, we find no significant difference performance between our bucket-basee method and the island- based approach. Given this, we opt for the simpler bucket-based strategy.
2.4.3 REWARD
In this subsection, we detail the computation of the execution time-based reward function, which serves dual purposes: (1) guiding parameter updates in reinforcement learning and (2) constructing effective prompts. Given reference CUDA implementation qa from PyTorch with successful exe- cution time lg, and a generated code candidate d with execution time la we define the single-run speedup score as:
bg
single run (d)
td
More details for denoising the rewards are shown in Appendix A.6
For RL training, we adopt the Group Relative Policy Optimization (GRPO) strategy (Shao et al. Please refer to Appendix A.3 for details.
2024
EXPERIMENTS AND ANALYSIS
3.1
BERNELBENCH AND EVALUATION
Our evaluation is conducted on the KernelBench dataset (Ouyang et al. 2025). The KernelBench Dataset contains a collection of 250 PyTorch workloads designed to evaluate language models' abil- ity to generate efficient GPU kernels. The dataset is structured across three hierarchical levels based on computational complexity: Level contains 100 tasks with single primitive operations (such as convolutions, matrix multiplications, activations, and normalizations), Level 2 includes 100 tasks with operator sequences that can benefit from fusion optimizations (combining multiple operations like convolution + ReLU + bias), and Level 3 comprises 50 full ML architectures sourced from popular repositories including PyTorch, Hugging Face Transformers, and PyTorch Image Models (featuring models like AlexNet and MiniGPT). Each task in the dataset provides a reference Py- Torch implementation with standardized input/output specifications, enabling automated evaluation of both functional correctness and performance through wall-clock timing comparisons. The dataset represents real-world engineering challenges where successful kernel optimization directly trans- lates to practical performance improvements. Throughout this paper, evaluation benchmark. KernelBench recognized as a challenging benchmark the community (Ouyang et al. 2025), with even the best current L.LMs improving fewer than 20% of tasks.
KernelBench as the
5

---

Published as a conference paper at ICLR 2026
Successt
Speedupt >1.01x out of total
Configuration
Method
Mean
Max
75%
30%
25%
total
All
3.12
120
2.25
1.42
1.17
249/250
226/250
Level
2.78
65.8
1.75
1.28
1.12
99/100
80/100
Default
Level 2
3.55
120
2.05
1.39
1,20
100/100
98/100
Level 3
2,96
24,9
2.60
1.94
1.42
50/50
48/50
All
2.77
69.0
2.55
1.72
1.14
249/250
203/250
Level
3.04
59.7
2.71
1.99
1.41
99/100
89/100
Torch Compile
Level
2,91
69.0
1.99
1.55
1.10
100/100
78/100
3evel 3
1,98
8.57
2.28
1.68
1.00
50/50
36/50
All
2.88
80.1
2.48
1.67
1.13
249/250
200/250
Level
3.38
55.3
3.02
2.29
1.61
99/100
90/100
Torch Compile RD
Level 2
3.00
80.1
2.06
1.54
1.10
100/100
79/100
Level 3
1.62
8.67
1,76
1.13
0.991
50/50
31/50
All
2.81
97.9
1.83
1.20
0.954
249/250
147/229
1 Level
3.18
59.6
2.09
1.38
1.04
99/100
68/88
CUDA Graph
1 Level
2.84
97.9
1.55
1,08
0.932
100/100
53/94
Level
2.06
24.6
1.74
1.08
0.887
50/50
26/47
Performance comparison across different configurations on KernelBench on A100. RO Reduce Overhead. Success and Speedup indicate the number of successful benchmarks out of the total for each level. Note that for CUDA Graph, the total benchmark count differs from the dataset/data-subset size, some original reference code in KernelBench cannot be successfully transformed into the corresponding CUDA Graph implementations.
Table 1:
For each task with reference implementation g. we evaluate the performance of generated CUDA code using a similar protocol to training: We execute both q and in randomized order within fixed time budget of 20 minutes per task. The number of execution rounds varies across tasks due to differences in individual runtimes. The final evaluation score for d is computed as the av- erage speedup ratio across all execution rounds within the allocated time window. Unsuccessful implementations receive a score of zero. The metrics we report include speedup statistics (mean, maximum, and 75th, 50th, and 25th percentiles), success rate, and percentage of improvements.
3.2 COMPARISON SETUPS
To perform comprehensive evaluation on the generated code, we perform the following compar- isons:
Default
This compares the CUDA-LI generated code with the reference code by KernelBench.
II Torch Compile
This compares the CUDA-LI generated code with the reference code enhanced by compile with default settings. forch.comple applies graph-level optimizations including operator fusion, memory planning, and kernel selection to accelerate PyTorch models through just- in-time compilation.
III) Torch Compile Reduce Overhead This compares the CUDA-LI generated code with the ref- erence code enhanced by torch.compile with reduce-overhead mode enabled. This mode minimizes the compilation overhead by caching compiled graphs more aggressively and reducing recompila- tion frequency, making it particularly suitable for inference workloads with static shapes.
IV) CUDA Graph Since KernelBench does not provide official CUDA Graph implementations, we employ Claude 4 to generate CUDA Graph-augmented code for each reference implementation. CUDA Graphs capture a series of CUDA kernels and their dependencies into single graph structure that can be launched with minimal CPU overhead, eliminating the need for repeated kernl launch commands and significantly reducing CPU-GPU synchronization costs. Specifically, we provide Claude 4 with the reference code and request the addition of CUDA Graph optimizations. The gen- erated output is then evaluated for correctness. the code fails validation, we iterate by providing Claude 4 with both the original reference code and the previous erroneous outputs, requesting rected version. This iterative process continues for up to 10 attempts until the generated code passes all correctness checks. We release the CUDA Graph codes for KernelBench to the community, pro-

---

Published as a conference paper at ICLR 2026
Success
Speedupt
Methods
Model
Mean
Max
75%
50%
# out of 250
>1.01 # out of 250
Llama 3,1-405B
0.23
3.14
0.63
0
0
68
DeepSock-V3
0.34
2.96
0.76
0
0
99
Vomilla
DeepSock-R1
0.88
14.4
1,00
0.75
0
179
18
OpenA1-O1
0.73
12.4
1.00
0.55
141
Llama 3,1-B05B
1.18
18.4
1.03
1.00
1.00
247
88
DeepSeek-V3
1.32
52.4
1.32
1.03
1.00
247
113
Evolve
DeepSeck-RI
44.2
1.17
1.00
248
OpenA1-O1
1.35
63.9
1,38
1.16
1.00
247
158
Stage
32.7
1.00
0.96
240
Stage 1+2
1,36
48.3
1.09
1.00
247
175
Stage I+2+GRPO
2.41
84.6
1.83
1.33
1.11
247
207
CUDA-LI
stages random
2.14
64.5
1.62
1.21
1.09
241
186
island
3,21
126
2.21
1.40
1.16
249
223
bucket
3.12
120
2.25
1.42
1.17
249
226
Table 2: Model performances on KernelBench All Level.
viding researchers and practitioners with ready-to-use optimized implementations that can serve as strong baselines for future performance studies and benchmarking efforts.
3.3 MAIN RESULTS ON KERNELBENCE
The experimental results in Table T demonstrate CUDA-LI'' optimization effectiveness across dif- ferent baseline configurations on KernelBench. CUDA-LI achieves substantial performance im- provements over the Default baseline with 3.12x average speedup and 120x maximum gains. Against Torch compilation baselines, CUDA-LI delivers moderate but consistent improvements with 2.77-2.88x mean speedup ratios, while demonstrating 2.81 x mean improvement over CUDA Graph baseline with notable 97.9x maximum gains.
Across difficulty levels, CUDA-LI's optimization effectiveness varies by task complexity. For Level (single operations), CUDA-LI achieves moderate improvements ranging from 2.78-3.38X over different baselines. Level (operator sequences) shows CUDA-LI's strongest performance with 3.55x improvement over Default baseline. Level 3 (complex MI. tasks) reveals interesting baseline-dependen effectiveness: CUDA-LI achieves 2.96x improvement over Default baseline, but shows reduced effectiveness against Torch compilation baselines (only 1.62-1.98x improve- ments), suggesting these configurations provide stronger baseline performance for complex opera- tions.
3.4 BASELINE COMPARISON
We compare the results with the following three groups of baselines:
Vanilla Foundation Models: To establish baseline performance benchmarks, we evaluate OpenAI- o1, DeepSeek-R1, DeepSeek-V3, and Llama 3.1 3.1.405B Instruct (denoted by OpenAl-ol-vanilla: (oapSeek-RR-vaiilll DeepSeek-Ve -V3-aiik-V- and Llama 3.1-405B-vanille model optimize the reference CUDA code. The generated CUDA code is directly used for eval- uation without further modification. For each task, we repeat this process 5 times and report the best.
by prompting each
We implement evolutionary LLM strategies where, given a set of previous codes, we sample up to 4 high-performing kernels based on evaluation scores. The key difference is that the model only performs contrastive analysis without updating model parameters. We adopt the island strategy for code database construction and sampling, as suggested in Novikov et al. (2025). We conduct experiments on DeepSeck-R1, OpenAl-ol and and Llama 3.1-405B, denoted as DeepSeek-RI -R-eeolve, OpenAl-ol-evolve, DeepSeek- BonoSSeek----evor V3-evolve and Llama 3.1-405B-evolve
Evolutionary LLM:
Different combinations of CUDA-LI components and variants:
stage1: Uses only the outcome from the first stage with supervised fine-tuning applied
stagel+2: Applies only the first two stages without reinforcement learning
stagel+2 + GRPO: Replaces the contrastive RL with a vanilla GRPO strategy, without
comparative analysis

---

Published as a conference paper at ICLR 2026
Success
Speedup
Configuration
GPU Device
Mean
Max
75%
50%
25%
out of 250
>1.01x
A100
3.12
120
2.25
1.42
1.17
249
226/250
3090
2.51
114
1.57
1.18
242
201/250
Defauit
H100
3.85
368
1.76
1.32
1.09
250
218/250
226/250
1,40
3.13
182
1,88
1,31
1.08
248
215/250
A100
2.77
69.0
2.55
1.72
1.14
249
203/250
3090
2.58
73.2
2.23
1.50
1.00
242
177/250
Torch Compile
H100
2.74
49.7
2.83
1.92
1.11
250
195/250
H20
2,89
49.4
3,21
2.04
1.19
247
209/250
L40
2.85
96.9
2.43
1.82
1.13
248
199/250
A100
2.48
1.67
200/250
3090
2.61
72.9
2.29
1.48
1.00
242
172/250
Torch Compile RO
H100
2.78
1.61
187/250
H20
2.82
52.1
3,18
1.64
1,06
247
192/250
90.9
2.54
1.72
A100
2.81
97.9
1.83
1.20
0.954
229
147/229
3090
3.34
156
1.94
1.28
0.997
206
148/206
CUDA Graph
H100
2.23
70.1
1.60
1.04
0.838
222
119/222
H20
2.20
64.6
1.69
1.09
0.854
229
133/229
1,40
3.98
275
1.83
1.16
0.862
224
137/224
Table 3:
Performance comparison across different configurations and GPU devices KernelBench. RO = Reduce Overhead. Speedup is defined as value exceeding 1.01x.
random sampling: Replaces the bucket sampling strategy with simple random sampling of exemplars
island sampling: Adopts an island-based sampling strategy Novikov et al. 2025 where examples are distributed across different islands, prompts are constructed using exemplars from the same island, and newly generated examples are added to that island. After fixed number of iterations, examples half of the inferior islands are eliminated and examples from superior islands are copied replace them.
Results are shown in Table 2 As observed, all vanilla foundation models perform poorly on this task. Even the perforrffo models, such as DeepSeek-RI and OpenAl-ol, achieve speedups over the reference kernels fewer than 10% tasks, while Llama 3.1BB05B optimizes only 2.4% of tasks. This confirms that vanilla foundation models cannot be readily applied CUDA optimization due to their insufficient grasp of CUDA programming principles and optimization techniques.
We observe significant performance improvements introduced by the Evolutionary LLM models compared to vanilla foundation model setups. despite sharing the same parameter sets. All Evolve models achieve speedups in over 70% of tasks, with DeepSeek-FR reaching 72.4% success rate. This demonstrates that leveraging contrastive analysis, which exploits the model's general reason- ing abilities, more effective than direct output generation. The superiority of evolutionary LLM over vanilla LLM also provides evidence that contrastive RL should outperform non-contrastive RL approaches like vanilla GRPO, as relationship between evolutionary and vanilla LLMs parallels that between contrastive and non-contrastive RL methods.
When comparing the different combinations of CUDA-LI components, we observe a progressive increase in speedup rates from stagel (SFT only) 22.4% to stagel+2 (SFT + self-supervised) at 66%, and further to stage1+2+GRPP at 88.4%. This demonstrates the cumulative benefits of each training stage in improving model performance.
All Rased approaches significantly outperform evolutionary LLM baselines with fixed model parameters, with the best RL methods achieving over 95% speedup rates compared to 72.4% for the best evolutionary approach. This demonstrates the necessity of model parameter updating for achieving optimal performance in CUDA optimization tasks.
8

---

Published as a conference paper at ICLR 2026
3.5 GENERALIZATION OF A A100-OPTIMIZED KERNELS TO OTHER GPU ARCHITECTURES
Even without being specifically tailored to other GPU architectures, we observe significant per- formance improvements across all tested GPU types, with mean speedups ranging from 2.38x to x. H100 achieves the highest mean speedup (3.85x) with exceptional maximum gains (368 x), while A 100 PCIe and L.40 demonstrate strong performance with mean speedups of 3.12x and 3.13x respectively. L.40 shows the second-highes, maximum speedup (182x) among all GPUs. The con- sumer RTX 3090 achieves a competitive mean speedup of 2.51 x, while H20 shows moderate per- formance with 2.38 mean speedup. Notably, A100 maintains the highest 75th percentile (2.25 50th percentile (1 (1.42x), and 25th percentile (1.17x) values, indicating more consistent optimization performance on the target architecture.
The success rates remain high across all architectures (242-250 out of 250), with H100 achieving perfect success (250/250), validating that CUDA optimization techniques can generalize across dif- ferent GPU architectures. Speedup achievement rates (>1.01 x) vary by architecture, with H20 and A100 showing the highest effectiveness (226 and 226 successful optimizations respectively), while RTX 3090 demonstrates good performance with 201 successful optimizations.
These results demonstrate that while A A000ooptmmized kernels transfer to other GPUs with varying degrees of effectiveness, the optimizations achieve substantial improvements across architectures. H100's exceptional performance suggests strong compatibility with the optimization techniques, while A A100's consistent percentile performance validates the target architecture optimization. The varying maximum speedups (63.7x to 368 x) across GPUs indicate architecture-specffe optimiza- tion potential, suggesting that dedicated optimizations for each GPU type would further enhance performance. We plan to release kernels specifically trained for different GPU types in an updated version of CUDA-LI.
4 RELATED WORK
4.1
RI-AUGMENTED LLMS FOR CODE OPTIMIZATION
Starting this year, there has been a growing interest using LLM or RL-augmented LLM models for code optimization, including recent work on compiler optimization (Cummins et al. 2025) and
assembly code optimization (Wei et al., 2025a which use speed and correctness as RL. training rewards. Other more distant related is software optimization that scale RL-based LLM reasoning for software engineering (Wei et al. 2025b). Regarding CUDA optimization, the only work that comprehensively delves into KernelBench is from Lange et al. (2025), which uses a meta-generation faoge
procedure that successfully optimizes 186 tasks out of 250 tasks in KernelBench with a medium speedup of 34%. Other works remain in preliminary stages, including Chen al. (2025), which has optimized 20 GPU kernels selected from three different sources: the official NVIDIA CUDA Samples, LeetGPU, and KernelBench using proposed feature search and reinforcement strategy; and an ongoing tech report (Schulman et al. 2017) that optimizes 4 kernels.
4.2 EVOLUTIONARY LLMS
Evolutionary large language models Zhang et al. 2024: Liu ct al. 2024b Romera- Pomera-Paredes et al., 2024 Novikov al. 2025 Wei et al. 2025a Dat al.) [025 Lee et al.) (025) represent paradigm shift in automated algorithm discovery, exemplified by systems such as Google DeepMind's Al- phaEvolve (Novikov et al. 2025) and FunSearch (Romera- (Romera-Pareeees al. 2024).
5 CONCLUSION
In this paper, we propose CUDA-LI, pipelined system for CUDA optimization powered by con- trastive RL. CUDA-LI achieves significant performance improvements on the CUDA optimization task, delivering an average speedup of x3.12 (median x 1.42) over the default baseline across all 250 CUDA kernels of Kemnelennch, with peak speedups reaching x 120 on A100. Against other base- lines, CUDA-LI demonstrates x2.77 over Torch Compile, x2.88 over Torch Compile with reduce overhead, and x2.81 over CUDA Graph implementations. CUDA-LI can independently discover CUDA optimization techniques, learn to combine them strategically, and more importantly, extend the acquired CUDA reasoning abilities to unseen kernels with meaningful speedups. We hope that CUDA-LI would open new doors for automated optimization of CUDA, and substantially promote GPU efficiency and alleviate the rising pressure on GPU computing resources.

---

Published as a conference paper at ICLR 2026
REFERENCES
Thomas Bäck and Hans-Paul Schwefel. An overview evolutionary algorithms for parameter optimization. Evolutionary computation, 1(1):1-23, 1993.
Wentao Chen, Jaace Zhu, Qi Fan, Yehan Ma, and An Zou. Cuda-Ilm: Lims can write efficient cuda kernels. arXiv preprint urXiv:2506. 09092, 2025.
Chris Cummins, Volker Seeker, Dejan Grubisic, Baptiste Roziere, Jonas Gehring, Gabriel Synnaeve, and Hugh Leather. Llm compiler: Foundation language models for compiler optimization. In Proceedings of the 34th ACM SIGPLAN International Conference on Compiler Construction, pp- 141-153, 2025.
Pham Vu Tuan Dat, Long Doan, and Huynh Thi Thanh Binh. Hsevo: Elevating automatic heuristic design with diversity-driven harmony search and genetic algorithm using llms. In Proceedings of the AAAI Conference on Artificial Intelligence, volume 39, pp. 26931-26938, 2025.
Aaron Grattafiori, Abhimanyu Dubey, Abhinav Jauhri, Abhinav Pandey, Abhishek Kadian, Ahmad Al-Dahle, Aiesha Letman, Akhil Mathur, Alan Schelten, Alex Vaughan, al. The Ilama 3 herd of models. arXiv preprint arXiv:2407.2 2024.
Daya Guo, Dejian Yang, Haowei Zhang, Junxiao Song, Ruoyu Zhang, Runxin Xu, Qihao Zhu, Shirong Ma, Peiyi Wang. Xiao Bi, et al. Deepseek-rl: Incentivizing reasoning capability in Ilms via reinforcement learning. arXiv preprint arXivv2501.12945 2025.
Aaron Hurst, Adam Lerer, Adam Goucher, Adam Perelman, Aditya Ramesh, Aidan Clark, AJ Os- trow, Akila Welihinda, Alan Hayes, Alec Radford, et al. Gpt-4o system card. arXiv preprint arXiv:2410.21276, 2024.
Aaron Jacch, Adam Kalai, Adam Lerer, Adam Richardson, Ahmed El-Kishky, Aiden Low, Alec Helyar, Aleksander Madry, Alex Beutel, Alex Carney, et al. Openai ol system card. arXiv preprint arXiv:2412.16720 2024.
Albert Jiang, Alexandre Sablayrolles, Antoine Roux, Arthur Mensch, Blanche Savary, Chris Bam- ford, Devendra Singh Chaplot, Diego de las Casas, Emma Bou Hanna, Florian Bressand, et al. Mixtral of experts. arXiv preprint arXiv:2401.04088. 2024.
Robert Tjarko Lange, Aaditya Prasad, Qi Sun, Maxence Faldor, Yujin Tang, and David Ha. The ai cuda engineer: Agentic cuda kernel discovery, optimization and composition. Technical report, 2025.
Kuang-Huek Lee, Ian Fischer, Yuch-Hua Wu, Dave Marwood. Shumeet Baluja, Dale Schuurmans, and Xinyun Chen. Evolving deeper I1m thinking. arXiv preprint arXiv:2501.09880 2025.
Aixin Liu, Bei Feng, Bing Xue, Bingxuan Wang, Bochao Wu, Chengda Lu, Chenggang Zhao, Chengqi Deng, Chenyu Zhang, Chong Ruan, et al. Deepseek-v3 technical report. arXiv preprint arXiv:2412.194377 2024a.
Fei Liu, Xialiang Tong, Mingxuan Yuan, Xi Lin, Fu Luo, Zhenkun Wang. Zhichao Lu, and Qingfu Zhang. Evolution of heuristics: Towards efficient automatic algorithm design using large language model. arXiv preprint arXiv:2401.0200 2024b.
Niklas Muennighoff, Zitong Yang, Weijia Shi, Xiang Lisa Li, Fei-Fei Li, Hanna Hajishirzi, Luke S. Zettlemoyer, Percy Liang. Emmanuel J. Candes, and Tatsunori Hashimoto. sl: Simple test- time scaling. ArXiv, abs/2501.19393, 2025. URL lattp:://apt www.nticscoolerr......ccc.
CorpusID: 2D:27607969
Alexander Novikov, Ngån Vü, Marvin Eisenberger, Emilien Dupont, Po-Sen Huang, Adam 7solt Wagner, Sergey Shirobokov, Borislav Kozlovskii, Francisco JR Ruiz, Abhas Mehrabian, et al. Alphaevolve: A coding agent for scientific and algorithmic discovery. arXiv preprint arXiv:2506.1 13131, 2025.
Team OLMo, Pete Walsh. Luca Soldaini, Dirk Groeneveld, Kyle Lo. Shane Arora, Akshita Bhagia, Yuling Gu, Shengyi Huang, Matt Jordan, et al. 2 olmo 2 furious. arXiv preprint arXiv:2501. 2.0602500.00666.. 2024.
10

---

Published as a conference paper at ICLR 2026
Anne Ouyang. Simon Guo, Simran Arora, Alex L Zhang, William Hu, Christopher Ré, and Azalia Mirhoseini. Kernelbench: Can Ilms write efficient gpu kemels?? arXiv preprint arXiv:2502.10511, 2025.
Bernardino Romera-Paredes, Mohammadamin Barckatain, Alexander Novikov, Matej Balog, M Pawan Kumar, Emilien Dupont, Francisco JR Ruiz, Jordan S Ellenberg. Pengming Wang, Omar Fawzi, et al. Mathematical discoveries from program search with large language models. Nature, 625(7995) 2024.
John Schulman, Filip Wolski, Prafulla Dhariwal, Alec Radford, and Oleg Klimov. Proximal policy optimization algorithms. arXiv preprint arXiv: (Et:::770006344 2017.
Zhihong Shao, Peiyi Wang, Qihao Zhu, Runxin Xu, Junxiao Song, Xiao Bi, Haowei Zhang, Mingchuan Zhang, YK Li, Y Wu, et al. Deepseekmath: Pushing the limits of mathematical reasoning in open language models. arXiv preprint arXiv:2402.033: 03300, 2024.
Zhang Shengyu, Dong Linfeng, Li Xiaoya, Zhang Sen, Sun Xiaofei, Wang Shuhe, Li Jiwei, Runyi Hu, Zhang Tianwei, Fei Wu, et al. Instruction tuning for large language models: A survey. arXiv preprint arXiv:2308. 10792, 2023.
Gemini Team, Rohan Anil, Sebastian Borgeaud, Jcan-Baptiste Alayrac, Jiahui Yu, Radu Soricut, Johan Schalkwyk, Andrew Dai, Anja Hauth, Katie Millican, et al. Gemini: a family of highly capable multimodal models. arXiv preprint arXiv:2312. /805, 2023.
Gemma Team, Thomas Mesnard, Cassidy Hardin, Robert Dadashi, Surya Bhupatiraju, Shreya Pathak, Laurent Sifre, Morgane Rivière, Mihir Sanjay Kale, Juliette Love, et al. Gemma: Open models based on gemini research and technology. arXiv preprint arXiv:2403 08295, 2024.
Shuhe Wang, Shengyu Zhang, Jie Zhang. Runyi Hu, Xiaoya Li, Tianwei Zhang, Jiwei Li, Fei Wu, Guoyin Wang, and Eduard Hovy. Reinforcement learning enhanced Ilms: survey, arXiv preprint 1rXi::2412.10400, 2024.
Anjiang Wei, Tarun Suresh, Huanmi Tan, Yinglun Xu, Gagandeep Singh, Ke Wang, and Alex Aiken. Improving assembly code performance with large language models via reinforcement learning. arXiv preprint arXiv:2505. 11.022555,1110,,, 2025a.
Jason Wei, Xuezhi Wang, Dale Schuurmans, Maarten Bosma, Fei Xia, Ed Chi. Quoc V Le, Denny Zhou, et al. Chain-of-thouuht prompting elicits reasoning in large language models. Advances in neural information processing systems, 35:24824 258377 2022.
Yuxiang Wei, Olivier Duchenne, Jade Copet. Quentin Carbonneaux, Lingming Zhang, Daniel Fried, Gabriel Synnaeve, Rishabh Singh, and Sida I Wang. Swe-rl: Advancing Ilm reasoning via rein- forcement learning on open software evolution. arXiv preprint arXiv:2502. 18449, 2025b.
Ronald J Williams. Simple statistical wradient-following algorithms for connectionist reinforcement learning. Machine learning, 8:229-255 256, 1992.
An Yang, Anfeng Li, Baosong Yang, Beichen Zhang, Binyuan Hui, Bo Zheng, Bowen Yu, Chang Gao, Chengen Huang, Chenxu Lv, et al. Qwen3 technical report. arXiv preprint arXiv:2505. (51.:2555009888, 2025.
Rui Zhang, Fei Liu, Xi Lin, Zhenkun Wang, Zhichao Lu, and Qingfu Zhang. Understanding the importance of evolutionary search in automated heuristic design with large language models. In International Conference on Parallel Problem Solving from Nature, pp. 185-202. Springer, 2024.
11

---

Published as a conference paper at ICLR 2026
THE USE OF LARGE LANGUAGE MODELS (LLMS)
We used a large language model (i.e., ChatGPT 03) only as a general-purpose assist tool for minor English grammar corrections during manuscript preparation. The L.M had role in research ideation, methodology, experimental design, data collection, analysis, interpretation, or substantive writing beyond copyediting at the sentence level. All scientific content, claims, and conclusions are the authors' own. The authors take full responsibility for all contents written under their names, including any text that may have been edited with the assistance of an LLM. LLMs are not eligible for authorship or contributorship on this work.
A DETAILS FOR CUDA-LI
A.1 DATA COLLECTION FOR SFT
To expand the model's exposure to CUDA patterns, we begin with data augmentation based reference code from 250 tasks in KernelBench, which provides the official implementations used in PyTorch. To generate executable and correct CUDA code efficiently, we leverage six existing LLM models: GPT-4o, OpenAI-ol, DeepSeek-RI, DeepSeek V3, Llama 3.1-405B Instruct, and Claude 3.7. For each model, we construct prompts using the one-shot strategy, where the prompt contains the reference code (denoted by i E [1, 250) and asks the LLM to generate an alternative speedup implementation. We employ multiple models to maximize the diversity of successful CUDA code generation. The detailed prompt structure is provided in Table 4 For each of the six models, iterate through 250 tasks. Each task allows up to 20 trials and terminates early we successfully collect 2 trials that are both executable and correct. Notably, some tasks may fail to produce any successful code across all trials. The successful code is denoted by dr where j E (1.2..... n, and ny denotes the number of successful code snippets for the reference code d:- Through this process, we collected 2,105 successful CUDA code snippets. Now we have collected the dataset D =(g (
A.2 PsuDO CODE FOR STAGE2: SELF-SUPERVISED LEARNING
Self-supervised Learning Algorithm
1: Initialize finetuned model Mp after SFT stage with parameters 6.n
2: for io to Mitrations do
3: Generate batch of CUDA codes C, (ci,-.,;a) G using model M,-1
4: Evaluate each e E for:
5:
1. Executability (compiles and runs)
6:
2. Correctness (produces expected output)
7:
Filter successful codes: Chursess
{e E Ceeecuuuablle A correct)
8:
if Craceess o 9 then
9:
Compute gradient update V using Caucoess
10:
Update model: 8; + 0:-1 + 778
11:
clse
12:
0: - ---1 (no update)
end if
end for
return Final improved model My
Table 4: Self-supervised learning for cuda optimization in Stage 2.
A.3 DETAILS FOR RL-TRAINING
Specifically, for each reference prompt containing selected exemplars as shown in Table 9 we sample G code outputs from the current policy Told. denoted as {dl .2.....de de)- Let r (r1, ... , ,G) represent the reward scores associated with the generated code samples. Different from standard GRPO training, rewards are smoothed to mitigate the reward hacking issue; the de- tails of this approach will be elaborated in Section A.5 Further, GRPO, rewards normalized
12

---

Published as a conference paper at ICLR 2026
within each group using:
A r; mean(r)
(3)
std(r)
The complete GRPO objective optimizes the policy model by maximizing:
[d:l
To di Pi.
(0) ( Pi), (dla)
min
G
Tulle (di,elg, di,
i=1
t=
To To(di,loa di,ct) Toal (di,e de,
clip
- e, + +
3D [Tol]
(4)
where:
is the policy model being optimized
Tead is the old policy model from the previous iteration
8 is the parameter for clipping
3 is the KL penalty coefficient that controls deviation from the reference policy
Dr.. denotes the KL. divergence between the current and reference policies
We refer readers to (Shao al. 2024) for details of GRPO. Model parameters are optimized using the GRPO objective, with contrastive prompts that incorporate comparative examples. This con- cludes our description of contrastive RI..
A.4 CONTRASTIVE- RL's ADVANTAGES OVER EVOLUTIONARY LLM APPROACHES
Contrastive-RL draws inspiration from broad range of literature, including evolutionary algorithms (Bäck & Schwefel, (993) and their applications to LLMs (Liu et al. 2024b) Romera Comera-Paredes, et al.) Novikov al. 2025 Wei et al. 2025a), where multiple solution instances with associated fit- ness scores are presented to LLMs to analyze performance patterns and generate improved solutions.
However, Contrastive-RL improves evolutionary LLM approaches in several critical aspects:
Model Adaptation vs. Fixed-Model Reasoning: Contrastive-RL. employs gradient- based pa- rameter updates to continuously enhance model capabilities, whereas evolutionary LLM approaches rely exclusively on in-context learning with static parameters. This fundamental architectural dif- ference endows Contrastive-RL with substantially greater representational capacity and task adapt- ability. Evolutionary LLM methods are fundamentally limited by the frozen foundation model's initial knowledge and reasoning abilities, while Contrastive-Rl progressively refines the model's domain-specific expertise through iterative parameter optimization. From this perspective, evolu- tionary L.M approaches be viewed degenerate only the Fixed-Parameter Solution Optimization component while omitting the Foundation Model Enhancement mechanism. This theoretical relationship explains why Contrastive-RL consistently outperforms evolutionary approaches: it leverages both optimization dimensions simultancously rather than constraining itself a single fixed-capacity search space.
of Contrastive-RL. that implements
Scalability and Generalization:
Contrastive-RL demonstrates superior scalability by training a single specialized model capable of handling diverse CUDA programming tasks and generating var- ious types of optimized code. In contrast, evolutionary LLM approaches typically require separate optimization processes for each distinct task domain, limiting their practical applicability and computational efficiency.
A.S MITIGATING REWARD HACKING RL TRAINING
Reinforcement learning is notorious for exhibiting reward hacking behaviors, where models ex- ploit system vulnerabilities to achieve higher rewards while generating outputs that deviate from the intended objectives. particularly challenging aspect of these pitfalls is that they cannot be anticipated prior to training and are only discovered during the training process. During our initial training procedure, we identified the following categories reward hacking behaviours:
Improper Timing Measurement. KernelBench measures execution time by recording timing
events on the main CUDA stream:
13

---

Published as a conference paper at ICLR 2026
start_event. record (
2 model (*inputs)
3 end_event record (endiggnaaaammmeeststeee torch. cuda. synchrcnize (device=device)
However, RL-generated code exploits this by creating additional CUDA streams that execute asynchronously. Since KemelBench only monitors the main stream, fails capture the tual execution time of operations running on parallel streams. This vulnerability is significant: in our initial implementation, we find that 82 out of 250 (32.8%) RL-generated implementa- tions exploit this timing loophole to appear faster than they actually are, leading to an overall speedup of 18x. To address this issue, prompt engineering alone insufficient. The evalua- tion methodology should be modified to synchronize all CUDA streams before recording the end time, ensuring accurate performance measurement across all concurrent operations as follows:
start_events record cecustoo_mooder stream)
custon_mode/ (*inputs)
3 # Wait for all model streams to complete before recording end event
4 i mustom__ontaal seataii_ne__streanns
5
for stream in uustonnmmodelsssrremm:::
6
customnmodellstreee record (custon ret.eorrccccttttoooooooeeaeeeee 8 torch. cuda. cuda. synchronize (device=devicee
(stream)
7
Lazy Evaluation Another important hacking strategy lazy evaluation, detected by community users on GitHub: Calling custom_mode costoo_model("inuus doesn't ensure the output is actually material- ized/computed. The computation is actually executed the correctness check phase when calling the torch.allclose() function, allowing it to pass the correctness check.
class LazyMatmul (torch. Tensor)
2
def _new____________________________ ________ll__ (cls, A, B):
3
result
wwwe__aaalllaaaassaaaaaaa (cls, torch. (0))
4
result.A
result.B
result._shape (A.size B.size (1))
7
return result
8
def (selr)))
"""Trigger actual computation"" return torch.matmul se,, self.B)
2
3
class ModelNew (an. Kodule) def forwardself., A, B) return LazyMatmul (A, B) Returns lazy object
5
To mitigate this issue, we enforce validation check that ensures materialize functions are called before ending the time measurement, which involves checking the following conditions: the output must be tensor, must be standard torch.Tensor (not subclass), must be on the correct device, must have allocated memory, and the corresponding storage must be valid.
1 # Check 1: Must be tensor
if not isinstance (out, torch.Censez Densor)
3 return False, f'(prefix) is not a tensor: (type (out) }"
4
5
Pl Check 2: Must be standard torch. Tensor, not subclass
if type _aat)____m_____ not in ['Tensor', 'Parameter' return False, ("(prffix) is [type (out) name_ not standard
7
8
# Check 3: Must be on correct device
0 if out device device: return False, 1" f"(prefix) on wrong device: (eviccevice) (expected (device))"
1
2
3 # Check 4: have allocated storage
14

---

3. Extended Measurement Window: We conduct multiple evaluation rounds with prede- fined running time of 30 minutes per candidate. This adaptive approach yields between several of thousands 1M rounds depending on individual kernel execution times. Published as a conference paper at ICLR 2026
14 storage size out unttuped_storrae () .size() if storage_size 0: return False, f" [p(effiix) has no allocated storage (likely lazy)
7
8 # Check 5: Storage pointer must be valid
9 ptr out.data_ptr()
if ptr 0: return False, f"[prefix) storage pointer is null (likely lazy)
Hyperparameter Manipulation: In KernelBench, each computational task is associated with specific hyperparameters, including batch_size, dim. in_features dimension, out_feature dimension, scal I ing_factor, and others. The RL agent learned to exploit these parameters by generating code that artificially reduces their values, thereby achieving superfi- cial speedup improvements that do not reflect genuine optimization performance.
Result Caching: The RL agent developed strategies to cache computational results across eval- uation batches based on input addresses. When another input's address matches a cached one, it returns the cached output. In theory, this should not pass correctness validation because the cached output differs from the expected one. However, given that correctness validation checks whether the difference each position between the reference output and custom code output below a certain threshold, there a few cases where it is able to squeeze past the correctness bar. The following code snippet gives an illustration:
cache_key x.caata_ptr
2 # Check if result is cache
3 is cache_ key in self.
4
return self.cacher cache k
A.6 REWARD DENOISING
We observe a significant variance in tu measurements for identical implementations d, which intro- duces noise reward estimation. This noise is particularly detrimental to RL training stability. To address these challenges, we implement the following robust measurement strategies:
Each evaluation runs on an exclusively allocated GPU. Shared GPU usage leads to significantly higher variance in timing measurements, even when memory and compute utilization appear low.
Dedicated GPU Allocation:
2. Paired Execution with Order Randomization: For fair comparison, each evaluation round executes both the reference q and candidate d implementations. Crucially, we ran- domize the execution order within each round to account for GPU warm-up effects, where subsequent runs typically benefit from cache warming.
4. Bucketized Variance Control: We partition all Scoresingle (d) measurements into buckets and compute bucket-wise averages. Evaluations with bucket variance ex- ceeding 0.005 are discarded.
5. Robust Central Tendency: The final reward uses the median of bucket averages, which proves more stable than the mean against outlier effects:
rd) median ((Buckets 1T=1)
(5)
6. Conservative Rounding: We apply conservative rounding speedup ratios (i.e., Score(d)
), truncating to two decimal places while biasing toward unity (e.g., 1.118 + 1.11, 0.992 1.00).
7. Striet Verification Protocol:
Despite these precautions, we still occasionally observe spu- rious large speedups due to GPU turbulence. For any candidate showing either:
Absolute value of speedup > 3, or Speedup exceeding twice the previous maximum we perform verification on different GPU of the same type. The result is accepted only if the verification measurement differs by < 10% from the original.
15

---

Published as a conference paper at ICLR 2026
A.7 TOWARDS ROBUST REWARD DESIGN AND TRAINING PROCEDURES
To mitigate reward hacking, we implement the following strategies during training:
A reward checking model When there is significant leap in reward, adversarial model in- tervenes to determine whether the code exploits the reward system. We use DeepSeek-RI R1 for this purpose and find that it successfully identifies reward hacking above over 60% of the time.
Hacking-case database We maintain a dynamic hacking-case database that is updated whenever a new reward hacking behavior is detected. The reward checking model leverages this database for detection: given newly generated code snippet to examine, we retrieve the three most similar cases from the database and include them context for the reward checking model's input.
Reward smoothing Sharp reward increases are smoothed reduce their magnitude, preventing the RL agent from over-prioritizng any single high-reward solution, whether legitimate or not:
T nnnomalizedd
(6)
T ssmmorch
clip( Fnommalzzed k: k:)
where μ and a are the mean and the mean and standard deviation of the reward distribution, respec- tively. k is hyperparameter that controls the clipping threshold 1.5, think achieving a 1.5x speedup over the official PyTorch implementation already represents significant optimization performance.
B DISCOVERED CUDA OPTIMIZATION TECHNIQUES
An analysis of optimization strategies commonly employed in enhanced CUDA implementations veals interesting patterns. Through GPT-4o-based technical term extraction and frequency analysis, identified the most prevalent optimization techniques:
Memory Layout Optimization, which ensures data is stored in contiguous memory blocks;
Memory Access Optimization, which arranges data access patterns to maximize memory bandwidth and minimize latency through techniques like shared memory usage, coalesced global memory access, and memory padding;
Operation Fusion, which combines multiple sequential operations into a single optimized kernel execution;
Memory Format Optimization, which aligns data layout with hardware memory access patterns;
Memory Coalescing, which optimizes CUDA kernel performance by ensuring threads in the same warp contiguous memory locations;
Warp-Level Optimization, which leverages the parallel execution of threads within warp (typically 32 threads) efficiently perform collective operations;
Optimized Thread Block Configuration, which carefully selects grid and block dimen- sions for CUDA kernels to maximize parallel execution efficiency and memory access pat- terns;
Shared Memory Usage, enables fast data access by storing frequently used data in cache accessible by all threads within thread block;
Register Optimization, which stores frequently accessed data in fast register memory to reduce latency and improve computational throughput:
Stream Management, which enables parallel execution operations for improved GPU utilization.
Tables
I1
13 and
14 present detailed CUDA optimization techniques with accompanying code
examples.
C CASE STUDIES
Table the KernelBench tasks that achieved the highest speedups. We examine these some of them in detail and perform an ablation study the applied CUDA optimization techniques, showing how much each technique contributes the final speedup.
16

---

Published as a conference paper at ICLR 2026
Level
Task
Task Name
Speedup
2
33_Conv3d _ Baripooo__MM___CC_ _an__llmpppp_a_m Clamp_ Dropout
120.3
1
12. Matmul_witt diagonal
witholigooo diagonal_matrices matrices
64.4
2
80
80_ L___em___aa Seak__Sutrrct_GGLL
31.3
1
9_ 9_Tall_skimny_maari httpr_ttttiii_illliiitttt
24.9
3
VisionAttention
24.8
2
96
96 CrnnTTaaspose3d Multiply _ax__Maaaa_aaaaa GilobalAvgPoo Clamp
16.2
2
66 Matmul Mropout MeattMMeen_Ssaaass Softmax
14.5
Matmul, for sym.strrmmmmm..mmm mattricss
3
43
43. MinGPTCaaualAAtenntion
13.1
3
44
44. MiniGPTBlock
10.5
Table 5: KemelBench Tasks Ranked by CUDA-L1 Acceleration (Top 10)
C.I DIAG(A) B: 64x FASTER
We first examine the code for level I, task 12, which performs matrix multiplication between a diagonal matrix (represented by diagonal elements) and a dense matrix, both with dimen- sion N=4096. The reference code is as follows where _iiit__ function of the class is omitted:
class Kodel (nn. Module); del forward 3, B):
2
A: (N,) ) 1D tensor shape N
4
(N, X) 2D
ahape x (N, N) creates diagonal matrix from
5
toach.diag
6
Reault: (N, N) (N, 8) - (N, K)
return torch,diag (A)
Let's see the optimized code by CUDA-LI:
1 class Model (nn, Kodule) def foreard (sel f, B): return unsqueeze
The optimized implementation leverages PyTorch's broadcasting mechanism perform diagonal matrix multiplication efficiently. first reshapes the diagonal vector from shape (N, to (N, using unsqueeze (1), transforming it into column vector. Next, it utilizes PyTorch's is auto- matic broadcasting to multiply each row of matrix B by the corresponding element of A, where the (N, 1) shaped tensor is implicitly expanded proach completely avoids creating the full N x N diagonal matrix, which would be sparse and memory-intensive. The key benefits of this technique are substantial: it requires only O(1) extra memory instead of O(N2) for storing the diagonal matrix, reduces computational complexity from O(N2M) operations for full matrix multiplication to just O(NM) element- operations, leading to 64x speedup.
match the (M, dimensions of B. This ap-
What makes this particularly valuable that RL can systematically explore the vast space of equiva- lent implementations. By exploring semantically equivalent implementations, RL learns to identify patterns where computationally expensive operations can be replaced with more efficient alterna- tives. The power of RL extends beyond simple algebraic simplifications and it can uncover sophis- ticated optimizations such as: replacing nested loops with vectorized operations identifying hid- den parallelization opportunities discovering memory-effiiien mathematical reformulations finding non-obvious algorithmic transformations that preserve correctness while improving performance What makes this particularly valuable is that RL systematically explore the vast space of equiv- alent implementations something that would be impractical for human engineers to do manually.
C.2 LSMT: 3.4x FASTER
Now let's look at classical neural network algorithm LSTM (level 3, task 35), which CUDA-11 achieves a speedup of 3,4x. By comparing the reference Py Torch implementation with the opti- mized output, we identified the following optimization techniques:
CUDA Graphs, which captures the entire LSTM computation sequence (including all layer operations) into replayable graph structure, eliminating kernel launch overhead by record-
17

---

Published as a conference paper at ICLR 2026
Contiguration
Graphe
Memory Contigaity
Spendap
CLDA + Memery + Stadic
3.4+
LSTM compataion
CUDA - Memory
2.96
Manery allecatios
Menury layost
CUDA Only
2.77x
Momary ovethedd
Memary Statk
1,0%
Kemel lannch oveitead
Momocy Only
1.00
Kemel braoh ovetheed
Only
mechhad
Handline
1,00
Kemol launch mentaad
Table 6: Speedup achieved by different CUDA optimization techniques on LSTMs.
ing operations once and replaying them with minimal CPU involvement for subsequent executions.
2. Memory Contiguity, which ensures all tensors maintain contiguous memory layouts through explicit .contiguous) calls before operations, optimizing memory access patterns and improving cache utilization for CUDA kernels processing sequential data.
3. Static Tensor Reuse, which pre-allocates input and output tensors during graph initializa- tion and reuses them across forward passes with non-blocking copy operations, eliminating memory allocation overhead and enabling asynchronous data transfers.
Tablee]repreeentss the results for 8 different optimization combinations across the three optimization techniques above. As can be seen, CUDA Graphs is essential for achieving any meaningful speedup in this LSTM model. All configurations with CUDA Graphs achieve 2.77x-3.42x speedup, while all configurations without achieve only 1.0x (no speedup). The combination of all three techniques provides the best performance at 3.42x, demonstrating that while CUDA Graphs provides the ma- jority of the benefit ( 81% total speedup), the additional optimizations contribute meaningful improvements when combined together.
C.3 3D TRANSPOSED CONVOLUTION: 120x FASTER
We examined the code for Level 2, Task 38, which implements a sequence of 3D operations: trans- posed convolution, average pooling, clamping, softmax, and element-wise multiplication. By com- paring the reference PyTTrch implementation with the CUDA-LI optimized output, we identified the following optimization techniques applied by CUDA-LI:
Mathematical Short-Cireuis which detects when min_ value equals 0.0 and skips the en- tire computation pipeline (convolution, normalization, min/clamp operations), directly turning zero tensors since the mathematical result is predetermined.
2. Pre-allocated Tensors, which creates zero tensors of standard shapes during initializa- tion and registers them as buffers, eliminating memory allocation overhead during forward passes for common input dimensions.
3. Direct Shape Matching, which provides a fast path for standard input shapes by immedi- ately returning pre-allocated tensors without any shape calculations, bypassing the compu- tational overhead entirely.
4. Pre-computed Parameters, which extracts and stores convolution parameters (kernel size, stride, padding, dilation) during initialization, avoiding repeated attribute lookups and tuple conversions during runtime.
Table 7 represents the results for 16 different optimization combinations across the four optimiza- tion techniques above. As can be seen, mathematical short-cireuit is essential for this task, where all configurations with mathematical short-cireuit achieve 28.6x+ speedup, while all configurations without it achieve only 1.0x (no).
The fact that CUDA-LI identified this precise optimization strategy demonstrates the power of re- inforcement learning in navigating complex optimization spaces. While a human developer might intuitively focus computational optimizations (like parallel algorithms) memory layout im-
18

---

Published as a conference paper at ICLR 2026
Coafgaratie
Pre-slacttad Temers
Skaps
Prococppatd Paran
Spuadg
Bedlio
Shape
Proilee Only
Table 7: Speedup achieved by different CUDA optimization techniques the Conv3d task.
provements (like tensor pre-allocation), RL discovered that the mathematical properties of the op- eration completely dominate performance. This discovery is particularly impressive because: RL able to find this non-obvious solution: The 120x speedup from exploiting the mathematical short- circuit is counterintuitive as most developers would expect to optimize the convolution kernel or memory access patterns for such compute-heary operation, This shows how RL can discover optimal solutions that challenge conventional wisdom in deep learning optimization. Where hu- man intuition might suggest "optimize the convolution algorithm first," CUDA-L1 learned through empirical evidence that "recognize when computation can be entirely skipped" yields dramatically better results. The agent's ability identify that min(x, 0) followed by clamp(O, 1) always produces zeros demonstrates how RL can uncover mathematical invariants that humans might overlook in complex computational pipelines.
19

---

Published as a conference paper at ICLR 2026
D PROMPT USED IN THE PAPER
Data Augmentation Prompt - Used in Supervised fine- tening
Task for CUDA Optimization
You are an expert in CUDA programming and GPU kernel optimization. Now you're tasked with developing nigh-performance cuda implementation of Softmax. The im- plementation must:
Produce identical results the reference PyTorch implementation.
Demonstrate speed improvements on GPU.
Maintain stability for large input values.
Reference Implementation (exact copy)
import torch
2
import torch. an as
class Model (nn. :on.Med
nnm
6 Simple model that performs Softmax activation.
def init init_ (self):
super (Model, self) .ini_iiit__ ()
def forward x: torch. ('ensor) => torch.Tensoo
Applies Softmax activation the input tensor. Args:
15
x (torch. Tensor): Input tensor of shape (batch_size, num_features).
Returns:
torch. Tensor: Output tensor with Softmax applied, same shape as input.
18
return torch. softmax (x, dim=l)
batch size 16
dim 16384
def get_inputs ():
rand (batch_sste size, dim) return [x]
def got_init_iaputell
return # No special initialization inputs needed
Table 8: Prompt illustration for data augmentation in Section ??. For each KernelBench task (soft- max shown here for illustration), the prompt is fed to each of six LLM Mondel-GGPT-40, OpenAI- ol, DeepSeek-R1, DeepSeck V3, Llama 3.1BB Instruct, and Claude 3.7 Sonnet-to generate alternative CUDA implementations.
20

---

Published as a conference paper at ICLR 2026
CUDA Optimization Task Prompt
Used in Contrastive-RL
Task for CUDA Optimization
You are CUDA programming expert specializing in GPU kernel optimization. Given a reference CUDA implementation, your objective is create an accelerated version that maintains identical functionality. You will receive previous CUDA implementations accompanied by their performance metrics. Conduct a comparative analysis of these implementations and use the insights to develop optimized and correct CUDA code that delivers superior performance.
Reference Code
globai_ void kernel_vl (float input, float output, int
// Baseline implementation
Previous Cuda Implementations with Scores
1 // codei (scorel)
2 global void kernel_v (float* input, float* output, int
N)
6 // code2 (score2)
global__ void kernel_v2 (float* input, float* output, int
Generation Protocol
You MUST use exactly two hash symbols (#Hf) the beginning of each section. ## Performance Analysis: Compare code snippets above and articulate on
1. Which implementations demonstrate superior performance and why?
2. What particular optimization strategies exhibit the greatest potential for im- provement?
3. What are the primary performance limitations in the implementation?
4. What CUDA-specific optimization techniques remain unexploited?
5. Where do the most significant acceleration opportunities exist?
Algorithm Design: Describe your optimization approach Code Implementation: Provide your improved CUDA kernel
Requirements and Restrictions
## Critical Requirements
1. Functionality must match the reference implementation exactly. Failure to do
so will result score of O.
2. Code must compile and run properly modern NVIDIA GPUs
## Key Restrictions
1. Do not cache reuse previous results the code must execute fully on each run.
2. Keep hyperparameters unchanged (e.g., batch size, dimensions, etc.) speci- fied in the reference.
Table Prompt structure for CUDA optimization task showing reference implementations and their performance scores used in Con
21

---

Published as a conference paper at ICLR 2026
E CASE STUDY: CODE SNIPPETS BEFORE AND AFTER OPTIMIZATIONS
Tech Dese
Before optimization
After optimization
Memory Layout
Optimization
-Non-contignous
Ensuring contigouus layout
.h-python
"Fytonn
Memory Layout
(R
natrix nultiply optamized 1N,
Optimization
night
Igunes remory
iguens nemory leyout
ensures data
patterne
stored in con- liguous max-
recuen
cononiiuuaa 1f cods_contigucco
cononiiguess 1e Bite, contigueusci
imize cache
efficiency and
reduce memory
lotency
GPU
computations.
Memory Con-
lescing
Uncoalesced memory aecess
Coulesced memory access with loop unrolling
Meuda
Memory coa-
globel wold uncoslesced sernel iflost
global
vold coslesced kernel (fTosts ingut,
lescing optimizes
input, float outpet)
flcaz+ outputi
GPU
inz tid thrcacl6x.Kj
Lnt rid EazcadIdk.x/
inz stride 2lo.kDi8.X/
int batch_1dk blockIdx.Xi
ing threads
// Rach thread accesses non-contigiom nemory looat one
// Base pointers for this batch Iten conat Flost batch tngur - inpot 1dx 1024)
warp acvess contiguous mem- locations, reducing memory transacioon and incressing hand- width utlizatioe.
fox (int 1 Dj 1 1024) 1+-) {
output(tid 1 stride] Irgut[tid stride] 2.0f;
Eloate batca_catpat. outpet batch_1dx 1324;
// Rach thread progesses cont Igueos nerory chanks
Ipzzgna unroll 4
for Hin I 0j 1 < 1024; 1 16J
batch output H1 batch Inpot/51 2.05i batch output [1-1] batch Inputile1]
2.0f;
Nateh_outgus [112] batch 2.05;
// note contigueus accesses
batch output I 11-15 batch (npot(1415) 2.0f;
Warp-Level
Oplimizations
Kach thread independently caleulates min value
=iiing warp-level operations reducctio
"ennd
Warp-Leved
gloha woid derueerne' pernee__effor _before (conet fllows inpat, flost output, int wize)
Optiiiztioons
globel 2_ void min_keroelaafe Sourat Floate
leverage the
iths blockidk.x blockDin.x
input, floats ausput, mixe; intt bisckidx.x bisckuin.x ChreadIdk..r
CUDA execution
threadidx.x;
model where
size)
execute
for (int Dp 1 DEPrI, 1nD &
Flost min_val tetoe;
Int Jare_ "hreed"dde.x & 32; // Thread' posision within warp inb wasp_id shheadd.dk.. 322 // Katp namber vithin the
in groups
min_val tiitniaaa,,, input (adx alae;)2
32 (warps)
allel efficiency
output [1dx] xin_val)
float min_val leel0;; 1idx sizc!
thoogg collaho-
ntive operatioss
17 Each thread Daring ittl lint sin_val nin (nin.vai, input [idx sizeli;
and memory
access paitems..
17 darp-level parallel reduction using lint offset offset Dj offses
other
shri_dowr dowl_ pyno (ORtrcteer.,
sin_val, offset)y
nin_val nia (nin_val, othes)y
in sarp
result
0) outpot (blocetkkk :tleckDin.w/933
warp_id] min_valr
Table 10: (Part 1) Code snippets before and after optimizations.
22

---

Published as a conference paper at ICLR 2026
Toch + Dese
Before optimization
After optimization
Memory Hierarchy Opti
mization
Using glohal memory dircctly
Using mensory hicrarchy (shared, constant, reglsters)
'oude
globbll void
clow: sudepttiiiaeeee weight 13x3x517 for weights
Memory Hierarchy Opti-
depthvi www.eeeeabbieeee aereaeeeeeeeeeeeee harrsuuammeee.. inddptinieed
11 Constant
misstion invoves
conec Cloate input, corst floate
conscant float o_pointaloe_ weight (3x64))
gically
depthalee weight, const floste
levels
memory
pointwise
global vold
constant menoey) to mini-
(registers shanel memory.
floats outpet, /+ other paranctors +f)
int ext_y bisecidk.y Blockbin.y
ccest floate Lnput, floace outpet, /4 other paratetara +/1
mise global memsey access
threadicx.y? cut blockidk.x
latency and maximize data
nevory inguc
padding
reas..
thread dirostly foz fint for corputation
global
[S8AR60_ (EN_STRDDE):
oz ost_chanmeiar oct!!
// Collaborative leading inpet data sharec nezcey
tloat seault 0.01:
for (int is G ie in_ehannelay iemt)
// [shazed sasory loading coda....
synothreede
Flowe depthvise seress___eee result 0.0ry // Gixest global maresy accea0 Icz each kersel clemenc for lint ky 01 ky < kyee!
1/ Segistor caching for intorseciata reselts float wapchrioe_resuttsiP // Store in regleters
soe (ins kx 97
(act))
ioys out_y atride xy
// Cosputo usine shared nanory and constant netory
gedstnng
inunhemeels;
s0 padsineg
Fully unrolled convolation using shared
height
in_x [suniesk) I
shared ingut [c] paa__ (sa_x_
inpuc (Bacch_id
nLchhanela iel
I/ (sore unrolled consation......
sttheight s1L inukidth in
depthtvssn www.v//tarseetaaaccom [c] marg /7 Store in register
degthvise se geight to
ky
k
outpus values zegisters sccput_cache[3E).
// Compute
pointwlee weight channels t91p
1initz for oxchees depthviss resulbel0)
c_pointxi weight 3 01
oucput (I (Hbatcc idx out_channels ool out beight out_y? out_vidth oat_x] - cessity
depcawite_ reouttssl
o.gointvise. bointw.:e_veighhEf
3 11 dustawisa c_
// Conlesoed to globel menory joupput_id (stsat_
Asynchromee Execution
Sequential
Asynchrunuu
Asynchranous
"Pythen
<>pytoon
CUDA
forward (self,
forward xi
opensioos quenel exccuted concumendy on
operations exeoute defauls streen,
desicated conpute
sequentiall
r.corgote
separate stremss. euabling overtappung computation with nemory tmndaos for improved GPU utilization..
fe.f.coav 3o1 result
priority searppicii
Execute opezations asynchrenecs.y in the custon strean
kith torch.cuda rean r.ooputte stream) result = se' f.. trrited bed ouda fonaard k.6typs)
Contzel returns imndiacely shile computation continues in auckerouns. suturn zesult
Table 11: (Part 2) Code snippets before and after optimizations.
23

---

Published as a conference paper at ICLR 2026
Tech
Before optimizatioo
optimizadion
Memory Access Optinization
Naive matris unlliplication with
Using shared lilnng and register blocking
Menuyy Acces Opinniza-
coude
tin CUDA improves
// Jefore optimii zet in Haive natrix
globa void nataul wot.oniiieeeeeeee delieeddtoon chloate
porfemmnas by equaisnng data
muttippication kon with poor
cloatr B, cinat ind x, ant N, int
asse patienss to uscinia:
global vold sattnal naive (floste h, fleace
// Block index and chread Lodex
maimuse
R, floate Cy int Hy inc M, int 10
const Int bx blocktdk.
memory lneryy through
tilns, codlescing,
tut row bi blocktdryy x bitcknsmyy
thizeddios.37
const Int by coeat inc LX Ehzoadl d6. coest lnc ty thzeadldk.22
shared
int. col - bloskIdk. blockDin.2 threadicx.2g
// Outpus positions cocat ins EaN by 051
tzow se B1 0.052
Atsov Blecl
// Reaister accumlattion sun30 0.02, sun01 0.0r, float suni0 0.02, euxll 0.ori
FLeow coll
// Shazed narery tiles with padcing avoid bank confilors saared #loat Aa(0]1331;
XAACRC cloat [e(22]1331
1/ loop over tiles 02 tile + 211 / 32, rttilel loz
CeLLaberattie loading tilos 1nto remory
Aatty]fex] ALEON tile
[tx)) 0.02;
if (eol < N s6 tile 32 by < K) Relsyllex] Biesi X tile 52 tyl2
else Reftyllex] 0.0%
Lynccheeadd O2
products using
Epragns unrol1 for si
+48
float Ae Aettylikis
floct al
Aa [ty 4] [K])
float 20 Ba [K] TEXTS
Float b1 Batk] TEX 1611
sun20 +- 40 bby
eumb1 a0 -h2
zumi bog
sumil bi;
12 Meitee resuica to globel menory
11 leox N 64 ead < a) Cleow M col)
cunD0)
1f Irow N 66 ost 16 M1 CFrow
csl 161 smm01;
if leoN < H ss col BI CI + 2011 suni.y
teoM
16 Cloco
4) col 16] suril;
Operation
Sepurate operations
Fused operations JIT
Operatio Puson comtines miliiple consecutive operations sindle uptiniuel kenel roduce memory uatsfess and ingime comprseime ellicinnyy on CUDA dosins.
ryython
fozward (aeif, xi:
Btozeh. acript
poolx_o
pose proczso bx, subcract_yvon): dis=!]
oelf sub cestrratte.eeed -1s
torch, naxit, dinelito1
torahh at gnggnddde)
cotuza Lax din-1)(0]
dof forvazd tseld, x):
P.aax P_sax_paci3ddis,
kernel pize- www..pool..s.....oooooooooooosss pot.soooooo_Leereessss.....
stritiemeet proi_
zuturn sest. process relffeeotrrcc..eeeii ei. 1.11 101
Table 12: (Part 3) Code snippets before and after optimizations.
24

---

Published as a conference paper at ICLR 2026
Toch + Desc
Before optimization
After optimization
Optimized Thrrad Block
Configuration
Busic throad block onliguration
Candulll luned thread black configuration
Pychca
"Python
Qptnized Block
block din (16, 161 Birgle equare thread
block din = 132, 31 Restangular block aptimizes For nateix multiplicatio
bicek
selecting back
qiid 1633
(nath.cell DJ 7 16), nath.cesl os
cridLsin = coit 221, nath. IN B)
dimensiore CUDA
maximize poralletio.
efficiny..
kecnel dgradggridd slockwsioocs_ding
kernei igricegrin blosteblonk.cim axga cata_ptri), Bidata_ptr(1,
and compunsionnll thmugh-
arga dt (.
por based oz the hootware achteeaure and algeritim
O. ptrll,
C.data pert), RC)
...
characerisiis.
Brunchles Implementatiio
Boachlos implencrtatin
eeplaces conditional stacmmnnts with mathemiical opecations avnid branch divergenen and improws GPP per peffrmmnce.
cotppe (-1.0, mninnet.rr,
out purpuu
clse outpus
Sharee Momory Uongs
Fach thnoad reatss disgonal element from glotal memory
Using shared eathe diegonal elements
Shard memory it CUD "tuddd
allows thrads with the
global ucld
globa wild
some block ahfennll dum
diag_nttaulxxeree wtwp_ueceecccccpiiicccooooo inperra
[oss_nattalkkeree optirized floar Au const floats floats C, inc Me lct 80
dal, roducing global memory
float. & coast float. B, floct> C, int
accesies impoovin petfor
fir algutillms
int zov bLockIdk.y biackoan.y
int HI I
cocat ins BLCCK-1REE3
8:
puleesa.
threadres.
354296 fleat ALatared (Labarrd,BA0cc__S5 F3 /7 Shared netory for disgonel elemente
tot co blockkdr.x blockhim,x
threedTex.2gg
Fox birectax.y Blocknin.y threacIdk.y)
diagonal
bleekIde.x
C[row col) Alvon] B[zov
1/ diogcnal alenenta lato abared ronce pe in blackt
it jthreadise, 6s row A_sharedfthhreaadd.... Alzow)
Lyaethheoda (1; // Ensure all throada aes the leaded volues
it frox
Use cached diaçonal elerant teon renory
atared coll;
Minimal Sunchronization
Defallt synchontsation behavlor
Minimal Synchmontration
Minined Sunchmnistion zt-
sx-pytson
dianse cectiedd by mininticing the zumber gunctronnnoon
def coxward (selif
forwará
Rach CUDA operation inplicitly
Create dedicated
ponns berwen CPU and CIPL operations allovins wsya- chropous exccution Blough dedicased sutums.
syschroniizz
kita tazch, cuta, www..uut....aaaaaooo. wtw.ennaele.ecooooooooooo fae.. worppcccaaaaaeeeoooooooe Operationa Esynchzonnusly ly in this
x.contiquaual II
reselz wcw..comw. cco.__trarspesccdf
strear
return reedit
x_opt inized contiiguu torme
tasch, chance ch jast
result
c.orv transposeldi optinired! Tegitel synchroni 7at only happens Vhen sesult used zocurn zesult
Table 13: (Part 4) Code snippets before and after optimizations.
25

---

Published as a conference paper at ICLR 2026
Refore qptimizationn
optinization
Thread Currsening
processes
processes
Thread Cousseting is
optimization techiique whzre each threat processes miltipe dats dencua issizdd of jus one. itereaiing arihmeio mirraity and maucing theaad
Ceacuro plzey
tx ieature_size
stomickkdigoe Teature_ett size
Saaturessiee +nn festore_slze
AnRignLvNt x_xal);
sealar_t x_ve12
"isture_oiee threacs_xi]
feature_alae
atsmcckddtsviadjk feature_pize_pediee d), assign_val *.valll;
atomicicdtieladd festure stre.pieeeeeeeeee thrmede sexign.ve x_val2);
rensining (Eageraalze
festure_size:
threada_xi
scalar_t X_vai rax_
testure_size tnstrre akdd:iiddisv feature size padded d1, aesign vel valli
Asynchronces Execution
Sequntial exeuiiion
- -Ayynchremus exccation with cindom stream
Sythea
"'yytoon
Asyadhmauus
der forward (eeci, 812
det forvect (se1s, x):
opeatims
Operations execute ie the cecauis stzean,
Create dedicated compute strean
blocking sequeatially
self cerggaat screar
curentl
result self, conon trorsgoseedde
torch.can wtt.hhccaaaa.aaaaa.oooo sts.ssreppprroorttteell
Righ
asabling eeetapping
return reeolt
priority at strees
with inneer
inproved GPU uilliaation.
Exccute operations asynchronocsly
wtw.t..pppinieeeeuuee x.dtype)
Contzel raturaa innediately while conputation continues In background veruen sesult
Table 14: (Part 5) Code snippets before and after optimizations.
26

---

Published as a conference paper at ICLR 2026
F CASE STUDY: COMPARING REFERENCE CODE AND CUDA-LI OPTIMIZED NEURAL NETWORK IMPLEMENTATIONS
F.1 LSTMS
Table 15: Reference code and CUDA-LI generation for LSTM class
LSTM Reference Code Simple baseline implementation
import torch import torch. nn as nn
class Model (nn. Module)
def init _ (self, input_sized hidder_size, num_ lamera output_size, dropout=0.0): 0):
Initialize the LSTM model. ***
super (ModeL, seit)._
10
# Initialize hidden state with random values
self.h0 borch.randh (num__ayeer, batch_size, hidden_size))
12
self.c0 corch.randa (num__aaers) batch_size
hidden_ (idden_sizz))
self.1stm = M..LSMiinuuuu_ieee hidden_size
num layers, batch_ first=True, dropout-dropout,
bidirectional=False))
14
self.fc nn. nn.Linear (hidden_size, output_size)
15
def forward (self, x):
18
Forward pass through the LSTM model.
19
self.b0 h0 self.h0. to (x. device)
self.c0 self.h0 .to (x.device) BUG: This should be self.co.to (xvieiice)
22
A Forward propagate LSTM out, state = feif.istm (x, (self.hC, self.c0)) shape of out: biatch_size, seq_length, sidden_size)
26
$ Decode the hidden state of the last time step out fo [i, -1, :l) # shape of out: (batch_size, output_size)
20
return state[0]
30
# Test code batch size 10
23
sequence sequence_Lengthh length 512
34 input_size 128
35 hidden_size = 256
nun_layers =
output_siza 10
38
dropout 0.0
39
def get set_inpus(() return storch.randh tbatch_size, sequen.ce_Length. input_size)
42
43 def get_init pet_init_inputset
return [input_size, hidden_aize, num_layers, output_size, dropout
27

---

Published as a conference paper at ICLR 2026
LSTM Fully Optimized Code All optimizations enabled (3.4x faster)
import torch
import torch.nn import torch.cuda cuda
class ModelNew (nn. Module):
def iiitt (self, input_size, hidder size, num layers, output_size drepout=0.0):
".n Initialize the LSTM model with three core optimization techniques.
Color coding:
11
BLUE: CUDA Graphs optimizaiion
GREEN:
Memory Contiguity optimization
13
ORANGE: Static Tensor Reuse optimization
super (ModelNew, self) (i.f)_initt_
16
17
# Initialize hidden states as buffers
self. registrresttrrbteee ('',', torch. randn ( (num (num_layers, batch_size, hidden_size)) sel ru c'',', torch. randn (num_layers, batch_size, hidden_size))
20
# Use PyTorch's optimized LSIM implementation
self 1stm In.LSTM
Lnput_size=lnpu__size,,,
24
hiddee_sile-hh.den_si..,,,
rum_layyrrs-nu__layezs, batch_first=Irie
dropout-dropout, bidireationai=Fallse
28
29
self. nn.iinear midden__size, output_sizer
CUDA GRAFHS: Variables for graph capture and replay self. None
35
self grw.h_araaprryy False
36
self. self.input_shappe None
STATIC TENSOR REUSE: Pre-allocates tensors for graph execution
39
self.stati self.static_inputt - None
self.static_outee None
42
CUDA GRAPHS: Streams for graph operations self. self.graph_strean = None
Track we' re running on CUDA
46
self seltt.s_cudd__vaiiioo = torch. cuda, is available ()
def
(aelf):
49
CUDA GRAFHS: Initialize CUDA stream for graph operations***
if self. grt.ghsrrh_sstaaan is None:
self. self.graph_streaat cuda. Stream()
52
28

---

the captured Published as a conference paper at ICLR 2026
def _capture_graph (seLf, result):
CUDA GRAPES: Capture the computation graph for replay
STATIC TENSOR REUSE: Create static tensers for graph capture
57
STATIC TENSOR
Clone tensors for static
allocation
19
selfistatic. sstaticinnnutttttt = x.clone() ()
sellfsstaiccooue = result.cloner
62
CUDA GRAPHS: Capture the computation graph with torch. wwd.e stream (self. graph__steam) self. cuda. CUDAGrupaa () with cuda graph (self.graph):
Operat ons to capture in the graph static_out, self. 1.ff..smmms) (self..taatic_inuus (self.h0, self.c0))
67
09
MEMORY CONTIGUITY: Ensure contiguout momory layout static_last static_out -1, :l elnen
72
self.statio www.potttuooo.....o...o (stattc
73
Wait for graph capture to complete wyrch.cuu..ssnccro....
76
77
Mark graph as ready for use self graph_ready True
def stwaddard
x):
" Standard forward pass with memory contiguity optimization""
MEMORY CONTIGUITY: Ensure input contiguous if not x.is_ sris__cotiigous ():
84
= x. contiqueus ()
Forward pass through (self.h0, self.c0)) out,
88
MEMORY CONTIGUITY: Make last output contiguous for optimal memory access
91
last_out out 1:, -1, :l conttiuuoss ()
return self fe (Last_oout)
95
def forward(self, x):
96
Forward pass through the LSTM model with three optimization techniques.
98
99
Opt imizatioo flow:
1. CUDA CRAPHS: Check is graph (fast path)
101
2. StatiC TENSOR REUSE: Use pre-allocated tensors
for graph replay
29

---

Published as a conference paper at ICLR 2026
3. MEMORY CONTIGUITY: Ensure optimal memory layout " Chroughout
104
CUDA GRAPHS: Fast path use captured graph if available
106
(x. isccd and sel g.raph_readd and x.shape self input shape)
STATIC TENSOR REUSE: Copy to pre-allocated tonsor with nen-blocking transfer
109
self.static__nnpt.cco copy tor_blocking-trux)
111
CUDA GRAPHS: Replay the captured graph sel ....ph....rppppp replay
112
Return the output from static buffer return wtw.ssstttccoutuu.ccloou
115
116
117
I Standard execution path
with torch.no_grad (): result self. staff_.sttndddddoooorardl
120
CUDA GRAPHS: Initialize graph on first CUDA input x.is x.is_cuda cuda and self cuda available and not self.graph ready:
122
123
try:
Store the current input shape self input_shape x.shape
126
CUDA GRAPHS: Initialize CUDA resources self www.tiaaliieeuuaaeeeeeeeeeee
129
CUDA GRAPHS + STATIC TENSOR REUSE:
Capture the graph
131
celf_ccaptue__graph (x, result)
132
133
except Exception as 0:
If graph capture fails, continue without
it
135
self self.graph_ready False
136
return resu_t
139
# Hyperparameters from the zeference implementation
140
batch_size - 10
141
sequence_length 512
input_size 128
hidden_siae
256
num Lum_layees 144
145
owtputtssiitt size = 10
146
dropout = 0.0
def get_iaputs(}- 148
149
return [torch. randn (batch_aize, sequencce_ennthh, input_size)
def get init inputs() 152
return linput_size, hidden_sied num layers, sstput_size, dropout]
30

---

Published as a conference paper at ICLR 2026
153
Example usage demonstrating the three techniques name main
155
156
import time
157
print
BLUE: CUDA Graphs optimization")
159
print ("
GREEN: Memory Contiguity optimizaticn")
print
ORANGE: Static Tensor Reuse optimization")
161
print ("=" 60)
162
4 Create model
model ModelNew (*get_ inpuustttttttttiiieeee model.eval ()
166
167
$ Test input
get_inputs| () [0]
170
Move to GPU if available if carch.cuda.c is isavaiiiaa :): model = model. ouda
171
cudu
175
print ("Running on CUDA all three optimizat ions active")
First captures graph
178
print (n First forward pass: Capturing CUDA
graph...")
with srrch.no_gradil
180
output model (x)
181
print Output shape: poutput.shapes
182
print (f" Graph ready:
Subsequent
captured graph
185
print ("An Subsequent passes: Using captured graph
with")
186
print static tensor reuse and memory contiguity")
warmup
for in range (10)
190
with torch ()) model (x)
191
Measure performance
torch. cuda, synchronize
195
start_event torch. cuda torch.cuda. Event (enable_ (imbgggrettnn
196
ead_event = torch. cuth. (enable weabbblettiinng-rrte
n_runs 100
100
start_event. record with torch.no [r
200
for in range (n runs):
output model (x) und_event..rcord
203
204
205
torch. cuda. worch.cuda.syncheeee... synchronize
avg_time start_event wtw.eeeeeeeellaseddddee.e.ee wwwpsedd time (endeeeent)))) runs
208
print (E" PPerforraae:: (avg_time: 3f) ms per forward pass")
31

---

Published as a conference paper at ICLR 2026
209
print (E"
Expected speedup: -3.42x with all optimizations")
210
211
else:
print ("And Running CPU only memory
contiguity active")
213
print (" (CUDA graphs and static tensor reuse require GPU) ")
with torch.no_grad
216
output model (x)
217
print (f"An Output shape: loutput. shape) ))
32

---

Published as a conference paper at ICLR 2026
F.2 3DCONV
Table 16: Reference code and CUDA-LI generation for Conv3D class
Conv3D | Reference Code Simple baseline implementation
import torch
import torch.on nn
class Model (nn Modull)
Model that performs 3D convolution, applies Group Normalization, minimum, clamp, and dropout.
.*.
(att_(sef,, in_channels, out out_channels, kernel_size, groups, min_value, max_va max_value, dropout_p): super (Model, seif) iniit_t self.conv cn..oonvdiinn kernel_size)
out_channels,
11
self.norm self self.dropout nn nn.Dropout (dropout_p)
CaroupNorm(gropppp out out_channela)
12
self self.min_value min_value
delf.max_value
value
15
16
def (sel, x):
17
x self.conv (x)
self. (x)
norch.mint (x, torch. (wwlffmmnnn.eeeeeeoooo... meolf.min_vaauc)
20
x forch.clamp (x, min-self..inn_aillll
max-self. )
self.dropout (x)
return
24 # Hyperparameters 35 batch_size - 128
in_channels out_channels 16
depth, height, width 16, 32,
29 size =
groups =
min_value 0.0
max_value 1.0
33
dropout_p 0.2
def get_inputs() ():
return [torch. randn (batch_aize, in_charnels, depth, beight, width)]
37
def get init inputs()
return lin channels, out channels, kernel size, groups, min_value, max_value, dropout_p]
Conv3D Fully Optimized Code All optimizations enabled (120x faster)
import torch
import torch.nn
# Hyperparameters sattch_size = 128
in_channels
out out_channels
depth, height, width 16, 32, 32
33

---

Published as a conference paper at ICLR 2026
kernei_size = groups =
11 value 0.0
12 max_value 1.0
13
dropout_P 0.2
class ModelNew (nn.Module):
16
def (sit_,seef, in_channels, out out_charrels, sernel_size, groups, min_value, max_value, dropout_p): super (ModelNew, .i_f)__iiit_
17
$ Store the original Layers for parameter compatibility self.conv (nn out_channels, kernel_size)
19
20
self.norm
(GroupNo:m(groupp)) out_channels)
self drooout nn Dropout (dropout_p)
melf.min_vaiie min_value
_aa__va nax_value
24
self.dropout = dropout_P
TECH Mathematical Short-Cireuit Optimization
27
# Detects whon min_value=O, . to skip entire computation
28
self. welffuue_opptine use_optiimiee__ptt
(min_value == 0.0)
30
TRCH 4: Pre-computed Convolution Parameters
31
$ Extract and store cony parameters once during
initialization
if isinstance (kernel size, int): self.kernel (xernel_size, kernel_size, kernel_size)
else:
self.kerner size kernel_size stride sel seif.conv.stride self paddinn conv, cow..addiin
37
38
self.dilation = self. www..coovv..l....
40
TECH 4: Pre-compute output dimensions for
standard input
41
self. colffouu_depth = (depth + 2 self self.padding -
self.dilation[0]] keelf..erme] 1) 1) // self.
42
self out_ffeutthhhtttt (heighh + + self.padding(1]
self.dilations
keelf..kermi_ssielll 1) 1)
// self. + +1
self. out_wiottwidth (width
self.padding(2)
self.ddiation1221 (self. kere.. 1) // self. +
44
# Standard output shape for the default batch size
self staadard_ (batch_size, out_channels,, self. bout_dephhh self.out height,, self.cut. width)
47
TECH Pre-allocated Zero Tensors Create zero tensors once avoid allocation overhead if self wtt..uusooootddeedpeeeee (' z'zzero_ooo
50
torch.zeros (arerrsssss
stype=torch.flcat3333
53
persistent=walse)
54
relitregister_biie buffer ('zero zero 2etteeeooooutttie_e_ee output_floatlll
torch. zeros saaass shape,
persistent=Palse)
34

---

Published as a conference paper at ICLR 2026
selt.register_ wwt.reertttttooottttttoooo "eto_ootpit__tilat11"
torch. zeros (se..ttaaa.. dtype=terch..stlcaa1)) peraistent=Palse)
59
def calculate_owe pap (self, input_shape); """Calculate the output shape of the convolution
operation. ***
63
aetch_sized _ d, h, w = input_shape
TECH Use precomputed parameters # Avoid repeated attzibute lookups out_d = ((d 2 self self.padding[O) - self.dilation (self.kernel_siag)
66
1) // self.stride [0])
out (th 2 seli.padding [1] self.dilation |1] (self.kermel_sine size[1] 1) 1) // self.stride [1])
+ +1
out_w ((W 2 self.padding self.dilation[2] (self.
// sclf.stride
+ 1
70
return (batch_size, wwl.f.covv.ott_caeee out_d, out_h, out_x)
73
def forward x):
TECH Mathematical Short-Cireuit Main optimization
73
skip all computation when we know result will be zeros
if not self.use_ov wttpuuse_ptiiiee_pppppep
Standard path for non-optimized
78
self.conv (x)
79
x = self.norm(x)
= torch. (x, torch.tensor (self.min miin_vluee,
device=x.aevicel)
torch.cl an (x, min-self mm fii__vaaee,,
max-sel mat__mxxxlaaaaaa
82
-
self.dropout (x)
return
Optimized path when min_vallt value 0.0 Since min(x, 0) followed by clamp (C, 1) always produces 2cros
86
88
TECH 3: Direct Shape Matching Fast path for standard input dimensions if x, ahappe -- (batch_size, in_channels, depth, height, width):
91
TECE 2: Use pre-allocated tensors
Return pre-allocated zeros matching input dtype
if x.dtype torch.float.33.
return scif..cco__ottpit__falaattt
95
elif x. x.dtype == torch.
96
return celf..ec:_outpuu_fllooo
elif x.dtype =- torch.bfloatie
return sei F.rerro_oupptt_bflott
99
else:
100
Fallback for other dtypes
return torch. zeros (self. standard. shape,
device=x.devie stype=x.ctype
else:
35

---

Published as a conference paper at ICLR 2026
103
# For non-standard input shapes, calculate output shape
104
aaaaluulate_utppii shape (x. shape) return torch, zeros (output_shape, deviceex.device. dtype-x.ctype)
pntput_shapp
105
def get ():iipuuss():
108
return [torch. randn (batch (batch_size, size, in channels, depth, height, width)]
109
def get_init_inputs//
111
return [in_channels, out_channels, kernel_size, groups, min_valus, max_value, drepout_p]
112
Color Legend:
114
TECH Mathematical Short-Cireuit (Blue) Skips computation when min value00
TECH Pre-allocated Tensors (Purple) Pre-allocates zero tensors
TECH 3: Direct Shape Matching (Green) Fast path for standard shapes
117 TECH 4:
Pre-computed Parameters (Orange) Pre-computes conv paraneters
36


Reviews:
[{"review_id": "TO2Byi7bcn", "reviewer": "Reviewer_paTv", "summary": "The paper introduces CUDA L1, a three stage system (SFT followed by self supervised filtering followed by contrastive RL) to optimize CUDA code, where the RL stage injects scored prior variants into the prompt to force comparative reasoning and then updates the policy with a GRPO style objective. On KernelBench (250 tasks), the authors report 3.12× mean speedup over the original Pytorch implementation and consistent improvements over torch.compile variants.", "questions": "•\tWhat are the full training hyperparameters and compute for each stage? \n\n•\tWhat fraction of improvements stems from kernel level vs graph level vs framework level optimizations? \n\n•\tCan you include TVM/Ansor/Triton comparisons?", "limitations": "", "rating": 2, "confidence": 3, "soundness": 2, "presentation": 2, "contribution": 2, "strengths": "•\tStrong speedups on a community benchmark and across GPUs. \n\n•\tClear prompt/exemplar design and informative case studies.", "weaknesses": "•\tOne of my concerns with the paper is that it does not guarantee the correctness of optimized code. Optimizing code to generate highly optimized CUDA code is well studied in the compiler literature, and many techniques that guarantee the generation of highly optimized, correct code exist. The fact that this proposed method does not guarantee the correctness of the generated code is an important limitation.\n\n•\tAutomatic code optimization is well studied in compilers. While the comparison with Pytorch torch.compiler is interesting, the paper should also discuss and compare to other state-of-the-art compilers such as TVM, AutoTVM, Ansor, and Triton, which are the de facto baselines for automated GPU optimization.\n\n•\tSeveral wins stem from graph level optimizations (CUDA Graphs, static reuse, stream control) rather than low level kernel transformations. This is fine in practice but should be made explicit. Consider reporting the proportion of gains due to kernel level vs. graph level vs. PyTorch level changes.\n\n\n•\tContrastive prompts with scored exemplars are compelling, but the policy optimization is standard GRPO. Please clarify more what is algorithmically new.\n\n•\tThe authors compare to the Pytorch implementations of KernelBench, but it is unclear what backend library is used to accelerate the Pytorch operators. Are the operators accelerated on GPU ? Which library is used to accelerate the operators? Specifying this is very important since this constitutes the baseline and the difference in performance between the different variants is significant.\n\n\n•\tThe paper lacks critical training and compute details (GPU hours, steps, batch sizes). Without these, reproducing the reported speedup is unlikely.", "quality": null, "clarity": null, "significance": null, "originality": null, "strengths_and_weaknesses": "", "was_revised": false, "final_justification": "", "rebuttal": "", "claims_and_evidence": "", "methods_and_evaluation_criteria": "", "theoretical_claims": "", "experimental_designs_or_analyses": "", "relation_to_broader_scientific_literature": "", "essential_references_not_discussed": "", "other_strengths_and_weaknesses": "", "other_comments_or_suggestions": "", "questions_for_authors": "", "summary_of_paper": "", "summary_of_recommendation": "", "technical_quality": "", "clarity_of_presentation": "", "potential_impact": "", "robotics_focus": "", "quality_of_limitations": "", "questions_for_rebuttal": "", "reasons_to_accept": "", "reasons_to_reject": "", "questions_to_authors": "", "extra_scores": {}}, {"review_id": "BpVjP65Wsq", "reviewer": "Reviewer_rehX", "summary": "This paper tackles CUDA code optimization by fine-tuning a large language model through a sequence of supervised fine-tuning (SFT), self-supervised learning, and contrastive reinforcement learning. The authors first construct a dataset for SFT by prompting diverse LLMs with reference PyTorch code from the KernelBench dataset to produce optimized code variants of the original dataset. They then SFT the model on this corpus and subsequently apply self-supervised learning; the model generates code and successfully executed code samples are fed back for further training. Finally, they employ contrastive reinforcement learning driven by code performance scores to refine the model’s preferences toward faster implementations. The resulting model achieves a 3.12× speedup over the KernelBench code.", "questions": "- Is the reward function defined solely in terms of execution time? Are there any components accounting for memory usage or other resource efficiency metrics?\n\n- Why was the KernelBench dataset chosen for this study? KernelBench is designed to evaluate whether an LLM can generate CUDA code from given PyTorch implementations, not to optimize the PyTorch code itself. Therefore, comparing CUDA-L1’s results with the vanilla KernelBench code and reporting speedups appears to be an inappropriate or uninformative evaluation.", "limitations": "", "rating": null, "confidence": 4, "soundness": 1, "presentation": 2, "contribution": 2, "strengths": "The authors proposed a contrastive reinforcement learning strategy that incorporates reward scores into prompts, potentially enhancing the LLM’s ability to generate optimized code on KernelBench tasks.", "weaknesses": "- LLM is both trained and evaluated solely on the same reference code from KernelBench without splitting the training/test dataset. We cannot verify that the finetuned model can create faster optimized versions of code for other tasks or datasets.\n\n- The comparison between CUDA-L1 and the baseline is unfair. CUDA-L1 is fine-tuned on the code that is known to be successfully compiled and executable, benefiting from multiple inference rounds, whereas other baseline LLMs are limited to only five attempts. Since CUDA-L1 requires additional training costs, a proper assessment of the training method should compare it against alternative fine-tuning strategies, not just untuned baselines.\n\n- There is no direct evaluation of the generated CUDA kernels; only the PyTorch code from the KernelBench is measured. Although the end-to-end evaluation may implicitly include the performance of the CUDA kernels, it is not possible to isolate and assess their contribution. The performance improvements observed could instead originate from optimizations in the surrounding Python code (e.g., removing Python overhead). Consequently, the experiment does not provide evidence that the fine-tuned model can optimize CUDA kernels themselves.\n\n- The only novel element in the finetuning pipeline is prompting with a reward score. The methodology does not encode any prior knowledge of CUDA (or PyTorch) into the LLM. Therefore, it seems largely unrelated to CUDA optimization, contrary to its title.\n\n- The authors use terminology that could mislead readers. The second stage of the fine-tuning pipeline should not be described as self-supervised learning but rather as supervised fine-tuning, since failed trials are excluded from the training data. Because this process depends on feedback from an external execution environment, it is more appropriate to regard it as a form of labeling. Moreover, the fine-tuned LLM does not generate optimized CUDA kernels; it produces optimized PyTorch code written in Python, which is fundamentally different from generating efficient CUDA implementations.", "quality": null, "clarity": null, "significance": null, "originality": null, "strengths_and_weaknesses": "", "was_revised": false, "final_justification": "", "rebuttal": "", "claims_and_evidence": "", "methods_and_evaluation_criteria": "", "theoretical_claims": "", "experimental_designs_or_analyses": "", "relation_to_broader_scientific_literature": "", "essential_references_not_discussed": "", "other_strengths_and_weaknesses": "", "other_comments_or_suggestions": "", "questions_for_authors": "", "summary_of_paper": "", "summary_of_recommendation": "", "technical_quality": "", "clarity_of_presentation": "", "potential_impact": "", "robotics_focus": "", "quality_of_limitations": "", "questions_for_rebuttal": "", "reasons_to_accept": "", "reasons_to_reject": "", "questions_to_authors": "", "extra_scores": {}}, {"review_id": "PayCGwPvAn", "reviewer": "Reviewer_1yB6", "summary": "This paper introduces CUDA-L1, a multi-stage training pipeline designed to enhance large language models’ (LLMs) capability to generate efficient CUDA kernels. The proposed framework comprises three stages:\n- Supervised Fine-Tuning (SFT): The model is trained on synthetic data generated by various state-of-the-art LLMs to increase its exposure to diverse CUDA programming patterns.\n- Self-Supervised Learning: The model is further refined by leveraging its own successful generations as additional training data, thereby improving the overall success rate.\n - Contrastive Reinforcement Learning (RL): The model is trained with contrastive rewards applied both at the parameter level and within the LLM context, targeting execution efficiency.\nThe first two stages primarily aim to maximize correctness and success rate, while the final stage focuses on improving kernel execution speed.\n\nEmpirical evaluations demonstrate that CUDA-L1 achieves significant acceleration and near-optimal success rates compared to both human-designed and LLM-based compiler baselines. Moreover, the method shows strong generalizability across different GPU architectures.", "questions": "- How is the mean acceleration metric in Table 1 computed? Is it calculated as the mean of per-task acceleration ratios, or as the ratio of mean run times between CUDA-L1 and the baseline? The former could potentially exaggerate the metric, as individual acceleration ratios may vary widely (e.g., from 0 to 100).\n- When transferring the model to other GPUs, such as H100, is it necessary to re-run the training pipeline, or can the trained model generalize directly?\n- Given that the model is primarily trained in the context of A100 GPUs, does it demonstrate the ability to adapt to new hardware functionalities, such as Tensor Memory Access (TMA) units in H100?\n- Are there any plans to open-source the model or training pipeline, allowing the community to build upon this work?", "limitations": "", "rating": 8, "confidence": 3, "soundness": 3, "presentation": 4, "contribution": 3, "strengths": "- The paper is well-written and easy to follow, presenting all critical information clearly in the main text.\n- The proposed approach, particularly the contrastive RL stage, is well-motivated. Training the LLM to reason about previous code performance is a thoughtful design, and ablation studies demonstrate its effectiveness.\n- Outstanding performance: Empirical results show near 3× improvement on KernelBench with a near-perfect success rate, highlighting the method’s substantial practical impact.\n\n- Comprehensive and transparent analysis: The evaluation and supporting analyses are detailed and informative, with some examples:\n  - Detailed performance comparison: Main results include detailed distributions across different difficulty levels in KernelBench.\n  - Ablation studies: Systematic studies illustrate the contributions of each stage in the training pipeline. Comparisons with straightforward alternatives (e.g., evolutionary methods or simple RL such as Stage 1+2+GRPO in Table 2) confirm the superiority of the proposed pipeline.\n  - Cross-hardware evaluation: Experiments on diverse hardware demonstrate consistent success rates, emphasizing the generalizability advantage of LLM-based methods over traditional human-designed compilers.\n  - Transparency: The appendix provides a lot of details, e.g. examples of reward hacking, showing the model attempting to manipulate profiler timers via asynchronous streams.\n  - Case studies: Illustrative examples in the appendix demonstrate how the model improves CUDA code generation in practice.", "weaknesses": "- It is not entirely clear how the model is evaluated. Specifically, does the evaluation involve iterative refinement over multiple reasoning steps, or is it conducted in a one-shot generation manner? Clarifying this would help assess the robustness of the reported results.\n\n- In Table 1, the average acceleration of nearly 3× appears inconsistent with the P50 speedup of only ~1.2× for the CUDA Graph rows. This raises concerns that the reported average may be skewed by a few poor baseline implementations, resulting in an inflated speedup. Additionally, for Level 2 and Level 3 tasks, which involve more complex operations, the P50 and speedup ratios do not seem to approach a neutral baseline. However, I do not have a better idea on how to present the distribution of acceleration to be more convincing.", "quality": null, "clarity": null, "significance": null, "originality": null, "strengths_and_weaknesses": "", "was_revised": false, "final_justification": "", "rebuttal": "", "claims_and_evidence": "", "methods_and_evaluation_criteria": "", "theoretical_claims": "", "experimental_designs_or_analyses": "", "relation_to_broader_scientific_literature": "", "essential_references_not_discussed": "", "other_strengths_and_weaknesses": "", "other_comments_or_suggestions": "", "questions_for_authors": "", "summary_of_paper": "", "summary_of_recommendation": "", "technical_quality": "", "clarity_of_presentation": "", "potential_impact": "", "robotics_focus": "", "quality_of_limitations": "", "questions_for_rebuttal": "", "reasons_to_accept": "", "reasons_to_reject": "", "questions_to_authors": "", "extra_scores": {}}, {"review_id": "UIFfzOzHCt", "reviewer": "Reviewer_hU7V", "summary": "This paper proposes CUDA-L1, a comprehensive framework for automating the optimization of CUDA kernels using large language models (LLMs) enhanced through contrastive reinforcement learning. The framework targets the challenge of improving GPU computing efficiency by enabling LLMs to autonomously generate optimized CUDA implementations that are both correct and performant.\nCUDA-L1 adopts a three-stage training pipeline—supervised fine-tuning with data augmentation, self-supervised refinement, and contrastive reinforcement learning—to progressively enhance the model’s ability to reason about, generate, and evaluate CUDA code. Through comparative learning, the model integrates execution feedback directly into its reasoning process, allowing it to learn from past examples of successful and unsuccessful optimizations.\nThe system is evaluated on a large and diverse set of CUDA kernels across multiple GPU architectures. The paper further contributes a set of CUDA Graph implementations for KernelBench, enhancing community baselines.", "questions": "- Given that the paper's core contribution is framed as \"CUDA optimization\" but the evidence suggests most improvements come from higher-level framework patterns, how do you distinguish between these two distinct types of optimization in your evaluation methodology?\n- Since CUDA evolves over time, how would CUDA-L1 handle updates to the CUDA toolkit? Would it require full retraining, or does the model have mechanisms to adapt to new features incrementally?\n- The paper builds upon deepseek-v3-671B. How critical is this specific large-scale model to CUDA-L1’s performance? Have experiments been conducted with smaller open-source models, and if so, how does performance degrade?\n- The case studies show the model discovering both high-level algorithmic improvements and low-level CUDA implementation optimizations. Does CUDA-L1 include an explicit mechanism to decide the level of abstraction for optimization, or does this behavior emerge naturally from the RL process?\n- To aid reader understanding, could the authors include a concrete example (e.g., a filled-in mid-training prompt) in the appendix? While Table 9 outlines the general format, an actual example with real—possibly truncated—code snippets and corresponding performance scores would better illustrate how comparative exemplar guidance works in practice.\n- Will the training code, reward-checking implementation, or exemplar-guided prompts be released to facilitate replication by other groups?", "limitations": "", "rating": 4, "confidence": 4, "soundness": 2, "presentation": 3, "contribution": 2, "strengths": "- The extensive evaluation across all 250 KernelBench kernels—covering multiple GPU architectures and varying complexity levels—provides strong empirical validation of the proposed approach.\n- The enrichment of the KernelBench dataset with CUDA Graph implementations represents a valuable and reusable contribution to the research community.\n- The identification of non-obvious optimization techniques (e.g., mathematical short-circuiting) highlights the potential of reinforcement learning to uncover previously unexplored performance improvements.", "weaknesses": "- The paper conflates CUDA kernel optimization with PyTorch-level program optimization. Case studies such as LSTM and 3D convolution reveal that the “optimized” code primarily alters high-level operator invocation and execution patterns, leaving the underlying CUDA kernels unchanged and thus creating a misleading impression of kernel-level innovation.\n- While the paper demonstrates impressive success and speedup rates, it lacks an in-depth analysis of cases where the optimization failed or underperformed.\n- The work does not discuss the computational resources required to train the CUDA-L1 system. Given the multi-stage training and RL fine-tuning involved, such costs could be substantial.\n- The paper focuses heavily on empirical results but provides limited theoretical explanation for why the contrastive-RL approach outperforms alternatives.\n- Although the paper claims that the reward checking model “successfully identifies reward hacking over 60% of the time,” it does not quantify how much this detection reduced actual reward hacking, nor whether the system can recognize previously unseen hacking strategies.\n- While the internal baselines and ablations are comprehensive, the work lacks direct comparison with external, published state-of-the-art methods.", "quality": null, "clarity": null, "significance": null, "originality": null, "strengths_and_weaknesses": "", "was_revised": false, "final_justification": "", "rebuttal": "", "claims_and_evidence": "", "methods_and_evaluation_criteria": "", "theoretical_claims": "", "experimental_designs_or_analyses": "", "relation_to_broader_scientific_literature": "", "essential_references_not_discussed": "", "other_strengths_and_weaknesses": "", "other_comments_or_suggestions": "", "questions_for_authors": "", "summary_of_paper": "", "summary_of_recommendation": "", "technical_quality": "", "clarity_of_presentation": "", "potential_impact": "", "robotics_focus": "", "quality_of_limitations": "", "questions_for_rebuttal": "", "reasons_to_accept": "", "reasons_to_reject": "", "questions_to_authors": "", "extra_scores": {}}]

Author rebuttal:

Decision: Accept (Poster)