Triton GPU编程语言历史 | AI生成和翻译

Home 2026.06

问题:请介绍一下 GPU 编程语言 Triton 的历史。

回答:

Triton 最初是 Philippe Tillet 在哈佛大学的博士工作。基础论文为 《Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations》(MAPL 2019)。核心论点:GPU 内核编写处于一个尴尬的极端——CUDA 提供了控制力,但迫使你手动管理共享内存、内存合并和 warp 级调度;高层 DSL 隐藏了硬件细节,但无法表达新颖操作。Triton 的赌注在于一种 基于 tile 的抽象:你以数据块(tiles)为单位进行编程,编译器负责处理 tile 内部的内存布局、合并和 SM 调度。

大致时间线:

值得内化的编程模型洞察:在 CUDA 中,你以 线程 粒度进行推理(每个线程计算一个元素,你编排网格)。在 Triton 中,你以 块/tile 粒度进行推理——单个 @triton.jit“程序实例”操作整个 tile,你通过 tl.load/tl.store 加上掩码来表达加载/存储范围。编译器决定线程映射。这就是为什么一个 Triton 的 matmul 只需约 30 行代码,而一个具有竞争力的 CUDA 版本需要数百行——你把 warp 级调度和共享内存分阶段的工作委托给了编译器。

一个最小化的内核来锚定心智模型——向量加法:

import triton
import triton.language as tl
import torch

@triton.jit
def add_kernel(x_ptr, y_ptr, out_ptr, n, BLOCK: tl.constexpr):
    pid = tl.program_id(0)                      # 我负责哪个 tile
    offs = pid * BLOCK + tl.arange(0, BLOCK)    # 这个 tile 的元素索引
    mask = offs < n                             # 保护尾部未对齐部分
    x = tl.load(x_ptr + offs, mask=mask)
    y = tl.load(y_ptr + offs, mask=mask)
    tl.store(out_ptr + offs, x + y, mask=mask)

def add(x, y):
    out = torch.empty_like(x)
    n = x.numel()
    grid = lambda meta: (triton.cdiv(n, meta["BLOCK"]),)
    add_kernel[grid](x, y, out, n, BLOCK=1024)
    return out

注意缺失了什么:没有 threadIdx,没有 __shared__,没有显式的内存合并。tl.arange 配合 BLOCKmask 就是整个分块的故事。BLOCK: tl.constexpr 是编译期特化的参数——Triton 会对其进行自动调优。

如今 Triton 的地位:当 cuBLAS/cuDNN/FlashAttention 尚未覆盖你需要的操作时,Triton 是机器学习中事实上的 内核编写层。FlashAttention-2/3、融合 MoE 内核(与你对 DeepSeek v4 MoE 的探索直接相关——专家路由的 gather/scatter 和 grouped GEMM 正是人们在 Triton 中编写的那种操作),以及现代训练栈中大多数自定义融合操作,都是 Triton 或与 Triton 相关的。它之所以战胜“直接写 CUDA”的战略原因在于:它将内核逻辑与硬件特定的调度解耦,因此同一个内核可以在不同 SM 世代乃至不同供应商之间自动调优。

对于你个人的发展轨迹而言:Triton 是目前最具杠杆效应的 GPU 技能,因为(1)它是 torch.compile 的基础,(2)它是无需 CUDA 的 NVIDIA 锁定就能编写 MI300X 内核的实用路径,(3)用 Triton 编写融合注意力或 MoE 内核,能让你学到内存层级推理(SRAM 分块、重新计算与物化),这种推理可以迁移到理解 为什么 FlashAttention 在数学层面有效。

参考文献:

我根据已有信息而非搜索给出了这段历史,因此具体日期请视为近似值——特别是 MLIR 重写和 ROCm 成熟的时间线在多个版本间有所变化。如果你需要,我可以拉取精确的逐版本更新日志或 AMD 后端的当前状态。


Back Donate