ForgeKernel: Self-Evolving Kernels Catching and Surpassing Experts and Agents

arXiv:2611.04998v1 [cs.DC] 3 Nov 2026

ForgeKernel: Self-Evolving Kernels Catching and Surpassing Experts and Agents Yaojian Chen∗

Dazhi Peng∗

Yiyun Zheng∗

Tsinghua University Beijing, China yj-chen21@mails.tsinghua.edu.cn

ModelBest Inc. Beijing, China pengdazhi@modelbest.cn

ModelBest Inc. Beijing, China zhengyiyun@modelbest.cn

Zhui Zhu

Shengye Xiang

Leshan Li

Tsinghua University Beijing, China z-zhu22@mails.tsinghua.edu.cn

Tsinghua University Beijing, China xcy23@mails.tsinghua.edu.cn

Tsinghua University Beijing, China lils25@mails.tsinghua.edu.cn

Zhen Li

Xianglong Deng

Haojun Sun

ModelBest Inc. Beijing, China lizhen@modelbest.cn

University of Chinese Academy of Sciences Beijing, China dengxianglong02@gmail.com

ModelBest Inc. Beijing, China sunhaojun@modelbest.cn

Qingfeng He

Shangzhan Li

Yifei Shen

Tsinghua University Beijing, China heqf21@mails.tsinghua.edu.cn

Harbin Institute of Technology Harbin, China szli@ir.hit.edu.cn

ModelBest Inc. Beijing, China shenyifei@modelbest.cn

Lin Gan

Yuxuan Li†

Xu Han†

Tsinghua University Beijing, China lingan@tsinghua.edu.cn

ModelBest Inc. Beijing, China yxuanl1995@gmail.com

Tsinghua University Beijing, China han-xu@tsinghua.edu.cn

Zhiyuan Liu† Tsinghua University Beijing, China liuzy@tsinghua.edu.cn

Abstract

We present ForgeKernel, a kernel-generation system for continued optimization. It is designed to accumulate and reuse knowledge across runs while enabling continued evolution of kernel architectures. It combines two mechanisms. First, it analyzes the hardware behavior of a high-performance reference and distills its validated architectural decisions into reusable knowledge of hardware mechanisms rather than imitating its source code. Second, it self-evolves by sealing a kernel that passes correctness and performance validation and using it as the reference for the next optimization run, allowing the optimization results of one run to inform subsequent runs and providing a basis for revisiting and adjusting the kernel architecture. Across 27 operators and 29 input configurations on NVIDIA Hopper and Blackwell GPUs, ForgeKernel matches or beats its expert reference on 25 of them, for a +10% geometric mean. On Ascend 910C, on the public case set of all 8 level-4 operators in the CANN-Bench Season 1 final

Purpose-built kernel agents can now write high-performance GPU kernels. As base models grow stronger, though, even a naive agentic loop beats human experts on many operators, so both the purpose of a kernel agent and the way it should be designed need rethinking. Yet existing kernel agents generally lack cross-run knowledge accumulation and reuse, as well as continued evolution at the architectural level: each run starts from the same initial knowledge and merely outputs the best-performing kernel found during that search, without retaining the knowledge gained through optimization for subsequent runs; meanwhile, once the kernel architecture is fixed early, later iterations are largely confined to local optimization within it. ∗ The first three authors contributed equally to this work. † Corresponding authors.

Preprint, arXiv 2026. 1

Yaojian Chen, Dazhi Peng, Yiyun Zheng, Zhui Zhu, Shengye Xiang, Leshan Li, Zhen Li, Xianglong Deng, Haojun Sun, Qingfeng He, Shangzhan Li, Yifei Shen, Preprint, 2026, arXiv Lin Gan, Yuxuan Li, Xu Han, and Zhiyuan Liu

leaderboard, ForgeKernel leads the per-operator public best, with a +25.1% geometric-mean lead, while its full-suite submission ranks first on that benchmark’s A3 leaderboard. As of 2026-10-30, public submissions generated by ForgeKernel rank first on all three CANN-Bench 2.0 solution leaderboards (Ascend 910C, 950DT, and 950PR) and on the SOL-ExecBench FlashInfer-Bench track.

(a) One goal-mode run

Keywords: GPU kernels, code generation, large language models, self-evolution, performance optimization

hardware ceiling, unknown

× the architecture settles early local tuning fills the rest of the run and cannot change it

Introduction

pretraining

“kernel complete”

Knowledge is discarded

promote

Modern machine-learning systems increasingly demand high-performance GPU kernels. New operators and model architectures, including sparse attention indexers [8], gated delta networks [28], and latent-attention variants [7], continue to appear and further expand the computational patterns that require optimization. Each such pattern is a new optimization task: work partitioning, dataflow, and pipeline organization have to be determined again for the concrete operator, input shape, and hardware. Faced with this continuing demand for kernel development, agents that automatically write and optimize operators have drawn attention. Prior work has built agents dedicated to kernel optimization [4, 9, 10, 14, 25, 26], organizing generation, evaluation, and iterative improvement into a dedicated workflow, and has surpassed expert implementations on some tasks. These workflows still have two structural limits. On tasks that fall short of the expert reference, optimization traces show that architectural decisions, such as work partitioning, dataflow, and pipeline organization, are often fixed in the early attempts. Subsequent iterations tune within that architecture and rarely revise these decisions from performance feedback. In addition, producing a highperformance kernel in a single run does not mean that its validated findings have become cross-run knowledge that can guide later architectural design and implementation. Overall, an operator-specific agent should close these two limits. First, performance feedback should revise architectural decisions, rather than leaving optimization inside an early choice of work partitioning, dataflow, and pipeline organization. Second, validated findings should become crossrun knowledge that guides later architectural design and implementation. Guided by these design objectives, we introduce ForgeKernel, a self-evolving operator agent system (Figure 1). To address these two problems, ForgeKernel starts optimization from architectural design rather than existing code, and uses a hardware knowledge base to retain and reuse measurementvalidated optimization experience. For architectural design, ForgeKernel uses a highperformance reference implementation as a design reference and starting point. It separates architecture from coding, with a dedicated architecture agent reviewing and adjusting

gen 0 = sealed ref. gen 1 gen 2 roofline ✓ each accepted kernel is the next reference halts when a generation’s gain falls below the noise floor

run 1 run 2 run 3

× a run leaves a kernel, not knowledge every generation pays again for what the last one learned

forge prescribe / harvest

1

Architecture settles

discarded

device time, faster ↓

the architecture is fixed early, nothing outlives the run

op A op B op C õ knowledge base

✓ each generation starts where the last ended findings organized by GPU, shared across operators

(b) Self-evolving operator agent both loops continue; the hardware stops them, not the run

Figure 1. Single-run optimization and cross-run selfevolution. architectural decisions so that optimization is not confined to code-level tuning within a fixed architecture. Building on this separation, self-evolution enables continual reflection on the architecture as a whole and iterative improvement, reducing the risk of becoming trapped in local optima. For knowledge accumulation, ForgeKernel distills mechanisms discovered during optimization and validated through independent experiments on the target hardware into crossrun knowledge that guides architectural design and implementation in subsequent tasks. Reference updates provide new performance baselines, while accumulated knowledge provides reusable evidence for optimization decisions; together, they support continued optimization. We evaluate ForgeKernel on NVIDIA Hopper and Blackwell GPUs across 29 test configurations covering 27 operators. ForgeKernel matches or exceeds reference performance on 25 configurations, achieving a geometric-mean speedup of 1.10× over reference implementations from expert libraries and other kernel-generation systems. On Ascend, ForgeKernel outperforms the per-operator public best on the public case set of all 8 level-4 operators in the CANNBench Season 1 final leaderboard. In experiments comparing against kernels generated by a general-purpose coding agent, ForgeKernel achieves performance gains of up to 9.5%. In addition, knowledge-provenance records demonstrate reuse of 2

ForgeKernel: Self-Evolving Kernels

Preprint, 2026, arXiv

validated optimization knowledge across runs and operators, while knowledge-base ablations show that this knowledge helps improve generated kernel performance on some evaluated tasks. As external validation beyond these controlled experiments, as of 2026-10-30, public submissions generated by ForgeKernel rank first on all four leaderboards: the three CANN-Bench 2.0 solution leaderboards (Ascend 910C, 950DT, and 950PR) and the SOL-ExecBench FlashInfer-Bench track [2, 20]. The ForgeKernel source code is publicly available at https: //github.com/OpenBMB/ForgeKernel. This paper makes the following contributions:

memory access overhead, and hardware resource utilization, so different implementations of the same operator can exhibit substantial performance differences. Kernel optimization involves both revising these structural design choices and refining instructions, synchronization operations, and resource configurations within a given structure. When performance is limited by work partitioning, data layout, or pipeline organization, local parameter tuning alone may be insufficient to eliminate the bottleneck, requiring the corresponding design choices to be reconsidered. Kernel optimization therefore involves not only parameter search but also analysis and modification of the implementation structure.

• We propose a kernel optimization method centered on architectural decisions. It separates architectural design from code implementation, guides implementation through explicit architectural decisions, and revises these decisions based on performance feedback. When optimization encounters a structural bottleneck, the method revisits work partitioning, dataflow, or pipeline organization rather than only tuning locally within a fixed architecture. Across runs, it also reassesses an existing implementation’s architecture rather than directly carrying it forward. • We propose a measurement-validated mechanism for accumulating and reusing knowledge across runs. Hardware mechanisms discovered during optimization are validated through independent probes on the target GPU before being stored in a knowledge base shared across operators and input shapes. This knowledge guides architectural design and implementation planning in subsequent runs, converting findings from individual runs into reusable knowledge. • We implement ForgeKernel, a self-evolving operator optimization system that integrates architectural revision, reference updates, and knowledge reuse. Evaluation on NVIDIA Hopper and Blackwell GPUs and Ascend demonstrates performance gains over expert and generated references, including a 1.10× geometricmean speedup on NVIDIA GPUs. As of 2026-10-30, submissions generated by ForgeKernel rank first on all four leaderboards: the three CANN-Bench 2.0 solution leaderboards (Ascend 910C, 950DT, and 950PR) and the SOL-ExecBench FlashInfer-Bench track.

2

Background

2.1

Kernel Design

2.2

Kernel Agent

As large language models (LLMs) improve in code generation and reasoning, researchers have begun using specialized LLM-based agents to write and optimize kernels. These systems are referred to as kernel agents. Given an operator, an input configuration, and a target GPU, a kernel agent generates candidate implementations, compiles them and checks their correctness, measures execution performance, and iteratively revises the code based on feedback to search for a functionally correct and high-performance kernel.

3

Design

ForgeKernel separates architectural planning from kernel implementation, coordinating an architecture agent and a coding agent through design plans and measurement feedback. Figure 2 shows the system architecture. 3.1

Overview

Kernel optimization requires both structural design and detailed implementation. Structural design determines how computation and data movement are organized on the target hardware; implementation turns those choices into correct, executable code. Combining both responsibilities in one agent risks allowing compilation, debugging, and local edits to dominate the design context and anchoring further optimization to the current implementation. ForgeKernel separates these responsibilities so that architectural choices can be reconsidered from measurement evidence without making the existing code structure the starting point for every revision. The architecture agent owns design decisions and plan revisions. Given the operator definition, input configuration, target hardware, and reference evidence, it determines computation organization, dataflow, scheduling, and hardware mapping, and specifies what evidence would validate these choices. Its output is a milestone plan. The coding agent owns realization of that plan: it implements the required computation, data access, and synchronization, handles compilation and debugging, and performs local optimization within the

GPU kernel design requires mapping an operator’s computation onto the parallel execution and memory structures of the target hardware. For a given operator and input configuration, designers must determine how to partition computation across thread blocks and threads, lay out data across memory levels, and pipeline data transfers with computation [17–19, 21, 23]. These choices jointly affect data reuse, 3

Yaojian Chen, Dazhi Peng, Yiyun Zheng, Zhui Zhu, Shengye Xiang, Leshan Li, Zhen Li, Xianglong Deng, Haojun Sun, Qingfeng He, Shangzhan Li, Yifei Shen, Preprint, 2026, arXiv Lin Gan, Yuxuan Li, Xu Han, and Zhiyuan Liu

> Reference kernel

analyse: NCU / timeline probes

¨ ARCH AGENT certified recipes

Reference analysis evidence

Milestone construction

õ Knowledge base (hw_kb)

Ð CODING AGENT

M0

M1

M2

Implement decision row

...

M0: architecture decisions

fail

1

harvest: probe-certified

candidate, timed vs ref in one GPU job pass

Milestone gate

3

Ç Shipped kernel

plateau → re-analyse + re-distil 2

a shipped kernel becomes the next optimization run’s reference

Figure 2. ForgeKernel architecture. plan’s constraints. Changes to those architectural constraints are handled through plan revision by the architecture agent. The interface between the agents is a design plan in one direction and implementation feedback with measurement evidence in the other. The coding agent receives design requirements, interfaces, and acceptance conditions, but not the reference kernel implementation. The execution harness enforces this boundary, checks correctness, measures performance, and validates the mechanisms required by the plan. These checks determine whether implementation can advance or needs further repair; when optimization stalls or the plan cannot proceed, the trajectory and measurements return to the architecture agent for reassessment. This separation assigns architectural reasoning, code realization, and measured acceptance to distinct parts of the system while retaining a feedback loop between them. Sections 3.2 and 3.3 describe the two agents in detail. Selfevolution extends the loop through plan revision and reference updates (§3.4), while the hardware knowledge base retains validated mechanisms for reuse (§3.5). 3.2

layouts. Together, these choices specify how computation and data movement interact. Architectural decisions must state the design choices, supporting evidence, and expected effects, distinguishing direct observations from hypotheses that remain to be tested. The reference implementation’s measured performance determines relative performance targets, while its resource usage and execution characteristics inform subsequent mechanism validation. Knowledge entries used in the design are explicitly cited in the plan.

Milestone planning. The architecture agent organizes architectural decisions and subsequent optimization tasks into an ordered milestone sequence 𝑀0, 𝑀1, . . . , 𝑀𝑘 . The initial milestone 𝑀0 requires the implementation of the key structural decisions above to produce a verifiable initial implementation; subsequent milestones perform fine-grained tuning within that structure. Each milestone specifies implementation tasks, decision rationales, expected costs, and prerequisite dependencies, together with correctness, performance, and mechanismvalidation conditions. Dependencies between decisions determine the implementation order, and decisions that must be implemented together to take effect form a single implementation unit. In particular, 𝑀0 focuses on correctness and architectural constraints without imposing an absolute performance threshold; subsequent milestones are assessed against their respective optimization objectives. Milestones describe tasks through design requirements and interface contracts, without including the operator’s implementation code.

Architecture Agent

The architecture agent combines operator requirements, reference measurements, and existing hardware knowledge to formulate architectural decisions and milestone plans. Feedback-driven design revision is described in §3.4. Architectural decisions. The architecture agent first determines the candidate implementation’s computation organization, dataflow, and mapping to the target hardware. Decisions cover kernel decomposition, data dependencies, and intermediate-result storage, as well as computational tiling, thread roles, register budgets, and shared-memory 4

ForgeKernel: Self-Evolving Kernels

3.3

Preprint, 2026, arXiv

Coding Agent

without requiring a new external expert implementation. This implementation participates in the subsequent run as a measurable black box, with its source code unavailable to that run’s agents. The architecture agent analyzes the updated reference and uses it to formulate architectural decisions and performance targets. The coding agent generates a new implementation according to the corresponding milestone plan, while validated knowledge is retained and reused through the knowledge base. Cross-run iteration terminates when the performance gain in a subsequent run is confirmed not to exceed the measurement noise threshold, or when the budget is exhausted.

The coding agent translates design requirements into an operator implementation according to the architecture agent’s milestone plan. It follows decision dependencies and consults the knowledge entries cited in the plan to implement the computation, data access, and synchronization required by the current milestone. The harness checks candidate correctness, performs paired timing of the candidate and its declared reference within the same measurement job when that reference can run alongside the candidate, and collects the execution and resource characteristics needed for mechanism validation. Reference-free black-box arms use separate post-stop confirmation channels instead. The harness determines whether a candidate passes acceptance based on correctness-test results and the milestone’s performance and mechanism-validation conditions. For acceptance items that require a performance improvement, the measurements described above determine whether the gain exceeds the measurement noise threshold. Mechanism validation checks the execution or resource characteristics specified in the plan to determine whether the implementation satisfies the corresponding design constraints and exhibits the expected behavioral changes. After acceptance, the coding agent advances to the next milestone; otherwise, it continues modifying the implementation based on the check results. If optimization stagnates or acceptance cannot proceed, the system returns the implementation trajectory and measurement evidence to the architecture agent for the revision process described in §3.4. 3.4

3.5

Knowledge Validation and Reuse

The hardware knowledge base stores hardware mechanisms distilled during optimization and diagnosis, organized by target hardware architecture. Each entry records the mechanism, applicability conditions, trigger symptoms, a minimal runnable example, and known limitations. A candidate mechanism is admitted after passing independent probe validation on the target hardware and providing supporting measurement evidence. During subsequent optimization, the architecture agent selects relevant knowledge based on the task’s measured characteristics and the entries’ applicability conditions, and explicitly cites the entries in milestones for the coding agent to implement the corresponding mechanisms. Milestone acceptance still checks their effects in the current task. The system records the creation and use of entries to support traceability and reuse across runs.

Self-Evolution

Self-evolution encompasses plan revision within an optimization run and reference updates across runs. When optimization stagnates, the executable optimization decisions in the current plan are exhausted without reaching the target, or milestone acceptance cannot proceed, the system triggers reassessment of the current design by the architecture agent. The architecture agent combines implementation trajectories, measurement evidence, and supplementary probe experiments to revise architectural decisions and the corresponding milestone plan. The revised plan is re-executed according to decision dependencies. Local adjustments resume from the relevant tasks; structural changes to computation organization or dataflow restart implementation from the earliest affected decision, including 𝑀0 when necessary. Subsequent implementation and measurement evaluate each revision. If no performance improvement can be attributed to the change, the change is reverted and its failure conditions are recorded to inform subsequent diagnosis. Beyond revising the current plan, the system can use an accepted implementation as the reference for the next optimization run, extending optimization through its own outputs

4

Evaluation

This chapter evaluates the performance of ForgeKernelgenerated implementations and the roles of architecture planning, feedback-driven revision, and knowledge reuse. §4.1 describes the setup; §4.2 reports public-benchmark results and comparisons with expert libraries and artifacts from other generation systems; §4.3 evaluates self-evolution; §4.4 examines the individual designs; and §4.5 uses GQA and DSA cases to examine the roles of the architecture agent and self-evolution in architecture changes. 4.1

Experimental Setup

Hardware and tasks. NVIDIA experiments run on Hopper H100 80 GB (SXM, 132 SMs) and Blackwell B200 (148 SMs) nodes in a shared cluster, covering attention, matrix multiplication, convolution, GDN, and DSA. Ascend experiments run on 910C; the per-operator comparison covers all 8 level-4 operators in CANN-Bench 1.1.2. Coverage is reported separately for public leaderboards and per-operator experiments. Input shapes and data types are listed in Appendix C. 5

Yaojian Chen, Dazhi Peng, Yiyun Zheng, Zhui Zhu, Shengye Xiang, Leshan Li, Zhen Li, Xianglong Deng, Haojun Sun, Qingfeng He, Shangzhan Li, Yifei Shen, Preprint, 2026, arXiv Lin Gan, Yuxuan Li, Xu Han, and Zhiyuan Liu

Comparison targets. Per-operator experiments use expert library implementations and kernels released by other generation systems as references. Libraries include FlashAttention4, DeepGEMM, cuBLAS, cuDNN, FlashInfer, FlashMLA, and FlashQLA; generation systems include K-Search, Kernel Design Agents, and CAKE. Reference versions are pinned in the execution environment image and re-timed in our environment. Comparisons with generation systems evaluate their released artifacts. Versions and entry points are listed in Appendix B. Models and run protocol. The arch agent and coding agent of ForgeKernel (§3.1) use Claude Fable 5.1 and Claude Opus 5, respectively. By default, ForgeKernel receives a sealed reference that can be run, timed, profiled, and disassembled, without source access. The knowledge base accumulates across runs, so its initial state can differ between runs. Performance targets and budgets are set per task; runs end when the target is reached, the budget is exhausted, or the run is stopped manually. Expert-library experiments target parity, while self-evolution typically targets +5% over the preceding generation; specific targets and stopping reasons follow the individual run records. Main performance results select the best optimization run per configuration. The full system and each variant in the module ablation are run separately for that experiment. Correctness validation. Candidates must pass the official or benchmark tolerance test associated with their operator source before performance measurement. Operators covered by FlashInfer-Bench, including the operators in the SOL-ExecBench FlashInfer-Bench collection, use FlashInferBench’s tolerance tests; FlashAttention operators use the official FlashAttention tolerance tests; other operators likewise use the tolerance criterion of their source benchmark or implementation. The harness also rejects candidates that call the reference implementation directly or delegate computation to it. The source-specific criteria are summarized in Appendix A. Timing method. On NVIDIA hardware, we measure device execution time under cold-𝐿2 cache conditions. Average device time per call is the total GPU time over repeated calls divided by the call count. The candidate and reference use the same timing protocol and cache conditions. Results are based on repeated measurements. Speedup is 𝑆 = 𝑇ref /𝑇candidate , where the numerator and denominator are the device execution times of the reference and candidate, respectively; performance gain is (𝑆 − 1) × 100%. Ascend experiments use the hosted evaluation harness of CANN-Bench, with device execution time recorded in msprof’s kernel_details. Baselines, hardware performance bounds, and scoring rules follow the corresponding benchmark version. Scores are computed by the hosted service; results are aggregated separately for each platform. The /goal configuration. The /goal agent follows the goaldriven setting of JustGoal [5], using Claude Fable 5.1 without

an operator-specific workflow and with a 24-hour budget per task. Prompts use a fixed format specifying the operator signature, input shapes, target hardware (one H100), and optimization goal; copying code from the web is prohibited. No existing implementation is supplied as a reference. Generated implementations are subsequently evaluated against expert libraries. The objective has no preset speedup ceiling. Optimization state is not retained across sessions, and acquired knowledge is not transferred between runs. Selected artifacts serve as initial references for some self-evolution experiments in §4.3. 4.2

Performance Comparison

Overall performance on public benchmarks. As of 2026-10-30, public submissions generated by ForgeKernel rank first on all three CANN-Bench 2.0 solution leaderboards— Ascend 910C, 950DT, and 950PR—and on the SOL-ExecBench FlashInfer-Bench collection leaderboard [2, 20]. Table 1 reports coverage, speedups, and ranks. Speedups use each benchmark’s own baseline and aggregation rule. Ranks follow the official scores reported below. For the detailed historical result retained in this paper, on the CANN-Bench Season 1 A3 (Ascend 910C) 1.1 leaderboard, ForgeKernel covers all 53 operators across levels 1–4, including elementwise operations, normalization, matrix multiplication, convolution, attention, and fused operators. It passes all 1,060 standard-set cases and 4,185 of 4,240 hidden-set cases. The two sets score 4362.7547 and 4203.3563, respectively, yielding a first-place total of 8566.1110. Geometricmean speedups over the benchmark reference are 1.759× on the standard set and 1.854× on the hidden set; the leaderboard reports a combined baseline speedup of 1.835×, computed as their weighted geometric mean with weights of 20% and 80%, respectively [2]. On SOL-ExecBench, the FlashInfer-Bench collection covers normalization, matrix multiplication, GQA/MLA attention, and MoE kernels. With evaluation stack v1.1 on NVIDIA B200, ForgeKernel solves all 26 kernels, spanning 660 workloads, and achieves a SOL Score of 0.815204. Beyond this first-place score, Fast1 = 647/660, or 98.03%, shows that the generated implementations outperform the scoring baseline on almost all workloads. The leaderboard reports an Avg Speedup of 5.54× over that baseline, providing a complementary measure of acceleration magnitude rather than only the fraction of workloads improved [20]. Here, Fast1 is the fraction of workloads strictly faster than the baseline, aggregated across kernels. Operator-level performance comparison. Beyond the overall public-benchmark results, we evaluate representative operators to compare ForgeKernel with expert-library implementations, artifacts from existing generation systems, and per-operator public-best submissions. 6

ForgeKernel: Self-Evolving Kernels

Preprint, 2026, arXiv

Table 1. Public-benchmark results. Benchmark

Coverage

Hopper H100

Attention

Speedup Rank

CANN-Bench v1.1 Ascend 910C

53 operators Standard: 1,060/1,060 Hidden: 4,185/4,240

1.84×

SOL-ExecBench NVIDIA B200

26/26 kernels 660 workloads

5.54×

GQA fwd (16:1)

1

+0.3%

MHA bwd

+5.1%

GQA bwd (8:1)

+2.7%

GEMM

1

FP8 GEMM

+0.8%

FP8 grouped GEMM

+0.3%

FP16 GEMM

Across 22 configurations on Hopper and Blackwell, ForgeKernel matches or exceeds the expert implementation in 18, with another 3 within 1% of the reference. The geometric-mean performance gain across all configurations is +1.6% (Figure 3). Each bar in Figures 3 and 4 represents the best optimization run for an operator, input shape, and hardware configuration, using the measurement protocol in §4.1. The vertical zero line marks parity; the Hopper and Blackwell panels in Figure 3 share a scale. Beyond expert libraries, ForgeKernel matches or exceeds all 7 published kernels from K-Search [3], Kernel Design Agents [13], and CAKE [30], with a geometric-mean gain of +39.4% (Figure 4). Hopper references are K-Search’s FlashInfer-Bench and TriMul implementations; Blackwell references are KDA-v0.5 [32] and CAKE implementations. We further evaluate ForgeKernel on Ascend 910C, using the public case sets of all 8 level-4 operators in the CANN-Bench Season 1 final leaderboard. The ForgeKernelgenerated implementations outperform the per-operator public best, with a geometric-mean performance gain of +25.1% (Figure 5). Each bar reports the percentage performance gain over the per-operator public best, computed from the ratio of their speedups against the same baseline. The zero line marks parity; the final row reports the geometricmean speedup converted to a percentage gain. Public-best submissions come from the same final leaderboard and are not optimization inputs. The benchmark evaluates 20 public cases per operator, weighting compilation, accuracy, and hardware-attainable performance by 0.2, 0.3, and 0.5, respectively. 4.3

+0.9%

Convolution Conv2D fwd

− 0.5%

Conv2D weight grad

+1.0%

Conv2D input grad

+0.7%

Conv3D fwd

+7.6%

DSA chain DSA indexer fwd

+2.9%

DSA indexer top-k DSA sparse attn fwd

+2.2%

− 0.5%

DSA sparse attn bwd

+1.4%

DSA score recompute

+6.2%

DSA indexer bwd

+2.5%

GDN GDN prefill fwd

− 0.2%

GDN prefill bwd

+1.1% Blackwell B200

Attention MHA bwd MLA prefill (sparse)

+2.5%

− 1.3%

Attn + RoPE + cast

+0.0%

GEMM FP4 grouped GEMM

+0.6%

−4%

0%

4%

8%

Gain over the sealed expert library

Figure 3. Performance gains over expert libraries on Hopper and Blackwell.

Self-Evolution Experiments

This section evaluates reference updates across runs: an existing implementation becomes the next run’s reference. In the default ForgeKernel protocol, this reference is measurable but source-hidden, and the arch agent analyzes it to develop a new plan for the coding agent to implement. The DeepSeekV4.1 case below uses a source-visible generated reference and a single coding agent, as specified in its case description. We evaluate two paths. In the first, an initial implementation is generated without any existing implementation as a reference; ForgeKernel then uses that artifact as its reference for further optimization. In the second, ForgeKernel first uses

an expert-library implementation as its reference to generate generation 0, then uses the preceding generation as the reference for self-evolution. Generation 0 is the output of the first optimization run. Table 2 reports initial performance, subsequent round gains, and final cumulative gains. Initial performance is expressed as a percentage gain relative to the expert library; final gains use the same expert baseline. Round 1 is relative to the initial implementation, and round 2 to the round 1 output. Final gains for the DeepSeek-V4.1 case are measured directly against FlashMLA; the remaining final cumulative 7

Yaojian Chen, Dazhi Peng, Yiyun Zheng, Zhui Zhu, Shengye Xiang, Leshan Li, Zhen Li, Xianglong Deng, Haojun Sun, Qingfeng He, Shangzhan Li, Yifei Shen, Preprint, 2026, arXiv Lin Gan, Yuxuan Li, Xu Han, and Zhiyuan Liu Hopper H100

of the GQA experiment beginning with reference-free generation.

Ref.

GQA decode (paged)

K-Search

MLA decode (paged)

K-Search

MLA prefill (paged)

K-Search

TriMul

K-Search

Blackwell B200

Ref.

DSA sparse attn (paged)

KDA

GDN prefill (varlen)

KDA

KDA prefill (recurrent)

CAKE

+24.5%

Self-evolution on a complex fused operator. We examine self-evolution on a complex fused operator using DeepSeek-V4.1 fused sparse-attention prefill. Its fused structure provides a demanding test of ForgeKernel’s capability. The experiment uses a single coding agent (Claude Opus 5.5, high effort) for approximately 24 hours per generation on B200. The workload fixes 𝑠𝑞 = 4096, ℎ𝑞 = 128, 𝑠 kv = 32768, and topk = 640, with an attention sink. FlashMLA is used for evaluation only; the agent cannot inspect its source or run it during generation and optimization. The first generation is produced in /goal mode. The next generation can read the first generation’s source, but not its optimization notes. The initial gain of +5.3% over FlashMLA is the median of three independent same-card retests. The firstround gain +10.0% and final gain +15.5% are the separate medians of four acceptance measurements of the final implementation against the initial implementation and FlashMLA, respectively. All four final acceptances pass the five official correctness cases and the in-place-overwrite re-check. The task-specific benchmark reports cold-𝐿2 device time. The operator combines Q RoPE, sparse MQA attention with an attention sink, conjugate output RoPE, and FP8 casting. The CUDA C++ implementation fuses these operations into one kernel on B200. Its architecture uses a persistent pair of CTAs to process each query token, cooperatively gathers KV data, and assigns separate warp roles to data movement, matrix multiplication, softmax, and the epilogue. Specific architecture changes during cross-run optimization and their decision processes are examined in §4.5.

+170.6% +146.1% +8.7%

+7.0% +5.8% +0.4%

0%

50%

100%

150%

Gain over the sealed system kernel

Figure 4. Performance gains over published kernels from other generation systems. Ascend 910C GQA

+41.8%

GMM-SwiGLU-Quant

+32.2% +44.8%

GRU +9.8%

LSTM

+26.2%

MHA MLA

+12.4%

MLA-Prolog

+10.1%

SparseFA

+28.5%

+25.1%

Geo. mean

0%

20%

40%

60%

Gain over the public best

4.4 Figure 5. Performance gains over public-best submissions on Ascend 910C.

Ablation Study

This section examines architecture planning and revision, disassembly, and the knowledge base. We evaluate within-run plan revision separately from cross-run reference updates, whose results appear in §4.3. The experimental conditions and results follow.

gains are estimated by compounding the initial and subsequent speedups. A dash denotes a round that was not conducted. For the path beginning with reference-free generation, starting from the /goal implementations of GDN prefill forward and GQA attention forward, ForgeKernel obtains further gains of +9.5% and +5.0%, respectively. Combined with the initial implementations’ performance relative to expert libraries, the cumulative gains are approximately +100% and +29%. For the path beginning with an expert-library reference, the initial optimization is followed by one or two selfevolution rounds. GQA attention forward involves three ForgeKernel runs in total: the initial run produces generation 0, and the next two produce generations 1 and 2, with cumulative gains over the expert library of approximately +10% and +15%, respectively. This sequence is independent

Architecture planning and revision. Table 3 compares the full system with configurations that individually remove the milestone template, plan revision, SASS disassembly, or the knowledge base. Each cell reports the performance gain over the expert reference from one optimization run. Retaining only the initial plan without revision reduces performance on all three tasks relative to the full system, supporting the role of feedback-driven revision on these tasks. Removing the milestone template substantially reduces GDN prefill-forward performance, with a smaller change for DSA sparse attention forward. The attention-backward implementation without the template passes output-error checks but lowers internal precision (marked † ). The DSA run without the knowledge base uses different inputs, and its result is reported separately. 8

ForgeKernel: Self-Evolving Kernels

Preprint, 2026, arXiv

Table 2. Initial performance and subsequent gains along two optimization paths. Initial performance vs. expert

Operator

Round 1 gain vs. initial

Round 2 gain vs. round 1

Final gain vs. expert

Reference-free generation, then optimization (ForgeKernel optimizes the generated implementation) GDN prefill fwd GQA fwd (16:1) DeepSeek-V4.1 fused sparse-attention prefill

+82.6% +22.8% +5.3%

+9.5% +5.0% +10.0%

— — —

+100% +29% +15.5%

— — +4.2%

+7% +3% +15%

Expert-reference optimization, then self-evolution (ForgeKernel throughout) −0.5% +1.1% +0.3%

DSA sparse attn fwd GDN prefill bwd GQA fwd (16:1)

Table 3. Performance gains over expert references in component ablations.

Configuration Full ForgeKernel No milestone template First milestone only No SASS (NCU kept) No knowledge base

Attn bwd

GDN prefill fwd

DSA sparse fwd

Geo. mean

+0.6% +16.2%† −1.0% −10.2% +0.3%

−1.5% −16.2% −4.7% +16.2% −16.4%

+2.0% +1.4% −1.2% −5.1% +6.2%

+0.4% −0.4% −2.3% −0.3% −3.8%

self-evolution case focuses on using an accepted implementation as the reference for a subsequent run, with renewed analysis and design to explore further architectural improvements. These roles are complementary: the former organizes architecture changes, while the latter makes an existing optimized artifact the starting point for further architecture exploration. Architecture-agent-driven replanning: GQA.. The first GQA attention-forward self-evolution round illustrates how the architecture agent coordinates interdependent changes. The initial implementation assigns one query block to each thread block (CTA), which then exits. Rather than proposing only local code changes, the architecture agent changes this to persistent execution, so that the same CTA processes multiple query blocks. It organizes cross-block state management and Q-data prefetching into a staged plan for the coding agent to implement. The coding agent first implements the persistent loop and cross-block state management, obtaining a local performance gain of approximately +0.8%. It then adds Q double buffering to load the next query block’s data early, obtaining a further local gain of approximately +6.1%. Both steps undergo correctness checks and separate performance measurements; their local gains cannot be directly added. The round ultimately improves performance over the preceding generation by +9.9% (Table 2), including other changes; this gain cannot be attributed entirely to the two steps above. The case highlights how separating architecture from coding supports coordinated planning: the initial structural change delivers only a modest gain, but enables subsequent pipeline optimization, with the coding agent implementing the interdependent changes in dependency order.

Disassembly information. Removing SASS while retaining NCU profiling reduces performance on attention backward and DSA sparse attention forward, whereas GDN prefill forward exceeds the full system. Knowledge base. Knowledge reuse is widespread (Figure 6a). Nearly all configurations in the SM90 provenance cohort adopt external knowledge. The most common adoption range is 6–10 distinct entries, while the highest reaches 26–30: operators typically draw on multiple entries, and the extent of reuse varies substantially, consistent with the differing ablation effects. Reuse also crosses operator boundaries in both directions (Figure 6b). In the DSA–GQA and GEMM–MHA cases, new mechanisms developed by a knowledge recipient are adopted in subsequent runs of the original contributor, establishing cross-operator knowledge feedback. The knowledge base therefore serves as a shared, evolving collection of hardware mechanisms. Counting rules and the detailed source– recipient matrix are provided in Appendix E. 4.5

+7.2% +2.0% +9.9%

Case Studies: Two Mechanisms for Architecture Changes

Architecture reassessment through self-evolution: DSA.. The DSA sparse attention-forward case examines how an existing optimized artifact becomes the starting point for a new round of architecture analysis. The initial

The two cases illustrate the roles of the architecture agent and self-evolution in architecture changes. By separating architecture from coding, the architecture agent explicitly plans and revises interdependent architectural decisions. The 9

Yaojian Chen, Dazhi Peng, Yiyun Zheng, Zhui Zhu, Shengye Xiang, Leshan Li, Zhen Li, Xianglong Deng, Haojun Sun, Qingfeng He, Shangzhan Li, Yifei Shen, Preprint, 2026, arXiv Lin Gan, Yuxuan Li, Xu Han, and Zhiyuan Liu (a) External knowledge adoption

26–30 entries 21–25 entries

16–20 entries

8.7% 8.7%

(b) Reciprocal knowledge reuse

Fixed-container reuse Elementwise Handoff synchronization latency scheduling

0–5 entries 17.4%

DSA sparse attn fwd

DSA sparse attn fwd

8.7% TMA epilogue stores Accumulator writeback

13.0% 11–15 entries

GQA bwd (8:1)

43.5%

FP16 GEMM

6–10 entries

WGMMA deferred wait FP16 GEMM

MHA bwd

Figure 6. External-knowledge adoption and reciprocal reuse. implementation uses non-persistent execution with one resident CTA per SM. After this implementation becomes the reference, the next round obtains approximately 4.5% gain within the original architecture, after which further adjustments provide no effective improvement. Updating the reference does not require retaining the preceding generation’s architecture, allowing the system to reassess its computation organization rather than only tune parameters within that structure. Through diagnostic experiments, the architecture agent finds that each CTA’s query load, initial KV fills, and final output store lack computation from adjacent work items to overlap with. A probe removing output stores changes performance by approximately 6.0%, helping isolate uncovered overhead; this probe is diagnostic, not a complete candidate implementation. The system therefore adopts a persistent loop, overlapping one work item’s output store with the next item’s loads (Figure 7), and ultimately obtains a gain of +7.2% over the preceding generation (Table 2). This case illustrates the role of reference updates in self-evolution: retaining an existing artifact as a performance reference while reopening optimization through renewed diagnosis and architecture analysis, rather than fixing its structure as a constraint on subsequent runs.

5

gen-0 (sealed): non-persistent, grid 8192, 1 CTA/SM load

KV trips (MMA)

store

KV trips (MMA)

load

store

CTA boundary: no MMA in flight

gen-1: persistent loop, grid = #SMs (D0.1 reopened) load

KV trips (MMA)

store load

item n’s write-out overlaps item n+1’s Q load and first KV fills KV trips (MMA)

store

saved time on one SM

Figure 7. DSA sparse attention-forward architecture change (schematic, not to scale). and compiler for GPU kernels, enabling developers to express and implement computation partitioning and data organization [23, 24]. Beyond supporting developers in expressing designs, another line of work explores automating design choices. For example, Mirage jointly searches algebraic and scheduling transformations through a unified representation spanning kernels, thread blocks, and threads [27]. LLM-based operator generation and optimization. Prior work improves operator generation through feedbackdriven iteration, evolutionary search, and model training. GEAK generates Triton kernels through a feedback-driven reasoning loop [25]. STARK and Astra combine multi-agent collaboration with search, profiling, and planning [9, 26]; AlphaEvolve combines LLMs with automated evaluation for evolutionary search [16]. These methods focus on code iteration within a single run, without treating architectural decisions as separately revisable from measurements or using accepted kernels as references for a new round of architectural analysis. Kevin incorporates multi-turn implementation and execution feedback into reinforcement learning, while CUDA-L1

Related Work

Operator optimization and programming infrastructure. High-performance operator implementations require computation and data access to be organized according to computational characteristics and target hardware. FlashAttention uses tiling to reduce device-memory accesses in attention [6], while FlashInfer provides customizable, efficient attention implementations for LLM inference serving [29]. Translating such designs into efficient code requires appropriate programming abstractions and compiler support: CUTLASS provides composable high-performance computational components, and Triton provides a programming language 10

ForgeKernel: Self-Evolving Kernels

Preprint, 2026, arXiv

Cross-hardware design reuse. Hardware mechanisms in the knowledge base depend on specific architectures and applicability conditions; transfer to other hardware requires adaptation and validation. Future work will investigate mechanism abstraction, adaptation, and validation for cross-hardware reuse while retaining target-specific optimization.

uses contrastive reinforcement learning to improve CUDA optimization [1, 12]. Training changes the model that emits code, not the revision of architectural decisions during optimization. EGG fixes the algorithmic structure before hardware tuning, without revisiting it based on subsequent measurements [10]. CuBridge adapts expert attention kernels to target operators through an executable intermediate representation, retaining the expert implementation’s structure [15]. AKO optimizes kernels through a multi-agent workflow, with a human orchestrator redirecting stalled searches [14]. Structural choices or search redirections in these workflows are not explicit architectural decision revisions triggered by measurements. Another line of work retains prior experience and uses it to guide later generation. KernelClaw stores reusable expert optimization knowledge in long-term memory and uses short-term memory to avoid repeated backtracking within a run, replacing the model’s implicit heuristics with expert knowledge [22]. KLineage progressively simplifies expert implementations through validated steps, distilling each step into a skill with explicit applicability conditions for subsequent code generation [31]. Both retain knowledge or skills obtained from the expert side and use them to guide code generation. They do not store hardware mechanisms validated by independent probes during optimization as knowledge shared across operators, and later runs do not cite such entries when setting architectural decisions. ForgeKernel separates architectural design from code implementation, organizing milestone plans around explicit architectural decisions and revising those decisions in response to measurement feedback. When structural bottlenecks limit performance, the system revisits computation partitioning, dataflow, or pipeline organization rather than only tuning within a fixed architecture. For knowledge reuse, ForgeKernel distills hardware mechanisms during optimization, validates them through independent probes on the target hardware, and stores them in a knowledge base shared across operators and input shapes to guide architectural design and implementation planning in subsequent runs. Existing implementations can also become new references: their architectures are reassessed through measurement rather than carried forward by directly inheriting their source code.

6

7

Conclusion

We presented ForgeKernel, a self-evolving agent for kernel optimization. ForgeKernel separates architectural design from code implementation and revises work partitioning, dataflow, and pipeline organization based on performance feedback. Across runs, it uses validated kernels as new references and accumulates hardware mechanisms validated by independent probes in a shared knowledge base. Reference updates enable further architectural reassessment, while knowledge reuse carries optimization experience across runs and operators; together, they support continued self-evolution. With self-evolution and knowledge reuse, ForgeKernel achieves a geometric-mean speedup of 1.10× over reference implementations from expert libraries and other kernelgeneration systems across 29 configurations covering 27 operators on NVIDIA Hopper and Blackwell GPUs. On Ascend, it outperforms the per-operator public best on the public case set of all 8 level-4 operators in the CANN-Bench Season 1 final leaderboard, with a geometric-mean performance gain of +25.1%. Public submissions generated by ForgeKernel rank first on all three CANN-Bench 2.0 solution leaderboards (Ascend 910C, 950DT, and 950PR) and on the SOL-ExecBench FlashInfer-Bench leaderboard as of 2026-10-30.

References [1] Carlo Baronio, Pietro Marsella, Ben Pan, Simon Guo, and Silas Alberti. 2025. Kevin: Multi-Turn RL for Generating CUDA Kernels. arXiv:2507.11948 [2] CANNBench. 2026. CANNBench Public Leaderboard. https:// cannbench.com/leaderboard. CANN-Bench 2.0 solution snapshot as of 2026-10-30: ForgeKernel ranks first on the Ascend 910C, 950DT, and 950PR leaderboards. Historical Season 1, version 1.1, A3 (Ascend 910C) result: ForgeKernel: total score 8566.1110; standard/hidden scores 4362.7547/4203.3563, passed cases 1060/1060 and 4185/4240, geometric-mean speedups 1.7591600842 and 1.8543285162. Accessed 2026-10-30. [3] Shiyi Cao, Ziming Mao, Joseph E. Gonzalez, and Ion Stoica. 2026. KSearch: LLM Kernel Generation via Co-Evolving Intrinsic World Model. arXiv:2602.19128 [4] Terry Chen, Bing Xu, and Kirthi Devleker. 2025. Automating GPU Kernel Generation with DeepSeek-R1 and Inference Time Scaling. NVIDIA Developer Blog. https://developer.nvidia.com/blog/automating-gpukernel-generation-with-deepseek-r1-and-inference-time-scaling/. [5] Yaojian Chen, Dazhi Peng, Zhui Zhu, Leshan Li, Haojun Sun, Yiyun Zheng, Qingfeng He, Shangzhan Li, Xianglong Deng, Yifei Shen, Shengye Xiang, Zhen Li, Lin Gan, Yuxuan Li, Xu Han, and Zhiyuan Liu. 2026. JustGoal: With Only a Goal, the AI Baseline Is Above the Expert Kernels. arXiv:2610.36434

Limitations and Future Work

Architecture exploration and local optima. ForgeKernel revisits existing architectures through explicit planning and feedback-driven revision, but exploration is limited by the agents’ ability to generate candidates, and optimization may settle at local optima. Future work will investigate planning strategies that broaden the architecture search space and increase candidate diversity while preserving autonomous exploration and generalization. 11

Yaojian Chen, Dazhi Peng, Yiyun Zheng, Zhui Zhu, Shengye Xiang, Leshan Li, Zhen Li, Xianglong Deng, Haojun Sun, Qingfeng He, Shangzhan Li, Yifei Shen, Preprint, 2026, arXiv Lin Gan, Yuxuan Li, Xu Han, and Zhiyuan Liu [6] Tri Dao, Daniel Y. Fu, Stefano Ermon, Atri Rudra, and Christopher Ré. 2022. FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness. In Advances in Neural Information Processing Systems (NeurIPS). [7] DeepSeek-AI. 2024. DeepSeek-V2: A Strong, Economical, and Efficient Mixture-of-Experts Language Model. arXiv:2405.04434 [8] DeepSeek-AI. 2025. DeepSeek-V3.2-Exp: Boosting Long-Context Efficiency with DeepSeek Sparse Attention. https://github.com/deepseekai/DeepSeek-V3.2-Exp/blob/main/DeepSeek_V3_2.pdf. Tech report, September 2025. [9] Juncheng Dong, Yang Yang, Tao Liu, et al. 2025. STARK: Strategic Team of Agents for Refining Kernels. arXiv:2510.16996 [10] Yaochen Han, Ke Fan, Hongxu Jiang, et al. 2026. EGG: An ExpertGuided Agent Framework for Kernel Generation. arXiv:2606.26758 [11] Kimi Team. 2025. Kimi Linear: An Expressive, Efficient Attention Architecture. arXiv:2510.26692 Introduces Kimi Delta Attention (KDA). [12] Xiaoya Li, Albert Wang, Guoyin Wang, Jiwei Li, and Chris Shum. 2026. CUDA-L1: Improving CUDA Optimization via Contrastive Reinforcement Learning. In International Conference on Learning Representations (ICLR). arXiv:2507.14111. [13] Sihao Liu, Ligeng Zhu, et al. 2026. Kernel Design Agents (KDA): An Agent-Centric Workflow for High-Performance CUDA Kernels. https://github.com/NVlabs/kda. NVIDIA and MIT HAN Lab; software release. [14] Mao Luo, Hongbin Li, Feng Lin, Hanling Yi, and Zhe Huang. 2026. Agentic Kernel Optimization: Generating State-of-the-Art GPU Kernels Without Hand-Written CUDA. arXiv:2608.14560 [15] Xing Ma, Yangjie Zhou, Wu Sun, et al. 2026. CuBridge: An LLM-Based Framework for Understanding and Reconstructing High-Performance Attention Kernels. arXiv:2605.05023 ACL 2026. [16] Alexander Novikov, Ngân Vũ, Marvin Eisenberger, et al. 2025. AlphaEvolve: A coding agent for scientific and algorithmic discovery. arXiv:2506.13131 [17] NVIDIA Corporation. 2022. NVIDIA H100 Tensor Core GPU Architecture Whitepaper. https://resources.nvidia.com/en-us-hopperarchitecture/nvidia-h100-tensor-c. Accessed 2026-09. [18] NVIDIA Corporation. 2024. NVIDIA Blackwell Architecture Technical Overview. https://resources.nvidia.com/en-us-blackwell-architecture. Accessed 2026-09. [19] NVIDIA Corporation. 2026. Parallel Thread Execution ISA, Version 9.4. https://docs.nvidia.com/cuda/parallel-thread-execution/. wgmma.mma_ async, cp.async.bulk.tensor (TMA), and tcgen05 sections. Accessed 2026-09. [20] NVIDIA Research. 2026. SOL-ExecBench: FlashInfer-Bench Leaderboard. https://research.nvidia.com/benchmarks/sol-execbench/ leaderboard/collection/4/B200. B200 collection leaderboard, evaluation stack v1.1. Submission 76892 (ForgeKernel): SOL Score 0.815204, Fast1 647/660, solved 26/26, Avg Speedup 5.5403. Accessed 2026-10-30. [21] Jay Shah, Ganesh Bikshandi, Ying Zhang, Vijay Thakkar, Pradeep Ramani, and Tri Dao. 2024. FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-precision. arXiv:2407.08608 [22] Qitong Sun, Jun Han, Tianlin Li, et al. 2026. KernelClaw: A MultiAgent Framework for GPU Kernel Optimization with Dual-Memory Mechanisms. arXiv:2603.10085 [23] Vijay Thakkar, Pradeep Ramani, Cris Cecka, Aniket Shivam, Honghao Lu, Ethan Yan, Jack Kosaian, Mark Hoemmen, Haicheng Wu, Andrew Kerr, et al. 2026. CUTLASS: CUDA Templates for Linear Algebra Subroutines. https://github.com/NVIDIA/cutlass. Version 4.8; includes CuTe layout algebra and the Python CuTe-DSL. Accessed 2026-09. [24] Philippe Tillet, Hsiang-Tsung Kung, and David 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 (MAPL). DOI: 10.1145/3315508.3329973.

[25] Jianghui Wang, Vinay Joshi, Saptarshi Majumder, et al. 2025. GEAK: Introducing Triton Kernel AI Agent & Evaluation Benchmarks. arXiv:2507.23194 [26] Anjiang Wei, Tianran Sun, Yogesh Seenichamy, et al. 2025. Astra: A Multi-Agent System for GPU Kernel Performance Optimization. arXiv:2509.07506 [27] Mengdi Wu, Xinhao Cheng, Shengyu Liu, et al. 2025. Mirage: A Multi-Level Superoptimizer for Tensor Programs. In 19th USENIX Symposium on Operating Systems Design and Implementation (OSDI). arXiv:2405.05751. [28] Songlin Yang, Jan Kautz, and Ali Hatamizadeh. 2024. Gated Delta Networks: Improving Mamba2 with Delta Rule. arXiv:2412.06464 ICLR 2025. [29] Zihao Ye, Lequn Chen, Ruihang Lai, Wuwei Lin, Yineng Zhang, Stephanie Wang, Tianqi Chen, Baris Kasikci, Vinod Grover, Arvind Krishnamurthy, and Luis Ceze. 2025. FlashInfer: Efficient and Customizable Attention Engine for LLM Inference Serving. In Proceedings of Machine Learning and Systems (MLSys). Best Paper Award. arXiv:2501.01005. [30] Zihao Ye, Yingyi Huang, Hongyi Jin, Bohan Hou, Junru Shao, Zhongming Yu, Jinqi Chen, Meghan Cowan, Shiyi Cao, Shanli Xing, Hanfeng Chen, Vinod Grover, Tianqi Chen, and Luis Ceze. 2026. CAKE: Compiler-Agent Co-Design for Frontier Kernel Evolution. arXiv:2608.12629 [31] Shuoming Zhang, Qiuchu Yu, Yangyu Zhang, et al. 2026. KLineage: Recovering the Missing When of Kernel Optimization by Deoptimizing Experts. arXiv:2605.28213 [32] Dongyun Zou, Yixin Dong, Junxian Guo, et al. 2026. KDA-v0.5 Kernels for the MLSys 2026 FlashInfer AI Kernel Generation Contest. https://github.com/mit-han-lab/mlsys2026-flashinfer-contestsolution. Speedups over the contest winners re-measured by the authors under the flashinfer-bench CUPTI protocol, not by the contest evaluator. Accessed 2026-09.

A

Correctness Judges

Correctness validation first follows the official or benchmark tolerance test associated with each operator source. Operators covered by FlashInfer-Bench, including those in the SOL-ExecBench FlashInfer-Bench collection, use FlashInferBench’s tolerance tests; FlashAttention operators use the official FlashAttention tolerance tests; other operators use the criterion associated with their source. Within this source mapping, the concrete judges include DeepGEMM’s aggregate criterion for GEMMs, a relative max-abs bound for attention gradients, and a scaled tolerance plus a cosine bound for the GDN and KDA (Kimi Delta Attention [11]) state kernels, whose small outputs would pass a fixed absolute tolerance trivially. The GDN judges further add correlated-key and long-memory inputs that catch state-decay shortcuts.

B

Reference Versions and Entry Points

Library references are FlashAttention-4 (CuTe DSL), cuDNN 9 with cuDNN-frontend 1.28, DeepGEMM, FlashMLA, FlashInfer and FlashQLA, and are measured at these entry points: FlashAttention-4’s forward and _flash_ attn_bwd, FlashInfer’s fused chunk_gated_delta_rule, FlashQLA’s TileLang backward, DeepGEMM’s fp8_fp4_ mqa_logits, FlashMLA’s flash_mla_sparse_fwd, and 12

ForgeKernel: Self-Evolving Kernels

Preprint, 2026, arXiv

cuDNN-frontend’s CuTe DSL radix top-𝑘, sparse_attention_ backward, indexer_backward (three kernels summed), and the sparse_attn + sparse_indexer score-recompute composition. CANN-Bench 1.1.2 supplies the Ascend baselines; its public repository releases through v1.0.0 as of this writing.

C

Campaign Shapes

Table 4 lists the problem shape of every task and optimization run of the main paper.

D

Per-Task Public-Benchmark Speedups

Tables 5 and 6 list all 53 CANN-Bench Season 1 A3 operators and all 26 SOL-ExecBench FlashInfer-Bench kernels, respectively. All values are speedup ratios against the corresponding benchmark baseline; values below one indicate slowdowns and are retained. For CANN-Bench, we report the standard-set, hidden-set, and combined speedups from the Season 1 final results, checked on 2026-10-15. The combined column uses the benchmark’s weighted geometric mean with standard/hidden weights 0.2/0.8. Reported ratios are rounded to three decimals; they do not imply that all hidden cases passed. GRU and LSTM use hardware-bound-derived baselines rather than reference implementations [2]. For SOL-ExecBench, we use ForgeKernel’s per-kernel leaderboard entries on B200 with evaluation stack v1.1, accessed on 2026-10-30. Each ratio is the geometric-mean speedup across that kernel’s workloads against the official Scoring Baseline, not the Reference Implementation or SOL Bound. The geometric mean across all 26 kernel ratios is 5.54× [20].

E

Knowledge-Reuse Provenance

Figure 8 aggregates the SM90 run snapshot by source and adopting operator configuration. Each cell counts distinct external entries; repeated use across runs is deduplicated. White cells have no qualifying record, and gray cells denote self-reuse, which is excluded. All 23 recipient configurations are retained, including the zero-adoption configuration; source rows without any qualifying outgoing reuse are omitted. The cohort follows the original provenance plot: MHA and GQA attention variants are separated, and MHA forward is not a recipient because it is outside that plot’s evaluated configuration set. This is a provenance cohort, not a census of every evaluated task. Entry origins come from recorded birth cells, falling back to the earliest run of the originating operator. The original run-group ordering and temporal exclusions are retained. Adoption uses the milestone-citation and decision-delivery snapshot, rather than mere retrieval. Where commit lineage is available, delivery is checked against the best correct implementation; otherwise the extractor falls back to process 13

records. Thus the snapshot does not independently certify final-code retention for every entry, and zero denotes no qualifying recorded adoption. The two examples in Figure 6(b) show selected entries from the recorded exchanges, not an exhaustive list or a causal performance attribution.

Yaojian Chen, Dazhi Peng, Yiyun Zheng, Zhui Zhu, Shengye Xiang, Leshan Li, Zhen Li, Xianglong Deng, Haojun Sun, Qingfeng He, Shangzhan Li, Yifei Shen, Preprint, 2026, arXiv Lin Gan, Yuxuan Li, Xu Han, and Zhiyuan Liu

Table 4. Input shapes for the evaluated tasks and optimization runs. Operator

GPU

Shape

/goal agent ceiling attn_fwd_gqa fa_bwd_gqa gdn_prefill_fwd gdn_prefill_bwd dsa_indexer_fwd dsa_indexer_topk dsa_sparse_fwd dsa_sparse_bwd dsa_score_recompute dsa_indexer_bwd

H100 H100 H100 H100 H100 H100 H100 H100 H100 H100

B=2 S=4096 H_Q=16/H_KV=1 D=128 bf16 causal B=10 N=4096 H_Q=16/H_KV=2 D=64 fp16 causal N=8192 h_qk=16/h_v=32 d=128 bf16 N=8192 h_qk=16/h_v=32 d=128 bf16 s_q=4096 s_kv=32768 H=64 D=128 fp8 logits fp32 [4096 x 32768] k=1024 s_q=4096 s_kv=32768 h=128 d=512 topk=1024 s_q=32768 s_kv=8192 h=128 d=512 topk=1024 s_q=32768 s_kv=8192 topk=1024 s_q=32768 s_kv=8192 h=64 d=128 topk=1024

ForgeKernel optimization runs attn_fwd (GQA 16:1) H100 attn_bwd (MHA) H100 attn_bwd (GQA 8:1) H100 fp8_gemm H100 fp8_grouped_gemm H100 gemm (fp16) H100 conv2d_fwd H100 conv2d_wgrad H100 conv2d_dgrad H100 conv3d_fwd H100 dsa_indexer_fwd H100 dsa_indexer_topk H100 dsa_sparse_fwd H100 dsa_sparse_bwd H100 dsa_score_recompute H100 dsa_indexer_bwd H100 gdn_prefill_fwd H100 gdn_prefill_bwd H100 gqa_paged_decode H100 mla_paged_decode H100 mla_paged_prefill H100 trimul H100 attn_bwd (MHA) B200 fp4_grouped_gemm B200 sparse_mla_prefill B200 fused_attn_rope_cast B200 dsa_sparse_attn_paged gdn_prefill_varlen recurrent_kda_prefill

B200 B200 B200

B=2 S=4096 H=16/1 D=128 bf16 causal B=10 N=4096 H=16 D=64 fp16 causal B=10 N=4096 H=16/2 D=64 fp16 causal M=N=K=4096 blk 128 G=8 N=K=4096 blk 128 M=N=K=8192 fp16 N=128 C=K=256 28x28 3x3 s1 N=128 C=K=256 28x28 3x3 s1 N=128 C=K=256 56x56 3x3 s2 N=8 C=K=256 8x28x28 3x3x3 s_q=4096 s_kv=32768 H=64 D=128 fp8 4096x32768 fp32 k=1024 s_q=4096 s_kv=32768 h=128 d=512 topk=1024 s_q=32768 s_kv=8192 h=128 d=512 topk=1024 s_q=32768 s_kv=8192 topk=1024 s_q=32768 s_kv=8192 H=64 D=128 topk=1024 N=8192 h=16/32 d=128 bf16 N=8192 h=16/32 d=128 bf16 workload 7a7fc288 (61535 pages) workload 5bef8d88 (B=64; 75145 tokens) workload fe63f292 (B=18; 15883 q) S=1024 D=384 H=128 (case 6) B=8 N=4096 H=16 D=128 bf16 causal G=8 M=4096/group N=K=4096 s_q=4096 s_kv=32768 H_Q=128 H_KV=1 D=512 topk=1024 bf16 attn_sink s_q=4096 s_kv=32768 H_Q=128 H_KV=1 D=512 rope=64 topk=1024 n_wv_group=16 bf16->fp8 attn_sink q_ norm workload 564007ac (T=8; 8462 pages) workload a9540651 (T=8192 N=20) T=8192 H=96 K=V=128 bf16

14

ForgeKernel: Self-Evolving Kernels

Preprint, 2026, arXiv

Table 5. Per-operator CANN-Bench Season 1 A3 speedups (×) against the benchmark baseline on Ascend 910C. Operator

Level

Standard

Hidden

Combined

Gelu SwiGlu ForeachNorm Sigmoid Exp MaskedScale Mish ForeachAddcdivScalar GridSampler3D Cummin Gcd DynamicQuant CrossEntropyLoss ApplyRotaryPosEmb GroupNorm Softmax RmsNorm Scatter Gather UnsortedSegmentSum ResizeBilinear ArgMax ApplyAdamW Maximum NMS MoeGatingTopKSoftmax AddRmsNormDynamicQuant

L1 L1 L1 L1 L1 L1 L1 L1 L2 L2 L2 L2 L2 L2 L2 L2 L2 L2 L2 L2 L2 L2 L2 L2 L3 L3 L3

0.910 2.113 2.088 0.972 2.717 2.652 2.369 1.804 1.343 1.285 5.704 2.270 2.253 2.510 1.589 1.840 1.709 1.028 2.512 1.585 5.962 2.921 1.965 1.570 3.191 1.931 2.459

3.263 4.013 3.707 4.679 4.624 4.837 3.965 4.997 0.861 1.032 3.296 5.754 1.691 2.216 2.132 2.380 3.032 0.363 2.407 1.669 10.397 2.926 4.378 4.117 0.981 1.478 3.250

2.527 3.530 3.305 3.417 4.157 4.289 3.577 4.076 0.941 1.078 3.678 4.777 1.791 2.272 2.011 2.261 2.703 0.447 2.428 1.652 9.302 2.925 3.730 3.395 1.242 1.559 3.074

Operator

Level

Standard

Hidden

Combined

DepthwiseConv2D TopK MoeFinalizeRouting Dilation2D MhcSinkhorn StridedSlice DequantSwigluQuant WeightQuantBatchMatmul EngramGateFusion Conv3DBackpropFilter Unique GroupedMatmul Conv2D AdaptiveAvgPool3D MoeReRouting Transpose ROIAlign QuantMatmul GroupedMatmulSwigluQuant LSTM MLA GRU MlaProlog MHA SparseFlashAttention GQA

L3 L3 L3 L3 L3 L3 L3 L3 L3 L3 L3 L3 L3 L3 L3 L3 L3 L3 L4 L4 L4 L4 L4 L4 L4 L4

1.905 1.286 2.306 2.086 3.188 3.667 2.512 1.299 0.988 1.758 1.487 1.868 1.019 3.478 2.015 1.213 0.842 1.146 2.418 0.450 1.139 0.546 1.562 1.198 1.461 1.184

1.143 0.855 3.229 1.190 2.603 3.958 3.744 2.143 0.968 1.104 0.589 2.406 1.078 3.074 3.201 1.913 0.631 1.808 1.967 0.147 1.649 0.218 1.198 1.112 0.207 1.423

1.266 0.927 3.019 1.331 2.711 3.898 3.457 1.939 0.972 1.212 0.709 2.287 1.066 3.151 2.918 1.746 0.669 1.651 2.050 0.184 1.532 0.262 1.263 1.128 0.306 1.371

Table 6. Per-kernel SOL-ExecBench FlashInfer-Bench speedups against Scoring Baseline on B200 (evaluation stack v1.1). Speedup ( × )

Kernel

1.4651 4.4455 1.8506 2.4275 2.3825 1.4514 1.3595 1.5171 1.1265 1.2387 1.3481 16.1650 9.9521 159.3385 133.7091 17.5878 16.0272 31.9883 93.2019 28.1198 2.3360 2.3813 4.8512 2.4907 2.7274 3.6834

001_fused_add_rmsnorm_h2048 002_fused_add_rmsnorm_h4096 003_fused_add_rmsnorm_h7168 004_gemm_n128_k2048 005_gemm_n256_k7168 006_gemm_n2048_k4096 007_gemm_n4096_k4096 008_gemm_n4096_k14336 009_gemm_n5120_k2048 010_gemm_n6144_k4096 011_gemm_n28672_k4096 012_gqa_paged_decode_h32_kv4_d128_ps1 013_gqa_paged_decode_h32_kv8_d128_ps1 014_gqa_paged_prefill_causal_h32_kv4_d128_ps1 015_gqa_paged_prefill_causal_h32_kv8_d128_ps1 016_gqa_ragged_prefill_causal_h32_kv4_d128 017_gqa_ragged_prefill_causal_h32_kv8_d128 018_mla_paged_decode_h16_ckv512_kpe64_ps1 019_mla_paged_prefill_causal_h16_ckv512_kpe64_ps1 020_moe_fp8_block_scale_ds_routing_topk8_ng8_kg4_e32_h7168_i2048 021_rmsnorm_h128 022_rmsnorm_h512 023_rmsnorm_h1536 024_rmsnorm_h2048 025_rmsnorm_h4096 026_rmsnorm_h7168

15

Yaojian Chen, Dazhi Peng, Yiyun Zheng, Zhui Zhu, Shengye Xiang, Leshan Li, Zhen Li, Xianglong Deng, Haojun Sun, Qingfeng He, Shangzhan Li, Yifei Shen, Preprint, 2026, arXiv Lin Gan, Yuxuan Li, Xu Han, and Zhiyuan Liu

Source operator → adopting operator Conv2D fwd DSA indexer fwd DSA indexer top-k DSA sparse attn bwd DSA sparse attn fwd FP16 GEMM FP8 GEMM FP8 grouped GEMM GDN prefill fwd GQA bwd (8:1) GQA fwd (16:1)

2D

Co

Co

nv

Co

nv

2D fw d inp nv u tg 2D ra we d igh tg r Co ad nv DS 3D A fw ind d ex DS er bw A ind d ex DS er A ind fw DS d ex A er sc or t o er p-k DS ec A om sp p ar ut se DS e at A tn sp bw ar se d at tn f w FP d 16 GE MM FP FP 8G 8g E ro MM up ed GE GD MM N pr efi ll b GD wd N pr efi ll f GQ wd A GQ bw A d( de 8:1 co de ) ( p GQ ag ed A ) fw d( 16 :1) MH A b wd ML ML A A de de co co de de ML (p A ag pr ed efi ) ll ( pa ge d) Tri Mu l

MHA bwd

Distinct external knowledge entries 0

3

6

9

Figure 8. Sources and recipients of reused hardware knowledge.

16

Related documents

• MetaKernelBench: Measuring GPU Kernel Knowledge Transfer Beyond Code#1200665
• KernelZero: Co-Evolving Proposer and Coder for Continuously Improved GPU Kernel Generation#1117786
• RaMP: Runtime-Aware Megakernel Polymorphism for Mixture-of-Experts#1536120
• KernelOPT: Dispatch-Aware Agentic Search for GPU Kernel Optimization#1078130
• FastKernels: Benchmarking GPU Kernel Generation in Production#1186467