WorksComputer Science
ForgeKernel: Self-Evolving Kernels Catching and Surpassing Experts and Agents
Abstract
Purpose-built kernel agents can write high-performance GPU kernels, but existing systems generally lack cross-run knowledge accumulation and architectural evolution. Each run starts from the same initial knowledge, and later iterations largely optimize within an architecture fixed early in the search. We present ForgeKernel, a kernel-generation system that accumulates and reuses knowledge across runs while enabling continued evolution of kernel architectures. It analyzes the hardware behavior of a high-performance reference and distills validated architectural decisions into reusable knowledge of hardware mechanisms rather than imitating source code. It then seals a kernel that passes correctness and performance validation and uses it as the reference for the next optimization run, enabling architectural decisions to be revisited. Across 27 operators and 29 input configurations on NVIDIA Hopper and Blackwell GPUs, ForgeKernel matches or beats its expert reference on 25 configurations, with a geometric-mean performance gain of 10%. On Ascend 910C, 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 gain of 25.1%, while its full-suite submission ranks first on the 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.
Introduction
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 high-performance 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 cross-run 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 measurement-validated optimization experience.

Figure 1. Single-run optimization and cross-run self-evolution.
For architectural design, ForgeKernel uses a high-performance reference implementation as a design reference and starting point. It separates architecture from coding, with a dedicated architecture agent reviewing and adjusting
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 cross-run 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 CANN-Bench 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 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:
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× geometric-mean 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.
Background
Kernel Design
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, 18, 19, 21, 23]. These choices jointly affect data reuse, 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.
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.
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.

Figure 2. ForgeKernel architecture.
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 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. Self-evolution extends the loop through plan revision and reference updates (§3.4), while the hardware knowledge base retains validated mechanisms for reuse (§3.5).
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 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 . The initial milestone 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 mechanism-validation 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, 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.
Coding Agent
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.
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 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 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.
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.
Evaluation
This chapter evaluates the performance of ForgeKernel-generated 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.
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.
Comparison targets. Per-operator experiments use expert library implementations and kernels released by other generation systems as references. Libraries include FlashAttention-4, 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 § 9.
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 FlashInfer-Bench’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 § 8.
Timing method. On NVIDIA hardware, we measure device execution time under cold- 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 , where the numerator and denominator are the device execution times of the reference and candidate, respectively; performance gain is .
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 goal-driven setting of JustGo [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.
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, 950D1T, 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. Geometric-mean 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, Fast₁ = 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, Fast₁ 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.
| Benchmark | Coverage | Speedup | Rank |
| CANN-Bench v1.1<br>Ascend 910C | 53 operators<br>Standard: 1,060/1,060<br>Hidden: 4,185/4,240 | 1.84× | 1 |
| SOL-ExecBench<br>NVIDIA B200 | 26/26 kernels<br>660 workloads | 5.54× | 1 |
Table 1. Public-benchmark results.
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 Figure 3 and Figure 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.

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

Figure 4. Performance gains over published kernels from other generation systems.
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 ForgeKernel-generated 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 geometric-mean 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.

Figure 5. Performance gains over public-best submissions on Ascend 910C.
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 DeepSeek-V4.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 gains are estimated by compounding the initial and subsequent speedups. A dash denotes a round that was not conducted.
| Operator | Initial performance<br>vs. expert | Round 1 gain<br>vs. initial | Round 2 gain<br>vs. round 1 | Final gain<br>vs. expert |
| Reference-free generation, then optimization (ForgeKernel optimizes the generated implementation) | ||||
| GDN prefill fwd | +82.6% | +9.5% | — | +100% |
| GQA fwd (16:1) | +22.8% | +5.0% | — | +29% |
| DeepSeek-V4.1 fused sparse-attention prefill | +5.3% | +10.0% | — | +15.5% |
| Expert-reference optimization, then self-evolution (ForgeKernel throughout) | ||||
| DSA sparse attn fwd | −0.5% | +7.2% | — | +7% |
| GDN prefill bwd | +1.1% | +2.0% | — | +3% |
| GQA fwd (16:1) | +0.3% | +9.9% | +4.2% | +15% |
Table 2. Initial performance and subsequent gains along two optimization paths.
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 self-evolution 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 of the GQA experiment beginning with reference-free generation.
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 , , , and , 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 first-round 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-L2 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.
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.
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.
| Configuration | Attn bwd | GDN prefill fwd | DSA sparse fwd | Geo. mean |
| Full ForgeKernel | +0.6% | −1.5% | +2.0% | +0.4% |
| No milestone template | +16.2%† | −16.2% | +1.4% | −0.4% |
| First milestone only | −1.0% | −4.7% | −1.2% | −2.3% |
| No SASS (NCU kept) | −10.2% | +16.2% | −5.1% | −0.3% |
| No knowledge base | +0.3% | −16.4% | +6.2% | −3.8% |
Table 3. Performance gains over expert references in component ablations.
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.

Figure 6. External-knowledge adoption and reciprocal reuse.
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 § 12.
Case Studies: Two Mechanisms for Architecture Changes
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 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.
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 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.

Figure 7. DSA sparse attention-forward architecture change (schematic, not to scale).
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
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 feedback-driven 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 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.
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.
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.
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 kernel-generation 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.
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 FlashInfer-Bench’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.
Reference Versions and Entry Points
Library references are FlashAttention-4 (CuTe DSL), cuDNN 9 with cuDNN-front-end 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 cuDNN-frontend’s CuTe DSL radix top-k, 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. 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.
Campaign Shapes
Table 4 lists the problem shape of every task and optimization run of the main paper.
| Operator | GPU | Shape |
| * /goal agent ceiling* | ||
| attn_fwd_gqa | H100 | B=2 S=4096 H_Q=16/H_KV=1 D=128 bf16 causal |
| fa_bwd_gqa | H100 | B=10 N=4096 H_Q=16/H_KV=2 D=64 fp16 causal |
| gdn_prefill_fwd | H100 | N=8192 h_qk=16/h_v=32 d=128 bf16 |
| gdn_prefill_bwd | H100 | N=8192 h_qk=16/h_v=32 d=128 bf16 |
| dsa_indexer_fwd | H100 | s_q=4096 s_kv=32768 H=64 D=128 fp8 |
| dsa_indexer_topk | H100 | logits fp32 [4096 x 32768] k=1024 |
| dsa_sparse_fwd | H100 | s_q=4096 s_kv=32768 h=128 d=512 topk=1024 |
| dsa_sparse_bwd | H100 | s_q=32768 s_kv=8192 h=128 d=512 topk=1024 |
| dsa_score_recompute | H100 | s_q=32768 s_kv=8192 topk=1024 |
| dsa_indexer_bwd | H100 | s_q=32768 s_kv=8192 h=64 d=128 topk=1024 |
| ForgeKernel optimization runs | ||
| attn_fwd (GQA 16:1) | H100 | B=2 S=4096 H=16/1 D=128 bf16 causal |
| attn_bwd (MHA) | H100 | B=10 N=4096 H=16 D=64 fp16 causal |
| attn_bwd (GQA 8:1) | H100 | B=10 N=4096 H=16/2 D=64 fp16 causal |
| fp8_gemm | H100 | M=N=K=4096 blk 128 |
| fp8_grouped_gemm | H100 | G=8 N=K=4096 blk 128 |
| gemm (fp16) | H100 | M=N=K=8192 fp16 |
| conv2d_fwd | H100 | N=128 C=K=256 28x28 3x3 s1 |
| conv2d_wgrad | H100 | N=128 C=K=256 28x28 3x3 s1 |
| conv2d_dgrad | H100 | N=128 C=K=256 56x56 3x3 s2 |
| conv3d_fwd | H100 | N=8 C=K=256 8x28x28 3x3x3 |
| dsa_indexer_fwd | H100 | s_q=4096 s_kv=32768 H=64 D=128 fp8 |
| dsa_indexer_topk | H100 | 4096x32768 fp32 k=1024 |
| dsa_sparse_fwd | H100 | s_q=4096 s_kv=32768 h=128 d=512 topk=1024 |
| dsa_sparse_bwd | H100 | s_q=32768 s_kv=8192 h=128 d=512 topk=1024 |
| dsa_score_recompute | H100 | s_q=32768 s_kv=8192 topk=1024 |
| dsa_indexer_bwd | H100 | s_q=32768 s_kv=8192 H=64 D=128 topk=1024 |
| gdn_prefill_fwd | H100 | N=8192 h=16/32 d=128 bf16 |
| gdn_prefill_bwd | H100 | N=8192 h=16/32 d=128 bf16 |
| gqa_paged_decode | H100 | workload 7a7fc288 (61535 pages) |
| mla_paged_decode | H100 | workload 5bef8d88 (B=64; 75145 tokens) |
| mla_paged_prefill | H100 | workload fe63f292 (B=18; 15883 q) |
| trimul | H100 | S=1024 D=384 H=128 (case 6) |
| attn_bwd (MHA) | B200 | B=8 N=4096 H=16 D=128 bf16 causal |
| fp4_grouped_gemm | B200 | G=8 M=4096/group N=K=4096 |
| sparse_mla_prefill | B200 | s_q=4096 s_kv=32768 H_Q=128 H_KV=1 D=512 topk=1024 bf16 attn_sink |
| fused_attn_rope_cast | B200 | 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 |
| dsa_sparse_attn_paged | B200 | workload 564007ac (T=8; 8462 pages) |
| gdn_prefill_varlen | B200 | workload a9540651 (T=8192 N=20) |
| recurrent_kda_prefill | B200 | T=8192 H=96 K=V=128 bf16 |
Table 4. Input shapes for the evaluated tasks and optimization runs.
Per-Task Public-Benchmark Speedups
Table 5 and Table 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.
| Operator | Level | Standard | Hidden | Combined | Operator | Level | Standard | Hidden | Combined |
| Gelu | L1 | 0.910 | 3.263 | 2.527 | DepthwiseConv2D | L3 | 1.905 | 1.143 | 1.266 |
| SwiGlu | L1 | 2.113 | 4.013 | 3.530 | TopK | L3 | 1.286 | 0.855 | 0.927 |
| ForeachNorm | L1 | 2.088 | 3.707 | 3.305 | MoeFinalizeRouting | L3 | 2.306 | 3.229 | 3.019 |
| Sigmoid | L1 | 0.972 | 4.679 | 3.417 | Dilation2D | L3 | 2.086 | 1.190 | 1.331 |
| Exp | L1 | 2.717 | 4.624 | 4.157 | MhcSinkhorn | L3 | 3.188 | 2.603 | 2.711 |
| MaskedScale | L1 | 2.652 | 4.837 | 4.289 | StridedSlice | L3 | 3.667 | 3.958 | 3.898 |
| Mish | L1 | 2.369 | 3.965 | 3.577 | DequantSwigluQuant | L3 | 2.512 | 3.744 | 3.457 |
| ForeachAddcdivScalar | L1 | 1.804 | 4.997 | 4.076 | WeightQuantBatchMatmul | L3 | 1.299 | 2.143 | 1.939 |
| GridSampler3D | L2 | 1.343 | 0.861 | 0.941 | EngramGateFusion | L3 | 0.988 | 0.968 | 0.972 |
| Cummin | L2 | 1.285 | 1.032 | 1.078 | Conv3DBackpropFilter | L3 | 1.758 | 1.104 | 1.212 |
| Gcd | L2 | 5.704 | 3.296 | 3.678 | Unique | L3 | 1.487 | 0.589 | 0.709 |
| DynamicQuant | L2 | 2.270 | 5.754 | 4.777 | GroupedMatmul | L3 | 1.868 | 2.406 | 2.287 |
| CrossEntropyLoss | L2 | 2.253 | 1.691 | 1.791 | Conv2D | L3 | 1.019 | 1.078 | 1.066 |
| ApplyRotaryPosEmb | L2 | 2.510 | 2.216 | 2.272 | AdaptiveAvgPool3D | L3 | 3.478 | 3.074 | 3.151 |
| GroupNorm | L2 | 1.589 | 2.132 | 2.011 | MoeReRouting | L3 | 2.015 | 3.201 | 2.918 |
| Softmax | L2 | 1.840 | 2.380 | 2.261 | Transpose | L3 | 1.213 | 1.913 | 1.746 |
| RmsNorm | L2 | 1.709 | 3.032 | 2.703 | ROIAlign | L3 | 0.842 | 0.631 | 0.669 |
| Scatter | L2 | 1.028 | 0.363 | 0.447 | QuantMatmul | L3 | 1.146 | 1.808 | 1.651 |
| Gather | L2 | 2.512 | 2.407 | 2.428 | GroupedMatmulSwigluQuant | L4 | 2.418 | 1.967 | 2.050 |
| UnsortedSegmentSum | L2 | 1.585 | 1.669 | 1.652 | LSTM | L4 | 0.450 | 0.147 | 0.184 |
| ResizeBilinear | L2 | 5.962 | 10.397 | 9.302 | MLA | L4 | 1.139 | 1.649 | 1.532 |
| ArgMax | L2 | 2.921 | 2.926 | 2.925 | GRU | L4 | 0.546 | 0.218 | 0.262 |
| ApplyAdamW | L2 | 1.965 | 4.378 | 3.730 | MlaProlog | L4 | 1.562 | 1.198 | 1.263 |
| Maximum | L2 | 1.570 | 4.117 | 3.395 | MHA | L4 | 1.198 | 1.112 | 1.128 |
| NMS | L3 | 3.191 | 0.981 | 1.242 | SparseFlashAttention | L4 | 1.461 | 0.207 | 0.306 |
| MoeGatingTopKSoftmax | L3 | 1.931 | 1.478 | 1.559 | GQA | L4 | 1.184 | 1.423 | 1.371 |
| AddRmsNormDynamicQuant | L3 | 2.459 | 3.250 | 3.074 |
Table 5. Per-operator CANN-Bench Season 1 A3 speedups () against the benchmark baseline on Ascend 910C.
| Kernel | Speedup () |
001_fused_add_rmsnorm_h2048 | 1.4651 |
002_fused_add_rmsnorm_h4096 | 4.4455 |
003_fused_add_rmsnorm_h7168 | 1.8506 |
004_gemm_n128_k2048 | 2.4275 |
005_gemm_n256_k7168 | 2.3825 |
006_gemm_n2048_k4096 | 1.4514 |
007_gemm_n4096_k4096 | 1.3595 |
008_gemm_n4096_k14336 | 1.5171 |
009_gemm_n5120_k2048 | 1.1265 |
010_gemm_n6144_k4096 | 1.2387 |
011_gemm_n28672_k4096 | 1.3481 |
012_gqa_paged_decode_h32_kv4_d128_ps1 | 16.1650 |
013_gqa_paged_decode_h32_kv8_d128_ps1 | 9.9521 |
014_gqa_paged_prefill_causal_h32_kv4_d128_ps1 | 159.3385 |
015_gqa_paged_prefill_causal_h32_kv8_d128_ps1 | 133.7091 |
016_gqa_ragged_prefill_causal_h32_kv4_d128 | 17.5878 |
017_gqa_ragged_prefill_causal_h32_kv8_d128 | 16.0272 |
018_mla_paged_decode_h16_ckv512_kpe64_ps1 | 31.9883 |
019_mla_paged_prefill_causal_h16_ckv512_kpe64_ps1 | 93.2019 |
020_moe_fp8_block_scale_ds_routing_topk8_ng8_kg4_e32_h7168_i2048 | 28.1198 |
021_rmsnorm_h128 | 2.3360 |
022_rmsnorm_h512 | 2.3813 |
023_rmsnorm_h1536 | 4.8512 |
024_rmsnorm_h2048 | 2.4907 |
025_rmsnorm_h4096 | 2.7274 |
026_rmsnorm_h7168 | 3.6834 |
Table 6. Per-kernel SOL-ExecBench FlashInfer-Bench speedups against Scoring Baseline on B200 (evaluation stack v1.1).
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 [20].
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.

Figure 8. Sources and recipients of reused hardware knowledge.
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
References
- [1]Carlo Baronio, Pietro Marsella, Ben Pan, Simon Guo, and Silas Alberti. 2025. Kevin: Multi-Turn RL for Generating CUDA Kernels. arXiv:2507.11948arxiv.org/abs/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. K-Search: LLM Kernel Generation via Co-Evolving Intrinsic World Model. arXiv:2602.19128arxiv.org/abs/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-gpu-kernel-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 Yaojian Chen, Dazhi Peng, Yiyun Zheng, Zhui Zhu, Shengye Xiang, Leshan Li, Zhen Li, Xianglong Deng, Haojun Sun, Qingfeng He, Shangzhan Li, Yifei Shen, Lin Gan, Yuxuan Li, Xu Han, and Zhiyuan Liu Preprint, 2026, arXivarxiv.org/abs/2610.36434
- [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).arxiv.org/abs/2205.14135
- [7]DeepSeek-AI. 2024. DeepSeek-V2: A Strong, Economical, and Efficient Mixture-of-Experts Language Model. arXiv:2405.04434arxiv.org/abs/2405.04434
- [8]DeepSeek-AI. 2025. DeepSeek-V3.2-Exp: Boosting Long-Context Efficiency with DeepSeek Sparse Attention. https://github.com/deepseek-ai/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.16996arxiv.org/abs/2510.16996
- [10]Yaochen Han, Ke Fan, Hongxu Jiang, et al. 2026. EGG: An Expert-Guided Agent Framework for Kernel Generation. arXiv:2606.26758arxiv.org/abs/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, Guoyang 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.arxiv.org/abs/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.14560arxiv.org/abs/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.arxiv.org/abs/2605.05023
- [16]Alexander Novikov, Ngân Vũ, Marvin Eisenberger, et al. 2025. AlphaEvolve: A coding agent for scientific and algorithmic discovery. arXiv:2506.13131arxiv.org/abs/2506.13131
- [17]NVIDIA Corporation. 2022. NVIDIA H100 Tensor Core GPU Architecture Whitepaper. https://resources.nvidia.com/en-us-hopper-architecture/nvidia-h-100-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, Cameron Bixner, Ying Zhang, Vijay Thakkar, Pradeep Ramani, and Tri Dao. 2024. FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-precision. arXiv:2407.08608arxiv.org/abs/2407.08608
- [22]Qitong Sun, Jun Han, Tianlin Li, et al. 2026. KernelClaw: A Multi-Agent Framework for GPU Kernel Optimization with Dual-Memory Mechanisms. arXiv:2603.10085arxiv.org/abs/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, Hs iang-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.DOI
- [25]Jianghui Wang, Vinay Joshi, Saptarshi Majumder, et al. 2025. GEAK: Introducing Triton Kernel AI Agent & Evaluation Benchmarks. arXiv:2507.23194arxiv.org/abs/2507.23194
- [26]Anjiang Wei, Tianran Sun, Yoqesh Seenichamy, et al. 2025. Astra: A Multi-Agent System for GPU Kernel Performance Optimization. arXiv:2509.07506arxiv.org/abs/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.arxiv.org/abs/2405.05751
- [28]Mingyi Yang, Jan Kautz, and Ali Hatamizadeh. 2024. Gated Delta Networks: Improving Mamba2 with Delta Rule. arXiv:2412.06464 ICLR 2025.arxiv.org/abs/2412.06464
- [29]Zhao 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]Zhao Ye, Ying 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.28213arxiv.org/abs/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-contest-solution. Speedups over the contest winners re-measured by the authors under the flashinfer-bench CUPTI protocol, not by the contest evaluator. Accessed 2026-09.