Paper deep dive
D2K-Bench: Can LLM Agents Turn Expert Designs into Efficient GPU Kernels?
Daifeng Li, Huiqiang Jiang, Chengruidong Zhang, Wei Wu, Xudong Guo, Jianhong Tu, Jianwei Zhang, Binhang Yuan, Dayiheng Liu
Intelligence
Status: not_run | Model: - | Prompt: - | Confidence: 0%
Entities (0)
Relation Signals (0)
No relation signals yet.
Cypher Suggestions (0)
No Cypher suggestions yet.
Abstract
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-Bench, 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\times$ to $2.49\times$. 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.
Tags
Links
- Source: https://arxiv.org/abs/2610.03226v1
- Canonical: https://arxiv.org/abs/2610.03226v1
Trouble viewing inline? Open PDF directly →
Full Text
104,301 characters extracted from source content.
Expand or collapse full text
1]HKUST 2]Alibaba Group 3]USTC D2K-Bench: Can LLM Agents Turn Expert Designs into Efficient GPU Kernels? Daifeng Li Huiqiang Jiang Chengruidong Zhang Wei Wu Xudong Guo Jianhong Tu Jianwei Zhang Binhang Yuan Dayiheng Liu Affiliation: [ Affiliation: [ Affiliation: [ 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-Bench, 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×1.69× to 2.49×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: https://github.com/QwenLM/D2K-Bench 1 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. 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, 5, 3]. 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 expert-assisted 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-Bench, 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 implementation 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×1.69× to 2.49×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.1 Benchmark Task Selection D2K-Bench 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 orthogonalization 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 files for the Muon task. Each matrix is normalized in FP32 and then cast to FP16 to initialize 0=FP16((∥F+ϵ))X_0=FP16\! ( G( _F+ε) ). The kernel then conducts five polynomial Newton-Schulz updates with FP16 iterates. For a wide matrix, each update step, indexed by k=0,…,4k=0,…,4, forms the Gram matrix k=kk⊤R_k=X_kX_k and computes k+1=(ak+bkk+ckk2)kX_k+1=(a_kI+b_kR_k+c_kR_k^2)X_k, where I is the identity matrix and ak,bk,cka_k,b_k,c_k are the supplied coefficients. The normalization uses the Frobenius norm ∥F _F and a small positive constant ϵε. One workload contains 32 BF16 matrices of shape 4096×110084096×11008: each rectangular iterate has this shape, whereas its Gram matrix is 4096×40964096×4096. The direct PyTorch computation rebuilds the Gram matrix and writes a new rectangular iterate at every step, repeatedly processing the larger matrix. 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. 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 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 k=ak+bkk+ckk2P_k=a_kI+b_kR_k+c_kR_k^2. In exact arithmetic, since k+1=kkX_k+1=P_kX_k, steps 00 through j can be composed into =j⋯0Q=P_j·sP_0. The resulting rectangular matrix is then computed as j+1=0X_j+1=QX_0 in a single matrix multiplication, avoiding explicit computation of the intermediate rectangular matrices 1,…,jX_1,…,X_j. Meanwhile, the Gram matrix is updated directly as k+1=kkk⊤R_k+1=P_kR_kP_k , rather than recomputed as k+1k+1⊤X_k+1X_k+1 . This algorithmic transformation preserves the result in exact arithmetic while replacing repeated rectangular-matrix 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=2k=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 =a2+b2+c22Q=a_2I+b_2R+c_2R^2 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 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 SperfS_perf (Section 3.2), and define the LLM judge scores for independent designs, JSdesignJS_design, and generated implementations, JSimplJS_impl (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. 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 JSimplJS_impl. The evaluator measures SperfS_perf. 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. 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. Performance is measured relative to the first-correct baseline: Fable-5’s first correct submission in Claude Code without design guidance. Let iW_i denote the workloads of task i, tb,i,wt_b,i,w the baseline latency, and tm,c,i,wt_m,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: Sperfm,c,i=GMw∈i(tb,i,wtm,c,i,w),if checks and timing pass,0.5,otherwise.S_perf^m,c,i= cases _w _i ( t_b,i,wt_m,c,i,w ),&if checks and timing pass,\\[4.0pt] 0.5,&otherwise. cases (1) The model performance score over the task set T is Sperfm,c=GMi∈(Sperfm,c,i),||=26.S_perf^m,c=GM_i (S_perf^m,c,i ), |T|=26. (2) Here, GMGM 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×2× performance. Overall SperfS_perf 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,wt_m,c,i,w in Equation 1 and then applying Equation 2; the baseline has Sperf=1.0S_perf=1.0 by definition. The gain from guidance is the ratio of guided to unguided SperfS_perf. 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-Bench. The score uses an analytical estimate tSOL,i,wt_SOL,i,w of the minimum execution time of the PyTorch reference computation, derived from hardware capabilities. For an evaluated implementation with latency ti,wt_i,w, we compute a score for each workload: si,w=tb,i,w−tSOL,i,wti,w+tb,i,w−2tSOL,i,w.s_i,w= t_b,i,w-t_SOL,i,wt_i,w+t_b,i,w-2t_SOL,i,w. (3) The SOL score’s intended interpretation within [0,1][0,1] requires tb,i,w>tSOL,i,wt_b,i,w>t_SOL,i,w and ti,w≥tSOL,i,wt_i,w≥ t_SOL,i,w. However, D2K-Bench 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 Qwen3.8-Max-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 JSimplm,c,i∈[0,100]JS_impl^m,c,i∈[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 JSimplm,c,i=0JS_impl^m,c,i=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 JSdesignm,i∈[0,100]JS_design^m,i∈[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: JSdesignm=1||∑i∈JSdesignm,i,JSimplm,c=1||∑i∈JSimplm,c,i.JS_design^m= 1|T| _i JS_design^m,i, _impl^m,c= 1|T| _i JS_impl^m,c,i. (4) 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 SperfS_perf is a geometric 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 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 SperfS_perf weights models equally. PyTorch Baseline w/o Design Guidance w/ Design Guidance Expert Model SperfS_perf Fable-5 Correct SperfS_perf Correct SperfS_perf SperfS_perf GPT-6-Astra 0.0560.056 1.001.00 26 2.682.68 26 3.333.33 2.332.33 Claude-Opus-4.8 0.0560.056 1.001.00 26 1.381.38 26 2.312.31 2.332.33 GPT-5.6-Sol 0.0560.056 1.001.00 26 1.311.31 26 2.012.01 2.332.33 GLM-5.3 0.0560.056 1.001.00 21 1.181.18 25 1.341.34 2.332.33 Kimi-K3 0.0560.056 1.001.00 22 1.161.16 25 1.381.38 2.332.33 Overall 0.0560.056 1.001.00 121 1.461.46 128 1.951.95 2.332.33 Table 2: Alignment of LLM judge scores with measured kernel performance. All values aggregate 26 tasks per model. Averages are computed before rounding; JSdesignJS_design and JSimplJS_impl are rounded to integers. Proposal w/o Design Guidance w/ Design Guidance Model JSdesignJS_design L1 L2 L3 JSimplJS_impl SperfS_perf L1 L2 L3 JSimplJS_impl SperfS_perf GPT-6-Astra 7373 68.3868.38 68.8168.81 66.9666.96 6868 2.682.68 80.3880.38 79.5879.58 77.3877.38 7979 3.333.33 Claude-Opus-4.8 6767 65.0465.04 64.0464.04 60.4660.46 6363 1.381.38 72.0872.08 68.6568.65 67.9667.96 7070 2.312.31 GPT-5.6-Sol 7272 63.5063.50 63.5863.58 58.8858.88 6262 1.311.31 72.5072.50 72.8872.88 65.6965.69 7070 2.012.01 GLM-5.3 5959 43.4243.42 45.0845.08 46.4246.42 4545 1.181.18 67.6567.65 66.0066.00 61.1261.12 6565 1.341.34 Kimi-K3 6363 48.7348.73 48.2748.27 48.7748.77 4949 1.161.16 66.5066.50 64.6964.69 61.6961.69 6464 1.381.38 Average 6767 57.8257.82 57.9557.95 56.3056.30 5757 1.461.46 71.8271.82 70.3670.36 66.7766.77 7070 1.951.95 Comprehensive guidance improves both correctness and SperfS_perf in Table 1. Across 130 model–task pairs, correctness increases from 93.1% to 98.5%. Overall SperfS_perf increases from 1.46 to 1.95, a 33.9% improvement. All five models improve their SperfS_perf under guidance. The expert implementation achieves Sperf=2.33S_perf=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 SperfS_perf is the geometric mean speedup over all 26 tasks. Their geometric mean speedup increases from 1.69×1.69× to 2.49×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 SperfS_perf 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×1.69× to 2.49×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 JSdesignJS_design for independently proposed designs, JSimplJS_impl for final implementations, and the measured performance score SperfS_perf. 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, JSimplJS_impl follows the broad model ordering of SperfS_perf. 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. (i) Independent design scores also reflect broad performance differences. Higher JSdesignJS_design 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 Figure 3: Expert-relative performance under cumulative guidance. Results are available for all three models on all 26 tasks under each input configuration. 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. Values use Equations 1 and 2 with T equal to all 26 tasks and expert latency in the numerator, and are multiplied by 100. 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 token-efficiency results are reported in Appendix A.4. Summarized Answer to RQ3. Averaged across the three models, adding L1 yields the largest incremental gain, while L2 and L3 provide additional improvements. 5 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 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-Bench, 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] A. Ouyang, S. Guo, S. Arora, A. L. Zhang, W. Hu, C. Ré, and A. Mirhoseini (2025) KernelBench: can LLMs write efficient GPU kernels?. In Proceedings of the 42nd International Conference on Machine Learning, Proceedings of Machine Learning Research, Vol. 267, p. 47356–47415. External Links: Link Cited by: §A.3, §1, §1, §5. [2] S. K. S. Hari, V. Balaji, S. Damani, Q. Huang, and C. Kozyrakis (2026) Improving efficiency of GPU kernel optimization agents using a domain-specific language and speed-of-light guidance. arXiv preprint arXiv:2603.29010. External Links: Link Cited by: §1, §1, §5. [3] G. Oliaro, Y. Fu, M. Jiang, O. Lu, J. Wang, Z. Jia, H. Zhang, and S. Rajbhandari (2026) FastKernels: benchmarking GPU kernel generation in production. arXiv preprint arXiv:2605.23215. External Links: Link Cited by: §1, §1, §5. [4] T. Dao, D. Y. Fu, S. Ermon, A. Rudra, and C. Ré (2022) FlashAttention: fast and memory-efficient exact attention with IO-awareness. In Advances in Neural Information Processing Systems, Vol. 35, p. 16344–16359. External Links: Document, Link Cited by: §1, §2.1. [5] E. Lin, S. Modi, S. K. S. Hari, Q. Huang, Z. Ye, N. Qin, F. Zhou, Y. Zhang, J. Wang, S. Damani, D. Peri, O. Xie, A. Kane, M. Maor, M. Behar, T. Cao, R. Mehta, V. Singh, V. Sharma Mailthody, T. Chen, Z. Ye, H. Chen, T. Chen, V. Grover, W. Chen, W. Liu, E. Chung, L. Ceze, R. Bringmann, C. Zeller, M. Lightstone, C. Kozyrakis, and H. Shi (2026) SOL-ExecBench: speed-of-light benchmarking for real-world GPU kernels against hardware limits. arXiv preprint arXiv:2603.19173. External Links: Link Cited by: §1, §3.2, §5. [6] Y. Shui, C. Ma, H. Xu, S. Wen, and Y. Wang (2026) Harness engineering for LLM-driven GPU kernel generation. arXiv preprint arXiv:2607.17979. External Links: Link Cited by: §1, §5. [7] W. Kwon, Z. Li, S. Zhuang, Y. Sheng, L. Zheng, C. H. Yu, J. E. Gonzalez, H. Zhang, and I. Stoica (2023) Efficient memory management for large language model serving with PagedAttention. In Proceedings of the 29th Symposium on Operating Systems Principles, p. 611–626. External Links: Document, Link Cited by: §2.1. [8] L. Zheng, L. Yin, Z. Xie, C. Sun, J. Huang, C. H. Yu, S. Cao, C. Kozyrakis, I. Stoica, J. E. Gonzalez, C. Barrett, and Y. Sheng (2024) SGLang: efficient execution of structured language model programs. In Advances in Neural Information Processing Systems, Vol. 37, p. 62557–62583. External Links: Document, Link Cited by: §2.1. [9] Y. Yu, P. Zang, C. H. Tsai, H. Wu, Y. Shen, J. Zhang, H. Wang, Z. Xiao, J. Shi, Y. Luo, W. Zhang, C. Men, G. Liu, and Y. Lin (2026) Towards automated kernel generation in the era of LLMs. arXiv preprint arXiv:2601.15727. External Links: Link Cited by: §5. [10] J. Li, S. Li, Z. Gao, Q. Shi, Y. Li, Z. Wang, J. Huang, H. Wang, J. Wang, X. Han, Z. Liu, and M. Sun (2025) TritonBench: benchmarking large language model capabilities for generating Triton operators. In Findings of the Association for Computational Linguistics: ACL 2025, Vienna, Austria, p. 23053–23066. External Links: Document, Link Cited by: §5. [11] Z. Wen, Y. Zhang, Z. Li, Z. Liu, L. Xie, and T. Zhang (2025) MultiKernelBench: a multi-platform benchmark for kernel generation. arXiv preprint arXiv:2507.17773. External Links: Link Cited by: §5. [12] G. K. Kumar, Y. Perlitz, C. Lammie, A. Giovannini, and K. Hose (2026) DataKernelBench: can LLMs optimize database queries on GPUs?. arXiv preprint arXiv:2608.25061. External Links: Link Cited by: §5. [13] R. T. Lange, Q. Sun, A. Prasad, M. Faldor, Y. Tang, and D. Ha (2025) Towards robust agentic CUDA kernel benchmarking, verification, and optimization. arXiv preprint arXiv:2509.14279. External Links: Link Cited by: §5. [14] Y. Zhang, P. Yu, J. Wang, M. (. Fan, J. Reed, A. Mirhoseini, and W. Su (2026) KernelBench-Verified: do LLM-generated kernels actually beat PyTorch?. arXiv preprint arXiv:2607.16241. External Links: Link Cited by: §5. [15] H. Wang, J. Zhang, K. Jiang, H. Wang, J. Chen, and J. Zhu (2026) KernelBenchX: a comprehensive benchmark for evaluating LLM-generated GPU kernels. arXiv preprint arXiv:2605.04956. External Links: Link Cited by: §5. [16] S. Younesian, W. Ouyang, S. Rafati, M. Rezagholizadeh, S. Zhou, J. Liu, Y. Liu, Y. Yang, H. Li, Z. Liu, D. Li, V. Appia, Z. Gu, and E. Barsoum (2026) AgentKernelArena: generalization-aware benchmarking of GPU kernel optimization agents. arXiv preprint arXiv:2605.16819. External Links: Link Cited by: §A.3, §5. [17] S. Williams, A. Waterman, and D. Patterson (2009) Roofline: an insightful visual performance model for multicore architectures. Communications of the ACM 52 (4), p. 65–76. External Links: Document, Link Cited by: §5. [18] J. Huang, Z. Wen, T. Xu, M. Yan, X. Xia, and Z. Liu (2026) RealisticTritonBench: a benchmark for Triton-kernel generation in real-world AI frameworks. arXiv preprint arXiv:2608.12004. External Links: Link Cited by: §5. [19] S. Xing, Y. Zhai, A. Jiang, Y. Dong, Y. Wu, Z. Ye, C. F. Ruan, Y. Huang, Y. Zhang, L. Yin, A. Bayyapu, L. Ceze, and T. Chen (2026) FlashInfer-Bench: building the virtuous cycle for AI-driven LLM systems. In Proceedings of Machine Learning and Systems, Vol. 8, p. 2016–2064. External Links: Link Cited by: §5. [20] L. Yang, Y. Wang, S. Liang, L. Yang, D. Ying, C. You, R. Zhang, L. Wang, Y. Yu, G. Yang, and L. Zhang (2026) Are LLM-generated GPU kernels production-ready? a trace-driven benchmark and optimization agent. arXiv preprint arXiv:2607.14541. External Links: Link Cited by: §5. [21] J. Jaber and O. Jaber (2026) AutoKernel: autonomous GPU kernel optimization via iterative agent-driven search. arXiv preprint arXiv:2603.21331. External Links: Link Cited by: §5. [22] H. Li, K. Man, P. Kanuparthy, H. Chen, W. Sun, S. Tallam, C. Zhu, K. Zhu, and Z. Qian (2025) TritonForge: profiling-guided framework for automated Triton kernel optimization. arXiv preprint arXiv:2512.09196. External Links: Link Cited by: §5. [23] A. Wei, T. Sun, Y. Seenichamy, H. Song, A. Ouyang, A. Mirhoseini, K. Wang, and A. Aiken (2025) Astra: a multi-agent system for GPU kernel performance optimization. arXiv preprint arXiv:2509.07506. External Links: Link Cited by: §5. [24] Y. Han, M. Guo, Z. Liu, W. Chen, and S. Hu (2026) Making LLMs optimize multi-scenario CUDA kernels like experts. arXiv preprint arXiv:2603.07169. External Links: Link Cited by: §5. [25] Q. Sun, J. Han, T. Li, Z. Tang, S. Chen, F. Yang, A. Liu, X. Liu, and Y. Liu (2026) KernelSkill: a multi-agent framework for GPU kernel optimization. arXiv preprint arXiv:2603.10085. External Links: Link Cited by: §5. [26] G. Zhang, S. Zhu, A. Wei, Z. Song, A. Nie, Z. Jia, N. Vijaykumar, Y. Wang, and K. Olukotun (2026) AccelOpt: a self-improving LLM agentic system for AI accelerator kernel optimization. In Proceedings of Machine Learning and Systems, Vol. 8, p. 541–568. External Links: Link Cited by: §5. [27] S. Guo, M. Lin, and T. Yang (2026) DRTriton: large-scale synthetic data driven reinforcement learning for Triton kernel generation. arXiv preprint arXiv:2603.21465. External Links: Link Cited by: §5. [28] M. Spoczynski, D. Fleischer, M. Berchansky, G. B. Stan, S. Guskin, W. Xu, A. Siemieniuk, and A. Heinecke (2026) Xe-Forge: multi-stage LLM-powered kernel optimization for Intel GPU. arXiv preprint arXiv:2605.26118. External Links: Link Cited by: §5. [29] N. Wiedemann, Q. Leboutet, M. Paulitsch, D. Wofk, and B. Ummenhofer (2026) KernelFoundry: hardware-aware evolutionary GPU kernel optimization. arXiv preprint arXiv:2603.12440. External Links: Link Cited by: §A.3, §5. [30] X. Zhu, S. Peng, J. Guo, Y. Chen, Q. Guo, Y. Wen, H. Qin, R. Chen, Q. Zhou, K. Gao, Y. Wu, C. Zhao, and L. Li (2026) QiMeng-Kernel: macro-thinking micro-coding paradigm for LLM-based high-performance GPU kernel generation. In Proceedings of the AAAI Conference on Artificial Intelligence, Vol. 40(34), p. 29168–29176. External Links: Document, Link Cited by: §5. [31] J. Ragan-Kelley, C. Barnes, A. Adams, S. Paris, F. Durand, and S. Amarasinghe (2013) 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, p. 519–530. External Links: Document, Link Cited by: §5. [32] B. Hagedorn, A. S. Elliott, H. Barthels, R. Bodík, and V. Grover (2020) Fireiron: a data-movement-aware scheduling language for GPUs. In 2020 International Conference on Parallel Architectures and Compilation Techniques, p. 71–82. External Links: Document, Link Cited by: §5. [33] P. Tillet, H. T. Kung, and D. Cox (2019) 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, p. 10–19. External Links: Document, Link Cited by: §5. [34] L. Wang, Y. Cheng, Y. Shi, Z. Mo, Z. Tang, W. Xie, T. Wu, L. Ma, Y. Xia, J. Xue, F. Yang, and Z. Yang (2026) TileLang: bridge programmability and performance in modern neural kernels. In International Conference on Learning Representations, Vol. 2026, p. 73099–73134. External Links: Link Cited by: §5. Appendix A Supplementary Material A.1 Benchmark Composition Table 3: Operator categories and representative design content in D2K-Bench. GEMM denotes general matrix multiplication. Operator category Tasks Representative design content Attention and recurrent operators 7 Sparse attention, decoding with cached state, composable recurrences, and compressed state Mixture of experts and distributed execution 8 Grouped GEMM, routing, and distributed dispatch and combine Indexing, sampling, and control flow 6 Index sharing, exact sampling, compact decision state, and verification Quantization, training, and structural operators 5 Low-precision representation, clustering, optimizer updates, and fusion A.2 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,wt_b,i,w/t_m,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 SperfS_perf 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 Model Correct tasks Speedup Correct tasks Speedup GPT-6-Astra 26 2.68×2.68× 26 3.33×3.33× Claude-Opus-4.8 26 1.38×1.38× 26 2.31×2.31× GPT-5.6-Sol 26 1.31×1.31× 26 2.01×2.01× Expert 26 2.33×2.33× 26 2.33×2.33× Thresholded speedup coverage. Following the thresholded metric in KernelBench [1], we additionally report an adapted fastpfast_p relative to the same Fable-5 first-correct baseline used in our performance evaluation. For threshold p, let m,cC_m,c contain the tasks for which model m under configuration c produces a correct submission with valid latency measurements on every workload. We define fastp,m,cFable=1||∑i∈(i∈m,c∧Sperfm,c,i>p),||=26.fast^Fable_p,m,c= 1|T| _i 1\! (i _m,c\ \ S_perf^m,c,i>p ), |T|=26. (5) 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 fastp,m,cFablefast^Fable_p,m,c remains all 26 tasks regardless of correctness. We report fastp,m,cFablefast^Fable_p,m,c as a secondary thresholded summary; performance score SperfS_perf remains the main benchmark metric. KernelBench emphasizes fast1fast_1, while later evaluations commonly report fast1fast_1 and fast2fast_2 [29, 16]. We therefore report p∈1,1.5,2p∈\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 fast1Fablefast^Fable_1 fast1.5Fablefast^Fable_1.5 fast2Fablefast^Fable_2 GPT-6-Astra w/o guidance 96.2% 88.5% 57.7% w/ guidance 96.2% 92.3% 73.1% Claude-Opus-4.8 w/o guidance 61.5% 38.5% 23.1% w/ guidance 88.5% 65.4% 50.0% GPT-5.6-Sol w/o guidance 65.4% 38.5% 26.9% w/ guidance 88.5% 53.8% 34.6% GLM-5.3 w/o guidance 65.4% 38.5% 19.2% w/ guidance 61.5% 46.2% 34.6% Kimi-K3 w/o guidance 65.4% 26.9% 11.5% w/ guidance 73.1% 50.0% 30.8% Overall w/o guidance 70.8% 46.2% 27.7% w/ guidance 81.5% 61.5% 44.6% 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.5p=1.5, the guided results also preserve the performance score SperfS_perf ordering: GPT-6-Astra, Claude-Opus-4.8, GPT-5.6-Sol, Kimi-K3, and GLM-5.3. A.4 Token Efficiency Figure 4: Claude-Opus-4.8 optimization progress over the same 26 tasks. The performance score Sperf′S _perf uses a 0.10.1 fallback. 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.312.31, compared with 1.381.38 for the unguided curve, under the same 350-turn budget. We use Sperf′S _perf, adapting Equation 1 by replacing the 0.50.5 fallback with 0.10.1 and using the geometric mean over all 26 tasks in Equation 2. The 0.10.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.10.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.11 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. 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 Uncached input tokens (M) 0.041 0.108 0.149 Cache-write tokens (M) 37.261 48.084 85.345 Cache-read tokens (M) 402.034 633.801 1,035.835 Output tokens (M) 7.898 8.085 15.983 API cost (US$) 631.56 820.10 1,451.66 Guidance improves final SperfS_perf while increasing token usage and estimated API cost. Appendix 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∈ R^[…]× M× K\), normalize independently in fp32, store the iterate in fp16, and apply the five coefficient triples below in order. \[ X_0=fp16\! ( fp32(G) \|fp32(G)\|_F+10^-7 ), \] \[ A_i=X_iX_i , X_i+1=a_iX_i+ (b_iA_i+c_iA_i^2 )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. ## 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 floating-point 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" ], "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 = ( (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.", "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 pre-restart 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=X^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 Gram-like 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 composition-commit 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 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=X^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" ], "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". - **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" ], "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" , "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" , "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" , "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 Appendix 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. Brace-delimited fields are filled by the release scorer before transmission. Both assessments assign continuous scores in [0,100][0,100] at L1, L2, and L3 and use the equal arithmetic mean of the three level scores. The corresponding scores are JSdesignJS_design and JSimplJS_impl, 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. 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" ] 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 100 -- every required design property is complete and coherent substantial_design 75 -- all main properties are present; only secondary detail is missing partial_design 50 -- a coherent direction, but at least one required property is missing isolated_fragment 25 -- isolated relevant ideas without the end-to-end consequence absent 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" ] 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× 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.260.26 for the row-to-column ratio, which is 0.370.37–0.390.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) _a_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, 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× 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=39JS_design=39. Table 7: Muon implementation judge scores and performance for Claude-Opus-4.8 on this task. SperfS_perf is the geometric mean speedup over the first-correct baseline across the three Muon workloads. Input configuration L1 L2 L3 JSimplJS_impl SperfS_perf w/o design guidance 40 54 62 52 7.22 w/ design guidance 62 70 86 73 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=52JS_impl=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=73JS_impl=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.