D2K-Bench: Can LLM Agents Turn Expert Designs into Efficient GPU Kernels?

D2K-B ENCH: Can LLM Agents Turn Expert Designs into Efficient GPU Kernels? Daifeng Li1,2 , Huiqiang Jiang2 , Chengruidong Zhang2 , Wei Wu3 Xudong Guo2 , Jianhong Tu2 , Jianwei Zhang2 , Binhang Yuan1 , Dayiheng Liu2

arXiv:2610.03226v1 [cs.LG] 2 Oct 2026

1

HKUST, 2 Alibaba Group, 3 USTC

Abstract GPU kernels generated by large language model (LLM) agents can remain less efficient than expert implementations, but runtime alone does not reveal how the gap relates to design discovery and implementation. We introduce D2K-B ENCH, a diagnostic benchmark of 26 tasks and 85 workloads that measures how effectively agents translate expert design guidance into efficient GPU kernels. The guidance covers L1: high-level algorithmic insights, L2: dataflow design, and L3: low-level optimization tricks, including dependencies among these levels. Pairwise runs with and without guidance share task descriptions, workloads, tools, hardware, and a 350-turn budget. Complementary assessments examine independently proposed designs and the design properties implemented in generated code. Across five models on NVIDIA B200 GPUs, guidance raises correctness over 130 model–task pairs from 93.1% to 98.5% and increases the Performance Score over all 26 tasks from 1.46 to 1.95. For the three frontier models with correct submissions on all 26 tasks in both runs (GPT-6-Astra, Claude-Opus-4.8, and GPT-5.6-Sol), geometric mean speedup increases from 1.69× to 2.49×. Across all five models, the mean combined implementation score increases from 57 to 70 out of 100. These results show the value of expert design guidance while identifying design properties that remain unimplemented.

§ Code:

1

https://github.com/QwenLM/D2K-Bench

Introduction

In modern machine learning systems, GPU kernels usually execute the core computations of machine learning workloads, and their efficiency significantly affects training and inference latency, throughput, and GPU operating cost. Recently, large language model (LLM) agents have offered an effective way to automate GPU kernel development through code generation, compilation, testing, profiling, and iterative optimization [1, 2]. Although these development loops can improve performance, the generated GPU kernels could still be less efficient than expert implementations [3]. Concretely, the LLM-generated kernels and expert implementations could differ in their algorithms, dataflow, and low-level optimization techniques. In this paper, we investigate how effectively LLM agents can translate expert design guidance into efficient GPU kernel implementations. Understanding the performance gap between LLM-generated kernels and expert implementations helps target improvements in automated kernel development. The accurately identified gap can inform whether agents need better support for discovering effective designs, implementing and tuning those designs, or both. Algorithmic and dataflow choices can reduce computation and memory traffic, but their benefits depend on the generated code preserving these savings while using GPU resources effectively. Recognizing a promising design therefore does not ensure that an agent can achieve its performance potential. Useful diagnostics should connect an agent’s design choices, their implementation in code, and measured runtime to identify limitations that warrant further investigation. Without such evidence, additional tuning may leave an inefficient algorithm or dataflow unchanged, while further design guidance may fail to resolve implementation limitations.

1

However, assessing how effectively agents translate expert design guidance into efficient kernels is difficult since the execution time alone does not clearly reveal which design choices the generated code implements or how well it implements them. An agent must identify an effective algorithm and dataflow, then implement those choices through correct code, GPU execution mechanisms, and tuning. For example, FlashAttention combines online normalization with tiled computation to avoid storing the full attention matrix [4]. Discovering this design reveals an opportunity to reduce memory traffic; implementing the design further requires coordinated computation and data movement, careful GPU resource management, and tile-size tuning. Runtime alone cannot reveal where an agent falls short. Assessments of independently proposed designs and executable source provide complementary evidence about design discovery and implementation. Existing benchmarks evaluate kernel correctness and performance in different ways: KernelBench evaluates speedup over PyTorch, SOL-ExecBench uses analytical hardware-performance estimates, and FastKernels evaluates generated kernels against production implementations [1, 3, 5]. Prior work also studies optimization guidance and generated code: KernelBench examines hardware knowledge and optimization examples, Hari et al. [2] use Speed-of-Light guidance and LLM-based integrity review, and Shui et al. [6] compare expertassisted and autonomous kernel development. On the other hand, we focus on a diagnostic approach that can connect task-specific design guidance, pairwise comparisons with and without guidance, and structured assessments of independently proposed designs and generated code. This combination makes the assessed design properties explicit and relates their implementation to correctness and runtime. We introduce D2K-B ENCH, a benchmark that connects performance comparisons with assessments of design discovery and implementation. Expert design guidance and pairwise comparisons with and without guidance make both guidance benefits and unimplemented design properties observable. We make the following concrete contributions: Contribution 1. We construct a targeted diagnostic benchmark of 26 tasks and 85 workloads, pairing task

descriptions with selected expert references and task-specific design guidance (Section 2). The guidance covers L1: high-level algorithmic insights, L2: dataflow design, and L3: low-level optimization tricks, and records dependencies among these levels. This structure captures algorithmic transformations, data representation and movement, and GPU execution mechanisms within a common framework for guidance and assessment. Contribution 2. We develop an evaluation protocol that compares kernel performance and design implementa-

tion with and without expert design guidance (Section 3). The runs in each pair use the same task description, workloads, Triton implementation language, tools, and hardware; guidance provides explanations without expert source code or prescribed tuning parameters. An LLM judge assesses independently proposed designs and executable source evidence, with runtime hidden for source assessment and credit for correct alternative designs. Source comparisons use all 26 tasks in both runs, providing a diagnostic of implemented properties on these tasks and the shortfall that remains under guidance. Contribution 3. We characterize the benefits and limits of expert design guidance across five models on NVIDIA

B200 GPUs (Section 4). Guidance raises correctness over 130 model–task pairs from 93.1% to 98.5%. The Performance Score is computed over all 26 tasks for each model and increases from 1.46 to 1.95. For GPT-6-Astra, Claude-Opus-4.8, and GPT-5.6-Sol, which have correct submissions on all 26 tasks in both runs, geometric mean speedup increases from 1.69× to 2.49×. Across all five models, the mean combined implementation score increases from 57 to 70 out of 100, while guided code still leaves assessed properties unimplemented. Additional analyses examine cumulative guidance in a fixed level order and an illustrative implementation. The evaluations separately assess agents’ independently proposed designs and their ability to implement supplied expert designs in executable code, and relate these assessments to kernel performance.

2

Benchmark Design and Construction

We introduce how we select the benchmark tasks (Section 2.1), construct the inputs provided to agents (Section 2.2), and derive expert design guidance for each task (Section 2.3).

2

2.1

Benchmark Task Selection

D2K-B ENCH contains 26 GPU kernel optimization tasks. Each task includes at least two different configurations; we use the term workload to refer to one task with a particular configuration. The benchmark includes 85 workloads in total. The tasks and workloads are carefully selected by experts from systems and specialized implementations such as vLLM [7], SGLang [8], and FlashAttention [4]. We select tasks whose expert-written implementations reduce computation or memory traffic through algorithmic transformations or dataflow changes, with low-level kernel optimizations providing further performance gains. The resulting suite spans common kernel development scenarios in LLM training and inference, including attention, mixture of experts, and quantization. Appendix Table 3 summarizes the tasks by operator category along with the representative expert design ideas in each group. For each task, we include one expert implementation over the task’s corresponding workloads, i.e., the expert implementation is used across all workloads for one task. We validate the outputs and track the state updates against the PyTorch reference under the task’s numerical tolerances and verify that execution invokes the selected B200 expert kernels.

2.2

Task Specification

For each task, we construct a set of shared task files that specify what the agent must implement and how the task will be evaluated. These files include a task description, workload definitions, an input generator, a PyTorch reference, implementation requirements, and a solution.py template with an unimplemented run(...) entry point. The task description defines the required outputs, numerical tolerances, and state updates; workload definitions specify tensor shapes, data types, and layouts. We derive the task description and workloads from APIs, documentation, tests, and representative model settings. The implementation requirements mandate custom kernels for the target operator and prohibit calls to the PyTorch reference or existing implementations of that operator. Example task of Muon orthogonalization kernel. We enumerate one concrete example, i.e., the Muon orthogonal-

ization that isolates the matrix orthogonalization used by the optimizer, transforming an input gradient matrix G into an update direction. Figure 1(a) provides an illustrative summary of these shared  task filesfor the Muon task. Each matrix is normalized in FP32 and then cast to FP16 to initialize X0 = FP16 (∥G∥GF +ϵ) . The kernel then conducts five polynomial Newton-Schulz updates with FP16 iterates. For a wide matrix, each update step, 2 indexed by k=0, . . . , 4, forms the Gram matrix Rk = Xk X⊤ k and computes Xk+1 = (ak I + bk Rk + ck Rk )Xk , where I is the identity matrix and ak , bk , ck are the supplied coefficients. The normalization uses the Frobenius norm ∥G∥F and a small positive constant ϵ. One workload contains 32 BF16 matrices of shape 4096×11008: each rectangular iterate has this shape, whereas its Gram matrix is 4096×4096. The direct PyTorch computation rebuilds the Gram matrix and writes a new rectangular iterate at every step, repeatedly processing the larger matrix.

2.3

Expert Design Guidance

Expert design guidance connects algorithmic choices to their implementation. In the guided configuration, this guidance is made available to the agent alongside the shared task files defined in Section 2.2. Expert code supplies the design choices; corresponding technical documentation and comments clarify their intent. We organize the guidance into three levels: • L1: High-level algorithmic insights. Such guidance states the algorithmic transformation, why the proposed algorithmic insight would preserve the operator’s required behavior, and the computation or data movement it could optimize. • L2: Dataflow design. The guidance specifies where variables should be produced, stored, and consumed, together with their layouts and lifetimes. The guidance also identifies intermediate variables that can be reused or avoided. • L3: Low-level optimization tricks. The guidance explains how execution mechanisms implement the design, such as optimizing pipelining to overlap data movement with computation and warp specialization to 3

(a) Shared Task Files

(b) Expert Design Guidance

Task Description

L1 High-Level Algorithmic Insights

Compose updates on smaller square matrices R and Q instead of updating rectangular X after every step. R = XXᵀ ,P(R) = a I + bR + cR² X′ = P(R)X R′ = P(R)RP(R)ᵀ

Normalize each matrix in FP32; run five FP16 polynomial updates; return the input shape and dtype.

Workloads and Input Generator

batch=32; M=4096; N=11008; steps=5 G = torch.randn(batch, M, N, dtype=torch.bfloat16, device="cuda") return [G, steps]

Apply Q only at restart and final boundaries.

L2 Dataflow Design

Keep rectangular X fixed within each segment while smaller square states R and Q evolve. Build R Evolve R, Q Update X At restart, rebuild R from X and reset Q.

PyTorch Reference

X = normalize_fp32(G).half() for a, b, c in coefficients: # 5 steps R = X @ X.transpose(-2, -1) X = a*X + (b*R + c*(R@R)) @ X return X.to(G.dtype)

L3 Low-Level Optimization Tricks

Use square-state kernels within each segment. Apply the rectangular product only at restart and final boundaries. Compute one triangle of each symmetric product. Reuse each accumulator for both stores; write each diagonal tile once.

Solution Template (Agent Edit)

def run(G, steps=5): # FP32 normalization; five FP16 polynomial updates # implemented as custom Triton kernels on B200. raise NotImplementedError

Symmetric Products

Implementation Requirements

Use custom Triton kernels on NVIDIA B200. No result reuse or timed computation in PyTorch.

Triangular Computation

Figure 1 Example of the Muon orthogonalization task: (a) shared task files provided in both optimization runs, including the task description, workload definitions, an input generator, a PyTorch reference, a solution template, and implementation requirements; (b) expert design guidance additionally provided in the guided run at L1, L2, and L3. Expert implementation code is withheld.

assign different roles to warps. The guidance also specifies the required ordering and synchronization while leaving concrete tuning parameters unspecified. The guidance records dependencies among the three levels to explain how the design choices work together. These dependencies connect an algorithmic transformation to the data representation and movement it requires, and then to the GPU execution mechanisms that preserve its benefits. They specify the cross-level dependencies among algorithmic, dataflow, and low-level execution choices, including the required ordering and state transitions. Example expert guidance for Muon orthogonalization kernel. We provide the corresponding example of the expert

guidance for the Muon orthogonalization kernel. The Muon guidance replaces repeated updates of the large rectangular matrix with composition on smaller square matrices (Figure 1(b)). At L1, the guidance defines Pk = ak I + bk Rk + ck R2k . In exact arithmetic, since Xk+1 = Pk Xk , steps 0 through j can be composed into Q = Pj · · · P0 . The resulting rectangular matrix is then computed as Xj+1 = QX0 in a single matrix multiplication, avoiding explicit computation of the intermediate rectangular matrices X1 , . . . , Xj . Meanwhile, ⊤ the Gram matrix is updated directly as Rk+1 = Pk Rk P⊤ k , rather than recomputed as Xk+1 Xk+1 . This algorithmic transformation preserves the result in exact arithmetic while replacing repeated rectangularmatrix updates with composition on smaller square matrices. At L2, the dataflow keeps rectangular X fixed throughout each composition segment, while square states R and Q are updated and reused across consecutive polynomial steps. The updated rectangular matrix is materialized only at the segment boundary by applying Q to X. Before step k=2, the fixed restart applies the composition of steps 0 and 1 to X, rounds to FP16, rebuilds R = XX⊤ to limit numerical drift, and initializes a fresh Q = a2 I + b2 R + c2 R2 for steps 2 through 4. At L3, the kernel execution schedule respects these data dependencies: square-state matrix-product kernels execute within each segment, and rectangular products apply Q at restart and final boundaries after the corresponding 4

square-state updates are complete. To avoid redundant arithmetic, symmetric-product kernels compute one triangle and reuse each tile’s accumulator for the original and transposed stores, writing diagonal tiles only once. For each task, we prepare two LLM-judge prompts: one assesses an independently generated written design, and the other assesses generated implementation code. The written design is elicited in a separate agent run described in Section 3.3. Both judge prompts combine the task description, task-specific design criteria, and shared assessment instructions. These criteria form the design specification, describing the expert’s design choices, the costs they remove, and the requirements for combining them correctly. The design-assessment prompt includes the agent’s written design; the implementation-assessment prompt includes the three levels of guidance, the expert implementation, and the candidate source code. The scoring procedure is described in Section 3.3.

3

Evaluation Protocol

We describe the agentic evaluation pipeline with and without expert design guidance (Section 3.1), define the measured performance score Sperf (Section 3.2), and define the LLM judge scores for independent designs, JSdesign , and generated implementations, JSimpl (Section 3.3).

3.1

Agentic Evaluation Pipeline

For each model and task, we conduct paired agentic optimization runs with and without expert design guidance. Each run follows the same pipeline: task initialization, iterative kernel optimization and evaluation, and final submission selection. Figure 2 illustrates the agent loop and its interaction with the evaluator sandbox. Task initialization. Each agentic execution starts with an agent workspace containing the shared task files

defined in Section 2.2. The guided run additionally receives the expert design guidance from Section 2.3, covering all three levels, together with instructions. This additional input is represented by the dashed branch in Figure 2. Expert source code is unavailable to the agent in both runs. The two runs otherwise use the same implementation language, tools, hardware, and optimization budget. Iterative kernel optimization and evaluation. The agent develops candidate implementations by editing the

supplied solution.py template shown in Figure 1(a). The template specifies the task-specific signature of run(...) and leaves the function body unimplemented; the agent implements this function to produce the required outputs and state updates. The agent interacts with the execution sandbox through Model Context Protocol (MCP) tools: read and write support file inspection and implementation editing, bash executes compilation and profiling commands, and submit invokes the evaluator. On submission, the evaluator loads the candidate file and calls run(...) with inputs prepared for each workload according to the task description. The agent uses the returned correctness and timing feedback, together with profiling tools such as NVIDIA Nsight Compute (ncu), to diagnose bottlenecks and revise its implementation. This process forms an iterative loop of implementation, compilation, testing, profiling, and revision. Final submission selection. Both runs use the same fixed turn limit and selection procedure. One agent

turn consists of a model response requesting one or more tool calls, execution of those calls, and return of their results. A candidate is eligible for final selection only if the candidate passes the correctness and implementation checks and has valid latency measurements on every workload of the task. For each eligible candidate, we compute the geometric mean of its speedups across the task’s workloads, as defined in Section 3.2. Among eligible candidates produced within the turn limit, we select the one with the highest geometric mean speedup as the final submission.

3.2

Objective Evaluation Metrics

We include an objective evaluation module (shown in Figure 2) to check correctness and implementation validity, measure latency, and compute performance scores for submitted implementations.

5

Agent Loop

MCP Tools

With or without Design Guidance

Expert Design Guidance L1 High-Level Algorithmic Insights

Profile Find Bottlenecks

read/write Files

Objective Evaluation

bash Compile / Profile

Validity Correctness

submit Evaluate Test

Revise

Timing

Feedback

L2 Dataflow Design L3 Low-Level Optimization Tricks

Evaluator Sandbox

LLM Judge Independent Design Proposal

Implement Kernel

Generated Code

Qwen3.8-Max-0902 L1 / L2 / L3 Criteria

JSdesign

JSimpl

Figure 2 Evaluation workflow. The dashed branch supplies L1–L3 guidance in the guided run. MCP calls and feedback use separate arrows. The LLM judge assigns an implementation score at each level; the arithmetic mean of these scores is JSimpl . The evaluator measures Sperf . Correctness and implementation validity. A task is scored as correct only if its submission satisfies both the

correctness criteria and implementation requirements on every workload. Violations of the implementation requirements count as incorrect even when numerical outputs are correct. Correctness checks cover numerical outputs, updates to mutable inputs, repeated execution on new inputs, and, where applicable, distributed behavior. We check compliance with the custom-kernel requirements in Section 2.2 using abstract syntax tree (AST) analysis, static keyword scans, and inspection of libraries linked during compilation. For distributed tasks, we inspect communication traffic to verify that the required communication occurs. Manual review and unit tests check compliance with the full task description, including all required outputs and state updates. Execution latency timing. We measure GPU execution time for one call to the candidate run(...) entry point,

including its kernels and required communication. Input construction, compilation, and evaluator setup are excluded from the timed region. The evaluator controls input preparation and state restoration to provide equivalent inputs to all implementations. For distributed tasks, the reported latency is the maximum across participating ranks. All implementations of a given workload use the same timing procedure. For tasks with very short kernel execution times, such as decoding, we capture repeated calls to run(...) in a CUDA Graph that takes at least one second to execute and replay the graph five times. For other tasks, we use CUDA events to time batches of repeated calls to run(...), with each batch lasting at least one second. In both modes, the reported latency is the arithmetic mean of repeated per-call measurements. Performance scoring. We evaluate correctness on all 26 tasks for each model and input configuration. Per-

formance is measured relative to the first-correct baseline: Fable-5’s first correct submission in Claude Code without design guidance. Let Wi denote the workloads of task i, tb,i,w the baseline latency, and tm,c,i,w the latency of the final submission from model m under input configuration c (guided or unguided). For a task to pass evaluation, the final submission must pass the correctness and implementation checks and have valid latency measurements on every workload. The task performance score is defined as:    tb,i,w GM , if checks and timing pass, w∈Wi tm,c,i,w Sm,c,i (1) perf =  0.5, otherwise. The model performance score over the task set T is   m,c,i Sm,c , perf = GMi∈T Sperf 6

|T | = 26.

(2)

Here, GM denotes the geometric mean. Averaging first within tasks and then across tasks gives each task equal weight. We use 0.5 as a fixed multiplicative penalty. In log space, a task assigned this penalty contributes the same loss as the gain from an equally weighted task with 2× performance. Overall Sperf is the geometric mean of the five model scores for a given configuration. We omit indices when the model and configuration are clear from context. The PyTorch reference and the expert implementation selected in Section 2.1 are scored by substituting their latencies for tm,c,i,w in Equation 1 and then applying Equation 2; the baseline has Sperf = 1.0 by definition. The gain from guidance is the ratio of guided to unguided Sperf . To express performance as a percentage of expert performance, we replace baseline latency in the numerator of Equation 1 with expert latency and multiply the aggregate by 100. Following the SOL score definition in SOL-ExecBench [5], we observe a limitation when applying the score to D2K-B ENCH. The score uses an analytical estimate tSOL,i,w of the minimum execution time of the PyTorch reference computation, derived from hardware capabilities. For an evaluated implementation with latency ti,w , we compute a score for each workload: si,w =

tb,i,w − tSOL,i,w . ti,w + tb,i,w − 2tSOL,i,w

(3)

The SOL score’s intended interpretation within [0, 1] requires tb,i,w > tSOL,i,w and ti,w ≥ tSOL,i,w . However, D2K-B ENCH permits algorithmic and dataflow changes that can alter the computation and memory traffic assumed by the bound, so the SOL bound need not remain valid for a generated implementation. For Muon, transformed implementations can execute below the SOL bound derived from the PyTorch reference. The resulting scores for individual workloads exceed one in some cases and, when the denominator changes sign, become negative; we observe negative values for GPT-6-Astra, GPT-5.6-Sol, and Claude-Opus-4.8. Thus, we do not report the SOL score.

3.3

LLM Judge Scores for Designs and Implementations

Correctness and runtime alone are not able to completely reveal which design choices an agent implements or can propose independently. We therefore use Q WEN 3.8-M AX-0902 as an LLM judge (Figure 2) to assess generated code and independent design proposals against the expert-derived design specification in Section 2.3. Implementation assessment. For each task and input configuration, the judge assesses the generated code at L1 (algorithmic insights), L2 (dataflow design), and L3 (low-level optimization), as defined in Section 2.3. Each level is scored from 0 to 100 using five anchors: absent (0), isolated fragments (25), partial implementation (50), substantial implementation (75), and complete implementation (100). Intermediate scores are permitted, and every score must be supported by exact code passages demonstrating the implemented design choices. Runtime measurements are hidden from the judge. Alternative implementations receive credit when they satisfy the assessed design intent, even if they differ from the expert code. Functional correctness is evaluated separately in Section 3.2. The implementation judge score JSm,c,i impl ∈ [0, 100] is the arithmetic mean of the three level scores for model m, configuration c, and task i. To include all 26 tasks in implementation comparisons, we assign JSm,c,i impl = 0 to an incorrect final submission and set its three level scores to zero. This adjustment uses the separate correctness evaluation, not a runtime-based judgment. Independent design assessment. To assess design formulation, we elicit a written proposal in a separate run without expert guidance or code generation, using the same model configuration and maximum turn budget as kernel generation. The judge evaluates the proposal against task-specific design criteria at L1, L2, and L3. The design judge score JSm,i design ∈ [0, 100] is the arithmetic mean of the three level scores. It has no configuration index because the proposal is elicited only without guidance. Reported metrics. Model judge scores are arithmetic means over all 26 tasks:

JSm design =

1 X m,i JSdesign , |T |

JSm,c impl =

i∈T

1 X m,c,i JSimpl . |T |

(4)

i∈T

We aggregate each level in the same way and report implementation scores separately for guided and unguided runs. Overall judge scores are arithmetic means across the five models, whereas overall Sperf is a geometric 7

mean. We omit indices when context suffices. Section 4.2 tests whether these runtime-blind assessments align with measured kernel performance. The scores describe design quality and implementation completeness; they do not directly estimate speedup.

4

Experiments

Our experiments assess the benefits and limits of expert design guidance across five models and 26 kernel tasks. We use a budget of 350 agent turns per task (see Appendix A.2 for the experimental setup). We aim to answer three key research questions: RQ1. How does the comprehensive expert design guidance at all three levels affect correctness and end-to-end kernel performance? (Section 4.1); RQ2. How closely do LLM judge scores align with measured end-to-end kernel performance? (Section 4.2); and RQ3. How does guidance at L1, L2, and L3 affect kernel performance? (Section 4.3).

4.1

RQ1: Correctness and Performance Gains with Expert Guidance

Comprehensive guidance improves both correctness and Sperf in Table 1. Across 130 model–task pairs, correctness increases from 93.1% to 98.5%. Overall Sperf increases from 1.46 to 1.95, a 33.9% improvement. All five models improve their Sperf under guidance. The expert implementation achieves Sperf = 2.33 on the same 26 tasks. GPT-6-Astra, Claude-Opus-4.8, and GPT-5.6-Sol are correct on every task in both configurations, so their Sperf is the geometric mean speedup over all 26 tasks. Their geometric mean speedup increases from 1.69× to 2.49×. For GLM-5.3 and Kimi-K3, tasks without a correct submission or valid timing contribute 0.5. Appendix A.3 reports the individual speedups of the three models that are correct on all 26 tasks. Summarized Answer to RQ1. Comprehensive guidance improves both correctness and performance across all

five models: correctness rises from 93.1% to 98.5%, and Sperf over all 26 tasks rises from 1.46 to 1.95. For the three models that are correct on every task, geometric mean speedup increases from 1.69× to 2.49×.

4.2

RQ2: Alignment of LLM Judge Scores with Measured Performance

We assess whether the LLM judge scores in Section 3.3 align with measured end-to-end kernel performance despite the judge having no access to runtime measurements. Table 2 compares JSdesign for independently proposed designs, JSimpl for final implementations, and the measured performance score Sperf . Each value aggregates the same 26 tasks for a model; implementation and performance scores are reported separately for each guidance configuration. Our key findings are: (i) Implementation scores align with performance and guidance gains. In both configurations, JSimpl follows the broad model ordering of Sperf . Guidance also increases both scores for every model. Thus, despite having no runtime access, the judge captures both performance differences across models and the direction of improvement under guidance. (ii) Independent design scores also reflect broad performance differences. Higher JSdesign is generally associated with stronger unguided performance. Appendix C.3 presents a concrete Muon case illustrating the design and implementation assessments. Summarized Answer to RQ2. Runtime-blind implementation scores align with broad performance differences

and guidance gains. Independent design scores also reflect broad performance differences. This alignment is approximate and demonstrated across model aggregates.

4.3

RQ3: Effects of Design Guidance Levels

We evaluate four fixed input configurations in separate optimization runs: no guidance, L1, L1+L2, and L1+L2+L3. The guidance supplied in each run remains fixed. We select GPT-6-Astra, GPT-5.6-Sol, and Claude-Opus-4.8 for cumulative ablations because they produce correct implementations on all 26 tasks; all three also have correct submissions and valid timing on every workload under every ablation configuration. These cumulative configurations preserve dependencies among design decisions. Figure 3 compares their expert-relative performance using the same timing procedure and expert reference implementations as Table 1.

8

Table 1 Correctness and performance of final submissions on 26 tasks. Correct counts are out of 26 tasks for each model and 130 model-task pairs overall. Overall Sperf weights models equally. PyTorch

Baseline

Model

Sperf

Fable-5

Correct

w/o Design Guidance Sperf

Correct

w/ Design Guidance Sperf

Expert Sperf

GPT-6-Astra Claude-Opus-4.8 GPT-5.6-Sol GLM-5.3 Kimi-K3

0.056 0.056 0.056 0.056 0.056

1.00 1.00 1.00 1.00 1.00

26 26 26 21 22

2.68 1.38 1.31 1.18 1.16

26 26 26 25 25

3.33 2.31 2.01 1.34 1.38

2.33 2.33 2.33 2.33 2.33

Overall

0.056

1.00

121

1.46

128

1.95

2.33

Table 2 Alignment of LLM judge scores with measured kernel performance. All values aggregate 26 tasks per model. Averages are computed before rounding; JSdesign and JSimpl are rounded to integers. Proposal Model

w/o Design Guidance

w/ Design Guidance

JSdesign

L1

L2

L3

JSimpl

Sperf

L1

L2

L3

JSimpl

Sperf

GPT-6-Astra Claude-Opus-4.8 GPT-5.6-Sol GLM-5.3 Kimi-K3

73 67 72 59 63

68.38 65.04 63.50 43.42 48.73

68.81 64.04 63.58 45.08 48.27

66.96 60.46 58.88 46.42 48.77

68 63 62 45 49

2.68 1.38 1.31 1.18 1.16

80.38 72.08 72.50 67.65 66.50

79.58 68.65 72.88 66.00 64.69

77.38 67.96 65.69 61.12 61.69

79 70 70 65 64

3.33 2.31 2.01 1.34 1.38

Average

67

57.82

57.95

56.30

57

1.46

71.82

70.36

66.77

70

1.95

Values use Equations 1 and 2 with T equal to all 26 tasks and expert latency in the numerator, and are multiplied by 100.

Summarized Answer to RQ3. Averaged across the three mod-

els, adding L1 yields the largest incremental gain, while L2 and L3 provide additional improvements.

5

w/o design guidance

Expert-relative performance (%)

Across all three models, performance increases at every step of the guidance sequence. For GPT-6-Astra, most of the total gain comes from adding L1 guidance, while GPT-5.6-Sol and Claude-Opus-4.8 gain more from adding L3 than from adding L2. This pattern highlights the value of guidance about GPU execution mechanisms even after algorithm and dataflow choices have been supplied. Together, the three models show how guidance at each level supports the translation of expert designs into efficient implementations. We also observe that guidance increases token usage, but the guided optimization trajectory reaches a higher final retained score; the tokenefficiency results are reported in Appendix A.4.

L1

L1+L2

L1+L2+L3

150

131

140 143

115

100

99

Expert = 100%

56

69

74

86

82 86 59

50

0

GPT-5.6Sol

Claude-Opus4.8

GPT-6Astra

Figure 3 Expert-relative performance under cumulative guidance. Results are available for all three models on all 26 tasks under each input configuration.

Related Work

GPU kernel generation benchmarks. GPU kernel generation benchmarks vary in task and hardware coverage

and in the reliability of their evaluation [9]. KernelBench [1] established the task of producing a correct GPU kernel that runs faster than the PyTorch reference; TritonBench [10] focuses on Triton operators, MultiKernelBench [11] spans GPUs, NPUs, and TPUs, and DataKernelBench [12] extends the domain to database queries. Evaluation methods also differ. Robust-kbench [13] and KernelBench-Verified [14] strengthen verification and hidden testing, while KernelBenchX [15] and AgentKernelArena [16] analyze correctness, hardware efficiency, and agent generalization. SOL-ExecBench [5] evaluates kernel performance against analytically derived

9

hardware bounds [17]. FastKernels [3] emphasizes production interfaces, and RealisticTritonBench [18] uses framework pull requests and end-to-end tests. FlashInfer-Bench [19] studies serving workloads, while Atrex-Bench [20] samples operators and shapes from production inference traces. Kernel optimization by LLM agents. Agentic systems optimize kernels through iterative feedback, reusable

knowledge, and hardware guidance. AutoKernel [21] and TritonForge [22] organize compilation, testing, profiling, and revision, while Astra [23] and CUDAMaster [24] add multi-agent or hardware-aware workflows. KernelSkill [25] provides reusable optimization knowledge. AccelOpt [26] and DRTriton [27] learn from prior implementations or synthetic data, while Xe-Forge [28] stages search under hardware constraints. KernelFoundry [29] and QiMeng-Kernel [30] use evolutionary search or hierarchical planning. Hari et al. [2] allocate search effort using a compact domain-specific language and Speed-of-Light guidance, and Shui et al. [6] show that expert directions can improve final implementations in a B200 contest harness. A complementary line of work develops programming systems that make computation, data movement, and scheduling choices explicit, including Halide [31], Fireiron [32], Triton [33], and TileLang [34].

6

Conclusion

We propose D2K-B ENCH, a new benchmark that measures how effectively LLM agents translate expert design guidance into efficient GPU kernels. The benchmark pairs 26 tasks and 85 workloads with guidance on high-level algorithmic insights, dataflow design, and low-level optimization tricks. Its evaluation combines pairwise performance comparisons with assessments of independently proposed designs and generated code. Across five models on NVIDIA B200 GPUs, guidance raises correctness over 130 model–task pairs from 93.1% to 98.5%. The Performance Score is computed over all 26 tasks for each model and increases from 1.46 to 1.95. Guidance also raises the mean combined implementation score from 57 to 70. Guidance improves performance, but agents still leave assessed design properties unimplemented.

References [1] Anne Ouyang, Simon Guo, Simran Arora, Alex L. Zhang, William Hu, Christopher Ré, and Azalia Mirhoseini. KernelBench: Can LLMs write efficient GPU kernels? In Proceedings of the 42nd International Conference on Machine Learning, volume 267 of Proceedings of Machine Learning Research, pages 47356–47415. PMLR, 2025. URL https://proceedings.mlr.press/v267/ouyang25a.html. [2] Siva Kumar Sastry Hari, Vignesh Balaji, Sana Damani, Qijing Huang, and Christos Kozyrakis. Improving efficiency of GPU kernel optimization agents using a domain-specific language and speed-of-light guidance. arXiv preprint arXiv:2603.29010, 2026. URL https://arxiv.org/abs/2603.29010. [3] Gabriele Oliaro, Yichao Fu, May Jiang, Owen Lu, Junli Wang, Zhihao Jia, Hao Zhang, and Samyam Rajbhandari. FastKernels: Benchmarking GPU kernel generation in production. arXiv preprint arXiv:2605.23215, 2026. URL https://arxiv.org/abs/2605.23215. [4] Tri Dao, Daniel Y. Fu, Stefano Ermon, Atri Rudra, and Christopher Ré. FlashAttention: Fast and memory-efficient exact attention with IO-awareness. In Advances in Neural Information Processing Systems, volume 35, pages 16344–16359, 2022. doi: 10.52202/068431-1189. URL https://proceedings.neurips.cc/paper_files/paper/ 2022/hash/67d57c32e20fd0a7a302cb81d36e40d5-Abstract-Conference.html. [5] Edward Lin, Sahil Modi, Siva Kumar Sastry Hari, Qijing Huang, Zhifan Ye, Nestor Qin, Fengzhe Zhou, Yuan Zhang, Jingquan Wang, Sana Damani, Dheeraj Peri, Ouye Xie, Aditya Kane, Moshe Maor, Michael Behar, Triston Cao, Rishabh Mehta, Vartika Singh, Vikram Sharma Mailthody, Terry Chen, Zihao Ye, Hanfeng Chen, Tianqi Chen, Vinod Grover, Wei Chen, Wei Liu, Eric Chung, Luis Ceze, Roger Bringmann, Cyril Zeller, Michael Lightstone, Christos Kozyrakis, and Humphrey Shi. SOL-ExecBench: Speed-of-light benchmarking for real-world GPU kernels against hardware limits. arXiv preprint arXiv:2603.19173, 2026. URL https://arxiv.org/abs/2603.19173. [6] Yue Shui, Chenyu Ma, Hangfei Xu, Shengzhao Wen, and Yanpeng Wang. Harness engineering for LLM-driven GPU kernel generation. arXiv preprint arXiv:2607.17979, 2026. URL https://arxiv.org/abs/2607.17979. [7] Woosuk Kwon, Zhuohan Li, Siyuan Zhuang, Ying Sheng, Lianmin Zheng, Cody Hao Yu, Joseph E. Gonzalez, Hao Zhang, and Ion Stoica. Efficient memory management for large language model serving with PagedAttention. In

10

Proceedings of the 29th Symposium on Operating Systems Principles, pages 611–626, 2023. doi: 10.1145/3600006. 3613165. URL https://doi.org/10.1145/3600006.3613165. [8] Lianmin Zheng, Liangsheng Yin, Zhiqiang Xie, Chuyue Sun, Jeff Huang, Cody Hao Yu, Shiyi Cao, Christos Kozyrakis, Ion Stoica, Joseph E. Gonzalez, Clark Barrett, and Ying Sheng. SGLang: Efficient execution of structured language model programs. In Advances in Neural Information Processing Systems, volume 37, pages 62557–62583, 2024. doi: 10.52202/079017-2000. URL https://proceedings.neurips.cc/paper_files/paper/2024/hash/ 724be4472168f31ba1c9ac630f15dec8-Abstract-Conference.html. [9] Yang Yu, Peiyu Zang, Chi Hsu Tsai, Haiming Wu, Yixin Shen, Jialing Zhang, Haoyu Wang, Zhiyou Xiao, Jingze Shi, Yuyu Luo, Wentao Zhang, Chunlei Men, Guang Liu, and Yonghua Lin. Towards automated kernel generation in the era of LLMs. arXiv preprint arXiv:2601.15727, 2026. URL https://arxiv.org/abs/2601.15727. [10] 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. In Findings of the Association for Computational Linguistics: ACL 2025, pages 23053– 23066, Vienna, Austria, 2025. Association for Computational Linguistics. doi: 10.18653/v1/2025.findings-acl.1183. URL https://aclanthology.org/2025.findings-acl.1183/. [11] Zhongzhen Wen, Yinghui Zhang, Zhong Li, Zhongxin Liu, Linna Xie, and Tian Zhang. MultiKernelBench: A multiplatform benchmark for kernel generation. arXiv preprint arXiv:2507.17773, 2025. URL https://arxiv.org/abs/ 2507.17773. [12] Gokul Karthik Kumar, Yotam Perlitz, Corey Lammie, Andrea Giovannini, and Katja Hose. DataKernelBench: Can LLMs optimize database queries on GPUs? arXiv preprint arXiv:2608.25061, 2026. URL https://arxiv.org/abs/ 2608.25061. [13] Robert Tjarko Lange, Qi Sun, Aaditya Prasad, Maxence Faldor, Yujin Tang, and David Ha. Towards robust agentic CUDA kernel benchmarking, verification, and optimization. arXiv preprint arXiv:2509.14279, 2025. URL https: //arxiv.org/abs/2509.14279. [14] Yunxiang Zhang, Ping Yu, Jianyu Wang, Max (Xiangjun) Fan, Julian Reed, Azalia Mirhoseini, and Will Su. KernelBench-Verified: Do LLM-generated kernels actually beat PyTorch? arXiv preprint arXiv:2607.16241, 2026. URL https://arxiv.org/abs/2607.16241. [15] Han Wang, Jintao Zhang, Kai Jiang, Haoxu Wang, Jianfei Chen, and Jun Zhu. KernelBenchX: A comprehensive benchmark for evaluating LLM-generated GPU kernels. arXiv preprint arXiv:2605.04956, 2026. URL https: //arxiv.org/abs/2605.04956. [16] Sharareh Younesian, Wenwen Ouyang, Sina Rafati, Mehdi Rezagholizadeh, Sharon Zhou, Ji Liu, Yue Liu, Yuchen Yang, Hao Li, Ziqiong Liu, Dong Li, Vikram Appia, Zhenyu Gu, and Emad Barsoum. AgentKernelArena: Generalizationaware benchmarking of GPU kernel optimization agents. arXiv preprint arXiv:2605.16819, 2026. URL https: //arxiv.org/abs/2605.16819. [17] Samuel Williams, Andrew Waterman, and David Patterson. Roofline: An insightful visual performance model for multicore architectures. Communications of the ACM, 52(4):65–76, 2009. doi: 10.1145/1498765.1498785. URL https://doi.org/10.1145/1498765.1498785. [18] Jinjun Huang, Zhongzhen Wen, Tongtong Xu, Meng Yan, Xin Xia, and Zhongxin Liu. RealisticTritonBench: A benchmark for Triton-kernel generation in real-world AI frameworks. arXiv preprint arXiv:2608.12004, 2026. URL https://arxiv.org/abs/2608.12004. [19] Shanli Xing, Yiyan Zhai, Alexander Jiang, Yixin Dong, Yong Wu, Zihao Ye, Charlie F. Ruan, Yingyi Huang, Yineng Zhang, Liangsheng Yin, Aksara Bayyapu, Luis Ceze, and Tianqi Chen. FlashInfer-Bench: Building the virtuous cycle for AI-driven LLM systems. In Proceedings of Machine Learning and Systems, volume 8, pages 2016–2064, 2026. URL https://proceedings.mlsys.org/paper_files/paper/2026/hash/ 37e44c4b5321605735be9761f9b758fc-Abstract-Conference.html. [20] Lingyun Yang, Yuxiao Wang, Shenghao Liang, Linfeng Yang, Daocheng Ying, Chunbo You, Rui Zhang, Luping Wang, Yinghao Yu, Guodong Yang, and Liping Zhang. Are LLM-generated GPU kernels production-ready? a trace-driven benchmark and optimization agent. arXiv preprint arXiv:2607.14541, 2026. URL https://arxiv.org/abs/2607. 14541. [21] Jaber Jaber and Osama Jaber. AutoKernel: Autonomous GPU kernel optimization via iterative agent-driven search. arXiv preprint arXiv:2603.21331, 2026. URL https://arxiv.org/abs/2603.21331.

11

[22] Haonan Li, Keyu Man, Partha Kanuparthy, Hanning Chen, Wei Sun, Sreen Tallam, Chenguang Zhu, Kevin Zhu, and Zhiyun Qian. TritonForge: Profiling-guided framework for automated Triton kernel optimization. arXiv preprint arXiv:2512.09196, 2025. URL https://arxiv.org/abs/2512.09196. [23] Anjiang Wei, Tianran Sun, Yogesh Seenichamy, Hang Song, Anne Ouyang, Azalia Mirhoseini, Ke Wang, and Alex Aiken. Astra: A multi-agent system for GPU kernel performance optimization. arXiv preprint arXiv:2509.07506, 2025. URL https://arxiv.org/abs/2509.07506. [24] Yuxuan Han, Meng-Hao Guo, Zhengning Liu, Wenguang Chen, and Shi-Min Hu. Making LLMs optimize multi-scenario CUDA kernels like experts. arXiv preprint arXiv:2603.07169, 2026. URL https://arxiv.org/abs/2603.07169. [25] Qitong Sun, Jun Han, Tianlin Li, Zhe Tang, Sheng Chen, Fei Yang, Aishan Liu, Xianglong Liu, and Yang Liu. KernelSkill: A multi-agent framework for GPU kernel optimization. arXiv preprint arXiv:2603.10085, 2026. URL https://arxiv.org/abs/2603.10085. [26] Genghan Zhang, Shaowei Zhu, Anjiang Wei, Zhenyu Song, Allen Nie, Zhen Jia, Nandita Vijaykumar, Yida Wang, and Kunle Olukotun. AccelOpt: A self-improving LLM agentic system for AI accelerator kernel optimization. In Proceedings of Machine Learning and Systems, volume 8, pages 541–568, 2026. URL https://proceedings.mlsys. org/paper_files/paper/2026/hash/0f8426558905746fc38da5e335700aec-Abstract-Conference.html. [27] Siqi Guo, Ming Lin, and Tianbao Yang. DRTriton: Large-scale synthetic data driven reinforcement learning for Triton kernel generation. arXiv preprint arXiv:2603.21465, 2026. URL https://arxiv.org/abs/2603.21465. [28] Marcin Spoczynski, Daniel Fleischer, Moshe Berchansky, Gabriela Ben-Melech Stan, Shira Guskin, Weilin Xu, Adam Siemieniuk, and Alexander Heinecke. Xe-Forge: Multi-stage LLM-powered kernel optimization for Intel GPU. arXiv preprint arXiv:2605.26118, 2026. URL https://arxiv.org/abs/2605.26118. [29] Nina Wiedemann, Quentin Leboutet, Michael Paulitsch, Diana Wofk, and Benjamin Ummenhofer. KernelFoundry: Hardware-aware evolutionary GPU kernel optimization. arXiv preprint arXiv:2603.12440, 2026. URL https: //arxiv.org/abs/2603.12440. [30] Xinguo Zhu, Shaohui Peng, Jiaming Guo, Yunji Chen, Qi Guo, Yuanbo Wen, Hang Qin, Ruizhi Chen, Qirui Zhou, Ke Gao, Yanjun Wu, Chen Zhao, and Ling Li. QiMeng-Kernel: Macro-thinking micro-coding paradigm for LLM-based high-performance GPU kernel generation. In Proceedings of the AAAI Conference on Artificial Intelligence, volume 40(34), pages 29168–29176, 2026. doi: 10.1609/aaai.v40i34.40155. URL https://ojs.aaai.org/index.php/ AAAI/article/view/40155. [31] Jonathan Ragan-Kelley, Connelly Barnes, Andrew Adams, Sylvain Paris, Frédo Durand, and Saman Amarasinghe. Halide: A language and compiler for optimizing parallelism, locality, and recomputation in image processing pipelines. In Proceedings of the 34th ACM SIGPLAN Conference on Programming Language Design and Implementation, pages 519–530, 2013. doi: 10.1145/2491956.2462176. URL https://doi.org/10.1145/2491956.2462176. [32] Bastian Hagedorn, Archibald Samuel Elliott, Henrik Barthels, Rastislav Bodík, and Vinod Grover. Fireiron: A data-movement-aware scheduling language for GPUs. In 2020 International Conference on Parallel Architectures and Compilation Techniques, pages 71–82, 2020. doi: 10.1145/3410463.3414632. URL https://doi.org/10. 1145/3410463.3414632. [33] Philippe Tillet, H. T. Kung, and David Cox. Triton: An intermediate language and compiler for tiled neural network computations. In Proceedings of the 3rd ACM SIGPLAN International Workshop on Machine Learning and Programming Languages, pages 10–19, 2019. doi: 10.1145/3315508.3329973. URL https://doi.org/10.1145/ 3315508.3329973. [34] Lei Wang, Yu Cheng, Yining Shi, Zhiwen Mo, Zhengju Tang, Wenhao Xie, Tong Wu, Lingxiao Ma, Yuqing Xia, Jilong Xue, Fan Yang, and Zhi Yang. TileLang: Bridge programmability and performance in modern neural kernels. In International Conference on Learning Representations, volume 2026, pages 73099–73134, 2026. URL https://proceedings.iclr.cc/paper_files/paper/2026/hash/ 76fb92288bf90360c527efb0d1c2aba6-Abstract-Conference.html.

12

A

Supplementary Material

A.1

Benchmark Composition

Table 3 Operator categories and representative design content in D2K-B ENCH. GEMM denotes general matrix multiplication. Operator category

Tasks

Attention and recurrent operators Mixture of experts and distributed execution Indexing, sampling, and control flow Quantization, training, and structural operators

A.2

Representative design content

Sparse attention, decoding with cached state, composable recurrences, and compressed state Grouped GEMM, routing, and distributed dispatch and combine

7 8 6

Index sharing, exact sampling, compact decision state, and verification Low-precision representation, clustering, optimizer updates, and fusion

5

Experimental Environment

We evaluate kernels on an x86_64 Linux system with eight NVIDIA B200 GPUs connected via NVLink and NVSwitch. We use Python 3.12 in Docker containers based on nvcr.io/nvidia/pytorch:26.06-py3. CUDA Toolkit 13.0 and a prerelease PyTorch 2.9 build (cu130) were installed separately for the experiments. We use Claude Code CLI 2.1.197 for Claude-Opus-4.8, GLM-5.3, Kimi-K3, and Fable-5, and Codex CLI 0.147.0 for GPT-6-Astra and GPT-5.6-Sol. The LLM judge, Qwen3.8-Max-0902, runs through Codex CLI.

A.3

Speedup for Fully Correct Models

As an additional analysis, we report speedup for models that produce a correct submission with valid latency measurements on every workload of all 26 tasks under both configurations. We first take the geometric mean of tb,i,w /tm,c,i,w over workloads within each task and then the geometric mean over all 26 tasks. Table 4 reports this speedup for GPT-6-Astra, Claude-Opus-4.8, GPT-5.6-Sol, and the expert implementations. The performance score Sperf in Table 1 covers all 26 tasks for every model. Table 4 Speedup over all 26 tasks relative to the first-correct Fable-5 baseline for models with correct submissions and valid latency measurements on every workload under both configurations. The expert implementations are shared across configurations, so the same result appears in both column groups. w/o Design Guidance

w/ Design Guidance

Correct tasks

Speedup

Correct tasks

Speedup

GPT-6-Astra Claude-Opus-4.8 GPT-5.6-Sol

26 26 26

2.68× 1.38× 1.31×

26 26 26

3.33× 2.31× 2.01×

Expert

26

2.33×

26

2.33×

Model

Thresholded speedup coverage. Following the thresholded metric in KernelBench [1], we additionally report

an adapted fastp relative to the same Fable-5 first-correct baseline used in our performance evaluation. For threshold p, let Cm,c contain the tasks for which model m under configuration c produces a correct submission with valid latency measurements on every workload. We define  1 X  fastFable 1 i ∈ Cm,c ∧ Sm,c,i |T | = 26. (5) p,m,c = perf > p , |T | i∈T

Thus, the metric measures the percentage of all tasks with a correct final submission whose task-level geometric mean speedup over Fable-5 exceeds p. The denominator of fastFable p,m,c remains all 26 tasks regardless 13

of correctness. We report fastFable p,m,c as a secondary thresholded summary; performance score Sperf remains the main benchmark metric. KernelBench emphasizes fast1 , while later evaluations commonly report fast1 and fast2 [16, 29]. We therefore report p ∈ {1, 1.5, 2}, including 1.5 as an intermediate threshold.

Table 5 Fable-relative thresholded speedup coverage. Entries give the percentage of tasks with a correct final submission whose task-level speedup exceeds each threshold. Overall includes all 130 model–task pairs. Model

Configuration

fastFable 1

fastFable 1.5

fastFable 2

GPT-6-Astra

w/o guidance w/ guidance w/o guidance w/ guidance w/o guidance w/ guidance w/o guidance w/ guidance w/o guidance w/ guidance

96.2% 96.2% 61.5% 88.5% 65.4% 88.5% 65.4% 61.5% 65.4% 73.1%

88.5% 92.3% 38.5% 65.4% 38.5% 53.8% 38.5% 46.2% 26.9% 50.0%

57.7% 73.1% 23.1% 50.0% 26.9% 34.6% 19.2% 34.6% 11.5% 30.8%

w/o guidance w/ guidance

70.8% 81.5%

46.2% 61.5%

27.7% 44.6%

Claude-Opus-4.8 GPT-5.6-Sol GLM-5.3 Kimi-K3 Overall

In the Overall row, guidance increases the percentage of model–task pairs that are correct and exceed the corresponding speedup threshold at all three thresholds. At p = 1.5, the guided results also preserve the performance score Sperf ordering: GPT-6-Astra, Claude-Opus-4.8, GPT-5.6-Sol, Kimi-K3, and GLM-5.3.

A.4

Token Efficiency

We examine token efficiency by jointly considering the performance during optimization and the token usage of the two input configurations. In Figure 4, the guided trajectory reaches a final score of 2.31, compared with 1.38 for the unguided curve, under the same 350-turn budget. We use S′perf , adapting Equation 1 by replacing the 0.5 fallback with 0.1 and using the geometric mean over all 26 tasks in Equation 2. The 0.1 fallback applies only until a task’s first eligible submission; we then use its highest eligible score so far. The smaller fallback reduces early decreases when a task’s first eligible score replaces the fallback; scores below 0.1 are used as measured. The additional token usage is concentrated in cached inputs (Table 6): cache-read tokens increase by 57.6% and cache-write tokens by 29.0% across 26 tasks. API cost rises from US$631.56 to US$820.10 (+29.9%). Cached input processing accounts for 97.3% of the estimated API cost increase, while output tokens rise only 2.4%; the token totals combine guidance text and interaction history.1

w/o design guidance

w/ design guidance

2.31

2.5 2.0

1.38

S0perf 1.5 1.0 0.5 0

0

50

100

150

200

250

300

350

Turns

Figure 4 Claude-Opus-4.8 optimization progress over the same 26 tasks. The performance score S′perf uses a 0.1 fallback.

1 API estimates use standard Claude-Opus-4.8 prices of US$5/6.25/0.50/25 per million uncached input/cache-write/cache-read/output tokens, assuming five-minute cache writes and no discounts or surcharges. Source: https://platform.claude.com/docs/en/ about-claude/pricing, September 21, 2026.

14

Table 6 Claude-Opus-4.8 token usage and estimated API cost across 26 tasks per input configuration. M denotes million tokens. Metric

w/o Design Guidance

w/ Design Guidance

Total

0.041 37.261 402.034 7.898 631.56

0.108 48.084 633.801 8.085 820.10

0.149 85.345 1,035.835 15.983 1,451.66

Uncached input tokens (M) Cache-write tokens (M) Cache-read tokens (M) Output tokens (M) API cost (US$)

Guidance improves final Sperf while increasing token usage and estimated API cost.

B

Generation Inputs and Prompts

The generation appendix presents the inputs supplied for the Muon orthogonalization kernel and the prompt used for the independent design run. The ordering follows Figure 1 and Section 2.2 of the main paper. Both input configurations receive the first six shared task files; the input configuration with design guidance additionally receives the expert design guidance defined in Section 2.3. In the reproduced files, K in the task description and input generator denotes the same trailing dimension as N in the workload configuration and guidance.

B.1

Muon Orthogonalization Kernel

Task Description. The task description defines the operator, interface, numerical semantics, fixed five-step

coefficient sequence, and completion requirements. # Five-step polynomial orthogonalization map ## Interface The returned tensor has the same shape and dtype as \(G\). The ‘steps‘ argument is retained for call compatibility and must equal 5. ## Semantics This benchmark isolates the matrix orthogonalization primitive used inside an orthogonalizing matrix optimizer. It does not include momentum, weight decay, learning-rate scaling, distributed communication, or parameter mutation. For every matrix in a tensor \(G\in\mathbb R^{[\ldots]\times M\times K}\), normalize independently in fp32, store the iterate in fp16, and apply the five coefficient triples below in order. \[ X_0=\operatorname{fp16}\!\left( \frac{\operatorname{fp32}(G)} {\|\operatorname{fp32}(G)\|_F+10^{-7}} \right), \] \[ A_i=X_iX_i^\top,\qquad X_{i+1}=a_iX_i+\left(b_iA_i+c_iA_i^2\right)X_i. \] Leading batch dimensions are independent. The coefficient sequence is: | \(i\) | \(a_i\) | \(b_i\) | \(c_i\) | |---:|---:|---:|---:| | 0 | 7.892582874424408 | -20.38301394587957 | 13.555306149406924 | | 1 | 3.911484868135431 | -2.5464635929060884 | 0.4268988319673074 | | 2 | 3.760657955697423 | -2.512819018216563 | 0.4323647349070073 | | 3 | 3.160399673686287 | -2.149649518898498 | 0.3996366907664389 | | 4 | 2.1910971618617303 | -1.441662010214663 | 0.328146487623155 | This finite polynomial map is the observable operator. It is an orthogonalization surrogate and does not promise an exact polar factor.

15

## Workloads Registered workloads use ‘[batch, M, K]‘. The workload entries in ‘config.json‘ are the registered evaluation workloads for this task. Additional correctness checks may vary inputs within ‘task_contract‘ without adding entries to the benchmark workload list. ## Correctness The returned tensor is compared against the reference with effective ‘rel_norm_rtol = 0.012‘ as the primary gate (‘atol = 0.006‘, ‘rtol = 0.02‘). ## Peak Memory Auxiliary GPU storage includes every live output, temporary tensor, and private persistent workspace beyond the distinct storages of the public inputs. Its required space complexity is ‘O(batch_size * M * N + batch_size * min(M, N)^2)‘, where ‘N‘ is the trailing dimension written ‘K‘ above. ## Execution Requirements Each registered matrix shape, batch, dtype, and five-step count may use a specialized kernel; matrix values remain runtime data. Perform the FP32 Frobenius normalization, FP16 iterate conversion, and all five distinct coefficient triples in their declared order for every matrix in the batch, then return the complete tensor. Numerical tolerance covers floatingpoint reassociation only; it does not permit fewer steps, repeated coefficients, early convergence, sampled rows or columns, or a cached result.

Workload Definitions. The config.json file contains the registered workload definitions. { "benchmark_id": "five_step_polynomial_orthogonalization", "fixed_params": { "steps": 5, "coefficients": "fixed_five_step_coefficients" }, "workloads": [ { "label": "production_b32_m2048_n5464", "batch_size": 32, "M": 2048, "N": 5464, "steps": 5, "specialization_params": { "batch_size": 32, "M": 2048, "N": 5464, "steps": 5, "dtype": "bf16" } }, { "label": "production_b32_m4096_n11008", "batch_size": 32, "M": 4096, "N": 11008, "steps": 5, "specialization_params": { "batch_size": 32, "M": 4096, "N": 11008, "steps": 5, "dtype": "bf16" } }, { "label": "production_b8_m7168_n18432", "batch_size": 8, "M": 7168, "N": 18432, "steps": 5, "specialization_params": { "batch_size": 8, "M": 7168, "N": 18432, "steps": 5, "dtype": "bf16"

16

} } ], "input_rules": { "static_inputs": [ "steps" ], "dynamic_inputs": [ "G" ] }, "task_contract": { "registered_specializations": [ { "batch_size": 32, "M": 2048, "N": 5464, "steps": 5, "dtype": "bf16" }, { "batch_size": 32, "M": 4096, "N": 11008, "steps": 5, "dtype": "bf16" }, { "batch_size": 8, "M": 7168, "N": 18432, "steps": 5, "dtype": "bf16" } ], "required_outputs": [ "five_step_polynomial_result" ], "preserve_semantics": true }, "tolerance": { "rel_norm_rtol": 0.006, "atol": 0.003, "rtol": 0.01, "note": "divergence comes from FP32 accumulation order in the iteration; rel_norm_rtol is the primary gate (grader x2 = 0.012 effective)" } }

Input Generator. The get_inputs.py file constructs inputs for the selected workload. import torch

def get_inputs(config=None): batch = 32 M = 2048 K = 5464 steps = 5 if config is not None: M = config.get("M", M) batch = config.get("batch_size", config.get("batch", batch)) K = config.get("K", config.get("N", K)) steps = config.get("steps", steps) torch.manual_seed(int((config or {}).get("seed", 42))) G = torch.randn(batch, M, K, dtype=torch.bfloat16, device="cuda") return [G, steps]

PyTorch Reference. The reference states the executable functional semantics against which candidate outputs

are checked. import torch

_COEFFICIENTS = (

17

(7.892582874424408, -20.38301394587957, 13.555306149406924), (3.911484868135431, -2.5464635929060884, 0.4268988319673074), (3.760657955697423, -2.512819018216563, 0.4323647349070073), (3.160399673686287, -2.149649518898498, 0.3996366907664389), (2.1910971618617303, -1.441662010214663, 0.328146487623155), )

def run(G, steps=5): if G.ndim < 2: raise ValueError("G must end in two matrix dimensions") if steps != 5: raise ValueError("this operator is defined for exactly five steps") output_dtype = G.dtype X = G.float() X = ( X / (torch.linalg.vector_norm(X, dim=(-2, -1), keepdim=True) + 1e-7) ).to(torch.float16) for a, b, c in _COEFFICIENTS: A = X @ X.transpose(-2, -1) X = a * X + (b * A + c * (A @ A)) @ X return X.to(output_dtype)

Solution Template. The agent edits the supplied solution template. The staged filename is solution.py; the

template contains only the required entry point. def run(G, steps=5): raise NotImplementedError

Implementation Requirements.

The implementation requirements define the admissible implementation route and prohibit delegation to an existing implementation of the target operator. requirements: - The timed core computation must be implemented with custom Triton kernels. - Do not call or wrap any pre-existing implementation of the target operator, any equivalent fused/composite operator, or the upstream expert API. - PyTorch may be used only for interface glue, tensor allocation, and shape/stride handling. It may not perform the timed core computation. - Every output and caller-owned tensor mutation must be computed from the current invocation’s inputs. - Do not cache, freeze, precompute, or reuse any value derived from dynamic inputs.

Expert Design Guidance.

Only the input configuration with design guidance receives the expert design guidance file. The expert design guidance records the L1 high-level algorithmic insights, L2 dataflow design, L3 low-level optimization tricks, and the cross-level dependencies among the three levels. Expert source code is not included in the generation workspace. { "cross_level_rationale": [ { "bottleneck": "Even a highly tuned direct Newton-Schulz loop applies every polynomial step to the full rectangular update.", "level_1_high_level_algorithmic_insights": "Compose polynomial actions on the smaller Gram side and defer their action on X.", "level_2_dataflow_design": "Let X persist unchanged while smaller square R/Q states span a composition segment.", "level_3_low_level_optimization_tricks": "Use square-state symmetric GEMM and GEMM-add kernels for interior phases, and invoke rectangular GEMM only at restart and final commit boundaries.", "cost_reduction": "Replaces several large rectangular GEMMs with smaller square products and few rectangular commits.", "conditions": "Required order: orient the Gram to min(M,N), form R, compose each segment’s polynomial actions into Q while advancing only square state, and apply Q to X only at the restart and final boundaries. Invalid partial states: a larger-side Gram, a new rectangular X after every coefficient, or Q composition followed by per-step X updates. Completion evidence: interior steps materialize only square R/Q state and exactly the declared rectangular commit boundaries survive." }, { "bottleneck": "Unbounded indirect Gram composition can accumulate numerical error despite saving rectangular work.", "level_1_high_level_algorithmic_insights": "Required continuation of the main chain: commit Q and rebuild the Gram at the prescribed restart.",

18

"level_2_dataflow_design": "Terminate the old Q/R lifetime, transfer ownership to committed X, and seed a fresh segment from its Gram.", "level_3_low_level_optimization_tricks": "The restart is a hard representation boundary, not an optional extra update interleaved with stale state.", "cost_reduction": "Retains most deferred-composition savings while controlling drift with one deliberate rectangular boundary.", "conditions": "Required order: finish the pre-restart Q, commit it to X, round/store the new iterate as specified, rebuild R from that committed X, discard old Q/R, and begin the next segment. Invalid partial states: carrying prerestart Q/R forward, rebuilding R from stale X, or placing an extra rectangular update inside the segment. Completion evidence: every post-restart square product is derived only from the newly committed X." } ], "level_1_high_level_algorithmic_insights": [ { "name": "Compose polynomial steps on the smaller Gram side before touching the rectangular update", "algorithmic_change": "A strong direct implementation can use tuned low-precision GEMMs, smaller-side Gram orientation, and symmetric publication, but still writes a new full rectangular X after every Newton-Schulz step. After Frobenius normalization, form the smaller symmetric Gram R. For P_k(R)=a_k I+b_k R+c_k R^2, compose Q<-P_k(R) Q while advancing R<P_k(R) R P_k(R); apply Q to X only at the prescribed restart and final boundaries.", "correctness_basis": "For X with R=X^T X in the tall case (or R=XX^T in the wide case), every odd polynomial update of X can be written as X times a polynomial in R (or that polynomial times X). Polynomial actions therefore compose on the smaller side before one rectangular application.", "performance_mechanism": "Replaces repeated large rectangular GEMMs across the iteration with smaller square Gram products and only occasional rectangular applications.", "dataflow_consequences": "The normalized rectangular X persists across a segment while R and Q evolve on the smaller side; X changes only when the composed action is committed.", "constraints": "Orientation follows the smaller dimension and the coefficient sequence is fixed by the public operator; matrix tiling remains tunable." }, { "name": "Restart at the fixed boundary before coefficient 2", "algorithmic_change": "After composing coefficients 0 and 1, apply the accumulated Q to X before coefficient 2, store the committed iterate in FP16, recompute the Gram matrix from that value, discard the old Q and R, and begin the new segment with coefficient 2.", "correctness_basis": "In exact arithmetic, the restart evaluates the composition of coefficients 0 and 1 before continuing with coefficients 2 through 4. The FP16 commit and Gram rebuild match the fixed restart used by the expert implementation; floating-point reassociation remains subject to the task tolerance.", "performance_mechanism": "Keeps most steps on the small side without paying the numerical and conditioning cost of an unbounded composition, using only the necessary extra rectangular boundary.", "dataflow_consequences": "Q and R have segment lifetimes; both die at restart, while X inherits the committed result and becomes the source of the next Gram.", "constraints": "The restart occurs after coefficients 0 and 1 and before coefficient 2; it is fixed rather than triggered by measured numerical state. Lower-level GEMM scheduling is not fixed." } ], "level_2_dataflow_design": { "overview": "Normalize each input matrix independently by its FP32 Frobenius norm, cast X to FP16, and orient the polynomial work onto the smaller matrix dimension. Build a symmetric Gram matrix and evolve a composed square action Q together with the Gram state through coefficient steps. Keep rectangular X unchanged within each segment. After coefficients 0 and 1, apply Q to X at the fixed restart before coefficient 2, rebuild the Gram from the committed FP16 X, and begin a fresh segment. Apply the final Q after coefficient 4. Symmetric products materialize both matrix triangles from one triangular computation, and no sequence of per-step rectangular update tensors is produced.", "parallelism_strategies": [ { "name": "Process independent matrices through batched small-side symmetric products", "description": "Parallelize independent input matrices and square Gram/operator tiles, exploiting symmetry for Gramlike products and using dense rectangular products only at segment boundaries.", "available_parallelism": "Batches and output tiles are independent; only polynomial step order and segment boundaries are sequential." } ], "stages": [ { "stage_id": "normalized_small_side_state", "operation": "For each independent matrix, compute its FP32 Frobenius norm, normalize it, cast X to FP16, choose the smaller side, and form the initial symmetric Gram state.", "outputs": "Normalized FP16 rectangular X and smaller-side symmetric R. R is a full symmetric square matrix produced from triangular work.", "data_layout": "X is rectangular; R is a full symmetric square matrix formed from triangular work.", "owner_and_residency": "The orthogonalization call retains X across segment boundaries and retains R only within the current segment.", "consumers_and_lifetime": "Polynomial composition consumes R; X is consumed only at Gram construction and compositioncommit boundaries.", "performance_condition": "For the registered wide-matrix workloads, the MxM left Gram is the smaller side." }, { "stage_id": "gram_polynomial_composition", "operation": "Evaluate each coefficient step with square products, update R as required for the next step, and compose the corresponding action into Q without touching rectangular X.", "outputs": "Segment-local square action Q and Gram state R. Q is updated for every coefficient; R is advanced only

19

when another coefficient in the same segment requires it.", "data_layout": "Q and R are smaller-side square matrices.", "owner_and_residency": "The current polynomial segment owns both matrices in global storage between square products.", "consumers_and_lifetime": "The next coefficient consumes them; at restart or final commit Q is applied to X and the segment state dies.", "performance_condition": "No rectangular X update is issued for an interior step of a composition segment." }, { "stage_id": "composition_commit_and_restart", "operation": "Apply Q on the orientation-correct side of X. At the fixed restart before coefficient 2, rebuild R from the committed FP16 X, discard the old Q and R, and initialize the new segment from coefficient 2. At the final boundary, return the updated tensor.", "outputs": "Committed rectangular X and, at the restart, a fresh symmetric Gram state R derived from that X.", "data_layout": "A rectangular matrix and a symmetric restart Gram matrix.", "owner_and_residency": "X inherits the committed storage; new R transfers to the next composition segment.", "consumers_and_lifetime": "A restarted X feeds the next segment and the final X becomes public output; the old Q/R die at commit.", "performance_condition": "Rectangular applications occur only at the required restart and final boundaries." } ], "edges": [ { "from": "input", "to": "normalized_small_side_state", "value": "batched rectangular input matrices", "storage_and_lifetime": "The per-matrix norm is temporary; normalized X persists to Gram construction.", "readiness": "The Frobenius reduction includes every matrix element." }, { "from": "normalized_small_side_state", "to": "gram_polynomial_composition", "value": "normalized X and initial smaller-side Gram", "storage_and_lifetime": "X remains unchanged while square R/Q states evolve.", "readiness": "The Gram orientation matches the smaller dimension." }, { "from": "gram_polynomial_composition", "to": "composition_commit_and_restart", "value": "composed square action Q", "storage_and_lifetime": "One square operator crosses the segment boundary instead of multiple rectangular iterates.", "readiness": "The segment’s ordered coefficient steps are complete." }, { "from": "composition_commit_and_restart", "to": "gram_polynomial_composition", "value": "fresh Gram after restart", "storage_and_lifetime": "Old R and Q die; only committed X and rebuilt R continue.", "readiness": "The restart boundary has applied Q to X and recomputed the Gram." }, { "from": "composition_commit_and_restart", "to": "output", "value": "final orthogonalized matrix", "storage_and_lifetime": "Only the final rectangular tensor survives.", "readiness": "The final composed action has been applied on the orientation-correct side." } ], "restrictions": [ "A new full rectangular X tensor for every Newton-Schulz step.", "A Gram matrix formed on the larger matrix dimension.", "Independent evaluation of every polynomial action against X instead of composition in Q.", "Computing both symmetric triangles as separate matrix products.", "Carrying pre-restart Q or R into the next segment." ], "adoption_criteria": [ "Each registered wide-matrix workload (M<N) uses the smaller left Gram R=XX^T.", "Interior polynomial steps update only square R/Q state.", "Q is applied to X after coefficients 0 and 1 at the fixed restart before coefficient 2, and again after coefficient 4 at the final boundary.", "Restart recomputes Gram from committed X and clears old composition state.", "Symmetric products avoid duplicated off-diagonal computation while producing both triangles." ] }, "level_3_low_level_optimization_tricks": [ { "name": "Compute one Gram-matrix triangle and publish both symmetric halves", "supports": [ "Compose polynomial steps on the smaller Gram side before touching the rectangular update", "normalized_small_side_state", "gram_polynomial_composition"

20

], "challenge": "Gram and symmetric polynomial products are full square outputs, but computing mirrored off-diagonal tiles twice wastes nearly half of the product work.", "strategy": "Schedule only one triangular set of square output tiles. After an off-diagonal tile is complete, publish its values to both the direct and transposed output locations; diagonal tiles publish once. Keep the product accumulator live through both stores.", "pitfall": "A generic dense GEMM computes both mirrored halves independently, while a one-triangle store leaves later dense consumers without a complete matrix.", "constraints": "Tile dimensions, triangular traversal order, and store mechanism remain tunable." }, { "name": "Use square-state kernels inside each composition segment", "supports": [ "Compose polynomial steps on the smaller Gram side before touching the rectangular update", "gram_polynomial_composition", "composition_commit_and_restart" ], "challenge": "The advantage disappears if an interior coefficient phase issues or republishes a rectangular X update.", "strategy": "Within a segment, execute the ordered R and Q updates with square symmetric GEMM or GEMM-add kernels. Issue a shape-correct rectangular GEMM only to apply Q at the fixed restart or final boundary; do not update X during an interior coefficient phase.", "pitfall": "Updating X between coefficient steps restores the repeated large rectangular GEMMs the algebra was designed to remove.", "constraints": "Exact tile dimensions, warp-group roles, asynchronous transfer mechanism, pipeline depth, and persistent scheduling policy remain open." } ] }

B.2

Independent Design Proposal Prompt

The prompt below specifies the criteria for an independent design proposal and the required model_design.json output. The brief is staged as CLAUDE.md for all models; this release filename is not model-specific. The five task inputs named in the design brief and design_schema.json remain separate files that the agent reads with read-only workspace tools. The run receives no solution template or expert design guidance and is instructed to produce only model_design.json through submit_design, without writing or executing kernel code. The delimiter lines distinguish the system message from the design brief in this appendix. ===== SYSTEM MESSAGE ===== You are producing a closed-book GPU-kernel DESIGN deliverable. You have no GPU and must not write or run kernel code. Use the read-only workspace tools to read CLAUDE.md (the design brief) and every other public workspace file, including design_schema.json. When the design is complete, call submit_design with the complete model_design.json as a JSON string that validates against design_schema.json. Do not return a text-only final answer. ===== DESIGN BRIEF STAGED AS CLAUDE.md ===== # Mode 1 – author the design Read the workspace first: ‘task_description.md‘, ‘reference.py‘, ‘get_inputs.py‘, ‘config.json‘, ‘implementation_requirements.yaml‘, and ‘design_schema.json‘. The first five are the shared task inputs; the last is the output schema. The target GPU is **NVIDIA B200 (Blackwell), sm_100**, the same target used for kernel generation. There is no reference design to consult and no way to measure anything in this mode. Produce one file, ‘model_design.json‘, valid against ‘design_schema.json‘. Nothing else is graded. You are not being asked for a plan you could write for any kernel. A design that would read the same for a different operator has said nothing. Every claim should be one that this operator’s shapes, dtypes, and declared ranges make true, and false for some other operator. ## The angles to think along These are the fields the schema requires. They are the angles because each one is a question that can be answered wrongly. ### Cross-level rationale A chain is one causal story from a cost to its removal, carried through all three levels. For each chain state: - **bottleneck** – what actually dominates, at the registered shapes, and not in general. - **level_1_high_level_algorithmic_insights** – what changes in the mathematics. - **level_2_dataflow_design** – what that change does to which values exist, where, and for how long. - **level_3_low_level_optimization_tricks** – what the execution has to do to keep the benefit. - **cost_reduction** – the specific cost that no longer exists. Name it, do not say "overhead".

21

- **conditions** – the required order, the partial states that are invalid, and the evidence that the chain is complete. This is the field that decides whether the chain is a design or a wish: if a chain can be half-adopted and still look adopted, say what distinguishes them. ### Level 1 – High-level algorithmic insights Changes to the algebra, computation domain, recurrence, reduction, representation, or precision semantics. Per item: - **algorithmic_change** – the change itself. - **correctness_basis** – why the result is still exact, or under exactly which condition the difference is admissible. "Within tolerance" is not a basis; the condition is. - **performance_mechanism** – why this is faster, in terms of a resource. - **dataflow_consequences** – what level 2 now has to be. - **constraints** – what is left open to the implementation, deliberately. ### Level 2 – Dataflow design Specify where values are produced, stored, and consumed; their layouts and lifetimes; and which intermediate values are reused or avoided. Per design: - **overview** – who owns what, and what the operator looks like as a whole. - **parallelism_strategies** – the exploitable parallelism, and how much of it exists **at the registered shapes**. If a decomposition exposes fewer concurrent units than the machine has, that is a property of the design and belongs here with the arithmetic. Each entry gives its **name**, **description**, and **available_parallelism**. - **stages** – per stage: operation, outputs, data layout, owner and residency, consumers and lifetime, and the performance condition that must hold for the stage to be worth having. - In **data_layout**, describe indexing, shape, packing, or thread/register mapping. Put dtype and result completion in **outputs**, ownership and storage domain in **owner_and_residency**, and handoff requirements in edge **readiness**. State explicitly when a layout is unspecified. - **edges** – per edge: **from**, **to**, **value**, **storage_and_lifetime**, and **readiness**. - **restrictions** – values that must never be materialized. Each entry costs something: it forbids a structure. Only forbid what you would still forbid after measuring, and if a prohibition depends on a shape or a device property, say so in the entry rather than stating it unconditionally. - **adoption_criteria** – how a reader decides, from an implementation, whether this design was adopted. Each check should be decidable by looking at code, not by asking the author. ### Level 3 – Low-level optimization tricks Execution mechanisms that realize the algorithm and dataflow design, such as software pipelining that overlaps data movement with computation; warp or warp-group specialization and roles; producer-consumer scheduling; load balancing; required ordering and synchronization; layout transition instructions; and architecture-specific mechanisms. Do not repeat Level 2 value placement, storage residency, layout, or lifetime descriptions as standalone Level 3 ideas. Per item: - **supports** – which level 1 or level 2 element this exists to preserve. - **challenge** – the thing that makes it hard. - **strategy** – what the execution does. - **pitfall** – what a straightforward implementation does instead, and what that costs. - **constraints** – what stays tunable. State what must be *measured* rather than assumed. If a choice rests on a hardware rate, say which measurement would refute it. ## What is not credit - Naming a mechanism without the condition under which it holds. - A prohibition with no cost attached, or one that would forbid the fastest legal structure. - A bottleneck claim that does not reference this operator’s shapes. - A ‘conditions‘ entry that only restates the chain.

Design Output Schema. The separate schema defines model_design.json; its inherited Expert Design

Guidance title reflects reuse of a structural schema. { "$schema": "https://json-schema.org/draft/2020-12/schema", "title": "D2K-Bench Expert Design Guidance", "description": "Expert design annotations organized into High-level algorithmic insights, Dataflow design, and Low-level optimization tricks. The annotations describe algorithmic transformations, data representation and movement, and kernel execution mechanisms, together with dependencies among levels. Implementation and tuning remain agent decisions.", "type": "object", "required": [ "level_1_high_level_algorithmic_insights", "level_2_dataflow_design", "level_3_low_level_optimization_tricks", "cross_level_rationale"

22

], "additionalProperties": false, "properties": { "level_1_high_level_algorithmic_insights": { "type": "array", "minItems": 1, "description": "Task-specific changes to the algebra, computation domain, recurrence, reduction, representation, or precision semantics.", "items": { "type": "object", "required": [ "name", "algorithmic_change", "correctness_basis", "performance_mechanism", "dataflow_consequences", "constraints" ], "additionalProperties": false, "properties": { "name": { "type": "string", "minLength": 3 }, "algorithmic_change": { "type": "string", "minLength": 30 }, "correctness_basis": { "type": "string", "minLength": 20 }, "performance_mechanism": { "type": "string", "minLength": 20 }, "dataflow_consequences": { "type": "string", "minLength": 20 }, "constraints": { "type": "string", "minLength": 15 } } }, "title": "High-level algorithmic insights" }, "level_2_dataflow_design": { "type": "object", "description": "Dataflow design specifies where values are produced, stored, and consumed; their layouts and lifetimes; and which intermediate values are reused or avoided.", "required": [ "overview", "parallelism_strategies", "stages", "edges", "restrictions", "adoption_criteria" ], "additionalProperties": false, "properties": { "overview": { "type": "string", "minLength": 80 }, "parallelism_strategies": { "type": "array", "minItems": 1, "items": { "type": "object", "required": [ "name", "description", "available_parallelism" ], "additionalProperties": false, "properties": { "name": { "type": "string" },

23

"description": { "type": "string", "minLength": 20 }, "available_parallelism": { "type": "string", "minLength": 10 } } } }, "stages": { "type": "array", "minItems": 1, "items": { "type": "object", "required": [ "stage_id", "operation", "outputs", "data_layout", "owner_and_residency", "consumers_and_lifetime", "performance_condition" ], "additionalProperties": false, "properties": { "stage_id": { "type": "string", "pattern": "^[a-z][a-z0-9_]*$" }, "operation": { "type": "string", "minLength": 20 }, "outputs": { "type": "string", "minLength": 10 }, "data_layout": { "type": "string", "minLength": 15, "description": "Element indexing, shape, packing, or thread/register mapping stated in the source design. State explicitly when unspecified. Put dtype and result completion in outputs, logical ownership and storage domain in owner_and_residency, and handoff conditions in edge readiness." }, "owner_and_residency": { "type": "string", "minLength": 15 }, "consumers_and_lifetime": { "type": "string", "minLength": 15 }, "performance_condition": { "type": "string", "minLength": 20 } } } }, "edges": { "type": "array", "minItems": 1, "items": { "type": "object", "required": [ "from", "to", "value", "storage_and_lifetime", "readiness" ], "additionalProperties": false, "properties": { "from": { "type": "string" }, "to": { "type": "string" },

24

"value": { "type": "string" }, "storage_and_lifetime": { "type": "string", "minLength": 15 }, "readiness": { "type": "string", "minLength": 15 } } } }, "restrictions": { "type": "array", "minItems": 1, "items": { "type": "string" } }, "adoption_criteria": { "type": "array", "minItems": 1, "items": { "type": "string" } } }, "title": "Dataflow design" }, "level_3_low_level_optimization_tricks": { "type": "array", "minItems": 1, "description": "Low-level optimization tricks are execution mechanisms that preserve the benefits of the algorithm and dataflow design. This level may cover software pipelining, warp or warp-group specialization and roles, producer-consumer scheduling, synchronization and ordering, load balancing, layout-transition instructions, and architecture-specific mechanisms when source-backed. Value placement, storage residency, layouts, and lifetimes belong to dataflow design; exact schedule and tuning parameters remain open.", "items": { "type": "object", "required": [ "name", "supports", "challenge", "strategy", "pitfall", "constraints" ], "additionalProperties": false, "properties": { "name": { "type": "string", "minLength": 3 }, "supports": { "type": "array", "items": { "type": "string" }, "minItems": 1 }, "challenge": { "type": "string", "minLength": 20 }, "strategy": { "type": "string", "minLength": 25 }, "pitfall": { "type": "string", "minLength": 20 }, "constraints": { "type": "string", "minLength": 15 } } }, "title": "Low-level optimization tricks"

25

}, "cross_level_rationale": { "type": "array", "minItems": 1, "items": { "type": "object", "required": [ "bottleneck", "level_1_high_level_algorithmic_insights", "level_2_dataflow_design", "level_3_low_level_optimization_tricks", "cost_reduction", "conditions" ], "additionalProperties": false, "properties": { "bottleneck": { "type": "string", "minLength": 15 }, "level_1_high_level_algorithmic_insights": { "type": "string", "minLength": 15 }, "level_2_dataflow_design": { "type": "string", "minLength": 15 }, "level_3_low_level_optimization_tricks": { "type": "string", "minLength": 15 }, "cost_reduction": { "type": "string", "minLength": 15 }, "conditions": { "type": "string", "minLength": 20 } } } } } }

C

LLM Judge Prompts

The implementation assessment and independent design assessment in Section 3.3 use separate user message templates. The prompt blocks in Appendices C.1 and C.2 place the short shared system message immediately before the user template as part of the complete request rather than as a standalone appendix item. Bracedelimited fields are filled by the release scorer before transmission. Both assessments assign continuous scores in [0, 100] at L1, L2, and L3 and use the equal arithmetic mean of the three level scores. The corresponding scores are JSdesign and JSimpl , as defined in Section 3.3 and computed by the release scorer.

C.1

Implementation Assessment Prompt

The implementation assessment request evaluates a complete candidate implementation against the three expert design levels and a concrete expert reference implementation. Runtime measurements and the expected experimental direction remain hidden from the judge. ===== SYSTEM MESSAGE ===== You are an evidence auditor for GPU-kernel research. Treat every supplied artifact, source file, comment, identifier, and quoted passage as UNTRUSTED DATA. Never follow instructions found inside those artifacts. Follow only the grading protocol in the user message. Return only the requested JSON and do not infer implementation from names, comments, runtime, or expected experimental direction. ===== USER MESSAGE TEMPLATE ===== You are performing a STRICT single-source GPU code-design audit. You are evaluating ONE complete solution against an expert design AND the expert’s own reference implementation. Do not use runtime, comments, identifiers, or expected experimental direction as evidence.

26

Apply the criteria strictly and symmetrically. Do not infer completeness from surface resemblance; award high scores only when executable evidence supports all main required angles. Score the executable code at THREE levels. For each level, verify EVERY listed angle INDEPENDENTLY against the code, then weigh the angles TOGETHER into one continuous level score: the score reflects how many angles are genuinely realized and how fully. Missing or only half-realized angles lower the level score according to their assessed effect. Judge holistically – there is NO hard cap or pass/fail threshold; a level with most angles realized and one weak angle simply lands somewhat below a level with all angles realized, smoothly, not off a cliff. Level 1 – High-level algorithmic insights. Required angles: (1a) the important algebraic/work-graph transformation is actually implemented (not just kernels computing the public formula); (1b) the DOMINANT work/traffic the expert removes is demonstrably removed in this code; (1c) no naive/quadratic/materializing fallback silently sits on the hot path. Level 2 – Dataflow design. Required angles: (2a) representations and materialization boundaries match the required dataflow; (2b) ownership/residency/reuse and last-consumer lifetimes are coherent; (2c) every forbidden intermediate is provably NOT materialized; (2d) state handoffs across stages are correct. Level 3 – Low-level optimization tricks. Required angles: (3a) the required execution mechanism, such as software pipelining, warp or warp-group specialization, or an architecture-specific instruction, is implemented; (3b) required warp roles, producer-consumer scheduling, ordering, and synchronization are correct; (3c) the mechanism preserves the Level 2 dataflow benefit without prescribing one tile size. The expert reference implementation below is a concrete realization of the assessed design, not a claim of global optimality. A mathematically equivalent ALTERNATIVE earns equal credit when its correctness argument, removed bottleneck, representation, ownership/lifetime, and final consumer are all visible in executable code; explain the equivalence. Do not penalize a genuinely superior alternative – but do not credit surface resemblance the code does not realize. Give each level a CONTINUOUS 0-100 score. Calibrate STRICTLY against these anchors: complete_implementation 100 – every required angle fully realized to the expert reference’s completeness substantial_implementation 75 – ALL required angles present; only a secondary refinement is missing partial_implementation 50 – coherent but at least one REQUIRED angle is missing or only half-realized isolated_fragment 25 – isolated relevant fragments; the end-to-end consequence is absent absent 0 – naive/incompatible path or no relevant mechanism Reserve 75+ for code that genuinely realizes all main angles, and 90+ only for completeness that matches the expert reference. Code that realizes only some angles receives a score reflecting the realized and missing mechanisms, without a fixed cap. Assign a numeric score at every level. The scorer computes JS_impl = (L1 + L2 + L3) / 3. All three levels have equal influence. EVIDENCE IS MANDATORY AND MACHINE-VERIFIED. For every level score you MUST cite 1-4 exact executable passages. Each citation needs an inclusive line range and a quote copied CHARACTER-FOR-CHARACTER from those exact lines of the numbered SOLUTION source (leading indentation may be trimmed, nothing else). A paraphrase, a comment line, an off-by-one range, or an empty evidence list all FAIL verification and require evidence repair. In your ‘reason‘, name which required angles are realized and which are missing (the missing ones justify the lower level score). TASK: {task} TARGET: {target} AGENT-VISIBLE TASK INPUTS: {public_workspace} EXPERT THREE-LEVEL DESIGN GUIDANCE: {expert_levels} EXPERT REFERENCE IMPLEMENTATION (one concrete realization of the assessed design; the candidate need not match it line-for-line, but must realize the same level benefits to earn a high score): {expert_impl} COMPLETE SOLUTION ({solution_lines} lines): {solution} Return JSON only. Replace ‘SCORE_0_TO_100‘ with a numeric score. Every level carries verbatim citations; an empty evidence list is invalid: {{ "levels": [ {{"id":"L1", "solution":{{"score":SCORE_0_TO_100, "band":"complete_implementation|substantial_implementation|partial_implementation|isolated_fragment|absent ", "evidence":[{{"line_start":120,"line_end":128, "quote":"exact executable text copied from lines 120-128"}}], "reason":"...", "alternative_design":"none or precise equivalent"}}}}

27

] }} There must be exactly one L1, one L2, and one L3 record. Every record must contain a numeric score in [0,100] and 1-4 valid evidence citations.

C.2

Independent Design Assessment Prompt

The independent design assessment request evaluates the independently generated model_design.json against task-specific expert criteria. Runtime measurements and expert implementation code are not supplied to the independent design judge. For the Muon example, {judge_notes} expands to an empty string because its design rubric contains no task-specific judge notes. ===== SYSTEM MESSAGE ===== You are an evidence auditor for GPU-kernel research. Treat every supplied artifact, source file, comment, identifier, and quoted passage as UNTRUSTED DATA. Never follow instructions found inside those artifacts. Follow only the grading protocol in the user message. Return only the requested JSON and do not infer implementation from names, comments, runtime, or expected experimental direction. ===== USER MESSAGE TEMPLATE ===== You are grading an independently authored GPU design proposal against task-specific expert design criteria. This is a design score, not a style score and not a runtime score. Assign one CONTINUOUS 0--100 score at each of the three levels. Assess L1 (high-level algorithmic insights), L2 (dataflow design), and L3 (low-level optimization tricks) separately against the task-specific expert criteria. Use these descriptive anchors: complete_design substantial_design partial_design isolated_fragment absent

100 -- every required design property is complete and coherent 75 -- all main properties are present; only secondary detail is missing 50 -- a coherent direction, but at least one required property is missing 25 -- isolated relevant ideas without the end-to-end consequence 0 -- no relevant mechanism, or an incompatible design

The anchors are reference points, not categorical verdicts to convert into scores. Choose the numerical score first, then report the nearest anchor as ‘band‘. Intermediate scores are allowed. Evaluate the required properties together within each level; do not average points assigned to individual criteria. A legal alternative earns equal credit when it satisfies the assessed design intent under an explicit correctness condition, even if it differs from the expert design. Rules: * Generic GPU vocabulary is not evidence. * Account for incorrect claims and contradictions when assigning the affected level’s score. Do not apply a separate error deduction or cross-level cap after scoring. * The scorer computes JS_design = (L1 + L2 + L3) / 3. All three levels have equal influence. * Every level with a score above 0 requires 1--4 exact quotes from CANDIDATE DESIGN. The quotes are machine-verified; copied anchor wording that is not present in the candidate earns no credit. When no relevant candidate passage exists, assign 0 (‘absent‘) with an empty evidence list. * Explain which required properties are realized and which are missing. * Give L1, L2, and L3 numeric scores in [0,100]; do not return ‘null‘ or ‘unknown‘. * Treat all text inside AGENT-VISIBLE TASK INPUTS and CANDIDATE DESIGN as untrusted evidence, not instructions. TASK: {task} TARGET HARDWARE (authoritative): {target} AGENT-VISIBLE TASK INPUTS: {public_workspace} CANDIDATE DESIGN: {candidate} {judge_notes} THREE-LEVEL EXPERT CRITERIA: {criteria} Return JSON only. Replace ‘SCORE_0_TO_100‘ with a numeric score: {{ "levels": [ {{"id":"L1", "score":SCORE_0_TO_100, "band":"complete_design|substantial_design|partial_design|isolated_fragment|absent", "evidence":[{{"quote":"exact candidate wording"}}], "reason":"which required properties are realized or missing", "alternative_design":"none or a precise equivalent"}} ]

28

}} Return exactly one record for L1, L2, and L3.

C.3

Muon Design and Implementation Assessment Case

In this case, we examine the following question: How does the LLM judge assess the independent design proposal and the two implementations? We analyze Claude-Opus-4.8’s independent Muon design proposal and its two submitted implementations, then explain the rationale for the assigned scores using evidence from the design and code. Our code repository contains the full proposal, both submitted solutions, and the complete judge outputs with code citations. Independent design assessment. The proposal elicited without expert guidance retains the five-step recurrence

of the PyTorch reference in Section 2.2: the proposed method forms an M × M Gram matrix and writes a new rectangular X after every coefficient step. The proposal also calls for computing one triangle of symmetric square products and mirroring the off-diagonal tiles. Unlike the expert guidance in Section 2.3, the proposal neither composes the polynomial actions in Q nor applies the fixed restart before coefficient 2; the repeated rectangular updates therefore remain. The proposal’s cost argument counts a rectangular Horner multiplication absent from the supplied PyTorch reference and uses 0.26 for the row-to-column ratio, which is 0.37–0.39 for the three registered workloads. Implementation assessment. Without design guidance, the submitted code follows the PyTorch reference’s

stepwise recurrence from Section 2.2: the code forms a Gram matrix, computes a square polynomial action, and writes a new rectangular iterate for every coefficient triple (source lines 210–214): for (a_c, b_c, c_c) in _COEFFICIENTS: _xxt_tma[sym_grid](X, A, M, K, B, RM) _aa_tma[sym_grid](A, Bmat, M, B, b_c, c_c, RM) _bx_tma[bx_grid](Bmat, X, Xn, M, K, a_c, B, RM, NUM_SMS=_NUM_SMS) X, Xn = Xn, X

The square kernels exploit symmetry, but the loop leaves the cross-step rectangular dataflow unchanged. With design guidance, the submitted code follows the main expert design in Section 2.3: the code composes the polynomial actions in a square operator Q while advancing the Gram state, then applies Q to the rectangular input once at the end (source lines 315–325): X0 = normalize(G) # [B, M, K] fp16 A = bmm_sym(X0, X0, transB=True) # X0 @ X0^T -> [B, M, M] (symmetric) Q = None for i, (a, b, c) in enumerate(_COEFFICIENTS): A2 = bmm_sym(A, A) # symmetric P = poly(A, A2, a, b, c) Q = P if Q is None else bmm_sym(P, Q) # Q symmetric (product of commuting sym P’s) if i < 4: T = bmm_sym(P, A) # P@A symmetric (P is a polynomial in A) A = bmm_sym(T, P) # (P@A)@P symmetric X = bmm(Q, X0, out_dtype=out_dtype) # [B, M, K], cast fused into epilogue

The guided implementation follows the square composition strategy but omits the expert’s fixed restart before coefficient 2. Here A is the square Gram state R of Section 2.3, and Q is the composed action Q. The guided implementation uses the left Gram, which is the smaller side for the registered Muon workloads. The guided implementation’s square kernel publishes an off-diagonal tile and its transpose from the same accumulator (source lines 208–215): # store (m, n) c_ptrs = C + bid * scb + (om[:, None] * scm + on[None, :] * scn) tl.store(c_ptrs, acc16, mask=(om[:, None] < M) & (on[None, :] < M)) # store transpose at (n, m) when off-diagonal if pid_n < pid_m: acc_t = tl.trans(acc16) ct_ptrs = C + bid * scb + (on[:, None] * scm + om[None, :] * scn) tl.store(ct_ptrs, acc_t, mask=(on[:, None] < M) & (om[None, :] < M))

LLM judge assessment. Both assessments use the numerical anchors absent (0), isolated fragments (25), partial (50), substantial (75), and complete (100), with intermediate scores permitted. For the independent proposal,

29

L1 receives 32, between isolated fragments and partial: triangular Gram products provide an algorithmic saving, but the central Q composition and fixed restart are missing, and the claimed dominant saving is measured against work absent from the PyTorch reference. L2 receives 40, below partial: normalization, the M × M Gram (the smaller side for these workloads), and buffer reuse are described, but X is written at every step and no segment-local R/Q lifetime or restart handoff is proposed. L3 receives 45, between isolated fragments and partial design: triangular scheduling and mirrored stores are specified, but square-state execution across a composition segment is not. The arithmetic mean of these scores is JSdesign = 39. Table 7 Muon implementation judge scores and performance for Claude-Opus-4.8 on this task. Sperf is the geometric mean speedup over the first-correct baseline across the three Muon workloads. Input configuration

L1

L2

L3

JSimpl

Sperf

w/o design guidance w/ design guidance

40 62

54 70

62 86

52 73

7.22 11.95

For the unguided implementation, L1=40 lies between isolated fragments and partial implementation: the cited loop uses symmetric square work but has no composed Q and retains five rectangular updates. L2=54 is near partial: buffer reuse is coherent, but X is materialized at every step. L3=62 remains below substantial: mirrored stores are present, but the rectangular iterate still streams every step. Their mean is JSimpl = 52. For the guided implementation, L1=62 credits the square Q composition and removal of repeated rectangular updates, but not the fixed restart before coefficient 2. L2=70 credits an unchanged X0 and evolving square states, but lacks the restart handoff. Although the guided implementation omits the fixed restart before coefficient 2, the implementation passed numerical correctness checks on all three registered workloads. L3=86 credits triangular computation, mirrored stores, and square-state interior work. Their mean is JSimpl = 73. The higher implementation score with design guidance is consistent with the lower measured latency across all three Muon workloads, illustrating the LLM judge’s effectiveness in identifying design differences relevant to performance in this case.

30

Related documents

• MetaKernelBench: Measuring GPU Kernel Knowledge Transfer Beyond Code#1200665
• D2K-Bench: Can LLM Agents Turn Expert Designs into Efficient GPU Kernels?#1187904
• KLineage: Recovering the Missing When of Kernel Optimization by Deoptimizing Experts#1141552
• KernelZero: Co-Evolving Proposer and Coder for Continuously Improved GPU Kernel Generation#1117786
• KernelOPT: Dispatch-Aware Agentic Search for GPU Kernel Optimization#1078130