返回文档库
harness_evolve/notes/ref35_kernelbench.md

KernelBench: Can LLMs Write Efficient GPU Kernels?

一句话总结：斯坦福/普林斯顿造的一个"让 LLM 给 PyTorch 算子写更快的等价 GPU kernel（CUDA/Triton）"的评测基准——给一段 PyTorch 参考实现（Model 类），让模型输出一个自定义 kernel 版本（ModelNew），自动验证正确性 + 加速比；核心指标 fast_p = "既正确、又比 PyTorch 快 p 倍"的任务占比；结论很清醒：前沿推理模型（o1、R1）开箱即用也只在 <20% 的任务上追平 PyTorch，且难度随 p 上调而陡增——写高性能 kernel 对当前 LLM 仍是硬骨头。

TL;DR 速览

tags: #benchmark #GPU-kernel #CUDA #代码生成 #性能优化 #operator-fusion #test-time-refinement #fast_p #AI-for-systems

related: [[ref20_alphaevolve]]（进化式搜索优化 kernel/算法，KernelBench 是其天然试炼场）· [[ref27_sia-self-improving-ai]]（把 CUDA kernel 优化当自改进任务）· [[ref25_learning-to-discover-at-test-time]]（测试时发现，GPU kernel 属其发现目标之一）· [[ref09_meta-harness]]（harness/scaffold 对 agent 分数的巨大影响，正对应本文迭代精修反馈的增益）· [[ref31_re-bench]]（ML R&D agent 基准，含 GEMM/kernel 类任务，与本文互为参照）· [[ref30_paperbench]]（同期"AI 工程能力"评测，判分哲学互补）

摘要

高效的 GPU kernel 对构建高性能 ML 架构至关重要，但编写它们耗时且需要大量专业知识；因此我们探索用语言模型自动化 kernel 生成。我们提出 KernelBench，一个开源框架，在精心挑选的 250 个 PyTorch ML 工作负载上评测 LM 编写快速且正确的 kernel 的能力。KernelBench 代表一个真实的工程环境，在该基准上取得进展直接转化为更快的实用 kernel。我们提出一个新评测指标 fast_p，衡量"既功能正确、又相对 baseline 提供大于可调阈值 p 的加速"的生成 kernel 占比。

跨多种 SOTA 模型和测试时方法的实验表明：前沿推理模型开箱即用表现最好，但整体仍不足——在不到 20% 的情形下才能追平 PyTorch baseline。虽然我们证明利用执行与 profiling 反馈做迭代精修能改善结果，但 KernelBench 仍是一个有挑战性的基准，其难度随加速阈值 p 的提高而增加。

1 动机：为什么 GPU kernel 优化难，又为什么要自动化

论文的立意锚在一个真实的工程痛点上：AI 靠高效 GPU kernel 才能拿到高性能、低成本、低能耗，但写 kernel 一直很难。有两条相互叠加的趋势让这个痛点越来越尖锐：

  1. 架构爆炸：ML 架构出现了"寒武纪大爆发"（Transformer、SSM、RWKV…），但它们的可用实现常常远未跑满硬件潜力。
  2. 硬件爆炸：AI 加速器越来越多（NVIDIA 各代、TPU、Cerebras、Graphcore、Groq…），每款规格和指令集都不同，把算法从一个平台移植到另一个平台是巨大的痛点。

论文用一个极具冲击力的例子把这件事说透：FlashAttention kernel——Transformer 2017 年提出，直到 2022 年（5 年后）才有第一版高性能 attention kernel；而 NVIDIA Hopper GPU 发布后，又花了两年才把这个算法迁移到新硬件（FlashAttention-3）。"从一个好算法到一个好 kernel"之间隔着数年的专家人力——这正是作者想问的：语言模型能不能帮忙写正确又优化的 kernel？

[!TIP] 什么是 GPU kernel？为什么写它难？ GPU kernel 是一段在 GPU 上大规模并行执行的底层函数——它显式地把计算铺开到成千上万个线程上，并手动管理 GPU 的内存层级（寄存器→共享内存→全局 HBM，速度依次变慢、容量依次变大）。写好一个 kernel 要同时照顾： - 并行分解：怎么把一个大张量运算切成 block/thread 的网格； - 内存搬运：怎么把数据从慢的全局内存挪进快的共享内存/寄存器、并让访问"合并（coalesced）"； - 硬件特性：用不用 tensor core（专门做矩阵乘的硬件单元）、warp 级指令（如 wmma）、异步拷贝等； - 数值正确：并行归约/同步点一旦写错，结果就悄悄错了。

这就是为什么 PyTorch 底层调用的是 NVIDIA 闭源的、被专家手工调到极致的 kernel（cuBLAS、cuDNN）——它们是很强的 baseline。让 LLM 生成开源 kernel 去追平甚至超过它们，天然就难。

[!TIP] CUDA / Triton / ThunderKittens / CUTLASS：写 kernel 的几种抽象层次 AI 工程师写 kernel 时有一整套工具，从底到高： - PTX：NVIDIA GPU 的"汇编"，最底层、最难写（DeepSeek 用过）。 - CUDA-C：C++ 扩展，直接写 __global__ kernel 函数——本文让 LLM 主要生成的就是内联 CUDA。 - CUTLASS：NVIDIA 的 C++ 模板库，为线性代数（GEMM）提供高性能构件。 - Triton：一个 Python 式的中间语言 + 编译器，让你写"分块（tiled）"的神经网络计算，更容易用上 tensor core；FlashLinearAttention 等库基于它。 - ThunderKittens：斯坦福自家的库，用"简单、快、可爱"的 tile 抽象帮人写快 kernel。

本文的 baseline 实验让 LLM 生成原始 CUDA；但作者在未来工作里明确建议：换成 Triton/CUTLASS 这类更高层抽象，可能让 LLM 更容易用上 tensor core，是降低生成难度的一条路。

为设计这个 benchmark，作者列了环境应满足的三条： - 自动化 AI 工程师工作流：模型有完全自由决定优化哪些算子、怎么优化； - 支持多样的算法/编程语言/硬件平台； - 易于同时程序化评测性能与功能正确性，并能捕获生成 kernel 的 profiling/执行信息。

KernelBench 就是对这三条的回应。相比现有 LLM 代码生成任务（HumanEval 式只考"对不对"），写 kernel 需要海量且多样的信息（编译器反馈、profiling 指标、硬件规格与指令集、硬件效率技巧），这也是本文区别于普通 code-gen benchmark 的根本。

2 相关基准与工作

本文的定位可以一句话概括：现有 kernel 方案要么"靠专家手写"（cuDNN/CUTLASS/Triton），要么"编译器只给一小片自动优化"（torch.compile）；现有 HPC 代码生成评测要么"翻译任意 C++→CUDA"、要么"生成 GEMM 这种著名 kernel"——KernelBench 独占"从真实现代 ML 工作负载里策展 250 个多样 kernel、且很多没有现成人写实现"这个格子，因此解题直接对真实 DL 有益。

[!TIP] kernel 库与编译器（KernelBench 想超越的现状） 作者从 automation（自动化）/ breadth（覆盖广度）/ performance（性能） 三个维度评价现有 kernel 编程方案： - cuDNN / CUTLASS / Apple MLX：硬件专用、性能极高，但要大量专家人力（automation 低）。 - ThunderKittens / Triton：帮研究者写一批又快又对的 kernel（breadth 好），但仍需人来编程。 - torch.compile / FlexAttention：编译器自动优化，但只覆盖很窄的一片优化（narrow slice，如规则化的 fusion）。

KernelBench 问的是一个更雄心的问题：LLM 能不能对一大批多样 AI 工作负载自动生成高性能 kernel——同时要 automation、breadth、performance 三者。

[!TIP] LLM 做性能优化代码生成（本文所属的研究线） - 功能正确类：过去一年很多工作让 LLM 做算法编程（AlphaCode）、解 GitHub issue（SWE-agent / SWE-bench）、领域专用编程（DS-1000）——但这些只关心代码对不对、能不能跑。 - 算法效率类：后续工作（ECCO、Performance-Aligned LLMs）开始探索 LLM 产出渐进/算法效率更好的解。 - KernelBench 的差异：它关心 wall-clock 效率——生成的是高性能计算（HPC）代码，要理解底层硬件特性、设备指令集、并行处理器的性能特征。这比"渐进复杂度更优"更贴近工程现实（一个 O(n) 但访存糟糕的 kernel 可能比 O(n log n) 但访存好的还慢）。

[!TIP] 既有 HPC 代码生成评测（对照） - C++→CUDA 翻译类（CodeRosetta、BabelTower）：让 LLM 把任意 C++ 代码翻译成 CUDA。 - 著名底层 kernel 类（比较 Llama-2/GPT-3 生成 GEMM；RE-Bench 的 kernel 任务）：只生成 GEMM 这类广为人知的低层 kernel。 - KernelBench 的差异：策展 250 个来自真实现代 DL 工作负载的多样 kernel，其中很多没有现成的人写实现——换句话说，解 KernelBench 的题立刻对真实 DL 工作负载有益，而不是做一道有标准答案的练习。见 [[ref31_re-bench]]（RE-Bench 也含 kernel 类任务，但任务自包含、有 scoring function，两者互补）。

3 基准设计

这是 benchmark 论文的核心。KernelBench 的设计可拆成四块：任务格式 → 任务选择（三级分类）→ 评分（正确性 + 性能 + fast_p 指标）→ 硬件评测协议。先看总览图。

Figure 1: KernelBench 任务→生成→评测的三段式流程

Figure 1 逐元素解读（全文的骨架图，读懂它就读懂了 KernelBench 的任务闭环）： - 左侧输入（两块）： - 上「Task Instruction」：给模型的指令——"给定下面的模型架构，替换 PyTorch 算子以获得加速……通过生成内联嵌入的自定义 CUDA……" - 下「Reference Torch Module」：一个 class Model(nn.Module)，其 forward 里调用一串 torch 算子。这就是任务的"题面"。 - 中间「Language Model」：模型读入指令 + 参考模块，生成…… - 中下「Generated Torch + Compile Inline CUDA」：一个 class ModelNew(nn.Module)，其 forward 调用 torch 算子 + 自定义 CUDA kernel（图中绿色「Custom Kernel CUDA」块 + 「Driver Code」驱动代码）。注意模型不仅要写 kernel 本身，还要写把 kernel 集成进 PyTorch 的周边胶水代码（通常用 load_inline 内联编译）。 - 右侧自动评测（两块）： - 上「Evaluate Correctness」：用随机输入分别跑 Model.forward 和 ModelNew.forward，检查输出是否匹配（✅/❌），重复 n_correct 次。 - 下「Measure Performance」：把 ModelNew.forward（custom cuda）与 Model.forward（eager mode / torch.compile）用秒表 wall-clock 对比，重复 n_trial 次。 - 闭环含义：一次评测 = 模型读题写 kernel → 正确性验证 → 性能计时，完全程序化、无需人工判分。

3.1 任务格式（Task Format）

[!NOTE] 一个具体任务长什么样（附录 A，Level 1 matmul with large K） 参考实现极简：Model.forward(A, B) 就 return torch.matmul(A, B)，配 M=256, N=256, K=131072（一个 K 维超大的矩阵乘），get_inputs() 返回两个 torch.randn 张量。模型要输出的 ModelNew 则用 torch.utils.cpp_extension.load_inline 内联一段 matmul_kernel 的 CUDA-C 源码 + 一个 matmul_cuda wrapper，在 forward 里调用它。注意：这个示例输出是最朴素的三重循环 matmul（没用 tensor core、没 tiling），它能跑对但通常打不过 cuBLAS——这正是 Level 1 的难点写照。

3.2 任务选择：三级分类（250 题）

250 道任务按"含几个 PyTorch 原语操作（primitive operations / library functions）"分三级：

Level 题数 内容 代表任务 baseline 为何难打
Level 1 100 单个原语算子 卷积、矩阵-向量/矩阵-矩阵乘、losses、activations、layer norm PyTorch 底层调高度优化的闭源 kernel（cuBLAS 等），LLM 难超越；但若成功，开源 kernel 是有影响力的替代品
Level 2 100 算子序列（3–6 个原语） 如 conv + ReLU + bias 的组合，可融合成单个 kernel 编译器工具（torch.compile）很擅长 fusion，LLM 要超它不易；但 LLM 或许能提出比编译器规则更复杂的算法
Level 3 50 完整 ML 架构 AlexNet、MiniGPT 等（来自 pytorch / huggingface 等热门仓库） 现代模型规模大，峰值性能 kernel 需要编译器范围之外的算法级改写（如当年 Transformer→FlashAttention 花了 5 年）

[!TIP] 什么是 operator fusion（算子融合）？为什么它是 Level 2 的核心考点？ GPU 有"少量快内存 + 大量慢内存"。假设你要连续做 matmul → ReLU → bias，朴素做法是：算完 matmul 把结果写回慢的全局内存，再读回来做 ReLU 写回去，再读回来加 bias……每一步都在慢内存上来回搬数据。 Fusion 就是把这几个算子合并进一个 kernel：数据一次性加载进快内存后，在快内存里连做 matmul+ReLU+bias，只写回一次结果——大幅减少慢内存 I/O。 Level 2 专门考这个：它给的就是"能被融合成单个 kernel"的算子序列。挑战在于 torch.compile 这类编译器已经很擅长规则化 fusion，所以 LLM 要么融得更彻底、要么想出编译器规则外的算法（如 online softmax）才能赢。论文观察：LLM 在 Level 2 上仍"融合不足"，把很多本该融合的机会留在了桌上。

[!NOTE] 作者反复强调一个设计原则：每道题都含"一组有意义的 AI 原语操作或架构"，使得 LLM 在该题上的成功能直接带来真实世界影响——这不是人造练习题，而是策展自真实 DL 工作负载。

3.3 评分：正确性 + 性能，与核心指标 fast_p

KernelBench 是只评测（evaluation-only） 基准——不提供标准答案 kernel（因为设想用户会在各种硬件/输入类型/工作负载上评测，包括全新平台）。但它天生可自动验证。

评测两个轴： 1. 正确性（Correctness）：比对 Model 输出与 ModelNew 输出，每题用 5 组随机输入逐一比对（num_correctness = 5）。 2. 性能（Performance）：用重复试验对比 Model 与 ModelNew 的 wall-clock 执行时间，抵消计时波动。

[!TIP] 为什么用"随机输入 × 5"而不是形式化证明正确性？（讲透 + 数据） 作者引用停机问题（Halting Problem）：一般地判定两个程序是否等价是不可判定的（要检查所有输入的行为，而判定程序是否停机本身不可判定）。所以实践中用随机测试近似——对 AI kernel 尤其有效，因为它们控制流简单、焦点在数值正确。 为什么恰好 5 组？ 是"抓错能力"与"效率"的权衡。作者做过实验：100 个生成 kernel 里，50 个全对（5/5 且 100/100）、19 个数值不匹配（0/5 且 0/100）、4 个形状不匹配、10 个运行时错、17 个编译错。关键观察：那些错的 kernel 是 0/5 且 0/100——没有出现"部分正确"，说明 kernel 要么全对要么全错，5 组随机输入足以高概率抓住错误。（作者也承认：更系统的正确性评测、乃至形式化验证，是未来方向。）

[!TIP] 性能怎么测才可靠？（附录 B 的协议） 全部评测默认在裸金属 NVIDIA L40S GPU（Ada Lovelace，48GB HBM，300W）上，Python 3.10 / PyTorch 2.5.0+cu124 / CUDA 12.4。计时协议： - 测 nn.module.forward 的 wall-clock；确保 GPU 上无其他 CUDA 进程； - warmup 3 次，再用 torch.cuda.Event 计时 100 次（num_profile=100），取均值（也记 max/min/std）； - 变异系数 CV = std/mean 始终 < 3%，所以用均值做对比是稳的； - 加速比 speedup = T_Model / T_ModelNew（如 Model 2ms、ModelNew 1ms → 2× 加速）。

核心指标 fast_p 的定义：

\[ \text{fast}_p \;=\; \frac{1}{N}\sum_{i=1}^{N} \mathbb{1}\!\left(\text{correct}_i \;\wedge\; \{\text{speedup}_i > p\}\right) \]
符号 含义
\(N\) 任务总数（如某个 level 的题数，或全部 250）
\(\text{correct}_i\) 第 \(i\) 题生成的 kernel 是否功能正确（对 5 组随机输入都匹配参考输出）
\(\text{speedup}_i\) 第 \(i\) 题的加速比 = \(T_{\text{Model}} / T_{\text{ModelNew}}\)（PyTorch wall-clock ÷ 生成 kernel wall-clock）
\(p\) 可调加速阈值——只有加速比严格大于 \(p\) 才算过
\(\mathbb{1}(\cdot)\) 指示函数：括号内为真取 1、否则取 0
\(\text{fast}_p\) "既正确、加速比又 > \(p\)"的任务占比

[!IMPORTANT] 把 fast_p 讲透 + 数值举例 逐项物理含义：fast_p 用一个合取（AND） 把"正确性"和"性能"两个轴同时卡住——一个 kernel 必须先正确、再够快才被计入。这正是性能代码生成的真实目标：算错的快 kernel 毫无价值，算对但更慢的 kernel 也没优化意义。 - 两个特殊值： - \(p=0\)：fast_0 退化为纯正确率（任何加速比 > 0 都成立，等价于"只要正确"）——衡量"有多少 kernel 是对的，不管快慢"。 - \(p=1\)：fast_1 = 正确且加速比 > 1×，即"正确且不慢于 PyTorch"——本文的主用阈值。 - 数值举例：某模型在 Level 1（100 题）生成的 kernel 中，60 个正确（→ fast_0 = 60%）；这 60 个里只有 15 个加速比 > 1×（→ fast_1 = 15%）；只有 2 个加速比 > 2×（→ fast_2 = 2%）。可见 fast_p 随 p 增大而单调不增——p 越高越难，这就是"难度随 p 上调而增加"的机制。 - 为什么可调 p 很有用：① 随着 kernel 生成方法进步，可以上调 p 让 benchmark 不饱和（"beyond PyTorch baseline"）；② 训练时用 p < 1 也有价值——因为 PyTorch 依赖复杂优化 kernel，哪怕只匹配它一部分性能也算有益，p<1 给了这种"部分信用"。

[!TIP] fast_p 的两个变体：fast_p@k 与 fast_p@N（测试时方法用） 为评测"重复采样"和"迭代精修"这两种测试时方法，作者扩展了指标： - fast_p@k（重复采样）：抽 k 个样本时，至少有一个正确且加速 > p 的任务占比。 - fast_p@N（迭代精修）：到第 N 轮时，截至该轮生成的最好 kernel正确且加速 > p 的任务占比。 两者都是"给更多预算（k 个样本 / N 轮反馈）后能解多少题"的度量。

4 评测结果

4.1 One-shot baseline：开箱即用远不足

设置：给模型一个 prompt，含一个 PyTorch Model→ModelNew 的示例（示例极简，只有一个 add 算子，见附录 C.1）+ 待优化的目标 Model，用贪心解码（temperature=0） 生成，在 L40S 上测 fast_p。

主结果（Table 1，fast_1）：LM 生成的 kernel 平均在不到 20% 的任务上超过 PyTorch Eager。

模型 L1 (vs Eager) L2 L3 L1 (vs compile) L2 L3
GPT-4o 4% 5% 0% 18% 4% 4%
OpenAI o1 10% 24% 12% 28% 19% 4%
DeepSeek V3 6% 4% 8% 20% 2% 2%
DeepSeek R1 12% 36% 2% 38% 37% 2%
Claude 3.5 Sonnet 10% 7% 2% 29% 2% 2%
Llama 3.1-70B 3% 0% 0% 11% 0% 0%
Llama 3.1-405B 3% 0% 2% 16% 0% 0%

关键读表：推理模型（o1、R1）明显领先，尤其 R1 在 Level 2（fusion 题）拿 36%；但没有模型在任何 level 稳定超过 ~25%，Level 3（整模型）几乎全线个位数。Llama 系无论 70B 还是 405B 都近乎全 0——模型大小不保证 kernel 能力。（注：torch.compile baseline 在 Level 1 有时比 Eager 还慢，因为小 kernel 的可复现运行时开销，故正文主用 Eager。）

4.2 正确性：错误模式分析

Figure 2: kernel 失败模式——执行失败 vs 功能正确性

Figure 2 逐元素解读：横条形图，每行一个模型，横轴是"在全部 250 题中错误的百分比"，颜色把错误分成两大类： - 左半「Execution Failure」（执行失败，暖色系）：NVCC Error（编译器错，黄）、Python Error（浅黄）、CUDA Memory（内存违规，粉）、Runtime Error（运行时错，红）——即 kernel 根本没跑成。 - 右半「Functional Correctness」（功能正确性，冷色系）：Output Shape Mismatch（形状不匹配，深紫）、Output Value Mismatch（数值不匹配，浅蓝）——即 kernel 跑成了但算错了。 - 右端 ✓：绿色部分是正确的 kernel。

关键信息：推理模型（o1、R1）的执行失败明显更少（暖色段短），但所有模型的功能正确性都差到类似程度（冷色段都很长）。作者据此拆解：o1/R1 的总错误率更低（<55%）主要是因为它们编译错误更少，而不是因为它们算得更对——功能正确性是所有模型共同的、更深的瓶颈。作者的假设：CUDA 是低资源语言，只占 The Stack v1.2 代码语料的 0.073%，训练数据稀缺是根因。同时观察到一个权衡：模型越是尝试复杂优化 / 冷门硬件指令（如 tensor core wmma），越容易产生错误 kernel。

4.3 性能：加速比分布与 fast_p 随 p 的变化

Figure 3: fast_p 随加速阈值 p 增大的分布（三个 level）

Figure 3 逐元素解读（本文的招牌指标图，一张图看懂 fast_p 的难度曲线）：三个子图对应 Level 1/2/3，横轴是加速阈值 p（0 到 3），纵轴是 fast_p score，每条线是一个模型。读图要点： - 每条线从左（p=0，即纯正确率 fast_0）到右（p 增大）单调下降——这就是"难度随 p 增加"的可视化。 - 左端 p=0：Level 1 最高的 R1/o1 约 0.5–0.65（正确率），Level 2/3 稍低。 - p=1 处（正确且 >1× 加速）：所有 level 都 < 15% 的 kernel 能超过 PyTorch。 - 推理模型（o1 黄、R1 浅黄）的线整体压在其他模型上方——它们不仅正确率高，提供加速的能力也更强。 - Level 2 的曲线在中段（p≈1–2）比 Level 1 更"挺"——因为 Level 2 有更多 fusion 带来的真实加速机会；Level 3 则大多数模型在 p>1 后迅速趋零。

补充（附录 B.3，Figure 7 的箱线图，未嵌入）：Level 1 和 Level 3 的正确 kernel 中位加速比 < 1（即多数正确 kernel 反而更慢），Level 2 中位数略高于 1；Level 1 有最显著的离群点（个别 >10× 加速）。

4.4 硬件泛化性

one-shot baseline 不对硬件做任何假设，那生成的 kernel 跨 GPU 泛化吗？发现：Level 1 上跨 GPU 加速比相似（在 L40S 上超过 Eager 的 kernel 在其他 GPU 上也大致超过）；但 Level 2 上跨 GPU 波动大——R1 生成的 kernel 在 L40S 上 fast_1=36%，换到 A10G 却是 47%。这说明 one-shot 生成的 kernel 不一定跨硬件泛化，也引出了 §5.2 的"喂硬件信息"实验。

4.5 测试时方法之一：重复采样（Repeated Sampling）

设置：每题从 LLM 高温采样多个 kernel，用 fast_p@k 衡量（Figure 4，此处未嵌入）。结果：随 k 增到 100，fast_1@k 在三个 level 上都提升（DeepSeek-V3、Llama-3.1-70B）。最亮眼：DeepSeek-V3 在 Level 2 用 k=100 达到 fast_1=37%，而 one-shot 只有 4%。机制：高温采样探索解空间，增加"生成无错误 + 更好优化 kernel"的机会。但有天花板：若模型对某题的固有解出概率极低，加大采样预算收效甚微——例如 DeepSeek-V3 对 Level 1 的一组 34 个卷积变体，即使 100 次采样也从未生成过任何正确解。

4.6 测试时方法之二：迭代精修（Iterative Refinement）——最重要的增益

这是本文关于"如何提升 LLM kernel 能力"最有价值的一节。KernelBench 环境天生适合收集编译器反馈、执行错误、profiler 计时作为 ground-truth 信号喂回给模型。

Figure 5: KernelBench 框架让模型在迭代精修中接收并利用反馈

Figure 5 逐元素解读（迭代精修的反馈管线）：一条从左到右的流水线，每一关都可能把反馈打回给模型： - Problem → Model → Generated Kernel：模型先生成一个 kernel。 - 第一关「Compiles?」（能编译吗）：若否，把 NVCC 编译错误信息打回「Feedback Context」（左下黄块）。 - 第二关「Correct?」（对吗）：若否，把 Output Mismatch（输出不匹配） 信息打回。 - 第三关「Performant?」（快吗）：把 Execution Time + PyTorch Profiler（算子级计时分解）打回。 - 「Feedback Context」→ 回到 Model：模型带着上一轮生成 G + 编译/执行反馈 E +（可选）profiler 反馈 P，生成下一轮改进的 kernel。 - 含义：这三关正好对应 kernel 的三个失败层次——编译 → 正确 → 性能，反馈让模型逐层自我纠错。

每轮多轮对话：给模型上一轮生成 G + 编译/执行反馈 E +（可选）profiler 输出 P，跑 N 轮，测 fast_p@N。

Figure 6: 迭代精修中 G / G+E / G+E+P 三种反馈的 fast_1@N 轨迹（DeepSeek-R1 Level 2）

Figure 6 逐元素解读（本文测试时方法的头条结果图）：横轴是迭代轮数 N（1→10），纵轴是 fast_1 Score (%)，三条线对应三种反馈组合： - 黄线「Last Generation G only」：只给上一轮代码——从 ~27% 缓慢爬到 ~41% 就平台化（模型看不到错在哪，只能靠自己回顾代码）。 - 绿线「G + Execution Result E」：加上编译/执行反馈——爬到 ~62%，明显高于只给 G。 - 蓝线「G + E + Profiler P」：再加 profiler 计时——爬得最高，~72%，且前几轮就快速拉升。 - 关键信息：反馈信号越丰富，迭代精修越有效——从"只回顾代码"到"加执行反馈"到"再加 profiler"，fast_1 单调抬升。这与 [[ref09_meta-harness]] 的核心论点（更丰富的诊断反馈 → 更好的自改进）在 kernel 场景上完全共振。

Table 2（N=10 轮 fast_1，三种反馈 vs 两种测试时方法）核心数字：迭代精修在跨模型/level 上一致提升。最亮眼 DeepSeek-R1 Level 2：36% → 72%（G+E+P）。三个 level 的整体提升（引言里的招牌数字）：12%/36%/12% → 43%/72%/18%。

[!NOTE] 迭代精修 vs 重复采样（Table 2 的对照，同样 10 次调用预算） - 两种方法都显著超过 one-shot baseline；迭代精修在 6 个 case 里的 5 个更有效。 - 执行反馈 E 尤其擅长修执行错误：DeepSeek-R1 在 Level 1/2 能在 10 轮内对 >90% 的题生成一个功能可跑的 kernel（Table 9）。 - 但剩下的错几乎都是"功能不正确"——因为正确性反馈比执行错误信息粒度粗（"输出不匹配"不像编译错那样直接指出哪行错）。 - 一切受制于基座模型质量：R1 能持续从 E/P 获益，而 V3/Llama-70B 不总能——弱模型给了反馈也未必会用。

4.7 测试时方法之三：喂硬件知识与硬件规格

设置：在 prompt 里加（1）硬件效率技巧的 in-context 示例（GeLU 用 fusion、matmul 用 tiling、一个极简 FlashAttention 展示共享内存 I/O 管理），或（2）硬件规格（GPU 型号、内存大小/带宽、各精度 TFLOPS、寄存器数、共享内存容量……）+ 硬件概念（thread/warp/thread-block/SM 定义）。

结果（有点反直觉）： - 喂 in-context 示例反而降低整体 fast_1——因为 LLM 被激发去尝试更激进的优化策略，导致更多执行失败（o1 的生成平均长 25%）。但在正确的解里能看到有趣优化：o1 在 77% 的 Level 1 GEMM 变体上用了 tiling（虽因缺 tensor core 仍慢于 Eager），在 Level 2 的 11 道题上用了激进的共享内存 I/O 管理并超过了 Eager。 - 喂硬件规格效果有限：Llama/V3 基本没变化；但 o1 和 R1 的一部分生成用了硬件专用指令——R1 对约 50% 的 Level 1 矩阵乘题尝试生成 wmma（warp 矩阵乘累加）指令，虽大多编译失败；正确的里 R1/o1 每 level 产 1–3 个 ≥2× 加速的离群点。 - 总结：模型很少生成真正为底层硬件优化的 kernel——喂 few-shot 示例比喂硬件规格更能让模型调整策略，但整体仍是巨大 headroom。

4.8 有趣的成功 kernel（定性，附录 D）

论文展示了几个 LLM 生成的、显著超过 PyTorch 的 kernel，分三类：

类别 例子 加速 做对了什么
算法优化 torch.diag(A) @ B（对角阵×矩阵，L1 P11，Claude-3.5） 13× 识别出对角矩阵不必显式构造——直接把向量 A 的每个元素乘到 B 的对应行，避开加载对角阵的零元素
算子融合 GeLU（L1 P87，DeepSeek-V3） 2.9× 把 GeLU 的多步计算融进单 kernel，还做了常量折叠（预算 sqrt(2/π)）
Softsign x/(1+|x|)（L1 P29，Claude-3.5） 1.3× 融合成单 kernel
matmul+divide+sum+scale（L2 P13，Claude-3.5） 2.6× 把 4 个算子融进一个 kernel
硬件特性 Cosine Similarity Loss（L1 P96，o1） 2.8× 用共享内存做归约，减少冗余全局内存访问——含同步点和归约，人类都难写对
Triplet Margin Loss（L1，o1） 2.0× 用共享内存

[!IMPORTANT] 但作者明确指出：没有找到任何成功使用 tensor core 指令的例子——而 tensor core 对 AI 性能至关重要。这是当前 LLM kernel 能力最大的空白之一，也是 headroom 的核心所在。同时"LLM 把很多 fusion 机会留在了桌上"（Level 2 融合不足）。

5 局限与讨论

个人思考

与其他论文的关联（放进本项目坐标系）

方法论启示（可迁移到"造性能类评测"的通用思路）

  1. "正确 × 效率"的双判指标（fast_p）是性能代码生成评测的正确范式——只测正确性（HumanEval 式）会漏掉"对但慢"的一大类失败；用一个合取把两轴卡死、再用可调阈值 p 控制难度/防饱和，是很干净的设计。任何"生成更优解"的任务（不只是 kernel）都能套这个模板。
  2. evaluation-only + 随机测试验证：不提供标准答案、用"随机输入 × N 组比对参考实现"做正确性 gate——既回避了"程序等价不可判定"，又让 benchmark 可在任意硬件/新平台复用。这是"没有 ground-truth 但有 reference implementation"场景的通用招法。
  3. 反馈的层次化（编译→正确→性能）：Fig 5 把失败分成三关、每关喂不同粒度反馈，是"给 agent 设计反馈通道"的好范式——它解释了为什么 profiler 反馈（最细粒度）带来最大增益。
  4. 难度随阈值单调可调 = 抗饱和：多数 benchmark 会被刷爆，fast_p 通过上调 p / 上调 baseline 让 benchmark 随方法进步而自动变难，值得任何长期 benchmark 借鉴。

在我的工作中能怎么用

开放问题 / 疑问

局限性（对基准本身）