🤖 AI

KernelArc: A Multi-Agent Framework for GPU Kernel Optimization

KernelArc is a multi-agent framework that autonomously optimizes GPU kernels through parallel, strategy-specialized agents coordinating via shared memory and deterministic guards, achieving top rankings on the SOL-ExecBench leaderboard for diverse workloads on NVIDIA H100 and B200 GPUs.

Original authors: Joyjit Kundu, Ben Stoffelen, Kaili Wang, Peter Vrancx, Ludovic Denoyer

Published 2026-08-19

📖 6 min read🧠 Deep dive

CC BY 4.0

Original authors: Joyjit Kundu, Ben Stoffelen, Kaili Wang, Peter Vrancx, Ludovic Denoyer

Original paper licensed under CC BY 4.0 (http://creativecommons.org/licenses/by/4.0/). ✨ This is an AI-generated explanation of the paper below. It is not written or endorsed by the authors. For technical accuracy, refer to the original paper. Read full disclaimer

Modern computers rely on specialized chips called graphics processing units, or GPUs, to handle the massive calculations required by artificial intelligence. These chips are incredibly powerful, but they are also complex machines with many different parts that must work in perfect sync. To get the most out of them, engineers must write tiny, highly specific programs called kernels that tell the chip exactly how to move data and perform math. For years, this has been a job for human experts who spend countless hours tweaking code, balancing memory use, and timing operations to squeeze out every bit of speed. As these chips have grown more sophisticated, the manual work has become harder, and the gap between what the hardware can do and what the software achieves has widened.

Researchers have begun using large language models—computer programs trained on vast amounts of text—to help write and improve this code. These models can suggest changes, test them, and learn from the results, acting like an automated engineer. However, a single automated agent often gets stuck. It might find a good solution and keep polishing it, missing other, better approaches that lie in a different direction. It is like a hiker who finds a nice path up a mountain and keeps walking up it, never realizing that a different trail nearby leads to a much higher peak. To solve this, a team of researchers at IMEC has developed a new system called KernelArc, which uses a group of these automated agents working together to explore many different paths at once.

The KernelArc system operates by assigning different strategies to several agents running in parallel. Instead of one agent trying to solve a problem alone, each agent in the group focuses on a different angle of attack, such as changing how data is stored, altering the mathematical precision used, or fusing multiple steps into one. These agents do not share their entire thought process or every step they take, which would be too messy and confusing. Instead, they share only their final conclusions: what worked and what failed. They write these results into a shared memory space that acts like a bulletin board. If one agent discovers a trick that speeds up a calculation, it posts the result. Other agents can read this post and use the insight to guide their own work, avoiding dead ends and building on success.

To ensure that the agents do not waste time on broken code, the system includes a strict, automated referee. This referee runs every new code suggestion through a series of tests to check if it is correct and how fast it is. If a suggestion fails the test, it is discarded immediately. If it works but is not faster than the current best version, it is kept as a backup but does not replace the leader. Only when a new version is both correct and faster does it become the new standard. This process allows the team of agents to explore a wide range of possibilities without getting lost in the details of every single attempt. The system also has a safety mechanism that kicks in if an agent gets stuck on a plateau, where no further improvements can be found. When this happens, the system forces the agent to try a completely different approach, ensuring the search continues to move forward.

The researchers tested this system on two of the most advanced GPUs available, the NVIDIA H100 and the B200. They focused on a set of standard tasks used to measure performance, which include operations like matrix multiplication, attention mechanisms used in language models, and various forms of data fusion. In one specific test involving a single agent working alone with a detailed guide, the system managed to reach a speed of 766 teraflops, which is about 3.2 percent faster than the best existing library for that specific task. This showed that a single agent could go very deep into a specific problem if given a clear path. However, when the researchers switched to the multi-agent KernelArc system to tackle a broader set of tasks, the results were even more impressive. The system produced custom implementations for various complex operations, including specialized attention mechanisms and fused layers for large language models.

On a public leaderboard that ranks the speed of these operations across many different shapes and sizes of data, the KernelArc submissions took first place in several categories. For one task involving attention and residual addition, a single agent got stuck at a performance score of 0.441. The multi-agent system, by sharing insights and exploring different directions, broke through that barrier and reached a score of 0.481. In another test involving a complex attention task, the multi-agent system achieved a speedup of nearly 291 times compared to a standard reference implementation, and over 143 times faster than a highly optimized baseline. These results suggest that having multiple agents share their findings allows the system to find better solutions faster than a single agent working alone, especially when the problem space is large and varied.

The study also looked at how the system behaves when the number of agents and the amount of shared memory change. They found that giving the agents a shared memory where they could post their successes and failures helped them reach stronger results within a fixed number of attempts. When the memory was unlimited, the system performed the best, reaching a speedup of over 290 times compared to the starting point. This indicates that the ability to learn from the collective experience of the group is a key factor in success. The researchers noted that the value of each feature, such as the shared memory or the strategy specialization, depends on the specific task and the stage of the search. Sometimes the agents need to explore broadly, and other times they need to refine a specific solution.

This work does not claim to have solved all problems in computer optimization, nor does it suggest that human engineers are no longer needed. The results are specific to the tasks and hardware tested, and the system still relies on a human-designed framework to guide the agents. However, the findings demonstrate that a coordinated group of automated agents can explore a wider range of solutions than a single agent. By sharing only their most valuable conclusions, these agents can avoid repeating mistakes and build on each other's discoveries. This approach offers a promising way to handle the increasing complexity of modern computing, where the gap between hardware potential and software performance continues to grow. The system shows that with the right coordination, automated tools can help unlock the full power of the chips that drive the next generation of artificial intelligence.

Technical Summary: KernelArc

Problem Statement

Extracting peak performance from modern GPU accelerators (specifically NVIDIA Hopper and Blackwell architectures) is increasingly complex due to the need to coordinate high-throughput primitives—such as Warp Group Matrix Multiply Accumulate (wgmma.mma_async), Tensor Memory Accelerator (TMA) transfers, and NVFP4 formats—across memory hierarchies, register pressure, and synchronization barriers. While Large Language Model (LLM) agents have shown promise in automating kernel optimization (e.g., AutoKernel), existing single-agent systems face limitations. A single agent traverses a focused, human-designed "playbook" sequentially; when it encounters a performance plateau, it advances only one incumbent under a single local history. This approach may over-refine a specific design before moving on, failing to explore diverse algorithm families or data layouts within a fixed computational budget. Consequently, single agents may mistake a search-trajectory ceiling for a hardware ceiling.

Methodology

KernelArc introduces a multi-agent framework designed to broaden exploration and escape local plateaus through coordinated, strategy-specialized parallel search. The architecture relies on four core mechanisms:

Strategy-Specialized Parallel Agents: Instead of a single agent, KernelArc deploys multiple agents, each seeded with specific "strategy skills" (e.g., library selection, memory coalescing, compute optimization, fusion, precision, reduction, scheduling, and custom SM100 kernels). These skills act as overlapping search lenses rather than mutually exclusive paths, preventing agents from collapsing onto the first plausible solution family.
Conclusions-Only Shared Memory: Agents exchange information via a shared memory store that is strictly limited to validated conclusions. This store contains:
- Wins: Successful optimizations with speedup metrics, shape context, and a reflection on why it worked.
- Traps: Dead ends with error descriptions and reflections on why they failed.
- Exclusions: The memory explicitly excludes iteration counters, heartbeats, and progress logs to prevent context pollution. A configurable retention horizon (e.g., keeping the top 16 wins/traps per kernel type) ensures the memory remains compact and relevant.
Deterministic Benchmark Guard: A separate, deterministic process mediates all kernel changes. The agent proposes code, but the guard owns the decision-making regarding correctness and performance. The guard performs:
- Cascade Gate: Cheap pre-benchmark validation (syntax, entry points, portability).
- Benchmarking: Execution via the SOL-ExecBench bridge under a GPU lock to measure aggregate latency.
- State Management: Decisions to KEEP (improvement), ACCEPT (near-best working point), REVERT (regression), or REJECT (failure).
Plateau-Triggered Drafting: Agents track consecutive non-improving benchmarks. Upon reaching a threshold ( $r_{draft}$ ), an agent is triggered to "draft" a fundamentally different algorithm, DSL family, or data layout rather than continuing local refinements. Agents can also inspect a read-only "campaign leaderboard" of sibling solutions to identify stronger incumbents.

The system separates generative reasoning (LLMs) from stateful evaluation and coordination (deterministic code), ensuring that the search process remains robust and reproducible.

Key Contributions

The KernelArc Architecture: A novel multi-agent system featuring strategy-specialized parallel agents coordinated by conclusions-only memory, a deterministic guard, and plateau-triggered drafting. It distinguishes between reusable conclusions (LLM-visible) and operational state (deterministic code).
Evidence for Shared Multi-Agent Search: Empirical trajectories demonstrate that sharing validated conclusions allows the system to broaden search beyond isolated local trajectories, reaching stronger incumbents within a fixed candidate budget.
Efficient Kernels Across Diverse Axes: The framework successfully generated implementations spanning PTX-assisted custom BF16 GEMM, static cuBLASLt configuration tables, fused Mixture-of-Experts (MoE) backward, shape-gated decoder-layer fusion, native NVFP4 attention, and paged prefill attention.

Experimental Results

The framework was evaluated on NVIDIA H100 and B200 GPUs using the SOL-ExecBench (Speed-of-Light Execution Benchmark), which aggregates performance across multiple input shapes.

Single-Agent Baseline (GEMM Study): In a controlled, playbook-guided study on a fixed 4096×4096 BF16 GEMM shape on H100, a single agent reached 766 TFLOPS (3.2% above the matched cuBLAS baseline of 742 TFLOPS). This demonstrated the depth achievable by a focused agent but highlighted the limitations of a single trajectory.
Multi-Agent Performance (SOL-ExecBench):
- On the L1-030 task (attention output projection with residual addition), a single agent plateaued at a SOL score of 0.441. The multi-agent KernelArc system escaped this local minimum, reaching 0.481, securing the first-place rank on the public leaderboard snapshot (July 30, 2026).
- Across four workload categories (L1, L2, Quantization, FlashInfer), KernelArc submissions achieved first-place rankings on representative tasks, including L1-030, L2-025 (MoE backward), Q-031 (NVFP4 grouped-query attention), and FI-014 (paged prefill attention).
Ablation Studies (FI-014): On a fixed budget of 100 candidates, the Multi-unbounded configuration (two agents with shared memory and no retention cap) achieved a geometric-mean speedup of 2.04× over the Single-agent configuration and 1.59× over the Multi-bounded configuration. This supports the hypothesis that shared memory helps agents avoid repeated dead ends and redirect search toward stronger solutions.

Significance and Claims

The paper claims that KernelArc validates the central motivation that shared multi-agent search can broaden exploration and reach stronger incumbents within a fixed candidate budget, a capability that sequential single-agent systems lack.

The authors emphasize that the value of specific coordination features (e.g., memory retention, drafting triggers) is context-dependent, varying by kernel type and optimization stage (early exploration vs. late refinement). The system does not claim to universally outperform vendor libraries in all scenarios; rather, it demonstrates that a structured, measurement-gated multi-agent approach can effectively navigate the fragmented search space of modern GPU optimization, producing state-of-the-art results across diverse operator families and precision formats. The results suggest that measured coordination is a viable path for autonomous kernel optimization, though the paper notes that isolating the causal effect of individual mechanisms remains an open question for larger controlled studies.