KernelArc：一个用于GPU内核优化的多智能体框架

# 多智能体GPU内核优化框架
来源：https://arxiv.org/html/2608.17071
## KernelArc：面向GPU内核优化的多智能体框架
CCS：软件及其工程 自动化编程  
CCS：计算机系统组织 并行架构

Joyjit Kundu\*Ben Stoffelen Kaili Wang Peter Vrancx Ludovic Denoyer

###### 摘要\.

我们提出KernelArc，一个用于在异构工作负载中实现自主GPU内核优化的多智能体框架。策略专精的智能体并行运行，并通过仅共享结论的共享内存、确定性基准测试守卫以及具有平台期触发重起草机制的只读跨智能体状态进行协调。我们使用SOL-ExecBench基准测试中具有类别代表性的工作负载，在NVIDIA H100和B200 GPU上评估了KernelArc。最终实现覆盖了自定义BF16 GEMM、静态cuBLASLt专家API配置表、融合混合专家反向传播、形状门控解码器层融合、原生NVFP4分组查询注意力以及分页预填充注意力。在2026年7月30日记录的公开SOL-ExecBench排行榜快照中，这些提交在具有代表性的L1、L2、量化和FlashInfer任务中排名第一。这些轨迹支持本文的核心动机：共享的多智能体搜索能够在固定的候选预算内拓宽探索范围并达到更优的基准，而各协调特性的价值取决于具体内核和优化阶段。

###### 关键词：

GPU内核优化，多智能体优化，LLM智能体，自主研究

11脚注文本：邮箱：joyjit\.kundu@imec\.be## 1\.引言

从现代GPU加速器中提取峰值性能正变得愈发困难。Hopper和Blackwell架构提供了高吞吐量原语——包括Warp组矩阵乘累加（wgmma\.mma\_async）、张量内存加速器（TMA）传输、异步屏障、SM100指令以及NVFP4格式——但要获得有效的加速，需要在布局、内存层次结构、寄存器压力、同步和启动开销等方面协调这些机制。

LLM智能体为手动性能工程提供了一种可编程的替代方案。它们可以对内核进行分析、提出结构性转换、编辑CUDA/PTX或特定领域语言代码、选择并自动调优cuBLASLt、cuDNN和CUTLASS等库、选择精度和量化格式，并通过正确性测试和基准测试形成闭环。现有的单智能体系统如AutoKernel（jaber2026autokernel (https://arxiv.org/html/2608.17071#bib.bib1)）表明此闭环可行。然而，其成功依赖于经过精心设计的人工操作手册和足够的顺序预算来遍历这些步骤。即使单智能体在平台期做出调整，它也仅在单一局部历史下推进一个基准设计，因此固定预算覆盖的算法族更少，可能在继续推进之前过度优化某个设计。

我们引入KernelArc，一种面向这种碎片化搜索场景的多智能体架构。策略专精的智能体同时探索不同的优化方向。它们仅通过共享内存交换经过验证的结论，该内存的保留期限可配置，同时一个确定性守卫负责管理正确性、基准测试以及保留/回滚决策。只读的跨智能体状态暴露了更强的同级解决方案以供检查，而平台期触发器会要求智能体尝试不同的算法、特定领域语言族或数据布局，而不是继续局部优化。观察到的轨迹与共享的预期作用相符：配置的系统突破了单智能体的平台期，而固定预算的消融实验倾向于共享内存的多智能体配置。我们将其作为系统级模式报告，因为每个协调特性的价值可能取决于具体内核以及搜索是处于早期探索、平台期突破还是后期局部优化阶段。这种分工将生成性推理交由LLM处理，而有状态评估和协调则交由确定性代码处理。

在八小时的实际时钟预算下，使用详细的针对Hopper架构的通用矩阵乘法（GEMM）优化手册，单智能体自主研究循环在一个固定形状上达到了766TFLOPS（BF16），比活动期间测量的匹配cuBLAS基线高出3.2%。这证明了在一条狭窄的手册引导路径上具有深度；它并未确立一个比cuBLAS更快的通用GEMM内核，也未体现优化跨多个形状聚合分数的能力。要达到766TFLOPS，需要一系列正确的底层Hopper优化，包括WGMMA/TMA流水线、同步、寄存器压力管理和后处理调度。为了研究更广泛的场景，我们使用光速执行基准测试（SOL-ExecBench）（solexecbench2025 (https://arxiv.org/html/2608.17071#bib.bib2)），其SOL分数聚合了跨输入形状的性能（越高越好）。在L1-030（一个在16个形状上评估的带有残差相加的注意力输出投影）上，顺序运行在SOL 0.441处进入平台期，而配置的KernelArc系统达到了SOL 0.481。111在此快照之后，2026年8月8日记录的最新L1-030运行排名第三，达到了SOL 0.490500，延迟为0.036174毫秒。这支持了完整的共享多智能体设计，而无需每个协调特性在搜索的每个点都发挥作用。

我们在SOL-ExecBench（solexecbench2025 (https://arxiv.org/html/2608.17071#bib.bib2)）所有四类工作负载的任务上评估了KernelArc架构，涵盖了常见的神经网络算子。结合单智能体GEMM配置，最终实现包括PTX辅助的自定义BF16 GEMM实现（与匹配的cuBLAS基线进行基准测试）、静态cuBLASLt配置表、融合混合专家（MoE）反向传播、形状门控解码器层融合、原生NVFP4注意力和分页预填充注意力，展示了跨优化轴的迁移能力，而非单一编码的内核族。详细的内核和公开排名结果见附录A（https://arxiv.org/html/2608.17071#A1）。

贡献\.

1. \(1\) KernelArc架构：策略专精的并行智能体通过具有可配置保留期限的仅共享结论内存进行协调，辅以确定性基准测试守卫、平台期触发重起草和只读跨智能体状态；可重用的结论对LLM保持可见，而尝试计数、停止阈值、重启动逻辑和内部排行榜发布则保留在确定性代码中。
2. \(2\) 共享多智能体搜索的证据：观察到的轨迹支持本文的核心动机，即共享经过验证的结论可以拓宽搜索范围，超越孤立的局部轨迹，同时留待解答在哪些内核场景下哪些协调特性最为重要。
3. \(3\) 跨不同优化轴的高效内核：KernelArc产生的实现范围从PTX辅助BF16 GEMM和静态cuBLASLt配置表到融合MoE反向传播、形状门控解码器层融合、原生NVFP4注意力和分页预填充注意力；提交的SOL-ExecBench内核在公开排行榜快照中获得了第一名。222公开排行榜排名记录于2026年7月30日。

## 2\.相关工作

#### 基于LLM的内核生成\.

多个系统使用LLM进行一次性或迭代式内核生成。KernelBench（ouyang2025kernelbench (https://arxiv.org/html/2608.17071#bib.bib3)）提供了超过250个GPU内核，用于评估LLM生成的代码。AutoComp（hong2025autocomp (https://arxiv.org/html/2608.17071#bib.bib4)）针对跨张量加速器的可移植LLM驱动内核优化，将硬件特定的优化智能体与针对NVIDIA GPU、TPU、AWS Trainium、Gemmini、RISC-V向量处理器和Apple Silicon等目标的评估后端配对。AutoKernel（jaber2026autokernel (https://arxiv.org/html/2608.17071#bib.bib1)）引入了一个自主智能体循环：分析完整工作负载以识别瓶颈，根据阿姆达尔定律对其排序，并通过五阶段正确性测试工具和六层优化手册迭代优化Triton或CUDA C++内核。在其H100 FP16内核基准测试中，AutoKernel报告了在RMSNorm上比PyTorch eager模式最高快5.29倍，在softmax上快2.82倍。KernelArc研究了一个结构相似的观察-编辑-评估-保留循环作为其单智能体边界情况，然后为GEMM研究添加了Hopper特定的指导，并为更广泛的工作负载添加了共享多智能体搜索编排。

#### 基于搜索的代码优化\.

Weco首次引入了AIDE，这是一个将机器学习工程构建为代码优化、并将反复试错构建为对候选解决方案进行*树搜索*的系统，从有前景的节点进行分支并剪枝失败方案（jiang2025aide (https://arxiv.org/html/2608.17071#bib.bib5)）。Weco的第二篇报告将这一思路扩展到递归自我改进，使用外部自主研究循环来重写和评估内部循环的智能体工具本身（weco2026aide2 (https://arxiv.org/html/2608.17071#bib.bib6)）。KernelArc是互补的：它保持智能体工具、守卫和启动器固定，并研究在严格的正确性和基准测试约束下，共享状态和策略专化如何影响GPU内核搜索。

#### 多智能体优化架构\.

现有系统涵盖了扁平角色专化、层次化归档、束搜索和通用编排。Astra（wei2025astra (https://arxiv.org/html/2608.17071#bib.bib7)）使用五个角色专化的智能体，没有持久记忆；AKO4X（ako2026 (https://arxiv.org/html/2608.17071#bib.bib8)）使用跨会话归档和死胡同目录；AccelOpt（zhang2025accelopt (https://arxiv.org/html/2608.17071#bib.bib9)）结合了宽度为4的束搜索和经验蒸馏；CORAL（qu2026coral (https://arxiv.org/html/2608.17071#bib.bib10)）提供了工作树隔离、共享状态和心跳触发的调整。它们报告的加速比在任务、硬件或预算之间无法直接比较。KernelArc结合了策略专化、具有可配置保留期限的仅共享结论内存、外部确定性守卫以及平台期重起草。

#### 自动调优与编译器优化\.

传统的自动调优器和编译器栈，例如TVM（chen2018tvm (https://arxiv.org/html/2608.17071#bib.bib11)）及其Ansor自动调度器（zheng2020ansor (https://arxiv.org/html/2608.17071#bib.bib12)），在编译器定义的调度或中间表示空间内进行优化，包括平铺、循环展开、向量化、张量化和算子融合等选择。这些空间功能强大，但通常不支持任意源代码级重写、特定于硬件的控制流重构或底层的CUDA/PTX惯用写法（如手动Warp专化）。KernelArc则直接搜索内核源代码。

## 3\.方法论

### 3\.1\.测量门控的内核优化

自主内核优化从种子实现\(k_0\)、工作负载集\(W\)、目标硬件/软件环境\(H\)和确定性评估器开始。搜索目标是

\(1\)k^{\star}=\arg\max_{k\in\mathcal{K}}S(k;W,H)\quad\text{s.t.}\quad C(k,w)=1\;\;\forall w\in W,\)
其中\(\mathcal{K}\)是源代码级实现空间，\(C\)是正确性谓词，\(S\)是基准测试分数（例如给定问题所有工作负载的几何平均延迟的倒数）。约束\(C(k,w)=1\)要求候选\(k\)在工作负载情况\(w\)下通过评估器的正确性检查；无论速度如何，任何\(w\in W\)不通过的候选都是不可行的。此处，\(\mathcal{K}\)包括结构性变更——例如用供应商库调用替换框架算子、更改内核特定领域语言、融合后处理或引入Warp专化——以及来自固定调度模板的参数。

智能体观察当前实现和测量结果，形成瓶颈假设，编辑内核，并将候选提交给评估器。改进基准的正确候选被保留为接受的最佳候选；未改进的候选不会提升该最佳，尽管KernelArc可以暂时保留接近最佳的横向候选作为工作点。评估器必须在目标环境中运行，因为编译器版本、库启发式方法、时钟、缓存状态和工作负载聚合都会影响候选排序。

### 3\.2\.带守卫的单智能体优化循环

KernelArc的单智能体边界情况遵循一个测量门控循环，结构上类似于先前的自主内核优化系统（如AutoKernel（jaber2026autokernel (https://arxiv.org/html/2608.17071#bib.bib1)））：观察基准，形成瓶颈假设，编辑一个内核，运行正确性和时序检查，然后仅在测量到改进时才提升接受的最佳。此循环用于研究深度：它在单一局部历史下优化一个基准，并展示一条聚焦的轨迹能走多远。因此，其结论与工作负载、形状、调用模式、模型主干和预算紧密相关。跨越更多形状、算子和搜索方向的覆盖需求，促使了KernelArc共享多智能体设计的产生。

### 3\.3\.从顺序平台期到设计需求

当目标不再是单一的固定形状算子，而是多形状的SOL-ExecBench任务时，顺序搜索可能成为搜索过程的瓶颈。对于L1-030（Blackwell B200上带有残差相加的注意力输出投影），其提交在16个输入形状上评分；因此一个好的搜索必须结合每个形状的配置选择、融合、精度选择以及可靠的冷缓存测量。

#### 为何单一轨迹会进入平台期\.

不同的L1-030形状可能倾向于不同的库配置、精度模式和融合策略。顺序智能体一次只能调查一个活跃方向，并且可能在其预算的相当一部分中，在饱和后重新访问附近的配置。因此，其局部历史可能将搜索轨迹的上限误认为是硬件上限，特别是当其他实现族尚未被探索时。

这种失败产生了四个系统需求：多样化优化方向、在不转移完整历史的情况下转移经过验证的经验、保持评估和协调状态确定性，以及提供从局部平台期明确退出的路径。

### 3\.4\.KernelArc架构

KernelArc将需求映射为可配置机制：专精的智能体使搜索多样化；内存提炼出具有可调整保留期限的可转移结论；级联门和守卫负责验证、评估以及保留/回滚状态；而平台期重起草可以可选地使用从守卫归档中派生的只读跨智能体状态。图1（https://arxiv.org/html/2608.17071#acmlabel1）展示了评估的架构。第4\.3节（https://arxiv.org/html/2608.17071#S4.SS3）评估的是完整配置，而非单个机制的边际贡献。

图1\. KernelArc架构。策略技能为智能体组合提供种子；智能体通过共享内存回忆和记录经过验证的结论，然后将候选提交给确定性级联门和守卫。守卫在SOL-ExecBench上对候选进行基准测试，返回REJECT/KEEP/ACCEPT/REVERT操作和STOP反馈，归档保留的最佳变体，并通过内部活动排行榜暴露守卫最佳，以供只读的跨智能体检查。此活动排行榜不同于用于外部排名的公开NVIDIA SOL-ExecBench排行榜。策略技能为一组智能体提供种子。智能体通过共享内存交换成功经验和陷阱，将