10周玩转Triton编译器：从零解剖GPU编程与编译原理

我一直觉得，编译原理是一门被学院派教学耽误的硬核知识。很多人一听到“编译器”三个字，第一反应就是那本厚厚的龙书，以及里面晦涩的正规表达式、LR分析表和寄存器分配算法。这种劝退感我太懂了，因为我一开始也是这么被吓跑的。直到我遇到 OpenAI 开源的 Triton 编译器，才发现原来编译器完全可以换一种学法：不啃龙书，而是找一台“活”的编译器，把它拆开揉碎，从用起来到改起来，一步步刨根问底。

Triton 是一个面向 GPU 编程的领域专用语言（DSL）和编译器，最初是为了让深度学习算子开发变得更简单而设计的。你只需要写看起来很像 Python 的核函数（kernel），它就能帮你自动完成线程调度、显存访问优化、寄存器分配等工作，最终生成高效的 GPU 机器码。但 Triton 真正的妙处在于，它的编译链路极其完整且模块清晰：从 Python AST 到 Triton IR，再到 MLIR、LLVM IR、PTX，最后落到 SASS。这意味着，一个普通开发者只要花心思，就能在一门真实运行的编译器上，把教科书里前端、中端、后端那套东西全过一遍，而且完全不需要啃龙书。

这篇文章想分享的就是我总结出来的 10 周“刨根问底”计划，我给它起了个外号叫“活体编译器解剖课”。它适合什么人群呢？一类是想系统学编译器但被理论劝退的开发者，另一类是用过 PyTorch 或 CUDA、想深入理解 GPU 算子底层原理的算法工程师。我会从为什么选 Triton 说起，再到 10 周学习路径的拆解、具体实操怎么做，最后把踩过的坑和排查方法全列出来。保证每一条都是我在 Linux 环境下真实跑过、验证过的经验。

1. 为什么选 Triton 而不是先啃龙书

1.1 龙书不是不好，而是它把你按在原地不动

先声明一下，我不是否定龙书的价值。编译原理是计算机科学的地基，龙书里讲的前端词法分析、语法分析、语义分析，中端的数据流分析、循环优化，后端的指令选择、寄存器分配，这些概念到今天依然是所有编译器的骨架。问题是，龙书的知识密度太高，而且大多数人在读它的时候，并没有一个真实的编译需求悬在头上。没有需求牵引，知识就是一盘散沙，读完后端的寄存器分配，你根本不知道它到底解决了什么问题，过两周全忘了。

我自己就是典型反例。学生时代啃了三个月的龙书，能做题，能画状态转换图，甚至能默写 LR(1) 分析表构造步骤。但你让我写一个能用的词法分析器？我第一反应还是掏出 flex。这种“学了一堆概念但做不出东西”的状态，其实非常打击人。所以后来我带团队带新人，一直坚持一个原则：想学编译原理，先找一台真正在用的编译器，又要小、又要完整、又要好读，最好还有很强的现实价值。Triton 刚好完美命中这三个条件。

1.2 Triton 的独特定位：小到能读懂，大到能实战

Triton 不像 LLVM 和 GCC 那样是几十万行到几百万行代码的庞然大物，它的核心代码量相对可控，尤其是 Python 侧的前端和驱动部分，一个人通读下来并非不可能。但它又不是玩具。OpenAI 用它支撑了 ChatGPT 训练推理里的很多高性能算子，社区里还有 FlashAttention、FlashDecoding 等一系列里程碑式的工作是基于 Triton 实现的。换句话说，你学它不只是“学一个玩具编译器”，你学的是当前 AI Infra 领域最前沿的实用工具。

更重要的是，Triton 的编译过程非常透明。你可以非常轻易地把一个 kernel 从 Python 源码一路“解剖”到 PTX 甚至 SASS。这种透明性是学习编译器设计思想最好的土壤。我第一次运行 triton.compile 并看到它输出的 TTIR、TTGIR、LLVM IR 一长串中间表示时，脑子里很多模糊的概念瞬间就串起来了。原来一个程序在进入后端之前，会被拆成这么多个层次，每一层都有自己的抽象和目标。

1.3 10 周计划的总体思路：把编译过程变成一条可走的“探险路线”

这 10 周不是让你读源码读到吐，而是把 Triton 的编译过程当成一条主线，逐步往里走，每一周都有明确目标和产出。我把它分成三个阶段：

• 第一阶段（第 1~3 周）：建立整体认知，把 Triton 用起来，搞清楚“用户写的东西最后变成了什么”。
• 第二阶段（第 4~7 周）：攻克编译器前端，看 Python AST 如何一步步变成 Triton IR，理解类型推导、AST 语义分析这些概念在真实代码里长什么样。
• 第三阶段（第 8~10 周）：转战后端，探索 MLIR 方言、Pass 管线、GPU 调度和代码生成，真正理解“为什么同一个 kernel 在不同 GPU 上性能差异那么大”。

整个计划只有一个核心方法论：每次读到新概念，都要立刻回到一个具体的 kernel 上，用代码验证它、改变它、观察结果变化。这样你学到的不是孤立的知识点，而是可迁移的思考方式。

2. 10 周实战计划的进阶拆解

2.1 前三周：跑通工具链，建立 IR 的“体感”

我见过很多人学编译原理，第一个坎不是概念难，而是工具链没跑通。所以这个计划的前三周，我建议把所有精力花在“让 Triton 跑起来”和“学会看 IR”上。

第 1 周的任务很简单：安装 Triton，跑通第一个向量加法 kernel，理解 grid 和 block 这两个核心概念。这里要提醒一点，Triton 的安装版本和 PyTorch 的 CUDA 版本必须匹配，否则你会被各种“非法内存访问”折磨到怀疑人生。我的建议是用官方预编译包，不要一上来就自己从源码编译。源码编译涉及 LLVM 版本、CUDA Toolkit、Ninja 一堆依赖，第一周就搞这个，大概率会劝退。

第 2 周可以开始“玩 IR”。Triton 有一个非常友好的特性：编译结果里每个阶段的中间表示都能导出来看。你可以写一个最简单的 kernel，然后用 triton.compile 拿到它的 TTIR 和 TTGIR，一行一行对照 Python 源码去理解。第 3 周再做个小项目：手写一个 softmax kernel，和 PyTorch 的 F.softmax 做性能对比。通过这个对比，你会第一次直观感受到“编译器的调度优化到底是什么”，也能对 grid、block、num_warps 这些参数产生直观认识。

2.2 中间四周：深入 Python 前端，搞懂类型推导和 AST 语义

很多人以为“编译器前端”就是写正则表达式和递归下降解析器，但实际上，Triton 这种嵌入式 DSL 的前端是另一条路：它直接在 Python AST 上做语义分析和类型推导。这个过程非常精彩，也非常值得研究。

第 4 周，我会带你读python/triton/language/core.py和python/triton/language/semantic.py。你会发现，tl.load、tl.store、tl.arange这些看起来普通的函数，其实并不是真正执行的操作，而是构建了 IR 节点。这就像是在玩积木：你写一行 Python，它其实是在搭一棵语法树。第 5 周深入triton.compiler的compile函数，看它如何处理 Python 函数对象、如何用inspect拿到源码、如何把 AST 转换成 Triton IR。第 6 周和第 7 周可以做两个横向对比：第一个是自定义一个tl操作符（或者扩展一个现有操作符），第二个是给 Triton 前端写一个很小的 AST 变换插件，比如自动把某种模式的乘法替换成移位操作。这些任务做完，你对编译器前端的理解会超过很多只啃龙书的人。

2.3 最后三周：攻后端，理解 GPU 编译器调度和代码生成

后面的内容是整个计划里最有挑战也最有收获的部分。Triton 的后端构建在 MLIR 之上，GPU 相关的优化集中在 TritonGPU dialect 里。这个概念我不要求你在第 8 周的一开始就能完全理解，但它很重要。

第 8 周，从 MLIR 的基本概念入手，理解方言（dialect）、操作（operation）、Pass 这些名词。Triton 把ttir（Triton IR）转换成ttgir（TritonGPU IR）的过程，就是典型的“IR 降级”（lowering）。第 9 周，聚焦在third_party/nvidia/backend/compiler.py，看 LLVM IR 是怎么生成的，PTX 又是怎么来的。第 10 周，我会建议你做一个完整的 profiling 任务：分别在不同 GPU 上跑相同的 kernel，打开TRITON_PRINT_AUTOTUNING=1观察自动调优过程，并且手动指定不同num_warps和num_stages，分析结果差异。做完这个任务，你会对 GPU 的线程模型、共享内存、寄存器压力有非常深刻的动手认知。

3. 实操过程与核心环节实现

3.1 环境准备与第一个可解剖的 kernel

先交代我的实验环境，方便你对齐。我用的是 Ubuntu 22.04，Python 3.10，CUDA 12.1，PyTorch 2.1.0，Triton 2.1.0。这个组合不是必须的，但如果你完全没接触过 Triton，建议别用太新的版本，社区里坑更少。

安装这一步其实很直白：

pip install triton

装完以后，先写一个最小的验证程序。我给你看的第一份代码，不是能跑就行，而是要留出后续“解剖”的空间：

import torch import triton import triton.language as tl @triton.jit def add_kernel(x_ptr, y_ptr, z_ptr, n_elements, BLOCK_SIZE: tl.constexpr): pid = tl.program_id(axis=0) block_start = pid * BLOCK_SIZE offsets = block_start + tl.arange(0, BLOCK_SIZE) mask = offsets < n_elements x = tl.load(x_ptr + offsets, mask=mask) y = tl.load(y_ptr + offsets, mask=mask) tl.store(z_ptr + offsets, x + y, mask=mask) def run_add(x, y): z = torch.empty_like(x) n_elements = x.numel() grid = (triton.cdiv(n_elements, 1024),) add_kernel[grid](x, y, z, n_elements, BLOCK_SIZE=1024) return z

先从@triton.jit说起。这个装饰器是 Triton 前端的起点，它把下面的 Python 函数变成一个可以被编译器处理的 DSL 函数。当我们调用add_kernel[grid](...)时，Triton 并不是立即执行函数体，而是先把函数体编译成一个 GPU kernel。这个“编译”动作发生在首次调用时，后续会走缓存。

BLOCK_SIZE: tl.constexpr是 Triton 里一个很优雅的设计：constexpr参数在编译期就必须确定，因为它决定了线程块的大小，直接影响后续所有 IR 的构建。在 Python 里理解这个很简单——你没法在运行时才决定一个数组的长度，它必须在编译期就是固定的。

tl.program_id(axis=0)得到的是当前程序实例的 ID。想象 GPU 被分成很多个块，每个块执行同样一份代码，program_id就是告诉你“你现在是第几块”。tl.arange(0, BLOCK_SIZE)会在编译期生成一个从 0 到 BLOCK_SIZE 的整数序列，它对应的是一组并行的线程索引。这跟你手动写 CUDA 时计算threadIdx.x + blockIdx.x * blockDim.x是一个意思，但 Triton 帮你把这些细节藏起来了。

3.2 从 Python 到 TTIR：源码解剖的第一站

现在到了核心环节。我要用一个小脚本把 add_kernel 的中间表示导出来，这个过程我建议你亲手跑一遍，因为它是你建立“IR 体感”最重要的一步：

import triton from triton.compiler import compile # 获取经过 jit 封装的 kernel 函数对象 kernel = add_kernel # 构造一个当前 kernel 的编译 key 信息 src = triton.compiler.ASTSource( fn=kernel.fn, signature="*fp32,*fp32,*fp32,i32,constexpr", # 三个指针 + int + constexpr constexprs={"BLOCK_SIZE": 1024}, ) compiled = compile(src, target=("cuda", 0))

实际上更常用的方式是用kernel.warmup或者直接跑一次run_add，然后在环境变量里打开 dump 开关：

TRITON_ALWAYS_COMPILE=1 TRITON_KERNEL_DUMP=1 python my_add.py

TRITON_KERNEL_DUMP=1会把编译过程中产生的所有中间文件输出到当前目录。你会在里面看到.ttir、.ttgir、.llir、.ptx等文件，按 kernel 名和编译选项命名。

拿 TTIR 文件来说，你会发现 add_kernel 里的循环被消除了，因为tl.arange的展开发生在编译期。这是一个非常好的“编译器到底在干什么”的入门例子。TTIR 里你还会看到tt.load、tt.store这样的操作，它们和前端tl.load一一对应。这个阶段还没有做任何 GPU 相关的优化，IR 里的维度信息还是“块级别”。

3.3 深入 TTGIR：一次让你对 GPU 调度改观的过程

继续看 TTGIR（TritonGPU IR），会发现信息量和 TTIR 完全不是一个档次。它已经包含了 GPU 线程布局的信息，比如每个操作被分配到哪个线程、用什么方式做数据重排。

这里我强烈建议你刻意做一个小实验：把同一份 kernel 分别用num_warps=4和num_warps=8编译，然后对比 TTGIR 里的差异。你会发现 thread layout 的部分完全不同，这就是调度器的功劳。很多人以为 GPU 编程只要会写 CUDA 就够了，但实际上，同样一段逻辑，线程怎么排、数据怎么铺，性能能差好几倍。Triton 把这个过程自动化了，而读懂 TTGIR 能让你理解它究竟是怎么自动化的。

调试 TTGIR 时有个小技巧，我是在一次偶然中发现的：直接用 Python 打印 compiled kernel 的属性。在triton.compiler.CompiledKernel里有个asm字典，你可以这样看：

compiled = add_kernel.warmup( torch.empty(1024, dtype=torch.float32, device="cuda"), torch.empty(1024, dtype=torch.float32, device="cuda"), torch.empty(1024, dtype=torch.float32, device="cuda"), 1024, BLOCK_SIZE=1024, grid=(1,), ) print(compiled.asm.keys()) print(compiled.asm["ttir"]) print(compiled.asm["ttgir"])

warmup会触发编译但不真正执行，非常适合拿来做 IR 审查。这些输出的可读性非常好，MIT 许可证下你也可以直接把它复制到文本编辑器里慢慢研究。我当时就是把这些 IR 打印出来贴在办公桌前面，每天看一点，一周后 TTGIR 对我来说就不再是天书了。

3.4 用tl.dot撬动矩阵乘法，理解优化的关键战场

向量加法只是开胃菜。如果你想在 10 周内真正“刨根问底”，矩阵乘法是绕不开的核心场景。Triton 的tl.dot专门用于矩阵乘法的硬件加速，它背后映射到 GPU 上的 Tensor Core。

我建议你第 5 周左右写一个matmul_kernel，然后观察它在 TTGIR 里是如何被布局的。你会发现tl.dot会触发一个独立的tt.dot操作，这个操作会要求操作数遵循特定的布局，比如 MMA（matrix multiply-accumulate）布局。如果布局不匹配，编译器会插入tt.convert_layout操作做数据重排。看懂了这些，你就理解了为什么有些矩阵乘法实现快，有些慢。不是算法复杂度不同，而是在硬件上的数据摆放方式不同。

实际的矩阵乘法核心长这样（简化版）：

@triton.jit def matmul_kernel( a_ptr, b_ptr, c_ptr, M, N, K, stride_am, stride_ak, stride_bk, stride_bn, stride_cm, stride_cn, BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr, ): pid = tl.program_id(axis=0) num_pid_m = tl.cdiv(M, BLOCK_SIZE_M) num_pid_n = tl.cdiv(N, BLOCK_SIZE_N) pid_m = pid // num_pid_n pid_n = pid % num_pid_n offs_m = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M) offs_n = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N) offs_k = tl.arange(0, BLOCK_SIZE_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 acc = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32) for k in range(0, K, BLOCK_SIZE_K): a = tl.load(a_ptrs) b = tl.load(b_ptrs) acc = tl.dot(a, b, acc) a_ptrs += BLOCK_SIZE_K * stride_ak b_ptrs += BLOCK_SIZE_K * stride_bk offs_cm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M) offs_cn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N) c_ptrs = c_ptr + offs_cm[:, None] * stride_cm + offs_cn[None, :] * stride_cn c_mask = (offs_cm[:, None] < M) & (offs_cn[None, :] < N) tl.store(c_ptrs, acc, mask=c_mask)

这个例子值得你反复琢磨，尤其是tl.arange和[:, None]在广播维度上的使用。它们决定了每个线程块内的“视图”，也就是从全局内存里看到的一块子矩阵。理解了这个 padding 和 mask 的应用场景，你就掌握了大算子开发的通用套路。

3.5 动手改造 IR：在 Pass 里做一次“小动作”

学编译器如果不亲手改一次 Pass，总觉得差了点什么。Triton 的 Python 侧给了我们一个必然可行的入口：在编译时插入自定义 Pass。虽然 Triton 的 pass 管线是用 MLIR 的 C++ 实现的，但你可以借助triton.compiler.compiler的流程，自己编写一个“复制优化 pass”的 Python 版本，专门处理 AST 阶段的改写。

一个最简单的实验是，把tl.load的 mask 参数做一个自动填充。例如，当你没显式指定 mask 时，自动推断访问范围并生成对应的 mask。这个过程在语义分析阶段完成，实现起来不需要深入了解 MLIR 的 API，只需要在semantic.py里加一个 AST 改写步骤。这个操作看起来小，但跑通它，你就明白了“编译器插桩”是一种什么样的体验。我在第 6 周做了类似的事情，整个下午都在调试 AST 的 Visiter，最后真的把一个小优化加进去跑通了，那种成就感堪比第一次让 CUDA kernel 跑出正确结果。

4. 常见问题与排查技巧实录

4.1 版本不匹配导致的“非法内存访问”陷阱

这是所有 Triton 新手最容易踩的坑，没有之一。Triton 对 CUDA 版本和 PyTorch 版本非常敏感。如果你用 PyTorch 编译时带的 CUDA 是 11.8，但当前系统默认的nvcc是 12.3，编译出的 PTX 可能无法在当前驱动上加载，表现就是程序直接崩溃或者报CUDA error: illegal memory access was encountered。

我自己的排查方法是三步走：先确认torch.version.cuda和triton.__version__是否匹配；再检查nvidia-smi里的驱动版本是否支持你用的 CUDA；最后在代码里加一个torch.cuda.synchronize()，把异步错误固定到具体行。这三个步骤能解决至少 80% 的“感觉是 WSL 不稳定”的问题。需要注意的是，不要光看本地编译结果，GPU 驱动和 CUDA runtime 的兼容矩阵一定要先查清楚。

4.2 类型推导报错：小改动引发的前端连锁反应

Triton 前端做类型推导时，严格得让人崩溃。最常见的报错是TypeError: unsupported operand type(s) for +: 'Tensor' and 'NoneType'之类。这时候不要慌，先检查是不是某个变量在某个分支里没有定义，或者某个 mask 的 shape 和你 load 的数据不匹配。

我第二次写 softmax 时就遇到过：tl.load(ptr, mask=mask, other=-float('inf'))，我写成了other=float('-inf')，结果语义分析报错。原因很简单——Triton 的类型检查认为other的值需要在编译期就能确定类型，而-float('inf')这种表达式返回的是 Python 的 float，它期望的是 tensor 标量。改成other=-float('inf')其实底层逻辑没变，但 Triton 的前端解析器能处理。这种细节，只有你真正写了才会知道。后面我总结出一个习惯：所有tl.load的other参数都明确写成一个常量或tl.zeros形状匹配的张量，避免类型模糊。

4.3 打开调试开关，让编译器“开口说话”

写 Triton kernel 时，最难的不是功能正确，而是性能不达预期时不知道瓶颈在哪。这里我分享三个环境变量，都是我自己实测过最有效的：

• TRITON_ALWAYS_COMPILE=1：强制每次运行都重新编译，避免缓存导致你改了源码但跑的还是旧版。调试时一定要开，否则你会被“改了没效果”折磨到疯。
• TRITON_PRINT_AUTOTUNING=1：打印自动调优过程，能清楚看到它在尝试哪些num_warps、num_stages，以及每个配置下的耗时。对理解 Triton 的启发式搜索很有帮助。
• MLIR_ENABLE_DUMP=1：如果编译过程中 MLIR 层报错，这个开关可以把 MLIR 的运行日志打出来，定位到是哪个 pass 出了问题。

我自己的习惯是，写一个新 kernel 时，先开着TRITON_ALWAYS_COMPILE=1跑一遍，等逻辑确认没问题了，再关掉它测真实性能。

4.4 问题速查表：我遇到过的 10 个典型报错

下面这张表是我个人在 10 周学习过程中遇到的典型问题，整理出来供你参考。这些报错信息会因版本不同略有差异，但排查思路是通用的。

4.5 学习方法层面的三个独家心得

前面讲的是技术排错，但我觉得这 10 周学习过程中最值钱的其实是学习方法上的调优。分享三个我复盘后认为最关键的心得：

第一，每学一个新的 IR，一定要用“翻译”的方式把它变成人话。我读 TTIR 时，会在旁边用中文写注释，把tt.load(%ptr, %offset, %mask)翻译成“从 %ptr 这个地址，按 %offset 这个偏移量去读，只有 %mask 为真的位置才有效”。这样翻译过几个算子之后，再看那些 IR 文件就非常有亲切感。

第二，尽量用性能对比来驱动源码阅读。不要单纯为了读代码而读代码。比如我读完tl.dot相关源码后，立刻写了一个朴素的matmul和一个用tl.dot的版本，在 A100 上对比 FLOPS。当看到 10 倍以上的性能差距后，我就特别想搞清楚背后的原因，读起调度代码来也格外有动力。

第三，把 Git 历史当成教科书。Triton 更新很快，很多设计决策在源码注释里根本没有，但你在 commit message 和 PR 描述里能看到。比如你疑惑为什么某个 pass 要那么写，去git log里翻一翻，常常能找到当年修复的性能问题或者 bug 背景。这种“考古式”学习，比看任何文档都有用。

5. 10 周之后：接下来还能怎么深入

如果你按照前面的计划走完 10 周，基本上对 Triton 的整个编译流程已经有了完整的动手认知。但学习这件事，最怕的是拿到一点成果就停下来。我个人实际体验下来，真正让我从“会用 Triton”变成“敢改 Triton”的，是在计划结束后的两三周内做的几件事，这里一并分享给你。

第一件是尝试为一个小众 GPU 写后端。Triton 的抽象比较干净，你可以照着third_party/nvidia的目录结构，用最小的代价把 kernel 编译到一个新的 target 上。不需要真的能跑，哪怕只是接到 MLIR 的 CPU 执行引擎上，也能极大加深对“后端”这个概念的理解。

第二件是自己做一套 IR 可视化工具。官方有triton-viz和 timemory 的部分工具，但自己写一个只在你的实验 kernel 上工作的可视化脚本，对理解 kernel 的执行行为帮助特别大。我当时写了一个把 TTGIR 转成 dot 格式的小工具，配合 Graphviz 看图，很多布局迁移问题一眼就明白了。

第三件是源码贡献。Triton 社区的 issue 区里有很多good first issue标签的任务，大多涉及错误处理、边界测试和文档补全。修一个小 issue 的过程，就是一次完整的高质量代码阅读训练。我修复过第一个和文档相关的 PR，虽然只是核对了一处公式，但为了确认它是否正确，我几乎把相关源码都读了一遍，收获远超这个 PR 本身的价值。

回看这 10 周的经历，我最深的体感是：编译原理不是“啃”出来的，是“用”出来的。你不需要先变成理论大师再动手，只需要选一台像 Triton 这样小而美的现代编译器，顺着它的源码往下挖，挖到哪里算哪里。有的地方你会卡住，有的地方你会恍然大悟，但每走一步，你对“程序到底是怎么变成机器指令的”这个概念，都会比前一天更清晰。这个过程带来的底层认知升级，才是这 10 周最值钱的东西。