把 vLLM 的 FlashAttention-3 前向 kernel 再跑快一截,26 组配置都要打平,正确性不能退——这是 GPU 算子优化里的硬骨头。论文做了一次对照:一边是按主流配方搭到极致的自主智能体,一边是只留最小 harness、靠人提问来驱动。前者停在 0.48 倍,后者做到 1.13 倍。

满配版:给智能体一整条流水线

满配的那一方叫 Arsenal,是论文自己搭的基线,沿 AVO、KDA 这条演化搜索路线做。每一步都塞给模型:按 NCU 瓶颈标签选出的 16 个优化杠杆之一、解析过的性能诊断、反汇编的指令混合差异、BM25 检索来的 PTX ISA 与 CUDA 指南文档、从 DeepGEMM/ThunderKittens/CUTLASS 挖出的参考提示(磁盘上 1.6 GB)、15 条种子知识笔记,外加 sanitizer 与禁止符号的反作弊检查。每步 prompt 中位数 22.6K 字符,37% 是预取材料。

结果是:133 个模型驱动步,API 花费 660 美元,62.4 小时里跑了 44.7 小时才到最好成绩。最好的几何平均 0.48 倍(摘要写作 0.50 倍),最差配置 0.24 倍,26 组配置里打平的有 0 组,目标未达成。冠军在第 35 步到达 0.48 倍后再没动过,剩下 98 步原地踏步。期间开过 37 个分支,36 个从同一步分叉,3 次人工回滚丢掉 117 步。

分水岭:82% 对 26%

论文落在两个数上。满配方 prompt 里有 37% 是预取参考材料,可它读到的工具结果文本中,82% 花在重读自己的 kernel 源码上;提问方什么都不预取,这个比例只有 26%。

工具用量上提问方反而更多:原始输出跑了 429 次 profiler、372 次 SASS/PTX 检查、377 次构建配 937 次基准跑,还派了 48 个子智能体。满配方则是每步一次解析好的 profiler、92 次文档检索、0 个子智能体。全程 16 个会话、11 532 次工具调用(8 791 次在主会话,2 741 次在子智能体),平均每条人工消息对应约 71 次工具调用。

论文的观察很直接:提问没有减少工具使用,而是把工具用在对的地方。

那 17 个问题长什么样

162 条人工消息按主行为分类:指令 58 条、提问 44 条、批准 30 条、纠正 19 条、目标 11 条;63 条是疑问句式。提问分八类:测量校验 13 条、边界差分 10 条、基准有效性 7 条、差距分诊 7 条、流程 6 条、质疑上限 5 条、策略选项 5 条、假设校验 3 条。

论文挑了四个案例,形态比数量更能说明问题:

  • 「什么配置在拖低几何平均?」agent 以为该调 hd128 路径,实际 hd96、hd192、hd256 全掉在 0.01 倍的标量回退路径上。删掉回退后,4 小时内几何平均从 0.44 倍拉到 0.96 倍。
  • 「这个上限是硬件的性质,还是我们设计的性质?kernel 之外的时间花在哪?」agent 认定 split-KV 解码「在 harness 规则内已到真实天花板」。host 侧 profile 抓出每次调用 2-5 μs 的张量填充,删掉后 split-KV 从 0.87 倍到 0.97 倍,几何平均从 1.02 倍到 1.13 倍。
  • 「我们的场景覆盖了算子真实被调用的形状和特性吗?」agent 一直以为 40 个场景全对就等于完成。搬来完整测试网格后,58 个里只对 52 个,sink 和左填充输出差了 495 倍和 103 倍。改完 58 个全对,新网格几何平均 1.10 倍。
  • 「你为什么觉得目标达成了?我们还远低于打平。」Stop hook 的判定此前已接受「目标达成」,agent 承认自己在第一个有效杠杆后就宣布胜利。fp8 hd256 解码从 0.20 和 0.34 变成 0.66 和 0.77。

这些问题都不含答案。人没写一行 kernel 代码,63 条疑问句里只有 17 条提到注意力概念,没有一条指定 CUDA 构造。两个 kernel 共 13.8K 行,全是 agent 写的。

交出来的东西

fa3_hopper.cu(3.6K 行)和 fa4_blackwell.cu(10.2K 行)以 Apache 2.0 发布,附带评测器、验证 Blackwell 原语的 PTX 探针、问题库和全部 trace。fa3_hopper.cu 在 182 个 Hopper 场景上全对,1.16 倍 FA3,对最快可用实现(含 cuDNN)1.10 倍,182 个里 135 个达到或超过打平;fa4_blackwell.cu 在 154 个可跑场景全对,1.46 倍 FA4,对最快可用实现 1.07 倍,124 个里 98 个打平以上,最差 0.78 倍;在 SOL-ExecBench 的生产形状上是 1.10 倍 FA4。

作者自己列的折扣也得照抄:剔除 fp8-KV 解码这一类 claim 形状后,Blackwell 全集几何平均变成 1.37、1.32、1.08 倍;edge 集是短 kernel,部分 1.5 到 2 倍来自上游 dispatch 开销;Hopper 侧在 fp8 dense prefill、softcapped prefill、head dim 96 和长滑窗上仍是 0.34 到 0.79 倍;Blackwell 侧 fp8 输入和 192 及以上的 head dim 还没做(182 个里 28 个);两个 kernel 都只覆盖前向。

这份对照该怎么读

论文把限定写得很清楚:每个 arm 只跑一次,前后端不同(满配方走 Claude Code 无头模式,提问方一半 Claude Code 一半 Cursor),受控组对着不受控组;KernelSeek 在此前两天已为 FA3 的 40 场景测试网格写过 kernel,所以它 0.009 倍的起点低估了领先;68.0 小时里有 26.7 小时是空闲间隔。

更要紧的是论文承认还没做的那个实验:问题库能不能替掉人。设计是四组对照(满配循环、裸 agent 加 harness、裸 agent 加从问题库取问的提问器、人提问),锁住模型和预算——写下来了,还没跑。72 条从公开资料挖来的问题也还没指挥过 agent。

也别把结论读成「工具流没用」。别的论文在别的任务上有产出:AVO 报告在 B200 注意力上连续演化 7 天,比 cuDNN 高出最多 3.5%、比 FA4 高出最多 10.5%;Cursor 与 NVIDIA 的多智能体系统在 3 周里优化 235 个 Blackwell kernel,报告 38% 几何平均提升。任务和评测器不同,不能直接横比,但说明这篇量到的是一个具体失灵,不是一条通用判决。

所以呢

如果你的智能体正卡在某个性能目标上,先别急着加检索源和知识库。先看它读的东西有多少是自己的产出——满配那个是 82%。再问一句:有没有哪个问题,能把它挪到一个它没检查过的层,或者逼它为已经宣布的上限辩护?论文的结论句挺直接:深度算子优化里,稀缺的不是信息,是方向。全部数据与评测器在 arXiv:2611.03517,文中提到的多智能体系统出自 Cursor 与 NVIDIA 的公开报告。