Meta 的 TLX 处在 CUDA 与 Triton 之间。这层中间层已经用于生产

2026 年 11 月 10 日，Meta Engineering 发表文章，称 Triton Low-Level Extensions，简称 TLX，已经在推荐训练、推荐推理和大规模语言模型预训练的生产环境中运行。署名是 Hongtao Yu、Peng Chen、Daohang Shi、Nick Riasanovsky、Lei Wang 和 Partha Kanuparthy。5 月发布在 arXiv 上的论文《TLX: Hardware-Native, Evolvable MIMW GPU Compiler for Large-scale Production Environments》说明了编程模型。可核对的主张，并不是 Meta 从头写了一门新的 GPU 语言，也不是每个内核都已摸到峰值。它是：CUDA 的线程级控制与 Triton 的块级抽象之间缺了一层，也就是 Meta 称为 MIMW 的 warp group 编排，正是这层让 Hopper 和 Blackwell 的硬件特性能够按生产节奏用起来。

TLX 实际给 Triton 加了什么

Triton 是一种嵌在 Python 里的语言，用来在"一块工作"的层面上写 GPU 内核，而不是按单个线程来写。文章称，在 Meta，它已经成为把 PyTorch 模型变成 GPU 代码的常规组成部分。图 1 显示，在 10 个月里，Triton 内核数量大约增长了 3.8 倍，CUDA 约为 2.0 倍，两条曲线都以第 1 个月为 1.0。这是公司报告的内核数量，不是性能分数。

TLX 并不取代 Triton。程序员仍用熟悉的 Triton 张量运算来写计算；需要明确控制执行、数据搬运或同步时，才使用 TLX。用文章的说法，Triton 回答要做什么样的计算，TLX 回答如何编排它。

文章列出的五类扩展都是具体的。Warp group 编排让程序员给相互配合的 warp group 分配角色，例如一个负责搬数据的生产者，和一个负责跑张量核运算的消费者。同步在 warp group 这一级表达，使用 mbarrier 这类硬件原语。内存编排直接提供对 Tensor Memory Accelerator（TMA）、tensor memory（TMEM）以及异步张量核操作的访问，从而让数据搬运和运算可以重叠。Cluster 执行让多个线程块通过片上共享内存相互配合。可选的 layout 意味着编译器会推断默认数据布局，只有在某种布局更快时，程序员才需要指定。

5 月的论文写明，代码开源在 https://github.com/facebookexperimental/triton，也就是 Meta 的实验性 Triton 仓库。我们核验时，该网址返回了公开页面。11 月的文章归在 Open Source 分类下，但正文本身没有给出仓库链接，因此开源主张按论文来表述。

Hopper 和 Blackwell 上缺失的那一层

新近的 NVIDIA GPU 要求程序员同时协调若干专用单元，而不是停留在写出矩阵运算。Hopper 引入了 warp specialization 和 TMA，后者是一条异步搬数据的硬件通路。Blackwell 又增加了 TMEM、cluster 级执行、成对线程块执行，以及更多异步机制。文章举例说，一个现代的 FlashAttention 内核，可能把一个 warp group 专门用于 TMA 传输，另一个用于张量核计算。Warp group 是一组 warp，在 NVIDIA 硬件上通常是 128 个线程，可以在内核里承担不同角色。

CUDA 能够表达这些结构，但往往需要大量针对特定架构的编排代码。Triton 写起来更快，却把许多执行关系藏在编译器的选择后面。作者把这层缺失的接口称为 Multiple Instructions, Multiple Warps，简称 MIMW。5 月的论文使用更短的展开 Multi-Instruction, Multi-Warp。两种表述里，编排单位都是 warp group，处在 CUDA 的线程与 Triton 的块之间。这些硬件特性本来就在。论点是：当唯一的用法要么是直接写 CUDA，要么是等编译器赶上时，它们很少被用到。

生产中的案例

文章称，TLX 已经在 Meta 的生产环境中运行了几个月，并给出三处部署。

第一处是在 NVIDIA Blackwell 上训练 Generative Ads Model，简称 GEM。作者称，GEM 比典型的语言模型训练更难优化，因为输入形状更小，每层有 10 个以上的归一化，模块类型也不均一，这些都会限制模型浮点利用率，也就是实际用到的峰值运算吞吐比例。他们称，最新版本的 GEM 训练在数千块 Blackwell GPU 上达到了约 10^24 次训练浮点运算的量级，端到端利用率为 20% 到 25%。他们在 Blackwell 上评估了 CUTLASS、CuTeDSL 和 TLX，并认为对大多数内核而言，TLX 是生产效率与性能的最佳平衡。对于 GEM 中的 Generalized Dot Product Attention，简称 GDPA，他们报告通过把多种类似注意力的运算并进同一个 GPU 内核，前向最高加速 2 倍。内部工具 Agent for Compute Efficiency，简称 ACE，用于自动做分析与调参，它是这套栈的一部分，而不是文章的主旨。

第二处是推荐推理中的 In-Kernel Broadcast Optimization，简称 IKBO，运行在 Hopper 上。用户特征在所有候选上是相同的，却会在每个批次里被复制许多次。在此前的优化之后，文章称，一个按 warp 分工、多阶段的 TLX 线性压缩核，在 NVIDIA H100 上带来 4 倍加速，并把计算收进一个持续运行的单一内核。与 FlashAttention 结合后，IKBO 被报告为在共同设计的模型上，把计算密集的净延迟最多降低三分之二，并用于 Meta 的 Adaptive Ranking Model，覆盖 GPU 和 MTIA，也就是 Meta Training and Inference Accelerator。

第三处是一套基于 TLX 的分布式分组矩阵乘，已部署在 Meta 的大规模语言模型预训练中。它加速混合专家模型所需的分组矩阵乘，并把通信做进内核，已在大型机群上持续运行。文章没有给出这个内核的加速倍数，只谈了生产使用和稳定性。

作者还称，TLX 的性能通常与手工调优的 CUDA 相当，但迭代更快。他们写，IKBO 的实现用了数周而不是数月。这是公司自报的工期，不是受控实验。

这些数字不能证明什么

GEM 上 20% 到 25% 的端到端利用率，仍然只是峰值的一小部分。文章把它写成一个异构推荐模型上得来不易的工作点，而不是证明整套训练栈已经接近硬件上限。Meta 也说，对大多数内核而言，TLX 是生产效率与性能的最佳平衡，这并不等于在与 CUTLASS 或 CuTeDSL 的每场竞赛中都取胜。

2 倍和 4 倍都是针对具名内核的"最高"结果，分别是 GDPA 和线性压缩核，而不是对每一个 GEM 或推荐模型的端到端加速。部署到机群的混合专家内核，没有公开的加速倍数。Triton 内核数量增长 3.8 倍，是 Meta 内部的数量统计，不能证明那些内核更快。ACE 也不是对 TLX 的独立评估。

更公允的读法

更公允的读法，是生产节奏加上具体内核的加速，而不是声称 TLX 赢得每场内核竞赛。如果一层中间层能让工程师在 Triton 的 Python 模型里表达 warp group 角色、TMA 传输和 TMEM，并且这层已经在 GEM、一套排序模型和预训练机群上运行，那么"缺了一层"这个判断就是在做实事。若 IKBO 用了数周而不是数月的说法属实，这也是论据的一部分。作者自己也说，一个只快一点、却要花三倍时间来写的内核，在机群规模上是净损失。

诚实的界限就在这一判断旁边。端到端利用率仍然不高，头条加速也只连着特定内核。Meta 自己在 GEM 上仍把 CUTLASS 和 CuTeDSL 留在工具箱里，并把 TLX 说成大多数这类工作负载的合适抽象，而不是全部。

谁该在意，接下来看什么

在 Hopper 或 Blackwell 上编写生产内核的团队应当关心，想把同一套想法用到 AMD MI300 或 MI350x 上的团队也应当关心。文章称，对更多架构的支持，包括 MTIA，正在推进，作者还希望让 TLX 成为 torch.compile 的一个后端。未来六个月里，真正的检验是：这个实验性 Triton 仓库是否仍是 Meta 以外的人用得上的公开路径；那条 torch.compile 路径是否以他人可运行的形式出现；以及那些具名的生产内核，是否会配上可公开比较的基线，而不是停留在"最高可达"的倍数上。

收束

TLX 是一层扩展，把 warp group 编排做成 Triton 里的一等能力，而且 Meta 称这层已经用于生产。接下来要看的，不是内核基准上的通吃。而是当下一世代 GPU 到来时，这层中间层是否仍然好用，而不至于迫使团队退回针对特定架构的 CUDA，或再去等待编译器。