Kevin: Multi-Turn RL for Generating CUDA Kernels
Authors: Carlo Baronio
Pietro Marsella
Ben Pan
Simon Guo
Silas Alberti
Venue: ICLR 2026 Poster (2026)
Keywords: multi-turn
RL
GPU kernel
code generation
Primary area: foundation or frontier models, including LLMs

Abstract: Writing GPU kernels is a challenging task and critical for AI systems' efficiency. It is also highly iterative: domain experts write code and improve performance through execution feedback. Moreover, it presents verifiable rewards like correctness and speedup, making it a natural environment to apply Reinforcement Learning (RL). To explicitly incorporate the iterative nature of this process into training, we develop a flexible multi-turn RL recipe that addresses unique challenges encountered in real-world settings, such as learning from long trajectories and effective reward attribution across turns. We present Kevin - K(ernel D)evin, the first model trained with multi-turn RL for CUDA kernel generation and optimization. In our evaluation setup, Kevin shows significant gains over its base model (QwQ-32B), improving correctness of generated kernels (in pure CUDA) from 56% to 82% and mean speedup from 0.53x to 1.10x of baseline (PyTorch Eager), and surpassing frontier models like o4-mini (0.78x). Finally, we study its behavior across test-time scaling axes: we found scaling serial refinement more beneficial than parallel sampling. In particular, when given more refinement turns, Kevin shows a higher rate of improvement.

Published as a conference paper at ICLR 2026
KEVIN: MULTI-TURN RL FOR GENERATING CUDA KERNELS
Carlo Baronio" Stanford University Cognition AI carlo@coonitiion.ai
Pietro Marsella" Stanford University
Ben Pan" Stanford University Cognition AI ben cenncognition.dc
marsel aaeelll@sstinnn fa@stanfordeddu
Simon Guo
Silas Alberti Cognition AI
Stanford University
simongudsstaaford.eeet
silas@cogmition.co
ABSTRACT
Writing GPU kernels a challenging task and critical for Al systems' efficiency. It is also highly iterative: domain experts write code and improve performance through execution feedback. Moreover, presents verifiable rewards like cor- rectness and speedup. making it natural environment to apply Reinforcement Learning (RL). To explicitly incorporate the iterative nature of this process into training, we develop flexible multi-turn RL recipe that addresses unique chal- lenges encountered in real-world settings, such as learning from long trajectories and effective reward attribution across turns. We present Kevin K(ernel D)evin, the first model trained with multi-turn RL for CUDA kernel generation and op- timization. In evaluation setup, Kevin shows significant gains over its base model (QwQ-32B), improving correctness of generated kernels (in pure CUDA) from 56% 82% and mean speedup from 0.53x to 1. 1.10x of baseline (Py'oocch Fager), and surpassing frontier models like 04-mini (0.78x). Finally, we study its behavior across test-time scaling axes: we found scaling serial refinement more beneficial than parallel sampling. In particular, when given more refinement turns, Kevin shows a higher rate of improvement.
INTRODUCTION
2025) in domain-specific Writing efficient GPU kernels (Dao et al. 2022 Zhao et al. 2025 Ye et al.
languages: CUDA. Triton, ThunderKittens, CUTLASS, etc. (Nickolls al. 2008 Tillet al. 2019 Spector et al. 2024 NVIDIA Corporation| 2025) is critical for enabling Al systems' efficiency at scale, yet remains dilficult and costly due to the deep domain expertise required. This has led to surge of interest in exploring how Large Language Models (LLMs) could help generate GPU kernels (Ouyang et al. 2025 Li et al. 2025 2025) using agentic systems (Damani et al., 2024 Chen et al. 2025 METR| 2025 Lange et al. 2025) Google DeepMind, (025) that leverage extensive test-time compute. These inference-basee
approaches are inherently limited by the base models' capability in this domain. On the other hand, the presence of verifiable rewards in the form of correctness and speedup against a reference implementation makes reinforcement learning (RL) a natural approach. This leads our investigation: How can we train model using RL solve the real-world engineering task of CUDA kernel generation?
GPU kernel generation emphasizes not just functional correctness, but more importantly performance distinguishing this code optimization problem from binary-reward tasks that involve passing unit tests (Jimenez al. 2024) producing acceptable proof (Zheng et al.) 2022). Since speedup is continuous goal, performance engineers take an iterative approach: they conduct many rounds of optimization based on previous kernel code, its execution result, and timing profiles. Hence, arriving at an optimized solution relies on multiple turns conditioned on previous exccution feedback. In
"qqual contribution

---

Published as conference paper at ICLR 2026
Training
Turn
Turn 2
Turn 3
Turn
Step Single Trajectory
Instruction
Instruction
K1
CoT' Evel
Instruction
K2 ColT Eval
Context Construction
Instruction
K1
Col'+ Eval
K2
Coll Eval
V
Inference
Model
Model
Model
Model
V
Kernel (K1)
Kernel 2 (K2)
Kernel 3 (K3)
Kernel (K4)
Execution & Evaluation
Correct: X
Correct: E
Correct: E
Speedup: 1.6x *
Correct:
Speedip:: N/A
Speedup: 4x
Speedup: N/A
Scoring
S1
Per-Turn Reward
R
S1
S2
S3
S4
R
S2
S3
S4
R
S3
S4
R
S4
Figure I:
Within each training step, the model iteratively generates, executes, and refines kernels over multiple turns. Kernels are rewarded individually, based both on their performance and their contribution to subsequent speedups: K1, for example, while incorrect, leads to both correct, slow kernel, K2, and correct, performant kernel, K3, and should thus be rewarded accordingly. This setup enables Kevin learn advanced code generation strategies that span multiple turns. Note: CoT' is the summarized chain of thought (CoT)).
contrast, popular RL methods to train LLMs verifiable rewards (Shao et al., 2024 Lambert et al., 2025) rely on the outcome reward single turn ("single-turn RL training") ) We hypothesize that explicitly incorporating successive turns of code generation, exccution, and feedback into each RL training step ("multi-turn RL training") better mirrors the iterative nature of kernel development, helping the model to learn more advanced code strategies that span multiple refinement turns.
We design a simple yet effective multi-turn RL training recipe, shown Figure that addresses the key challenges of training for CUDA kernel generation and optimization:
1. Long trajectories lead sparse rewards and context explosion. To improve sample efficiency, we split trajectories and use each turn as an individual training sample. To address context explosion from long CoTs while preserving reasoning information, we summarize CoTs of prior turns.
2. Finding optimal solution may require rewarding suboptimal kernels that eventually lead to more performant ones. Therefore, we study approaches to aggregate intermediate rewards across turns, finding configuration that balances the correcteess-eefforaance trade-off.
3. Reward hacking is prevalent as kernel generation is open-ended, real-world engi- neering task: e.g. the model can trick the evaluation harness, lazily copying the reference implementation instead of actually implementing kernels. To prevent this, we analyze the model's failure modes and enforce strict rule-based checks.
Enabled by our multi- murti-turn RL training method on 180 KernelBench tasks from Level 1 and 2, we present K(ernll D)evin, the first RL-trained model generate CUDA kernels. We compare Kevin and other models in our evaluation setting on representative KernelBench eval set. Kevin improves upon its base model QwQ-32B, (Team] 2025db both correctness (56% - 82%) and mean speedup of generated kernels (in pure CUDA): from 0.53x 1.10x over PyTorch Eager, while outperforming frontier models like OpenAl o4-mini (0.78x).
We then study the characteristics of Kevin in test-time scaling setting, comparing it to single-turn RL-baseline. We systematically compare the benefits of scaling along two axes of test-time compute: sequentially with more refinement turns (Ehrlich et al. 2025 Wang et al.) 2025a) or in parallel with more trajectories (Brown et al. 2024) (nell et al. (024). D. our setting, we find that sequential scaling much more effective, highlighting the importance of iterating upon execution feedback. We observe
2

---

Published as a conference paper at ICLR 2026
that the model trained with multi-turn RL exhibits better scaling characteristics with more refinement turns, compared to the base model and the single-turn RL baseline. Our core contributions include:
1. We design an effective yet flexible multi-turn RL training strategy that significantly improves model's capability on CUDA kernel generation. This strategy addresses challenges that arise in real- wollwwrldddd settings, and may be applicable to other environments that benefit from iterative optimizations.
2. We found multi-turn is more effective both for training and inference through systematic ablations: the multi-turn trained model outperforms the single-turn trained model across different evaluation setups. Furthermore, we found multi-turn inference is more effective across both models under fixed inference budget.
3. Kevin exhibits strong test-time scaling frends on both serial and parallel axes, with faster rate of improvement than its single-turn RL counterpart and its base model, while maintaining exploration capacity.
BACKGROUND AND RELATED WORK
2.1 LLM FOR GPU KERNEL OPTIMIZATION
There has been surge of interest in exploring how to leverage LLMs to generate GPU kernels (NVIDIA 2025), driven by the high cost and the long engineering cycles required to develop them (c.g. 2 years for efficient FlashAttention (Dao 2023) port after Hopper GPU relcase). However, frontier models underperform on representative benchmarks like KernelBench (Ouyang et al., 2025) and TritonBench (Li et al. 2025), likely due to GPU code being severely underrepresented in the training data (CUDA, for example, accounts for less than 0.1% of pretraining data in the Stack (Kocetkov al. 2022 Li al., 2023)). Collecting more expert- written code is expensive, only limited number of developers are able to implement high-quality kernels. To tackle this task, there has been an explosion of agentic systems (Damani al. 2024 Chen al. 2025 METR 2025 with
custom workflows and evolutionary search methods (Lange et al. 2025 Google DeepMind 2025).
Yet these approaches typically incur high inference cost e.g. $15 per kernel (Lange et al. 2025). Improving the base L.MM's boost the efficiency for downstream agentic workflows.
ability is therefore essential and could significatlly
2.2 RL OPTIMIZATION FOR LLMS TARGETING VERIFIABLE DOMAINS
2024 have been shown significantly enhance I.MMs' performance on verifiable domains ( (Lambert et al. 2025) such as math (Tem
Reinforcement Learning techniques like GRPO (Shao et al.
2025c Wang al. 2025b) and competitive programming (Team 2025d Luo et al. 2025a These approaches can be further adapted for real-world software tasks, using fine-grain unit tests (Liu et al.
2023) or comparisons between code edits (Wei et al. 2025) outcome rewards. Existing methods for code optimizations where objective concerns performance beyond correctness have been largely confined to supervised fine-tuning (Waghjale et al., (024)) and imitation learning (Shypula
et al. 2024). highlighting Kevin's 's RL approach a novel contribution for this setting.
Given that tasks like performance optimization or long-horizon planning require multiple sequences of interrelated actions, several works (Goldie et al. 2025 Cao et al. 2025) Wang et al., 2025c Zhou
(tt al., 2024) Zhhang* al., 2025) have explored RL training for multi-tun optimizations beyond optimizing for outcome from a single turn. Specific for the code setting, RLEF (Gehring al. 2025)
frames code generation as a multi-turn RL task: the model is allowed a fixed number of refinements
turns and assigned a single binary pass/fail reward for final generation training with such an approach might present
issues. Unlike RLEF, which assigns rewards only at the final turn, our multi-turn RL framework for Kevin trains every regardless of how optimal the code is, and optimizes for performance beyond just correctness. It is worth noting that Kevin's multi-turn RL training could be viewed as a variant of Meta-Learning (Xiang et al. 2025 Duan et al. 2016) or In-Context Reinforcement Learning (Nie et al.) 2024 Tajwar et al. 2025 Schmied et al. 2025). where the focus is to improve solution quality during test-time with feedback (Qu et al. 2025): but adapted in a novel way to the challenging real-world setting of GPU kernel generation and code optimization.
3

---

Published as a conference paper at ICLR 2026
3 TASK AND BASELINE
3.1 ENVIRONMENT AND EVALUATION
We use KernelBench (uuyag al. 2025). popular dataset for evaluating the LLMs' ability to generate CUDA kernels for deep learning workloads in PyTorch. We chose 180 of both 100 Level 1 problems (basic operators: convolutions, matrix multiplies, loss functions, etc.) and 100 Level problems (sequences of operators with fusion opportunities) training environments. Since current KernelBench does not provide a train-test split, we construct 80 additional novel tasks following same methodology (sce Appendix (A). We build the evaluation set by combining 80 newly created tasks with the 20 remaining original KernelBench tasks, for a total of 100 evaluation tasks.
Each KernelBench task consists of generating CUDA kernel given PyTorch reference implementa- tion, which is used to evaluate correctness and speedup. In our setup, we evaluate the model-generated verify the output is in the correct format (ensure resultant code is only im-
kernels as follows: plemented with inline CUDA) and check for reward hacking (Section 6.2). We then evaluate the kernel for compilation, runtime errors, and correctness. If the implementation is correct, profile the kernel for its runtime.
3.2 KERNEL SCORE DESIGN
As we are concerned both with correctness and speedup, we assign a score for each kernel evaluation result that effectively balances the correctness-performaate trade-off.
T Taseline
S = 0.3 (correct) +
1 [correat)
T Kermel
Correctness is checked against the reference program when tested with randomized inputs; speedup computed as the ratio between Py PyTocch baseline time and kernel runtime. We experimented with various weights of correctness and speedup, finding this configuration through ablations models ranging from 7B to 32B.
In addition, explored rewarding intermediate objectives (successfully compile execute), yet this caused model to over-optimize for intermediate steps (e.g. generating kernels that only compile, but not necessarily correct). We also experimented with length penalty the response, as suggested by Team (0255), but found that it degrades our model's performance during training.
3.3 SINGLE- TURN TRAINING
We apply GRPO (Saa et al. 2024) to train the model on kernel generation without iterating on external feedback ("single-tury training). In each training step, sample 16 responses per task and assign the evaluated score as the reward for each kernel. We compute the GRPO loss according to (Shao et al. 2024). which updates the policy by maximizing the following objective:
Janro EBqq P(Q), for Tow. (Olg)]
To (o 0..
To(0E,19, 01, e, +c) Ar. BDxx z (mol (Te)
min
A clip
Oi
To (0.,:99,0,.c
To (oi,lg, 0i,5t)
(I)
where
A mean(r) and P3 is the score of specific kernel. sud(r)
We choose Owen Ow0- 32B
Team 2025d) as base model. See Appendix B.6 for the rationale.
4

---

Published as a conference paper at ICLR 2026
Following Yu et al. 2025), we apply Clip-Higher. We
Single-Turn Training Reward
sample with temperature = 0.9 for both training and inference. We set the KL coefficient to 0 to allow the model to deviate freely from the base policy, following Luo al. 2025a).
Rowart
0.6
Reward
0.4
We observe that reward plateaus after 100 gradient steps, likely because single-turn training prevents the model from refining its kernels. Many generated kernels are nearly correct-often syntax or compilation fix away-but still receive 0 reward, discouraging the model from producing them. Similarly, the correct kernels do not achieve high speedup, the model optimizes for correctness rather than attempting risky approach. We address these limitations through multi-turn training.
0.3
0.1
0
20
40
60
90
Gradient Step
Figure 2: Reward plateaus during single-turn training.
4 MULTI-TURN TRAINING
In each multi-turn training step:
1. For each task, we sample m parallel trajectories of n sequential refinement turns each. Within each turn, the model generates: (1) cfain-of-thhuuht reasoning. (2) the kernel itself wrapped in code block, and (3) a summary of the chain of thought (including the main implementation techniques used). We construct the training sample of that turn by including all of (1), (2). (3) in the response, over which we compute the loss. However, in subsequent turns, only (2) and (3) are kept in the context of the model.
2. We construct the context of sample by including the history of previous responses, which include generated kernels along with their summarized CoTs, and evaluation feedback.
3. We evaluate the generated kernel and compute its score as shown in Section 3.2 The reward of each turn (CoT + response) is the discounted sum of current and subsequent scores, which we elaborate in Section 4.3
4. For each task, normalize the rewards across the mn samples for advantage calculation. Then we compute the GRPO loss over the entire batch.
4.1 MANAGING CONTEXT
Reasoning models generate long CoTss, especially for complex tasks like kernel generation. Including all CoT's causes the context to grow rapidly, reaching 50-100k tokens within a few turns, surpassing the model's context length. To prevent context explosion, discard CoT's of previous turns; yet to preserve information regarding the reasoning process, we ask the model to summarize the changes applied. This summary, along with the generated kernels and evaluation results, is passed to subsequent turns.
4.2 TRAINING ON EVERY REFINEMENT TURN
reasonable approach for multi-turn RL to treat an entire n-turn trajectory as a single training sample with the final outcome reward, as in RLEF (Gehring et al., 2025). However, this approach cannot attribute credit to individual turns within the trajectory, which we motivate could be important in the kemel setting in Figure 1. Instead, we split each n-turn trajectory into individual training examples. Each training sample corresponds to a single turn within the trajectory and includes: (1) complete CoT and kernel from that turn, and (2) the context history from previous turns (summarized COT, kernels, and evaluation results). This turn-by-tun formulation enables direct reward assignment to each turn, substantially improving sample efficiency.
4.3 REWARD AGGREGATION AND DISCOUNTING
We initially explored two naive strategies for multi-turn credit assignment. The greedy approach assigns to each turn its corresponding kemel score, while the outcome-based approach assigns to all
5

---

Published as a conference paper at ICLR 2026
turns the best score in the trajectory. The former failed to reward early suboptimal turns that lead to performant kernels later, while the latter ignores individual contributions and is sample inefficient.
Our method balances both approaches by aggregating the future kernels scores with discount factor. We conduct ablations on the reward formulation. For score aggregation, we can either take the sum R2 = Σi= "r, or maximum R3 = m { (r, over future scores. Sum favors generating multiple good kernels, while max prioritizes achieving one perfoornng kemell. We evaluate both forms with y = 0.4 and = 0.8.
Experiments show that sum with 7 0.4 scales best over 8 turns, though max performs better with 0.8 with fewer turns. We decide to use the sum reward formulation with discount factor 7y = 0.4.
Best@16 Performance Refinement
12
3.0
Correctness (9)
60-
Performance
50-
20-
10
Refinement Turns
9um
by
Refinement Turns
Figure 3: Sum with y = 0.4 exhibits the best scaling behavior. We evaluate models trained with different reward formulations under 16 parallel trajectories and 8 refinement turns.
4.4 MULTI-TURN TRAINING BEHAVIOR
Multi-Turn Training Reward
0.7
Rooorzi
For multi-tun turn ablations and training runs, we train to 80 gradient steps; within each step, for each task, sample 16 parallel trajectories and conduct 4 refinement turns.
Reward
0.4
Each batch contains tasks. (See Appendices BB55 for detailed hyperparameters C.I for training statistics)
0.2
Unlike single-turn training, reward now steadily increases. We also observe response-length behaviors similar to Luo al. (2025): the response length initially decreases, and then it starts increasing again as the model attempts more sophisticated solutions; we extend the max response length from 16K to 22K tokens at gradient step 60.
0.0
Gradient Step
Figure 4: Reward climbs steadily for multi-turn training.
5 EVALUATION
As kernel generation challenging task, models are often given extensive test-time compute to tackle it. At inference, employ multiple parallel trajectories, each made up of several serial turns.
We mark a trajectory correct if contains at least one correct kernel. Its performance the speedup of the fastest kernel (within the trajectory) over the PyTorch Eager reference (speedup of Ox if no kernel is correct). We also consider the fast, metric, introduced by Ouyang et al. (2025). which is a binary indicator for whether trajectory contains correct kernel with performance of p more. To aggregate a metric across & parallel trajectories for a given task, we compute: best@k, the maximum for that metric across all trajectories; avg@k, k, the average value across trajectories.

---

Published as a conference paper at ICLR 2026
5.1 RESULT ON KERNELBENCH EVAL SET
We compare Kevin against frontier models and the single-turn RL baseline on our aforemeenioned KemelBench eval of 100 tasks (Section B.T]. with 16 parallel trajectories, serial refinement turms. As shown in Table [ Kevin achieves a higher performance than its single-turn trained counterpart and other frontier models, demonstrating significant improvement from its base model (QWQ-32B) Qualitatively, Kevin is able to more effectively implement more aggressive optimizations across several turns (see Appendix for examples); see Appendix for additional evaluation details.
Model
Correctness
Performance
fast
fast
best@ 16 avg@
best@I 16
avg@16
best@ 16
avg@ 16
best @ 16
avg 16
Kevin (Multi-Turn)
82%
46 %
1.10x
0.40x
43%
15%
20%
6%
Single-Tum RL
82%
45%
0.85x
0.35x
43%
16 %
16%
4%
Qwen QwQ-32B
56%
11%
0.53x
0.08x
23%
3%
10%
1%
OpenAl o4-mini
38%
22%
0.78x
0.27x
21%
7%
13%
6%
OpenAl 03-mini 27%
8%
0.30x
0.08x
9%
2%
4%
2%
Table I: Kevin (multi-turn RL) outperforms other models in correctness and performance. We evaluate on 100 unseen KernelBench tasks with 16 parallel trajectories and refinement turns.
5.2 SCALING REFINEMENT TURNS
Leveraging execution feedback is crucial at test time (Ehrlich et al. 2025 Wang et al. 2025a). Thus, evaluate how Kevin scales with additional refinement turns. As shown Figure the single-turn model achieves slightly better performance with 1 turn, as its training objective optimizes for single attempt. However, when given more refinement turns, the multi-turn trained model achieves significantly higher performance, with its curve showing the highest slope. This shows that multi-turn training enhances the model's 's ability to refine and optimize kernels over turns.
5.3 SCALING PARALLEL SAMPLES
We study how be performance scales when increasing the number parallel trajectories k. while fixing the number of serial refinements turns. Prior work for RLVR on math problems (uue et al. 2025) found that RL training limits models' exploration capacity, leading to worse best@7 metrics than the base model at large As shown in Figure the performance curve of the single-turn RL model presents lower slope compared to the base model, possibly hinting at this phenomenon. In contrast, our model trained with multi-turn RL achieves a higher slope compared to both the single-turn counterpart and the base model, suggesting that multi-turn training could maintain model's exploration capacity while improving model's performance.
5.4 PARALLEL VS SEQUENTIAL SCALING
As scaling test-time compute through parallel sampling (Snell et al. 2024) and sequential iterative refinement (Ehrlich al. 2025) are both beneficial, we want systematically compare their effec- tiveness for kernel generation. To investigate, we evaluate 3 inference-tine configurations under the same total inference call budget (128 kernels): 128 trajectories with turn, 32 trajectories with turns, and 16 trajectories with 8 turns. As Table shows, allocating more refinement turns during test-time consistently better across various models, with 16 trajectories and turns being optimal.
As Section 5.1 shows, multi-turn outperforms single-turn training when evaluated in multi-turn inference setting. But since single-turn training optimizes for single-turn performance, natural question arises: does the single-turn trained model perform better by generating more single-turn responses in parallel? In Table we observe that in single-turn inference setting with 128 parallel trajectories, the single-turn model achieves slightly better performance than the multi-turn model. However, when given more refinement turns at inference, the performance and correctness improve for all models. This strengthens the case for training model that could use feedback effectively across
7

---

Published as a conference paper at ICLR 2026
Best@816 Performance vs Refinement Turns
Best@k Performance (B refinement turns)
Kavin (tul-Tum SingleTer 1..
IMuts-Tum Sengle-tum
1,0
+-
OwD 525
0.5
performance
Performance
0.4
16
Refinement Turns
k (parallel trajectories)
Figure Kevin effectively leverages multiple turns. We evaluate the above checkpoints under the same environment with parallel trajectories and 8 refinement turns.
Figure 6: Multi-turn training maintains explo- ration capacity. Refinement turns are fixed to 8, and best k performance is computed with the es- timator according to Chen et al. (202T)
multiple turns. Moreover, the multi-turn trained model achieves significantly higher performance, with faster improvement rates compared to the single-turn trained model at test-time.
Inference Config
Performance
Correctness
Model
Total
# Traj
# Turns
best @# traj
best@# # traj
Multi-Tune RL
128
16
8
1.10x
82.00%
Multi-Turn RL
128
32
4
1.02x
83.00 %
Multi-Turn RL
128
128
0.65x
76.00%
Single-Turn RL
128
16
8
0.85x
82.00 %
Single-Turn RL
128
32
4
0.81x
79.00%
Single-Turn RL
128
128
1
0.70x
73.00%
QwQ-32B
128
16
8
0.53x
57.00%
QwQ-32B
128
32
4
0,47x
52.00%
QwQ-32B
128
128
1
0.42x
54.00%
Table Multi-turn inference with 16 trajectories and 8 turns is our most optimal setup, when comparing inference configurations and their performance (x speedup) and correctness rates.
Not Okay Ratio vs Step
DISCUSSION
Clipped Not Okay Ratio
6.1 MODEL INSTABILITY
0.5
As prior RLVR work (Team et al. 2025 on QwQ 323 has shown, maintaining RL training stability is recurring challenge. In our multi- turn setting, we notice distinctive patterns of in- stability, and develop a proxy signal that guides mitigation strategies. Specifically, we observe that training for longer often causes generation of repetitive and nonsensical outputs ("junk"). In the multi-turn case, junk first appears in the final turn and gradually spreads to earlier turns, leading to model collapse.
0.4
Ratio
0.3
0.2
0.1
0.0
40
50
55
65
70
Step
Figure 7: "Not Okay Ratio" foresees model in- stability. Here the proxy signal appears roughly 15 steps earlier than junk, which is indicated by the response "Clipping Ratio" metric (Luo al. 20256).
We identified proxy signal, which we call the "Not Okay Ratio". QwQ-32B always
8

---

Published as a conference paper at ICLR 2026
begins its chain of thought with "Okay, but after 80 gradient steps, the model
begins with erratic variants like "Okay Amigos, I need to optimize this 3D tensor-matrix mult pultipl cat on"" and "Okay Holy crap, I need to get this code optimmi medzzdd ; tracking the "Not Okay Ratio" offers reliable early proxy for model instability and well precedes junk.
As detailed in Appendix [P. after attempting mitigations such KL penalty, found that using constant-length normalization in the GRPO loss (Liu et al. 2025), together with gradient-norm clipping at 0.05, successfully delayed the onset of junk responses beyond 100 gradient steps.
6.2 REWARD HACKING
We observe forms of reward hacking, as model capabilities fall short of task difficulty (Amodei al., 2016). Concretely, when a weaker model such as PepSeeee-R-Distiiee Tails to produce the correct CUDA kernels, it resorts to directly copying the reference implementation, inheriting from it, or wrapping
try-except statements. With a stronger prior like QwQ-32B, the model only fuses simple operators (ReLU, Max) and leaves key operators unmodified (in Py Torch). We address these issues by imposing stricter format checks that assign 0 reward to responses with any PyTorch functional operators. We elaborate concrete examples in Appendix G
6.3 DATA DISTRIBUTION
We found it critical to have a balanced difficulty distribution across the dataset, so that on average each batch contains both easier and harder tasks. In one experiment with Deepsek-RR-Distil---ON
(DeepSeck-AI] (025). we trained on a subset of only easy tasks. The reward quickly plateaus as the model overfits to a single difficulty level. Training with stronger base model QwQ-32B and on both level and of KernelBench resolved the issue.
7
CONCLUSION
7.1 SUMMARY
We designed a multi-turn RL training recipe that addresses challenges when applied the real-world task of kernel generation: specifically, effective context management and credit attribution across every turn enable better sample efficiency. We also added safeguards against reward hacking, and experimented with approaches to constrain and predict instability.
We present Kevin, the first model trained with RL to generate CUDA kernels. Evaluated on an unseen evaluation set, Kevin outperforms both its single-turn RL counterpart and frontier models, demonstrating that our training recipe enables the model to learn more effective refinement strategies. Multi-turn training also enables better test-time scaling, both when increasing sequential refinement and parallel sampling compute, while preserving the exploration capacity of the model.
7.2 LIMITATIONS
Our work is limited by the number of robust tasks in kernel generation (unlike math or general coding with thousands of readily available tasks). KernelBench contains only 250 tasks and requires substantial pre-processing (Appendix [A). Moreover, multi-turn RL is computationally expensive, even after extensive system optimization (Appendix C). as each rollout involves serial steps of reasoning inference, complex code generation, and careful kernel evaluation.
Nonetheless, we believe that showing significant performance gains in this domain, even under limited data and compute, highlights the effectiveness of our multi-turn training recipe. With more robust kernel environments, stronger model priors, and improved RL frameworks, we expect our method to scale accordingly.
We further note as KernelBench tasks are specified with pre-defined tensor input sizes, the speedups measure in Section 3.2 are only accurate for those dimensions and NVIDIA H200 GPUs.

---

Published as a conference paper at ICLR 2026
7.3 FUTURE WORK
We see several directions for extending our method. Incorporating a learned value network and PPO (Schulman et al. 2017) may improve baseline estimation. More sophisticated search methods (beam search, MCTS (Silver et al. (2017) may be applied at train and test time. Inspired by recent works (Sareen et al. 2025), the value network could also serve verifier for search at test-time.
Our nurl RL recipe demonstrates success in the real-world engineering task of GPU kernel generation. We hope our flexible design could be applicable to wider range of tasks with verifiable rewards and execution feedback across a trajectory. We believe explicitly training models to reason about complex tasks over multiple turns key step towards enabling autonomous Al systems.
ETHICS STATEMENT
This work introduces Kevin, multi-turn RL training method enhance LLM's ability specifically for the task of automatic kernel generation. Our research builds the publicly available model of QwQ-32B (Team, 2025db and KernelBench dataset (Ouyang et al. 2025). We document in-depth how the dataset and post-train the model.
Our work does not introduce new risks that are not already inherent in the underlying base model. We do not involve any human subject nor do we make comparison with human kernel engineers in our study, as our baseline comparisons are against the PyTorch framework (Ansel et al. 2024), following the evaluation methodology proposed in KernelBench.
REPRODUCIBILITY STATEMENT
Training Recipe:
We cover various challenges encountered during training in detail and propose effective mitigation: covering training stability F avoiding reward hacking G and careful considera- tions for RL design 4 with ablation studies. We elaborate how conduct dataset processing and provide a comprehensive set of hyper-parameters for our final model Appendix[B.5]
Computational Requirements: Each of our multi-turn training runs (for ablations and the final run) requires 650 H200 hours. As discussed in Appendix C we take steps to improve the training efficiency of this complex multi-turn RL pipeline with in-the-loop kernel profiling- We elaborate
the computation cost and step time in Appendix C.I and specifically in Table 4
Hardware Specifications: We conduct all of our RL training, evaluation, and inference on the NVIDIA H200 platform. All of our kernel runtime measurement and baseline are specific to PyTorch 2.6 and H200 hardware.
Model Weights: Model weights will be released as open source and will be accessible to ensure reproducibility.
Evaluation: For our result, we compare our methods with other models that are either released open- source (QwQ-32B) or using a fixed version of the cloud API endpoints (o4-ming -202 2022-00-1111 04-16,
10

---

Published as a conference paper at ICLR 2026
REFERENCES
Dario Amodei, Chris Olah, Jacob Steinhardt, Paul Christiano, John Schulman, and Dan Mané. 60.b/666606665
Concrete problems in ai safety, 2016. URL htt.tttts////......ooo
Jason Ansel, Edward Yang, Horace He, Natalia Gimelshein, Animesh Jain, Michael Voznesensky, Bin Bаo, Peter Bell, David Berard, Evgeni Burovski, Geeta Chauhan, Anjali Chourdia, Will Constable, Alban Desmaison, Zachary VeVito, Elias Ellison, Will Feng, Jiong Gong, Michael Gschwind, Brian Hirsh, Sherlock Huang, Kshiteej Kalambarkar, Laurent Kirsch, Michael Lazos, Mario Lezcano, Yanbo Liang. Jason Liang, Yinghai Lu, C. K. Luk, Bert Maher, Yunjic Pan, Christian Puhrsch, Matthias Reso, Mark Saroufim, Marcos Yukio Siraichi, Helen Suk, Shunting Zhang, Michael Suo, Phil Tillet. Xu Zhao, Eikan Wang, Keren Zhou, Richard Zou, Xiaodong Wang, Ajit Mathews, William Wen, Gregory Chanan, Peng Wu, and Soumith Chintala. Pytorch 2: Faster machine learning through dynamic python bytecode transformation and graph compilation. Proceedings of the 29th ACM International Conference Architectural Support for Programming Languages and Operating Systems, Volume 2, ASPLOS '24, pp. 229-947, New York, NY, USA, 2024. Association for Computing Machinery. ISBN 9798400703850. doi: 10. 10.1145/3620665.366 36446366.666666 URL httth:tts: 0 0ttp:/220..145322666555566555 31113/332666666666
Bradley Brown, Jordan Juravsky, Ryan Ehrlich, Ronald Clark, Quoe V. Le, Christopher Ré, and Azalia Mirhoseini. Large language monkeys: Scaling inference compute with repeated sampling,
2024.
Shiyi Cao, Sumanth Hegde, Dacheng Li, Tyler Griggs, Shu Liu, Eric Tang, Jiayi Pan, Xingyao Wang. Akshay Malik, Graham Neubig, Kourosh Hakhamaneshi, Richard Liaw, Philipp Moritz, Matei Zaharia, Joseph E. Gonzalez, and Ion Stoica. Skyrl-vO: Train real-world long-horizon agents via reinforcement learning, 2025.
Mark Chen, Jerry Tworek, Heewoo Jun, Qiming Yuan, Henrique Ponde de Oliveira Pinto, Jared Kaplan, Harri Edwards, Yuri Burda, Nicholas Joseph, Greg Brockman, Alex Ray, Raul Puri, Gretchen Krueger, Michael Petrov, Heidy Khlaaf, Girish Sastry, Pamela Mishkin, Brooke Chan, Scott Gray. Nick Ryder, Mikhail Pavlov, Alethea Power, Lukasz Kaiser, Mohammad Bavarian, Clemens Winter, Philippe Tillet, Felipe Petroski Such, Dave Cummings, Matthias Plappert, Fotios Chantzis, Elizabeth Barnes, Ariel Herbert-Voss, William Hebgen Guss, Alex Nichol, Alex Paino, Nikolas Tezak, Jie Tang, Igor Babuschkin, Suchir Balaji, Shantanu Jain, William Saunders, Christopher Hesse, Andrew N. Carr, Jan Leike, Josh Achiam, Vedant Misra, Evan Morikawa, Alec Radford, Matthew Knight, Miles Brundage, Mira Murati, Katie Mayer, Peter Welinder, Bob McGrew, Dario Amodei, Sam McCandlish, Ilya Sutskever, and Wojciech Zaremba. Evaluating large language models trained on code, 2021. /Rt.tttps.///aaa..oooo//////////
Terry Chen, Bing Xu, and Kirthi Devleker. Automating gpu kernel generation with deepseek-rl and inference-time httppp///bbeaeeeeeerrroacccooooe
kernel- -eezeel-eneratiorr- www.hhhdeeeeeeeeeeeeee r1 and www.re-ceeeeeeee-eeeee-------ee orgereeeeeee-eatat---eeeeeeeeeee stw.tiiiiiiiinigg/////oo
February 2025. Accessed: 2025-05-15.
Sana Damani, Siva Kumar Sastry Hari, Mark Stephenson, and Christos Kozyrakis. Warpdrive: An agentic workflow for ninja gpu transformations. In Proceedings of the Machine Learning for Systems Workshop at NeurIPS 2024, 2024. htthttts: htt////ooossttsss..... assets/ Accessed: 2025-05-15.
pepers/neurrps3222///opp
Tri Dao. Flashattention-2: Faster attention with better parallelism and work partitioning. 2023, URL https://aacxivoorr/a///33111
Tri Dao, Daniel Y. Fu, Stefano Ermon, Atri Rudra, and Christopher Ré. Flashattention: Fast and memory-efficient exact attention with io-awareness, 2022. URL httpssss///aazi..oo///bbb 2205.14135
DeepSeek-AL. Deepseek-r1: Incentivizing reasonnng capability in llms via reinforcement learning, 2025. Utt/httpss://rrx...or//////////
Tulsee Doshi.
Gemini 2.5: Our most intelligent models are getting even bet-
ter. https://blog.googg http://loo..gggglee/ececcooooooo wwt.e//eeeloggy///eeeeooooo doooog/ledddeiinngggggeeeeeeee
hpdatess-o-20255, May 2025. Accessed: 2025-05-21.
11

---

Published as a conference paper at ICLR 2026
Yan Duan, John Schulman, Xi Chen, Peter L. Bartlett, llya Sutskever, and Pieter Abbeel. RI:: Fast reinforcement learning via slow reinforcement learning, 2016. URL https://arxiv.org//
abs/1611.02779
Ryan Ehrlich, Bradley Brown, Jordan Juravsky, Ronald Clark, Christopher Ré, and Azalia Mirhoseini. Codemonkeys: Scaling test-time compute for software engineering, 2025. URL https://
mox.v.org/abs//25000
Paul Gauthier. ol tops aider's new polyglot leaderboard, httDs://a.der..htt/2222. 21/polyglot.html December 2024. Accessed: 2025-04-16.
Jonas Gehring, Kunhao Zheng, Jade Copet, Vegard Mella, Quentin Carbonneaux, Taco Cohen, and Gabriel Synnaeve. Rlef: Grounding code Ilms in execution feedback with reinforcement learning. 2025. URL
abs////b///aa 0atpb///210000888
Anna Goldie, Azalia Mirhoseini, Hao Zhou, Irene Cai, and Christopher D. Manning. Synthetic data
generation & multi-step for reasoning & tool use, 2025. URL httppss//ars/vvoo////bbbb
2504.04736
Google Deep.indd. Alphaevolve: gemini-powered coding agent for designigg advanced algoriin May 2025. URL https: Dttp:://aeenniiooo/////oo///bbba
Accessed: 2025-05-15.
Jian Hu, Xibin Wu, Zilin Zhu, Xianyu, Weixun Wang, Dehao Zhang, and Yu Cao. Openrlhf: An casy-to-use, scalable and high-performance rlhf framework, 2024. URRppst:s//axxx...o
abs/2405.1114
Naman Jain, King Han, Alex Gu, Wen-Ding Li, Fanjia Yan, Tianjun Zhang, Sida Wang, Ar- mando Solar-Lezamm Koushik Sen, and Ion Stoica. Livecodebench: Holistic and contamination free evaluation of large language models for code, 2024. URI. httph:tt//saat http:::://aaii..rr/abbbb.
24C3.07974
Carlos E. Jimenez, John Yang, Alexander Wettig, Shunyu Yao, Kexin Pei, Ofir Press, and Karthik Narasimhan. Swe-bench: Can language models resolve real-world github issues?, 2024. URL
Denis Kocetkov, Raymond Li, Loubna Ben Allal, Jia Li, Chenghao Mou, Carlos Muñoz Ferrandis, Yacine Jernite, Margaret Mitchell, Sean Hughes. Thomas Wolf, Dzmitry Bahdanau, Leandro von Werra, and Harm de Vries. The stack: 3 tb of permissively licensed source code, 2022. URL
httpss:/aarxv..rrg/ 2221511.55/
Woosuk Kwon, Zhuohan Li, Siyuan Zhuang, Ying Sheng, Lianmin Zheng, Cody Hao Yu, Joseph E. Gonzalez, Hao Zhang, and lon Stoica. Efficient memory management for large language model serving with pagedattention. In Proceedings of the ACM SIGOPS 29th Symposium on Operating Systems Principles, 2023.
Nathan Lambert, Jacob Morrison, Valentina Pyatkin, Shengyi Huang. Hamish Ivison, Faeze Brahman, Lester James V. Miranda, Alisa Liu, Nouha Dziri, Shane Lyu, Yuling Gu, Saumya Malik, Victoria Graf, Jena D. Hwang, Jiangjiang Yang, Ronan Le Bras, Oyvind Tafjord, Chris Wilhelm, Luca Soldaini, Noah A. Smith, Yizhong Wang, Pradeep Dasigi, and Hannaneh Hajishirzi. Tulu 3: Pushing frontiers in open language model post-training, 2025. URL http:/s:///axiv.orrrrr
abs/ 0ess/211111124
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, 2025. URL
htttpss:/ppu..ssaaaaaaa//// pttp:/ttta..ppppp...pp..p
Accessed: 2025-05-15.
Jianling Li, Shangzhan Li, Zhenye Gao, Qi Shi, Yuxuan Li, Zefan Wang. Jiacheng Huang, Haojie Wang, Jianrong Wang, Xu Han, Zhiyuan Liu, and Maosong Sun. Tritonbench: Benchmarking large language model capabilities for generating triton operators, 2025. URL https:tpss./. htttps://arxiv.ors
abs/ abs/2502. 255./2550.145
12

---

Published as a conference paper at ICLR 2026
Raymond Li, Loubna Ben Allal, Yangtian Zi, Niklas Muennighoff, Denis Kocetkov, Chenghao Mou, Marc Marone, Christopher Akiki, Jia Li, Jenny Chim, Qian Liu, Evgenii Zheltomozhskii, Terry Yue Zhuo, Thomas Wang, Olivier Dehaene, Mishig Davaadorj, Joel Lamy-Poirier, João Monteiro, Oleh Shliazhko, Nicolas Gontier, Nicholas Meade, Armel Zehaze, Ming-Ho Yee, Logesh Kumar Umapathi, Jian Zhu, Benjamin Lipkin, Muhtasham Oblokulov, Zhiruo Wang. Rudra Murthy, Jason Stillerman, Siva Sankalp Patel, Dmitry Abulkhanov, Marco Zocca, Manan Dey, Zhihan Zhang, Nour Fahmy, Urvashi Bhattacharyya, Wenhao Yu, Swayam Singh, Sasha Luccioni, Paulo Villegas, Maxim Kunakov, Fedor Zhdanov, Manuel Romero, Tony Lee, Nadav Timor, Jennifer Ding, Claire Schlesinger, Hailey Schoelkopf. Jan Ebert, Tri Dao, Mayank Mishra, Alex Gu, Jennifer Robinson, Carolyn Jane Anderson, Brendan Dolan-Gavitt, Danish Contractor, Siva Reddy, Daniel Fried, Dzmitry Bahdanau, Yacine Jernite, Carlos Muñoz Ferrandis, Sean Hughes, Thomas Wolf, Arjun Guha, Leandro von Werra, and Harm de Vries. Starcoder: may the source be with you!, 2023. URL
Jiate Liu, Yiqin Zhu, Kaiwen Xiao, Qiang Fu, Xiao Han, Wei Yang, and Deheng Ye. Ritf Reinforce- 2ttpr////aab/330.0343333
ment learning from unit test feedback, 2023. URL
Zichen Liu, Changyu Chen, Wenjun Li, Penghui Qi, Tianyu Pang, Chao Du Wee Sun Lee, and Min Lin. Understanding rl-zero-like training: A critical perspective, 2025. URL https
//arxiv.org/ abs/ /
Michael Luo, Sijun Tan, Roy Huang, Ameen Patel, Alpay Ariyak, Qingyang Wu, Xiaoxiang Shi, Rachel Xin, Colin Cai, Maurice Weber, Ce Zhang, Li Erran Li, Raluea Ada Popa. and Ion Stoica. Deepcoder: A fully open-source 14b coder at 03-mini level. https Source-
2025a.
Notion Blog.
Michael Luo, Sijun Tan, Justin Wong, Xiaoxiang Shi, William Y. Tang, Manan Roongta, Colin Cai, Jeffrey Luo, Li Erran Li, Raluca Ada Popa, and Ion Stoica. Deepscaler: Surpassing ol-preview with a 1.5h model by scaling rl. /ttps: / boo.
RL
00L-1111002211600064443333333333 2025b. Notion Blog.
Mathematical
Association of America. American invitational mathematics examination
AIME. httpa::/maa.org//me http:////wmmmccccciiiiiceceeeeee competitions
mathematics sew.tttiiiieeemmmatiiiii.n.. aime February 2024. Accessed: 2025-09-24.
METR. Measuring automated kernel engineering, February 2025. URL http::/tt////mtt /ttt://mmmetrrrr//
blog/2025- 02 htt.::/www.....aaaao..ooooo.oan. 2025-05-15.
Accessed:
Daniel Nichols, Aniruddha Marathe, Harshitha Menon, Todd Gamblin, and Abhinav Bhatele. Hpe-coder: Modeling parallel programs using large language models. In ISC High Perfor- mance 2024 Research Paper Proceedings (39th International Conference). pp. 1-12. IEEE, May 2024. doi: 10.239 2239199is..222..100 10528999. URL http: http:://ox.do..orr//0233311 ESC.2024.10528999
John Nickolls, Ian Buck, Michael Garland, and Kevin Skadron. Scalable parallel programming with cuda. In ACM SIGGRAPH 2008 Classes, SIGGRAPH '08, New York, NY, USA, 2008. Association for Computing Machinery. ISBN 9781450378451. doi: 10. 10.1145//401 1401
Allen Nie, Yi Su, Bo Chang, Jonathan N. Lee, Ed H. Chi, Quoe V. Le, and Minmin Chen. Evolve: Evaluating and optimizing Ilms for exploration, 2024. URL https://arxiv.org/aaaa 2410.06238
NVIDIA. Gpu mode at nvidia gtc 2025, 2025. wtt.itttp.////wwyootbbbb.o/////t mdDVkBeFy9AJ Accessed: 2025-05-15.
NVIDIA Corporation. Cutlass: Cuda templates for linear algebra subroutines, May 2025. URL Accessed: 2025-05-15.
13

---

Published as a conference paper at ICLR 2026
Anne Ouyang, Simon Guo, Simran Arora, Alex L. Zhang, William Hu, Christopher Ré, and Azalia Mirhoseini. Kemellbench: Can llms write efficient gpu kernels?, 2025. URL https
//arxiv.org/ awtgg//aa///e http://wrrg/aaa//222.55555 220...001.
Yuxiao Qu, Matthew Y. R. Yang, Amrith Setlur, Lewis Tunstall, Edward Emanuel Beeching, Ruslan Salakhutdinov, and Aviral Kumar. Optimizing test-time compute via meta reinforcement fine- tuning, 2025. URL htt.:/ttpss./aaatt.oot///a/a//./
Kusha Sarcen, Morgane M Moss, Alessandro Sordoni, Rishabh Agarwal, and Arian Hosseini. Putting the value back in rl: Better test-time scaling by unifying Ilm reasoners with verifiers, 2025. URL
Thomas Schmied, Jörg Borncheei, Jordi Grau-Moya, Markus Wulfmeeie, and Razvan Pascanu. Lims are greedy agents: Effects of rl fine-tuning decision-making abilities, 2025. URL
John Schulman, Filip Wolski, Prafulla Dhariwal, Alec Radford, and Oleg Klimov. Proximal policy optimization algorithms, 2017. /tL/
Zhihong Shao, Peiyi Wang, Qihao Zhu, Runxin Xu, Junxiao Song, Xiao Bi. Haowei Zhang, Mingchuan Zhang. Y. K. Li, Y. Wu, and Daya Guo. Deepseckmath: Pushing the limits of mathematical reasoning in open language models, 2024. URL https:://arxiv.org///
2402.03300
Alexander Shypula, Aman Madaan, Yimeng Zeng, Uri Alon, Jacob Gardner, Milad Hashemi, Graham Neubig, Parthasarathy Ranganathan, Osbert Bastani, and Amir Yazdanbakhsh, Learning performance-imrrovie code edits, 2024. JBL/:ttp..//arxxxooo///////2/2/2
David Silver, Thomas Hubert, Julian Schrittwieser, Ioannis Antonoglou, Matthew Lai, Arthur Guez, Mare Lanctot, Laurent Sifre, Dharshan Kumaran, Thore Graepel, Timothy Lillicrap, Karen Si- monyan, and Demis Hassabis. Mastering chess and shogi by self-play with a general reinforcement learning algorithm, 2017. /ttps:///x//t
Charlie Snell, Jaehoon Lee, Kelvin Xu, and Aviral Kumar. Scaling Ilm test-time compute optimally be more effective than scaling model parameters, 2024. URL /ttp:://///////o.ogg/////// /abs/
22408.03314
Benjamin F. Spector, Simran Arora, Aaryan Singhal, Daniel Y. Fu, and Christopher Ré. Thun- derkittens: Simple, fast, and adorable ai kemels, 2024. URL
(.200.0399
Fahim Tajwar, Yiding Jiang, Abitha Thankaraj. Sumaita Sadia Rahman, Zico Kolter, Jeff Schnei- der, and Ruslan Salakhutdinov. Training generally curious agent, 2025. URL https
7/axxv.org/abbl hatrg//bb//// ht////b//2222222 2502.17533
Gemma Team.
Gemma 3 technical report, 2025a.
URL Attps://arxiv.org. /abs/
2503.1 1986
Kimi Team.
Kimi k1.5: Scaling reinforcement learning with llms, 2025b.
URL https://
arxiv.org
NovaSky Team. Sky-tl: Train your ol preview model within $450. https://novasky.
aiggihubbio/postt/kky... 2025c. Accessed: 2025-01-09.
Prime Intellect Team, Sami Jaghouar, Justus Mattern, Jack Min Ong, Jannik Straube, Manveer Basra, Aaron Pazdera, Kushal Thaman, Matthew Di Ferrante, Felix Gabriel, Fares Obeid, Kemal Erdem, Michael Keiblinger, and Johannes Hagemann. Intellect-2: A reasoning model trained through globally decentralized reinforcement learning, 2025. URI.
abs/ 2255/25097799
Qwen Team. Qwg-32b: Embracing the power of reinforcement learning, March 2025d. URL htttss:/qweet m.gi thub. wwwb..oobll.. h//l htt.i///olooo/////ooo -2bp//3////
14

---

Published as a conference paper at ICLR 2026
Philippe Tillet, H. T. Kung, and David Cox. Triton: intermediate language and compiler for tiled neural network computations. In Proceedings of the 3rd ACM SIGPLAN International Workshop on Machine Learning and Programming Languages, MAPL 2019, pp. 10-19, New York, NY, USA, 2019. Association for Computing Machinery. ISBN 9781450367196 doi: 10.1145/3315508.333 33329973.
Siddhant Waghjale, Vishruth Veerendranath, Zora Zhiruo Wang, and Daniel Fried. Ecco: Can code improve modelmeeeddl--geeettededdddo code efficiency without sacrificing functional correctness?, 2024. URL
https://arxiv.orgg /
Guanhua Wang, Heyang Qin, Sam Ade Jacobs, Connor Holmes, Samyam Rajbhandari, Olatunji Ruwase, Feng Yan, Lei Yang, and Yuxiong He. Zero++: Extremely efficient collective communi- cation for giant model training, 2023. URL
abs http:/orgg/abs/33..1.00..0 2306.10209
Xingyao Wang, Boxuan Li, Yufan Song, Frank F. Xu, Xiangru Tang, Mingchen Zhuge, Jiayi Pan, Yueqi Song, Bowen Li, Jaskirat Singh, Hoang H. Tran, Fuqiang Li, Ren Ma, Mingzhang Zheng, Bill Qian, Yanjun Shao, Niklas Muennighoff, Yizhe Zhang, Binyuan Hui, Junyang Lin, Robert Brennan, Hao Peng, Heng Ji, and Graham Neubig. Openhands: An open platform for ai software developers as generalist agents, 2025a. URL https: http//aaxvvvooo/o /
Yiping Wang. Qing Yang, Zhiyuan Zeng, Liliang Ren, Lucas Liu, Baolin Peng. Hao Cheng, Xuehai He, Kuan Wang, Jianfeng Gao, Weizhu Chen, Shuohang Wang, Simon Shaolei Du, and Yelong Shen. Reinforcement learning for reasoning large language models with one training example, 2025b. URL http://s///xxxv./g//////// 355///0../2222220000
Zihan Wang, Kangrui Wang, Qineng Wang, Pingyue Zhang, Linjie Li. Zhengyuan Yang. Kefan Yu, Minh Nhat Nguyen, Licheng Liu, Eli Gottlieb, Monica Lam, Yiping Lu, Kyunghyun Cho, Jiajun Wu, Li Fei- Lijuan Wang, Yejin Choi, and Manling Li. Ragen: Understanding self-evolution in Ilm agents via multi-turn reinforcement learning, 2025c. TRL/attpss:/aaxivaatt 2007
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 reinforcement learning on open software evolution, 2025. URL 2502.18441
Violet Xiang, Charlie Snell, Kanishk Gandhi, Alon Albalak, Anikait Singh, Chase Blagden, Duy Phung, Rafael Rafailov, Nathan Lile, Dakota Mahan, Louis Castricato, Jan-Philipp Franken, Nick Haber, and Chelsea Finn. Towards system 2 reasoning in Ilms: Learning how to think with meta ofain-of-thoughe 2025. URL http:: http:://///a.v.v./r//// /
Zihao Ye, Lequn Chen, Ruihang Lai, Wuwei Lin, Yineng Zhang. Stephanie Wang, Tianqi Chen, Baris Kasikci, Vinod Grover, Arvind Krishnamurthy, and Luis Ceze. Flashinfer: Efficient and customizable attention engine for IIm inference serving. arXiv preprint arXiv:2501.01005 2025.
URL httpts::://xxi..ogg/////o// http:///gggabb//1...1100. 22t.//100111011
Qiying Yu, Zheng Zhang, Ruofei Zhu, Yufeng Yuan, Xiaochen Zuo, Yu Yue, Tiantian Fan, Gaohong Liu, Lingjun Liu, Xin Liu, Haibin Lin, Zhiqi Lin, Bole Ma, Guangming Sheng, Yuxuan Tong, Chi Zhang, Mofan Zhang, Wang Zhang, Hang Zhu, Jinhua Zhu, Jiaze Chen, Jiangjie Chen, Chengyi Wang, Hongli Yu, Weinan Dai, Yuxuan Song. Xiangpeng Wei, Hao Zhou, Jingjing Liu, Wei-Ying Ma, Ya-Qin Zhang, Lin Yan, Mu Qiao, Yonghui Wu, and Mingxuan Wang. Dapo: An open- source 1lm reinforcement learning system at scale, 2025. URL 2503.14476
Yang Yue, Zhiqi Chen, Rui Lu, Andrew Zhao, Zhaokai Wang, Yang Yue, Shiji Song, and Gao Huang. Does reinforcement learning really incentivize reasoning capacity in Ilms beyond the base model?, 2025.
Chenggang Zhao, Liang Zhao, Jiashi Li, and Zhean Xu. Deepgemm: clean and efficient fp8 gemm ker- nels with ine-grained scaling. http::/s//.ttttacccoccoeeeeeeeee
2025.
Kunhao Zheng. Jesse Michael Han, and Stanislas Polu. Minif2f: a cross-system bencharkk for formal olympiad-level mathematics, 2022. URL httppsss/aarrttoorr//////1111111
15

---

Published as a conference paper at ICLR 2026
Yifei Zhou, Andrea Zanette, Jiayi Pan, Sergey Levine, and Aviral Kumar. Archer: Training lan- guage model agents via hierarchical multi-turn rl, 2024. URL attps://acx....or///aaa 24C2.19446
Richard Zhuang*, Trung Vu*, Alex Dimakis, and Maheswaran Sathiamoorthy. Improving multi-turn tool use with reinforcement learning. https://www. wwtp:/esppiellaaaiiog/mmopigiggo huiblog/mprovinggmtit-u----s learning 2025. Accessed: 2025-04-17.
16

---

Published as a conference paper at ICLR 2026
A KERNELBENCH MODIFICATIONS
We use KernelBench Ouyang et al. (2025) as our training environments. KernelBench is a popular benchmark for evaluating LLMs' ability to generate performant CUDA kernels for deep learning workloads in Py Torch. Each KernelBench task consists in generating CUDA kernel given Py PyTorch reference implementation, which is used evaluate correctness and speedup.
A.1 TASK IMPROVEMENTS
We identify several limitations in the original KernelBench and introduce targeted modifications to address them. These changes are crucial to mitigate reward hacking, shown in Section 6.2
We sand-boxed the kernel evaluation process that fatal emror, such as CUDA i illegal memory accesses, do not crash the RL training process.
A significant issue we noted in KernelBench was that for many tasks, the input tensors used measure performance are quite small. This causes kernel launch overhead to take up significant portion of the runtime. To address this, we enlarged the tensor dimensions of the affected tasks.
A sneakier bug in the KernelBench's evaluation harness caused the tested kernel to recycle the output tensor from the reference implementation (which was run immediately before) as its own tensor output. As result of this, kernel that only computes (correctly) portion of the output tensor would still pass the correctness check. We address this by running the tested kernel first and only after the reference implementation, thus avoiding this hack.
In the end, we chose a total 180 tasks as training environments, with 90 of the 100 Level problems and 90 Level 2 problems (sequences of operators with fusion opportunities).
A.2 CONSTRUCTION ADDITIONAL EVALUATION SET
Since current KernelBench does not provide train-test split, we construct 80 additional tasks following the same methodology that KernelBench was constructed.
KernelBench Level 2 is constructed by composing subset of PyTorch operators sequences of operators. Specifically, the PyTorch operators are categorized as:
Main operators:
Conv2d, Matmul, Gemm, BMM, Conv3d, ConvTranspose2d,
ConvTranspose3d.
Activations:
ReLU, gigmiid, Tanh, LeakyRelU, GELU, Swish Softmax, Mish, Hardtanh, HardSwish.
Element-wise operators: Add, Multiply, Subtract, Divide, Clamp, Scale, ResidualAdd.
Normalizations: BatchNorm, LayerNorm, InstanceNorm, GroupNorm.
Pooling: MaxPool, AvgPool, Globalagpoll..
Bias: BiasAdd.
Reductions: Sum, Mean, Max, Min, LogSumExp.
Others: ResidualAdd, Scaling.
To construct the additional eval set (unseen from train set), following the methodology from original KernelBench task construction:
1. We sample from the available operators listed above: main operator (computationally expensive), and 2-5 other operators.
2. We ask language model, namely Gemini 2.5-Flash (Doshil 2025), to generate Py Torch program that creates a kernel by combining these operators. We also ask it to generate sample tensor sizes for the task.
17

---

Published as a conference paper at ICLR 2026
3. We ensure this PyTorch program can be executed and has runtime NVIDIA H200 .ms, to avoid the runtime being dominated by kernel launch (CPU) overhead.
4. We make sure this PyTorch program (with the same sequence operators) is not present in existing KernelBench Level and programs.
We manually inspected all new task programs to ensure their validity. We build the evaluation set by combining our 80 newly created tasks with the 20 remaining original KernelBench tasks, for a total of 100 unseen evaluation tasks.
B ADDITIONAL DETAILS ON MULTI-TURN RL
Here elaborate on design choices for our RL Training as described in Section 3.3 and Section along with some ablation results.
B. MOTIVATION FOR TURN-W REWARD
In our multi-turn RL training setup, within each training step we have a trajectory with refinement turns. possible approach would be to compute the reward based on the kernel at the last turn, similar to what is used in RLEF (Gchring et al., 2025). However, for the GPU kernel optimization setting, using just the last kernel might not be optimal at times: for example, as shown earlier in Figure kernel 3 is correct but kernel 4 is incorrect as the model attempts more aggressive optimizations.
In this setting, computing reward based on the best kernel among the trajectory instead (max speedup) more natural choice. However, using only the max kernel score forces discard all turns in a trajectory after the max turn, possibly wasting a significant amount inference rollouts: In the previous example,
would have completely discard the reasoning trace, code, and evaluation for kernel 4. Thus, we arrived at our approach in Section 4.3 which uses discounted look-ahead max or sum, enabling more sample-cfficient training.
1.2
0.6
Reward
Roward
EMA
DWA
1.0
0.5
0.8-
Roward
Reward
0.6
0.3
0.4
0.2
0.2
0.1
0.0
0.0
70
Gradient Step
Gradient Step
Figure 8: Training reward with correctness weighting of 1, performance speedup weight- ing of I. Concretely, S 1 [conre) + Theainx 1 [correct)
Figure 9: Training reward with no correctness weighting, performance / speedup weighting of 1. (speedup is O if kernel is incorrect). Concretely,
Tiaue
S = (correct) Thoas
Tandiu:
B.2 WEIGHTING FOR SCORE
In Section 3.2
we explain our score design, which assigns a scalar value (score S) based on a kernel's correctness and speedup. We explore score design and how to balance the correctness-eeffommacce trade-off, after series of small-scale ablations on QwQ 32B (Team 2025d).
We decided weighting of 0. correctness and using speedup for performance (raw speedup itself, weighting), which is S 0.3 (conreat)
Trwetns
(comret)
Thane
18

---

Published as a conference paper at ICLR 2026
Here present some ablation studies with different weighting configurations for score design, particularly focusing on adjusting the weighing for correctness, in the context of single-turn RL (GRPO) training (as shown in Section 3.3). As show example in Figure 8 where set the weighting to 1.0 for correctness, the reward plateaus and eventually decreased; concretely, we observed that the model over-optimizes for generating correct kernels and does not explore speedup as much, causing the reward to plateau during training. In another experiment in Figure | we set the weighting to 0 for correctness, only rewarding the model for generating performant (and correct) kernels. We again observed the reward plateau. Thus, we hypothesize that it is still important to reward the model for correct kernels, as long as the correctness reward is not too significant, balancing the correctness- perfrremnneeeeeooeeeceeeee tradeoff.
B.3 NUMBER OF TRAJECTORIES DURING TRAINING
We vary the number of parallel trajectories during Multi Turn RL training (Section 4 using 64 parallel trajectories instead of 16 for each task during each training step. We note that best @ 16 correctness slightly increases, but overall performance does not show significant improvements. Due to the high-compute requirements of doing more generations during training, we chose to train with 16 parallel trajectories.
B.4 LENGTH PENALTY
We explore incorporating response length part of the reward design to incentivize the model to use its reasoning tokens more efficiently. We attempted a run using the length penalty from Kimi Team (202bb DeepSeek-R1 -eiee--R--Dsttt---errrrr 1BB. As shown in Figures I0 and [T found that the response length of the responses collapses, with the model no longer outputting CoT after 10 training steps, suggesting that the addition of length penalty is counterproductive for setting.
trainfrewwrd
train/campletion lenghh
Nom y
0.5
1000
Figure 10: Training Reward collapses when in- cluding length penalty as part of reward
Figure 11: Response length of generations col- lapses when including length penalty as part of reward.
B.5 DETAILED TRAINING HYPERPARAMETERS
Here the set of hyperparameter for our final Kevin training run.
Constant Tearning rate of 2e-6 with warmup ratio of 0.03
Max grad norm 0.05
KL coeff =
Temperature 0.9
Top p 0.95
Eps clip = 0.2 with clip high 0.28
Max prompt len 8192
Max generate len 22432
19

---

Published as a conference paper at ICLR 2026
B.6 BASE MODEL CHOICE RATIONALE
We experimented with several different base models, such as DeepSeek-Rl-Di.se -t--iiitti--eenni 7B DeepSeek-Al (2025). DeepSeek-R1-D35 stll-iistl---laammi
and Gemma27B-Instrin [Team (2025a). These models, however exhibit weak kernel writing priors, which causes the initial reward to be overly sparse for effective learning, making the model resort to reward hacking (see Section 6.2)
Specifically, QwQ-32B exhibits stronger priors in math and coding than other base models, as shown in the the table below. Given that gradient updates only occur if the model receives non-zero reward in at least one rollout for given task, and that for many tasks the model is unable to generate any correct solutions, the model only receives a very sparse signal for the few kernels that it's able to implement correctly. This sparse feedback is not sufficient for the model to learn. As a result, even after 15 steps, the reward shows no improvements. This behavior is reflected in the attempted training runs with these models: with the same recipe, there is improvement in performance even after dozens steps, with the reward staying flat.
Model
AIME 24
LiveCodeBench
SWE-Bench Verified
Aider Polyglot
QwQ-32B
79.5%
63.4%
41.3%
20.9% %
Gemma27BBInstten
25.3%
29.7%
N/A
4.9%
50.4%
39.6%
N/A
N/A
DeepSeckkRI-Disstil--eeeZ
55.5%
37.6%
N/A
N/A
Table 3: Base Model Comparison on Math and Coding Benchmarks
Here we also present the training reward curve for Gemma-27B. The model is unable to improve on the task since receives little reward signal throughout the training run.
train/reward
0.05
0.04
0.03
train/globall_ttg
12
Figure 12:
Training reward curve on Gemma-27B. With sparse reward signal which often leads to zero gradients, preventing model from improving on this task.
We thus choose Qwen QwQ-32B TTam 2025db our base model, which exhibits, among all the models have evaluated of comparable size, the strongest priors.
C RL INFRASTRUCTURE
Conducting RL training on a highly challenging task like GPU kernel generation is computationally expensive process, requiring full-policy updates sufficiently capable base model, as discussed in Section
Although few open-source RL frameworks existed when began this study, it is still difficult to support training in a kernel evaluation environment and including multiple turns within one training step. We built our training framework top of the OpenRLHF (Hu et al.) 2024) framework.
We use vLLM wwon et al. 2023) for inference and DeepSpeed Zero-3 (Wang et al., 2023 for offloading optimizer states.
20

---

Published as a conference paper at ICLR 2026
FORWARD
Eval
VLLM #1
Parsed?
VILNN #2
VELNH #3
Promtts
VLLM
Speedap
Reward
Figure 13: Overview of our RL Training infrastructure.
Each of the 8 GPUs handles the kernel generation and evaluation for one task. After the response generation finishes, each GPU offloads its vLLM engine to CPU memory and evaluates the kernels it generated. We run the evaluation and calculate reward and evaluation info. Each GPU then wakes up its corresponding vLLM engine and regenerates kernels.
We optimized our training infrastructure to co-locate vLLMM rollout engines, kernel execution envi- ronments, and DeepSpeed trainers on the same GPU device, so other small research teams with cluster could also experiment with our proposed method.
C.I TRAIN TIME STATISTICS
Here elaborate more the cost of multi-turn training. The nature of multi-turn RL requires multiple serial turns of parallel rollouts and kernel compilation/execction after each step, making the overall training process compute-intensive. intensive. To accurately measure kernel runtime, must clear the GPUs of any running processes and perform additional operations, such as warmup steps before profiling, which further limits the training speed. Here we show key training time statistics:
Configuration
Value
Gradient steps
80
Parallel trajectory rollouts
16
Refinement turns (serial env, interactions)
Gradient updates per batch
2 (1 on-policy, 1 off-policy)
Time for rollout + kernel execution (per step)
1.5 hours
Time for gradient update (2 steps)
0.5 hours
Base model
QwQ-32B
Table 4: Setup and Cost of multi-turn Training for Kevin 8xH200s.
Overall, one training run (Section takes 650 H200 hours, equivalent to around 3 days and 9 hours on single node of 8xH200s. However, believe concurrent and future systems projects (such as SkyRL by Cao et al. (2025)) will improve training efficiency, especially for roll-outs with complex interactions with environments. The demanding computational requirement of multi-turn RL is what leads us to focus on improving the sample efficiency of our method; specifically, we choose to train on every sample regardless of their performance and attribute credit effectively with our reward design.
D INFERENCE SETUP
Our prompt similar the prompt used in KernelBench (Ouyang al. (025). We this during training and test-time inference. In the first refinement turn, we add an example the inline CUDA format to the prompt but remove it afterwards.
21

---

Published as a conference paper at ICLR 2026
Below show how construct the context in the simplest case (of one turn, the base prompt). In the context, we present model the KernelBench task, instructions, and simple 1-shot example of CUDA add kernel (to inform model the desired format for response):
You are given the following architecture:
import torch
import torch.nn an as mn
class Model Module)
=**
Simple model that performs Layer Normalization.
def ininti(sel normalized shaee tuple)
**. Initializes the LayerNorm layer.
12
Args:
normalized_shar (tuple): Shape of the input tensor to be
normalized.
15
**"
16
super (Model, self) init)_ ()
self.in = La.eeeererrmmo htw.re/eeeeocoeeeeeeeeeeeeeeooe
def (seward(eelf x: torch.eenser
Nensor:
20
**"
21
Applies Layer Normal ization to the input tensor.
Args:
24
x (torch.Tensor): Input tensor of shape (*,
normalized_shape).
25
Returns:
corch,Tensors Output tensor with Layer Normalization applied, same shape as input.
28
**"
return Self.inis
31 Replace pyterch operators the given architecture with raw CUDA kernels, optimizing for performance on NVIDIA H200 (e.g. shared memory, kernel fusion, warp primitives, vectorization,.... and name your optimized output torch.utils. wtt.eextensioo.ooo__iiii allored architecture ModelNew. You for Parameter, containers, and init). The input and output have to be on CUDA device. Your answer must be the complete new architecture (no testing code, no other code): it will be evaluated and you will be given feedback iterating, trying to naximize the speodup. After your answer, summarize your changes a few sentences, Here is an example:
Use
torch.ns (except
its correctness and speedup you
keep
32
import torch. as from forch.utils epp_ extension import Load inline
35
36
Define the custom CUDA kernel for element-wise addition dementwiss_adddssrre
37
Vinclude <torch/extension.it Vinclude <cuda_runtiime.t
40
41
global__ void elementwise_ade ker_eeeeeddddeeeeeeeeeeeeeeell kernel (const float const float b,
floatx out, int size) {
int idx blockIdx.x
blockDim.x threadIdx.x;
43
if (idx size)
44
out out [idx] = a [idx] b[idx];
45
22

---

Published as a conference paper at ICLR 2026
torch: Tensor elementwise add_cuda seorrch:Teenso a, torch: : auto size a. a.numel ()
30
auto out = :2eros_ (a);
51
const int block_size 256;
const int num_blocks (size block_size 1) / block_sizey
55
block (a (tt.atttattaattt p (),
out. ptttcttaattrrt size);
57
return out;
58
61
elementwise add___dd__ppp______ http__pppp________eeeeeeeeee source
62
"torch: Tensor clement wespnttwssc_ddde sc_add__uud_ (torch: :Tnnorr a, torch: :Tnsorr b);"
65
Compile the inline CDDA code for element-wise addition el -ementwise_add = load_inline
66
name="e "ddameettttssead"",,,"__
cpp app_scurcesreeeeentttisses add cpp_ source,
69
cuca seurcesreeeeeeettiieeeeeeeeeeeee add scurce, scurce, functicns-("elementi stt.lemmenntis___d__ddddddl"
70
71
verbose-True,
extra_cfllas=="") extra [atra_ldfllas=(""]
74
75
76
class ModelNew (nn. Modull))
def Lniit_(seff -> Nene:
79
super self.elementwise_.ee - eleeetwise_ead
80
def forward(se1f, a, b):
return elemetteeet add. wwweeewwweeeeeeeeee.eee.e..eeeee add cuda (a, b)
For our multi-turn RL training (Section 4 and inference (Section 5), we provide model with the kernels, CoTs (summarized), and evaluation results of all previous turns in chronological order. We truncate the turns that do not inside the context window, starting from the earliest ones.
<Base prompt containing pytorch architecture and instruction>
Here are your previous attempts:
for each (i) previously generated kernel > <Frevicusly generated kernel Gli]>
<Summary of CoTli|>
<if parsing error>
12
Your previous answer failed to be parsed due to not adhering to the desired formatting. Here is the error message: <error, perrormesssage
<elif compilation
16
Your previous answer failed to compile. Here is the error message: <error_message
<elif error>
23

---

Published as a conference paper at ICLR 2026
20
Your previous answer compiled successfully had runtime errors. Here is the error message: <error_meseage>
21
<elif correctness error>
24
Your previous answer was incorrect. Here the error message: <error_message
25
<elif correct>
28
Your previous answer was correct but can be made faster. Here is the speedup you achieved relative to the baseline: <speedup>
Restart your reascning process and generate new, complete code.
ADDITIONAL. EVALUATIONS
Here we present some additional evaluation results for Section
E CONFIDENCE INTERVALS
We compute the confidence intervals of best 16 and avg@@ performance for the multi-turn and single-turn RL across runs, shown in Table These results show multi-turn RL has statistically significant improvement on both metrics and hence its effectiveness.
Model
Performance
best@ 16
avg@ 16
0.40 + 0.011 Single-turn RL 0.85 0.048 0.35 0.013
Multi-turn RL 1.10 + 0.099
Table Evaluation on our evaluation set across 5 runs with confidence interval. Multi-turn RL outperforms Single-turn RL on both @e 16 and @ 16 performance.
E.2 CHOICE OF BASELINE MODEL COMPARISON
Here elaborate the choice of model comparisons used for 5.1 notably against both Kevin's 's base model (QwQ-32B) and frontier reasoning models (o4-mini, 03-mini). To the best of our knowledge, are not aware of any model specifically "fine-tuned" for the CUDA context (efforts like Nichols et al. (2024) focus on OpenMP CPU code). CUDA, or GPU code in general, is extremely sparse in the pretraining corpus, only 0.073% of the Stack (Li al. 2023) code corpus; this makes approaches that depend on readily available data (such as "fine-tuning" difficult. Hence, this data challenge actually highlights the value of our RL-based approach, as we discussed Section [ We believe that the comparisons of Kevin against SoTA general-purpose L.MMs are fair and fitting, and actually demonstrate the advantage of our RL-based approach in this domain.
Our baseline comparisons, o4-mini and o3-mini, are frontier models that achieve SoTA on challenging code generation benchmarks. Specifically we use 04-mini 1i-miinn222:: 04-16 and o3 33-mini-- -001--205---- 10002-01--1 31. As shown below, o4-mini demonstrates significant lead over our base model QwQ-32B, especially on challenging real-world software tasks such as SWE-Ba Bench (Jimenez et al.] 2024) and Polyglot (Gauthier 2024). Hence, results in Section [5.] and Table [] showing Kevin ppsst-trained QwQ-32B with multi-turn RL) exceeding o4-mini should be noted as significant improvement and demonstrate method's effectiveness.
24

---

Published as a conference paper at ICLR 2026
Model
AIME 24
LiveCodeBench
SWE-Bench Verified
Aider Polyglot
QwQ-32B
79.5%
63.4%
41.3%
20.9%
04-mini 93.4%
74.2%
68.1 %
72.0 %
Table 6: o4-mini shows significant lead over DwQ-32B over a variety of reasoning, coding, and software engineering benchmarks (Mathematical Association America 2024, Jain et al. 2024 Jimenez et al. 2024 Gauthier 2024); Kevin is post-trained on QwQ-32B and shows improvement over both Q and 04-mini, as shown in Section 5.1
E.3 EVALUATION ON KERNELBENCE LEVEL 3
While we focus our training and evaluation mostly on KernelBench Level and 2 (Section 3 we were also curious and explore testing Kevin on KernelBench Level tasks. They are longer and more challenging (rather than single a few operators), requiring the end-to-end optimization of full model architectures, such as the Vision Transformer, and miniGPT attention blocks. Kevin trained using a subset (180) of the KernelBench Level and tasks (single and sequence of operators), and Level 3 tasks are completely unseen. We evaluate the multi-turn (Kevin), single-turn, and base model (VwQ-32B) on the 50 level tasks following the same evaluation setup as Section 5, As shown in the table below, multi-turn RL can also generate much faster kernels for these much more complex tasks over both single-turn RL and the base model.
We view Level 3 primarily an out-of- of-distribution test: these tasks involve full model architectures with much longer-horizon reasoning, and requiring both kernel generation and effectively dealing with long context. We do not train on any Level 3 tasks as the length of these programs would lead to context explosion (Section 4.1). Hence, our main analysis focuses on Levels and 2, which better focuses kernel generation performance with more controlled conditions.
Correctness
Performance
Method
best@1 16
avg 16
best 16
avg@
Multi-turn RL
36 %
11.75% %
0.41
0.08
Single-turn RL
36%
8.38%
0.36
0.06
QwQ-32B
4%
0.25%
0.04
0.002
Table 7: Multi-turn RL achieves improvements also on the completely unseen and more complex KernelBench Level 3.
E.4 EVALUATION AGAINST CORCILCOMPILE BASELINE
torch.compile uses various rules and heuristics to optimize the performance of PyTorch program, so in theory it is stronger baseline. However, it is a less reliable baseline than Py PyTorch eager. As pointed out in KernelBench ((Ouyang et al. 2025)), the runtime can vary significantly across hardware platforms). Nevertheless, to substantiate the validity of our results against potentially stronger baseline, we run additional evaluation here with the baseline computed by torch. torch.compile default Inductor backend.
Under both the eager baseline and the torch compile baseline, multi-turn's best performance exceeds both its single turn counterpart well as all other models evaluated. We note that torch large number of operators. Given
compile effective when the computation graph is composed that we focus optimizing the performance of either a single a few operators, the impact of torch compile tends to be small, and for some kernels even negative, given the additional runtime launch overhead that torch compile introduces.
25

---

Published as a conference paper at ICLR 2026
Correctness
Speedup (eager)
Speedup (compile)
Method
best@ 16
avg@@ 16
best@ 16
avg@ 16
be
avg@ 16
Multi-Turn RL
82%
46 %
1.10x
0.40x
1.04x
0.36x
Single-Tur RL
82%
45 %
0.85x
0.35x
0.91x
0.37x
QwQ-32B
56%
11 %
0,53x
0.08x
0.51x
0.08x
sonnet-4
71%
34%
0.69x
0.26x
0.71x
0.27x
o4-mini
38%
22%
0.78x
0.27x
0.61x
0.20x
gpt-4.1
36%
14%
0.25x
0.09x
0.24x
0.09x
03-mini
27%
8%
0.30x
0.08x
0.22x
0.07x
Table 8: Correctness and speedup comparison across models.
E.5 EVALUATION ON 100
In mentioned that speedups are only accurate for the pre-defined tensor dimensions and on NVIDIA H200 GPUs. To demonstrate the effectiveness of multi-turn RL transfers to other hardware architecture, we use the same evaluation setup on NVIDIA A100 GPUs, which use the Ampere architecture (instead of Hopper for H200s).
important to note that the models we evaluate were trained with execution feedback on H200s, but the speedup and correctness they see at inference-time are based on A 100s, where the hardware layout different. For example, Hopper GPUs have up to 228 KB of shared memory per SM, but Ampere GPUs have only 164 KB, so the trained models might think they have access to more shared memory than they actually do, leading to errors in the kernel refinement process. In spite of the difference in the hardware architecture during evaluation training, results show the same trend as H200s, that multi-turn RL achieves better best@ 16 and avg@ 16 correctness and speedup than single-turn RL and the base model.
Correctness
Performance
Method
best@ 16
avg@ 16
best@ 16
avg 16
Kevin-32B
79%
44%
0.88
0.33
Single Turn
78%
42%
0.71
0.29
QwQ-32B
53%
12%
0.45
0.08
Table 9: Multi-turn RL generalizes to other hardware architectures, also surpassing the single-turn RL and base models.
F TRAINING STABILITY
The analysis of the "not okay ratio" led believe that model instability caused the appearance of nonsensical and repetitive outputs. Therefore, we attempted runs where we enabled KL divergence penalty in the GRPO loss, which would penalize the model from deviating from the base policy too much. Following DeepScaleR (Luo et al., (2255)). we set the KL coefficient to 0.001 and attempted an ablation run. However, we found that the reward plateaus with KI. enabled, suggesting that the KI. penalty slows down learning. Thus we attempted other techniques of constraining the model from deviating into regions of instability, such as clipping the gradient norm aggressively which was effective in our setting,
We use refinement turns at train-time for efficient training. During test time, we can afford more extensive test-time compute, we evaluate on 8 turns instead of turns.
G REWARD HACKING
Here we present excerpts from generated kernels that show signs of reward hacking, previously mentioned in Section 6.2
26

---

Published as a conference paper at ICLR 2026
multi-turn training KL
0.7
0.6
70
Gradient
Figure 14: Adding a KL penalty slows down learning, Here we conduct an ablation with KL coefficient 3 0.001 versus 3 O. We see that the reward plateaus with KL enabled.
In the following example, the model simply copies the PyTTorch reference implementation, thus getting rewarded for generating a correct answer with 1.0x speedup. To prevent this, we modify our kernel evaluation environment so that it checks each generated kernel if contains instances of torch. nn or torch. nn. functicnal We assign reward of O to those.
class ModelReLl (Module)
def forward (self, x) relu = forch.nn.Reiu return (x)
Similarly, the model wraps incorrect implementation of the CUDA kernel in try-except statement and invokes the PyTorch implementation functions as a fallback. To prevent this, we assign a reward of kernels that contain try or except.
class ModclReIl (Module):
def (srward x) try:
A# CUDA impl implementation
except Except as e: print ("Custom ReIUU kernel failed to compile. Using defau t ReLU instead.")
relu torch.nn.ReUU return relu(x)
Here the model inherits from the reference implementation, bypassing the need for a CUDA imple- mentation. To prevent this, we assign a reward of to kernels that contain pass.
class ModelReil (Module)
def (self, x) \ CUDA implementation
class Mode ReLU (Model): pass
\# Model is the the reference implementation
H KERNEL HIGHLIGHT
As example, we look the task form evaluation set, Conv3D_GRINJ_TE Tanh 2anh_Clamp:
27

---

Published as a conference paper at ICLR 2026
import torch
import torch.an an as nn
class Model :on.Module)
=**
Model using Conv3d, GELU, Clamp, Tanh
def iseft_(seff, in chanhann,,,, out out_channels, kernel_size,
clamp_min, clamp_nax);
super (Model, self) (.ff)__nnne (nit_
10
self.conv nn. Conv3d chn_cannee,,,,,,, out_channels, kernel_size) self.clamp_min clamp_min self.clamp_max clamp_max
11
def forward(self, x)
15
Forward pass of the model.
Args:
19
X (torch. Tensor) Input tensor of shape (batch_size, in in__hannels, depth, height, width).
Returns:
22
torch. Output tensor shape size, out cht_channes,, feoth--kerel_sizee height-kernel_eleee widta-kkrrel_sieet))..
self. (x)
25
x torch.tane (x)
26
x = torch. torch.clamp (x, min=self ctresslffcaamppp.. min, max=self. clamp max)
= torch.nn nn. funatienall (xllll
The task made up of a computationally intensive function (3D convolution) and a series of lightweight functions (tanh, clamp, gelu).
Note: all the speedups reported below are measured over torch in eager mode, as 'torch.compil" support varies heavily across different platforms. When compiled using torch.compile, the kernel achieves a speedup of 1.32x over its eager counterpart. Thus, this trajectory achieves maximum speedup of 1.93x over eager and 1.46x over torch compile.
The model implements a correct solution in the first turn by fusing the kernels, but it only achieves speedup of 1.06x over the reference implementation.
import torch
from torch. utiiss hwtpssppppxxtttteeeeeeeee import load inline import math
class Model New Module))
def init init_nnnttt_____ (self, in_ehannels, out cht_nels,,,,, Bernel_size clamp_ min, clamp_max):
max)
super (ModeeNew, self) init_ ()
self.clamp_mi self self.clamp_max clamp_max
clamp_min
11
Initialize convolution weights and bias
12
self, weight = Barrch.nn.Paramete (torch. empty (out_channelss
in cn__hannels, kernel_size, kernel_size, kernel_size))
self.bias torch. torrhh....Parmeteertttnnnt
Neight initialization following PyTorch's is Conv3d default
15
torch.nn. init. aazchha.....t..aaammm.mmmmmm.oe.
a=math.sqrt
fan
cou__cclluuaa atnaat__fe ft f.wighht)
bound 1 nath (fan_in)
28

---

Published as a conference paper at ICLR 2026
terch.nn. www.iiiiii.iiii..... uniform (niifommmsel..biasss -bound, bound)
20
Define the fused convolution and activation CUDA kernel cenv_activation_see
21
=
#include storch/extension.
#include mrada_runtimg h>>
24
#include <math.h>
25
26
global_ void conv www__activatio__ketee (const float* input, const float* weights, const float* bias, float* output, int N, int CLiny int D, int H, int H, int Coout, int kD, int kh, int kw, float clamp_min, float
28
clamp_men max) I
int idx blockIdx.x blockDim.x threadIdx.x;
30
const int __out const int HLout = ki 1;
kW 1;
31
32
const int D_out = - kD + 1;
const int coements N CLout D Bouut Woout; if (idx >= total_elements) return;
36
// Decompose index into output coordinates int v_out idx & __ut;;
int b_out (idx W [_uut) _out;;
39
int dou (idx (W out H [out]] out,;
40
int f (idx / (W_out H_out D_out)] C_out;
41
int n idx / (C_ * D_out Houut w
float
44
for (int a 0; c C_in; ++c) {
45
for (int kd = 0; kd < kDj ++kd) {
46
for (int kh = 0; kh < kH; ++kh) {
for (int kw = 0; kw < kW; ++kw)
48
// Input tensor index calculation
49
int d_in d_out kd;
50
int h_in h_uut kh;
int wiin w_out kw;
int input_idx
C_in
53
54
d in * H
h iin * * +
w_in;
// Weight tensor index calculation
58
int weight_idx = C_in KD * KB
kw
59
* kD kH kw +
kd < kH kW
kh kW
62
kw;
63
sum += input linput_ idx] *
weigats [wwight__ddde idx];
66
}
67
sum += bias[f];
// Apply activations sequentially sum = tanh (sum)
71
72
sum = fmaxf (clamp_min, fminf (sum, clamp_max) )
float inner sum 0.044715f sum sum sum; inner 0.79788456f; // sqrt (2/pi) approximation sum = sum 0.55 * (1.05 tanh (inner)
75
76
29

---

Published as a conference paper at ICLR 2026
// Output tensor index calculation
int output_idx CLout * DLout Houu * RLout f D_out BLout woout +
81
82
dout Boout Mouut +
h_out R_out
w_out;
85
cutput [output_ idx] sum;
36
torch: Tensor aonvaactivttion_.o cuda (torch: :ensoor input, torch:: weights, :orshs bias, float clamp_min, float
89
clamp_max) {
// Dimension extraction
int N input.size (0);
92
int C_in input.size
93
int D input.sizc(2))
94
int B
int w input.
int C_out =eeights.siee(0)
97
int kD weights.size
98
int kH weights.sizee3) (3);
int kW weights.sizel (4);
101
// Output dimensiens
102
int D_out D kD + 1;
103
int __out и kHl + 1;
int W_out W kW + 1;
106
// Create output tenser auto output = torch: terphy:e C_out, D_out, __out, W_out input (ppinnoptiions))..
107
109
// Launch kernel int total_es ements tomat_slements const int block_size 256;
110
* CLout: + D_out * Houut + W_outz
111
const int grid_size -
block_size, 1)
block_sizer
113
conv activation wtw.aattttttooeeeeeeeeeeeeeeeeee kernelce<grid size, block
block_sizee>e size>>>
114
input. .anut.ddte data caaatt_prrfflltt//ot (),, weights. ptr<flost>> ptteitss (),
bias. data ptr<float>> (),
output, output.dated data_ptr<float>0) N, C_in, H, W, C_out, kD,
kH, kW, 116
clamp_min, clamp_max
117
);
118
return output;
120
121
122
conv conv_acte headerr :arch::Tenser conv_act :ornhs:lennor weights, inntivvaiennccuuee (torch: :Tensor input, bias, float clamp_min, float
conv_activation vat
124
clamp_ max);
125
127
Lead the CUDA kernel
128
sclf colf..cov_acttiattion = inaininline
129
nane="con__cciivationn
auda_oourceeecmm..aatiiiiii....e verbose=True
132
133
30

---

Published as a conference paper at ICLR 2026
def (srwa x)
137
return self.covv_ctivationn otttlia..oo__sttvattiioooooo (x, self self.weieg weight, self.clamp_max
self.blas,
On turn 2 attempts to improve the existing implementation, but fails due compilation
On turn 3, it attempts to modify the loop order to improve coalesced memory access, but this actually worsens performance, as the kernel now exhibits speedup of 0.61x. The model attempts this strategy again at turn 4, and is able to order the loops in a way that increases memory coalescing: this brings up the performance 1.19x. At turn 5 and 6, the model attempts some minor optimizations (such as precomputing certain indices, or deleting intermediate variables), which however has no measurable effect in performance.
On turn 7, the model figures out surprising optimization: it records the convolution's 's kernel size as a compile time constant, which allows it use # pragma unroll to unroll the loops needed to apply the convolution. This results in significant speedup of 1.93x over baseline.
import torch
from torch. www.ttillss..... www.lls.cp_eeeteeeee import load inline import math
class noce
def iin_i_ii____l_____ (self, in_channels, out_channels, kernel_size, al clamp_min, clamp_max):
super (MadelNew, self) init__ ()
self self.clamp_min = clamp_min
self.clanp max clamp max
11
Initialize convolution weights and bias self.weight =
12
(torch. (out_channels,
in channels,,,, kernel_size, kernel_size, kernel_size))
sclf.bias wwr.ch....aaammtrrrreee..oo.....
empty (out_channels))
Weight initialization following PyTorch's Conv3d default
15
terch.nn. in. www.iint.kimmmngmmiffrr... (self. a-math.sqrt (5)
16
fan_in, torch. nn. wtt.taan.iiinttiooo...... wtw.tt._csouuattetaaaaaaaaa
far_ (self. bound 1 math. (fan in)
18
torch.nn. nn. init (sw.dioo__mmeee_bbbii -bound, bound)
19
20
Define the fused convelution and activation CUDA kernel with optimizations
kernel size conv cew_activation.
kernel_st siie
22
sourcesoureeeeeeeee gunn #include storch/extensicn. Hinclude Scuda_runtime #include <math.h>
23
24
26
27
#define SERNEL_SI22 (kernel_size_vall _all
28
#define KERNEL_D [kernel_size_val]
#define KERNEL_ [kernel_size_val]
#define KERNEL [kernel_size_aal)
31
32
global void cony conv_activation a
33
const float+ ressrrcc___ input, const float+ restrict_
weights, const float= ressrric___ bias,
float restrict__ cutput,
35
int N, int c_in, int D, int Н, int W, int CLout,
36
float clamp clamp_min, float clamp_mase I|
int idx = blockIdx.x int Woout
blockDim.X threadIdx.x/
KERNEL 1;
39
const int Lout
KERNEL 1;
40
const int Doout = D KERNEL_D + 1;
const int
N Coout D_out E_out
W_outz
if (idx @otal_element eleeents) return;
31

---

Published as a conference paper at ICLR 2026
int w_out idx $ w_out;
45
int h_out (idx w_out) $ _ouut;
46
int d_out (idx / (R_out [Lout]] $ D_out;
int f (idx / [W_cut o_uut (_out)) $ C_out; int n idx / (C_out D_out Houut R_out);
49
50
float sum 0.0;
// Precompute constants for GELU const float sqrt_2_over__r 0.79788456f, const float apprcx_coeft 0.044715f;
55
for (int c 0; c C_c_nnn in; ++a) {I
// Compute input channel base pointer
const int input_c_offsed n C_in D H
* D * W;
59
const float input_e input input_c_offsct;
// Compute center position in the input channel const int center_offsed dut W h_out W
62
+ +_outt
63
const float* center = i + center_offsed
65
// Meight base for current output channel f and input channel c
66
const int weight_base f. CLin KERNEL_
KERNEL_H H KERNEL_W c KERNEL_D KERNEL_H KERNEL_W;
#pragma unrell
69
for (int kd = 0₃ kd < KERNEL_D; ++kd) [[
foragma unrell
for (int kh 0; kh < KERNEL_H; ++kh) 11
#pragma unroll
73
for (int kw - 0ş kw < <RRNRI_W> ++kx) 11
74
censt int offset - kd * H W + kh W +
kw
float input_val center [offset] const int wooff weight base +
76
KERNEL H KERNEL W kh
KERNEL W kw;
const float weight_val val = weights [w_ofE]>
sum input_val weight_val;
30
F
81
}}
1)
}
85
sum += pias[fl;
//
Apply activations with optimi zed order
{{
39
float tanh_val tanhf (sum)
90
float clamped fmaxf (clamp min, fminf staan__va,,
clamp_ max)) );
float inner clamped soprox_coefé clamped
clamped clampod;
92
inner += sqrt_ 2 __over_pii
93
sum = clamped 0.5f (1.0f tanhf (inner)
}}
96
// Output index calculation const int output_base n CLout D_out * BLout
97
M_out
D Dut oout Woout
32

---

Published as a conference paper at ICLR 2026
out H out W_out
100
houu Woout;
101
output [output_base + w_out] sumy
102
}}
torch: : Tensor conv_activation (torch: :ensor input, torch: Tensor weights, torch: Tensor bias, float clamp_min float
105
[aamp_max) II
int N iaput.sizen (
int C_in input si (1);
108
int D input. size (
109
int H input. size (3);
int W input. (
int C_out weights size (0)
112
113
int D_out D KERNEL_D 1;
134
int Bouut H KERNEL_ + 1;
int oou W KKEREE_ 1;
117
auto output torch:::eepppp :enptteemttyyyty ((O, C_out, D_out, _ Hhouth input.options ());
118
C_out D_out out W_uut; const int block size 256;
int total
120
121
const int grid_size - (total_elemen) block_size - 1)
block_sizer
conv activation wrw.aiionnkemeeekkgridd..d size, block sick_ssie>>>((cc (),
124
input. datu Saata_tt<ffllat/at (),, daaa derssdata_ptsfllat>>>> (),,
125
output. data__tr<fflaa>t N, c_in, D, B, W, C_out, clamp_min, clamp_max
127
);
128
return outputy
129
}} #B
130
132
conv conv_act caw_activation header с.*
133
torch: : Tensor conv conv_ activation_noae cuda (torch: :ensor input, torch:: weights, torch: Tensor bias, float clamp_min) float
clamp_max)
135
**"
136
Load the CUDA kernel with fast math optimization self. con__activaiinon Icad_inline(
139
name="cony.activation
140
header,
funciioonssttoonracactttttoooood (xtra__w@__fllag==---eee___ verbose-True
144
145
146
def forward x):
148
return self. cov wtw.riooo..cooeeetttttttooooooo (x, weight self.blas, self colf.clamp_min, celf.clamp_nan
In its final turn, the model attempts a more advanced implementation that further parallelizes the computation across kernels before performing warp-level reduction. However, it fails to implement the strategy correctly, due to applying the reduction across the wrong axis. We do note the model has shown success in implementing complex warp reductions in several other tasks.
33

---

Published as a conference paper at ICLR 2026
I ERROR CORRECTION BEHAVIOR
An important quality for kernel generation is to correct errors over turns. With multi-turn training. Kevin learns to effectively manage kernel writing across multiple turns, making more aggressive optimizations while correcting errors more effectively. This is shown by our results in Figure P where Kevin exhibits better scaling behavior across serial turns (compared to its single-turn counterpart). Here we focus on an example that illustrates error correction behavior.
We consider the KernelBench task of Casi meiimiaarityyooss (Level Task 97):
prediction [128, 4096] (128,
target [128, 4096]
def (searr predictions, targets); cosine_sim torch.nn. functional, wttp://wi....iiiennnnnnnaaatttte (pr.larriityppeeiitiiii... targets, dim-1) return torch.mean (1 cosine_sim)
Note: the speedups below calculated with respect to the eager baseline. For this kernel, torch compile causes a slight degradation in performance (0.96x speedup over eager), likely due to additional overhead introduced.
Test-time Trajectory from Single-Turn RL:
Turn 1: The model implements correct kernel the first try. The kernel is straightforward yet performant, obtaining 3.15x speedup over the reference implementation. This is expected as the single-turn model is trained to maximize single-turn correctness.
Turn 2: The model attempts to calculate the block size dynamically; however, it results in suboptimal GPU utilization, causing decrease runtime 2.19x.
Turns 3-5:
The model reverts the change from Turn 2, and adds minor improvements (such as performing the mean calculation the warp level rather than using single thread, tweaking the loading of shared memory); however, it has no significant impact on performance (the speedups range from 3.1 1x to 3. 3.19x).
Turns 6-8: The model attempts to tune the grid size and block size, though this fails to improve performance (the speedups ranging from 2.99x to 3.x)...
Test-time Trajectory from Multi-Turn RL:
Turn The model attempts a kernel similar the single-turn implementation. A suboptimal block size and a non-contiguous access, however, cause speedup of 1.74x, inferior to the single turn version.
Turn 2: The multi-turn model makes various optimizations (such as tweaking block and grid size, and avoiding recomputing certain values), but it forgets include required C++ header, causing compilation error.
The multi-turn model adds the correct C++ header and improves the memory coalescing. This, combined with the previous optimizations, causes speedup of 3.3
Turn 3:
Turn The model now attempts more aggressive optimization techniques to compute the sums within the kernel, namely using warp intrinsics to perform warp-level reductions. The implementation incorrect, however, as only a subset the threads for given warp execute the reduction, resulting in deadlock.
for (int s - 8; 8 >= 1; s >>= 1) {
if (thread_id s) float temp_ dot dct += temp_dot; float temp_p p_norm_sq += temp_p;
sh1f ddwn__wwweccccceeeeeeee (OxFFFFFET5, det, s)
shwl_ddww__sssshhee_oeo__posssss synessyddwsseeeeeeeeeeeeoe (OXFFFFFFEF, P_norm_sq, s);
float temp_t = ssffl schfl_down__ymee (OXFFFFFFEF, t_norn_sq, s); t_norm_sq temp_t;
34

---

Published as a conference paper at ICLR 2026
syncthreads
Turn 5:
The model then tries different strategy, implementing a simpler reduction that avoids using warp intrinsics. However, this proves less efficient than the previous methods, and achieves speedup of only 2.26x.
Turn 6: The model now attempts fix the errors present in the advanced warp level technique from turn 3, spending most of the chain-of thhuuhthtt in trying to address the previous error. This time, the model is successful with speedup of x.8lx.
for (int 16; >= 1; >>= 1) flcat temp
shfl down (OXFFFFFFFE, dot, s);
dot += temp;
temp = shfl_ddww_ssp_ (OXFFFEFFFF, p_normesg, 8);
P_nerm_sq += tomp;
temp shfl_downn (OXFFFFEFEF, t_norm_sq,
t_nerm_sq temp;
Turns 7, 8: The model tries few other optimizations, yet has significant impact on the kernel performance, resulting in slightly worse kernels. The best kernel thus achieved turn with speedup of 7.81x, compared to the best kernel found by single-turn (3.19x).
Comparing the trajectories of the multi-turn vs single-turn model, we notice that the single-turn model is generally careful in its optimizations and rarely deviates from the previous attempts. The multi-turn model attempts riskier strategies. While may result in wrong kernels sometimes, is more capable of recovering from these errors and eventually finding more performant ones.
35


Reviews:
[{"review_id": "O3yJ18x2TQ", "reviewer": "Reviewer_5ts9", "summary": "In this paper, the authors propose an RL training method to generate high-performance CUDA kernels by using multi-turn RL training, known as Kevin. In contrast with previous code-generation studies that focused only on correctness, Kevin starts from a reference PyTorch implementation and lowers it to CUDA, monitoring both correctness and performance. To accomplish this, the authors assign a score to each kernel that balances correctness and performance, and train an LLM using GRPO on samples without immediately iterating on external feedback, as is typical of single-turn training. Instead, during each training step, multiple responses are evaluated and assigned a reward. During multi-turn training, not only is the code used to stir the process, but the chain of thought as well. However, with the aggregation of so many instances, the CoT context becomes quite large, so the authors summarize the context from earlier runs. Two approaches to assigning scores were investigated: a greedy approach that assigns a kernel score to each turn and an outcome-based approach that assigns a score to all turns based on the best score in the trajectory. The authors propose a hybrid approach that balances the pros of these two credit assignment strategies and perform an ablation study to support the efficacy of this choice. The evaluation section illustrates the improved performance of multi-turn training compared to single-turn alternatives and other LLMs to generate code. The code generated by KL yields a higher rate of correctness and achieves improved performance in most tests.", "questions": "- Is it possible to add a graph to illustrate the issues related to smaller models and reward hacking in Appendix B.6? Not strictly necessary, but it would provide more context for readers who are not as familiar with these issues.", "limitations": "", "rating": 8, "confidence": 3, "soundness": 3, "presentation": 3, "contribution": 3, "strengths": "- This is my first time encountering multi-turn in the context of code generation. The results support its usage from both a correctness and performance perspective. \n- The authors investigated multiple aspects of the RL training procedure, from the generation of the training data, reward assignment, and the composition and structure of samples during training. All of these training hyperparameters were shown to have a notable impact on the quality of the result and should provide a useful data point for other researchers in the area, considering multi-turn RL.\n- The issue of context length control seemed to be an interesting, but maybe ephemeral problem, as the context lengths continue to grow. The solution they proposed did not seem detrimental to kernel improvement.\n- The tradeoff inference configurations, illustrated in Table 2, were an interesting display of the performance impact of trajectory vs turns on both the performance and correctness.", "weaknesses": "- I found the discussion of the choice of baseline model in Appendix B.6 to be insufficient from the reader's perspective. While it seems completely plausible for the largest model to have the best priors and smaller models more susceptible to reward hacking, it may be the case that certain updates to the reward function or training could alleviate these issues.\n- One of the major limitations, noted by the authors, is the limited number of robust tasks usable for training. With access to more tasks, the kernel generation capabilities of Kevin could be significantly higher.\n- All the figures are difficult to digest for people who struggle with differentiating colors. A change in the markers and/or line styles would make it easier to tell the differences at a glance.", "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": "VAbTAswhV5", "reviewer": "Reviewer_PBZv", "summary": "Kevin is a multi-turn reinforcement learning (RL) based framework to train LLMs to generate better GPU kernels. Authors propose a design for multi-turn RL with a reward function proportional to correctness and generated kernel's performance. Authors have used GRPO for training the model. During training authors sample m parallel trajectories with n iterative refinement steps. Authors also propose to compress CoTs from prior attempts to save on context length available for generation. Authors have demonstrated that single-turn RL saturates early on and does not benefit from the refinement aspects for test-time compute scaling. Authors also propose to training on each turn in the multi-turn realization. Authors aggregate the reward across multiple turns with two specific discount factors 0.4 and 0.8. Multi-turn training behaviour shows reward monotonically increasing on an average. Authors have used 180 tasks out of 200 (from level 1 and 2) tasks from kernelbench benchmark. Remaining 20 and additional author created 80 tasks are used for evaluation. Authors compare speedup across pytorch implementation provided in tasks. The paper also shows that test-time scaling trends outperform the baseline LLM scaling. Further, authors have shown that Kevin benefits sequential axis of test-time scaling than that of parallel.", "questions": "Please see weaknesses section.", "limitations": "", "rating": 2, "confidence": 4, "soundness": 1, "presentation": 1, "contribution": 2, "strengths": "- Demonstration of multi-turn RL for GPU kernel generation. \n- RL reward formulation as a function of correctness and speedup.\n- Effectiveness of multi-turn RL to improve inference time scaling trends.\n- Clearly comparing gains against single-turn RL.", "weaknesses": "- The evaluation methodology does not follow standard practices. Kevin trains on 180/200 examples from kernelbench evaluation benchmark. The authors must create their own dataset for training and then evaluate on kernelbench. This does not inspire fair comparison with existing approaches and goes against standard practices. \n- Paper does not specify where are the initial CoTs obtained from.\n- In section 4 line 237, the description is unclear. \n- In subsection 4.1: Summarizing all previous CoTs does not appear to be the best way to build context. There are several prior methods that have shown using prior top-k + bottom-k attempts to build context that results in better contrastive learning.\n- subsection 4.2 appears to be incomplete.\n- Performance is evaluated with Pytorch implementation as a baseline. However, even if the benchmark provides reference Pytorch implementations, these are not truly optimal implementations available today. I encourage authors to compare speedup against torch.compile as their baseline. This will show true improvements over manual kernel optimization which is a very time consuming and tedious task. This will justify all the effort and energy spent in training/inference/generation of LLMs.\n- There is no profiler feedback integrated. Neither during training nor during inference. The profiler feedback is very crucial which helps human experts to reason and write best strategies to improve kernel performance. \n-  Authors compare against o3/4-mini models which are not optimized for code. They should instead compare against GPT4.1 and claude-sonnet-4 which are optimized for code.\n- With best@16 Kevin shows 1.10x performance improvement. Which does not inspire any confidence in the efficacy of Kevin given that 10% gain could come from measurement errors. \n- There are some case studies in appendix showing speedup of 1.9x, 2+x, 3+x; all these speedups are over naive pytorch calls. torch.compile significantly optimizes a given pytorch code by performing operator fusion on the fly using a pool of manually optimized kernels. This should be the ideal baseline that AI solution must beat.", "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": "UTjOaJVNDy", "reviewer": "Reviewer_cuJN", "summary": "The paper proposes Kevin, a CUDA kernel generation model trained with multi-turn RL. Instead of producing one kernel in one shot, the model iteratively refines kernels over several turns, using runtime and profiling feedback from previous attempts. The training pipeline assigns discounted reward across turns and treats each refinement step as its own RL sample, with context summarization to keep prompts short. \nOn 100 held-out CUDA tasks (KernelBench style), Kevin matches a strong single-turn RL baseline in correctness (82% best@16) but shows somewhat higher peak runtime speedup (1.10× vs 0.85× best@16). \nThe paper argues that this proves multi-turn RL improves the ability to iteratively optimize kernel performance and scales better when you allow more refinement turns at test time.", "questions": "Q1: Multi-turn RL and single-turn RL have the same solve rate (82% best@16) and same fast1 rate (43%), and only a modest speedup gain (1.10× vs 0.85×). In what concrete situations does multi-turn RL succeed where single-turn RL fails (either solving a task it couldn’t solve, or finding a faster kernel it couldn’t reach)? Please quantify.\n\nQ2: Kevin is trained starting from a strong 32B base model. You mention weaker bases tend to hack rewards or collapse.\nCan you show any numbers (even partial) for smaller models, or describe concretely how and why training fails there?\n\nQ3: Your method bundles many heuristics (CoT summarization, discounted reward shaping etc). Can you provide ablations that isolate the effect of multi-turn RL itself versus these stabilizing heuristics? Without that, it’s hard to attribute the reported gains to “multi-turn” rather than to engineering tricks.", "limitations": "", "rating": 4, "confidence": 3, "soundness": 2, "presentation": 2, "contribution": 2, "strengths": "**Realistic problem setup.** The work models how GPU performance engineers actually iterate: propose kernel → profile → refine. The RL formulation explicitly credits early partial attempts that later lead to a fast kernel, instead of rewarding only final outputs. \n\n**Clear engineering advances.** They introduce (i) per-turn training on every refinement step to improve sample efficiency and (ii) discounted future-return style reward aggregation for credit assignment across turns. \n\n**Beating strong general models.** Kevin, post-trained from QwQ-32B, outperforms powerful proprietary baselines like o4-mini and o3-mini on this CUDA kernel benchmark, even though those models are otherwise stronger than QwQ-32B on standard coding and reasoning tasks. \n\n**Test-time scaling insight.** The paper studies how performance scales with more refinement turns and shows the multi-turn–trained model benefits more from extra turns than single-turn RL, suggesting better iterative optimization ability.", "weaknesses": "**Performance gains are marginal vs the main ablation.** Compared to the single-turn RL baseline, Kevin does not improve solve rate: correctness best@16 stays 82% vs 82%, and fast1 best@16 stays 43% vs 43%. The main improvement is higher best-case speedup (1.10× vs 0.85×), and average speedup only rises slightly (0.40× vs 0.35×). \nSo the method is not solving more tasks; it’s mostly squeezing somewhat better runtime on tasks that were already solvable.\n\n**Heavy system heuristics, limited theory.** The approach bundles many practical tricks (turn-by-turn training, summarizing CoT to avoid 50k–100k token contexts, discounted reward shaping). \nBut these are presented as an engineering recipe rather than a principled analysis of why multi-turn RL should outperform single-turn RL.\nDetailed ablation study would help us understanding real contribution of introducing multi-turn.", "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": "yEwQ7KBqZe", "reviewer": "Reviewer_dQdu", "summary": "This paper introduces \"Kevin,\" a model trained to generate and optimize GPU kernels using a multi-turn RL framework. The study think that the iterative nature of code optimization is best captured by a training process where the model receives and learns from execution feedback over multiple rounds. The core is a multi-turn RL recipe designed to overcome challenges inherent in long-horizon tasks. By explicitly modeling the \"generate -> execute -> get feedback -> refine\" loop, the training process aligns more closely with how human experts work. Evaluation results on the KernelBench dataset demonstrate that this approach is highly effective, showing that Kevin achieves significant improvements in both the correctness and performance of the CUDA kernels it generates compared to its base model and other strong baselines.", "questions": "Please check weakness.", "limitations": "", "rating": 6, "confidence": 3, "soundness": 2, "presentation": 3, "contribution": 3, "strengths": "1. Well-motivated multi-turn RL framework to solve the Kernel Bench CUDA task. And the approach is validated by strong empirical results. Not only surpasses its base model and a single-turn RL counterpart but also outperforms strong frontier models like OpenAI 04-mini on both correctness and performance metrics. \n2. Clearly demonstrate that sequential refinement is more compute-efficient than parallel sampling. \n3. Offers practical engineering solution, detailing robust solutions to reward hacking probelm.", "weaknesses": "1. The most significant weakness is the reliance on the relatively small KernelBench benchmark, which contains only around 250 tasks. The scarity of training data on this specific domain will raise my concern on the generalizability compared to other large language models. \n2. The paper mentioned that \"only accurate for those dimentions and NVDIA H200 GPUs\". The reported speedup performance improvement are highly specific to H200 evaluation environment within KernelBench. More analysis on the methods potential performance transfer to different hardware architectures or real-world workloads remains an open question, further validation on other benchmark environments are needed.", "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)