KernelZero: Co-Evolving Proposer and Coder for Continuously Improved GPU Kernel Generation
Changxin Ke, Rui Zhang, Zixiang Fang, Zhenghong Li, Yuanbo Wen, Jiashuo Shen, Shuo Wang, Jiaming Guo, Ling Li, Qi Guo, Yunji Chen
cs.LG, cs.SE
2026-09-27
A 7B Proposer-Coder pair co-evolves GPU kernels, reaching 75.8%/69.6% CUDA pass@1 on KernelBench L1/L2 (pass@10 100%/97%) and 77.2%/72.5% on Triton.
LLM kernel generators can already submit CUDA and Triton on KernelBench, but training is still expensive and the objective fights itself. cudaLLM spent 128 H100 GPUs on supervised training and still reached only 81% pass@10 on Level 1 with an 8B model. Static corpora sit at the wrong difficulty: too easy and RL never sees the failure modes that matter; too hard and the gradient is mostly noise.
Reward design tightens the knot. Correctness-only training yields conservative kernels that barely beat PyTorch. Pushing speedup too early breaks functional equivalence. KernelZero treats those as two coupled problems: keep generating modules that sit on the model's current failure boundary, and keep performance optimization gated until correctness is stable.
Two specialized 7B models start from Qwen2.5-Coder-7B. The Proposer is the question writer: it consumes a set of Torch APIs and emits a runnable nn.Module. The Coder is the kernel writer: it translates that module into CUDA or Triton. Each side is updated with its own reward while the other is frozen, four rounds of 20 steps each, 80 steps in total. Alternating updates keep the question distribution stable long enough for the Coder to learn, and give the Proposer a fixed Coder whose failures mark the frontier.
API lists are sampled from a co-occurrence distribution over real Torch modules: 4,000 combinations, half two-API and half three-API, covering 442 APIs. Every proposed module must pass three checks: an AST match against the requested APIs, an actual GPU run via getinitinputs / getinputs, and a NaN/Inf filter.
Frontier Reward, used in FO-GRPO, depends only on the Coder's mean correctness on that module. Payoff peaks when the success rate is near 0.5; trivial and impossible items score lower, and invalid modules get a flat -1. The Proposer is paid to hunt questions the Coder currently gets right about half the time.
The Coder is cold-started first. From cudaLLM's public 71,996 instances, gpt-oss-120B rewrites kernels with chain-of-thought under four compiler heuristics (Tiling, Fusion, Pipeline, Reordering). Only correct traces are kept: 42,454 CUDA and 59,998 Triton, then SFT. The four heuristics map to data reuse, operator fusion, overlapping memory with math, and access-pattern reordering.
CA-GRPO splits the two signals. If a rollout group's correctness is below threshold α (0.5), the reward is binary correctness. Once the group clears the gate, correct kernels also receive a min-max normalized speedup term with weight β=0.1. Incorrect kernels stay at 0, so a fast wrong kernel cannot buy a positive reward.
Training runs on A100-80GB. A Proposer update hosts a frozen Coder and the Keck execution service, 12 GPUs, about 3.1 hours per 20 steps. A Coder update uses 8 GPUs and about 4.5 hours. The Proposer sees 400 API-list prompts, group size 4, with 5 nested Coder samples per module. The Coder trains on 1,224 validated modules, group size 5.
Evaluation is one-shot KernelBench Level 1 (100 single operators) and Level 2 (100 fusion patterns). The harness is in-house Keck: one candidate per GPU, absolute and relative tolerance 1e-2, no cuBLAS/cuDNN on the CUDA track, and Triton must go through @triton.jit.
| Model | CUDA L1 pass@1 | CUDA L2 pass@1 | Triton L1 pass@1 | Triton L2 pass@1 |
| Claude-4.5-Sonnet | 76.4 | 66.5 | - | - |
| cudaLLM-8B | 72.5 | 67.4 | - | - |
| DeepSeek-V4-Pro | - | - | 49.7 | 63.4 |
| Dr.Kernel-8B (3-turn) | - | - | 34.4 | 73.0 |
| KernelZero-SFT-7B | 70.6 | 59 | 69.4 | 64.9 |
| KernelZero-7B | 75.8 | 69.6 | 77.2 | 72.5 |
On CUDA Level 1, KernelZero sits 0.6 points behind Claude-4.5 on pass@1 and reaches 100% pass@10 against Claude's 97%. On Level 2 it leads 69.6 vs 66.5 pass@1 and 97 vs 93 pass@10. Triton Level 1 pass@1 is 77.2 against DeepSeek-V4-Pro's 49.7 and GLM-5.2's 66.7. Level 2 is 72.5, 0.5 points under Dr.Kernel-8B's three-turn 73.0; Dr.Kernel does not report pass@5/10. CUDA generations cost about 2.8k/4.5k tokens, shorter than Kevin-32B's 9.2k/13.7k.
Speedup splits by backend. fast1@1 is the share of tasks where a single sample is both correct and has speedup above 1× versus PyTorch Eager. Triton Level 1 records 43.9 fast1@1 and 87 fast1@10; CUDA records 17.6 and 29. Level 2 is 65.4/96 for Triton and 2.4/12 for CUDA. Generated CUDA never calls vendor libraries, while the PyTorch reference dispatches conv and GEMM to cuDNN/cuBLAS.
Freezing the Proposer from step 40 leaves CUDA Level 2 pass@1 at 64.4; full co-evolution reaches 69.6. Triton Level 2 moves from 68.3 to 72.5. Level 1 gains are thin: 0.9 points on CUDA. In the α sweep, never turning on the speedup term (α=1.1) gets the highest pass@1 at 74.85% with only 9.75% fast1.5@1; α=0.9 reaches 11.00% on that speed metric. α=0.5 wins best-of-10 speedup on all six profiled operators, including Tensor-MM from 0.12× after SFT to 4.42×.
For operator and short-fusion tasks in KernelBench, a 7B model plus this RL loop already lands in the same band as Claude-4.5 and the specialized 8B models, with shorter rollouts than Kevin. The cold start still consumes cudaLLM's public corpus; the SFT checkpoint already has 70.6 CUDA Level 1 pass@1, and RL adds 5.2 points. Co-evolution pays off mainly on Level 2, not as a from-scratch kernel skill.
It will not replace production conv/GEMM on the CUDA path. Triton is closer to useful, because the language and compiler already absorb tiling and fusion. The proposer-coder loop is a reusable pattern for code generation that needs a moving difficulty frontier. The paper only measures it on KernelBench.
There is no standalone Limitations section.
KernelBench has 250 tasks across three levels; only the 200 Level 1 and Level 2 tasks are scored. Numbers come from Keck, not the original KernelBench scripts: isolated GPU runs, statedict alignment, 1e-2 tolerance. Comparisons with cudaLLM and Kevin are therefore not a strict same-harness ranking.
The abstract's claim of beating Claude-4.5 holds for CUDA pass@5/10 and Level 2 pass@1, not for Level 1 pass@1. Dr.Kernel is three-turn, KernelZero is one-shot; 72.5 vs 73.0 on Triton Level 2 is not a clean ranking. Weak CUDA Level 2 speedup is the no-library constraint hitting a vendor-library baseline. α=1.1 is more accurate; 0.5 trades some correctness for speed, and the paper does not plot a full Pareto. Training still needs 8-12 A100s plus live compile-and-run, not a zero-compute bootstrap.