TileRT 技术谱系：ThunderKittens、TileLang、MPK 与 Event Tensor，常驻 kernel 如何走向动态 megakernel

为什么 TileRT 需要一篇技术谱系补篇#

TileRT 主文 已经讲了它自己的系统设计：latency-first、execution boundary、persistent Engine Kernel、tile-level pipeline、warp/block/GPU specialization、vLLM PD 分离和 InferenceX 评测。Talk PDF 的后半部分还有一大块 related work：ThunderKittens、TileLang、MPK、Event Tensor。它们不是背景八卦，而是在回答同一个问题的不同层级：低 batch decode 太快之后，kernel 边界、同步、launch 和中间结果往返变成了主要瓶颈，GPU 程序应该怎样重写？

这条路线可以从小到大排列：

1. ThunderKittens：单 kernel 怎么用 tile、warp 专用化和 persistent grid 写得又快又不至于不可维护；
2. TileLang：如何把 tile 程序变成可编译、可调优、可跨硬件复用的 DSL；
3. MPK：如何把多个算子自动 mega-kernel 化，在一个长驻 kernel 里用 tGraph 和 in-kernel runtime 调度；
4. Event Tensor：当 shape、routing、sparsity 都到运行期才知道时，动态 megakernel 该如何表达依赖；
5. TileRT：把这些思想用在特定大模型 decode 上，为极致单用户延迟做定制化系统。

起点：kernel-per-operator 模型的边界税#

传统深度学习执行模型是 operator-per-kernel：一个算子一个 GPU kernel，host 按图顺序 launch，kernel 内部完成 load → compute → store，下一个 kernel 再从 global memory 读回中间结果。这套模型在训练和大 batch 推理里很好，因为每个 kernel 足够大，launch 和同步开销能被计算摊薄。

低延迟 decode 改变了这个前提。batch=1 或小 batch 时，单个算子的计算时间可能只有几微秒，甚至比一次 kernel launch 和同步还短。Talk 里 Event Tensor 部分给出一个典型量级：GPU kernel launch 通常约 5–10 μs，而某些轻量 kernel 本身可能只有约 2 μs。此时优化单个 GEMM 并不能解决端到端问题，因为真正慢的是边界。

边界税主要有四种：

• CPU launch：host 每次发起 kernel，CPU/GPU 控制路径进入关键链路；
• kernel barrier：kernel 边界强制 wave/CTA 收敛，无法跨算子做细粒度流水；
• intermediate store/load：中间结果写回 HBM，下个算子再读，片上局部性被清空；
• communication synchronization：多 GPU 之间的通信被粗粒度 layer/operator 边界切开，难以和计算重叠。

常驻 kernel / megakernel 化的目标就是把这些边界尽量拿掉。

ThunderKittens：把高性能 kernel 的手艺封装成 tile 原语#

ThunderKittens 的出发点很现实：FlashAttention 级别的 kernel 性能很好，但手写 CUDA 需要同时管理计算、memory layout、data movement、warp 协作、寄存器排列、shared memory bank conflict、异步流水线和 grid scheduling。高层框架好写但性能不足，手写 CUDA 性能强但复杂度高。ThunderKittens 想在两者之间找一个小而硬的抽象层。

它的核心抽象是 warp-level tile。普通 CUDA 初学者常把线程当成基本单位：每个 thread 算一个元素或几个元素。ThunderKittens 则把 tile 当成基本单位：一个 warp 或 warp group 协同持有一个寄存器 tile、一个 shared-memory tile，程序员写的是 tile load、tile compute、tile store，而不是每个线程对应哪个矩阵元素。

这带来两个好处：

1. 隐藏 per-thread data mapping 和 register permutation 的细节；
2. 暴露真正重要的硬件层级：register tile、shared-memory tile、Tensor Core tile、TMA/cp.async data movement。

从朴素 GEMM 到 Hopper-style GEMM#

Talk 用 GEMM 解释 ThunderKittens 的必要性。朴素 GEMM 让每个线程计算输出矩阵的一个元素；shared-memory tiled GEMM 让一个 CTA 协作计算一个输出 tile，把 A/B 的子块搬进 shared memory 复用；WMMA GEMM 让一个 warp 协作调用 Tensor Core，但未必显式优化 shared memory 复用。

Ampere-style GEMM 开始显式体现三段数据流：

1
HBM → Shared Memory: cp.async
2
Shared Memory → Registers: ldmatrix
3
Tile Computation: mma.sync

Hopper-style GEMM 又引入 WGMMA 和 TMA：warp group 协同执行更大的矩阵乘，TMA 负责更高效的异步张量搬运。此时程序员真正要编排的已经不是「线程 i 算元素 j」，而是「哪些 warp 负责搬，哪些 warp 负责算，tile 如何在 HBM/shared memory/register 之间流动」。ThunderKittens 正是把这个编排对象提升成语言原语。

Warp group asynchronous pipeline：专用化的第一层#

ThunderKittens 明确采用 producer-consumer pipeline：producer warps 负责数据搬运，consumer warps 负责 Tensor Core computation，两者通过异步流水重叠。CUTLASS 3.x 在 Hopper GEMM 中也采用类似 warp-specialized design。

这和本站 FlashAttention-3 讲过的思路一致：注意力 kernel 不再是 load 完再 compute，而是让数据搬运、softmax、GEMM 交错推进。关键收益不只是少等 HBM，而是让 warp scheduler 有真正可重叠的工作。若所有 warp 都执行相似指令流，它们很容易重新 phase-aligned，在同一个 barrier 一起停住；专用化让不同 warp 处于不同阶段，延迟更容易被藏起来。

ThunderKittens 的实验结论也说明抽象不必牺牲性能：单个 TK GEMM 可与优化过的 cuBLAS/CUTLASS 竞争；FlashAttention forward 接近 FlashAttention-3，backward 可快 10–40%，原因之一是 shared-memory stall cycles 少约 85%，bank conflict 控制更好。

Grid-level scheduling：persistent grid 与 cache-aware block scheduling#

ThunderKittens 还有一层 grid-level scheduling。Persistent grid 指启动常驻 thread blocks 到 SM 上，让每个 block 连续处理多个 output tiles，摊销 CTA launch 和调度开销。Cache-aware block scheduling 则按数据复用模式安排 block/tile 顺序，提高 L2 cache locality，减少不必要的 HBM traffic。

这对 TileRT 很重要，因为 TileRT 的 Engine Kernel 正是把 persistent 的思想从单 kernel 推到整张 decode 图：不是一个 CTA 多做几个 tile，而是整个 decode graph 长驻 GPU；不是只在一个 GEMM 内部做 producer-consumer，而是在层间、算子间、通信间做持续流水。

TileLang：把 tile 编程变成编译器问题#

ThunderKittens 让高性能 kernel 更容易写，但仍然偏向「人写出结构，库帮你处理细节」。TileLang 往前推进一步：把 tile-level programming 做成 DSL，让分块、存储放置、warp 分工、layout、software pipelining 这些选择进入编译器和自动调优空间。

TileLang 暴露的典型问题包括：

• tile size 多大；
• Q/K/V 这类 tile 应该放 shared memory 还是 registers；
• pipeline stages 取 1、2、3 还是更多；
• 哪些 warps 负责哪些 tile；
• register layout 如何分布到线程；
• data movement 如何与 compute overlap。

它提供的程序原语也很直接：T.alloc_shared 明确把 tile 放进 shared memory，T.alloc_fragment 放进 registers，T.copy 表达 tile-level data movement，T.gemm 表达 tile-level matrix computation。

Tile Recommendation：用成本模型搜索硬件感知配置#

TileLang 的第一项关键设计是 Tile Recommendation。它会探索一组 hardware-aware candidates，例如 block_M、block_N、Q 的 memory placement、pipeline stages、warp partitioning 等，然后用静态成本模型估计执行时间：

这里 $i$ 是 memory hierarchy level，例如 HBM、L2、L1；$j$ 是 compute unit，例如 Tensor Core、CUDA Core、SFU；$t_\text{intrinsic}$ 是具体指令开销。这个模型不需要精确模拟所有细节，但它把优化问题从「经验调参」变成「在硬件约束下比较候选」。

对 TileRT 来说，这种能力是底层基础。Engine Kernel 的 tile shape、片上驻留、pipeline depth、warp group 划分，如果全靠人手写，会很难扩展；如果能被 DSL/编译器表达，就有机会迁移到更多模型和更多硬件。

Tile Inference：layout 不是注释，而是正确性和性能#

TileLang 的第二项关键设计是 Tile Inference。一个逻辑访问如 $A[i,j]$，最终必须映射到具体线程、具体 shared memory 地址、具体 register 位置。这个映射决定性能，也决定是否正确。Tensor Core GEMM 对 operand layout 有严格要求，accumulator layout 也决定输出元素在每个线程寄存器中的位置。

Talk 把 layout inference 分成三类：

这和 OpenAI Jalapeño 里的 Gluon/Linear Layouts 有相通处：现代 kernel 的难点已经不是写一个 for 循环，而是确保张量元素与硬件资源之间的映射既正确又高效。TileLang 用推断和搜索处理这个问题，Gluon 用布局代数处理这个问题，二者都说明 layout 已经是一等编程对象。

Pipeline Inference：从 ThunderKittens 继承 warp specialization#

TileLang 还做 pipeline inference，使用类似 ThunderKittens 的 warp specialization。也就是说，编译器不只决定 tile 放哪，还决定数据搬运和计算如何重叠，哪些 warp 充当 producer，哪些 warp 充当 consumer。

Talk 给出的实验数字显示，TileLang 在 NVIDIA H100 上相对 Triton 平均 3.02× 加速；对比 ThunderKittens，GEMM 可达到 0.99–1.11× 性能，同时代码复杂度降低 77%；FlashAttention 最多 1.10× 加速，并把实现从 185 行降到 66 行。这组数字的意义不是「TileLang 永远比 Triton 快 3 倍」，而是说明 tile DSL 能在可编程性和性能之间找到很强的折中。

MPK：从单 kernel 到图级 mega-kernel#

ThunderKittens 和 TileLang 主要解决高性能 fused kernel 怎么写。MPK 的问题更大：能不能把整个 tensor program 中多个 operator 编译成一个 mega-kernel？

传统 baseline 是 kernel-per-operator，问题包括 kernel barrier 阻止跨 operator software pipelining，compute/communication overlap 粒度太粗，kernel launch overhead 累积。MPK 的做法是 SM-level mega-kernel execution：编译器把 tensor program 编译成 SM-level task graph（tGraph），再由 in-kernel parallel runtime 在 GPU 内部动态调度任务。

这比普通 operator fusion 更激进。普通 fusion 常把几个相邻 elementwise 或小算子合并；MPK 想把更大的图拆成 tile task，让任务在 SM 之间以事件驱动方式推进。

tGraph generation：operator decomposition 与 dependency analysis#

MPK 编译器先做 operator decomposition：把每个 operator 的输出张量划分成互不重叠的区域，每个区域对应一个 block-level tile task。只要任务之间没有依赖，它们就可以在不同 SM 上并行执行。

随后做 dependency analysis：分析 producer task 和 consumer task 之间的依赖，用事件表示「某个 producer 完成后，哪些 consumer 可以启动」。如果每对 producer-consumer 都单独建事件，事件数量会爆炸；因此 MPK 做 event fusion：

• Successor-Set Fusion：后继任务集合相同的事件可以合并；
• Predecessor-Set Fusion：前驱任务集合相同的事件可以合并。

这一步很关键。Megakernel 不是把所有代码塞进一个巨大函数就结束了，内部必须有可控的依赖表达，否则同步开销会在 kernel 内部重新长出来。

In-kernel runtime：scheduler SM 与 worker SM#

MPK 的 runtime 在 kernel 内部运行。Talk 中给出 H100 示例：128 个 SM 作为 worker，4 个 SM 作为 scheduler。Scheduler 从 activated events 队列中取事件，把依赖任务分发给 worker；worker 执行任务，完成后通知对应事件。

它支持两种 task launch 策略：

MPK 的 hybrid strategy 是：attention 这类数据相关 operator 用 JIT，其他可预测 operator 用 AOT。这样既保留动态性，又避免所有任务都等运行期调度。

实验上，MPK 与 vLLM、SGLang 对比，在 5 个 0.6B–30B 模型、batch size 1–16 上，H100 最高 1.7×，B200 最高 1.5×；小模型和 low-batch workload 收益更大。这个结果和问题定义一致：越小、越低 batch，边界税占比越高，mega-kernel 化收益越明显。

Event Tensor：动态 megakernel 的依赖抽象#

MPK 仍然有一个限制：tGraph 基本是静态任务图。不同 input shape 可能需要不同 task graph 或重新编译；MoE routing 这类 data-dependent control flow 更难表达。LLM decode 恰恰越来越动态：变长上下文、稀疏注意力、Top-K token selection、MoE routing、MTP accept/reject 都让执行路径在运行期才确定。

Event Tensor 要解决的是：megakernel 如何面对动态 shape 和数据相关依赖。它提出的核心抽象是把事件本身张量化、符号化、动态化。

第一类动态是 shape dynamism。例如 batch size $B$ 是符号维度，不同运行时形状可以由同一 Event Tensor graph 表达，不必为每个 shape 生成一张完全不同的图。

第二类动态是 data-dependent dynamism。任务依赖由运行时值决定：某个 task 完成后更新哪个 event，某个 event 触发后启动哪些 task、多少 task，都可以由运行时数据决定。

一个简单例子：partial sum 到 final sum#

Talk 用 partial sum 例子解释 Event Tensor。第一阶段有 $n \times 4$ 个 partial tasks，每个 $(i,j)$ 任务完成后更新事件 $E_i$。事件 $E$ 的 shape 是 $n$，wait_count=4。第二阶段有 $n$ 个 final sum tasks，每个 final_sum(i) 等到 $E_i$ 收到 4 次更新后启动。

这个例子的价值在于：依赖不是单个标量 barrier，而是张量化事件。不同 $i$ 可以独立推进；某个 $i$ 的 4 个 partial 完成后，它的 final sum 不必等其他 $i$。把这个思想放进 LLM decode，意味着不同 token、不同专家、不同注意力块可以用更细粒度事件推进，而不是整层整卡一起 barrier。

Compiler-runtime co-design#

Event Tensor 把工作拆给编译器和 in-kernel runtime：

• 编译器负责编译每个 tile task 的代码，生成 symbolic dependency rules，编码运行时依赖如何解析；
• in-kernel runtime 在 megakernel 内部运行，依据运行时值解析具体依赖，判断哪些 task ready，并把 ready tasks dispatch 给 worker SM。

这比 CUDA Graphs 更靠近计算内部。CUDA Graphs 的依赖粒度仍然是 kernel 或 graph node，DAG 捕获后整体相对静态；Event Tensor 把事件语义下沉到 megakernel 内部的 tile task 粒度。它不是减少 host launch 的一个包装，而是试图让动态控制流在 GPU 内部自洽地推进。

TileRT 与四个工作的关系#

现在回到 TileRT。TileRT 的 persistent Engine Kernel 可以看作这条路线的生产级定制版本。它并没有声称自己是通用 MPK 或 Event Tensor，而是选择少数模型，把 decode graph 静态编译成常驻引擎，手工/编译器协同决定 tile pipeline、warp/block/GPU specialization 和通信重叠。

对应关系如下：

因此 TileRT 的局限也更容易理解：它很快，因为它是单件定制；它难扩展，因为通用自动化工具链还没有完全成熟。ThunderKittens/TileLang 解决了 kernel 级可编程性，MPK 刚开始解决图级自动化，Event Tensor 刚开始解决动态性。TileRT 跑在这条曲线的最前面，但后面的基础设施还在追它。

为什么 GPU specialization 是 warp specialization 的自然外推#

Talk 里 TileRT 主体部分有一页很重要：GPU 0 做 Top-K token selection，然后把 indices 广播给 GPU 1–7 做 sparse MLA；后续 MoE block 又重新用全 8 卡。这看起来像系统工程细节，其实是 specialization 思想的外推。

最早的 specialization 是 warp 级：producer warp 搬数据，consumer warp 做 Tensor Core。接着是 block/CTA 级：不同 CTA 承担不同 tile 或 pipeline stage。再往上就是 GPU 级：某些 GPU 承担索引、路由、控制、广播，其他 GPU 承担密集 MLA 或 MoE。

这在 DeepSeek-V3.2/DSA 类 sparse attention 里尤其合理。稀疏索引打分和 Top-K 选择需要全局信息，但计算量不一定大；如果所有 GPU 都重复计算 full index scores，会浪费。让一个 GPU 负责 indexer，其余 GPU 专注 MLA，可以减少冗余，并让后续 MoE 阶段重新聚合全部设备。这种异构 worker 设计，本质上是 warp specialization 的系统级版本。

动态性是这条路线的最大挑战#

Talk 最后一页总结 challenge：runtime 才能确定的 shape、routing 和 sparsity 会破坏静态调度；runtime scheduling overhead 增大；动态依赖重新引入 barrier，破坏流水线。

这三条就是常驻 kernel 路线的核心难题：

1. 静态编译越彻底，越怕运行期形状变化；
2. 运行期调度越灵活，越可能把开销搬回关键路径；
3. 动态依赖若处理不好，kernel 内部会重新出现粗 barrier。

TileRT 用特定模型定制绕过了很多通用性问题；MPK/Event Tensor 则试图给出通用抽象。未来低延迟推理框架能否普及，取决于这三点能不能被编译器和 runtime 一起解决。

与本站已有文章的衔接#

读这条谱系，建议和本站几篇文章连起来看：

• GPU GEMM 优化系列：理解 tile、shared memory、register blocking、Tensor Core 和异步流水线；
• FlashAttention 系列：理解 IO-aware attention、online softmax、warp specialization；
• TileLang 完全拆解：理解 tile DSL、layout inference、software pipeline；
• Blink：理解把 CPU 从推理调度关键路径移除的另一种方式；
• TileRT 主文：理解这些 kernel/compiler 思想如何进入 8×B200 的生产低延迟 decode。

把它们连起来，会看到 2026 年推理系统的一条清晰趋势：优化对象正在从「单个算子」上升到「整条 token 生成流水线」。

小结#

TileRT PDF 里的 related work 不是附录，而是低延迟推理系统的技术地基。ThunderKittens 把 warp-level tile 和 producer-consumer pipeline 封装成可写的 kernel 抽象；TileLang 把 tile shape、memory placement、warp partitioning、layout 和 pipeline 变成编译器可搜索的问题；MPK 把多个算子编译成 SM-level tGraph，并用 in-kernel runtime 调度 worker/scheduler；Event Tensor 把动态 shape、routing、sparsity 收进事件张量，让动态 megakernel 有可表达的依赖语义。

TileRT 则把这条路线压进特定模型 decode，换取极致单用户延迟。它不是终点，而是样板：当每个 token 的计算越来越快，系统必须把边界、同步、launch、通信和中间结果搬运一起消掉。未来常驻 kernel 路线能走多远，取决于这套自动化工具链能否从单件定制，走向更多模型、更多 shape 和更多动态控制流。

参考资料#

1. 本站：TileRT 完全拆解
2. 本站：TileLang 完全拆解
3. 本站：GPU GEMM 优化完全拆解（一）
4. 本站：GPU GEMM 优化完全拆解（三）
5. 本站：FlashAttention 完全拆解（三）
6. 本站：Blink：把 CPU 请出推理关键路径