September 3, 2026 中 #LLM Agent #Reasoning #Operating Systems

HIERA 阅读笔记：让 LLM Agent 先学会「在哪一层优化」，再谈「怎么优化」——GPU 算子实现空间的分层规划

笔记日期： 2026-09-03 作者： Zhongzhu Zhou 阅读论文： HIERA: Workload-Aware Planning Across Implementation Spaces for GPU Kernel Optimization 论文作者： Jinghao Wang, Qiqi Gu, Chenpeng Wu, Jianguo Yao, Haibing Guan, Xijun Li（上海交通大学） arXiv： 2608.21157 (cs.DC)，提交于 2026-08-21 投稿/状态： 论文版权声明显示投给 AAAI 2027

1. 为什么是这篇，为什么是今天

今天是周四，按照我自己排的五方向轮换表，属于「ML Systems」方向：分布式训练、调度、编译优化、内存管理。往常这个方向我更常挑分离式架构、流水线调度、KV cache 放置这类偏”服务系统”的论文，HIERA 略有不同——它正好卡在”ML 系统”和”LLM Agent 写代码”的交叉点上，处理的是一个正在悄悄变成基础设施的问题：LLM Agent 能不能自动写出又快又对的 GPU kernel？如果能，它的搜索预算应该怎么分配才不浪费？

如果你手写过 CUDA kernel，或者看过自动调优器（Ansor、TVM）在巨大的搜索空间里死磕，你大概已经知道结论了：kernel 优化最难的部分往往不是”怎么让它更快”，而是”我到底应该在哪个抽象层次上干活”。 一个融合的逐元素+归约算子，可能用三行对 torch.compile 友好的 PyTorch 代码就是最快的；一个数据复用模式很特殊的 stencil 计算，可能只有手写共享内存 tiling 的原生 CUDA 才能跑到性能上限。HIERA 的核心洞察是：把”该在哪个层次干活”这个选择，做成 Agent 循环里一个显式的、可学习的、按任务动态决定的步骤，而不是工具里事先写死的固定承诺。

这篇笔记假设读者大致知道”GPU kernel”和”LLM Agent”是什么，但未必接触过自动调优框架（TVM/Ansor）、Agentic CUDA 优化流水线（CUDAForge、KernelBench-Caesar）、或者用强化学习训练的 kernel 生成器（CUDA-L1、Kevin）。第 2 节会先把这些背景铺垫好，再进入 HIERA 本身。

2. 前置知识

2.1 “GPU kernel” 是什么，为什么写快很难

GPU kernel 是一段在 GPU 的成千上万个并行线程上跑的小程序。当你在 PyTorch 里调用 torch.matmul(a, b) 时，底层会分发到 cuBLAS 这类经过深度调优的厂商库 kernel。当你写一个自定义算子——比如”乘法 + 加偏置 + GELU”融合起来——PyTorch 默认的 eager 执行模式可能会跑三个独立的 kernel，每两个 kernel 之间都要把中间结果写回、再读出显存。因为 GPU 的算力往往远超其显存带宽（这就是所谓”内存墙”，同一周出现的硬件综述 arXiv:2608.28048 对此有更系统的讨论），把多个算子融合成一个自定义 kernel、避免不必要的显存往返，往往能带来巨大加速——而且往往不需要减少任何算术运算量，收益纯粹来自”不写/不读中间结果”。

大致来说，实现一段 GPU 计算有三个”海拔高度”，这正是 HIERA 要显式做决策的那个轴：

高层框架算子（PyTorch 的 torch.nn.functional.*、torch.compile）：你用已有的构建块拼出计算。开发成本低，跨 GPU 型号可移植，但你继承了框架帮你做的一切融合/调度决策，而且没法表达框架本身不擅长处理的计算模式。
优化过的厂商/社区库（cuBLAS、cuDNN、CUTLASS、Triton 编译出的 kernel）：你依然调用一个预先写好、由专家精心调过的例程，但这个例程本身是专门针对 GEMM、卷积等一类”形状”精心打磨到接近峰值吞吐的。相比第一层，这是更窄、更专门化的计算类型集合。
自定义、手写（或 LLM 生成）的 CUDA kernel：你自己写原生 kernel——控制线程/block 布局、共享内存 tiling、warp 级原语、显存合并访问、Tensor Core mma 指令。灵活性和潜在性能都最高，但代码写错、编译不过、数值不对、或者比库基线还慢的风险也最高。

关键点在于：没有哪一层能在所有工作负载上都占优。 一个简单的逐元素激活函数，往往用纯 PyTorch 表达就是最快的（torch.compile 的框架级融合就够用了，手写 CUDA 不会带来收益反而带来真实风险）。一个大规模 GEMM，几乎总是 cuBLAS/CUTLASS 最快——从零重新实现矩阵乘法很难打得过几十年积累的厂商调优。但一个数据复用模式特殊的负载（比如复用结构不寻常的 stencil，或者自定义稀疏注意力模式）可能只有在手工调优的 CUDA 里才能触到真正的性能天花板，因为没有任何库原语能匹配它的确切访问模式。

2.2 KernelBench：这条研究线的标准评测基准

KernelBench（Ouyang 等，2025，ICML）是评测”LLM 能不能写出快的 GPU kernel”的标准评测框架。它包含 250 个任务，分成三个难度递增的层级：

Level 1（100 个任务）： 单个 PyTorch 算子（例如单独一个矩阵乘、单独一个 softmax）。
Level 2（100 个任务）： 融合的多算子模式（例如”卷积 → BatchNorm → ReLU”作为一个任务）。
Level 3（50 个任务）： 完整的模型级架构（例如一个小型 Transformer block 或 ResNet 阶段）。

对每个任务，Agent 拿到一个 PyTorch 参考实现，需要产出一个替代实现（可以自由混合自定义 CUDA / 库调用 / PyTorch），这个实现必须：(a) 功能正确——输出必须在数值容差范围内匹配参考实现；(b) 更快——在同一块 GPU 上跑基准测试时要快于参考实现。KernelBench 定义了三个几乎所有后续论文（包括 HIERA）都会报告的核心指标：

$\text{fast}_0$ ：至少产出一个有效（能编译、正确）实现的任务比例——衡量的是可靠性，不是速度。
$\text{fast}_1$ ：最优有效实现比参考实现更快（加速比 $>1\times$ ）的任务比例——这是”到底有没有真正带来收益”的核心指标。
$\text{fast}_2$ ：最优有效实现达到超过 $2\times$ 加速的任务比例——对”确实令人印象深刻的加速”设置了更严的门槛。

这三个指标为什么必须同时看，为什么 $\text{fast}_0$ 不是可以随便忽略的边角指标？原因是：更激进地探索原生 CUDA 之类空间的系统，往往会产出更多无效候选——LLM 生成 CUDA 代码一个非常常见的失败模式是：代码能编译，但悄悄产出了错误的数值结果，或者干脆运行时崩溃。一个系统完全可能在 $\text{fast}_2$ 上看起来很唬人，但同时在 40% 的任务上彻底失败（ $\text{fast}_0$ 很低）；HIERA 的核心论点恰恰是：“一律搜索原生 CUDA”这种朴素策略正好会掉进这个陷阱。

2.3 LLM 驱动 kernel 优化的三条脉络

要给 HIERA 定位，需要先了解它对比、并相对其定位的三个前人系统家族：

KernelBench-Caesar —— KernelBench 官方自带的迭代基线：生成一个候选、跑一下、把错误/性能反馈回去、循环。除了”生成模型想生成的任何东西”之外，不对实现层次做任何固定的事先承诺，纯粹靠执行反馈来引导。
CUDAForge（Zhang 等，2025）—— 一个 Agentic 的”Coder–Judge”框架，只在自定义 CUDA 实现范围内做迭代精炼，用编译器诊断信息、GPU 硬件规格、Nsight Compute（NCU）性能剖析指标作为反馈。这是一个很强的系统，但设计上它对每个任务都固定承诺”原生 CUDA”这一实现空间，不管这是不是该任务合适的海拔高度。
CUDA-L1（Li 等，2026）和 Kevin（Baronio 等，2026）—— 分别用强化学习（对比式 RL 和多轮 RL）来训练模型权重，使其偏好正确、快速的 CUDA 代码。这需要一次 RL 训练过程——成本高，而且得到的偏好被固化在模型参数里，推理时不可检查、不可调整。

HIERA 的定位是：训练无关（无需 RL、无需微调——用提示词/编排即可搭配任意现成基座 LLM），并且把实现空间本身的选择当作一个显式的规划决策，而不是像 CUDAForge 那样固定死，也不是像 KernelBench-Caesar 那样完全不加约束。

2.4 Nsight Compute（NCU）性能剖析，简述

NVIDIA Nsight Compute 是一个 kernel 级性能剖析工具，能报告硬件计数器：达成的占用率（occupancy）、相对峰值的显存吞吐、warp 停滞原因、指令构成、Tensor Core 利用率等。HIERA 的反馈回路用的是 NCU 的输出（而不只是墙钟时延），这样优化 Agent 能推理”为什么这个候选慢”——比如”显存带宽受限，只跑到峰值的 40%“vs”计算受限但占用率低”——而不是只看到一个笼统的时间数字。

3. 问题形式化

HIERA 用形式化语言给出了它的优化目标，值得仔细走一遍，因为这套记号会贯穿整个方法部分。

给定一个 kernel 优化任务 $\tau$ 、参考实现 $f_\tau$ 、目标硬件平台 $H$ 、有限搜索预算 $B$ （即 Agent 总共最多可以生成 $B$ 个候选实现），定义 $C_B(\tau)$ 为该预算内已探索的候选实现集合，令

\text{Valid}(x; \tau) \in \{0, 1\}

表示某个候选 $x$ 是否成功编译并且满足 $\tau$ 的接口/语义要求（即在任务的测试框架下，输出在数值容差范围内与参考实现一致）。

对任意有效候选，其相对参考实现的加速比定义为两个测得时延的比值：

s(x; \tau, H) = \frac{t(f_\tau; H)}{t(x; \tau, H)} \tag{1}

其中 $t(\cdot; H)$ 表示在硬件 $H$ 上测得的墙钟执行时延。直观理解：如果参考实现耗时 10 毫秒、你的候选耗时 5 毫秒，则 $s=2$ ，即 $2\times$ 加速。如果你的候选更慢（比如耗时 20 毫秒），则 $s=0.5<1$ ——“加速比”低于 1 实际表示变慢了，这也是为什么 KernelBench 的 $\text{fast}_1$ 指标要专门问” $s>1$ 是否成立”。

于是，整体优化目标是：在预算内找到最快的有效候选：

x_\tau^{*} = \arg\max_{x \in C_B(\tau)} s(x; \tau, H) \quad \text{s.t.} \quad \text{Valid}(x; \tau) = 1 \tag{2}

为什么这个形式化很重要，以及它刻意留白了什么： 注意公式 (2) 是一个带约束的 argmax——无效候选是被直接排除在考虑范围之外的，而不是通过某种软性损失项被”惩罚”。这是一个有实际后果的建模选择：它意味着一个找到一个”有效但普通”（比如 $s=1.05\times$ ）候选的搜索策略，会打败一个找到九个无效候选、外加一个”理论上会非常快但从没编译通过”候选的搜索策略。这种”有效但普通胜过无效但理论上很棒”的不对称性，正是 HIERA 作者论证”实现空间的选择”（它直接决定了产出有效候选的概率）应当被当作搜索的一等公民、而不是”生成任何东西看看能不能凑合”背后附带品的原因。

还有一点值得注意： $C_B(\tau)$ （已探索的候选集合）通常远小于所有可能实现构成的完整空间——对于 Level 3（完整小模型架构）任务，“用 CUDA/库/PyTorch 各种组合实现这个”的空间在组合意义上是巨大的，而 $B$ （HIERA 实验中通常是 18）相比之下微不足道。这正是本文论点的核心：当你的探索预算相对于空间大小如此之小时，你能做出的最高杠杆决策，就是选择把预算花在空间的哪个区域——这是一个规划问题，而不是搜索精炼问题，应当被当作规划问题来对待。

4. 方法

图 1：HIERA 对搜索空间的重新框定

Figure 1 (paper Fig.1): HIERA 把固定空间的 kernel 优化，重新框定为按工作负载自适应地在三个嵌套实现方案之间做选择。

图 1 是论文在引入任何具体机制之前给出的概念性主张：现有方法要么落在”高层算子/库优化”这个角（安全、可移植、但上限有限），要么落在”直接手写/生成 CUDA kernel”这个角（上限最高，但失败率也最高）。HIERA 的主张是：对于某个具体工作负载而言，正确的选择并不总是落在同一个角落——应该按任务、甚至按每一轮精炼步骤，根据工作负载特征和性能剖析反馈来动态选择。

4.1 整体架构——四个 Agent，一个反馈闭环

图 2：HIERA 流水线

Figure 2 (paper Fig.2): HIERA 的完整流水线——契约增强的任务规格、分层的搜索空间+优化方向规划、策略转译、候选生成，以及评估/剖析反馈。

HIERA 由四个协作组件构成（论文把中间两个称为”Agent”——实际上四个都是带不同角色提示词的 LLM 调用，由一个控制循环编排）：

契约增强的任务规格（Contract-Augmented Task Specification，预处理步骤，严格说不算”Agent”——是确定性的）。
搜索空间决策 Agent（SSDA）——决定在哪里搜索：选哪个实现空间，以及在该空间内选哪个优化方向。
策略规划 Agent（SPA）——把选定的方向转译为具体、可执行的优化策略（具体的代码变换）。
优化 Agent（OA）——实际写出实现该策略的候选 kernel 代码。

生成之后，候选会经过编译检查 → 正确性验证 → 性能评估 → NCU 剖析，结果反馈进下一轮的 SSDA 调用。下面把每个环节拆开讲清楚。

4.2 契约增强的任务规格——为什么要这么麻烦

这里有一个细微但重要的工程点，论文讲得很到位：如果你直接把 KernelBench 原始任务丢给 LLM，说”写一个更快的版本”，模型就得额外重建一堆和优化本身毫无关系的样板代码——宿主端 Python 封装、PyTorch C++ 扩展绑定、输入生成/测试框架代码、参考实现的准确调用约定。每一个花在重新生成这些样板代码上的 token，都是没有花在真正优化思考上的 token；而模型每偏离一次预期接口（参数顺序错了、返回类型错了、dtype 错了），就多一种和优化想法好坏毫无关系、却会导致候选编译失败或正确性检查失败的方式。

HIERA 的解法：对每个任务，冻结一组模型永远不需要（也不允许）重新生成的契约文件：

torch_demo.py —— PyTorch 扩展封装。
cpp_source.cpp —— 连接 Python 和 CUDA kernel 的 C++ 绑定源码。
groundtruth.py —— 正确性参考实现。
params_semantics.json —— 只针对 Level 2/3 任务——结构化描述参数角色、约束、参数间依赖关系（Level 1 任务省略这个文件，因为单个算子的参数语义可以直接从接口推断出来）。

给定这些固定的产物，LLM 的任务就收窄到只产出一个文件：cuda_source.cu，以结构化响应的形式给出候选实现。之后这个文件会自动和冻结的契约文件拼接起来，通过标准的 KernelBench 测试框架编译。

设计选择讨论——为什么有效、明显的替代方案是什么、什么情况下会失效：

为什么有效： 消除了整整一类”优化思路是对的，但周边配套写错了”的失败，也让模型的上下文预算走得更远，因为不用每轮都重新推导不变的样板代码。
明显的替代方案： 让模型每次都生成一个完整的、可独立运行的程序（KernelBench-Caesar 实际上就是这么做的）。实现起来更简单，但根据 4.6 节的消融实验，在更难的任务层级上，正确性会付出惨重代价。
什么情况下会失效： 这个契约的质量完全取决于 params_semantics.json 的覆盖度。如果某个 Level 2/3 任务自动推导出的参数语义有微妙错误（比如对广播规则的描述有误），契约本身就会变成一个正确性 bug 的来源，而模型完全没有办法察觉或者绕过它——这是论文没有讨论的一个失效模式（第 9 节会展开）。

4.3 跨粒度搜索空间规划——HIERA 的核心机制

这是 HIERA 的招牌想法，形式化如下。在每一轮精炼步骤 $t$ ，**搜索空间决策 Agent（SSDA）**从三个按许可度递增排列的嵌套实现空间中选择一个：

纯 CUDA——只允许自定义手写 CUDA kernel。
CUDA 库——额外允许调用优化过的厂商库（如 cuBLAS）。
CUDA 库 + PyTorch——进一步允许退回到高层 PyTorch 算子。

注意这个嵌套关系：从”允许什么”的角度看，空间 3 ⊇ 空间 2 ⊇ 空间 1，但这与 Agent 拥有的底层控制力恰好相反——空间 1 给予最大控制力但实现负担最重；空间 3 实现负担最小，但能压榨出的收益上限也被限制住了。

记 $t$ 时刻选定的空间为 $g_t$ 。这个选择之所以重要，是因为它在任何代码写出来之前就决定了优化 Agent 能提出什么样的方案——对于一个复合的 Level-3 工作负载，如果一开始就选择 $g_t=$ “CUDA 库 + PyTorch”，Agent 就不用把有限的预算花在从零用原生 CUDA 重新推导算子依赖图和中间数据流上，而可以直接组合已知可靠的构建块，把自定义 CUDA 的精力留给性能剖析真正揭示出瓶颈的那一两个子算子。

4.4 领域知识引导的优化方向剪枝

在选定的空间 $g_t$ 内，HIERA 进一步通过从一个固定的、由前人 GPU 优化文献和专家实践中提炼出的五类常见性能瓶颈分类里，选择一个优化方向来收窄搜索：

C —— 控制流与边界特化：消除条件分支带来的 warp 分歧，对边界处理代码路径做特化（例如 stencil 中为内部 tile 和边缘 tile 分别写不同 kernel），让常见情形避开分支开销。
P —— 线程与 warp 级并行度：确保暴露出足够的独立工作来打满 GPU 的 SM——例如增大 grid/block 维度，重构循环以增加每个线程的可并行工作量。
M —— 显存事务效率：合并全局显存访问、避免共享内存 bank 冲突、将访问模式对齐到硬件事务粒度。
R —— 数据复用与数据搬运流水化：通过 tiling 最大化片上（共享内存/寄存器）数据复用，让计算和数据搬运重叠（双缓冲、类似 cp.async 的流水线）。
T —— Tensor Core 与指令流水线利用率：重构计算以命中硬件矩阵乘加（mma）单元，或改善指令级流水调度以避免停滞。

记 $D=\{C,P,M,R,T\}$ 。在步骤 $t$ ，SSDA 用任务规格 $\tau$ 、选定空间 $g_t$ 、当前候选 $x_t$ 、性能剖析反馈 $F_t$ （来自 NCU）、专家知识 $E$ 、固定评分准则 $R$ 给每个方向打分：

s_t = \text{SSDA}(\tau, g_t, x_t, F_t, E; R) = \big(s_t^d\big)_{d \in D} \tag{3}

即输出是一个五维向量，每个方向对应一个相关性分数。当前步骤选定的主方向就是取最大值：

d_t^{*} = \arg\max_{d \in D} s_t^{d} \tag{4}

设计选择讨论： 为什么用一个固定的五分类，而不是让模型用自由的自然语言描述它感知到的瓶颈？固定分类让方向选择这一步变得结构化、跨步骤可比较——SSDA 输出的五个分数本质上等价于一个轻量分类头，这一方面让决策可审计、可用统一 schema 记录下来，另一方面让下游策略规划 Agent 能够直接调用一个按这五个类别索引好的、专家整理的策略库（见下一节），而不需要去解析任意自由文本的瓶颈描述。明显的替代方案——自由形式的瓶颈诊断——表达力严格更强（某个工作负载可能真的被这五类分类无法干净捕捉的东西所限制，例如指令缓存压力或跨 block 归约的原子操作争用），但牺牲了这种结构化。论文明确承认这个分类”是操作层面的，而非穷尽的”，也就是说，作者承认这是对一个更丰富瓶颈空间的实用性五分桶近似，刻意用完备性换取下游策略检索的可靠性。

4.5 策略规划 Agent——从方向到可执行策略

给定主方向 $d_t^{*}$ ，**策略规划 Agent（SPA）**产出一个具体、结构化的策略：

\pi_t = \text{SPA}(\tau, g_t, x_t, F_t, d_t^{*}, E), \quad \pi_t \in \Pi_{g_t, d_t^{*}} \tag{5}

这里 $\Pi_{g_t, d_t^*}$ 表示在当前实现空间和方向约束下有效的策略空间——例如”为了利用 Tensor Core mma 而重构”这个策略（ $d^*=T$ ）只有在 $g_t=$ 纯 CUDA 或 CUDA 库时才有意义（Tensor Core 内联指令不是从纯 PyTorch 调用里能触发的东西）。这种把策略空间按（空间, 方向）做类型约束的设计，正是防止 SPA 提出一个 OA 在结构上根本无法执行的方案的关键。

4.6 优化 Agent——生成下一批候选

最后，**优化 Agent（OA）**拿到策略 $\pi_t$ 并将其应用，生成下一批候选 kernel：

X_{t+1} = \text{OA}(\tau, g_t, x_t, F_t, \pi_t) \tag{6}

注意这里产出的是一个候选集合 $X_{t+1}$ （不是单个候选）——HIERA 的实验采用”种群式”循环，每轮生成多个候选（主对比实验中每轮 6 个），存活下来的候选作为下一轮的父代。

4.7 反馈驱动的评估循环——逐步展开的伪代码

把 4.2–4.6 节整合起来，下面是完整的 HIERA 循环，逐步写出（这是我根据论文分散在多个小节的文字描述自行重构的，不是论文原文的逐字算法块，因为论文本身没有给出单独一个算法框）：

算法 1：HIERA 优化循环
输入：任务 τ，参考实现 f_τ，硬件 H，总预算 B，
      轮数 R，每轮候选数 K（论文取 R=3，K=6，B=RK=18）
输出：预算内找到的最优有效候选 x*_τ

 1: spec  ← ContractAugment(τ)                 # 冻结封装/绑定/参考文件
 2: x_0   ← InitialCandidate(spec)              # 种子候选（例如朴素翻译）
 3: g_0   ← SSDA.select_space(spec, x_0, F=∅, E)   # 首次实现空间选择
 4: best  ← x_0（若 Valid(x_0; τ) 成立）否则为 NONE
 5: for t = 0 .. R-1:
 6:     s_t     ← SSDA.score_directions(spec, g_t, x_t, F_t, E; R)   # 公式 3
 7:     d*_t    ← argmax_d s_t[d]                                     # 公式 4
 8:     π_t     ← SPA.plan(spec, g_t, x_t, F_t, d*_t, E)              # 公式 5
 9:     X_{t+1} ← OA.generate(spec, g_t, x_t, F_t, π_t, K)            # 公式 6，共 K 个候选
10:     for X_{t+1} 中的每个候选 x:
11:         若 x 编译不通过：跳过（因编译失败丢弃）
12:         若 x 不满足 Valid(x; τ)：跳过（因结果错误丢弃）
13:         在 H 上测量 t(x; τ, H)                                    # 时延
14:         F_x ← NCU_profile(x, H)                                   # 硬件计数器
15:         计算 s(x;τ,H) = t(f_τ;H) / t(x;τ,H)                       # 公式 1
16:         若 best 为 NONE 或 s(x;τ,H) > s(best;τ,H)：
17:             best ← x
18:     x_{t+1} ← best-of-round(X_{t+1})        # 本轮最强存活候选作为新的父代
19:     F_{t+1} ← NCU_profile(x_{t+1}, H)       # 反馈带入下一轮 SSDA 调用
20:     g_{t+1} ← SSDA.select_space(spec, x_{t+1}, F_{t+1}, E)   # 每轮重新决定实现空间
21: 返回 best

有几点很容易在第一遍读论文时忽略，值得专门指出：

实现空间的决策是每一轮都重新做的（第 20 行），而不是一开始就定死。这意味着 HIERA 可以，比如说，一开始对一个 Level-3 任务选用更宽松的”CUDA 库 + PyTorch”空间，快速拿到一个能跑、能测的结果，然后随着剖析结果揭示出具体瓶颈子算子，逐步向”纯 CUDA”收窄——论文明确把这种”由粗到细”的策略框定为它预期的使用方式。
无效候选是被静默丢弃的，而不是被”打个折扣”或”解释一下”。 没有针对”接近正确但不完全对”的候选给予部分信用；根据公式 (2) 的形式化，只有有效候选才有资格比拼速度。
预算的核算完全基于候选数量（ $B=R\times K$ ），而不是墙钟时间——这一点很重要，因为第 6 节中”预算 $B\in\{1,6,12,18\}$ “的曲线，是从同一条生成轨迹里事后截取出来的，而不是针对每个预算独立重新跑一遍。

5. 实验设置

评测基准： KernelBench 全部三个层级（共 250 个任务：Level 1/2/3 分别为 100/100/50）。
基座 LLM： DeepSeek-V3.2、Qwen3.6-Plus、Gemini-3.6-Flash——三个不同的底层模型，用来验证 HIERA 的收益是否与具体模型无关（一个训练无关的方法，原则上应该能不经重新训练就迁移到不同基座模型上，这一点与需要重新训练的 CUDA-L1/Kevin 不同）。
基线： KernelBench-Caesar（无约束的迭代基线）、CUDAForge（固定纯 CUDA 空间的 Agentic 基线）、CUDA-L1（RL 训练的基线，因为不随基座模型变化，只报告一次）。
预算： $R=3$ 轮精炼，每轮 $K=6$ 个候选，单任务最大预算 $B=18$ 个候选；同时通过截取同一条轨迹，报告更小子预算 $B\in\{1,6,12\}$ 下的累积结果。
硬件： 每个任务使用单块 NVIDIA A100-PCIe-40GB（论文表 1 给出完整软件版本：CUDA 12.8、PyTorch 2.7.1、cuDNN 9.5.1）。
评估协议： 每个有效候选用 FP32 精度测试，3 次热身运行后跑 100 次测量运行；报告完整预算内速度最快的有效候选。

另外还有两个专门的对照实验用来隔离特定设计组件的贡献：

实现空间对比（RQ2，跨粒度）： 在 90 个抽样任务上（每个层级 30 个），把 HIERA 和三个固定空间变体——始终纯 CUDA、始终 CUDA 库、始终 CUDA 库 + PyTorch——进行对比，其余所有设置（基座模型用 Qwen3.6-Plus、契约规格、方向规划流程、预算）保持一致。这样就能隔离出自适应空间选择的边际价值，其余变量全部固定。
消融实验：“HIERA w/o Contract”（去掉冻结产物规格，强制模型重新生成样板代码）和”HIERA w/o Planning”（完全去掉分层的 SSDA/SPA 规划，替换成一个固定、无约束的优化提示词）。

6. 结果

6.1 RQ1——主对比：样本效率与整体质量

图 3：有限预算下的性能曲线

Figure 3 (paper Fig.3): 在 Qwen3.6-Plus 上，HIERA vs. CUDAForge vs. KernelBench-Caesar 的 fast0/fast1/fast2 随搜索预算 B 变化的曲线。

最亮眼的数字：在 $B=1$ 时（即只生成了一个候选，还没经过任何迭代精炼），HIERA 就已经达到 $71.6\%$ 的 $\text{fast}_0$ 和 $32.4\%$ 的 $\text{fast}_1$ ——分别领先 KernelBench-Caesar $+30.8$ 和 $+22.4$ 个百分点，领先 CUDAForge $+52.4$ 和 $+13.6$ 个百分点。这是一个引人注目的结果：它说明 HIERA第一次猜测，因为有空间和方向规划步骤的指导，已经比基线经过整整一轮迭代精炼后更有可能既有效又快。这一优势在 $B=18$ 时依然存在但差距有所收窄：HIERA 达到 $92.4\%/60.4\%$ ，对比 KernelBench-Caesar 的 $85.2\%/29.6\%$ 和 CUDAForge 的 $85.6\%/44.8\%$ 。

有意思的是，在更严格的 $\text{fast}_2$ 指标（ $>2\times$ 加速）上，HIERA 的领先幅度存在但更小，在大预算时 CUDAForge 甚至反超（见下方表 2 与图 4 的箱线图）——论文对这一细节做了诚实报告而没有回避，我在第 9 节批判性分析中会回到这一点，因为它让”HIERA 全面碾压”这个叙事变得没那么简单。

6.2 跨模型对比（表 2，复现并解读）

基座 LLM	方法	L1 fast0/fast1/fast2	L2 fast0/fast1/fast2	L3 fast0/fast1/fast2
DeepSeek-V3.2	HIERA	91/56/22	90/35/9	64/42/10
DeepSeek-V3.2	KernelBench-Caesar	77/17/3	80/34/5	46/12/2
DeepSeek-V3.2	CUDAForge	90/22/6	78/27/15	60/32/10
Qwen3.6-Plus	HIERA	97/68/18	99/62/20	70/42/10
Qwen3.6-Plus	KernelBench-Caesar	91/21/6	90/40/13	64/26/6
Qwen3.6-Plus	CUDAForge	90/47/18	88/45/28	72/40/8
Gemini-3.6-Flash	HIERA	95/63/20	96/50/14	68/40/12
Gemini-3.6-Flash	KernelBench-Caesar	80/19/4	85/32/11	52/14/2
Gemini-3.6-Flash	CUDAForge	92/45/16	91/40/20	70/34/8
（不区分模型）	CUDA-L1（RL 训练）†	74/19/10	81/36/13	72/50/6

†CUDA-L1 不随基座模型变化，只报告一次。

仔细读这张表（而不是只引用论文摘要里”27 组对比中 22 组最优或并列最优”的总结）会发现很多真实的层次感：HIERA 的优势在 $\text{fast}_1$ 上最大也最一致——在所有九个”层级×模型”组合上，它都是”是否真正比参考实现更快”这一指标上最强的系统。但在 $\text{fast}_2$ （ $>2\times$ 门槛）上，CUDAForge 在 DeepSeek-V3.2/L2（15 对 9）、Qwen3.6-Plus/L2（28 对 20）、Gemini-3.6-Flash/L2（20 对 14）上都明确胜出——而且每一次都恰好在Level 2上。CUDA-L1 尽管完全没有针对具体任务的规划，却在 Level-3 的 $\text{fast}_1$ 上以 50% 对 HIERA 最好的 42% 胜出。所以更公允的读法是：HIERA 是总体最好、最可靠的系统，但它并没有在每个层级的每个指标上全面碾压——RL 训练的 CUDA-L1 在最难的（ $\text{fast}_2$ 、Level 3）尾部指标上仍保留优势，很可能是因为它的训练过程专门针对”深度、激进的 CUDA 重写”这一目标做了优化，代价是牺牲了 HIERA 所提供的通用性。

6.3 RQ2——隔离自适应空间选择的价值

图 4：固定 vs. 自适应实现空间的加速比分布

Figure 4 (paper Fig.4): 在 90 个抽样的 KernelBench 任务上，三种固定实现空间与 HIERA 自适应选择的最优验证加速比分布（箱线图 + 散点）。

在我看来，这是全文信息密度最高的一张图，直接支撑了核心论点。解读这张箱线+散点图：

纯 CUDA（始终自定义 CUDA）拥有最高的天花板——观测到的最大加速比是 $9.32\times$ ，是整张对比图里最亮眼的单一数字——但均值只有 $1.00\times$ ，中位数只有 $0.62\times$ （也就是说，对于典型任务而言，“一律手写原生 CUDA”实际上相对 PyTorch 参考实现反而是净变慢！），方差也最大（1.79）。这就是支撑论文动机论断的实证证据：不加约束的低层搜索，会把预算耗在最终比什么都不做还差的实现上。
CUDA 库（允许调库）相比纯 CUDA 把方差大致砍半（降到 0.27），但均值/中位数仍然比较平庸。
CUDA 库 + PyTorch（最宽松、偏框架化）达到最低方差（0.26），均值/中位数提升到 $1.10\times$ / $1.01\times$ ，但观测到的上限只有 $2.98\times$ ——安全是以牺牲上限为代价换来的。
HIERA（自适应）在聚合统计量上做到了两全其美：在四者中均值最高（ $1.42\times$ ）、中位数最高（ $1.15\times$ ），同时相对纯 CUDA 把方差降低了 $23.4\%$ ——也就是说，它捕获了纯 CUDA 大部分的上限潜力（散点图中有若干点高于 $2\times$ ，最高接近 $5\times$ ），同时避开了困扰”始终纯 CUDA”策略的那条”比什么都不做还差”的尾巴。

设计选择讨论： 这张图是全文最有说服力的证据，值得说清楚为什么有说服力：这是一个受控实验（同样的 90 个任务、同样的基座模型、除了空间选择策略以外的一切都相同），所以差异可以完全归因于自适应性本身，而不是某个混杂因素（比如更强的基座模型、更大的预算）。这个消融实验的明显替代方案，是和一个完全不同的 Agent 框架做正面对比，但那样会把”自适应空间选择是否有效”和”这个框架的提示词工程/编排是不是整体更好”这两件事混在一起——这个受控的 90 任务对比精确隔离出了论文想要归功的那一个变量。

6.4 消融实验

图 5：消融契约增强规格

Figure 5 (paper Fig.5a)：随着任务复杂度从 Level 1 上升到 Level 3，去掉契约增强模板导致的有效性/加速比下降幅度递增。

去掉冻结的契约文件（“w/o Contract-Augmented Template”）在每个层级都会造成性能下降，但下降幅度随任务复杂度急剧放大：Level 1 的 $\text{fast}_0$ 从 97% 降到 86%（11 个百分点，尚可承受），但 Level 2 从 99% 骤降到 26%（73 个百分点的崩塌），Level 3 从 70% 降到 10%（60 个百分点的崩塌）。直觉解释：Level 1 任务是单一算子，即使没有冻结契约，需要重建的样板代码也不多；Level 2/3 任务需要从零重建复杂的算子依赖图和多文件绑定，而这恰恰是模型在没有固定框架的情况下最容易出错的地方。

图 6：消融分层搜索空间规划

Figure 6 (paper Fig.5b)：去掉基于 Agent 的搜索空间/方向规划，几乎不影响有效性（fast0），但会让真正获得加速的任务比例大幅崩塌（fast1/fast2）。

这个消融实验讲的是相反的故事，图 5 和图 6 的对比本身就很有启发性。去掉分层规划（替换成一个固定、无约束的优化提示词）几乎不影响 $\text{fast}_0$ （Level 2：99% → 96%，仅下降 3 个百分点）——模型仍然通常能产出某个有效的东西，即使没有引导。但 $\text{fast}_1$ 和 $\text{fast}_2$ 崩塌了：Level 2 的 $\text{fast}_1$ 从 62% 降到 6%，Level 1 的 $\text{fast}_1$ 下降 52 个百分点，Level 3 下降 38 个百分点。换句话说：规划的主要作用不是帮你产出能跑的代码——而是帮你产出既能跑、又真正更快的代码。 论文自己的总结抓得很准：“契约增强主要保住了可行性，而分层规划提升了加速效果”——这是两个针对不同失败模式的组件，这个结论令人满意且可证伪（消融实验的设计本来完全有可能揭示出这两个组件在做同一件事，但实证结果并非如此，这一点很好）。

6.5 RQ3——案例研究：科学计算中的 Stencil 计算

图 7：走出 ML 工作负载——一个二维盒式 Stencil

这是一个很好的通用性检验：论文完全跳出 KernelBench，去问 HIERA 的方法能否迁移到非 ML的 GPU 工作负载上。任务是：一个半径 $R=3$ （即 $7\times7=49$ 点邻域）的二维盒式 stencil，输入尺寸 $10240\times10240$ ，重复计算 $10{,}240$ 次，报告的是（摊薄后的）单步时延。选定的密集计算基线是 cuDNN 的 cudnnConvolutionForward，配置为单通道 $7\times7$ 零填充互相关运算，FP64 精度——这是用卷积来计算 stencil 的一种合法方式，也是一个真正强大、经厂商调优的参考实现（7.23 毫秒/步）。

搜索经过 5 轮，种群规模为 10（总预算 50 次候选评估——明显大于主 KernelBench 对比中使用的 $B=18$ ，因为这是一次单独的深度案例研究，而不是大范围的基准测试跑）。运行时间单调下降： $13.71 \to 13.00 \to 9.86 \to 7.39 \to 4.71$ 毫秒。候选直到第 5 轮才超过 cuDNN，最终落在 4.71 毫秒——相对 cuDNN 提升 $34.8\%$ ，即 $1.53\times$ 加速，相对第一个候选整体下降 $65.6\%$ 。

设计选择讨论——这个对比公平吗？ 论文在这里的表态明确且恰如其分——它把这个案例研究定性为”可行性证据，而非全面验证”，这是对的：这只是一个算子、一种配置、一块GPU。半径为 3 的 stencil 具有对共享内存 tiling 特别有利的具体结构（规则、小、对称的邻域），这恰好是自定义 CUDA 往往能击败通用密集卷积表达方式的地方；如果换成一个访问模式高度不规则或数据相关的 stencil，或者一个真正大半径、寄存器/共享内存压力成为硬约束的 stencil，会是一个截然不同（也更难）的测试。

7. 局限性（作者自述）

论文自己的”范围与局限性”部分写得相当直接，我先如实列出，第 9 节再补充我认为论文遗漏的部分：

单一 GPU 架构： 所有实验都在 NVIDIA A100 上完成；性能特征（占用率目标、显存带宽与算力之比、Tensor Core 代际差异）在 H100/B200 级别硬件上会有实质差异，论文并未声称已学到的方向分类或空间选择启发式规则不经重新验证就能直接迁移。
主基准只用了一种精度格式： KernelBench 评估全程用 FP32；stencil 案例研究用 FP64。真正主导 LLM 训练/推理工作负载的低精度格式（FP16/BF16/FP8）被明确标注为尚未探索。
没有多 GPU 评估： 每个任务都分配到单块 GPU；引入 GPU 间通信作为新的主导成本轴的多 GPU kernel 不在本文范围内。
候选生成的随机性： 论文提到时延测量（3 次热身 + 100 次测量运行）经过重复以控制硬件/测量噪声，但候选生成本身仍然是随机的——重新跑同一套流水线很可能产出不同的轨迹，论文没有报告端到端重复运行（区别于同一次运行内 90 个抽样任务之间的方差）的方差数据。
Stencil 案例研究只有一个算子/一种配置，明确定性为”可行性”证据，而非针对科学计算工作负载的、经过验证的通用能力。

8. 可复现性说明

基准/任务： KernelBench 公开可用（github.com/ScalingIntelligence/KernelBench）；250 个任务及其参考实现都是开放的。
基座模型： DeepSeek-V3.2 和 Qwen3.6-Plus 依发布条款可通过 API/开放权重获取；Gemini-3.6-Flash 仅限 API——要完全按位复现，需要拿到论文所用的确切模型快照（托管 API 上的模型会随时间漂移/弃用，这是任何针对商业 LLM 端点做基准测试的论文都会遇到的常见复现性隐患）。
硬件： 论文指定了 A100-PCIe-40GB、驱动 565.57.01、CUDA 12.8、PyTorch 2.7.1+cu126、cuDNN 9.5.1（表 1）——这种详细程度是良好实践，在可比硬件上做类似重跑应当是可行的。
未发布（截至本笔记撰写时）： 论文没有提及 HIERA 自身编排代码（SSDA/SPA/OA 的提示词、契约生成工具、params_semantics.json 的推导流程）的公开代码/工件发布——而这恰恰是本工作真正的创新产物。没有这些，独立复现只能依据论文的文字描述重新实现四 Agent 循环和五方向专家知识库（我为了写出上面的算法 1 就是这么做的）——可行，但并不轻松，而且几乎肯定无法和作者内部的提示词做到逐字一致。
采样温度： 为公平起见，所有方法统一固定为 0.3——如果要尝试重跑，这一点值得保留，因为 kernel 生成质量对采样温度很敏感，不同条件之间意外引入温度差异很容易带偏结果。

9. 批判性分析

(a) 本文特有的弱点与瑕疵。

“27 组对比中 22 组”这个头条式表述，部分掩盖了一个真实存在的、按层级细分的弱点。 正如第 6.2 节所梳理的，HIERA 在 $\text{fast}_2$ 上输给 CUDAForge，而且恰好都在 Level 2 上——在全部三个基座 LLM 上都是如此，不是偶然。Level 2（融合的多算子模式）大概正是”自适应地组合 CUDA 库并做轻量自定义融合”（HIERA 的方式）与”全力押注自定义 CUDA 融合”（CUDAForge 的方式）权衡最直接的层级，而数据显示，就”极致”（ $>2\times$ ）加速这一具体目标而言，坚决押注自定义 CUDA 仍然能赢。论文的汇总统计（“22/27”）在字面上没错，但只读摘要/结论的读者不会知道这个失利是系统性的、按层级分布的，而不是噪声。
五方向分类的覆盖度只是被断言，而没有被实测。 论文说 $\{C,P,M,R,T\}$ 分类”是操作层面的而非穷尽的”，这是一个诚实的免责声明，但没有任何实验去量化：SSDA 的实际瓶颈诊断（来自 NCU 剖析）有多频繁落在这五类之外，以及落在外面时会发生什么——SSDA 是不是只是挑一个契合度最差的类别凑数，从而悄悄拉低策略质量？这本可以是一个成本很低但很有启发性的实验：为一批任务记录原始的 NCU 停滞原因构成，检查它有多经常能干净地映射到这五个桶之一。
图 4 用来做核心跨粒度对比的”90 个抽样任务”只占完整 250 任务集的 36%，抽样过程（“每个层级各 30 个”）说是随机的，但论文没有给出具体的随机种子/任务 ID，这一点对精确复现图 4（区别于用完整 250 任务集的表 2 主结果）很重要。
全文没有任何置信区间或统计显著性检验。 表 2 和 RQ1/RQ2 图中的每一个数字，都是每个（方法、基座 LLM、层级）组合单次运行的点估计。考虑到论文自己在局限性部分承认候选生成是随机的，而且几份报告的方法间差距都在个位数百分点范围内（例如 DeepSeek-V3.2/L3：HIERA 64% vs. CUDAForge 60% 的 $\text{fast}_0$ ——只差 4 个百分点），到底哪些差异能在重复运行的方差分析中站得住脚，其实并不清楚。

(b) 作者低估或遗漏的局限性。

四 Agent 流水线本身的成本/时延从未被报告。 HIERA 每个任务的每一轮精炼，至少要调用一次 SSDA、一次 SPA、一次 OA，再加上 NCU 剖析开销——相比 KernelBench-Caesar 这种”生成并检查”的单调用基线，这是明显更高的每候选 LLM 推理成本。论文没有报告任何关于端到端耗时、总 token 成本、或者在匹配预算 $B$ 下相对基线的美元成本的数字。对于一篇整体前提就是”高效利用有限预算”的论文来说，把预算纯粹按”候选数量”衡量、却完全忽略四次 LLM 调用的流水线和单次调用基线之间截然不同的单候选成本，是一个真实的空缺——完全有可能在匹配算力/美元预算（而非匹配候选数量预算）的情况下，HIERA 的部分优势会缩水，因为它的每个候选生产成本更高。
契约增强消融和规划消融之间的交互没有联合测试。 我们知道单独去掉契约影响很大（图 5），单独去掉规划影响也很大（图 6），但论文没有报告”两者都去掉”的条件，所以无法判断它们的效果是可加的、次可加的（各自防止的失败之间有重叠），还是超可加的。这在已有的消融实验网格上只是补一行的成本。
论文没有讨论当 SSDA 的空间/方向选择本身就选错了会发生什么——比如对一个 Tensor Core 根本不适用的工作负载（例如形状不利于分块的场景），却坚持选择”纯 CUDA” + 方向”T”（Tensor Core）。系统能否在剩余的轮次里通过剖析反馈回路纠正过来，还是说一次糟糕的早期规划决策会永久性地浪费掉预算？考虑到轮数很少（ $R=3$ ），一次糟糕的早期决策完全有可能主导整条轨迹，但论文没有提供任何失败案例分析来检验这一点。

(c) 具体的、可操作的改进建议。

报告按成本归一化的对比。 在现有的候选数量轴之外，增加一个总 LLM token（或墙钟秒数，或 API 美元成本）的预算轴，用这个新轴重新画一遍图 3。这能直接回答”在匹配成本而不只是匹配候选数量的情况下，HIERA 的优势是否依然存在”——这是对任何想判断该框架是否值得在生产环境采用的人而言，最重要的缺失证据。
在图 5/6 的对比网格中加一列”同时去掉契约和规划”的联合消融——鉴于消融实验的基础设施显然已经具备，运行成本很低，能直接回答上面提到的可加性问题。
测量并报告 SSDA 的”方向分类误判率”：对一批任务，把 SSDA 选出的方向和原始 NCU 瓶颈分类做对比，报告排名最高的方向与实际主导的停滞原因相匹配的频率。这能用数据验证（或证伪）分类覆盖度的主张，而不只是定性的免责声明。
把 stencil 案例研究扩展到至少 3–5 个算子/配置，覆盖”这对手写 CUDA 有多友好”这条谱系上的不同点（例如改变 stencil 半径，尝试不规则/数据相关的访问模式，尝试一个预期库基线会获胜的工作负载）——一个未说明选取标准、单一的 $R=3$ stencil 案例研究，难免让人（合理地）怀疑它是不是因为提前知道效果好才被选中的。
报告重复运行的方差——哪怕只针对表 2 里某一个代表性的（基座模型、层级）单元格重复跑 3 次，也能让读者判断那些较小的差距（个位数百分点）到底是有意义的信号，还是噪声范围内的波动。

10. 结语

HIERA 的核心贡献是一次重新框定，而不是一个新的底层优化技巧：它论证——并用一个真正受控的消融实验（图 4）加以支撑——在哪个抽象层次里搜索这个选择本身，对于预算受限的 LLM kernel 优化 Agent 而言，是可用的最高杠杆决策——事实上比”在固定层次内怎么搜索”杠杆更高。再配合一个务实、低摩擦的方案来修复由样板代码引起的正确性失败（契约增强规格），这个训练无关的系统在几乎所有指标上都广泛超越了此前的训练无关基线，甚至在大多数指标上小幅领先一个需要 RL 训练的基线（CUDA-L1），却完全不需要更新任何权重。论文诚实列出的这些附加条件——在最严格加速比门槛上出现的、按层级分布的失利，四 Agent 流水线未被报告的成本开销，以及一个只有单一算子的通用性案例研究——意味着这篇论文最好被理解为对一个设计原则（把粒度当作一个决策，而非常量）的有力证据，而不是一个可以直接拿来生产使用的成品工具。如果你正在构建或评估 Agentic kernel 优化系统，可操作的启示很明确：不要把 Agent 硬编码死在某一个实现海拔高度上；如果你确实给了它自由选择的权利，就要衡量结果的分布（均值/中位数/方差），而不只是最好情形下的天花板——图 4 的箱线图之所以是全文最有用的一张图，正是因为它把这种”看分布而不只是看最好情形”的论证做得让人无法忽视。