英伟达重磅！CUDA 13.1改写GPU编程，AI性能狂飙！

GPU编程的演进，一直是科技界关注的焦点。特别是随着人工智能技术的飞速发展，如何高效、便捷地利用GPU强大的计算能力，成为了众多开发者和研究者面临的重要课题。在这个背景下，英伟达（NVIDIA）公司在GPU编程模型上的创新，无疑为业界带来了新的曙光。

GPU编程新范式：CUDA Tile的崛起

英伟达推出的CUDA Tile，正是一种基于GPU的编程模型，其核心目标是为英伟达Tensor Core（张量核心）提供更强的可移植性，进而释放GPU的峰值性能。值得一提的是，CUDA Tile的灵活性让开发者能够在其之上构建自己的领域特定语言（DSL）。新媒网跨境获悉，目前英伟达正在积极推动一项重要工作：将CUDA Tile作为OpenAI Triton的后端。

OpenAI Triton是一款开源的Python领域特定语言，专为编写深度学习（DL）GPU内核而设计。它支持分块计算（tiled computation），这项技术能够将数据和计算任务划分为更小的块，以提高效率。Triton内部集成了一个基于MLIR的编译器，能够生成PTX代码，这使得即便不具备深厚CUDA编程经验的研究人员，也能编写出高效的GPU代码。这一进步，无疑将大幅降低GPU高性能计算的门槛。

深入理解CUDA Tile与CUDA Tile IR

CUDA Tile对现有的CUDA编程模型进行了扩展，旨在为分块编程提供一流的原生支持。这项技术首次在CUDA 13.1版本中亮相，标志着GPU编程领域的一次范式转变。以往，开发者需要以SIMT（单指令多线程）模型为基础，考虑每一个独立线程的行为。而现在，瓦片（tile）模型允许开发者在更高的抽象层次上表达计算逻辑。

简而言之，开发者只需指定对数据块（瓦片）的操作，剩下的线程调度、硬件映射以及资源分配等底层细节，都将由编译器和运行时系统自动完成。这种设计不仅大大降低了编程的复杂性，同时也为编译器提供了更大的优化空间，能够实现更积极、更深层次的性能优化。

CUDA Tile IR则是基于MLIR的中间表示形式（IR）和编译器基础设施。它的发展遵循CUDA Tile IR规范，该规范详细定义了在英伟达GPU上进行瓦片式计算的形式语义、操作以及类型系统。这为瓦片计算提供了一个统一且标准化的底层接口。

Triton-to-TileIR：连接高级语言与下一代GPU模型

那么，Triton-to-TileIR究竟扮演着怎样的角色呢？它实际上是Triton的一个后端桥梁，使得Triton能够将目标代码编译为CUDA Tile IR，而非传统的PTX。通过这一扩展，Triton的编译器生态系统得到了进一步增强，开发者可以将用OpenAI Triton编写的GPU内核编译并执行到新引入的CUDA Tile IR后端上。

这项集成意义非凡，它将高级编程语言（如Triton）与英伟达下一代GPU编程模型紧密连接起来，为开发者提供了无缝利用现代硬件能力的新途径，且无需重写现有代码。随着GPU编程从传统的SIMT模型向瓦片式抽象的演进，这种集成使得开发者既能享受到Triton易用的Python语法，又能直接获得Tile IR对Tensor Core的原生支持和卓越的架构可移植性。

Triton-to-TileIR的出现，让更多人能够接触到这些前沿功能。值得注意的是，Triton本身就是一种基于瓦片（tile）的编程语言——开发者以数据块（瓦片）而非单个线程来表达计算，这与CUDA Tile IR的理念不谋而合。

这就提供了一条直接的后端编译路径：Triton-to-TileIR不再将Triton的瓦片级抽象编译成线程级的SIMT代码，而是保留瓦片级的语义，直接编译为原生理解瓦片粒度计算的CUDA Tile IR。

这意味着，现有的Triton用户社区无需学习新语言或重写现有代码，就能直接享受到CUDA Tile IR带来的诸多优势。只需简单地配置一个环境变量，就可以将编译流程从PTX后端切换到CUDA Tile IR后端，从而实现性能提升，并获得面向未来的架构兼容性。Triton用户甚至可以在其应用程序中，针对每个内核选择使用PTX后端还是CUDA Tile IR后端。

Triton-to-TileIR的开发路线图

作为triton-lang组织内部的一个孵化项目，Triton-to-TileIR正处于积极的开发阶段。其代码仓库作为一个重要的协作空间，用于实现和完善CUDA Tile IR后端，为未来可能集成到Triton主编译器中做好准备。新媒网跨境了解到，项目的开发路线图涵盖了多个技术工作流，其中包括：

核心转换基础设施： 实现MLIR方言转换模式，将Triton操作映射到对应的CUDA Tile IR操作。
测试与验证： 开发全面的测试套件，以验证转换的语义正确性，包括对控制流、内存访问模式以及数值精度等边缘情况的测试。
性能基准测试： 建立性能基准，比较通过Tile IR编译的内核与通过PTX编译的内核在各种操作（如矩阵乘法、卷积、逐元素操作、归约等）上的表现。
开源项目集成： 与开源社区进行协调，以期在诸如Helion等开源项目中更好地支持CUDA Tile IR后端。

如何使用Triton-to-TileIR

目前，Triton-to-TileIR仅支持基于源代码的编译。暂不提供预编译二进制文件，这意味着开发者需要在本地环境中从源代码构建该项目。

使用前提

CUDA版本： 必须为CUDA 13.1或更高版本。
GPU架构： 当前仅支持英伟达Blackwell系列GPU（例如，GeForce RTX 5080）。旧的GPU架构将在未来的CUDA版本中逐步得到支持。

从源代码构建

当满足所有前提条件后，即可克隆项目并从源代码进行构建：

# 克隆仓库
git clone https://github.com/triton-lang/Triton-to-tile-IR.git
cd Triton-to-tile-IR
# 构建并安装
# 具体构建指令可能因项目README而异
pip install -e .

请注意，详细的构建步骤可能有所不同。建议查阅Triton-to-TileIR项目的README文件和构建文档，以获取针对特定架构的配置、依赖管理和故障排除指南。

验证Tile IR编译

构建完成后，可以通过运行矢量加法教程来验证安装是否成功，并确认是否正在使用Tile IR后端：

# 进入教程目录
cd python/tutorials
# 启用Tile IR并运行矢量加法示例
export ENABLE_TILE=1 python 01-vector-add.py

当Tile IR后端处于活动状态时，Triton会缓存编译后的内核文件，其扩展名为.tileIR，而非SIMT后端常用的.cubin文件。您可以检查这些缓存文件：

# 查找Triton缓存目录（通常位于~/.triton/cache）

Triton-to-TileIR的局限性

尽管Triton-to-TileIR开启了GPU编程充满希望的新篇章，但作为一个相对早期开发阶段的项目，它仍然存在一些已知的局限性，例如部分操作暂不支持以及暂时的性能问题。

未支持的操作

并非所有Triton支持的操作都已在Tile IR后端中实现。随着CUDA持续发布新版本，Triton CUDA Tile IR后端的兼容性也将不断提高。

指针张量（Tensor-of-pointer）性能退化问题

在CUDA 13.1版本中，Triton中的“指针张量”（tensor-of-pointer）模式，即张量由描述内存访问模式的指针组成，在Tile IR后端上表现出次优的性能。这是一种暂时性的性能状况。对于受影响的工作负载，开发者可以采取以下策略：

临时回退到SIMT后端： 对于某些关键操作，暂时切换回SIMT后端。
等待未来的优化： 期待在项目后续版本中发布的优化通道。
重构代码采用TMA加载/存储API： 优化代码，使其采用TMA加载/存储（TMA load/store）API。

关于第三点，即重构代码采用TMA load/store API的说明：许多在内核中加载的张量都具有连续的瓦片以及明确的形状和步幅。因此，在内核内部实例化一个“指针张量”已不再是必需的。相反，这些布局信息可以直接传递给TMA load/store API，从而使Tile IR后端获得更好的性能。例如，典型的“指针张量”模式可能如下所示：

# 之前：指针张量样式
offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
offs_k = tl.arange(0, BLOCK_K)
a_ptrs = a_ptr + (offs_m[:, None] * stride_am + offs_k[None, :] * stride_ak)
b_ptrs = b_ptr + (offs_k[:, None] * stride_bk + offs_n[None, :] * stride_bn)
a = tl.load(a_ptrs)
b = tl.load(b_ptrs)

在这里，a_ptrs中的每个元素都是在内核中计算出的显式指针，尽管瓦片本身是连续的，并且其布局可以通过（shape、strides、block_shape）完全描述。使用TMA，相同的操作可以重写为：

desc_a = tl.make_tensor_descriptor(
    a, # 基指针
    shape=(M, K),
    strides=(stride_am, stride_ak),
    block_shape=(BLOCK_M, BLOCK_K) # 瓦片大小
)
desc_b = tl.make_tensor_descriptor(
    b,
    shape=(K, N),
    strides=(stride_bk, stride_bn),
    block_shape=(BLOCK_K, BLOCK_N)
)
offs_m = pid_m * BLOCK_M
offs_n = pid_n * BLOCK_N
a_tile = desc_a.load([offs_m, 0]) # [BLOCK_M, BLOCK_K]
b_tile = desc_b.load([0, offs_n]) # [BLOCK_K, BLOCK_N]
desc_c.store([offs_m, offs_n], acc) # TMA支持的存储

深入了解Triton-to-TileIR

Triton-to-TileIR项目代表了GPU编程发展中的一个重要里程碑，它成功地弥合了开发者生产力与硬件效率之间的鸿沟。通过使Triton易用且面向瓦片的编程模型能够以CUDA Tile IR虚拟指令集为目标，这项集成有望为机器学习从业者和GPU开发者带来卓越的性能、可移植性以及面向未来的就绪度。

对于已经在使用Triton的开发者而言，Tile IR后端提供了一条通往利用下一代GPU架构的途径，并且只需进行最少的代码更改。对于更广泛的GPU编程生态系统而言，这项合作展示了语言设计者与硬件供应商之间的战略伙伴关系如何能够产生复合效益——在不牺牲实现快速创新所需的高级抽象的情况下，使先进的硬件功能变得易于访问。

随着该项目从孵化阶段走向生产就绪，其将如何影响Triton的普及以及瓦片式GPU编程的整体发展轨迹，无疑将是一件令人着迷的事情。新媒网跨境认为，最终的成功衡量标准将很简单：那些缺乏GPU专业知识的研究人员，能否编写出在英伟达GPU上以接近最优性能执行的Triton代码？答案很快就会揭晓。

项目GitHub仓库地址：triton-lang/Triton-to-tile-IR

新媒网（公号: 新媒网跨境发布），是一个专业的跨境电商、游戏、支付、贸易和广告社区平台，为百万跨境人传递最新的海外淘金精准资讯情报。

本文来源：新媒网 https://nmedialink.com/posts/nvidia-rewrites-gpu-ai-coding-cuda-131.html