KernelBench 深度解读：LLM 能写出高效 GPU 内核吗

一句话总结：KernelBench 是目前最系统的”LLM 写 CUDA 内核”评测框架——250 个 PyTorch 任务、三级难度、硬件实测。结论很残酷：最强模型（DeepSeek R1）在单次生成下也只有不到 20% 的任务能跑赢 PyTorch，但迭代优化 + 反馈闭环能把 Level 2 的成功率从 36% 拉到 72%。

一、为什么需要 KernelBench

1.1 问题背景

高效的 GPU 内核是所有大模型推理和训练的性能基石。但写一个好的 CUDA kernel 需要：

• 深厚的硬件知识（内存层次、Warp 调度、Tensor Core 指令）
• 大量的调试时间（一个 kernel 调优可能需要数周）
• 对特定输入形状和计算模式的针对性优化

这和 LLM 的代码生成能力形成了一个天然的交叉点：能不能让 LLM 自动写 CUDA kernel？

1.2 现有评测的不足

在 KernelBench 之前，缺乏一个系统化的评测框架来衡量 LLM 的 kernel 编写能力。现有的代码生成基准（HumanEval、MBPP 等）都聚焦于通用编程，不涉及：

• 硬件感知的性能评估（不只是”能跑”，还要”跑得快”）
• 真实的 ML 工作负载（不是玩具题）
• 端到端的模型级别优化（不只是单个算子）

1.3 KernelBench 的定位

KernelBench 不是一个静态的 leaderboard，而是一个开源框架 + 环境，核心理念是：

• 模拟真实的 kernel 工程场景
• 基准上的进展可以直接转化为实际应用中更快的推理速度
• 支持迭代优化——LLM 可以利用执行反馈和性能分析工具（如 Nsight Compute）来改进生成的 kernel

---

二、数据集设计

KernelBench 包含 250 个 PyTorch ML 工作负载，分为四个难度级别：

2.1 Level 1 🧱 — 单内核算子（100 题）

基础构建块：卷积、矩阵乘法、层归一化、Softmax、各种激活函数等。

任务形式：给一个 PyTorch 算子（如 torch.nn.Conv2d），让 LLM 生成等价的 CUDA 实现。

考察能力：基本的 CUDA 编程、内存访问模式、线程块设计。

2.2 Level 2 🔗 — 简单融合模式（100 题）

多个算子的组合：Conv + Bias + ReLU、Matmul + Scale + Sigmoid、LayerNorm + Dropout 等。

核心挑战：算子融合（Kernel Fusion）——将多个操作合并为一个 kernel，消除中间结果的内存读写开销。

为什么重要：算子融合是 GPU 优化中最核心的技术之一，torch.compile 的主要优化手段就是这个。

2.3 Level 3 ⚛️ — 完整模型架构（50 题）

端到端优化整个模型：MobileNet、VGG、MiniGPT、Mamba 等。

核心挑战：需要理解整个模型的计算图，做全局性的优化决策，而不只是局部的算子替换。

2.4 Level 4 🤗 — HuggingFace 模型（20 题，附加）

直接优化 HuggingFace 上的真实模型，要求模型不仅会写 kernel，还能在软件库级别做集成修改。目前不提供标准评估。

设计亮点

设计决策	原因
只关注前向传播	简化评测复杂度，前向传播是推理的核心
固定输入形状	允许针对特定形状做极致优化（这也是实际工程中的常见做法）
PyTorch Eager 为基线	而非 torch.compile，因为 Eager 是最常见的开发起点
宽松的数值容限（1e-02）	允许精度变化和算法替代，不过度苛求 bit-exact

---

三、评估体系

3.1 核心指标：fast_p

KernelBench 引入了一个优雅的复合指标 fast_p，同时捕捉正确性和性能：

$\text{fast}_p = \frac{|\{\text{内核} : \text{功能正确} \wedge \text{加速比} > p\}|}{|\text{全部任务}|}$

• fast₀：只要求功能正确（加速比 > 0），等同于正确率
• fast₁：功能正确 且 比 PyTorch Eager 快（加速比 > 1）
• fast₂：功能正确 且 比 PyTorch Eager 快 2 倍以上

为什么 fast_p 好：传统的代码生成评估只看”对不对”，但 kernel 领域”对”和”快”是两码事——一个正确但慢 10 倍的 kernel 毫无意义。fast_p 通过调节 p 来考察不同级别的优化能力。

3.2 评估流程

输入: PyTorch 参考实现
     ↓
LLM 生成 CUDA kernel
     ↓
Step 1: 编译检查（语法是否有效）
     ↓
Step 2: 正确性验证
  · 5 组随机输入
  · 与 PyTorch 输出对比
  · 绝对误差 + 相对误差 ≤ 1e-02
     ↓
Step 3: 性能测试（仅对正确的 kernel）
  · 多次运行取均值
  · 计算加速比 = t_baseline / t_generated

3.3 测试配置

• 硬件：NVIDIA L40S GPU（论文主实验）；也测了 A100、H100
• 采样策略：
  • Greedy decoding（temperature=0，单次生成）
  • Pass@k（k 次采样中至少一次正确的概率）
  • 迭代优化（固定预算内利用反馈迭代改进）

---

四、实验结果

4.1 单次生成的 fast₁ 表现

这是最核心的结果——各模型在”一次机会”下能生成多少个比 PyTorch 快的正确 kernel：

模型	Level 1	Level 2	Level 3
DeepSeek R1	12%	36%	2%
OpenAI o1	10%	24%	12%
Claude 3.5 Sonnet	10%	7%	2%
DeepSeek V3	6%	4%	8%
GPT-4o	4%	5%	0%
Llama 3.1-405B	3%	0%	2%
Llama 3.1-70B	3%	0%	0%

几个关键发现：

1. 推理模型碾压非推理模型：o1 和 DeepSeek R1 远超 GPT-4o 和 DeepSeek V3，说明”深度思考”对 kernel 编写至关重要
2. Level 2 是 R1 的甜区：36% 的 fast₁ 远超其他所有模型在任何级别的表现，R1 对算子融合有独特的理解
3. Level 3 整体崩塌：端到端模型优化对所有模型都极其困难，最好也只有 12%（o1）
4. 开源模型大幅落后：Llama 系列几乎全军覆没，参数量增加（70B→405B）帮助不大

4.2 迭代优化的效果

固定 10 次调用预算，对比”重复采样”和”迭代优化（带反馈）“：

方法	Level 1	Level 2	Level 3
DeepSeek R1 单次	12%	36%	2%
DeepSeek R1 迭代+反馈	43%	72%	18%
DeepSeek V3 单次	6%	4%	8%
DeepSeek V3 迭代+反馈	19%	6%	14%

反馈闭环的威力：

• DeepSeek R1 在 Level 2 上从 36% → 72%，翻倍
• Level 1 从 12% → 43%，接近 4 倍提升
• 关键在于提供了两类反馈：
  • 执行反馈（E）：编译错误、运行时错误的具体信息
  • 性能分析反馈（P）：Nsight Compute 的 profiling 数据

这也是 KernelBench 最重要的启示：LLM 写 kernel 不应该是一次性的，而应该是 Agent 式的迭代循环。

4.3 错误模式分析

模型生成的 kernel 失败原因分为两大类：

执行失败（编译错误、CUDA 内存违规、运行时崩溃）：

• 推理模型（o1、R1）产生的执行失败 < 55%
• 非推理模型更容易犯语法级别的错误

功能正确性错误（能跑但结果错）：

• 所有模型的最大痛点
• 代码能编译、能运行、不报错，但输出的张量值是错的
• 这类错误很难通过简单的编译反馈来修复

4.4 性能异常值：LLM 的”灵光一现”

虽然整体表现不佳，但存在一些令人惊艳的异常值：

任务	模型	加速比	优化手段
对角矩阵乘法	多个模型	13x+	识别出无需构造完整对角矩阵的算法
Matmul+Div+Sum+Scale	Claude 3.5 Sonnet	3x	四个操作融合为一个 kernel
GELU 激活	DeepSeek R1	2.9x	融合 + 共享内存优化
Cosine Similarity Loss	-	2.8x	利用共享内存减少全局内存访问

启示：LLM 有能力做高层次的算法优化（比如识别对角矩阵乘法的特殊结构），这是传统编译器难以自动完成的。但 LLM 很少成功使用 Tensor Core 指令（wmma），这恰恰是 GPU 性能的关键。

4.5 硬件迁移性

同一个 kernel 在不同 GPU 上的表现波动很大：

• 在 L40S 上优化好的 kernel，换到 A100 或 H100 上性能可能大幅退化
• 当前模型不具备自动适配不同硬件架构的能力
• 这意味着”写一次、到处快”的目标还很远

---

五、生态与后续工作

KernelBench 发布后迅速催生了一系列后续工作：

5.1 MultiKernelBench（2025.07，南大 & 浙大）

解决的问题：KernelBench 只支持 NVIDIA GPU + CUDA，但实际硬件生态远比这丰富。

扩展：

• 285 个任务，覆盖 NVIDIA（CUDA）、AMD（HIP/ROCm）、华为昇腾（AscendC）、Intel（SYCL）、Google TPU（Pallas）
• 支持 Triton 和 TileLang 等 DSL
• 提出了 Category-Aware Prompting 策略

5.2 KernelSkill（2026.03，多Agent框架）

核心创新：用”专家优化技能”替代 LLM 的隐式启发式规则

• 多 Agent 协作 + 双层记忆（长期：可复用技能；短期：防重复回溯）
• 在 KernelBench Level 1-3 上达到 100% 成功率
• 平均加速：Level 1 5.44x、Level 2 2.82x、Level 3 1.92x
• 相比论文原始结果（R1 最好也只有 72% fast₁ at Level 2），这是一个巨大的飞跃

5.3 生态工具链

KernelBench 本身也在持续演进：

• 多后端支持：除了 CUDA，还支持 Triton、CuTe、TileLang、ThunderKittens
• AMD GPU 支持：通过 HIP 后端支持 ROCm ≥ 7.1
• HuggingFace 数据集托管：ScalingIntelligence/KernelBench v0.1

---

六、关键洞察与启示

6.1 对模型能力的启示

1. 推理能力是 kernel 编写的关键：o1/R1 >> GPT-4o/V3，说明写 kernel 需要深度推理而非模式匹配
2. 基座模型仍然重要：即使加了推理时计算，Llama 系列仍远不如 DeepSeek/OpenAI，说明预训练数据中的 CUDA 知识密度很关键
3. 正确性是最大瓶颈：不是写不出 kernel，是写不对。功能正确性错误比编译错误更难修复

6.2 对工程实践的启示

1. Agent 式迭代优于一次性生成：反馈闭环能带来 2-4 倍的提升，这说明 kernel 生成应该被设计为 Agent 工作流而非简单的 prompt → code
2. 算子融合是低垂果实：Level 2 的表现远好于 Level 1 和 3，说明 LLM 对”把多个操作合并”这个模式理解得最好
3. 硬件感知优化仍是盲区：Tensor Core、内存层次、Warp Scheduling 等硬件特定优化，LLM 几乎不会用

6.3 对我们数据团队的启示

KernelBench 的结论跟我们在数据生产中的观察高度一致：

• 反馈闭环 > 一次性生成：数据质检也应该是 Agent 迭代式的
• 推理能力决定上限：复杂的数据质量判断需要深度推理模型
• 领域知识密度决定基线：模型在 CUDA 上的表现取决于预训练数据中 CUDA 代码的质量和数量——类比到数据工作，模型对数据标准的理解取决于我们给它的上下文质量

---

七、技术细节附录

7.1 项目信息

项目	详情
论文	arXiv:2502.10517
会议	ICML 2025
团队	Stanford Scaling Intelligence Lab
代码	github.com/ScalingIntelligence/KernelBench
数据集	HuggingFace: ScalingIntelligence/KernelBench
许可	Apache 2.0
硬件	NVIDIA L40S（主实验），A100/H100（迁移性测试）

7.2 运行示例

# 安装
pip install kernelbench

# 单题评测
uv run python scripts/generate_and_eval_single_sample.py \
  dataset_src=huggingface level=2 problem_id=40 \
  server_type=google model_name=gemini/gemini-2.5-flash

# 批量生成
python scripts/generate_samples.py --level 1 --model deepseek-r1

# 评估
python scripts/eval_from_generations.py --level 1 --run_name my_run

7.3 参考文献

1. Ouyang et al., “KernelBench: Can LLMs Write Efficient GPU Kernels?”, ICML 2025
2. Wen et al., “MultiKernelBench: A Multi-Platform Benchmark for Kernel Generation”, arXiv:2507.17773
3. Sun et al., “KernelSkill: A Multi-Agent Framework for GPU Kernel Optimization”, arXiv:2603.10085
4. Stanford Scaling Intelligence Lab Blog: KernelBench 官方博客