Source-linked AI summary
KernelArc: A Multi-Agent Framework for GPU Kernel Optimization
Joyjit Kundu, Ben Stoffelen, Kaili Wang, Peter Vrancx, Ludovic Denoyer
TL;DR
GPU kernel optimization is a fragmented search problem in which sequential trajectories can plateau. KernelArc coordinates strategy-specialized agents around validated conclusions and measurement-gated benchmarking, and shared-memory multi-agent runs achieved 2.04× the Single geometric-mean speedup at roughly 100 candidates.
Problem
Modern GPU optimization requires coordinating hardware mechanisms across layouts, memory, registers, synchronization, and launch overheads, while sequential search can plateau before other implementation families are explored.
Method
KernelArc runs strategy-specialized agents in parallel, sharing validated conclusions under a deterministic benchmark guard with read-only sibling state and plateau-triggered strategy changes.
Results
At a roughly 100-candidate budget, Multi-unbounded achieved 2.04× the Single geometric-mean speedup across the ablation runs.
Takeaways & Limitations
The evidence supports shared multi-agent search for broadening exploration beyond isolated trajectories, while coordination features matter differently across kernels and optimization regimes.
Takeaways & Limitations
The case studies do not estimate a suite-wide win rate across all 235 problems, and reported rankings are time-specific.
Abstract
from arXiv · showhide
We present KernelArc, a multi-agent framework for autonomous GPU kernel optimization across heterogeneous workloads. Strategy-specialized agents run in parallel and coordinate through conclusions-only shared memory, a deterministic benchmark guard, and read-only cross-agent state with plateau-triggered drafting. We evaluate \kernelarc{} on NVIDIA H100 and B200 GPUs using category-representative SOL-ExecBench workloads. The resulting implementations span custom BF16 GEMM, static cuBLASLt Expert-API configuration tables, fused mixture-of-experts backward, shape-gated decoder-layer fusion, native NVFP4 grouped-query attention, and paged prefill attention. At the public SOL-ExecBench leaderboard snapshot recorded on July~30, 2026, these submissions ranked first on representative L1, L2, Quantization, and FlashInfer tasks. The trajectories support the paper's central motivation: shared multi-agent search can broaden exploration and reach stronger incumbents within a fixed candidate budget, while the value of individual coordination features depends on the kernel and optimization stage.
1 Introduction
Modern GPU performance engineering requires coordinating increasingly complex hardware mechanisms across memory, layouts, registers, synchronization, and launch overheads. KernelArc addresses this fragmented-search regime with parallel strategy-specialized agents, validated conclusions-only coordination, deterministic benchmarking, and cross-agent state.
- Motivation: Modern accelerators expose high-throughput primitives, but speedups require coordinating layouts, memory hierarchy, register pressure, synchronization, and launch overheads.The passage names wgmma.mma_async, TMA transfers, asynchronous barriers, SM100 instructions, and NVFP4 formats.
- Motivation: A single agent advances one incumbent under one local history, so a fixed budget may cover fewer algorithm families and over-refine a design before moving on.This limitation motivates broader concurrent exploration rather than relying only on local sequential pivots.
- KernelArc architecture: KernelArc runs strategy-specialized agents concurrently and coordinates them through validated conclusions-only memory, a deterministic benchmark guard, plateau-triggered drafting, and read-only cross-agent state.Shared-memory retention is configurable, while correctness, benchmarking, and keep/revert decisions are assigned to the deterministic guard.
- Evidence: 766 TFLOPS (BF16) was reached on one fixed shape, 3.2% above the matched cuBLAS baseline measured during the campaign, using an eight-hour single-agent loop.The passage characterizes this as depth along one narrow, Hopper-specific GEMM playbook-guided path rather than a general multi-shape conclusion.
- Evaluation: KernelArc evaluation spans all four SOL-ExecBench workload categories and implementations including custom BF16 GEMM, static cuBLASLt tables, fused MoE backward, decoder-layer fusion, NVFP4 attention, and paged prefill attention.The listed workloads demonstrate coverage across common neural-network operators.
2 Related Work
Related work covers LLM-based kernel generation, search-based code optimization, multi-agent optimization architectures, and compiler or auto-tuning approaches. KernelArc is positioned within these lines through its observe–edit–evaluate–retain loop, multi-agent orchestration, and source-level GPU optimization scope.
- LLM-based kernel generation: KernelBench provides more than 250 GPU kernels for evaluating LLM-generated code, while AutoComp targets portable optimization across diverse tensor accelerators.AutoComp pairs hardware-specific optimization agents with evaluation backends for NVIDIA GPUs, TPUs, AWS Trainium, Gemmini, RISC-V vector processors, and Apple Silicon.
- LLM-based kernel generation: KernelArc studies an observe–edit–evaluate–retain loop as its single-agent boundary case and adds Hopper-specific GEMM guidance plus shared multi-agent orchestration.The passage also reports 2.82× on softmax.
- Search-based code optimization: Weco frames machine-learning engineering as code optimization through AIDE’s tree search, then extends the approach with recursive self-improvement of the agent harness.KernelArc is described as complementary because it keeps the agent harness, guard, and launcher fixed.
- Auto-tuning and compiler optimization: TVM and Ansor optimize compiler-defined schedule or IR spaces using tiling, unrolling, vectorization, tensorization, and operator fusion.The passage contrasts these spaces with arbitrary source-level rewrites, hardware-specific control-flow restructuring, and low-level CUDA/PTX idioms such as manual warp specialization.
3 Methodology
KernelArc formulates GPU kernel optimization as deterministic, correctness-constrained search over structural and schedule-level implementations. Its methodology combines parallel strategy-specialized exploration, conclusions-only memory, deterministic evaluation and acceptance, and plateau escape through optional read-only cross-agent state.
- Optimization objective: The search starts from a seed implementation, workload set, target hardware/software environment, and deterministic evaluator, maximizing benchmark score subject to correctness across every workload.Candidates failing any workload are infeasible regardless of speed.
- Optimization objective: The candidate space includes structural transformations—library replacement, DSL changes, epilogue fusion, and warp specialization—as well as parameters from fixed schedule templates.This broadens search beyond tuning a single predetermined schedule.
- Iteration protocol and plateau escape: Plateau escape optionally lets agents inspect faster sibling archives read-only while recalling conclusions, selecting eligible strategies, editing independently, and submitting candidates through the guard.The architecture maps diversification, transferable conclusions, deterministic state, and explicit plateau escape to specialized agents, memory, guard mechanisms, and drafting.
- Strategy-specialized parallel search: Eight overlapping strategy skills seed parallel agents with distinct optimization lenses, while eligibility rules prevent irrelevant strategies from being applied to particular kernel types.Skills are intended to prevent concurrent agents from collapsing onto the first plausible implementation family.
- Conclusions-only memory: Shared memory stores reusable wins and traps rather than iteration traces, with configurable retention limiting stale or low-value context.The default configuration retains 16 wins per kernel type by descending speedup and 16 traps per kernel type in FIFO order.
- External deterministic guard: A deterministic guard validates, benchmarks, and accepts candidates: only correctness-passing improvements meeting the configured margin become KEEP states, while plateau, target, budget, or cap conditions emit STOP.The guard mediates all kernel changes, and agents cannot decide whether their own edits improve the incumbent.
4 Experimental Results
The experiments show how far a guided single agent can optimize Hopper GEMM, while SOL-ExecBench trajectories demonstrate that multi-agent search can escape plateaus and outperform single-agent and bounded coordination under a fixed candidate budget. Results also show that model choice affects optimization cost and that multi-agent gains remain exploratory and variable across trajectories.
- Motivating GEMM campaign: Throughput rose from 19.3 to 541 TFLOPS across WGMMA, asynchronous-copy, tile-scaling, and software-pipeline milestones.Warp specialization reached 617 TFLOPS, while epilogue and wait-depth changes reached 720 TFLOPS before campaign-specific refinements supplied the final 46 TFLOPS.
- Motivating GEMM campaign: 766 TFLOPS was reached after 17 accepted techniques, exceeding the 742 TFLOPS strongest cuBLAS end-to-end baseline under the same protocol.The single-agent campaign progressed from 4.6 to 766 TFLOPS during eight hours on Hopper.
- Motivating GEMM campaign: Kimi K3 exceeded 700 TFLOPS after a few dollars and plateaued near 718 TFLOPS, while Claude Opus 5 reached approximately 766 TFLOPS at greater cost.Other model backbones saturated lower, indicating a productive but narrow playbook whose cost and depth depend on the model.
- SOL-ExecBench trajectory: SOL 0.481 was the final L1-030 result after multi-agent search escaped the single-agent plateau at SOL 0.441 and improved through intermediate scores of 0.457 and 0.472.The gains came from library configuration, fusion, and measurement refinements rather than one large algorithmic substitution.
- Coordination ablation: At a roughly 100-candidate budget, Multi-unbounded achieved 2.04× Single and 1.59× Multi-bounded geometric-mean speedups, while Multi-bounded achieved 1.28× Single.Configurations overlapped early, and individual runs plateaued at different levels.
- Coordination ablation: The geometric standard-deviation factor was 1.67 for Multi-unbounded, compared with 1.35 for Multi-bounded and 1.44 for Single.The updated five-seed analysis strengthens the contrast but remains exploratory because early kept edits can steer later proposals into different local regions.
5 Analysis and Discussion
KernelArc’s analysis contrasts deep, GPU-specific guidance with broad multi-agent coverage: detailed playbooks can optimize fixed kernels, while shared strategies address operator- and shape-dependent search. The case studies remain limited by time-specific rankings, incomplete suite coverage, heterogeneous ablations, and feature effects that vary by kernel and optimization regime.
- Search phases and conclusions-only memory: Search evolves from diverse early exploration to viable implementation families and mature local refinements in configuration, launches, or data movement.These mature states are described as search-process fixed points.
- Playbook depth versus breadth: 766 TFLOPS shows that detailed, prescriptive guidance can optimize one fixed GEMM, whereas multi-shape tasks require broader operator- and shape-dependent coverage.Both one-agent and shared multi-agent playbooks guide search without supplying final implementations.
- Limitations: The case studies do not estimate a suite-wide win rate across all 235 problems, and their rankings are time-specific.The Hopper–Blackwell contrast also changes both workload and hardware.
- Limitations: The ablation uses few heterogeneous runs and evaluated candidates rather than wall time, tokens, or dollar cost.This limits interpretation of coordination overhead and efficiency.
- Limitations: Coordination features are difficult to isolate because their value may depend on the kernel and optimization regime, including diversification, plateau escape, or late-stage refinement.The passage identifies these as distinct optimization stages.
6 Conclusion
KernelArc frames GPU-kernel optimization as coordinated, measurement-gated search, combining strategy-specialized exploration with validated shared conclusions. The framework supports diverse workloads, including library-backed kernels, fused operators, quantized and paged attention, and custom GEMM development.
- Coordination: KernelArc combines strategy skills, a deterministic guard, and reusable win/trap memory to coordinate optimization through validated conclusions.These components let agents explore different optimization lenses while sharing measured results.
- Workload coverage: The framework supports library-backed kernels, fused operators, quantized attention, paged prefill attention, and custom GEMM development.These workload categories are represented across the evaluated SOL-ExecBench tasks.
- Ablation: The ablation study supports KernelArc’s coordination hypothesis at a roughly 100...The supplied passage truncates the quantitative ablation result, so no complete figure is reported here.
A Appendix: Cross-Category Kernel and Public-Ranking Details
The appendix details cross-category GPU kernel implementations and reports time-specific public leaderboard results. Solutions use specialized fusion, quantization, batching, and shape-gated execution strategies across L1, L2, Quantization, and FlashInfer tasks.
- Public-ranking snapshot: At the July 30, 2026 snapshot, L1-030, L2-025, L2-053, Q-031, and FI-014 ranked first5 on their respective tasks.Current public ranks as of Aug. 17, 2026 were L1-030: 3rd; L2-025: 1st; L2-053: 1st; Q-031: 1st; FI-014: 1st.
- L1-030: Attention output projection with residual: L1-030 aliases the destination to the residual tensor and uses cached per-M cuBLASLt Expert-API plans with heuristic fallback.The implementation avoids autotuning during timed evaluation and uses one measured heuristic-rank override when available.
- L2-025: Mixture-of-experts backward: L2-025 fused MoE backward combines counting-sort, padded expert-major layouts, TF32 batched GEMMs, gradient reuse, and fused activation backward.The final SOL was 0.535 (1st5, 16/16 shapes faster than baseline, 1.13× baseline, 291× reference).
- FI-014: Paged prefill causal GQA: FI-014’s final paged-prefill design combines lean Triton attention with a split-causal cuDNN path for large-cache workloads.The final SOL was 0.986 (1st5, 30/30 shapes faster than baseline, 143.8× baseline, 61,973× ref).