alphaFeedback

WorksComputer Science

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

AI contributions
Not recorded for this paper
View details
Ancillary data
1 code link
View

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.

Single-run optimization and cross-run self-evolution

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.

ForgeKernel architecture diagram

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 M0,M1,…,MkM_0, M_1, \ldots, M_k. The initial milestone M0M_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 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, M0M_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.

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 M0M_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 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-L2L2 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 S=Tref/TcandidateS = T_{\mathrm{ref}}/T_{\mathrm{candidate}}, where the numerator and denominator are the device execution times of the reference and candidate, respectively; performance gain is (S−1)×100%(S - 1) \times100\%.

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.

BenchmarkCoverageSpeedupRank
CANN-Bench v1.1<br>Ascend 910C53 operators<br>Standard: 1,060/1,060<br>Hidden: 4,185/4,2401.84×1
SOL-ExecBench<br>NVIDIA B20026/26 kernels<br>660 workloads5.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.

Performance gains over expert libraries on Hopper and Blackwell

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

Performance gains over published kernels from other generation systems

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.

Performance gains over public-best submissions on Ascend 910C

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.

OperatorInitial performance<br>vs. expertRound 1 gain<br>vs. initialRound 2 gain<br>vs. round 1Final 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 sq=4096s_q = 4096, hq=128h_q = 128, skv=32768s_{\mathrm{kv}} = 32768, and topk⁡=640\operatorname{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 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.

ConfigurationAttn bwdGDN prefill fwdDSA sparse fwdGeo. 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.

External knowledge adoption and reciprocal knowledge reuse; (a) External knowledge adoption: 0–5 entries, 6–10 entries, 11–15 entries, 16–20 entries, 21–25 entries, 26–30 entries; (b) Reciprocal knowledge reuse: DSA sparse attn fwd, GQA bwd (8:1), DSA sparse attn fwd; FP16 GEMM, MHA bwd, FP16 GEMM

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.

DSA sparse attention-forward architecture change, showing gen-0 non-persistent execution and gen-1 persistent loop with overlapping loads and KV fills

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.

OperatorGPUShape
* /goal agent ceiling*
attn_fwd_gqaH100B=2 S=4096 H_Q=16/H_KV=1 D=128 bf16 causal
fa_bwd_gqaH100B=10 N=4096 H_Q=16/H_KV=2 D=64 fp16 causal
gdn_prefill_fwdH100N=8192 h_qk=16/h_v=32 d=128 bf16
gdn_prefill_bwdH100N=8192 h_qk=16/h_v=32 d=128 bf16
dsa_indexer_fwdH100s_q=4096 s_kv=32768 H=64 D=128 fp8
dsa_indexer_topkH100logits fp32 [4096 x 32768] k=1024
dsa_sparse_fwdH100s_q=4096 s_kv=32768 h=128 d=512 topk=1024
dsa_sparse_bwdH100s_q=32768 s_kv=8192 h=128 d=512 topk=1024
dsa_score_recomputeH100s_q=32768 s_kv=8192 topk=1024
dsa_indexer_bwdH100s_q=32768 s_kv=8192 h=64 d=128 topk=1024
ForgeKernel optimization runs
attn_fwd (GQA 16:1)H100B=2 S=4096 H=16/1 D=128 bf16 causal
attn_bwd (MHA)H100B=10 N=4096 H=16 D=64 fp16 causal
attn_bwd (GQA 8:1)H100B=10 N=4096 H=16/2 D=64 fp16 causal
fp8_gemmH100M=N=K=4096 blk 128
fp8_grouped_gemmH100G=8 N=K=4096 blk 128
gemm (fp16)H100M=N=K=8192 fp16
conv2d_fwdH100N=128 C=K=256 28x28 3x3 s1
conv2d_wgradH100N=128 C=K=256 28x28 3x3 s1
conv2d_dgradH100N=128 C=K=256 56x56 3x3 s2
conv3d_fwdH100N=8 C=K=256 8x28x28 3x3x3
dsa_indexer_fwdH100s_q=4096 s_kv=32768 H=64 D=128 fp8
dsa_indexer_topkH1004096x32768 fp32 k=1024
dsa_sparse_fwdH100s_q=4096 s_kv=32768 h=128 d=512 topk=1024
dsa_sparse_bwdH100s_q=32768 s_kv=8192 h=128 d=512 topk=1024
dsa_score_recomputeH100s_q=32768 s_kv=8192 topk=1024
dsa_indexer_bwdH100s_q=32768 s_kv=8192 H=64 D=128 topk=1024
gdn_prefill_fwdH100N=8192 h=16/32 d=128 bf16
gdn_prefill_bwdH100N=8192 h=16/32 d=128 bf16
gqa_paged_decodeH100workload 7a7fc288 (61535 pages)
mla_paged_decodeH100workload 5bef8d88 (B=64; 75145 tokens)
mla_paged_prefillH100workload fe63f292 (B=18; 15883 q)
trimulH100S=1024 D=384 H=128 (case 6)
attn_bwd (MHA)B200B=8 N=4096 H=16 D=128 bf16 causal
fp4_grouped_gemmB200G=8 M=4096/group N=K=4096
sparse_mla_prefillB200s_q=4096 s_kv=32768 H_Q=128 H_KV=1 D=512 topk=1024 bf16 attn_sink
fused_attn_rope_castB200s_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_pagedB200workload 564007ac (T=8; 8462 pages)
gdn_prefill_varlenB200workload a9540651 (T=8192 N=20)
recurrent_kda_prefillB200T=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.

OperatorLevelStandardHiddenCombinedOperatorLevelStandardHiddenCombined
GeluL10.9103.2632.527DepthwiseConv2DL31.9051.1431.266
SwiGluL12.1134.0133.530TopKL31.2860.8550.927
ForeachNormL12.0883.7073.305MoeFinalizeRoutingL32.3063.2293.019
SigmoidL10.9724.6793.417Dilation2DL32.0861.1901.331
ExpL12.7174.6244.157MhcSinkhornL33.1882.6032.711
MaskedScaleL12.6524.8374.289StridedSliceL33.6673.9583.898
MishL12.3693.9653.577DequantSwigluQuantL32.5123.7443.457
ForeachAddcdivScalarL11.8044.9974.076WeightQuantBatchMatmulL31.2992.1431.939
GridSampler3DL21.3430.8610.941EngramGateFusionL30.9880.9680.972
CumminL21.2851.0321.078Conv3DBackpropFilterL31.7581.1041.212
GcdL25.7043.2963.678UniqueL31.4870.5890.709
DynamicQuantL22.2705.7544.777GroupedMatmulL31.8682.4062.287
CrossEntropyLossL22.2531.6911.791Conv2DL31.0191.0781.066
ApplyRotaryPosEmbL22.5102.2162.272AdaptiveAvgPool3DL33.4783.0743.151
GroupNormL21.5892.1322.011MoeReRoutingL32.0153.2012.918
SoftmaxL21.8402.3802.261TransposeL31.2131.9131.746
RmsNormL21.7093.0322.703ROIAlignL30.8420.6310.669
ScatterL21.0280.3630.447QuantMatmulL31.1461.8081.651
GatherL22.5122.4072.428GroupedMatmulSwigluQuantL42.4181.9672.050
UnsortedSegmentSumL21.5851.6691.652LSTML40.4500.1470.184
ResizeBilinearL25.96210.3979.302MLAL41.1391.6491.532
ArgMaxL22.9212.9262.925GRUL40.5460.2180.262
ApplyAdamWL21.9654.3783.730MlaPrologL41.5621.1981.263
MaximumL21.5704.1173.395MHAL41.1981.1121.128
NMSL33.1910.9811.242SparseFlashAttentionL41.4610.2070.306
MoeGatingTopKSoftmaxL31.9311.4781.559GQAL41.1841.4231.371
AddRmsNormDynamicQuantL32.4593.2503.074

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

KernelSpeedup (×\times)
001_fused_add_rmsnorm_h20481.4651
002_fused_add_rmsnorm_h40964.4455
003_fused_add_rmsnorm_h71681.8506
004_gemm_n128_k20482.4275
005_gemm_n256_k71682.3825
006_gemm_n2048_k40961.4514
007_gemm_n4096_k40961.3595
008_gemm_n4096_k143361.5171
009_gemm_n5120_k20481.1265
010_gemm_n6144_k40961.2387
011_gemm_n28672_k40961.3481
012_gqa_paged_decode_h32_kv4_d128_ps116.1650
013_gqa_paged_decode_h32_kv8_d128_ps19.9521
014_gqa_paged_prefill_causal_h32_kv4_d128_ps1159.3385
015_gqa_paged_prefill_causal_h32_kv8_d128_ps1133.7091
016_gqa_ragged_prefill_causal_h32_kv4_d12817.5878
017_gqa_ragged_prefill_causal_h32_kv8_d12816.0272
018_mla_paged_decode_h16_ckv512_kpe64_ps131.9883
019_mla_paged_prefill_causal_h16_ckv512_kpe64_ps193.2019
020_moe_fp8_block_scale_ds_routing_topk8_ng8_kg4_e32_h7168_i204828.1198
021_rmsnorm_h1282.3360
022_rmsnorm_h5122.3813
023_rmsnorm_h15364.8512
024_rmsnorm_h20482.4907
025_rmsnorm_h40962.7274
026_rmsnorm_h71683.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 5.54×5.54 \times [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.

Heatmap showing sources and adopting operators, with source operator labels Conv2D fwd, Conv2D input grad, Conv2D weight grad, Conv3D fwd, DSA indexer bwd, DSA indexer fwd, DSA indexer top-k, DSA score recompute, DSA sparse attn bwd, DSA sparse attn fwd, FP16 GEMM, FP8 GEMM, FP8 grouped GEMM, GDN prefill bwd, GDN prefill fwd, GQA bwd (8:1), GQA decode (paged), GQA fwd (16:1), MHA bwd, MLA decode, MLA decode (paged), MLA prefill (paged), and TriMul; adopting operator labels 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), and MHA bwd; color scale from 0 to 9

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. [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. [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. [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. [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. [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. [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. [7]DeepSeek-AI. 2024. DeepSeek-V2: A Strong, Economical, and Efficient Mixture-of-Experts Language Model. arXiv:2405.04434arxiv.org/abs/2405.04434
  8. [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. [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. [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. [11]Kimi Team. 2025. Kimi Linear: An Expressive, Efficient Attention Architecture. arXiv:2510.26692 Introduces Kimi Delta Attention (KDA).
  12. [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. [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. [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. [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. [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. [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. [18]NVIDIA Corporation. 2024. NVIDIA Blackwell Architecture Technical Overview. https://resources.nvidia.com/en-us-blackwell-architecture. Accessed 2026-09.
  19. [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. [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. [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. [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. [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. [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. [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. [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. [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. [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. [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. [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. [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. [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.

Paper details

vlnc:W_4v7pqak7
v1 · latestHistory (1)

published · Fresh — added to the Hub within the last 30 days

added ·

Usage

views · 7

downloads · 0

A view counts once per browser per day, a download each time the PDF is requested.

CC BY 4.0

Original source

  • Computer Science

Not recorded for this paper · View statement →

Looking for related papers…
Contents