一篇 8 月 14 日发布的论文，重新检查了已通过测试的 AI 生成 GPU kernel（在 GPU 上执行特定计算的程序）：作者用一套 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 个样本，被新验证器拒绝
审计对象是 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% 以容差无关的方式损坏。后者是论文给出的保守下限：即使不考虑需要设置容差的检查，也有这一比例的样本存在错误。
最常见的失败模式值得单独强调。各检查的失败率以适用该项检查的样本为分母，不能直接按全部 2638 个样本计算。按这一口径排序，第一是 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: 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 生成的进展，需要先确认生成结果是否正确。仅报告现有测试的通过率还不够，测试能覆盖哪些错误，同样需要说明。
