Turing

Introducing KernelQuest: Can agents write Triton kernels?

A benchmark for measuring whether AI agents can deliver correct, verifiable, expert-level GPU kernel optimization.

Turing Frontier Research Lab
Turing Frontier Research Lab
6 MIN READ07 Sep 2026
RESEARCH

The Turing Frontier Research Lab is launching KernelQuest, a benchmark designed to measure whether AI agents can perform the full job of GPU kernel optimization. KernelQuest contains 100 engineer-authored Triton tasks, each delivered as a live environment where an agent can profile a PyTorch workload, write and revise kernels, build timing probes, measure progress, and submit a final solution.

The benchmark asks a stricter question than whether an agent can beat a stopwatch. Can it produce a genuine Triton kernel that is correct, faster than a serious compiled baseline, and able to survive an adversarial verifier?

The Turing Frontier Research Lab built KernelQuest around a straightforward premise: a benchmark should reward real kernel optimization, not shortcuts that happen to score well. Our mission is to create frontier-calibrated data and benchmarks for the next generation of AI. We build benchmarks and training data that push frontier models forward in the domains where real capability matters. Kernel optimization is one of those domains because small differences at the kernel layer compound across every training step and every served token.

The challenge: speed is easy to claim and hard to verify

Kernel performance is where much of the cost of AI is decided. A kernel that leaves a third of the hardware idle makes training and inference more expensive. The distance between generic compiler output and expert-written code is routinely 2x or wider. Closing that gap requires architecture knowledge and a working method: profile, form a hypothesis, rewrite, measure, discard, and repeat.

The core evaluation problem is not simply code generation. A credible kernel result must satisfy three conditions at once:

  1. The output is numerically correct across meaningful input variation.
  2. The implementation is faster than the best generic baseline available for the workload.
  3. The speedup comes from the submitted kernel, not from reduced precision, cached outputs, graph capture, a compiler fallback, or a weakness in the harness.

Existing benchmarks have exposed why all three are necessary. KernelBench originally emphasized one-shot PyTorch-to-CUDA or Triton transpilation against eager PyTorch on a fixed input distribution. Meta's KernelBench-Verified added a TF32 baseline and hidden distributions, after which the best frontier model's geometric-mean speedup fell from 1.43x to 0.88x. Sakana AI's robust-kbench saw measured speedup fall from 3.13x to 1.49x after contaminated tasks were removed.

The lesson is not that agent-written kernels have no value. It is that a benchmark can reward the wrong behavior if its baseline is soft, its input distribution is narrow, or its stopwatch is trusted without structural and numerical checks.

Why we built KernelQuest

Recent benchmark work has improved the field. KernelQuest builds on those lessons, but uses a different anchor.

Benchmark family

Primary focus

What KernelQuest adds

KernelBench and KernelBench-Verified

PyTorch-to-CUDA or Triton transpilation, with stronger hidden tests in Verified

Five seeded inputs, the faster of eager and generic torch.compile as baseline, and a per-task target demonstrated by an engineer

TritonBench and GEAK-eval

Operator synthesis and execution correctness

Optimization against a live PyTorch workload, spanning operators, fusion chains, complete models, and large transformers

TritonGym

Multi-step agent loops with hand-tuned Triton oracles

End-to-end speedup over a compiled baseline, plus structural, freshness, precision, and degeneracy gates

robust-kbench

Hardening against known harness exploits

Task-specific adversarial hardening and direct evidence of how often apparently fast submissions fail verification

MultiKernelBench

Breadth across hardware targets and programming systems

One pinned NVIDIA L40S environment, traded for timing stability and reproducible per-task certification

AgentKernelArena (AMD, 2026)

Complete agent workflows across 196 HIP-to-HIP, Triton-to-Triton, and PyTorch-to-HIP tasks, with an unseen-configuration protocol

A human-certified performance bar, a continuous reward, and anti-hardcoding checks applied to every submission

FastKernels (Snowflake, 2026)

Production alignment across 46 representative architectures in a serving framework; the strongest agent reaches 0.94x against production baselines

A certified target for each single-GPU task and end-to-end whole-model timing against a warm-timed torch.compile build

SOL-ExecBench

Distance to analytically derived hardware ceilings

A reachable bar that a human Triton solution has already cleared under the same gates

The distinctive choice is human certification. Every KernelQuest target is the speedup an engineer-written Triton solution already reached on the same GPU through the same harness. The target is therefore hard, but reachable by construction. A miss measures distance from demonstrated expert performance rather than distance from an aspirational or theoretical number.

What KernelQuest measures

The 100 tasks cover 21 kernel families and four capability levels. They were selected for demonstrated headroom, diversity of computational structure, a distinct competency at each level, deterministic verification, and no overlap with tasks in public kernel-generation suites.

Level

Tasks

Capability under test

Median certified target

Tasks at or above 2x

L1: Single operator

30

Tiling, masking, coalescing, occupancy

4.14x

22

L2: Fused epilogue

24

Eliminating intermediates and fusing operation chains

2.26x

14

L3: Complete model

18

Profiling a published model and finding the bottleneck

1.40x

3

L4: Large transformer or MoE

28

Full-scale attention, routing, and algorithmic optimization

1.47x

6

The levels are not a simple difficulty ladder. L1 and L2 have more raw headroom, while L3 and L4 require diagnosis inside larger systems. Graph operations are the largest family because irregular memory access leaves room for improvement. L1 through L3 use fp32; the 28 large transformer and MoE tasks use deployment-aligned fp16.

Across all tasks, the median certified target is 1.78x, the mean is 3.47x, and 45 tasks require at least 2x. Nine targets are below 1.2x. Low targets are not automatically easy: whole-model baselines already benefit from graph-level compilation.

Quality is part of the benchmark, not a post-processing step

Each task ships with an instruction, a seeded PyTorch reference, a pinned environment with Triton 3.1.0, PyTorch 2.5.0+cu121 and CUDA 12.1, one NVIDIA L40S, and a one-hour agent budget. The certified solution and author-side verifier remain hidden, but the grading contract is stated up front.

Before performance counts, a submission must clear four gates:

Gate

Requirement

Structural

A genuine @triton.jit kernel performs real work in the graded path

Correctness

Outputs match on five independently seeded inputs across up to five distributions, with dtype, storage freshness, sensitivity, and parameter structure checked

Policy

Reference precision is preserved, with no TF32 or lower-precision casts on fp32 tasks and no CUDA-graph capture

Performance

Speed exceeds the faster of eager PyTorch and generic torch.compile, reaches the certified target, and, for the 97 tasks with a derived physical floor, remains above that floor

Timing is also treated as an experiment. Reference and candidate run in separate subprocesses, rounds alternate to reduce thermal drift, backend settings are re-pinned, every trial gets a fresh input, five warmups are discarded, and the result is the mean of 100 trials. Operator tasks use cold-cache timing, while whole-model tasks use warm timing to match how each is deployed. All 100 certified solutions re-cleared their targets, with a median margin of 7 percent.

This rigor changes the outcome. Of 1,200 agent runs, 565 were rejected before speed could count. More strikingly, 170 had already measured at or above their certified target. A stopwatch-only benchmark would have published those runs as wins.

One task shows why the gates matter

Consider the benchmark's Level 2 attention spotlight. The reference processes an 8 x 32768 x 512 fp32 tensor through self-attention, softplus, a 512-wide bias add, and softmin. The certified target is 1.63x. Profiling is not the hard part: every agent found within ten steps that the attention call consumed 99.8 percent of runtime.

Eleven of twelve runs reached the stopwatch, and all eleven beat the target. Only one passed.

The accepted run built an IEEE fp32 attention kernel and reached 1.70x. Among the rejected runs, seven ran dot products on tensor cores, and an eighth emulated fp32 on half-precision hardware. Two submissions also used stability shortcuts that passed the visible Gaussian examples but failed on hidden distributions. Some submissions triggered more than one failure category. One run replaced attention with a closed-form approximation and reported 27.86x, but broke on another distribution and ran below the hardware's physical floor. One never created a submission file.

The slowest timed result was the only valid result. That is exactly the behavior KernelQuest is built to reveal. Kernel skill was present across the field. The separating capability was respecting the arithmetic contract while optimizing inside it.

Leaderboard

Each system ran every task three times in its vendor scaffold, for 300 runs per system and 1,200 runs overall. These are model-and-scaffold results, not isolated model scores. The chart reports run-weighted pass rates with 95% Wilson intervals.

System

Scaffold

Pass rate, 95% CI

Tasks solved in at least 2 of 3 runs

Geo-mean speedup

Valid coverage

Mean wall-clock

Claude Opus 5

Claude Code

67.0% ±5.3

68

3.10x

93%

50.1 min

GPT-5.6 Sol

Codex

26.0% ±4.9

22

2.32x

71%

15.1 min

Gemini 3.7 Flash

Terminus 2

20.7% ±4.6

20

1.93x

63%

29.4 min

Grok 4.5

GrokBuild

15.3% ±4.1

14

2.19x

70%

37.8 min

Claude Opus 5 leads with a 67.0% run-weighted pass rate, versus 26.0% for the next-best system. It reaches the target at least once on 88 tasks and repeats that success in at least two of three runs on 68 tasks. The corresponding pairs for the other systems are 46 and 22, 37 and 20, and 30 and 14. Claude leads on both task coverage and repeatability, and its 67.0% pass rate still leaves substantial headroom.

Failure modes

Of the 813 non-passing runs, 565 failed a gate and 248 produced accepted work that remained below the target. The model-specific patterns are more useful than a single rank:

System / harness

Characteristic failure pattern

Claude Opus 5 (Claude Code)

44.4% of its failures were correct kernels that missed the bar; 18 runs downcast fp32 arithmetic

GPT-5.6 Sol (Codex)

Its largest mode was an unloadable solution at 25.7%; all 45 precision losses were TF32 opt-ins

Gemini 3.7 Flash (Terminus 2)

98 failures returned stale or constant output, 41.2% of its failed runs

Grok 4.5 (GrokBuild)

49 failures had no Triton kernel in the graded path; 39 captured CUDA graphs and five used another GPU backend

Scale changes the problem. At L1 and L2, losses are mostly about arithmetic, precision, and performance. At L3 and L4, delivery dominates: unloadable artifacts, missing Triton kernels, and timing sidesteps account for 62.2 percent and 65.5 percent of failures. This is actionable evidence. Better kernel agents need not only stronger scheduling ideas, but also verifier-equivalent testing, precision discipline, and a continuously valid artifact on disk.

KernelQuest also emits a continuous reward that records how much of the available target speedup a valid run captured. That signal distinguishes a correct near-miss from an artifact that never loaded, making the same tasks useful for both evaluation and training.

Building the evidence needed to push frontier models forward

The Turing Frontier Research Lab builds benchmarks where real capability matters and where the evaluation remains useful after the first leaderboard is published. KernelQuest fits that mission in two ways. First, its certified human targets make the gap to expert performance directly measurable. Second, its verifier records why an attempt failed, producing evidence that can guide training rather than a flat pass or fail label.

Kernel optimization affects the economics of training and inference, but it also exposes a broader frontier problem: an agent must understand the contract, use tools iteratively, resist invalid shortcuts, and deliver an artifact that survives independent verification. KernelQuest measures that complete loop.

For more details, see the KernelQuest research page. To request KernelQuest task samples or discuss the benchmark with the Turing Frontier Research Lab, contact us.

Turing Frontier Research Lab

Turing Frontier Research Lab builds the benchmarks, training sets, and training recipes that push frontier models forward in the domains where real capability matters. For more about us, visit https://www.turing.com/research.

Contact us

Learn more about our research and datasets.

Request samples