OpenAI久违发了篇「正经」论文：线性布局实现高效张量计算
仅用于站内搜索，没有排版格式，具体信息请跳转上方微信公众号内链接
机器之心报道
编辑：Panda
OpenAI发论文的频率是越来越低了。
如果你看到了一份来自OpenAI的新PDF文件，那多半也是新模型的系统卡或相关增补文件或基准测试，很少有新的研究论文。
至于原因嘛，让该公司自家的ChatGPT来说吧：「截至目前，OpenAI在2025年在arXiv上公开发布的论文数量相对较少，可能反映了其对研究成果公开策略的谨慎态度，可能出于商业保密或安全考虑。」
不过近日，OpenAI也确实发布了一份完全由自己人参与的、实打实的研究论文，其中提出了一种用于高效张量映射的统一代数框架LinearLayouts。这是一种使用二元线性代数而非比特表示（bitrepresentation）的张量布局的通用代数形式，解决了Triton等深度学习编译器中长期存在的难题。
论文标题：LinearLayouts:RobustCodeGenerationofEfficientTensorComputationUsing𝔽₂
论文地址：https ://arxiv. org/pdf/2505. 23819.pdf
要理解这项研究的意义，首先需要先理解一下什么是张量布局（tensorlayouts）。
简单来说：张量布局=逻辑张量与硬件资源（例如内存、线程、向量单元）之间的映射关系。下图给出了两个布局示例。
对于现代深度学习工作负载而言，所需要的张量布局需要满足几个要求：
高效（为了性能）。
灵活（以支持多种算子）。
可组合（为了变换和优化）。
然而，当前的布局系统却难以充分满足这些需求，而是往往：
需要根据实际需求设计，而且往往是硬编码的（需要手动编写规则）。
不可扩展（每一对布局都需要二次组合）。
容易出错，尤其是在像Triton这样的低层级的后端中——截至目前，Triton的GitHub库中提交的12%的Bug与布局有关。
另外，深度学习硬件（如GPU）的日益复杂也导致张量布局日益复杂。
例如，为了实现高效的矩阵乘法，英伟达在Ampere、Hopper和Blackwell等不同代际的GPU上采用了不同的使用TensorCore的布局，并且每种布局在使用不同数据类型时都有不同的变体。AMD和英特尔等其它GPU供应商在利用其类似TensorCore的技术进行加速时，也使用了不同的布局。因此，硬件架构的快速发展和多样化的深度学习模型需要一种新的张量布局建模方法。
为此，需要解决一些技术难题：
在将张量映射到硬件资源方面，需要一种通用且可组合的表示方法。
布局转换应该用统一的形式来表达，甚至需要包含诸如数据交换（dataswizzling）等复杂变换。
这种表示必须与低级硬件优化无缝集成，以确保高效的数据访问和计算。
不过，在介绍OpenAI这篇论文的贡献之前，我们需要先了解一些基础概念。
相关背景知识
GPU架构
在设计上，现代GPU的目标是通过包含多层硬件资源的分层执行模型来充分利用并行性。
其关键执行单元包括协作线程阵列(CTA)、Warp和线程。每个GPU线程都可以访问私有寄存器——这些寄存器提供最低延迟的存储空间，但容量有限。常规指令可以由各个线程独立执行。然而，某些特殊功能单元必须在更高的粒度级别上执行。
例如，英伟达的mma（矩阵乘法累加）指令利用TensorCore的方式是并行执行由各个Warp发出的多个乘加运算。而wgmma（Warp组矩阵乘法累加）等高级变体则是通过在多个Warp上同时执行矩阵乘法而对这些功能进行了扩展。AMD也引入了类似的原语，例如mfma（矩阵融合乘加）指令。
请注意，这些指令要求数据分布在线程和Warp之间，或者以特殊布局驻留在共享内存或特殊内存单元（例如Blackwell上的TensorMemory）中，才能产生正确的结果。
然而，这些布局通常不会为加载/存储等其他操作带来最佳性能，而且并非总是可以使用特定指令将数据直接从全局内存复制到特殊内存单元。
因此，通常必须对数据进行重新排列，以便将用于内存访问的布局转换为计算单元偏好的布局。
简而言之，要实现峰值性能，不仅需要利用这些专用单元，还需要精心设计张量布局和转换。
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布局并使用TensorCore和类似的单元。
𝔽₂数学基础
我们可将两个元素{0,1}的域表示为𝔽₂。在𝔽₂中，所有算术运算均以2为模执行。
例如，加法定义为，其对应于逻辑异或（XOR）。
而乘法定义为，对应于逻辑与（AND）。
在𝔽₂上，一个基本运算是矩阵乘法。令是元素在𝔽₂中的两个矩阵。乘积𝐶=𝐴𝐵的各个元素的定义为
这类似于标准矩阵乘法，不同之处在于所有算术运算都在𝔽₂中执行。
𝔽₂中的算术运算与二进制逻辑自然契合，使得该领域的运算在硬件实现中非常高效。因此，𝔽₂广泛应用于密码学和纠错码等领域。
传统布局
图2列出了Triton中所有可用的布局。
在最高层级，布局分为分布式（Distributed）布局和内存（（Memory）布局。前者是指张量元素分布在不同的执行单元中，而后者是指张量元素存储在特定的特殊内存中。
分布式布局又可进一步分为Blocked、Sliced、MMA和MMAInput布局等类型，而内存布局又可进一步分为Unswizzled和Swizzled布局。
Blocked布局通常用于连续的内存访问。MMA和MMA输入布局用于矩阵乘法运算（例如tt. dot）的输出和输入。MMA布局可以根据其映射到的硬件指令进一步分类，例如英伟达GPU上的mma和wgmma，或AMDGPU上的mfma。Sliced布局是从其父布局中提取一个维度，用作广播或某个归约运算的输出。
传统Triton布局系统要求每个布局定义自己的接口方法，例如每个线程的元素数量和连续元素的数量。此外，必须为每个布局显式实现对张量元素的索引以及布局之间的转换。这种方法导致布局构造和转换常出现bug。
LinearLayouts（线性布局）
下面将简单介绍线性布局的定义、一些基本的线性布局算子、创建各种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)=(0𝑏00,0𝑏10)。作为对比，𝑟_1元素的坐标中，𝑖的最后一位始终为0，而𝑗的最后一位始终为1。例如，𝑡_9的𝑟_1位于(2,3)=(0𝑏10,0𝑏11)。
此外，对于任何偶数线程𝑡_𝑘，𝑘的最后一位与𝑟_0中𝑗的倒数第二位匹配，𝑘的倒数第二位与𝑟_0中𝑗的倒数第三位匹配。例如，𝑡_10=𝑡_0𝑏1010的𝑟_0位于(2,4)=(0𝑏10,0𝑏100)。这种系统性对齐持续存在，表明二次幂结构足以清晰地决定了每个线程元素的分布。
综上所述，假设一个大小为8的向量𝑣表示一个Warp中线程的一个元素，其中前2位表示寄存器(Reg)，接下来的5位表示线程(Thr)，最后一位则表示Warp(Wrp)，则可以如此定义布局𝐴：
带标注的向量空间。如此，可为布局中的每一位（bit）分配标签。输入𝑣位于空间中，建模了Reg×Thr×Wrp的空间。输出𝑤遵循结构，表示逻辑张量(𝑖,𝑗)的两个维度。
定义与构造
定义1（线性布局/LinearLayouts）。线性布局的定义是在𝔽₂上的向量空间之间的线性映射。
定义2（组合/Composition）。给定𝔽₂上的向量空间𝑈、𝑉、𝑊以及线性布局𝐿₁:𝑈→𝑉和𝐿₂:𝑉→𝑊，它们的组合定义为：
将𝐿₁和𝐿₂表示为矩阵𝑀₁和𝑀₂，表示𝐿₂◦𝐿₁的矩阵由𝔽₂上的（逐标签）矩阵乘法𝑀₂𝑀₁给出。
定义3（积/Product）。给定𝔽₂上的两个向量空间𝑈和𝑉，定义它们的积为：
给定两个线性布局𝐿₁:𝑈₁→𝑉₁,𝐿₂:𝑈₂→𝑉₂，且𝑢1₁∈𝑈₁,𝑢₂∈𝑈₂，则定义它们的积为：
将𝐿₁和𝐿₂表示成矩阵𝑀₁和𝑀₂，则表示𝐿₁×𝐿₂的矩阵为：
定义4（左除/LeftDivision）。若矩阵𝑀具有如下结构，则矩阵𝑀左侧可被矩阵𝑀₁整除：
这里将左侧的除法记为。此运算可在线性布局中逐标签处理。
左除法可用于确定布局是否可以分解为满足高效硬件原语（例如ldmatrix）的较小布局，
定义5（右逆/RightInverse）。在𝔽₂上的满射线性布局𝐿:𝑈→𝑉具有一个右逆。
如果𝑀是𝐿的一个矩阵表示，其形状为𝑚×𝑛，则可将𝑀−1定义为𝑀𝑋=𝐼_𝑚的𝑛×𝑚最小二乘解，其中𝐼_𝑚是𝑚×𝑚的单位矩阵。具体来说，它可以通过对𝔽₂进行高斯消元法计算得出。
当需要从逻辑张量的坐标中恢复硬件索引时，需要使用求逆运算。
对线性布局的更详细完备性说明请访问原论文，其中涉及到说明分块布局、mma和wgmma的输入和输出布局、线性布局的slice、每个分布式布局、MMAswizzled布局、内存布局都是线性布局。另外，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个案例，因此他们使用了误差线（errorbars）来表示每个基准测试的最小和最大加速。
需要注意的是，由于硬件限制，并非所有基准测试都适用于每个平台。例如，某些基准测试需要仅在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在AMDGPU上实现的加速低于在英伟达GPU的。
对于OpenAIOpen的这个研究，你有什么看法呢？
©THEEND
转载请联系本公众号获得授权
投稿或寻求报道：liyazhou@jiqizhixin. com
