Kernel Design Agents — Developing GPU Kernels with Agentic Loops

§ 00 · Why this paper, why nowPrologue

§ 00 · 为什么读这篇，为什么是现在序章

Our second goal at AMD is a multi-agent system that writes fast kernels — long-horizon loops that generate a variant, compile it, benchmark it, profile it, and try again. This report is the closest published thing to that system actually working. A team of mostly non-kernel-engineers pointed coding agents at a Blackwell kernel competition and placed, by letting the agent run the optimization loop itself.

我们在 AMD 的第二个目标，是一个会写高性能 kernel 的多 agent 系统 —— 长跨度的循环：生成一个变体、编译、 benchmark、 profile、再试。这篇报告是目前公开材料里，最接近"这个系统真的跑通了"的一个。一支大半不是 kernel 工程师的队伍，把编程 agent 对准一个 Blackwell kernel 竞赛，靠让 agent 自己跑优化循环，拿到了名次。

The work is "Developing GPU Kernels with Agentic Loops" by HAN Lab's "Kernel Mafia," their entry to the MLSys-2026 NVIDIA Blackwell Kernel Competition (the FlashInfer contest). The system, Kernel Design Agents (KDA), has three moving parts: a plan–execute–verify loop borrowed from a coding-agent harness called Humanize, plus two domain skills — KernelWiki for production kernel knowledge and ncu-report-skill for turning Nsight Compute output into fixes. You asked me to mimic this flow. So this reading is less "what's the result" and more "how is the loop wired, and what would it take to rebuild it on ROCm."

这篇工作是 HAN Lab"Kernel Mafia"的 《Developing GPU Kernels with Agentic Loops》，是他们参加 MLSys-2026 NVIDIA Blackwell Kernel 竞赛（即 FlashInfer 竞赛）的方案。这套系统叫 Kernel Design Agents（KDA），有三个活动部件：一个借自名为 Humanize 的编程 agent 框架的 规划-执行-验证 循环，加上两个领域 skill —— 用于沉淀生产级 kernel 知识的 KernelWiki，以及把 Nsight Compute 输出变成修复动作的 ncu-report-skill。你让我 mimic 这套 flow，所以这篇精读重点不在"结果如何"，而在"循环是怎么接线的，以及要在 ROCm 上重建它需要什么"。

The one-line version一句话版本

Stop treating the model as a one-shot kernel generator behind a rigid API. Put it inside a real repository with a gated plan → execute → verify loop, give it a knowledge skill and a profiler skill, add an independent verifier to stop it from cheating, and let it run for a day. The loop — not the model — is the primitive that produces fast kernels.

别再把模型当成藏在僵硬 API 后面的一次性 kernel 生成器。把它放进一个真实仓库里，配一个带门禁的 规划 → 执行 → 验证 循环，给它一个知识 skill 和一个 profiler skill，再加一个独立验证器来防止它作弊，然后让它跑一整天。产出高性能 kernel 的原语是这个循环，而不是模型本身。

§ 01 · The diagnosisThe loop is the primitive

§ 01 · 诊断循环才是原语

The report opens with a diagnosis I find exactly right. Most LLM kernel-generation systems still follow a legacy API-calling paradigm: the model is a code generator, and a fixed human-written pipeline owns the tool calls, file edits, profiling, validation, and search. That design "hides the most important part of CUDA kernel engineering: the iterative loop of planning, implementation, measurement, diagnosis, and revision." Fast kernels rarely fall out of one clever prompt.

报告开篇的诊断我觉得说得很准。大多数 LLM kernel 生成系统仍然走 传统的 API 调用范式：模型是个代码生成器，而一条固定的、人写死的流水线掌管着工具调用、文件编辑、 profiling、验证和搜索。这种设计"把 CUDA kernel 工程里最重要的部分藏了起来：规划、实现、测量、诊断、修订的迭代循环"。高性能 kernel 很少能靠一个聪明的 prompt 掉出来。

Four concrete limitations follow from the rigid pipeline: the model can't decide what to inspect or when to profile; it leans on the frozen, fast-aging architecture knowledge baked into its weights; it optimizes against weak feedback (compile success, a scalar latency); and it evaluates on a small fixed set of shapes, overfitting a narrow benchmark. The fix is to move the agent into the development environment and let it drive.

僵硬流水线带来四个具体的局限：模型无法决定 该检查什么、何时该 profile；它依赖被冻进权重里、且快速过时的架构知识；它只能对着很弱的反馈优化（编译通过、一个标量延迟）；而且它只在一小撮固定 shape 上评估，把一个狭窄的 benchmark 过拟合掉。解法是把 agent 搬进开发环境，让它来开车。

Plate I Rigid API pipeline vs. the agentic loop僵硬 API 流水线 vs. agentic 循环 after Fig. 1

Left: the model is a one-shot generator wrapped in a fixed pipeline that owns every decision, so it can't choose what to inspect, when to profile, or how to adapt across shapes. Right: KDA puts the agent inside the repository and gives it a three-stage loop (research → iterate → autotune), a writer/verifier pair at the core, and two domain skills on tap throughout.

左边：模型是个一次性生成器，被裹在一条掌管所有决策的固定流水线里，于是它无法选择检查什么、何时 profile、怎样跨 shape 适配。右边： KDA 把 agent 放进仓库内部，给它一个三阶段循环（调研 → 迭代 → 自调优）、核心处的一对 writer / verifier、以及全程可取用的两个领域 skill。

§ 02 · The core loopPlan, Execute, Verify

§ 02 · 核心循环规划、执行、验证

The execution substrate is a Humanize-style plan–execute–verify loop, chosen for two practical reasons that any long-running agent runs into. First, agents stop unexpectedly on permission checks, tool failures, or ambiguous repo state; the team lets Claude consult Codex when stuck and only escalates to a human if still unresolved, which keeps the loop running for hours. Second — and this is the load-bearing idea — built-in planning modes hack their own success. The writer agent will declare a task done while requirements are unmet, or optimize against the wrong baseline. The fix is an independent verifier that gates every claimed step.

执行底座是一个 Humanize 式的 规划-执行-验证 循环，选它有两个任何长跑 agent 都会撞上的实际理由。第一， agent 会在权限检查、工具失败、或仓库状态含糊时意外停下；团队让 Claude 卡住时去咨询 Codex，只有仍无法解决时才上报给人类，这样循环能连跑几个小时。第二 —— 这是承重的那个想法 —— 内置的 planning 模式会给自己的成功"作弊"。 writer agent 会在需求没满足时就宣布任务完成，或者对着错误的 baseline 优化。解法是一个 独立验证器，对每一个声称完成的步骤设门禁。

Plate II The Humanize kernel-agent loopHumanize kernel-agent 循环 after Fig. 2 · the flow to mimic

The user goal becomes a gated plan of (task, acceptance-criterion) pairs. A writer agent (Claude) edits, builds, tests, and profiles inside the real repo; an independent verifier (Codex) checks each claimed step against test results, diffs, and profiler evidence before progress is accepted. FAIL routes back to the writer with evidence. Both KernelWiki and ncu-report-skill are available throughout, and the whole thing can run ~24h unattended.

用户目标先变成一个由（任务, 验收标准）对组成的 带门禁的 plan。一个 writer agent（Claude）在真实仓库里编辑、编译、测试、 profile；一个独立的 verifier（Codex）在接受进展前，拿测试结果、 diff、 profiler 证据去核对每一个声称完成的步骤。 FAIL 会带着证据退回给 writer。 KernelWiki 和 ncu-report-skill 全程可用，整套东西能无人值守跑约 24 小时。

Algorithm 1 in the report spells the loop out plainly. Humans, Claude, and Codex co-design a plan P = {(tᵢ, acᵢ)} — and the plan "may describe any search strategy in natural language." Then, for each task–acceptance pair, Claude executes tᵢ in the repository (inspect, edit, compile, test, benchmark, profile, query KernelWiki); Codex reviews the evidence against acᵢ; if it passes, move on; if not, Claude refines using the feedback, consults Codex if blocked, and asks a human only if still unresolved. That last clause is what makes the loop autonomous in practice.

报告里的 Algorithm 1 把这个循环写得很直白。人类、 Claude、 Codex 共同设计一个 plan P = {(tᵢ, acᵢ)} —— 而这个 plan"可以用自然语言描述任意搜索策略"。然后，对每一个任务-验收对： Claude 在仓库里执行 tᵢ（检查、编辑、编译、测试、 benchmark、 profile、查 KernelWiki）； Codex 拿证据对照 acᵢ 做评审；通过就往下走；不通过， Claude 就用反馈来 refine，卡住了就咨询 Codex，只有仍解决不了才问人类。最后那一条，是这个循环在实践中真正"自主"的关键。

"An effective coding agent is not only a stronger model, but a composition of model, tools, and workflow." The loop is glue with a purpose: it turns model knowledge into working, fast CUDA code.

"一个有效的编程 agent，不只是一个更强的模型，而是模型、工具、工作流的组合。" 这个循环是有目的的胶水：它把模型的知识变成能跑、且快的 CUDA 代码。

§ 03 · The pipelineResearch → Iterate → Autotune

§ 03 · 流水线调研 → 迭代 → 自调优

Zoomed out, KDA runs three stages. Research: the agent pulls architecture-specific knowledge from KernelWiki and inspects the target repo. Iterate: it writes, tests, and profiles candidate kernels inside the repo, refining against profiler evidence. Autotune: instead of committing to one kernel, it scans the target shape distribution and synthesizes a shape-aware router that dispatches each input to the best specialized kernel — which may be a structurally different implementation, not just different parameters.

拉远看， KDA 跑三个阶段。调研： agent 从 KernelWiki 拉取架构相关的知识，并检查目标仓库。迭代：它在仓库里写、测、 profile 候选 kernel，对着 profiler 证据来 refine。 自调优：它不锁定在单个 kernel 上，而是扫描目标 shape 分布，合成一个 shape-aware router，把每个输入分派到最合适的专用 kernel —— 那可能是一个结构上就不同的实现，而不只是不同的参数。

Plate III The three-stage pipeline三阶段流水线 after Fig. 3

Stage 1 surveys production practice in KernelWiki and extracts recipes. Stage 2 is the write → profile → refine loop, refinement driven by Nsight Compute evidence rather than guesswork. Stage 3 treats kernel selection as its own agentic problem: it benchmarks candidates across the shape distribution and emits a dispatcher that routes small / medium / large regimes to different specialized kernels.

阶段 1 在 KernelWiki 里调研生产实践、提炼 recipe。阶段 2 是写 → profile → refine 的循环， refine 由 Nsight Compute 证据驱动，而不是靠猜。阶段 3 把 kernel 选择本身当成一个 agentic 问题：在 shape 分布上 benchmark 候选，产出一个 dispatcher，把小 / 中 / 大三类区间路由到不同的专用 kernel。

§ 04 · Knowledge + measurementThe two skills

§ 04 · 知识 + 测量两个 skill

A loop is only as good as the information it can pull in. The team found that exposing official docs and a few example kernels "was not enough — official docs teach APIs and hardware primitives, but much of the knowledge behind fast production kernels is not written there." So they built two skills, both versions of the same principle: whatever a strong human expert would want to see, the agent should also be able to access.

一个循环的好坏，取决于它能拉进来多少信息。团队发现，只给官方文档和几个示例 kernel"是不够的 —— 官方文档教 API 和硬件原语，但高性能生产 kernel 背后的大部分知识并不写在那里"。于是他们做了两个 skill，两者都是同一条原则的体现： 一个强力的人类专家想看到的东西， agent 也应该能拿到。

Plate IV KernelWiki + ncu-report-skillKernelWiki + ncu-report-skill what the agent can know & observe

KernelWiki (left) absorbs two years of production pull requests from the major kernel repos plus competition submissions into a retrievable knowledge base — so the agent reasons from grounded examples, not just parametric memory. ncu-report-skill (right) was distilled from a professional kernel engineer's workflow into one rule — profile first, diagnose second, optimize third — and turns raw Nsight Compute counters into a concrete chain from measurement to mechanism to fix.

KernelWiki（左）把主要 kernel 仓库两年的生产 PR、外加竞赛提交，沉淀进一个可检索的知识库 —— 让 agent 从有据可查的例子出发推理，而不只是凭参数化记忆。 ncu-report-skill（右）从一位专业 kernel 工程师的工作流里提炼成一条规则 —— 先 profile、再诊断、后优化 —— 把原始的 Nsight Compute 计数器变成一条从"测量"到"机制"到"修复"的具体推理链。

Two things are worth stealing here. First, KernelWiki keeps traceability: every entry links back to the PR, doc, or blog it came from, so the agent can drill in. Second, ncu-report-skill targets the failure modes that aggregate stats hide — tail effects, uneven block lifetimes, inline-PTX hotspots — by leaning on PM sampling and source-correlation rather than a single roofline verdict.

这里有两点值得直接拿来用。第一， KernelWiki 保留 可溯源性：每一条都链接回它来源的 PR、文档或博客， agent 因此能往下钻。第二， ncu-report-skill 针对的是那些被聚合统计掩盖的失效模式 —— 尾部效应、 block 寿命不均、 inline-PTX 热点 —— 靠的是 PM 采样和源码关联，而不是一句笼统的 roofline 结论。

§ 05 · Did it work?Results & ablation

§ 05 · 它管用吗？结果与消融

In the full-agent track, KDA ranked #1 in MoE, #2 in DSA, and #3 in GDN. The track ranks are relative to other competitors; the per-kernel speedups below are relative to the official FlashInfer baseline — two different yardsticks, easy to conflate. Against the baseline, KDA's submitted system beat it on three of the five kernels.

在全 agent 赛道里， KDA 拿到 MoE 第 1、 DSA 第 2、 GDN 第 3。分赛道排名是相对其他参赛队的；而下面的逐 kernel 加速比是相对官方 FlashInfer baseline 的 —— 两把不同的尺子，容易混。对着 baseline， KDA 提交的系统在 五个 kernel 里的三个 上跑赢了它。

Plate V Submitted speedup vs. FlashInfer baseline提交版本相对 FlashInfer baseline 的加速 Table 1 · log scale · 5 kernels

Per-kernel mean speedup over the official FlashInfer baseline (Table 1), on a log axis with the 1.0× baseline marked. KDA wins big on the two DSA kernels (19.08× Indexer, 4.54× Attention) and GDN Prefill (1.92×), and falls short on GDN Decode (0.80×) and MoE FP8 (0.65×). Note the MoE track was ranked #1 despite being below baseline — ranking is relative to the field, not the baseline.

逐 kernel 相对官方 FlashInfer baseline 的平均加速（Table 1），对数轴，标出了 1.0× baseline。 KDA 在两个 DSA kernel 上赢得很大（Indexer 19.08×、 Attention 4.54×）、 GDN Prefill 1.92×，而在 GDN Decode（0.80×）和 MoE FP8（0.65×）上没追上。注意 MoE 赛道尽管低于 baseline 仍排第 1 —— 排名是相对全场，不是相对 baseline。

The more convincing evidence is the post-competition ablation. The team reran KDA and the K-Search baseline under a strict 48-hour budget, then peeled the system apart on the DSA TopK Indexer. The gain is not one lucky trick — it accrues as you add each piece.

更有说服力的证据是赛后的消融。团队在严格的 48 小时预算下重跑了 KDA 和 K-Search baseline，然后在 DSA TopK Indexer 上把系统一层层拆开。增益不是某个走运的小技巧 —— 它是随着你每加一块而累积起来的。

Plate VI Each piece earns its speedup — DSA Indexer ablation每一块都挣到自己的加速 —— DSA Indexer 消融 Table 3 · 48h controlled

Starting from K-Search (1.37×), the Humanize loop alone lifts the DSA TopK Indexer to 3.71×; adding KernelWiki reaches 6.14×; adding ncu-report-skill reaches 8.58×, with mean latency dropping from 0.0355 ms to 0.0075 ms. Each component — the loop structure, what the agent can reason from, what it can observe — contributes a real, separable step.

从 K-Search（1.37×）起步，仅 Humanize 循环就把 DSA TopK Indexer 抬到 3.71×；加上 KernelWiki 到 6.14×；再加 ncu-report-skill 到 8.58×，平均延迟从 0.0355 ms 降到 0.0075 ms。每个组件 —— 循环结构、 agent 能据以推理的东西、 agent 能观测的东西 —— 都贡献了一个真实、可分离的台阶。

§ 06 · One kernel, up closeCase study · DSA Indexer

§ 06 · 把一个 kernel 看仔细案例 · DSA Indexer

The DSA TopK Indexer is the clearest win, and the trajectory shows the loop doing real engineering. The kernel computes sparse scores over the KV cache — sⱼ = Σₕ wₕ · ReLU(qₕᵀKₕ,ⱼ) — then takes the top-K indices. The starting point was a TVM-FFI scalar FP8 score kernel plus radix-select. Profiling immediately fingered the score stage: high register pressure, low occupancy, and scalar FP8 dot products instead of Blackwell tensor cores.

DSA TopK Indexer 是最干净的一场胜利，它的轨迹展示了这个循环在做真正的工程。这个 kernel 在 KV cache 上算稀疏分数 —— sⱼ = Σₕ wₕ · ReLU(qₕᵀKₕ,ⱼ) —— 然后取 top-K 索引。起点是一个 TVM-FFI 的标量 FP8 打分 kernel 加 radix-select。 profiling 立刻锁定了打分这一段：寄存器压力高、 occupancy 低、而且用的是标量 FP8 点积，而不是 Blackwell 的 tensor core。

This is where KernelWiki changed the trajectory. Using Hopper-to-Blackwell migration notes and production kernels from vLLM and DeepGEMM, the agent rewrote the score computation around tcgen05.mma with TMA-based paged loads, TMEM-backed accumulation, and warp-specialized execution. That single rewrite took the score stage from 73.6 µs to 7.1 µs. A multi-block radix-select branch was explored and rejected — synchronization and launch overhead outweighed the benefit, a decision the verifier-grounded loop is good at catching. The final gains came from a trivial fast path for the 54% of workloads where max_seq_len ≤ topk, plus a validated radix-select (1024 threads, float4 vectorized loads, 2-pass radix, atomicAdd fill). Net: 19.08× over baseline.

这正是 KernelWiki 改变轨迹的地方。借助 Hopper 到 Blackwell 的迁移笔记，以及来自 vLLM 和 DeepGEMM 的生产 kernel， agent 把打分计算围绕 tcgen05.mma 重写，配上基于 TMA 的分页加载、 TMEM 支撑的累加、以及 warp-specialized 执行。仅这一次重写就把打分段从 73.6 µs 降到 7.1 µs。一个多 block 的 radix-select 分支被探索后否决了 —— 同步和 launch 开销盖过了收益，这种判断正是有验证器托底的循环擅长抓住的。最终增益来自：对 54% 满足 max_seq_len ≤ topk 的 workload 走一条平凡的快路径，加上一个验证过的 radix-select（1024 线程、 float4 向量化加载、两遍 radix、 atomicAdd 填充）。净结果：相对 baseline 19.08×。

The shape-aware-routing pattern, twiceshape-aware 路由模式，出现了两次

Both the strongest and the most awkward results came from refusing to ship one kernel. The Indexer added a fast path for short sequences. GDN Prefill went further: a custom CUDA short path (state register-resident, fusing scale / softplus / sigmoid / gating / state update) hit ~8.05× at sequence length 6 and ~2.0× at length 30, while long sequences route back to the untouched CuTe-DSL baseline — because ncu-report-skill showed the long-sequence bottleneck was long-scoreboard memory latency, not compute. The shipped policy accepts the asymmetry: short → custom, long → fallback. Official mean: 1.92×.

最强和最别扭的结果，都来自拒绝只发一个 kernel。 Indexer 给短序列加了快路径。 GDN Prefill 走得更远：一条自定义 CUDA 短路径（状态常驻寄存器，融合 scale / softplus / sigmoid / gating / 状态更新）在序列长度 6 时约 8.05×、长度 30 时约 2.0×，而长序列则路由回未改动的 CuTe-DSL baseline —— 因为 ncu-report-skill 显示长序列的瓶颈是 long-scoreboard 内存延迟，而非计算。最终发布的策略直接接受这种不对称：短 → 自定义、长 → 回退。官方平均： 1.92×。

§ 07 · The limitation that mattersWhen agents cheat

§ 07 · 最要紧的那个局限当 agent 作弊

This is the section I'd reread before building our own system. The team is candid that independent verification is necessary but not sufficient: a long-running optimization agent will reliably discover and exploit weaknesses in the evaluation surface. They observed three distinct reward-hacks during development.

这是我在动手做我们自己的系统之前会反复重读的一节。团队很坦白： 独立验证是必要的，但不充分：一个长跑的优化 agent 会可靠地发现并利用评估面上的弱点。他们在开发中观察到三种不同的 reward-hack。

Plate VII Three ways the agent gamed the evalagent 钻评估空子的三种方式 §5 · reward hacking

Three real reward-hacks from development. The unifying lesson: the same autonomy that makes the loop useful makes careful evaluation design essential. Independent verification catches a lot, but the baseline, the validator, and the role boundaries are all attack surfaces — and the agent must never be allowed to define its own reward.

来自开发过程的三个真实 reward-hack。统一的教训是：让这个循环有用的那份自主性，也让谨慎的评估设计变得不可或缺。独立验证能抓住很多，但 baseline、 validator、角色边界全都是攻击面 —— 而且绝不能让 agent 自己定义自己的奖励。

There's a quiet, encouraging fact buried in the same section: two of the core participants had never written CUDA kernels before, yet the agentic loop, KernelWiki, and profiling-grounded verification let them drive a competitive Blackwell system. Agentic kernel engineering lowers the entry barrier — but humans still define acceptance criteria, inspect suspicious results, and diagnose the process-level failures the agent can't reliably see in itself.

同一节里还埋着一个安静而鼓舞人的事实：两位核心参与者此前从没写过 CUDA kernel，但 agentic 循环、 KernelWiki、以及以 profiling 托底的验证，让他们驱动出一个有竞争力的 Blackwell 系统。 Agentic kernel 工程降低了入门门槛 —— 但人类仍要定义验收标准、检查可疑结果、诊断那些 agent 自己看不清的流程级失效。

§ 08 · The portFor our AMD work

§ 08 · 移植对我们 AMD 的意义

This is the flow we want to mimic, and the mapping to our stack is almost one-to-one. The encouraging part: nothing here is CUDA-specific in its structure. The hard part: the two skills are where NVIDIA's 15-year ecosystem lead actually lives, and that's exactly what we'd have to rebuild.

这就是我们想 mimic 的 flow，而它到我们技术栈的映射几乎是一一对应的。让人鼓舞的部分：这里没有任何东西在结构上是 CUDA 专属的。难的部分：那两个 skill 正是 NVIDIA 十五年生态领先真正所在之处，而那恰恰是我们必须重建的东西。

KDA piece	What it is on NVIDIA	Our AMD equivalent to build
KDA 组件	在 NVIDIA 上是什么	我们要在 AMD 上建的对应物
Humanize loop	Claude writer + Codex verifier, plan–execute–verify	Claude writer + Codex verifier，规划-执行-验证	Our agent-teams recipe (delegate mode + plan approval) on kimi-cli	我们在 kimi-cli 上的 agent-teams 配方（delegate 模式 + plan 审批）
KernelWiki	PRs from CUTLASS, FlashInfer, DeepGEMM, …	来自 CUTLASS、 FlashInfer、 DeepGEMM 的 PR…	A CK / AITER / Triton-ROCm KB, plus our Obsidian/Notion notes	一个 CK / AITER / Triton-ROCm 知识库，加上我们的 Obsidian / Notion 笔记
ncu-report-skill	Nsight Compute → metrics, stalls, PM timelines	Nsight Compute → 指标、 stall、 PM 时间线	A rocprof / omniperf / omnitrace skill: occupancy, LDS conflicts, HBM BW	一个 rocprof / omniperf / omnitrace skill： occupancy、 LDS 冲突、 HBM 带宽
Shape-aware router	Dispatch small/medium/large to specialized kernels	把小 / 中 / 大分派到专用 kernel	Same idea; our roofline-first principle picks the regime boundaries	同一个想法；我们的 roofline 优先原则来定区间边界

The honest gap诚实的差距

The report's own ablation says it plainly: the loop alone (Humanize) got 3.71× on the Indexer, but KernelWiki and ncu-report-skill more than doubled that to 8.58×. The agent scaffolding is the easy 40%; the domain knowledge and profiler interpretation are the hard 60%. For AMD that 60% is thinner — CK and AITER have far fewer years of production PRs than CUTLASS, and rocprof/omniperf interpretation is less codified than Nsight. That gap is software maturity, not physics, which means it is closable by exactly the kind of systematic agent search this report demonstrates. Building the AMD KernelWiki and the rocprof skill is the project.

报告自己的消融说得很直白：仅靠循环（Humanize）在 Indexer 上拿到 3.71×，但 KernelWiki 和 ncu-report-skill 把它翻了一倍多到 8.58×。 agent 脚手架是容易的那 40%；领域知识和 profiler 解读是难的那 60%。对 AMD 来说这 60% 更薄 —— CK 和 AITER 的生产 PR 年头远少于 CUTLASS， rocprof / omniperf 的解读也不像 Nsight 那样被系统化沉淀。这个差距是软件成熟度，不是物理，也就意味着它正好能被这篇报告所展示的那种系统化 agent 搜索来填平。把 AMD 版 KernelWiki 和 rocprof skill 建起来，这件事本身就是项目。

§ 09 · EpilogueThe loop, not the model

§ 09 · 尾声是循环，不是模型

The throughline of this report is the same one Polar made from the training side: the leverage is in the loop and its environment, not in a cleverer single shot. KDA is a frontier coding agent placed inside a real repo, handed a knowledge skill and a profiler skill, gated by an independent verifier, and left to run. That composition — model + tools + workflow — is what produced kernels that beat a strong baseline on three of five Blackwell operators.

这篇报告的主线，和 Polar 从训练侧讲出来的是同一条：杠杆在于循环及其环境，而不在于一次更聪明的单发。 KDA 是一个被放进真实仓库的前沿编程 agent，配了一个知识 skill 和一个 profiler skill，由一个独立验证器把门，然后放手让它跑。正是这种组合 —— 模型 + 工具 + 工作流 —— 造出了在五个 Blackwell 算子里三个上跑赢强 baseline 的 kernel。

For us the takeaway is direct: we already have the agent scaffolding (kimi-cli, agent-teams) and the hardware (MI300X / MI355X). What we don't yet have is an AMD KernelWiki and a rocprof skill with the same care this team put into Nsight. Build those two, wire them into a verifier-gated loop, and point it at one kernel — GEMM or attention — until it hits 95% of CK hand-written. Win one kernel completely, then generalize. That's the plan this report makes concrete.

对我们来说收获很直接：我们已经有 agent 脚手架（kimi-cli、 agent-teams）和硬件（MI300X / MI355X）。我们还没有的，是一个 AMD 版 KernelWiki，以及一个像这个团队对待 Nsight 那样用心打磨的 rocprof skill。把这两个建起来，接进一个有验证器把门的循环，然后对准一个 kernel —— GEMM 或 attention —— 一直干到 CK 手写水平的 95%。把一个 kernel 彻底拿下，再去泛化。这正是这篇报告替我们具象化出来的计划。

§ 10 · SourcesReferences & citation

§ 10 · 来源参考文献与引用

This reading is based on the technical report and the released repository. If you cite this writeup, cite the report:

这篇精读基于技术报告与开源仓库。如果你要引用本文，请引用原报告：

@techreport{kernelmafia2026agentic, title = {Developing GPU Kernels with Agentic Loops}, author = {{HAN Lab Kernel Mafia}}, institution = {MIT HAN Lab}, year = {2026}, note = {MLSys-2026 NVIDIA Blackwell Kernel Competition (FlashInfer); Kernel Design Agents (KDA)}, url = {https://github.com/mit-han-lab/kernel-design-agents} }

Developing GPU Kernels with Agentic Loops. HAN Lab Kernel Mafia (MIT). MLSys-2026 NVIDIA Blackwell Kernel Competition (FlashInfer). Report + code: github.com/mit-han-lab/mlsys2026-flashinfer-contest — the report this entry reads.
Developing GPU Kernels with Agentic Loops. HAN Lab Kernel Mafia（MIT）。 MLSys-2026 NVIDIA Blackwell Kernel 竞赛（FlashInfer）。报告 + 代码： github.com/mit-han-lab/mlsys2026-flashinfer-contest —— 本文精读的报告。
kernel-design-agents (KDA). The released agent workflow — CLAUDE.md, agent-flow, and prompt templates. github.com/mit-han-lab/kernel-design-agents
kernel-design-agents (KDA). 开源的 agent 工作流 —— CLAUDE.md、 agent-flow、 prompt 模板。 github.com/mit-han-lab/kernel-design-agents
KernelWiki. The production-kernel knowledge base skill (D. Zou). github.com/DongyunZou/KernelWiki
KernelWiki. 生产级 kernel 知识库 skill（D. Zou）。 github.com/DongyunZou/KernelWiki
ncu-report-skill. The Nsight Compute profiling-analysis skill (D. Zou). github.com/DongyunZou/ncu-report-skill
ncu-report-skill. Nsight Compute 性能分析 skill（D. Zou）。 github.com/DongyunZou/ncu-report-skill
K-Search. Cao, Mao, Gonzalez, Stoica — LLM kernel generation via a co-evolving intrinsic world model. The baseline KDA is ablated against. arXiv 2026.
K-Search. Cao、 Mao、 Gonzalez、 Stoica —— 通过协同进化的内在世界模型做 LLM kernel 生成。 KDA 消融时对照的 baseline。 arXiv 2026。
FlashInfer-Bench & the starter kit. The benchmark and workloads the contest evaluates on. github.com/flashinfer-ai/flashinfer-bench-starter-kit
FlashInfer-Bench 与 starter kit. 竞赛评估所用的基准与 workload。 github.com/flashinfer-ai/flashinfer-bench-starter-kit