一篇 8 月 14 日发布的论文，给"LLM 生成 GPU kernel"这个增速极快、正确率数字一路走高的方向泼了一盆冷水：作者用一套 12 项"契约级"门控的验证器，审计了 2638 个已经被一个公开系统自身测试流程判定为"正确"的机器生成 kernel，结果 39.5%（1043 个）以任何容差都无法辩解的方式损坏，62.1% 至少违反一项契约。而在该领域通行的标准验收测试下，这些坏 kernel 里有 1487 个会照常通过；反过来，被验证器放行却被标准测试拒绝的只有 14 个。
论文题为《A Contract-Grade Verifier for LLM-Generated GPU Kernels》，只有两名作者（Rishi Shah、Rishav Shrestha），发布于 arXiv（2608.12700）。它真正攻击的不是某一家公司的生成系统，而是所有这类系统共同依赖的那条正确性信号：一条几乎不探测任何东西的测试。
一个几乎不探测任何东西的验收标准
LLM 生成 GPU kernel 的评测（KernelBench、TritonBench 等）几乎都用一个标准判断"对不对"：在一个固定 shape 下，抽少量随机输入跑一遍，如果输出与参考实现对得上——用 torch.allclose 以 atol=rtol=10^-2 的宽松容差比较——就判定为正确。这套做法便宜、可复现，但它只探测了 kernel 行为里极窄的一片。
论文列举了一个 kernel 能通过这套测试、却仍然错误的方式：它可以在参考实现应当返回 NaN 或无穷大的位置返回一个普通数字；可以每次运行给出不同结果；可以在 batch 或序列长度一换就崩溃；可以在 fp16 下累加归约，而参考实现保持着 fp32 的累计和。这些问题没有一个会被"少数随机输入 + 单一 shape + 接近即通过"抓住。
这个观察并非首次提出。论文自己引用的另一篇工作《The Correctness Illusion in LLM-Generated GPU Kernels》（arXiv 2606.20128）曾在一个 24 个 kernel 的手工语料上论证过同一点。新论文的增量在于：它把这个问题第一次搬到了数千个已被真实系统验收的 kernel 上，并给出了一个可运行、可复现的量化答案。
十二项门控：把"接近即通过"变成"逐项契约"
验证器的做法是把"正确"拆成十二个必须满足的契约门控，每项都对着一个慢速高精度参考实现打分。其中几项不需要任何容差，因此不存在"把阈值调松一点就能解释掉失败"的余地：
| 门控 | 检查的性质 |
|---|---|
| CMP-01 | 在大量随机输入和对抗性输入（零、10⁶、10⁻⁶、次正规数、长序列）上都正确 |
| CMP-03 | 换 batch / 长度 / 宽度后仍正确，而不只是被测的那一个 shape |
| ORD-02 | 五次重复运行逐字节一致，且不别名共享缓冲区 |
| EXC-01 | 无穷大和 NaN 落在与参考完全相同的位置和符号上 |
| PRC-02 | 输入 fp16 时，内部仍保持 fp32 的累计和 |
| RES-02 | 编译出的 kernel 满足真实硬件限制（寄存器、共享内存、Tensor Memory 预算） |
其余门控覆盖梯度正确性（CMP-02）、归约求和的舍入误差界（ORD-01）、对抗性求和顺序（ORD-03）、fp32/fp16/bf16 三精度（PRC-01）、次正规数的 flush-to-zero 行为（EXC-02）和输出设备归属（RES-01）。论文强调每个容差都是推导出来而非拍脑袋定的——例如 ORD-01 的界来自浮点误差随归约长度 ∝√N 的累积规律——且每个门控都用固定种子重抽输入，避免"碰运气"式的裁决。
1487 对 14：方向性暴露了系统性盲区
审计对象是 Dr. Kernel / KernelGYM 发布的公开语料（hkust-nlp/drkernel-coldstart-8k，MIT 许可），共 8920 条监督微调轨迹，每条以替换某个 PyTorch 模块的 Triton kernel 收尾。作者只取了其中与 SSM 相邻的算子类（matmul、attention、softmax、scan、norm、conv、reduction），在 B200 上逐个沙箱运行得到 3134 个 kernel，排除工具链与编译产物后，留下 2638 个已被来源系统自身判定正确（final_speedup > 0）的样本。
核心结论：这 2638 个"验收通过"的样本里，62.1% 至少违反一项契约，39.5% 以容差无关的方式损坏。后者就是论文所说的"地板"——这是最保守的数字，因为它不受任何阈值辩论影响。
最常见的失败模式值得单独强调。按失败率排序，第一是 EXC-01（34.2%，868 个）：kernel 把参考实现返回的 NaN 或无穷大静默替换成一个普通数字。这不是舍入误差，而是把一次可检测的训练错误，变成不可检测的数据污染。紧随其后的是 CMP-01 数值错误（23.6%）、CMP-03 换 shape 即坏（18.1%）、PRC-02 丢失 fp32 累加器（13.8%）。
最能说明问题的是一张 2×2 差分表：把同一批 kernel 跑一遍 KernelBench 的标准检查（allclose atol=rtol=10^-2、五次随机输入、固定 shape），它放行 93.7%（2472 个）。其中"标准测试通过、验证器拒绝"的格子有 1487 个（占已验收集合的 56.4%）；反向"标准测试拒绝、验证器通过"只有 14 个（0.5%）。两者几乎只朝一个方向分歧——如果只是"验证器更严格"，分歧应当是双向的。这种单方向性正是"验收信号存在系统性盲区"的签名，而不是"两个检查器对'多接近算接近'有不同标准"。
四路自证：结论不是"检查器更严格"
这类审计最大的软肋，是"检查器只是比别人严格，外来 kernel 天生就过不了"。论文用四条可各自证伪的证据回应：
- 7/7 阳性对照：作者自己手写的 6 个 Mamba-3 Triton kernel 和一个原生 GDN backward 跑同一套验证器，全部通过全部容差无关通道。这个对照还真的抓出了作者自己的一个 bug——C5 kernel 从输入张量推断通道维度却没有核对卷积权重通道数，一行补丁（commit
7f46226）修掉后才得到 7/7。 - 阈值校准扫描：对每个带容差的门控注入已知幅度的扰动，确认"诚实噪声地板"和"真实错误起始点"之间有干净间隔。唯一偏薄的界是刻意放宽的 ORD-03（仅 1.2 倍间隔），其覆盖由 CMP-01 和 ORD-01 兜底。
- 98.5% 与 KernelBench 自身代码一致：在 1030 对样本上运行 KernelBench 自己的正确性代码（pin 到 commit
48642c5c），与作者复刻的判定 98.5% 吻合；15 个分歧里 7 个是作者复刻更严、8 个是旧版 allclose 的种子边界翻转。 - 分层人工审计：手查 31 个"标准测试通过、验证器拒绝"的争议样本，16 个确为损坏（容差无关核心）、8 个是真实但容差相关（宽松 allclose 藏掉的）、7 个被重分类为出范围并从严格地板中剔除。
稳健性也单独做了：在 torch 2.11 栈上重审 300 个 kernel 得到 68.6% 的违规率；在 Sakana CUDA Engineer 的 213 个原生 CUDA kernel 上，剔除以精度主导的失效后仍有 20.2% 的环境稳健残差（主要是 shape 刚性和次正规数处理）。
CAKE 与 RealisticTritonBench：正向路线和第三方佐证
同一天发布的另一篇论文《CAKE: Compiler-Agent Co-Design for Frontier Kernel Evolution》（arXiv 2608.12629）提供了对照的另一半：如果验证规则被内建到生成循环里，而不是事后补一个宽松检查，会怎样。CAKE 让 Agent 直接编写一种类型化、硬件显式的调度表示（CAKE IR），暴露 warp 角色、访存、同步与流水线，同时支持验证、成本建模和局部诊断；反复出现的失败会被固化成新的验证规则。在 B200 上的 Flash-KMeans 干净起步中，CAKE 候选在 8000 万 token 预算下跑到调优 FlashML 基线的 1.144 倍，而直接写 CUDA/PTX 只有 0.928 倍；其 Kimi Delta Attention 相对官方 FlashKDA 有 2.05 倍几何平均加速并通过端到端服务验证。它从 Ampere 到 Blackwell 都可用——但这属于作者自报，尚无第三方复现。
独立于以上两方，第三篇论文《RealisticTritonBench》（arXiv 2608.12004，已被 ASE 2026 接收）从另一个方向印证了"真实任务更难"：它从主流 AI 框架的真实 pull request 中抽取 Triton 生成任务，把生成的 kernel 放回原框架做端到端测试，并明确指出此前的 benchmark 依赖手写评测脚本，存在可被模型利用、绕过正确性检查的漏洞。其结论是：领先 LLM 在真实 Triton 任务上仍然吃力。
从业者该改什么，以及这篇论文没说的事
对依赖 AI 生成 kernel 的团队，可操作的结论很具体：在"输出接近参考"之外，验收标准至少应加入几项容差无关的契约——非有限值必须原样传播（EXC-01）、输出必须逐次确定（ORD-02）、换 shape 必须仍正确（CMP-03）。论文估算，这一小撮契约就能堵住大部分缺口，成本并不高。
边界同样要写清楚。其一，39.5% 是 Dr. Kernel / KernelGYM 这一个语料上的比例，不能直接外推为"所有 LLM kernel 的失败率"——尽管跨栈重审（68.6%）和第二个原生 CUDA 语料（20.2% 残差）表明这不是工具链或 Triton 特有的现象。其二，审计语料是纯前向的，12 项门控里 CMP-02（梯度）和 RES-02（资源元数据）因此空转，实际靠 10 项、其中 7 项承担主要判断。其三，论文的第二贡献——一个原生 Blackwell tcgen05 训练反向 kernel，覆盖 gated-linear-recurrence（GDN）家族并首次原生实现反向状态扫描阶段——本身是更慢的：比 fla Triton 库约慢 8 倍（L=512 时）到 78 倍（L=2048 时）。作者没有软化这个结构性劣势，而是把它当作"验证器公平"的又一证据：同一把尺子既能揭穿外来 kernel，也能精确说出自己 kernel 慢在哪里。
这条研究的核心判断可以压缩成一句：kernel 生成领域报告的所有进度，都建立在"正确性"这个前提之上，而那个前提目前由一个太弱的测试扛着。要修复的不是某一家系统的实现，而是所有评测共享的那条验收信号。