OpenAI Triton 1.0 发布

OpenAI 发布了 Triton 1.0，这是一种开源的类 Python 编程语言和编译器，旨在让研究人员无需具备深厚的 CUDA 经验即可编写高效的 GPU 代码。Triton 允许开发者使用显著更少的代码，实现与专家级调优的 CUDA 内核相当的硬件性能——例如，在 FP16 矩阵乘法方面可以达到与 cuBLAS 相当的性能。

解决 GPU 编程的复杂性

为了避免在原生框架算子中创建和移动临时张量所带来的性能损失，编写专门的 GPU 内核通常是必要的。然而，手动进行 GPU 优化非常困难，因为这需要开发者管理三个主要的架构组件：

DRAM: 内存传输必须合并为大型事务，以利用现代接口的总线宽度。
SRAM: 数据在复用前必须手动存入 SRAM，并进行管理以避免共享内存 Bank 冲突。
ALUs: 计算必须在流式多处理器 (SMs) 之间及内部进行仔细的划分和调度，以利用指令/线程级并行性和 Tensor Cores。

Triton 自动完成了内存合并、共享内存管理以及 SMs 内部的调度，同时将高层算法考量——如分块 (tiling) 和 SMs 间的同步——留给开发者。

Triton 编程模型

Triton 使用一种编程模型，其中内核被定义为带有装饰器的 Python 函数，并在实例网格 (grid of instances) 上启动。与 CUDA 或 Numba 使用的单指令多线程 (SIMT) 模型不同，Triton 通过对块 (blocks)（维度为 2 的幂次方的微型数组）的操作来暴露实例内的并行性。

关键技术区别

基于块的执行: 通过对块而非单个线程进行操作，Triton 抽象掉了 CUDA 线程块内的并发问题，包括共享内存同步和 Tensor Core 调度。
指针算术: Triton JIT 将输入视为指针而非张量，为处理像块稀疏张量 (block-sparse tensors) 这样复杂的数据结构提供了必要的底层内存访问控制。
融合内核 (Fused Kernels): Triton 简化了融合内核的创建。例如，在 Triton 中的融合 softmax 实现可以在整个归一化过程中将行保留在 SRAM 中，从而最大限度地提高数据复用。这种方法比等效的 PyTorch 实现效率高出多达 2 倍。

矩阵乘法性能

Triton 在矩阵乘法方面非常高效，而矩阵乘法是神经网络中的核心操作。它仅需大约 25 行 Python 代码，即可在 V100 tensor cores 上实现峰值性能。这种易用性允许开发者自定义包含融合变换（如 slicing 或 Leaky ReLU）的矩阵乘法内核，而无需具备卓越的 GPU 编程专业知识。

系统架构与编译器后端

Triton 的性能源于其以 Triton-IR 为核心的模块化架构，这是一种基于 LLVM 的中间表示，其中多维块是其一等公民。

编译流水线

Python AST: @triton.jit 装饰器遍历 Python 函数的抽象语法树 (AST)。
Triton-IR: 使用静态单赋值 (SSA) 构建算法将 AST 转换为 Triton-IR。
LLVM-IR & PTX: 编译器后端对 IR 进行简化和优化，自动实现并行化，并将其转换为 LLVM-IR，最终转换为可在 NVIDIA GPU 上执行的 PTX。

编译器优化

自动内存暂存 (Automatic Memory Stashing): 编译器分析计算密集型块级操作（例如 tl.dot）的操作数，并使用活跃度分析 (liveness analysis) 自动将数据暂存到共享内存中。
自动并行化: Triton 通过并发运行不同的内核实例来实现跨 SMs 的并行化，并通过在 SIMD 单元内划分块级操作的迭代空间来实现 SMs 内部的并行化。