KernelArc: A Multi-Agent Framework for GPU Kernel Optimization

Joyjit Kundu; Ben Stoffelen; Kaili Wang; Peter Vrancx; Ludovic Denoyer

cs.AIAug 17, 2026

KernelArc: A Multi-Agent Framework for GPU Kernel Optimization

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

Organizations: AILabs, Interuniversity Microelectronics Centre (IMEC)

Abstract

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. In the public SOL-ExecBench leaderboard snapshot recorded on August~20, 2026, KernelArc ranked first on every representative L1, L2, Quantization, and FlashInfer task evaluated. 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.

Figures & tables

Figure 1. KernelArc architecture. Strategy skills seed the agent portfolio; agents recall and record validated conclusions through shared memory, then submit candidates to a deterministic cascade gate and guard. The guard benchmarks candidates on SOL-ExecBench, returns REJECT/KEEP/ACCEPT/REVERT actions and STOP feedback, archives kept best variants, and exposes guard bests through an internal campaign leaderboard for read-only cross-agent inspection. This campaign leaderboard is distinct from the public NVIDIA SOL-ExecBench leaderboard used for external rankings. Strategy skills seed a portfolio of agents. Agents exchange wins and traps through shared memory, submit candidates to a cascade gate and guard, receive REJECT, KEEP, ACCEPT, REVERT, or STOP feedback, and use SOL-ExecBench measurements for validation. The guard may publish an internal campaign leaderboard for read-only cross-agent inspection.

Skill	Role
strat-library	cuBLASLt/cuDNN paths
strat-memory	coalescing, cache reuse
strat-compute	tensor-core use
strat-fusion	epilogue fusion
strat-precision	FP8/BF16 choices
strat-reduction	softmax, normalization

Table 1. KernelArc skills and meta-skills.

Tier	Curriculum Stage	Role in Campaign
1	Scalar warp-tiled CUDA baseline	Correct tiled starting point
2	WGMMA/TMA tensor-core pipeline	Main compute transition
3	Producer/consumer warp specialization	Overlap loads and compute
4	Persistent CTAs and raw mbarrier	Control residency and phases
5	Cluster-aware TMA multicast path	Shared-data reuse path
6	Epilogue and wait-depth refinements	Clean up completion path

Table 2. Condensed Hopper campaign curriculum. Rows follow the CUDA playbook dependency chain and summarize the campaign qualitatively, not as isolated ablations.

Figure 2. Motivating eight-hour single-agent progression: 4.6 to 766 TFLOPS across 17 accepted techniques on Hopper. The dashed line marks the strongest cuBLAS end-to-end baseline measured under the same protocol during the campaign (742 TFLOPS ). A line chart rises from 4.6 to 766 TFLOPS over 17 accepted kernel variants and crosses the 742-TFLOPS cuBLAS baseline near the end.

Figure 3. Cost–performance frontier for eight model backbones in the motivating GEMM study (best-so-far throughput versus cumulative cost; logarithmic cost axis). Claude Opus 5 attains the highest throughput (

{\sim}766\,\text{TFLOPS}

) but at the largest cumulative cost, while Kimi K3 reaches the near-cuBLAS regime most cheaply; the remaining backbones saturate at lower throughput. Eight best-so-far throughput curves versus cumulative dollar cost; Claude Opus 5 reaches the highest throughput, while Kimi K3 approaches cuBLAS at the lowest cost.

Figure 4. L1-030 SOL progression using Claude Opus 4.6. Single-agent search plateaus at 0.441279; multi-agent search reaches 0.546380, exceeding NVIDIA’s 0.500 baseline. A progression chart shows the single-agent score plateauing at 0.441279 and the multi-agent score increasing through successive submissions to 0.546380 and exceeding the NVIDIA baseline.

Problem	vs. PyTorch reference	vs. optimized baseline
L1-030	1.40 $\times$	1.06 $\times$
L2-025	291.44 $\times$	1.13 $\times$
Quant-031	1327.04 $\times$	43.79 $\times$
FI-014	61972.57 $\times$	143.78 $\times$

Table 3. Latency-ratio speedups for this work in the August 20, 2026 public NVIDIA SOL-ExecBench leaderboard snapshot 3 3 3 https://research.nvidia.com/benchmarks/sol-execbench/leaderboard . Values are computed as baseline latency divided by this work’s latency.

Figure 5. Fixed-budget optimization on FI-014 over five trajectories per configuration. (a) Faint step lines are individual incumbent-speedup trajectories; thick lines are category geometric means, and shaded regions are pointwise 95% trajectory-bootstrap intervals. (b) Candidate-100 outcomes, with individual trajectories overlaid on the geometric mean and 95% bootstrap interval (both panels use a log y-axis). Two panels compare single, multi-agent bounded-memory, and multi-agent unbounded-memory optimization. Individual trajectories and geometric-mean confidence bands are shown through 100 candidate positions; endpoint distributions show the unbounded-memory configuration with the highest geometric mean and greatest dispersion.

Configuration	Speedup ( $\times$ )	Latency (ms)
Single	142.6 [107.6, 190.8]	0.0132 [0.0116, 0.0153]
Multi-bounded	182.7 [144.5, 231.0]	0.0092 [0.0076, 0.0113]
Multi-unbounded	290.8 [197.3, 434.2]	0.0085 [0.0057, 0.0130]

Table 4. FI-014 ablation at a roughly 100-candidate budget. Speedups and latencies are geometric means with 95% trajectory-bootstrap intervals.

Appendix figures & tables1 asset

Supplementary material from the paper’s appendix.

Appendix

Figure 6. Baseline-normalized public SOL scores for the selected L1, L2, Q, and FI tasks (task SOL divided by NVIDIA baseline SOL). Values above 1.0 outperform the optimized baseline. Each of the four KernelArc submissions shown ranked first in the public leaderboard snapshot. 5 5 5 Public leaderboard snapshot recorded on August 20, 2026. A grouped bar chart shows KernelArc above the NVIDIA baseline on selected L1, L2, Quantization, and FlashInfer tasks.

GPU Kernel Optimization

GPU kernel optimization improves the speed and memory efficiency of computations by tailoring code to GPU hardware and workloads.

See all 205 papers

KernelClaw: A Multi-Agent Framework for GPU Kernel Optimization with Dual-Memory Mechanisms
Mar 10, 2026Qitong Sun, Jun Han, Tianlin Li +6GPU Kernel Optimization GPU Kernel Generation
School of Computer Science and Engineering, Beihang University, China · Zhejiang Lab, China · Nanyang Technological University, Singapore
AgentKernelArena: Generalization-Aware Benchmarking of GPU Kernel Optimization Agents
May 16, 2026Sharareh Younesian, Wenwen Ouyang, Sina Rafati +11GPU Kernel Optimization GPU Acceleration
AMD
daVinci-kernel: Co-Evolving Skill Selection, Summarization, and Utilization via RL for GPU Kernel Optimization
Jun 15, 2026Dayuan Fu, Mohan Jiang, Tongyu Wang +5RL for Code Generation GPU Kernel Optimization
1SII · 2SJTU · 3GAIR

Mar 10, 2026J/K move · Enter open · S save

KernelClaw: A Multi-Agent Framework for GPU Kernel Optimization with Dual-Memory Mechanisms

Qitong Sun, Jun Han, Tianlin Li, Zhe Tang+5

School of Computer Science and Engineering, Beihang University, China · Zhejiang Lab, China · Nanyang Technological University, Singapore

GPU Kernel Optimization GPU Kernel Generation Large Language Model-Guided Optimization Agent Skill Learning

Open paper AskPDF arXiv

Improving GPU kernel efficiency is crucial for efficient AI systems. Recent work has explored leveraging large language models (LLMs) for GPU kernel generation and optimization. However, existing LLM-based kernel optimization pipelines typically rely on opaque, implicitly learned heuristics within the LLMs to determine optimization strategies. This leads to inefficient trial-and-error and weakly interpretable optimizations. Our key insight is to replace implicit heuristics with expert optimization knowledge that is aware of task trajectories. Specifically, we present KernelClaw, a multi-agent framework with a dual-level memory architecture. KernelClaw operates by coordinating agents with long-term memory of reusable expert knowledge and short-term memory to prevent repetitive backtracking. On KernelBench Levels 1-3, KernelClaw achieves a 100% success rate and average speedups of 6.98x, 7.89x, and 1.94x over Torch Eager on Levels 1, 2, and 3, respectively, outperforming prior baselines. Code is available at https://github.com/0satan0/KernelMem/.

Abstract

GPU Kernel Optimization

Explore similar work

KernelClaw: A Multi-Agent Framework for GPU Kernel Optimization with Dual-Memory Mechanisms

AgentKernelArena: Generalization-Aware Benchmarking of GPU Kernel Optimization Agents

daVinci-kernel: Co-Evolving Skill Selection, Summarization, and Utilization via RL for GPU Kernel Optimization