代码仓:https://github.com/ict-agent/klineage

1 概述

KLINEAGE 不直接让大模型学习专家 Kernel 的最终代码,而是通过逆向去优化,恢复专家 Kernel 中隐含的优化步骤、适用条件和依赖关系,再将这些知识组织成可以复用的 Skill。

![image-20261014100912674](./论文分享:KLINEAGE RECOVERING THE MISSING When OF KERNEL OPTIMIZATION BY DEOPTIMIZING EXPERTS/image-20261014100912674.png)

LLM 往往知道应该采用哪些优化技术(What),却不知道这些优化应该在什么条件下应用(When)。

本篇论文解决的问题是:把专家代码中隐含的优化决策知识转化成可验证、可复用的条件化经验。

1.1 工作定位

维度 论文定位
输入 专家 Kernel、编译与测试环境、目标硬件和性能测量工具
离线产物 去优化谱系、具体正向变换、经过 roundtrip 准入的条件技能库
在线输入 目标算子、代码语言、硬件平台、当前 Kernel 和已实施优化
在线输出 在预算内找到的最佳已验证 Kernel
LLM 的角色 提出去优化、改写代码、抽象技能、在目标代码上实现技能并修复错误
核心贡献所在 专家经验的提取方式与条件化表示;论文明确不将其定位为新的搜索策略

实验覆盖 H100、RTX PRO 6000 Blackwell 和 Ascend 910B。每个平台分别从 GEMM、Conv2d、FMHA、GDN、Top-K 五类专家算子建立技能库;随后测试源任务重建、新算子复用和更广泛的 100 个融合任务。

1.2 评价

工作的价值在于把专家代码从**“可供模仿的最终答案”变成“可用于学习优化依赖和适用条件的证据”**。最有说服力的实验是:在其设置下,条件技能比直接提供同一专家源码更有效,并且能帮助未参与建库的算子。

但需要保留两条边界:第一,源任务上的表现包含专家知识提取优势,不能单独证明泛化;第二,三个平台分别建库,不能据此认定同一组 NVIDIA 技能已经自动迁移到 Ascend。

2 核心思想

![fig1-motivation](./论文分享:KLINEAGE RECOVERING THE MISSING When OF KERNEL OPTIMIZATION BY DEOPTIMIZING EXPERTS/fig1-motivation-1791988419700-3.png)

传统正向探索从简单 Kernel 出发,不断尝试优化。困难在于,模型可能无法在预算内发现高性能实现所需的复杂组合。

KLineage 从已知高性能终点出发:移除流水线、撤销向量化、将硬件 intrinsic 标量化等。每次退化都通过执行反馈检查,保留可以运行且满足验证条件的中间状态。随后,把这些中间状态组织成从简单实现走向专家实现的课程。

图 1 的三个面板表达了完整逻辑:左侧是正向搜索可能遇到的失败与局部最优;中间是专家 Kernel 的反向拆解及正向重建;右侧是条件技能的检索、实现和验证。这里的“谱系”是系统构造出的可执行路径,不是专家作者真实的历史开发过程。

2.1 把 What、When 和 How 分开

层次 含义 例子
What:intent 优化意图 用共享内存暂存可重复使用的数据块
When:precondition + scope 适用前提与验证范围 存在块内复用;容量可容纳;同步关系合法;已验证的平台与状态匹配
How:anchor + carrier 修改位置与实现提示 在加载/计算循环处插入暂存、同步及共享内存访问;提供伪代码或 diff 草图
effect / risk / evidence 预期作用、风险、实证 降低重复全局访存;可能增加占用或同步开销;记录实测与失败证据

3 创新点

3.1 通过去优化专家实现获取经验

现有前向经验记忆受到模型探索能力限制:模型没找到的优化结构,很难进入经验库。KLineage 直接从成功专家实现出发,通过验证门控的简化来产生中间状态,降低对正向发现能力的依赖。

3.2 把优化适用性作为一等信息

SkillCard 同时保存 intent、anchor、carrier、precondition、effect、evidence、risk、scope 和 verification log。与单纯的提示词、失败规则或最终代码片段相比,这一表示同时回答:

  1. 改哪里、实现什么?
  2. 当前状态是否满足前提?
  3. 预期改善什么、违反前提有什么风险?
  4. 哪些试验支持这项技能,在什么范围内验证过?

创新不在于单独增加一个字段,而在于用具体谱系和执行证据支撑这些字段。

3.3 通过 roundtrip 检验经验是否可被重新实现

LLM 对代码进行的总结可能听起来合理,却无法指导另一个实现过程。KLineage 要求候选技能至少在一次 held-out roundtrip 中被成功实现并复现预期效果,才进入可复用库。

这一机制测试的是“技能描述能否驱动重新实现”,比只检查原始专家代码正确更接近经验复用的目标。但一次成功依然只提供有限范围内的经验性支持。

3.4 依照当前状态检索和组合技能

检索考虑算子、语言、平台和已有优化。目标代码不必从 naive 实现开始,也可以是已有部分优化的版本。模型据此补全缺失优化或修复不恰当实现。

创新边界: 论文提供的是经验证的候选技能池;选择和落地仍由下游 Agent 执行。它没有证明全局最优搜索、完备的依赖规划或形式化的合法性保证。

3.5 与“能力类型化优化经验 IR”选题的关系

结合当前工作区的选题总结,KLineage 是一个需要直接对照的相关工作:它已经覆盖结构化经验、代码/IR 锚点、适用条件、风险和验证范围。仅以“用 IR 保存优化经验”或“增加平台标签”为创新,很难与之拉开距离。

可进一步研究的差异应落实到机制:例如把硬件能力要求变成可检查约束,把经验表示为具有明确语义的参数化变换,提供跨后端重绑定,以及在不满足条件时给出可解释拒绝。以上属于对用户选题的启示,不是本文已实现的机制。

4 方法

4.1 总体架构

![fig2-framework](./论文分享:KLINEAGE RECOVERING THE MISSING When OF KERNEL OPTIMIZATION BY DEOPTIMIZING EXPERTS/fig2-framework.png)

图 2 将系统分为三个部分:左侧离线提取专家经验,中间按目标上下文组织和检索 SkillCard,右侧由 Agent 进行代码实现与验证。底层支持不同 Agent harness 和硬件工具;正文提及 NVIDIA 的 ncu 和 Ascend 的 msprof。(§3.5,第 6 页)

总结如下:

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
flowchart TD
A[专家 Kernel] --> B[LLM 提出去优化并改写代码]
B --> C{编译与正确性检查<br/>并记录性能}
C -->|失败| D[保留失败证据<br/>修复或尝试其他简化]
D --> B
C -->|通过| E[保存简化状态与步骤]
E -->|继续拆解| B
E --> F[重建正向优化并执行验证]
F --> G[聚合具体变换<br/>抽象候选 SkillCard]
G --> H{held-out roundtrip<br/>复现预期效果}
H -->|成功| I[记录验证范围<br/>技能入库]
H -->|未成功| J[不准入可复用库]
I --> K[按算子、语言、平台<br/>已有优化检索]
T[新目标及当前 Kernel] --> K
K --> L[LLM 选择技能并实现]
L --> M{编译、正确性、性能验证}
M -->|通过且更优| N[更新最佳已验证 Kernel]
M -->|失败或不优| O[反馈与修复或继续搜索]
O --> L

4.2 引导式去优化

从 Kn=K∗K_n=K^* 开始,反向步骤为:

bi:Ki→Ki−1b_i:K_i\rightarrow K_{i-1}

每次移除或削弱一个优化,记录修改位置、动作类型和执行证据。拒绝的简化会成为后续 precondition/risk 的证据。动作词表是开放的,动作间顺序被视为软先验;违背预设顺序但验证通过的步骤仍可保留。

4.3 正向步骤重建

系统把接受的反向步骤恢复成:

τi:Ki−1→Ki\tau_i:K_{i-1}\rightarrow K_i

这里的逆向关系是语义上的。LLM 根据去优化动作与位置重新生成正向编辑,而不是机械地把 diff 倒过来;重新生成的结果再次通过编译、正确性与性能验证。最终构成:

L=(K0,τ1,K1,…,τn,Kn=K∗)L=(K_0,\tau_1,K_1,\ldots,\tau_n,K_n=K^*)

4.4 从具体变换抽象出技能

具有共同动作类别及相近修改位置/效果特征的变换被聚合。LLM 接收正向 diff、局部代码上下文、修改位置和性能变化,形成候选技能:

σ^=(ι,anchor,carrier,pre,effect,evidence,risk,scope,ver)\hat\sigma=(\iota,anchor,carrier,pre,effect,evidence,risk,scope,ver)

字段 含义
intent / ι\iota 抽象的优化意图或变换家族
anchor 优化所作用的代码/IR 构造与位置
carrier 可指导实现的 diff、伪代码、代码片段或注释
pre 适用的代码、布局、硬件等前提
effect 预期作用与性能效果
evidence 原始变换与验证证据
risk 前提被破坏时的风险及失败信息
scope 算子、语言、平台和先前优化的适用范围
ver roundtrip 等验证记录

一个 intent 可以在不同语言或平台拥有不同 carrier。

5 实验结果

5.1 实验配置与评分口径

项目 设置
平台 H100 PCIe / SM90;RTX PRO 6000 Blackwell / SM120;Ascend 910B
在线目标语言 NVIDIA:CUDA C;Ascend:AscendC
专家源算子 GEMM、Conv2d、FMHA、GDN、Top-K,每个平台分别建库
新算子 KDA、Sparse Attention、Top-P、Fused Add-RMSNorm、GQA
源算子建库与生成 Claude Code + Claude Opus 4.6;在线每算子 10 美元
迁移与广泛基准 Codex + DeepSeek-V4.1-Flash;新算子每项 2 小时,基准每题 30 分钟
广泛基准 NVIDIA:KernelBench Level 2,100 题;Ascend:相同任务的 MultiKernelBench Fusion 移植版本,100 题
加速比 参考延迟 / 生成 Kernel 延迟,越大越好

附录 D 的源算子基线比较说明:共用在线模型、预算、验证器和参考实现;正确性检查使用三个随机种子;triton.testing.do_bench 的 warmup=25、rep=100,单位为毫秒;每工作负载运行 5 次,报告 success-only geometric mean。这套重复与聚合口径不应未经说明直接套用到所有后续实验。

5.2 技能库规模与离线成本

平台 候选数 准入技能数 准入率 离线建库成本
H100 72 52 72.2% $131
RTX PRO 6000 45 36 80.0% $92
Ascend 910B 109 67 61.5% $197

5.3 Kernel 质量

![fig3-results](./论文分享:KLINEAGE RECOVERING THE MISSING When OF KERNEL OPTIMIZATION BY DEOPTIMIZING EXPERTS/fig3-results.png)

算子 H100:KLineage 相对专家参考 RTX PRO 6000:KLineage 相对专家参考
GEMM 0.99× 1.00×
Conv2d 1.73× 1.08×
FMHA 0.94× 0.99×
GDN 1.15× 1.08×
Top-K 1.12× 1.29×

论文报告,十个工作负载—平台组合全部正确,整体达到专家参考的 1.12×;AdaExplore 和 AccelOpt 的汇总结果分别为 0.25× 和 0.54×,在四个和两个组合上未找到正确 Kernel。

KLineage 接近或超过专家实现,并明显改善基线的失败问题,但并非每个算子都超过专家。H100 FMHA 仍慢约 6%,GEMM 基本持平。由于失败结果和 success-only 聚合的存在,汇总性能应与失败数量一起阅读,不能只凭三个总体数字判断全任务期望收益。

6 局限性

6.1 依赖高质量专家代码,知识覆盖受源任务限制

每个平台只从五类专家算子出发,主要涉及矩阵、注意力、规约及选择类模式。新算子和 100 题结果支持一定外推,但对复杂稀疏访存、动态形状、非规则 HPC 计算、通信 Kernel 与多节点应用的覆盖没有直接证据。

建议按优化机制及工作负载属性评估覆盖,而不只按算子名称统计;增加未见布局、规模和机制的分层测试。

6.2 分平台建库尚未验证跨架构经验迁移

H100、Blackwell 与 Ascend 各自建库。平台标签与多后端框架解决了经验组织问题,却没有提供硬件能力类型、跨后端参数约束映射或自动重绑定的实验证据。

若以跨架构迁移为后续研究,应将“目标平台独立建库”“直接迁移源卡”“能力约束重绑定”作为不同实验设置,并报告目标端调优成本和负迁移率。

6.3 前提和风险主要是经验性描述

LLM 根据有限代码状态及试验抽象条件,技能只需至少一次 roundtrip 成功即可准入。这不能保证所有满足自然语言条件的新输入都合法或更快。编译与随机测试也不能完整覆盖尾部、对齐、极值、数值容差和同步行为。

可补充多形状、多布局、不同数据分布与边界输入验证,将可检查的条件落实为约束,并明确记录数值契约与验证覆盖范围。

6.4 核心部件缺少充分消融

源码提示对比支持完整条件技能系统的价值,但正文及附录未给出充分的组件消融来隔离 precondition、risk、scope、anchor/carrier、谱系课程及 roundtrip 的作用。

建议在相同模型、相同预算和相近提示规模下,比较完整卡片、仅 intent、移除前提、移除验证、移除谱系顺序等设置,分别统计成功率、性能、验证次数和失败原因。