OpenAI久违发了篇「正经」论文：线性结构实现高效张量盘算

OpenAI久违发了篇「正经」论文：线性结构实现高效张量盘算

机械之心报道

编辑：Panda

OpenAI 发论文的频率是越来越低了。

如果你看到了一份来自 OpenAI 的新 PDF 文件，那大都也是新模型的系统卡或相关增补文件或基准测试，很少有新的研究论文。

至于原因嘛，让该公司自家的 ChatGPT 来说吧：「截至目前，OpenAI 在 2025 年在 arXiv 上果真宣布的论文数量相对较少，可能反应了其对研究结果果真战略的谨慎态度，可能出于商业保密或宁静考虑。」

不过近日，OpenAI 也确实宣布了一份完全由自己人加入的、实打实的研究论文，其中提出了一种用于高效张量映射的统一代数框架Linear Layouts。这是一种使用二元线性代数而非比特体现（bit representation）的张量结构的通用代数形式，解决了 Triton 等深度学习编译器中恒久保存的难题。

论文题目：Linear Layouts: Robust Code Generation of Efficient Tensor Computation Using ?论文地点：https://arxiv.org/pdf/2505.23819.pdf

要理解这项研究的意义，首先需要先理解一下什么是张量结构（tensor layouts）

简单来说：张量结构 = 逻辑张量与硬件资源（例如内存、线程、向量单位）之间的映射关系。下图给出了两个结构示例。

关于现代深度学习事情负载而言，所需要的张量结构需要满足几个要求：

高效（为了性能）。灵活（以支持多种算子）�？勺楹希ㄎ吮浠缓陀呕�。

然而，目今的结构系统却难以充分满足这些需求，而是往往：

需要凭据实际需求设计，并且往往是硬编码的（需要手动编写规则）。不可扩展（每一对结构都需要二次组合）。容易蜕化，尤其是在像 Triton 这样的低层级的后端中 —— 截至目前，Triton 的 GitHub 库中提交的 12% 的 Bug 与结构有关。

另外，深度学习硬件（如 GPU）的日益庞大也导致张量结构日益庞大。

例如，为了实现高效的矩阵乘法，英伟达在 Ampere、Hopper 和 Blackwell 等差别代际的 GPU 上接纳了差别的使用 Tensor Core 的结构，并且每种结构在使用差别数据类型时都有差别的变体。AMD 和英特尔等其它 GPU 供应商在利用其类似 Tensor Core 的技术进行加速时，也使用了差别的结构。因此，硬件架构的快速生长和多样化的深度学习模型需要一种新的张量结构建模要领。

为此，需要解决一些技术难题：

在将张量映射到硬件资源方面，需要一种通用且可组合的体现要领。结构转换应该用统一的形式来表达，甚至需要包括诸如数据交换（data swizzling）等庞大变换。这种体现必须与初级硬件优化无缝集成，以确保高效的数据会见和盘算。

不过，在介绍 OpenAI 这篇论文的孝敬之前，我们需要先了解一些基础看法。

相关配景知识

GPU 架构

在设计上，现代 GPU 的目标是通过包括多层硬件资源的分层执行模型来充分利用并行性。

其要害执行单位包括协作线程阵列 (CTA)、Warp 和线程。每个 GPU 线程都可以会见私有寄存器 —— 这些寄存器提供最低延迟的存储空间，但容量有限。通例指令可以由各个线程独立执行。然而，某些特殊功效单位必须在更高的粒度级别上执行。

例如，英伟达的 mma（矩阵乘法累加）指令利用 Tensor Core 的方法是并行执行由各个 Warp 发出的多个乘加运算。而 wgmma（Warp 组矩阵乘法累加）等高级变体则是通过在多个 Warp 上同时执行矩阵乘法而对这些功效进行了扩展。AMD 也引入了类似的原语，例如 mfma（矩阵融合乘加）指令。

请注意，这些指令要求数据漫衍在线程和 Warp 之间，或者以特殊结构驻留在共享内存或特殊内存单位（例如 Blackwell 上的 Tensor Memory）中，才华爆发正确的结果。

然而，这些结构通常不会为加载 / 存储等其他操作带来最佳性能，并且并非总是可以使用特定指令将数据直接从全局内存复制到特殊内存单位。

因此，通常必须对数据进行重新排列，以便将用于内存会见的结构转换为盘算单位偏好的结构。

简而言之，要实现峰值性能，不但需要利用这些专用单位，还需要精心设计张量结构和转换。

Triton 语言和编译器

Triton 是一种类似于 Python 的用于特定领域的语言，其设计目标是提供用于编写高性能深度学习原语的灵活接口。Triton 的编译器后端使用了 MLIR，支持多条理笼统表达。

究其焦点，Triton 内核遵循单程序大都据 (SPMD) 模型，其中盘算被划分为多个笼统的 Triton 程序实例。这种设计允许开发者主要关注 CTA 级别的并行性即可。在 Triton 中，「张量」一词指的是从原始 PyTorch 张量中提取的块，它们用作 GPU 核的输入和输出。

在编译历程中，Triton 的 Python 代码首先被翻译成 Triton 方言 (tt)，然后进一步翻译成 TritonGPU 方言 (ttg)。在此历程中，每个张量都与特定的结构相关联，以充分利用现代 GPU 上可用的硬件功效单位。例如，当遇到 dot 类算子（例如 tt.dot 和 tt.dot_scaled）时，会接纳 mma 结构并使用 Tensor Core 和类似的单位。

古板结构

图 2 列出了 Triton 中所有可用的结构。

在最高层级，结构分为漫衍式（Distributed）结构和内存（（Memory）结构。前者是指张量元素漫衍在差别的执行单位中，此后者是指张量元素存储在特定的特殊内存中。

漫衍式结构又可进一步分为 Blocked、Sliced、MMA 和 MMA Input 结构等类型，而内存结构又可进一步分为 Unswizzled 和 Swizzled 结构。

Blocked 结构通常用于连续的内存会见。MMA 和 MMA 输入结构用于矩阵乘法运算（例如 tt.dot）的输出和输入。MMA 结构可以凭据其映射到的硬件指令进一步分类，例如英伟达 GPU 上的 mma 和 wgmma，或 AMD GPU 上的 mfma。Sliced 结构是从其父结构中提取一个维度，用作广播或某个归约运算的输出。

古板 Triton 结构系统要求每个结构界说自己的接口要领，例如每个线程的元素数量和连续元素的数量。别的，必须为每个结构显式实现对张量元素的索引以及结构之间的转换。这种要领导致结构结构和转换常泛起 bug。

Linear Layouts（线性结构）

下面将简单介绍线性结构的界说、一些基本的线性结构算子、创立种种 Triton 结构以作为线性结构实例，以及应用于 Triton 的通用结构引擎。

一个示例

在 GPU 编程中，大大都参数都是 2 的幂：一个 Warp 由 32 或 64 个线程组成，一个 Warp 组包括 4 个 Warp，矩阵乘法内联函数（例如 mma 和 wgmma）要求 Tile 尺寸为 16 × ，其中 ≥ 1。

别的，在 Triton 的编程模型中，张量的维度以及与每个张量相关的结构子部分（例如每个线程的寄存器和线程数量）都被限制为 2 的幂。在图 1 中，结构 A 有一个 16 × 16 的张量，其使用了多个 2 × 2 的寄存器、4 × 8 的线程和 2 × 1 的 Warp。

由于这些量都是 2 的幂，因此使用其坐标的比特体现，可以直观地可视化结构 A 中元素的漫衍（如图 1 所示）。所有线程的寄存器 0 (_0) 都位于坐标 (, )，其中和的最后几位（bit）均为 0。例如，线程 _1 的 _0 位于 (0, 2) = (000, 010)。作为比照，_1 元素的坐标中，的最后一位始终为 0，而的最后一位始终为 1。例如，_9 的 _1 位于 (2, 3) = (010, 011)。

别的，关于任何偶数线程 _，的最后一位与 _0 中的倒数第二位匹配，的倒数第二位与 _0 中的倒数第三位匹配。例如，_10 = _01010 的 _0 位于 (2, 4) = (010, 0100)。这种系统性对齐连续保存，标明二次幂结构足以清晰地决定了每个线程元素的漫衍。

综上所述，假设一个巨细为 8 的向量体现一个 Warp 中线程的一个元素，其中前 2 位体现寄存器 (Reg)，接下来的 5 位体现线程 (Thr)，最后一位则体现 Warp (Wrp)，则可以如此界说结构：

当需要从逻辑张量的坐标中恢复硬件索引时，需要使用求逆运算。

对线性结构的更详细完备性说明请会见原论文，其中涉及到说明分块结构、mma 和 wgmma 的输入和输出结构、线性结构的 slice、每个漫衍式结构、MMA swizzled 结构、内存结构都是线性结构。另外，OpenAI 也在 Triton 说明了如何实现结构转换以及形状操作。

不但如此，OpenAI 体现，线性结构为在语言前端和编译器后端开发算法提供了结构化的基础。他们也在论文中给出了一些要害示例，这里就不过多展开。接下来简单看看新提出的线性结构的实际体现。

评估

OpenAI 将优化版 Triton（集成了基于线性结构的优化，即 Triton-Linear）与未集成这些优化的基准 Triton 进行了比较。Triton 和 TritonLinear 之间的主要区别如下：

Triton 使用古板的数据结构，不支持任意漫衍式结构的实用程序或它们之间的转换，因此容易泛起 bug。Triton 未接纳论文中描述的优化代码生成。例如，结构转换始终通过共享内存进行，对高效硬件原语的使用有限。

加入评估的硬件平台见表 1。

为了比较 Triton 和 Triton-Linear 的性能，该团队构建了一些合成微基准来进行测试，这方面的结果请会见原论文检察。这里仅看看它们在实际基准测试中体现。

在三个差别的平台上，OpenAI 运行了 TritonBench 中的 18 个基准测试。图 7、图 8 和图 9 中展示了 Triton-Linear 在三个平台上的性能提升。

由于每个基准测试包括多个输入，总计 420 个案例，因此他们使用了误差线（error bars）来体现每个基准测试的最小和最大加速。

需要注意的是，由于硬件限制，并非所有基准测试都适用于每个平台。例如，某些基准测试需要仅在 GH200 上才有的大型共享内存，而一些核使用的张量描述符依赖于 TMA 引擎，而 RTX4090 和 MI250 上均不支持 TMA 引擎。

可以看到，在 GH200 上，他们实现了 0.92 倍到 1.57 倍不等的加速，所有基准测试的平均加速均凌驾 1.0 倍。加速最显著的基准测试是 int4_gemm、ops_gemm 和 streamk_gemm。

可以视察到，高效的硬件原语（例如 ldmatrix 和 stmatrix）在这些核中被广泛用于结构转换以及共享内存的加载和存储操作。值得注意的是，layer_norm 实现了从 0.99 倍到 1.57 倍的加速 —— 在差别形状之间体现出了显著差别。关于某些输入形状，Triton-Linear 能够检测「等效」结构之间的转换，从而将转换历程降低为 no-op（无操作）。这种优化在旧版结构系统中无法实现，因为它无法直接比较差别类型的结构（例如，Blocked 结构和 Sliced 结构）。

在 RTX4090 上，新要领实现了 1.00 倍到 1.51 倍的加速。由于 mma (RTX4090) 和 wgmma (GH200) 指令之间的差别，他们在 template_attention 上实现了更高的加速。在本例中，tt.dot 运算的左操作数在循环外部界说，会重复从同一地点加载数据，因此 ldmatrix 和通例共享内存指令均可实现高吞吐量。虽然右操作数在每次迭代中都会更新，但 wgmma 会直接在共享内存中会见它，只有在 RTX4090 上，经过优化后，它才会被降级到 ldmatrix 中。因此，在 GH200 上实现的加速相对较低。在 MI250 上，新要领实现了 0.98 倍到 1.18 倍的加速。

总体而言，由于缺乏 ldmatrix 等高效的硬件原语，Triton-Linear 在 AMD GPU 上实现的加速低于在英伟达 GPU 的。

关于 OpenAI Open 的这个研究，你有什么看法呢

分享： 2025-06-06 23:27:37 共81款