先说结论：这次用 Agent 自动寻优 GPU kernel，从零搭建到冲上 NVIDIA kernel 榜单第 15 名，实际耗时不到 24 小时。整个过程没有手写一行调优代码，所有的 tile size、unroll 因子、共享内存策略、访存 pattern，全部由 Agent 在闭环里自己试出来的。这篇文章把完整的思路、工具链、踩坑记录和可以直接复用的代码框架写出来，给同样想做 kernel 自动调优或者 Agent 实操项目的朋友做个参考。

这个项目的核心价值在于：它把“调 kernel”这种极其依赖经验和耐心的活，拆成了一个可以自动化迭代的工程问题。你不需要在每一个算子上都投入人力去手搓汇编级别的优化，也不需要维护一套庞大的调参模板。Agent 负责生成候选、编译、跑分、淘汰、再生成，你只需要定义好搜索空间和奖励函数，剩下的交给它自己转。

如果你现在正在做算子库开发、性能优化、CUDA kernel 编写，或者对 Agent 在工程领域的落地方式感兴趣，这篇文章里会有不少可以直接抄作业的东西。下面我按实际推进的顺序，把这个项目拆开讲。

1. 项目目标：让 Agent 接管 kernel 调优的完整闭环

1.1 榜单规则与目标拆解

先说清楚“NVIDIA kernel 榜单”是什么。它本质上是一个公开的 kernel 性能评测平台，社区里会定期放出一批算子基准，参与者提交自己实现的 kernel，后台在统一型号的 GPU 上跑 benchmark，按吞吐量或者延迟排名。排名靠前的 kernel 会被更多人看到，所以这算是一个相对硬核的技术角力场。

榜单评测的算子没有固定种类，涉及矩阵乘、卷积、向量归约、FlashAttention 变体、Softmax、Gemm 等常见算子。每个算子有独立的排名，总榜则按一定规则把多个算子的性能加权合成。想要冲榜，光靠精通某一个算子是不够的，必须在多个算子上有平均之上的表现，这正好是“多点并行调优”的场景——很适合 Agent 来搞。

我把目标拆成三个层次：第一层是跑通流程，让 Agent 能生成、编译、提交、拿分数，这个半天内完成；第二层是让 Agent 学会根据分数反馈调整策略，而不是瞎猜，这个需要设计合理的搜索空间和反馈机制；第三层是把整个过程做成无人值守的，睡觉的时候它自己在跑，醒来直接看结果。实际执行下来，24 小时冲上第 15 名靠的也是这三层都打扎实了。

1.2 为什么不用人工调优而选 Agent

我见过太多 kernel 调优的日常：先看 profiling 报告，然后猜瓶颈是访存还是计算，改几个 macro，重新编译，跑一轮 benchmark，看曲线，再改。一个经验丰富的人调一个中等复杂度的算子，通常也需要三到五天才能逼近理论峰值，而且换个 GPU 型号就得重来。

Agent 的逻辑不一样。它不需要“理解”瓶颈，它只需要在搜索空间里高效地尝试。每次生成的 kernel 就是一个候选解，跑分就是适应度，分数低就淘汰，分数高就保留，再基于高分候选做交叉变异。这本质上是一个自动化的进化式搜索，只不过用 Agent 的代码生成能力替代了手动写变体。

这里有个很关键的点：Agent 生成 kernel 的能力上限，取决于你对搜索空间的约束。如果你给它完全自由的发挥空间，它可能会生成一堆编译不过的垃圾代码；如果你把 tile size、线程块数、寄存器分配、访存方式这些维度框成枚举值，让它在组合空间里搜索，效率会高非常多。这跟让人类新手直接写 kernel 是一样的逻辑，但在自动化的加持下，探索速度是人工的几十倍。

2. 工具链选型：Agent 框架、编译环境和评测闭环的搭建

2.1 硬件与基础环境

这次用的是一张 Ampere 架构的 GPU，显存 24GB 级别，驱动和 CUDA 版本分别为 535 和 12.2，编译工具链用的是 nvcc + gcc 的经典组合。为什么强调版本？因为 kernel 的编译行为在不同 CUDA 版本之间差异很大，比如 register allocation 的启发式规则、inline PTX 的支持程度，都会直接影响最终 binary 的性能。而 Agent 搜索出来的一些“偏门”写法，在较新的 CUDA 版本里才编译得动，老版本会在编译阶段就劝退一大半生成结果，导致搜索效率暴跌。

建议把环境写进 Dockerfile 里固化下来，别用宿主机环境直接跑。我自己第一次跑 Agent 的时候就是栽在环境漂移上：Agent 在容器里生成的一个 kernel 在宿主机上编译过了，但一跑就非法内存访问，排查了半天发现是宿主机驱动和容器内 CUDA 版本不匹配。把环境固化成镜像后，这种问题就再也没出现过。

2.2 Agent 框架选择

Agent 部分没有用特别重型的框架，核心就三层：一个 LLM 驱动的生成器、一个编译执行器、一个搜索控制器。生成器的职责是根据当前搜索状态和历史分数分布，生成下一批 kernel 变体；编译执行器负责把生成代码落盘、编译、运行 benchmark、收集指标；搜索控制器决定保留哪些候选、淘汰哪些、往哪个方向继续探索。

LLM 选用的是支持代码生成能力较强的开源模型，跑在本地。为什么不用更大规模的模型？核心原因是迭代延迟。在自动寻优场景里，每个 kernel 的编译和跑分本身需要几十秒到几分钟，而模型推理时间只占整个循环的小头。如果选一个推理速度很慢的大模型，生成一批候选就要等半天，24 小时根本跑不完几轮迭代。实测下来，一个中等规模模型的生成速度已经能喂饱整个搜索循环。

这里还有个小细节：生成 prompt 里必须带上最近十轮的分数变化趋势和保留候选的具体参数，这样 Agent 才能把握当前搜索方向。如果只是把随机噪声塞给它，生成的候选会和上一轮完全没有承接关系，搜索就退化成瞎蒙了。这个设计极大影响了收敛效率。

2.3 提交闭环与评测接口

榜单平台提供了 API 方式的提交入口，也支持本地评测后手动上传结果。为了省时间，我一开始就实现了本地的完整 benchmark 闭环，所有打分都在本地 GPU 上先跑一遍，Agent 只会在本地分数达到一定阈值之后才提交到平台。这样做有两个明显好处：一是避免把大量性能差的 kernel 提交上去浪费评测次数，二是本地评测可以带更细粒度的 profiling 信息，帮助 Agent 理解瓶颈。

评测接口这边有个坑要提一下：平台后台跑分时的 GPU 频率策略和本地不完全一样，尤其是长循环的 benchmark 容易触发降频，导致同一个 kernel 两次跑分能差出 10% 上下。我的处理方式是每次跑分前先跑一个固定的预热 kernel 来拉起频率，然后跑三遍目标 kernel 取中位数。这样虽然单次评测时间长了，但数据的稳定性高很多，Agent 的决策不会因为噪声刷掉了真正好的候选。

3. 自动寻优的核心策略：搜索空间的工程化

3.1 如何把“调 kernel”变成可搜索的组合空间

很多人第一次做自动调优时会犯一个错误：直接让 Agent 全文生成一个算子 kernel，然后期望它自己猜出所有优化参数。其实在现实里，一个算子的 kernel 结构是有基本范式的，比如矩阵乘有经典的 shared memory tile 分块 + 寄存器累加结构，attention 类算子有 online softmax 的固定流程。你不需要让 Agent 重新发明范式，你要做的是把范式里的可变参数暴露出来。

我把搜索空间分成了四个维度：tile 结构参数、线程组织参数、访存策略参数、编译期优化参数。tile 结构参数包括 BM、BN、BK 这些分块大小；线程组织参数包括线程束配置和每个线程负责的元素数；访存策略参数包括是否启用 vectorized load、是否使用 cp.async、L2 persistence 的开关；编译期优化参数包括 unroll 因子和#pragma unroll的展开策略。

这四个维度的组合空间大约有数万种，听起来很多，但 Agent 一轮可以并行生成几十个候选，并且可以基于已有的分数分布做贝叶斯式的偏好选择，所以实际搜索效率非常高。关键是不要让 Agent 在语法层面自由发挥，要把候选代码模板固化好，参数位置留成变量，Agent 只是填值并偶尔修改局部计算逻辑。

3.2 奖励函数与反馈链路的设计

自动寻优的另一个核心是奖励函数。如果只用“一次跑分”作为唯一的反馈信号，Agent 很难知道它生成的那个 kernel 到底是哪个决策起了作用。我的做法是拆细反馈：除了最终吞吐量，还把 kernel 的寄存器用量、shared memory 占用、指令数、bank conflict 次数、L2 命中率这些 profiling 指标打包传给 Agent。

这样 Agent 能看到的不只是“这个 kernel 慢”，还能看到“它慢是因为寄存器溢出导致 local memory 访问暴增”这种具体原因。下一轮生成的候选就会针对性地减少寄存器占用。实测下来，加上细化反馈后，搜索收敛的速度差不多快了一倍，而且跑出来的最终 kernel 的质量也更高，不会只找到局部最优就卡住。

反馈链路上还有个容易被忽略的环节：编译警告信息。很多 kernel 不是编译失败，而是带着警告通过编译，比如“variable may be used before initialization”这种，生成代码的人通常毫不在意，但这类警告有时候正是性能杀手的前兆。我让 Agent 把编译 warning 也作为文本反馈读入，它会自动学会规避那些产生不确定行为的写法。

3.3 核心算子优化案例：一次 Matrix Multiply 的完整调优

拿榜单里的矩阵乘算子举例。人工调优的常规路径是从 baseline 出发，先调 BM/BK/BN，再用 vectorized load 改访存，不行就换双缓冲流水线。Agent 的路径不太一样，它的第一轮候选是从我提供的基础模板出发，把参数打乱随机组合，生成 64 个变体并行编译跑分。这一轮跑完基本就能锁定一个不错的参数范围。

第二轮开始，Agent 会拿 profile 数据做决策：如果发现所有高分候选的 shared memory 占用都很高，它就倾向生成 shared memory 更大的变体；如果发现 bank conflict 出现频繁，它会在下一轮生成时主动修改数组 padding 的参数。我看到它生成的某个候选里，自动给矩阵 A 的 shared memory 数组加了一个 float 的 padding，bank conflict 从 15% 降到了 0.3%，吞吐直接拉了近 30%。这种优化人工也能想到，但通常是在 profiling 之后，而 Agent 在第三轮就自己试出来了。

最终这个矩阵乘 kernel 的得分从 baseline 的 3.8 TFLOPS 干到了 11.2 TFLOPS，在榜单该算子下排到了前二十。整个过程从第一轮生成到拿到这个结果，大概花了五个小时，其中有三个小时是在等待编译和跑分。这个产出速度对于传统人工调优来说是不可想象的。

4. 实操过程：24 小时的时间线、运行策略与关键节点

4.1 前 6 小时：搭建基线、定义搜索空间、跑通第一轮迭代

项目启动的第一件事并不是写复杂框架，而是先建立一套“结果衡量标准”。我在本地写了一个统一的基准脚本，把平台评测用的那几个算子都覆盖了，确定跑分位次、预热策略、重复次数这些能直接影响结果的标准流程。这个脚本是后面所有 Agent 决策的依据，务必先跑稳。

接下来是准备种子模板。每个算子我准备了一个能正确运行的 Naive 版本，和一个经过粗略手工优化的版本，后者作为搜索的起点。我没有花时间去做完整的性能分析，只保证算子在语义上是正确的，然后把可变参数暴露出来。这一步大概用了两小时。

跑通第一轮迭代是最关键的时间节点。因为框架里任何一环出错——Agent 生成代码的格式不对、编译脚本没捕获错误、跑分结果解析错位——都会让整个循环锁死。我先用 8 个候选跑了个小规模验证，确认输出结果和平台评测逻辑不会冲突，然后才把候选规模放大到 64 个。第一轮完整迭代跑完后，我把整个日志检查了一遍，确认分数分布、保留逻辑、淘汰逻辑都符合预期。

4.2 中段 12 小时：多卡并行、候选池管理、策略转向

进入稳定迭代后，瓶颈转移到了计算资源上。单个 GPU 同时在跑的 benchmark 会互相干扰，导致分数噪声暴涨。我的做法是用两张卡，一张跑搜索闭环，另一张留给“验证池”——凡是 Agent 觉得可以保留的候选，都会先在两卡上分别跑分交叉验证，确认没有卡间差异才放入保留池。

候选池的管理也是一个值得讲的地方。Agent 每轮生成的 64 个候选中，可能只有两三个能进入保留池，而保留池上限是 200 个。淘汰策略我不是按分数一刀切，而是同时考虑分数多样性和参数多样性。因为自动寻优最容易掉的坑就是过早收敛到局部最优，如果候选池里所有 kernel 的参数都很相似，那下一轮的搜索结果不会有本质突破。

所以在保留了高分候选的同时，我还让 Agent 每隔几轮强行生成一批“边缘探索”候选，也就是参数组合离现有池子最远的那些。这一策略在后半程挽救了整个项目——当主流候选卡的某个算子在 9 TFLOPS 附近时，一个之前被忽略的“边缘”参数组合突然在另一个算子上表现出了极强性能。如果没有多样性保护，我们大概率会错过这个突破点。

4.3 后 6 小时：合并冲刺、提交策略与榜单提升

最后六小时的主要工作不是继续撒网搜索，而是把已经找到的优质候选做细粒度合并。Agent 会尝试把两个不同算子上的最优参数拼到一起，看是否有冲突。很多参数在单个算子上表现为最优，但合并后因为共享内存总量超限而无法编译，这类冲突会自动退回次优组合。

提交策略上，我不会在 Agent 每找到一个高分候选就立刻提交到平台，而是攒够了 10 个左右的高分结果，统一走一次评测 API。这样即使其中有被平台评测方判定为无效的（比如存在未定义行为），也不至于把宝贵的评测次数浪费在过程上。榜单的刷新有延迟，最终在我最后一次提交完成后的半小时，排名落定在了第 15 名。

回看整个 24 小时的节奏，真正有效的“脑力劳动”其实只集中在前三小时和最后一小时的合并冲刺上，中间十几个小时完全是 Agent 在自动运转，我只需要偶尔盯一眼日志，处理一些环境层面的偶发问题。这也是我比较想传递的一个信号：Agent 不是取代你做所有事，而是把重复、高试错成本的那部分大量替代掉，让你把精力集中在策略设计和结果判断上。

5. 踩坑实录：Agent 自动寻优的常见问题与排查指南

5.1 生成代码的编译失败率过高

Agent 自动生成的 kernel 代码在编译阶段的失败率，第一版框架时期高得吓人，几乎有三分之一是编译不过的。后来排查发现，大部分失败都集中在两个原因上：一个是模板里的占位符被 Agent 改动时破坏了结构，比如把宏定义的括号搞乱了；另一个是 Agent 尝试访问了不存在的数组下标，因为模板里的数组维度是硬编码的。

解决办法是给 Agent 的 prompt 里加了明确的“代码结构约束”，不允许改动特定代码段以外的地方，并且每次生成完后用一个轻量 AST 解析器做结构校验，只要模板骨架变了就整批丢弃重新生成。这一刀切下去，编译失败率降到了 10% 以内。这个教训很重要：Agent 生成代码的自由度必须和脚手架约束成正比，自由多了就会生成不可控的东西，超出一个工程所能接受的容忍度。

5.2 偶发的非法内存访问与不确定行为

这类问题在 kernel 开发里最让人头疼，因为不是每次运行都会复现，偶尔会被评测平台判定为错误提交。Agent 生成代码的过程中，我在搜索空间里开放了某些过于激进的访存参数，比如跨越共享内存边界的 vectorized load 种类，结果就是某些特定形状下会有非法访问的潜在风险。

解决思路是引入运行时边界检查模式：专门有一轮搜索用于把当前候选在我们定义的“边界测试输入”上跑一遍，覆盖最大和最小的矩阵尺寸、对齐和不对齐的地址。这种模式跑得慢，但只在候选进入保留池前执行一次，成本可控。经历过一次评测平台打回无效提交后，我坚决把这个环节加进来了，宁可牺牲一点本地迭代速度，也不留这类隐蔽炸弹。

5.3 benchmark 分数波动导致的误判

自动寻优过程中最常见的隐蔽问题就是 benchmark 噪声。同一个 kernel 在同一个 GPU 上，不同时间跑出来的分能差 8% 甚至更多，原因是 GPU 的频率调度、显存温度、并发进程干扰。如果 Agent 根据带噪声的分数做决策，就会出现“上一轮被淘汰的候选，下一轮重新生成却能刷新纪录”这种荒谬情况，严重拖慢收敛速度。

解决思路分三层：第一层是每次跑分前预热 GPU 并固定频率，让 boost clock 稳定；第二层是每个候选重复跑三遍取中位数，排除单次运行的随机抖动；第三层是记录候选的“最佳分”和“中位分”两个指标，Agent 决策时优先用中位分，只有中位分优于当前保留池的候选才进入下一轮。实测下来，引入这三层后分数噪声的影响基本被压制住了。

5.4 问题速查表

这个表里的每一项都是真实发生过的，前两个尤其浪费时间，一旦出现往往意味着大批候选被废掉。如果你打算跑自己的自动寻优项目，建议把这个表贴在你的工作日志旁边。

5.5 搜索参数的日常调优心得

最后分享一个关于“Agent 自动寻优超参数”的经验：候选数量不是越大越好。我试过一轮 128 个候选，比 64 个候选多了不少并行度，但实际收敛速度并没有质变，原因是每轮迭代的时间拉长后，单位时间内的有效信息量几乎没变。64 个候选在一张 24GB 显卡上已经是性价比很好的平衡点。

另外就是迭代轮次之间的停顿时间。很多 Agent 框架默认允许连续生成下一轮候选，但实际执行时我故意在每轮之间增加一个短暂的冷却期，让 GPU 把上一轮 benchmark 的残留负载清掉，也避免显存碎片化。这个小改动提升了不少稳定性。

6. 这个项目后续还能怎么扩展

冲到第 15 名不是终点，这个方案还有不少可扩展的空间。比如现在搜索空间还是针对单个算子分别搜索的，下一步可以让 Agent 同时感知多个算子的评测结果，去搜索“跨算子的公共优化模式”。比如某些访存惯用手法在一个算子上有效，往往在另一个结构相似的算子上也有效，Agent 如果能自动迁移这种经验，后续新算子的调优就不再需要从零开始。

另一个方向是把 Agent 的搜索过程和 profiling 数据做更深的融合。目前反馈链路里已经有了寄存器用量和 bank conflict 信息，但还停留在文本层面。如果能把 profiling 的原始事件计数直接向量化后喂给搜索控制器，搜索结果的分辨率会细很多。这个改动工程上不难，主要是数据管线要梳理清楚。

对于想复制这个项目的人，我的建议是先别急着追求榜单名次，而是先跑通一个算子的闭环。选矩阵乘这种结构清晰、资料多的算子，把模板、搜索空间、跑分脚本都做好，确认 Agent 确实能在你的设定下找到比 baseline 更优的 kernel，再逐步放开到更多算子和更大榜单。这个思路适用于任何定位于“自动化工程优化”的项目，本质上都是在定义问题空间、构建反馈闭环、让智能体在其中高速迭代。

最后聊一点个人体会。做这种项目最大的收获不是那个名次，而是理解了“调优”这个活里到底哪些部分是经验直觉、哪些部分是纯粹的暴力搜索。有了 Agent 之后，暴力搜索的部分被大量接管，剩下的经验直觉反而变得更值钱了——比如搜索空间怎么定义、奖励函数怎么设置、什么时候该转向探索、什么时候该收敛打磨。这些决策质量直接决定了 Agent 最终能飞多高。如果你也在做类似的事，希望这篇拆解能帮你少走一些弯路，把更多精力放在真正能拉开差距的地方。