我一直觉得编译原理是一门被学院派教学耽误的硬核知识。很多人一听到“编译器”三个字第一反应就是那本厚厚的龙书以及里面晦涩的正规表达式、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 把ttirTriton IR转换成ttgirTritonGPU IR的过程就是典型的“IR 降级”lowering。第 9 周聚焦在third_party/nvidia/backend/compiler.py看 LLVM IR 是怎么生成的PTX 又是怎么来的。第 10 周我会建议你做一个完整的 profiling 任务分别在不同 GPU 上跑相同的 kernel打开TRITON_PRINT_AUTOTUNING1观察自动调优过程并且手动指定不同num_warps和num_stages分析结果差异。做完这个任务你会对 GPU 的线程模型、共享内存、寄存器压力有非常深刻的动手认知。3. 实操过程与核心环节实现3.1 环境准备与第一个可解剖的 kernel先交代我的实验环境方便你对齐。我用的是 Ubuntu 22.04Python 3.10CUDA 12.1PyTorch 2.1.0Triton 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(axis0) block_start pid * BLOCK_SIZE offsets block_start tl.arange(0, BLOCK_SIZE) mask offsets n_elements x tl.load(x_ptr offsets, maskmask) y tl.load(y_ptr offsets, maskmask) tl.store(z_ptr offsets, x y, maskmask) 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_SIZE1024) return z先从triton.jit说起。这个装饰器是 Triton 前端的起点它把下面的 Python 函数变成一个可以被编译器处理的 DSL 函数。当我们调用add_kernel[grid](...)时Triton 并不是立即执行函数体而是先把函数体编译成一个 GPU kernel。这个“编译”动作发生在首次调用时后续会走缓存。BLOCK_SIZE: tl.constexpr是 Triton 里一个很优雅的设计constexpr参数在编译期就必须确定因为它决定了线程块的大小直接影响后续所有 IR 的构建。在 Python 里理解这个很简单——你没法在运行时才决定一个数组的长度它必须在编译期就是固定的。tl.program_id(axis0)得到的是当前程序实例的 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( fnkernel.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_COMPILE1 TRITON_KERNEL_DUMP1 python my_add.pyTRITON_KERNEL_DUMP1会把编译过程中产生的所有中间文件输出到当前目录。你会在里面看到.ttir、.ttgir、.llir、.ptx等文件按 kernel 名和编译选项命名。拿 TTIR 文件来说你会发现 add_kernel 里的循环被消除了因为tl.arange的展开发生在编译期。这是一个非常好的“编译器到底在干什么”的入门例子。TTIR 里你还会看到tt.load、tt.store这样的操作它们和前端tl.load一一对应。这个阶段还没有做任何 GPU 相关的优化IR 里的维度信息还是“块级别”。3.3 深入 TTGIR一次让你对 GPU 调度改观的过程继续看 TTGIRTritonGPU IR会发现信息量和 TTIR 完全不是一个档次。它已经包含了 GPU 线程布局的信息比如每个操作被分配到哪个线程、用什么方式做数据重排。这里我强烈建议你刻意做一个小实验把同一份 kernel 分别用num_warps4和num_warps8编译然后对比 TTGIR 里的差异。你会发现 thread layout 的部分完全不同这就是调度器的功劳。很多人以为 GPU 编程只要会写 CUDA 就够了但实际上同样一段逻辑线程怎么排、数据怎么铺性能能差好几倍。Triton 把这个过程自动化了而读懂 TTGIR 能让你理解它究竟是怎么自动化的。调试 TTGIR 时有个小技巧我是在一次偶然中发现的直接用 Python 打印 compiled kernel 的属性。在triton.compiler.CompiledKernel里有个asm字典你可以这样看compiled add_kernel.warmup( torch.empty(1024, dtypetorch.float32, devicecuda), torch.empty(1024, dtypetorch.float32, devicecuda), torch.empty(1024, dtypetorch.float32, devicecuda), 1024, BLOCK_SIZE1024, 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操作这个操作会要求操作数遵循特定的布局比如 MMAmatrix 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(axis0) 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), dtypetl.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, maskc_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, maskmask, other-float(inf))我写成了otherfloat(-inf)结果语义分析报错。原因很简单——Triton 的类型检查认为other的值需要在编译期就能确定类型而-float(inf)这种表达式返回的是 Python 的 float它期望的是 tensor 标量。改成other-float(inf)其实底层逻辑没变但 Triton 的前端解析器能处理。这种细节只有你真正写了才会知道。后面我总结出一个习惯所有tl.load的other参数都明确写成一个常量或tl.zeros形状匹配的张量避免类型模糊。4.3 打开调试开关让编译器“开口说话”写 Triton kernel 时最难的不是功能正确而是性能不达预期时不知道瓶颈在哪。这里我分享三个环境变量都是我自己实测过最有效的TRITON_ALWAYS_COMPILE1强制每次运行都重新编译避免缓存导致你改了源码但跑的还是旧版。调试时一定要开否则你会被“改了没效果”折磨到疯。TRITON_PRINT_AUTOTUNING1打印自动调优过程能清楚看到它在尝试哪些num_warps、num_stages以及每个配置下的耗时。对理解 Triton 的启发式搜索很有帮助。MLIR_ENABLE_DUMP1如果编译过程中 MLIR 层报错这个开关可以把 MLIR 的运行日志打出来定位到是哪个 pass 出了问题。我自己的习惯是写一个新 kernel 时先开着TRITON_ALWAYS_COMPILE1跑一遍等逻辑确认没问题了再关掉它测真实性能。4.4 问题速查表我遇到过的 10 个典型报错下面这张表是我个人在 10 周学习过程中遇到的典型问题整理出来供你参考。这些报错信息会因版本不同略有差异但排查思路是通用的。报错关键词常见原因优先排查方向illegal memory access访存越界mask 没写对检查 offsets 范围、other参数TritonError: maximum block size exceededtl.arange长度超过编译上限降低BLOCK_SIZE或拆成多个循环Incompatible shapes in dottl.dot的两个操作数形状不匹配确保 K 维对齐检查广播维度UnknownTypeError类型推导失败某个变量类型不确定检查分支路径里变量是否一致确认constexprTriton assertion failedkernel 内部的tl.static_assert失败检查constexpr参数尤其是网格尺寸CUDA error: no kernel image available编译产物与当前 GPU 架构不匹配确认TORCH_CUDA_ARCH_LIST或 target 设置TritonError: invalid program_idgrid 维度设置错误检查grid的元组长度和tl.program_id(axis...)OutOfResources寄存器或共享内存超限减少num_warps简化 kernel 内临时变量LLVM ERROR: Cannot select后端不支持某些操作更新 Triton 版本或检查是否有不支持的 dtypePermission denied when writing dumpdump 目录没有写权限显式设置TRITON_CACHE_DIR到可写路径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 周最值钱的东西。
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 把ttirTriton IR转换成ttgirTritonGPU IR的过程就是典型的“IR 降级”lowering。第 9 周聚焦在third_party/nvidia/backend/compiler.py看 LLVM IR 是怎么生成的PTX 又是怎么来的。第 10 周我会建议你做一个完整的 profiling 任务分别在不同 GPU 上跑相同的 kernel打开TRITON_PRINT_AUTOTUNING1观察自动调优过程并且手动指定不同num_warps和num_stages分析结果差异。做完这个任务你会对 GPU 的线程模型、共享内存、寄存器压力有非常深刻的动手认知。3. 实操过程与核心环节实现3.1 环境准备与第一个可解剖的 kernel先交代我的实验环境方便你对齐。我用的是 Ubuntu 22.04Python 3.10CUDA 12.1PyTorch 2.1.0Triton 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(axis0) block_start pid * BLOCK_SIZE offsets block_start tl.arange(0, BLOCK_SIZE) mask offsets n_elements x tl.load(x_ptr offsets, maskmask) y tl.load(y_ptr offsets, maskmask) tl.store(z_ptr offsets, x y, maskmask) 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_SIZE1024) return z先从triton.jit说起。这个装饰器是 Triton 前端的起点它把下面的 Python 函数变成一个可以被编译器处理的 DSL 函数。当我们调用add_kernel[grid](...)时Triton 并不是立即执行函数体而是先把函数体编译成一个 GPU kernel。这个“编译”动作发生在首次调用时后续会走缓存。BLOCK_SIZE: tl.constexpr是 Triton 里一个很优雅的设计constexpr参数在编译期就必须确定因为它决定了线程块的大小直接影响后续所有 IR 的构建。在 Python 里理解这个很简单——你没法在运行时才决定一个数组的长度它必须在编译期就是固定的。tl.program_id(axis0)得到的是当前程序实例的 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( fnkernel.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_COMPILE1 TRITON_KERNEL_DUMP1 python my_add.pyTRITON_KERNEL_DUMP1会把编译过程中产生的所有中间文件输出到当前目录。你会在里面看到.ttir、.ttgir、.llir、.ptx等文件按 kernel 名和编译选项命名。拿 TTIR 文件来说你会发现 add_kernel 里的循环被消除了因为tl.arange的展开发生在编译期。这是一个非常好的“编译器到底在干什么”的入门例子。TTIR 里你还会看到tt.load、tt.store这样的操作它们和前端tl.load一一对应。这个阶段还没有做任何 GPU 相关的优化IR 里的维度信息还是“块级别”。3.3 深入 TTGIR一次让你对 GPU 调度改观的过程继续看 TTGIRTritonGPU IR会发现信息量和 TTIR 完全不是一个档次。它已经包含了 GPU 线程布局的信息比如每个操作被分配到哪个线程、用什么方式做数据重排。这里我强烈建议你刻意做一个小实验把同一份 kernel 分别用num_warps4和num_warps8编译然后对比 TTGIR 里的差异。你会发现 thread layout 的部分完全不同这就是调度器的功劳。很多人以为 GPU 编程只要会写 CUDA 就够了但实际上同样一段逻辑线程怎么排、数据怎么铺性能能差好几倍。Triton 把这个过程自动化了而读懂 TTGIR 能让你理解它究竟是怎么自动化的。调试 TTGIR 时有个小技巧我是在一次偶然中发现的直接用 Python 打印 compiled kernel 的属性。在triton.compiler.CompiledKernel里有个asm字典你可以这样看compiled add_kernel.warmup( torch.empty(1024, dtypetorch.float32, devicecuda), torch.empty(1024, dtypetorch.float32, devicecuda), torch.empty(1024, dtypetorch.float32, devicecuda), 1024, BLOCK_SIZE1024, 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操作这个操作会要求操作数遵循特定的布局比如 MMAmatrix 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(axis0) 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), dtypetl.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, maskc_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, maskmask, other-float(inf))我写成了otherfloat(-inf)结果语义分析报错。原因很简单——Triton 的类型检查认为other的值需要在编译期就能确定类型而-float(inf)这种表达式返回的是 Python 的 float它期望的是 tensor 标量。改成other-float(inf)其实底层逻辑没变但 Triton 的前端解析器能处理。这种细节只有你真正写了才会知道。后面我总结出一个习惯所有tl.load的other参数都明确写成一个常量或tl.zeros形状匹配的张量避免类型模糊。4.3 打开调试开关让编译器“开口说话”写 Triton kernel 时最难的不是功能正确而是性能不达预期时不知道瓶颈在哪。这里我分享三个环境变量都是我自己实测过最有效的TRITON_ALWAYS_COMPILE1强制每次运行都重新编译避免缓存导致你改了源码但跑的还是旧版。调试时一定要开否则你会被“改了没效果”折磨到疯。TRITON_PRINT_AUTOTUNING1打印自动调优过程能清楚看到它在尝试哪些num_warps、num_stages以及每个配置下的耗时。对理解 Triton 的启发式搜索很有帮助。MLIR_ENABLE_DUMP1如果编译过程中 MLIR 层报错这个开关可以把 MLIR 的运行日志打出来定位到是哪个 pass 出了问题。我自己的习惯是写一个新 kernel 时先开着TRITON_ALWAYS_COMPILE1跑一遍等逻辑确认没问题了再关掉它测真实性能。4.4 问题速查表我遇到过的 10 个典型报错下面这张表是我个人在 10 周学习过程中遇到的典型问题整理出来供你参考。这些报错信息会因版本不同略有差异但排查思路是通用的。报错关键词常见原因优先排查方向illegal memory access访存越界mask 没写对检查 offsets 范围、other参数TritonError: maximum block size exceededtl.arange长度超过编译上限降低BLOCK_SIZE或拆成多个循环Incompatible shapes in dottl.dot的两个操作数形状不匹配确保 K 维对齐检查广播维度UnknownTypeError类型推导失败某个变量类型不确定检查分支路径里变量是否一致确认constexprTriton assertion failedkernel 内部的tl.static_assert失败检查constexpr参数尤其是网格尺寸CUDA error: no kernel image available编译产物与当前 GPU 架构不匹配确认TORCH_CUDA_ARCH_LIST或 target 设置TritonError: invalid program_idgrid 维度设置错误检查grid的元组长度和tl.program_id(axis...)OutOfResources寄存器或共享内存超限减少num_warps简化 kernel 内临时变量LLVM ERROR: Cannot select后端不支持某些操作更新 Triton 版本或检查是否有不支持的 dtypePermission denied when writing dumpdump 目录没有写权限显式设置TRITON_CACHE_DIR到可写路径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 周最值钱的东西。