🔶 TIRx Harness：面向 Agentic GPU 编程的开放编译器 Harness

2026 年 9 月 29 日 • MLC Community

**TL;DR：**我们构建了 TIRx Harness：一个编译器 Harness，结合精简且稳定的编译器基础、知识库、工具和基准测试服务器，帮助 Agent 开发正确且高性能的 GPU Kernel。在已评估的工作负载中，Kimi Delta Attention（KDA）Kernel 相比 FlashKDA 的前向计算实现几何平均加速 2.94×，相比 Flash Linear Attention（FLA）的反向计算实现几何平均加速 6.84×。

开发高性能 GPU Kernel 是一个迭代式工程过程：探索实现方案、借助硬件知识和既有代码，并通过调试与测量持续改进结果。编程 Agent 已能自动完成这一过程中的重要部分，但在我们针对 KDA 等工作负载的实验中，它们的大量精力都花在优化本身之外的工作上：预测 Kernel 代码会如何被降级为硬件行为、诊断可能依赖时序或受到并发 GPU 活动干扰的测量失败，以及寻找与特定任务相关的实现或优化策略。因此，Token 效率急剧下降：Agent 的预算被用于消除优化周边的不确定性，而非探索优化本身，导致本可行的优化方案可能永远无法在实际可承受的搜索预算内被发现。

这些问题并不只是通过让 Agent 尝试更多变体就能解决的搜索问题。它们指向开发环境中缺失的组成部分：可预测的编程基础、能够超越单次执行进行推理的工具、可复用的实现知识，以及受控的评估机制。

我们将这些部分整体视为一个编译器 Harness：即围绕编译器基础、面向 Agent 的环境，其中包括领域特定分析与诊断、IR 特定知识语料库，以及基准测试和评估工具。随着 Agentic 编程将更多工作从编写代码转向塑造 Agent 的运行环境，我们预计这一层会变得愈发重要。这一框架受到多项同期 AI 导向系统设计工作的启发。其中，CAKE 明确提出了编译器与 Agent 的协同设计，而 OpenAI 的 Jalapeño 则体现了相关的全栈原则：将编程目标本身设计得足够清晰、可预测，从而让 AI 能够有效进行优化。

TIRx Harness 是这一理念在 GPU Kernel 开发中的一个实例。对我们而言，它始于一个核心问题：如何设计一个对于 Agentic 编程而言精简、稳定且可扩展的编译器基础？在这一基础周围，TIRx Harness 结合了知识库、正确性与性能工具以及基准测试服务器，为 Agent 提供更可靠的 Kernel 优化环境。在已评估的 KDA、MiniMax Sparse Attention（MSA）、Multi-head Latent Attention（MLA）和 Video Sparse Attention（VSA）系列中，报告的各系列几何平均加速比范围为 1.33× 至 6.84×；其中，KDA 前向计算相较 FlashKDA 达到 2.94×，反向计算相较 FLA 达到 6.84×。

围绕可靠迭代构建 TIRx Harness

这些组件解决的是同一个优化循环中不同来源的不确定性。TIRx Harness 提供面向编译器的环境，而 Agent 工作流则利用这些能力来组织搜索。我们围绕在 Agent 驱动 Kernel 优化中反复观察到的失败模式来设计每个组件。

TIRx Foundation：让编译器行为可预测

编程基础决定了 Agent 能否将它写出的内容直接关联到 GPU 实际执行的内容。高级抽象对于生产力很有价值，但在我们的实验中，它们可能使优化更难以推理。Agent 可能知道自己想尝试的执行策略，却不知道何种源程序会让编译器将其降级为该策略。漫长的编译路径还会在调试时引入另一种不确定性：当出现数值错误时，Agent 必须同时考虑自己编写的 Kernel 以及编译器对其施加的变换。

我们通过两种方式解决这些不确定性。首先，我们让 TIRx 接近硬件。其轻量级、PTX 层面的编程接口使 Agent 能够更直接地表达指令和执行策略，而不是依赖高级抽象被降级为其所需的行为。

其次，TIRx 有意保持精简。除硬件接口以外，它尽可能少地引入独立语义，从而让 IR 和编译路径保持小巧且稳定。在这样薄的抽象层下，PTX ISA 是指令语义的主要事实来源。这也使系统更易于扩展：暴露新的 PTX 能力通常只需要对编程接口进行轻量级补充，而不需要新增高级抽象或扩大编译器栈。我们仍保留循环和条件判断等基本结构，以便 Harness 的其余部分能够分析程序。

轻量级编程接口将 Kernel 操作连接到生成的代码。选择一行源代码即可查看其映射关系。

即使编程基础更简单，故障仍可能难以诊断。一个 Kernel 可能在多次运行中通过测试，随后却因间歇性的同步错误或数据竞争而失败，使 Agent 获得的信息几乎只有通过/失败信号。这促使我们为 Harness 建立了一项更广泛的原则：工具应当提供有关程序行为的有用反馈，而不仅仅是测试结果。

我们构建了同步和数据竞争分析工具，以直接检查并发行为。当 GPU 执行不可用时，数值模拟可通过在无 GPU 情况下模拟数值行为，提供另一种反馈来源。对于性能诊断，我们集成了包括 NCU 和 IKET 在内的既有分析工具。这些工具共同为 Agent 提供更多信息，帮助其决定下一步应调查和修改什么。

编译器分析会返回数值、同步和数据竞争诊断结果。选择一个组件即可查看其反馈。

知识库：复用已知思路，发现新思路

可预测的编程基础和更好的诊断能力，仍无法为 Agent 提供开发新 Kernel 所需的实现与优化知识。在早期运行中，Agent 花费了大量精力来学习编程模型，并重新发现有用的优化模式。

为避免在每次运行中重复学习这些思路，自然的方法是总结此前运行的经验。然而在我们的实验中，由 Agent 编写的摘要往往会过度拟合刚刚调试过的案例，并丢失有价值的实现上下文。因此，我们选择保留具体实现。

Kernel zoo 是这一知识库的核心。它包含超过 60 个 TIRx Kernel，为 Agent 提供具体起点和可复用的优化模式。但该 zoo 只能捕捉已在现有 Kernel 中出现的思路。PTX 文档则通过展示能够开启新优化方向的硬件能力来补充它。Kernel zoo 有助于复用已知思路；ISA 则有助于发现新思路。

Kernel zoo 提供可复用的实现；硬件与 ISA 文档揭示新的可能性。选择一个知识来源即可探索。

基准测试服务器：让测量结果可比较

正确性测试、耗时数据和性能分析决定候选方案是否被接受，以及 Agent 接下来会尝试哪种优化。因此，误导性的测量会影响许多次尝试，而非单次尝试。当 Agent 直接从自身环境发起 GPU 工作时，Harness 的两个目标会让可靠测量更难维持。第一是规模：我们希望许多 Agent 并发运行，但一旦它们共享同一块 GPU，耗时变化可能反映的是另一 Agent 的活动，而非候选方案自身的改进或退化；Agent 便无法将性能变化归因于自己的编辑。第二是平台覆盖：我们希望面向数据中心 GPU 之外的硬件，包括 Thor 等边缘设备，但这些平台并不适合作为编程 Agent 的宿主机，因此将 Agent 绑定到其优化目标机器会限制它可访问的硬件范围。

这两点都指向同一种设计：将评估与 Agent 分离。因此，TIRx Harness 采用了围绕 KCoral 构建的远程评估架构；KCoral 是我们的基准测试服务器（博客即将发布）。KCoral 管理 GPU；无论 Agent 何时需要测试正确性、评测 Kernel，还是收集性能分析数据，它都会向 KCoral 发送请求，而不是在本地启动工作。KCoral 集中调度所有请求，确保一次执行不会影响另一次执行，并在请求完成后返回测量结果与诊断产物。Agent 能够为每次测量获得受控且可信的环境，而 Harness 则可通过将新目标接入 KCoral 来扩展支持目标，无须让每台 Agent 宿主机都运行在这些硬件上。

KCoral 将 Agent 开发与 GPU 执行分离，并集中调度测量任务。选择一个组件即可了解其作用。

TIRx Harness 实现了什么

上述组件之所以重要，是因为它们改变了 Agent 在优化期间能够做的事情。除生成更多候选方案外，Harness 还能帮助 Agent 发现此前未曾考虑过的硬件能力、在 Kernel 之间迁移策略，并通过不止于通过/失败的反馈诊断不安全的候选方案。以下轨迹展示了这些能力在实践中的作用；下一节将评估最终 Kernel 的性能。

发现新方向

硬件文档能够扩展 Agent 的搜索空间，而不只是回答语法问题。在一条 MSA 稀疏预填充轨迹中，Agent 在 PTX ISA 中发现了 tcgen05.mma 输出 Lane 掩码，并将其转化为一项有用优化。该掩码并未减少 MMA 工作量；相反，它提高了观测到的 GPU 频率，带来 2–3% 的性能提升。

TIRx 接近硬件的编程接口也让 Agent 能够直接处理硬件资源之间的权衡。在 KDA 反向计算中，Agent 将转置、对角缩放和逐行点积转移到 Tensor Core 路径。对于点积，它计算完整的 64×64 矩阵后仅保留对角线，刻意增加 Tensor Core 工作量，以便从繁忙的计算 Warp 中移除共享内存部分结果和归约工作。这一看似违反直觉的权衡——在利用不足的硬件路径上执行更多算术工作，以缓解瓶颈路径——使延迟降低约 15%。

迁移策略

Kernel zoo 为 Agent 提供了可被适配而非照搬的具体策略。在 MSA 中，Agent 适配了 FlashAttention-4 的原生/软件混合指数计算策略，利用普通 FMA 工作来缓解原生 exp 路径压力，并将延迟改善约 2.9%。

在 KDA 前向计算中，Agent 迁移了 Gated DeltaNet（GDN）的分层逆分解方法，以小型块逆和基于 MMA 的合并操作替代串行的 32×32 递推，将延迟降低约 17%。这些案例表明了我们为何保留具体实现：它们揭示了可跨不同 Kernel 迁移的分解方法和数据流。

诊断不安全的候选方案

分析工具能够捕捉常规测试可能遗漏的正确性问题。在一条 KDA 轨迹中，一个候选方案通过了基准测试的正确性测试，但同步分析仍发现了潜在 Bug：一个被复用的 mbarrier 可能在较晚的消费者完成等待前一代状态之前，就推进至下一代。Agent 因而可以在这一依赖时序的问题实际在执行中暴露前，修复同步协议。

这些轨迹共同展示了 Harness 改变优化循环的三种方式：ISA 扩展搜索空间，Kernel zoo 提供可迁移策略，分析工具则将不安全的候选方案转化为具体诊断。随后，基准测试服务器提供受控的性能反馈，帮助决定保留哪些方案。

Agent 借助 Harness 取得了什么成果

评估旨在回答：上述能力是否能够转化为高性能 Kernel。我们评估了 KDA、MSA、MLA 和 VSA。本文报告的所有实验均在 NVIDIA Blackwell GPU 上运行，采用 Humanize 2 的 flame chase 工作流，优化期间禁用了 Web 访问。基线实现使用的是 2026 年 9 月 25 日至 26 日精选扫描中记录的版本。每个系列均与其各自优化后的参考实现进行比较，使用 GPU Kernel 时间而非端到端应用延迟。

在这些工作负载中，使用 TIRx Harness 的 Agent 生成了可与参考实现竞争、且通常更快的 Kernel。报告的各系列几何平均加速比范围为 1.33× 至 6.84×。图中展示了已评估配置的最小值至最大值范围，以及每个系列的几何平均值。

Agentic 编程领域正在快速演进；随着 Agent、工作流、编译器和参考实现的改进，基准结果可能在数天尺度上发生变化。因此，我们的目标并不是建立一种永久性的方案排名，而是展示 TIRx Harness 能够实现什么。由于不同系列使用的参考实现不同，这些加速比应在各自系列内解读，而不应作为跨系列排名。这些 Kernel 是

Sep 29, 2026 • MLC Community

TL;DR: We built TIRx Harness, a compiler harness combining a minimal stable compiler foundation, a knowledge base, tools, and a benchmark server to help agents develop correct, fast GPU kernels. On evaluated workloads, Kimi Delta Attention (KDA) kernels achieve geometric-mean speedups of 2.94× over FlashKDA (forward) and 6.84× over Flash Linear Attention (FLA) (backward).

Developing high-performance GPU kernels is an iterative engineering process: exploring implementations, drawing on hardware knowledge and existing code, and refining the result through debugging and measurement. Coding agents can already automate meaningful parts of this process, but in our experiments on workloads such as KDA, much of their effort went into work around the optimization itself: predicting how kernel code would be lowered to hardware behavior, diagnosing failures that could be timing-dependent or measurements distorted by concurrent GPU activity, and finding an implementation or optimization strategy relevant to the specific task. As a result, token efficiency drops sharply: the agent spends its budget resolving uncertainty around the optimization rather than exploring the optimization itself, so otherwise viable optimizations may never be reached within a practical search budget.

These were not simply search problems that could be solved by asking the agent to try more variants. They pointed to missing pieces in the development environment: a predictable programming foundation, tools for reasoning beyond a single execution, reusable implementation knowledge, and controlled evaluation.

We view these pieces together as a compiler harness: the agent-facing environment around a compiler foundation, including domain-specific analyses and diagnostics, an IR-specific knowledge corpus, and benchmarking and evaluation tooling. As agentic programming shifts more of the work from writing code to shaping the environment in which agents operate, we expect this layer to become increasingly important. This framing is inspired by several concurrent efforts in AI-oriented system design. Among them, CAKE makes compiler–agent co-design explicit, while OpenAI’s Jalapeño reflects a related full-stack principle: designing the programming target itself to be clear and predictable enough for AI to optimize effectively.

TIRx Harness is one instance of this idea for GPU kernel development. For us, that starts with a central question: how can we design a compiler foundation that is minimal, stable, and extensible for agentic programming? Around this foundation, TIRx Harness combines a knowledge base, correctness and performance tools, and a benchmark server, giving agents a more reliable environment for kernel optimization. Across the evaluated KDA, MiniMax Sparse Attention (MSA), Multi-head Latent Attention (MLA), and Video Sparse Attention (VSA) families, reported family-level geometric-mean speedups range from 1.33× to 6.84×; KDA forward and backward reach 2.94× over FlashKDA and 6.84× over FLA, respectively.

Building TIRx Harness around reliable iteration

These components address different sources of uncertainty in the same optimization loop. TIRx Harness provides the compiler-facing environment, while agent workflows use these capabilities to organize their search. We designed each component around a recurring failure mode we observed during agent-driven kernel optimization.

TIRx Foundation: Make the Compiler Predictable

The programming foundation determines how directly an agent can connect what it writes to what the GPU executes. High-level abstractions are valuable for productivity, but in our experiments they could make optimization harder to reason about. An agent might know the execution strategy it wants to try without knowing what source program will make the compiler lower to that strategy. A long compilation path adds another source of uncertainty during debugging: when a numerical failure appears, the agent must consider both the kernel it wrote and the compiler transformations applied to it.

We address these sources of uncertainty in two ways. First, we keep TIRx close to the hardware. Its thin PTX-level programming surface lets agents express instructions and execution strategies more directly, rather than relying on a high-level abstraction to lower into the behavior they want.

Second, TIRx is intentionally minimal. It introduces as little independent semantics as possible beyond the hardware interface, keeping the IR and compilation path small and stable. With such a thin abstraction layer, the PTX ISA serves as the primary source of truth for instruction semantics. This also makes the system easier to extend: exposing a new PTX capability usually requires only a lightweight addition to the programming surface, rather than a new high-level abstraction or a larger compiler stack. We still retain basic structure such as loops and conditionals so the rest of the harness can analyze the program.

A thin programming surface connects kernel operations to generated code. Select a source line to inspect its mapping.

Even with a simpler programming foundation, failures can still be hard to diagnose. A kernel may pass many runs and then fail because of an intermittent synchronization bug or data race, leaving the agent with little more than a pass/fail signal. This motivated a broader principle for the harness: tools should provide useful feedback about program behavior, not just test outcomes.

We built synchronization and data-race analyses to inspect concurrency behavior directly. When GPU execution is unavailable, numerical simulation provides another source of feedback by simulating numerical behavior without a GPU. For performance diagnosis, we integrate existing profiling tools including NCU and IKET. Together, these tools give agents more information to decide what to investigate and change next.

Compiler analyses return numerical, synchronization, and data-race diagnostics. Select a component to explore its feedback.

Knowledge base: reuse known ideas, discover new ones

A predictable programming foundation and better diagnostics still do not give agents the implementation and optimization knowledge needed for a new kernel. In early runs, agents spent substantial effort learning the programming model and rediscovering useful optimization patterns.

To avoid relearning these ideas in every run, a natural approach is to summarize lessons from previous runs. In our experiments, however, agent-written summaries often overfit the case that had just been debugged and lose useful implementation context. We therefore preserve concrete implementations instead.

The kernel zoo is the core of this knowledge base. With more than 60 TIRx kernels, it gives agents concrete starting points and reusable optimization patterns. But the zoo only captures ideas that have already appeared in existing kernels. PTX documentation complements it by exposing hardware capabilities that can open new optimization directions. The zoo helps reuse known ideas; the ISA helps discover new ones.

The kernel zoo supplies reusable implementations; hardware and ISA documentation expose new possibilities. Select a knowledge source to explore.

Benchmark server: make measurements comparable

Correctness tests, timings, and profiles decide whether a candidate is accepted and which optimization the agent tries next, so a misleading measurement affects many attempts rather than one. Two goals for the harness made such measurements harder to keep reliable when the agent launches the GPU work directly from its own environment. The first is scale: we want many agents running concurrently, but once they share a GPU, a timing change may reflect another agent’s activity rather than the candidate’s own improvement or regression, and the agent can no longer attribute a performance change to its edit. The second is platform coverage: we want to target hardware beyond datacenter GPUs, including edge devices such as Thor, but these platforms are not convenient hosts for a coding agent, so tying the agent to the machine it optimizes for would limit the hardware it can reach.

Both point to the same design: separate evaluation from the agent. TIRx Harness therefore uses a remote evaluation architecture built around KCoral, our benchmark server (blog coming soon). KCoral owns the GPUs, and whenever the agent needs one, whether to test correctness, benchmark a kernel, or collect a profile, it sends a request to KCoral instead of launching work locally. KCoral schedules all requests centrally so that the execution of one never affects another, and returns the measurements and diagnostic artifacts when the request completes. The agent gets a controlled, trusted environment for every measurement, and the harness gains new targets by attaching them to KCoral rather than by making every agent host run on them.

KCoral separates agent development from GPU execution and centrally schedules measurements. Select a component to explore its role.

What TIRx Harness enables

The components above matter because they change what an agent can do during optimization. Beyond generating more candidates, the harness can help agents discover hardware capabilities they had not considered, transfer strategies across kernels, and diagnose unsafe candidates with more than pass/fail feedback. The traces below illustrate these roles in practice; the next section evaluates the resulting kernel performance.

Discovering new directions

Hardware documentation can expand the agent’s search space, not just answer syntax questions. In one MSA sparse-prefill trace, the agent found tcgen05.mma output-lane masking in the PTX ISA and turned it into a useful optimization. The mask did not reduce the MMA work; instead, it increased the observed GPU frequency, yielding a 2–3% gain.

TIRx’s hardware-close programming surface also lets agents act directly on hardware-resource tradeoffs. In KDA backward, the agent moved transposes, diagonal scaling, and a row-wise dot product onto the Tensor Core path. For the dot product, it computed a full 64×64 matrix and kept only the diagonal, deliberately spending extra Tensor Core work to remove shared-memory partials and reduction work from busy compute warps. This counterintuitive tradeoff—doing more arithmetic on an underused hardware path to relieve the bottlenecked one—improved latency by about 15%.

Transferring strategies

The kernel zoo gives agents concrete strategies that can be adapted rather than copied literally. In MSA, the agent adapted FlashAttention-4’s mixed native/software exponential strategy, using ordinary FMA work to relieve the native exp path and improving latency by about 2.9%.

In KDA forward, the agent transferred Gated DeltaNet (GDN)’s hierarchical inverse decomposition, replacing a serial 32×32 recurrence with small block inverses and MMA-based merges and reducing latency by about 17%. These cases show why we preserve concrete implementations: they expose decompositions and dataflows that can transfer across different kernels.

Diagnosing unsafe candidates

Analysis tools can catch correctness issues that ordinary testing may miss. In one KDA trace, a candidate passed the benchmark’s correctness tests, but synchronization analysis still identified a latent bug: a reused mbarrier could advance to its next generation before a late consumer finished waiting on the previous one. The agent could then repair the synchronization protocol before this timing-dependent issue surfaced in execution.

Together, these traces show three ways the harness changes the optimization loop: the ISA expands the search space, the kernel zoo provides transferable strategies, and analysis tools turn unsafe candidates into concrete diagnoses. The benchmark server then provides controlled performance feedback for deciding what to keep.

What agents achieved with the harness

The evaluation asks whether the capabilities above translate into fast kernels. We evaluate KDA, MSA, MLA, and VSA. All experiments reported here were run on NVIDIA Blackwell GPUs using Humanize 2’s flame chase workflow, with web access disabled during optimization. The baseline implementations use the versions recorded in the September 25–26, 2026 curated sweep. Each family is compared against its own optimized reference implementation using GPU kernel time rather than end-to-end application latency.

Across these workloads, agents using TIRx Harness produced kernels that are competitive with—and often faster than—the reference implementations. Reported family-level geometric-mean speedups range from 1.33× to 6.84×. The figure shows the min–max range across evaluated configurations together with each family’s geometric mean.

The agentic programming landscape is evolving quickly, and benchmark results can change on the scale of days as agents, workflows, compilers, and reference implementations improve. Our goal is therefore not to establish a permanent ranking of approaches, but to demonstrate what TIRx Harness can enable. Because the reference differs across families, the speedups should be interpreted within each family rather than as a cross-family ranking. The kernels are available here.

Closing the loop

Kernel evolution can improve the harness as well as the kernel. When a run exposes a tool bug, unsupported behavior, or a missing check, its trace provides a concrete reproducer that can be used to improve the analysis tools. Successful kernels follow a different path: they enter the kernel zoo as better starting points and reusable optimization strategies for future runs.

This is how we think about self-improvement in TIRx Harness: useful outcomes from one evolution run become part of the environment available to the next. Rather than asking the agent to summarize what it learned, we preserve concrete artifacts—a validated tool improvement or a measured kernel implementation—that later runs can directly build on.

What’s next

We see three directions for the next stage of TIRx Harness. First, we want to evaluate how quickly the harness can adapt to new GPU architectures, including how rapidly agents can make use of new instructions and hardware features. Second, we want to push TIRx further toward the hardware, exposing lower-level mechanisms that give agents more direct access to performance opportunities. Third, we want to extend our correctness tooling to megakernels, where validating synchronization, memory accesses, and numerical behavior becomes harder as more computation is fused into a single program.

Acknowledgments

We thank the NVIDIA CAKE team, Kernel Design Agent (KDA) team and SOL-ExecBench team for helpful discussions and feedback throughout this work.

Getting started

To try TIRx Harness, start with the GitHub repository, which covers installation, running kernel optimization tasks, and using the harness’s analysis and remote execution tools. Our book, Agentic GPU Programming for MLSys introduces the main elements of agentic gpu programming and compiler harness.