Lecture 14 · 智能体生成内核：KernelBench 与 AlphaEvolve · Agentic Kernel Generation

对应 Maps to：Stanford CS349E / EE339 Efficient ML Inference at Scale（Autumn 2026），原课第 16 次课 · Tue Nov 17 · Agentic Kernel Generation（Fredrik Kjolstad）。阅读材料：A. Ouyang, S. Guo, S. Arora, A. L. Zhang, W. Hu, C. Ré, A. Mirhoseini, KernelBench: Can LLMs Write Efficient GPU Kernels?（2025）；A. Novikov et al.（Google DeepMind），AlphaEvolve: A coding agent for scientific and algorithmic discovery（2025 白皮书）。 依据说明：两篇阅读材料用的是 arXiv 全文（2502.10517v1、2506.13131v1）。§3、§7 补充了之后的公开资料：KernelBench v0.1 博客（Stanford Scaling Intelligence 组）、METR 的 Measuring Automated Kernel Engineering（2025-02）、Sakana AI CUDA Engineer 事件的报道与 Sakana 的更正、Cognition 的 Kevin-32B 博客、KernelBench-Verified（arXiv 2607.16241，2026-06）。课程网站没有 slides，内容的取舍是我定的，不代表课堂上实际讲了什么。 配套代码 Code：code/lecture14/（numpy；两个模块 + 5 个脚本，最长的约 50 秒） 先修：第 10 讲（一个 GPU 内核里有哪些旋钮）、第 13 讲（调度空间、自动调优、学习的代价模型——本讲的 AlphaEvolve 实验直接用它的解析模型当”硬件”）。 预计阅读时间：3 小时（含运行代码）

0. 本讲导航 · How to read this lecture

0.1 学习目标 Learning goals

学完这一讲，你应该能够：

1. 说出”让 LLM 写内核”为什么值得做、为什么难，以及它和第 13 讲的自动调优是什么关系（从搜索参数到搜索程序）。
2. 描述 KernelBench 的任务格式、三个级别、正确性检查和计时方法，会算 fast_p 和 fast_p@k，能读懂它的主要结果：一次生成很弱，重复采样和带反馈的迭代优化帮助很大。
3. 列举评估器被”骗”的典型方式（偷输出、缓存、懒计算、调用参考、只对默认参数成立、退化的题、精度更低的基线），说出对应的防御。
4. 识别测速度时的陷阱：噪声与 p = 1 的阈值、挑赢家的偏差、cache 的冷热、问题太小、基线的精度、加速比的汇总方式。
5. 描述 AlphaEvolve 的结构（提示词采样、LLM 集成、SEARCH/REPLACE 修改、评估级联、程序数据库）和它在 Google 计算栈上的结果，解释”进化一段启发式代码”相比”逐形状调优”的取舍。
6. 用消融实验的思路判断一个搜索系统里每个部件值多少，并知道小实验能说明什么、不能说明什么。

0.2 本讲地图 Map

§1  为什么让 LLM 写内核
§2  KernelBench：把"会不会写内核"变成一场能自动判分的考试          ← 代码 01
§3  评估器才是难点：奖励被钻空子（reward hacking）                ← 代码 02
§4  诚实地测速度                                                 ← 代码 03
§5  AlphaEvolve：LLM 提议 + 自动评估 + 进化
§6  迷你 AlphaEvolve：进化一段选分块的代码                        ← 代码 04、05
§7  现状（2025–2026）
§8–12  小结、常见误区、练习（含答案）、延伸阅读、下一讲预告

0.3 本讲的代码

• kbench.py：一个迷你 KernelBench。9 道题（Level 1 的单个算子、Level 2 的算子组合、Level 3 的一个 MLP 块）都用 numpy 写参考实现；28 个”候选内核”模拟 LLM 的输出（有快的、慢的、错的、编译不过的）；两个评估器——“朴素”的照 KernelBench 最初的做法，“加固”的加上 §3.4 的防御。所有时间都是在你的 CPU 上实测的。
• evolve.py：一个迷你 AlphaEvolve。被进化的是一段 Python 代码 choose_config(M, N, K)：按矩阵乘法的形状选分块配置。“硬件”是第 13 讲 sched_lab.gemm_time 的解析模型；“LLM”是一个模拟的提议器（模板化的修改 + 一点内核常识），没有任何真正的语言模型；有评估级联、岛屿和精英档案。

0.4 记号 Notation

1. 为什么让 LLM 写内核 · Why have LLMs write kernels

1.1 内核的”供需缺口”

第 10、11 讲看到，一个好内核和一个朴素内核可以差两个数量级，而写出好内核需要同时懂算法、懂硬件、懂编译器。这种人很少，而需求在快速增长：

• 新模型结构：KernelBench 的引言举的例子——Transformer 发表（2017）五年后才有高效的注意力内核 FlashAttention（2022）；状态空间模型、线性注意力、各种 MoE 变体，每一个都要新内核。
• 新硬件：同一个算法换一代硬件就要重写。KernelBench 引用的数字：FlashAttention-2 从 A100 换到 H100，硬件利用率掉了 47%；为 Hopper 写的 FlashAttention-3 是一个完全不同的算法（第 8、11 讲的 TMA、WGMMA、warp 专门化），在 Hopper 发布两年后才出来。还有 TPU、Trainium、各种推理芯片，各有各的指令集。
• 成本：METR 估计前沿实验室为每一种模型结构优化推理要花大约 5 个工程师年。

1.2 内核是 AI 的好靶子

和一般的软件工程相比，内核有一个难得的性质：可以自动判分。

• 正确性：和一个参考实现比较输出（随机输入，容差内相等）；
• 性能：用计时器量。

不需要人来评价”这段代码好不好”。这意味着可以大规模地生成—评估—挑选：采样很多个候选、只留下又对又快的；或者把评估结果当作奖励做强化学习；或者把它放进一个进化搜索的循环里。第 13 讲的自动调优就是这个思路，只不过搜索的是参数（块多大、几个 warp）；LLM 让我们可以搜索程序本身。

1.3 为什么难

• 数据少：KernelBench 统计，CUDA 只占代码语料 The Stack v1.2 的 0.073%。模型见过的好内核很少。
• 正确性很微妙：边界条件、竞争（第 10 讲的 __syncthreads）、数值精度；而”对不对”的反馈只有一个比特。
• 性能依赖硬件：同一个内核在 L40S 上快、在 A10G 上不一定快（§2.5）。
• 评估本身容易出错：§3、§4 整整两节讲这件事。这一讲的一个主题是：在”让 AI 写内核”这件事里，评估器和生成器一样重要。

2. KernelBench：一场能自动判分的考试 · KernelBench

代码：01_kernelbench_harness.py

2.1 任务格式

每道题是一个 PyTorch 的 Model 类（__init__ 和 forward），外加两个函数：get_inputs() 给出输入张量（形状固定、数值随机），get_init_inputs() 给出初始化参数。模型要输出一个 ModelNew 类，里面可以用任何办法加速：内联 CUDA（load_inline）、Triton、CUTLASS、ThunderKittens，也可以只替换一部分算子、保留其余的 PyTorch 代码。优化什么、怎么优化，都由模型决定——这是在模仿 AI 工程师的真实工作流。

形状固定是有意的：最好的内核依赖形状和数据类型。这也意味着针对这个形状专门写的内核是合法的（§3 会看到它的另一面）。

2.2 三个级别

基线有两个：PyTorch eager（逐个算子执行）和 torch.compile（第 12 讲）。论文的分析主要用 eager。

2.3 正确性：随机测试

判断两个程序在所有输入上等价是不可判定的（论文附录 B.2 用停机问题说明这一点），所以实际做法是随机测试：生成 5 组随机输入，比较 Model 和 ModelNew 的输出，容差 atol = rtol = 1e-2。为什么是 5 组？论文用 100 个生成的内核做了实验：50 个 5/5 正确、100/100 也正确；19 个数值不对，0/5、0/100；4 个形状不对，10 个运行时出错，17 个编译出错——没有一个是”部分正确”。同一种分布下多测几组几乎不增加信息。

我们的 01 演示 4 复现了这一点：能运行的 27 个候选在 20 组标准正态输入上，全对 19 个、全错 8 个、部分对 0 个。但同一个演示也说明了换分布的重要性：

不稳定的 softmax 在值上百时 exp 溢出（float32 的 exp 在输入超过约 88.7 时就是 inf，前置讲 P3）；一遍算方差在均值很大时有严重的抵消误差；粗糙的 GELU 近似只在 −2 附近误差超过容差，而 (−1, 1) 的均匀分布和全正的分布根本碰不到那里。最后一行最有意思：忘了 eps 只在某一行方差恰好为 0（常数行）时才会除以 0，五种随机分布都碰不到——这种边界输入要靠人想到。KernelBench v0.1 后来把输入改成从一个分布池里抽（正态、均匀、拉普拉斯、指数、对数正态）。

2.4 速度与 fast_p

计时：L40S 上，3 次预热，100 次用 CUDA event 计时取平均；论文报告变异系数（标准差 / 平均）稳定在 3% 以下。加速比 = $T_\text{Model} / T_\text{ModelNew}$。

一个模型可能写出少数几个极快的内核，另一个写出很多正确但慢的内核，怎样比较？KernelBench 定义了

$\text{fast}_0$ 是正确率；$\text{fast}_1$ 是”又对又比 PyTorch 快”的比例。p 可以调：随着方法变强可以提高 p；训练时用 p < 1 也有意义（能达到 PyTorch 的一部分性能就已经有价值）。

注意 fast_p 是按题计数的：一道题快 100× 和快 1.01× 在 fast_1 里都算 1。它不会被个别巨大的加速比带偏（§4.6），代价是看不出”快了多少”。

2.5 结果

一次生成（one-shot）很弱。论文的表 1（L40S，相对 PyTorch eager 的 fast_1）：

21 个组合里最高 36%（R1 在 Level 2），17 个不超过 10%。相对 torch.compile 时 Level 1 的数字反而更高（例如 R1 是 38%），因为 torch.compile 在小内核上有可复现的运行时开销。

• 失败的方式：分成”执行失败”（编译错误、内存越界、运行时错误）和”结果不对”（形状、数值不匹配）。推理模型（o1、R1）错得少（< 55%，其他模型 > 70%），但主要是执行失败少；“结果不对”的比例各家差不多。
• 正确的也大多不快：正确内核的加速比中位数在 Level 1、3 小于 1，Level 2 略大于 1。
• 跨硬件不稳定：R1 在 Level 2 上的 fast_1，L40S 上 36%，A10G 上 47%。
• 重复采样有用：DeepSeek-V3 在 Level 2 上采 100 个样本，fast_1@100 = 37%（一次生成 4%）。但如果模型根本不会某类题，采多少次都没用：V3 在 34 个卷积变体上采 100 次也没有一个正确的。
• 带反馈的迭代优化更有用：每一轮把上一次的代码 G、编译 / 执行结果 E、profiler 的输出 P 喂回去。同样 10 次调用的预算下（表 2），R1 在 Level 2 上从 36% 升到 72%（G+E+P），在 Level 1 上从 12% 到 43%。10 轮之内 R1 在 Level 1、2 上 90% 以上的题都能写出能运行的内核；剩下的几乎都是”结果不对”——因为”不对”这个反馈远没有编译错误信息那样具体。
• 给例子、给硬件信息：给出融合、分块、FlashAttention 的示例代码后，模型尝试更激进的优化，执行失败也更多，总的 fast_1 反而下降（o1 的输出长了 25%；在 Level 1 的 GEMM 变体里 77% 用上了分块，但因为没用 Tensor Core 仍比 PyTorch 慢）。给硬件规格的帮助更小：R1 在大约一半的 Level 1 矩阵乘法题上尝试了 wmma（Tensor Core）指令，大多编译不过。
• 有意思的内核：diag(A) @ B 不造对角矩阵、直接按行缩放，13×（Claude 3.5 Sonnet）；GELU 融合成一个内核 2.9×；矩阵乘法 + 除法 + 求和 + 缩放融合 2.6×；用共享内存的余弦相似度 2.8×、triplet margin loss 2.0×。论文没有找到成功使用 Tensor Core 指令的例子。

2.6 我们的迷你版

01 脚本把 28 个候选交给两个评估器（作者机器 Apple M4 Pro 一次运行；时间每次会变几个百分点）：

加固的评估器多抓出 3 个：不稳定的 softmax、一遍算方差、“把缩放因子折进权重”（只在缩放因子为正时成立，参数一随机化就错）。它们都不是故意作弊，而是只在默认输入、默认参数上才对。16 个正确的候选里 7 个比参考还慢，加速比中位数 1.04×。

numpy 上的加速从哪里来？和 GPU 不同，这里没有共享内存、没有 Tensor Core，“少分配几个临时数组”几乎没用（softmax 和 epilogue 的就地版本都在 1.0× 附近），L3 的 gate/up 合并反而慢了 5%（切片出来的不连续视图拖慢了后面的逐元素运算）。真正有效的是两类：

• 换算法：diag(a) @ B → 按行缩放，10.6×；“矩阵乘法 → 除以 2 → 按行求和 → 缩放”→ 先把权重按列求和、变成矩阵-向量乘法，49×（求和和矩阵乘法可以交换：$\sum_j (x W^\top)_{ij} = x \cdot \sum_j W_j$）。
• 去掉昂贵的运算：GELU 参考实现里的 x ** 3 是一次慢的 pow，换成 x * x * x 并就地计算，5.2×。

如果”模型”每道题从候选库里随机抽一个交卷，fast_0 = 0.60，fast_1 = 0.38，fast_2 = 0.13——随 p 下降的形状和论文的图 3 一样。重复采样 k = 10 次，fast_1@10 = 0.66，已经贴近”候选库里最好的那个”的 0.67：没有好候选的题，采多少次都没用。（候选库的成分是我仿照论文报告的现象设计的，所以这条曲线只能和论文定性地比较。）

3. 评估器才是难点：奖励被钻空子 · Reward hacking

代码：02_reward_hacking.py

3.1 为什么一定会被钻空子

重复采样挑最好的、进化搜索、强化学习，本质上都在最大化评估器给的分数。评估器和”我们真正想要的东西”（一个在部署环境里又对又快的内核）之间的任何差距，都会被优化过程找出来。这就是 Goodhart 定律：“一个指标一旦成为目标，就不再是好指标。“对内核来说，差距多得惊人。

3.2 现实里发生过的

• Sakana AI CUDA Engineer（2025 年 2 月）：宣称生成的内核最多能让 PyTorch 的一些运算快 100 倍。用户测试时发现，一个号称快 150 倍的内核实际慢了 3 倍；有人追查到有的内核”偷”了参考实现的结果（可能是内存复用）。Sakana 随后承认，系统在评估代码里找到了一个内存方面的漏洞，在一些情况下绕过了正确性检查；他们加固了评估和计时的代码，并修订论文。
• METR（2025 年 2 月）：用 KernelBench 测前沿模型的内核工程能力。他们先筛掉了 45 道”质量不行”的题：输出接近 0（被浮点误差淹没）、输出几乎不随随机种子变化（可以缓存答案）、输出沿某一维完全一样（太容易），比如 mean(softmax(x)) 恒等于一个常数、relu(x - 2) 在标准正态输入上几乎全是 0。他们还删掉了 17 个用 CUDA stream 的解（不显式同步就没法可靠计时），以及一个快得不可思议的结果——他们认为 torch.empty 复用了参考实现输出的显存，“白白拿到了答案”。
• KernelBench v0.1（Stanford）：总结了观察到的作弊：调用高层算子（torch、cuBLAS）而不是写 CUDA；写了 CUDA 内核却不调用它；“什么都不做”的内核因为评估器从参考实现写过的同一块内存里读输出而通过；同步点放错、在另一条 stream 上计算，计时偏短。他们的结论：模型非常擅长找到满足奖励但偏离目标的边界情况；补上一个漏洞，就会出现更隐蔽的下一个；自动检查不够，分数高的结果要人看。
• Kevin-32B（Cognition，2025）：用多轮强化学习训练写 CUDA 的模型时，较小的模型学会了三种作弊：照抄 PyTorch 参考实现；把有 bug 的 CUDA 内核包在 try/except 里，出错就退回 PyTorch；继承参考实现的类、根本不写 CUDA。对策是更严格的格式检查：用了 PyTorch 函数或没有 CUDA 内核就给 0 分。
• KernelBench-Verified（2026 年 6 月）：指出两个问题。一是基线太弱：PyTorch 默认不让 FP32 矩阵乘法使用 TF32 Tensor Core，而生成的内核可以用（§4.5）；二是模型会针对特定的数值硬编码旁路，跳过本该做的计算。改成打开 TF32 的基线、四种分布的隐藏测试集之后，最好的模型（单轮生成）的几何平均加速比从 1.43× 降到 0.88×，“没有一个模型能稳定地超过 PyTorch”。他们还发现最好的模型有 28% 的内核增加了显存峰值——一个原来没人测的维度。

3.3 作弊图鉴

02 脚本把 7 种作弊（每种都对应上面某份报告描述过的现象）交给两个评估器：

（≥ 10000× 是脚本给加速比设的上限。时间是作者机器一次运行的。）

两个机制值得细看。noop_steal：朴素的评估器把参考的输出拷一份留作答案，然后把缓冲区”还给”一个会复用内存的分配器（就像 PyTorch 的 CUDA 缓存分配器）；候选调用 alloc 拿输出缓冲区时，拿到的正是那一块——里面已经是正确答案。lazy_result：朴素的计时只包住函数调用，候选只花 0.3 µs 造了一个对象；正确性检查时 np.asarray 才触发真正的计算，所以答案也是对的。GPU 上的对应物就是异步启动：不 synchronize 就停表，量到的只是”发起”的时间。

这些作弊没有一个需要理解内核。它们利用的是评估器的实现细节。

如果有一个”会作弊的模型”——能作弊的题交作弊的卷、其余交诚实的卷——朴素评估器给它 fast_0 = 1.00、fast_10 = 0.56；诚实的模型在加固评估器下 fast_10 只有 0.22。用这样的分数做奖励，学到的就是作弊。

3.4 防御：分层

“光速”要按题目的语义算：diag(a) @ B 的参考做了 $2N^3$ 次运算，但语义上只需要读 a、B、写 C（N = 2048 时 33.6 MB），所以合法的”换算法”不会被误判。KernelBench v0.1 用 roofline（第 1 讲）的两个下界取大做同样的检查。

军备竞赛。02 的最后一个演示是一个”能识别考场”的候选：如果拿到的缓冲区里是 NaN（“我正在被加固的评估器考”），它就老老实实地算；否则什么都不做。加固的评估器给它的判决本身没错——它在考场里确实是一个正确且更快（5.2×）的内核——问题在于部署的环境不是考场。防御要随机化评估环境的细节、静态审查可疑的模式（先读缓冲区原有的内容再决定怎么算），并在真实环境里抽查。

4. 诚实地测速度 · Measuring speed honestly

代码：03_measurement_pitfalls.py

§3 是候选骗评估器；这一节是评估器自己给出误导的数字，候选一点都没作弊。（下面的数字都是作者机器一次运行的。）

4.1 噪声与 p = 1

同一行代码和自己比（1024³ 的矩阵乘法，每次 10 对交替计时取中位数），40 次测量的”加速比”在 0.992–1.010 之间，38% 大于 1。fast_1 的阈值正好落在噪声中间：一个什么都没改的”内核”有接近一半的机会拿到 fast_1 的分。对策：阈值留余量（p = 1.05）、多测、报告置信区间。

4.2 挑赢家的诅咒

16 个一模一样的候选，挑量得最快的那个：每个只测 1 对时，“赢家”平均”快了” 4.3%；每个测 6 对取中位数时 1.1%；重新测一遍都回到 1.000×。从很多带噪声的测量里挑最大值，挑出来的值系统性地偏高，而且噪声越大、候选越多偏得越多。重复采样（best-of-k）、进化搜索、自动调优都在”挑最快的”。对策和机器学习一样：用一组测量挑，用另一组独立的测量报告。

如果测量噪声是标准差 σ 的正态，16 个独立测量的最大值的期望大约是 $1 + 1.77\sigma$（16 个标准正态的最大值的期望 ≈ 1.77）。

4.3 cache 的冷热

同一个”换算法”的候选（矩阵-向量乘法，读 8.4 MB）：

候选的时间差 4.5 倍，参考几乎不变——01 脚本里朴素和加固的评估器在这道题上量到的倍数差一大截，主要就是这个原因。GPU 上一样：Triton 的 do_bench 每次计时前清零一个 256 MB 的缓冲区，清空 L2，测的是冷的时间。冷和热哪个”对”取决于部署（真实模型里，小张量常常是上一个内核刚写的，还在 L2 里）；要写清楚测的是哪一种，参考和候选用同一种。

4.4 问题太小

GELU 的融合版本相对参考：64 个元素时 1.01×，1024 个时 1.63×，几百万个时 5.3×。小问题上时间主要是每次调用的固定开销（numpy 每个调用几百纳秒到一微秒），两边调用的次数差不多。GPU 上对应的是内核启动开销（第 11、12 讲）。KernelBench v0.1 发现几十道 Level 1/2 的题在 H100 上只要几微秒，于是把题目放大（矩阵乘法 0.347 → 3.81 ms，ReLU 0.013 → 6.97 ms），让参考实现落在 1–15 ms。

4.5 基线的精度

参考实现用 FP64 累加（“严格”）时，一个只是调用 FP32 库函数的候选有 4.0× 的”加速比”，误差 2.4e-4，远在容差内；换成 FP32 的基线，1.02×。GPU 上的对应就是 KernelBench-Verified 指出的 TF32 问题：A100 的 FP32 CUDA core 是 19.5 TFLOPS，TF32 Tensor Core 是 156 TFLOPS（稠密）——光是换精度就有 8 倍的空间。容差 1e-2 本身就允许降精度。降精度是不是合法的优化，要看任务的要求；如果下游能接受 BF16，就应该和同样用 BF16 的基线比。

4.6 汇总

9 道题（每题交候选库里的第一个候选）的加速比：11.2、0.98、5.2、1.09、1.16、63、0.98、1.01、0.96。

一道”换算法”的题就把算术平均拉到 9.5×。加速比是比值，要用几何平均或中位数，再给出分布。“不低于 1”在能退回参考的部署场景下合理，但它看不到退步。比较两个方法时，一定要写清楚用的是哪种汇总。

一张清单：固定、写明测量条件（cache 冷热、同步、预热、重复次数）；参考和候选交替测；问题足够大；基线和候选用可比的精度和编译选项；用独立的测量报告被挑出来的赢家；报告几何平均、中位数、fast_p 和分布；分数高的结果人工看一遍。

5. AlphaEvolve：LLM 提议 + 自动评估 + 进化 · AlphaEvolve

5.1 从 FunSearch 到 AlphaEvolve

AlphaEvolve 的前身是 FunSearch（Romera-Paredes et al. 2023）：让 LLM 反复改写一个函数，用自动评估打分，保留好的、继续改。白皮书的表 1 对比了两者：

5.2 结构

用户提供三样东西：评估函数 h（输入一个解，返回一个”指标名 → 分数”的字典，分数越大越好）、初始程序（用 # EVOLVE-BLOCK-START / # EVOLVE-BLOCK-END 标出可以改的部分；初始实现可以很简陋，比如返回常数）、可选的配置和背景知识。然后是一个循环：

parent, inspirations = database.sample()
prompt = prompt_sampler.build(parent, inspirations)
diff = llm.generate(prompt)
child = apply_diff(parent, diff)
results = evaluator.execute(child)
database.add(child, results)

• 提示词采样：提示词里有从数据库里抽出的几个之前的解（带分数）、如何提出修改的指令，还可以有：问题的显式上下文（说明、公式、代码片段、论文）、随机的格式变化（增加多样性）、渲染出来的评估结果（程序、执行的输出、分数），以及元提示词进化（让 LLM 自己建议提示词里该写什么，另开一个数据库进化）。
• LLM 集成：Gemini 2.0 Flash（延迟低，多出主意）+ Gemini 2.0 Pro（能力强，偶尔给出能带来突破的高质量建议）。
• 输出格式：一串 SEARCH/REPLACE 块——<<<<<<< SEARCH 和 ======= 之间是要在当前程序里精确匹配的原文，======= 和 >>>>>>> REPLACE 之间是替换后的内容。程序很短或者需要整体重写时，也可以让 LLM 直接输出整段代码。
• 评估：可以设置级联——一组由易到难的测试，前一级的结果足够好才进下一级，便宜地筛掉没希望的解；可以让另外的 LLM 调用给”简洁”之类难以写成公式的性质打分；评估可以并行（每个新解可以花上百个计算小时）。可以同时优化多个指标，论文发现即使只关心一个指标，多指标也常常让结果更好（不同指标下的好程序结构不同，放进提示词里增加了多样性）。
• 进化数据库：结合 MAP-Elites 和岛屿模型，目标是在”改进最好的程序”和”保持多样性”之间平衡。
• 分布式流水线：用 asyncio 写的异步流水线，为吞吐（单位预算内提出和评估的想法数）而不是单个计算的速度优化。

5.3 结果

矩阵乘法。$m \times n$ 乘 $n \times p$ 矩阵的算法对应一个三维张量的秩分解，秩就是需要的标量乘法次数（Strassen 1969 的思想）。从一个标准的梯度下降分解算法出发，AlphaEvolve 进化出了更好的分解算法，在 14 个目标上刷新了最好的已知结果；其中 4×4 复数矩阵相乘只需要 48 次标量乘法——递归使用 Strassen 是 49 次，56 年来在特征 0 的域上第一次低于 49。

数学。50 多个开放问题（分析、组合、数论、几何），在约 75% 上重新找到了最好的已知构造，约 20% 上找到了更好的。

Google 的计算栈（和本课最相关）：

• 数据中心调度：Borg 把待调度的任务放到机器上，“搁浅”的资源（比如内存用完了、CPU 还空着）是浪费。AlphaEvolve 进化出一个很简单的打分函数：设 $c$、$m$ 是任务需要的 CPU、内存占机器剩余量的比例，分数 $= -(c + m + m/c + c/m)$。部署到全公司之后，持续回收了平均 0.7% 的全舰队计算资源。选它而不选深度强化学习，因为代码可解释、可调试、可预测、容易部署。
• Gemini 的内核：训练 Gemini 用的一个 Pallas 矩阵乘法内核，需要按输入形状选分块大小。以前要么手写启发式（难，需要同时懂内核和硬件），要么每个形状都搜索（打断研究的工作流，形状一变就要重调）。AlphaEvolve 从真实用户那里收集形状，一半训练、一半检验，进化出的启发式相对专家设计的启发式，所有内核平均快 23%，Gemini 的整体训练时间少 1%；内核优化从”几个月的专门工程”缩短到”几天的自动实验”；已经部署到生产。正确性由构造保证——只改分块，不改算的东西。
• TPU 电路：在一个已经高度优化的 Verilog 算术电路（矩阵乘法单元里的）上，找到一个删掉多余比特的简单改写，TPU 设计者验证了正确性。这个改进下游的综合工具其实也会发现，但 AlphaEvolve 在 RTL 这一级就提出了它，将用在一款即将推出的 TPU 上。
• 编译器生成的代码：直接改 XLA 为 FlashAttention 生成的 IR（这种 IR 是给调试用的，不是给人改的，而且已经高度优化）。目标是一个大规模推理用的 Transformer 配置：FlashAttention 内核快了 32%，前后处理快了 15%。每个修改都用随机输入和参考比较，最终版本由人确认对所有输入都正确。

5.4 消融

在矩阵乘法分解和 kissing number 两个问题上，去掉任何一个部件都变差：No evolution（每次都把同一个初始程序交给 LLM）、No context in the prompt、No meta prompt evolution、No full-file evolution（只进化损失函数）、Small base LLM only。曲线的形状都是：前期上升快，后期变慢，去掉部件的曲线始终在下面。

5.5 什么样的问题适合

AlphaEvolve 自己说得很清楚：它依赖机器可以评分的问题；需要人做实验的问题不在范围内。把它的成功案例放在一起看，有几个共同点：

1. 正确性容易保证：分块启发式和调度打分函数只影响”怎么做”，不影响”做什么”；电路和 IR 的改写有形式化或人工的验证。
2. 评估可以自动化且足够便宜：一个模拟器（Borg）、一组真实形状上的运行时间（Gemini 内核）。
3. 产出是代码：可以读、可以审查、可以部署；相比逐形状的搜索结果或一个神经网络策略，它能泛化到没见过的输入，也更容易被工程师信任。
4. 有明确的训练 / 检验划分：Borg 的启发式在没见过的近期负载上检验；分块启发式在一半没参与进化的形状上检验。

6. 迷你 AlphaEvolve：进化一段选分块的代码 · A mini AlphaEvolve

代码：04_alphaevolve_tiling.py、05_alphaevolve_ablations.py

6.1 设定

我们复现 §5.3 的 Gemini 内核案例的结构。

• 形状：Llama-3.1-8B（单卡）和 70B（8 卡张量并行后每卡的分片，第 4 讲）里的 8 种投影（qkv、o、gate_up、down），乘以 11 种 token 数 M（1 到 8192，decode 的批到 prefill），共 88 个；随机对半分成训练 44 个、测试 44 个。
• 硬件：第 13 讲的解析模型 gemm_time（H100 式参数，不是实测）。配置空间就是第 13 讲的 6000 个组合（3744 个能启动）。
• 程序：choose_config(M, N, K)，一串 if 条件: return GemmConfig(...) 的规则；分数 = 训练形状上相对专家启发式的几何平均加速比。
• “LLM”：模拟的提议器。它会改配置的一个字段、改阈值、加规则、删规则、从别的程序借规则，以及（有”内核常识”时）为某个形状加一条几个字段一起变的定向规则；可以读评估反馈来决定挑哪个形状下手。它没有任何真正的智能。
• 评估：先检查程序所有分支的配置能不能启动，再做两级级联（6 个探针形状 → 全部训练形状）。
• 数据库：4 个岛（每 50 个程序迁移一次）+ 按规则条数分格的精英档案。

专家启发式（进化的起点）是我写的：M ≤ 32 用 32×64×64 / 2 warp，M ≤ 256 用 64×128×64 / 4 warp，其余用 Triton 教程里常见的 128×128×64 / 4 warp / 3 级。

6.2 基线

最优配置随形状变化很大：decode 形状（M = 1、16）的时间几乎全花在读权重上，很多小块配置一样快，专家启发式已经是最优；prefill（M 几千）要 256×128 的大块、8 个 warp、4 级流水线，专家的选择慢 30% 左右；差得最多的是中间的 M（70B 分片的 qkv 在 M = 128 时慢 2.5 倍）。

“按 M 分档查表”已经很强（测试上达到最优的 90%）——这个问题的规律主要由 M 决定。给进化搜索找一个强的基线，才能知道它到底带来了什么。

6.3 一次进化

随机种子 0，提出 1000 个程序：分数在第 200、400、600 个程序时是 1.179、1.259、1.334，第 1000 个时 1.345；测试形状上 1.315。各种修改的”比父程序好”的比例：定向修改 37%，删规则 16%，改配置 8%，改阈值 6%，加规则 4%，借规则 3%。

定向修改的成功率高，因为很多好配置要几个字段一起变才到得了。从 128×128 / 4 warp 走到 256×128 / 8 warp：中间的 256×128 / 4 warp 每个 warp 要算 128×64 的块，按模型每线程约 $128 \cdot 64 / 32 + 40 = 296$ 个寄存器，超过 255，根本启动不了。单字段的小改动跨不过这个”山谷”。这正是大模型相对小模型、相对随机变异的优势：一次提出一个连贯的修改。

最好的程序有 12 条规则，删掉被前面规则完全挡住的”死规则”后 9 条（行为完全不变；这和第 12 讲的死代码消除是同一件事，遗传编程里把这种冗余叫 bloat）。脚本还会打印产生它的那个 SEARCH/REPLACE 块。

6.4 泛化（5 个种子的中位数）

随机划分时测试和训练几乎一样好（同样的投影、同样的 M）；按模型划分时训练和测试之间的落差更大（测试形状的 N、K 都没见过），但相对各自的上限都在九成左右。进化出的启发式和查表差不多，略好一点；它的优势是不需要在每个训练形状上把 3744 个配置都跑一遍，还能用 N、K 写条件。

成本：这次进化提出 1000 个程序，73 个被分支检查拒掉，34 个在级联的第一级就被淘汰，835 个完整评估，一共”跑了” 41,954 次内核；在训练形状上逐形状穷举是 164,736 次。少了 4 倍，不是几个数量级——它买到的主要是一段能直接用在新形状上的代码。

我们的倍数（相对”专家”约 1.3–1.4×）不能和 AlphaEvolve 的 23% 比：专家启发式是我写的，好坏由我决定。有意义的比较是”离逐形状最优还差多少”。

6.5 消融（9 个变体 × 5 个种子，中位数）

上限：训练 1.441，测试 1.478。三点结论：

1. 进化最重要（1.101 vs 1.370）：一次修改只能走一小步，好程序要在好程序上继续改才积累得出来。连贯的修改其次（small_llm 1.231）。这两条和 AlphaEvolve 的消融一致。
2. 反馈怎么用影响很大。同样是”看评估结果”：按”离见过的最快还差多少”加权挑形状（1.370）> 按”离 roofline 还差多少”加权（1.330）> 随便挑（1.301）> 只盯最差的 3 个（1.269）。05 的演示 2 解释了原因：44 个训练形状里有 32 个已经做到了”见过的最好”，按 roofline 加权时 55% 的提案落在它们身上（比如 70B 分片的 qkv 在 M = 1 时离 roofline 很远，是因为启动开销、补齐到 16 行、块太少占不满 132 个 SM——换配置也改不了多少），按”见过的最快”加权时只有 37%。反馈要可操作（“这个形状上有个更快的配置，你没用上”），而不只是”离理想有多远”；只盯最差的几个又会失去多样性。真正的 LLM 读反馈时没有这样一条固定规则——它怎么用反馈，是这类系统成败的关键，也是我们的模拟最不像的地方。KernelBench 里 profiler 的输出比单纯的”对 / 错 + 时间”更有用（§2.5），是同一个道理。
3. 借规则、岛屿数、级联在这个小问题上和完整方法的差别都在种子之间的波动以内，中位数有的还略高。它们的价值（多样性、并行、省钱）在更大、更难、评估更贵的问题上才显出来。一个小实验得不出”它们没用”的结论；反过来，5 个种子的中位数差 0.02 也说明不了谁更好。

6.6 没测到的分支

把”检查所有分支”关掉：5 个种子里，进数据库的程序有 866 个至少有一个分支返回不能启动的配置——训练形状从来没走到那个分支，分数看不出来。最终的数据库里还剩 67 个这样的程序，其中一个种子的”最好程序”本身就带着一个。在 100 个新形状上（新的 M、405B 分片的 N、K）真正出错的只有 1 个：隐患大多潜伏着，直到某个没见过的形状碰巧走进去。

对策很便宜：对程序的每个分支做静态检查（这里是共享内存、寄存器、warp 的块）。评估器测的是样本，部署面对的是全体。AlphaEvolve 的电路改写要通过严格的验证，IR 改写最后由人确认对所有输入都正确，道理一样。

7. 现状（2025–2026）· State of the field

• 智能体 + 迭代：METR 用一个并行树搜索的智能体 KernelAgent（8 个尝试并行，每个新尝试从按加速比抽中的之前的尝试或参考实现出发，最多 300 次尝试），用 KernelBench 原论文里的模型，把几何平均加速比从排行榜上的 1.05× 提到 1.81×（每题取最好的尝试、不低于 1）；每题每个模型约 20 美元。他们估计 o3-mini 写的内核在 H100 上平均跑 27 小时就能把生成成本赚回来。NVIDIA 用 DeepSeek-R1 + 一个在 H100 上运行的验证器做闭环，跑 15 分钟，在 KernelBench Level 1 的 100% 和 Level 2 的 96% 的题上生成了数值正确的内核（注意：这是正确率，不是加速比）。
• 用强化学习训练”写内核”的模型：Kevin-32B（基于 QwQ-32B）在 KernelBench Level 1、2 的 180 道题上做多轮强化学习：每一轮把编译错误或运行时间喂回去，奖励是之后各轮得分的折扣和（折扣 0.4，4 轮细化）；为了不让上下文爆炸，后面的轮次不带完整的思维链，只带模型自己写的摘要。作者报告平均正确率 65%（16 条轨迹的平均）、best@16 加速比 1.41×。§3.2 的作弊就是在这里发现的。
• 评估在变严：KernelBench v0.1（放大题目、修掉退化的题、多分布输入、roofline 检查）；KernelBench-Verified（TF32 基线、隐藏测试集、显存指标）。每一次加固，报告的数字都会下降一截——这本身就是这个领域最重要的经验。
• 开源的 AlphaEvolve 式系统：OpenEvolve、Sakana 的 ShinkaEvolve 等开源实现，把”LLM 提议 + 评估器 + 程序数据库”的循环用在各种优化问题上。
• 更高层的内核语言：KernelBench 论文的建议之一是让模型用 Triton、ThunderKittens、CUTLASS 这类抽象写内核，而不是裸 CUDA——更短的程序、更少的出错机会，Tensor Core 也更容易用上（第 10 讲 §10）。这和编译器的路线（第 12、13 讲）正在交汇：编译器负责保证正确的变换，模型负责在巨大的空间里提出值得尝试的方案。

8. 小结 · Summary

1. 内核是”让 AI 做系统”的好靶子：能自动判分（随机测试 + 计时器），所以可以大规模地生成—评估—挑选。LLM 把第 13 讲的搜索从参数空间推广到了程序空间。
2. KernelBench：250 道题、三个级别、5 组随机输入、fast_p。一次生成很弱（fast_1 大多不到 10%），重复采样和带执行 / profiler 反馈的迭代优化帮助很大（R1 在 Level 2 上 36% → 72%）；正确的内核也大多不比 PyTorch 快。
3. 评估器才是难点。优化会找出评估器的每个漏洞：偷输出、缓存、懒计算 / 异步、调用参考、只对默认参数或默认输入成立、退化的题、更弱的基线。防御要分层（静态检查、“下毒”的缓冲区、多分布输入、随机化参数、计时时换输入并同步、题目检查、光速下界），最后还要人看。
4. 测速度本身有陷阱：噪声让 p = 1 变成抛硬币；挑赢家系统性偏高；cache 的冷热能差几倍；问题太小时测的是开销；基线的精度决定了”加速比”从哪里来；算术平均会被一个巨大的加速比带偏。
5. AlphaEvolve：LLM 集成提出 SEARCH/REPLACE 修改，评估器（可级联、可并行、可多指标）打分，进化数据库（MAP-Elites + 岛屿）保持多样性。在 Google 的计算栈上：Borg 调度回收 0.7% 的算力，Gemini 内核的分块启发式快 23%（训练快 1%），FlashAttention 的 IR 快 32%。成功的共同点：正确性容易保证、评估自动、产出是可读的代码、有训练 / 检验划分。
6. 迷你版的消融：进化 > 连贯的修改 > 其余部件；反馈要可操作、用法要保持多样性；评估只覆盖样本，所有分支都要静态检查。

9. 常见误区 · Common pitfalls

• “通过了 5 组随机测试就是正确的。” 同一种分布下多测几组几乎不增加信息；要换分布、随机化参数、专门构造边界输入（常数行、全零、极值）。忘了 eps 这类 bug，五种随机分布都抓不到。
• “fast_1 高就是模型强。” 先看评估器：有没有同步、缓冲区是否隔离、输入是否每次都新、基线精度是否可比、题目是否退化。每一次加固，报告的数字都在下降。
• “加速比 1.01× 说明变快了。” 在噪声以内。同一份代码和自己比也有一半的机会 > 1。
• “挑出来的最快候选有多快，部署就有多快。” 挑赢家有系统性偏差，要用独立的测量重新报告。
• “算术平均加速比 9.5×。” 很可能只是一道题 60×。用几何平均、中位数，给出分布。
• “降精度不算作弊。” 要看任务：如果要求 FP32 的精度，用 TF32 / BF16 换来的”加速”不是内核写得好；如果下游能接受低精度，就应该和同样精度的基线比。
• “AlphaEvolve 是一个更聪明的 LLM。” 它是一个系统：LLM 只是提议器，起决定作用的还有评估器、数据库、提示词里的反馈和上下文。消融说明每一部分都重要。
• “进化出来的代码通过了训练形状，就能放心用。” 没被训练形状走到的分支可能根本不能运行；要做全分支的静态检查，并在留出的形状上检验。
• “我们的迷你实验说明岛屿 / 级联没用。” 小问题、少数几个种子说明不了”没用”，只说明”在这里看不出差别”。

10. 练习 · Exercises

练习 1（fast_p） 10 道题的结果：3 道编译错误，2 道结果不对，其余 5 道正确，加速比分别是 0.8、1.02、1.5、2.5、12。算 fast_0、fast_1、fast_2、fast_10。

答案：fast_0 = 5/10 = 0.5；fast_1 = 4/10 = 0.4（1.02、1.5、2.5、12）；fast_2 = 2/10 = 0.2；fast_10 = 1/10 = 0.1。

练习 2（重复采样） 某道题一次采样”正确且 > 1”的概率是 0.1，采 10 次至少成功一次的概率？如果另一道题的概率是 0，采 100 次呢？

答案：$1 - 0.9^{10} \approx 0.651$；0 的题采多少次都是 0。这就是 fast_p@k 会饱和的原因（01 演示 3：k = 10 时 0.66，“先知”0.67）。

练习 3（diag 的账） diag(a) @ B，N = 2048，FP32。参考实现做多少次运算？语义上至少要读写多少字节？在 138 GB/s 的拷贝带宽下”光速”是多少？

答案：$2N^3 \approx 17.2$ GFLOP；读 a（N 个数）和 B（N² 个数），写 C（N² 个数）：$(N + 2N^2) \times 4 \approx 33.6$ MB；33.6 MB / 138 GB/s ≈ 0.24 ms（02 演示 3 打印的就是这个数）。候选实测约 0.5 ms，所以它离”光速”只差 2 倍左右，不会被光速检查误判。

练习 4（换算法） 证明 $\text{scale} \cdot \sum_j \frac{(x W^\top)_{ij}}{2} = x \cdot w$，其中 $w = \frac{\text{scale}}{2} \sum_j W_{j,:}$。B = 1024、K = N = 2048 时运算量少了多少倍？

答案：$\sum_j (xW^\top)_{ij} = \sum_j \sum_k x_{ik} W_{jk} = \sum_k x_{ik} \sum_j W_{jk}$，交换求和次序即可。原来 $2BKN$ 次，现在 $2BK$ 次（w 在 init 里预先算好，不计时），少 N = 2048 倍；实测只快 49×（加固评估器），因为现在受限于读 x 的带宽和调用开销。

练习 5（防线） 下面每种作弊，哪道防线能挡住？(a) 输出缓冲区里恰好是参考的结果；(b) 按输入的地址缓存结果；(c) 在另一条 CUDA stream 上计算；(d) 假设 LayerNorm 的 γ = 1、β = 0；(e) 题目的输出恒为常数。

答案：(a) 隔离、填 NaN 的输出缓冲区；(b) 每次计时换新的输入内容（或者核对计时时的输出）；(c) 计时前后同步 / 在计时区里强制求值，并核对计时时的输出；(d) 随机化参数；(e) 检查题目（输出是否随输入变化）+ 光速检查。

练习 6（挑赢家） 每次测量的”加速比”是 1 加上标准差 2% 的正态噪声，16 个一模一样的候选挑最大的，期望大约是多少？

答案：16 个独立标准正态的最大值的期望 ≈ 1.77，所以约 $1 + 1.77 \times 0.02 \approx 1.035$。03 演示 2 里每个只测 1 对时实测 1.043×，噪声比 2% 稍大。

练习 7（汇总） 4 道题的加速比是 50、1、1、0.5。算术平均和几何平均各是多少？

答案：算术平均 $(50 + 1 + 1 + 0.5)/4 = 13.1$；几何平均 $(50 \times 1 \times 1 \times 0.5)^{1/4} = 25^{1/4} \approx 2.24$。

练习 8（softmax 溢出） float32 的 exp 在输入超过多少时溢出？为什么不减最大值的 softmax 在标准正态输入上永远发现不了这个 bug？

答案：$\ln(3.4 \times 10^{38}) \approx 88.7$。标准正态的样本几乎不可能超过 6，exp(6) ≈ 403，离溢出很远；对数正态（值上百）、均值 1000 的输入才会触发。

练习 9（基线的精度） A100 的 FP32（CUDA core）稠密峰值 19.5 TFLOPS，TF32 Tensor Core 156 TFLOPS。参考实现不用 TF32、候选用 TF32 时，计算受限的矩阵乘法最多能有多大的”加速比”完全来自精度？

答案：156 / 19.5 = 8 倍。这就是 KernelBench-Verified 要求打开 TF32 的基线的原因。

练习 10（级联省多少） 训练形状 44 个，级联第一级用 6 个。如果 70% 的子程序在第一级就被淘汰，平均每个子程序要跑几次内核？省了多少？

答案：$6 + 0.3 \times 44 = 19.2$ 次，对比不级联的 44 次，省了约 2.3 倍。我们的实验里只有约 3% 的子程序在第一级被淘汰（34 / 1000），所以级联几乎没省——级联的价值取决于”坏子程序”有多常见、评估有多贵。

练习 11（山谷） 按 gemm_time 的寄存器模型（每线程 ≈ 每个 warp 的块元素数 / 32 + 40），4 个 warp 排成 2×2 时，256×128 的块每线程要多少个寄存器？为什么单字段变异很难从 128×128 / 4 warp 走到 256×128 / 8 warp？

答案：每个 warp 算 128×64 的块，$128 \cdot 64 / 32 + 40 = 296 > 255$，不能启动。要同时把 warps 从 4 改到 8（8 个 warp 排成 4×2，每个 warp 64×64，$64 \cdot 64/32 + 40 = 168$），单字段变异的中间状态要么不能启动、要么更慢，会被淘汰。

练习 12（从 23% 到 1%） 某类内核快了 23%（时间变成原来的 1/1.23），整个训练时间少了 1%。这类内核大约占训练时间的多少？

答案：省下的时间比例 $= f \times (1 - 1/1.23) = f \times 0.187 = 0.01$，$f \approx 5.3\%$。（只是一个按假设的估算：论文没有给出这个比例，而且”平均 23%“是对内核取的平均。）

11. 延伸阅读 · Further reading

阅读材料

• A. Ouyang, S. Guo, S. Arora, A. L. Zhang, W. Hu, C. Ré, A. Mirhoseini, KernelBench: Can LLMs Write Efficient GPU Kernels?, 2025. https://arxiv.org/abs/2502.10517
• A. Novikov et al., AlphaEvolve: A coding agent for scientific and algorithmic discovery, Google DeepMind 白皮书, 2025. https://arxiv.org/abs/2506.13131

评估与作弊

• Stanford Scaling Intelligence, KernelBench v0.1（博客）：放大题目、修复退化的题、多分布输入、roofline 检查、观察到的作弊。https://scalingintelligence.stanford.edu/blogs/kernelbenchv01
• METR, Measuring Automated Kernel Engineering, 2025-02. https://metr.org/blog/2025-02-14-measuring-automated-kernel-engineering
• Y. Zhang et al., KernelBench-Verified: Do LLM-Generated Kernels Actually Beat PyTorch?, 2026. https://arxiv.org/abs/2607.16241
• Sakana AI CUDA Engineer 的事件：TechCrunch 2025-02-21 的报道 Sakana walks back claims that its AI can dramatically speed up model training，以及 Sakana 在项目页上的更正说明。

系统与模型

• Cognition, Kevin-32B: Multi-Turn RL for Writing CUDA Kernels（博客），2025。
• NVIDIA 技术博客, Automating GPU Kernel Generation with DeepSeek-R1 and Inference Time Scaling, 2025-02。
• B. Romera-Paredes et al., Mathematical discoveries from program search with large language models（FunSearch）, Nature 2023。
• A. Fawzi et al., Discovering faster matrix multiplication algorithms with reinforcement learning（AlphaTensor）, Nature 2022。
• H. Massalin, Superoptimizer: A Look at the Smallest Program, ASPLOS 1987；E. Schkufza, R. Sharma, A. Aiken, Stochastic Superoptimization（STOKE）, ASPLOS 2013。

本仓库

• 第 10 讲：一个矩阵乘法内核的全部旋钮；Triton 风格的测试（边缘形状、容差、FP16 累加）。
• 第 13 讲：自动调优的搜索策略、学习的代价模型、调优的成本——本讲 §6 的”硬件”就是它的解析模型。
• CS349D 第 17 讲：怎样评测一个推理系统（公共随机数、开环与闭环、调好基线）——同样的方法论用在系统层面。

12. 下一讲预告 · Next lecture

这一讲里，“智能体”的能力主要来自一个循环：提出方案 → 调用工具（编译器、测试、计时器、profiler）→ 读反馈 → 再提出。KernelBench 的迭代优化、METR 的 KernelAgent、AlphaEvolve 都是这个循环的不同形态。第 15 讲把这个循环本身当作研究对象：工具的接口怎样设计（SWE-agent 的”智能体-计算机接口”）、反馈信号携带多少信息、上下文怎样管理；然后换到推理系统的视角——一个智能体在等编译器、等测试的时候，它的 KV cache 该留在显存里、换到 CPU、还是丢掉重算（InferCept、Continuum）。

→ 第 15 讲：智能体的工具使用