写这篇东西的起因其实挺简单团队里来了个新人深度学习跑得很溜但一提到编译器就发怵桌上那本红黑封面的“龙书”翻了两个月还在第2章整天被词法分析和正则表达式按在地上摩擦。我说你别死磕了换个思路咱们直接用Triton编译器当教材边用边拆边拆边学。十周下来他不仅看懂了IR、认识了Pass还能给自己的算子写自动调优顺手还改了个小优化传给上游社区。这篇稿子就是把这十周的路线、思路和踩过的坑整理一遍给同样想在编译器上“刨根问底”但不想从龙书第一章开始硬啃的朋友做个参考。Triton是OpenAI开源的GPU编程语言与编译器主要面向深度学习场景里的高性能自定义算子。它在PyTorch生态里已经非常流行写kernel的体验比直接写CUDA C简单得多性能却基本能打平甚至反超手写CUDA。更妙的是它内部包了一条相对完整又贴近现代工业实践的编译流水线——Python前端、中间表示IR、多层Pass、LLVM后端、PTX生成——几乎是一个“最小可运行”的现代编译器样板。拿它当编译器启蒙教材既能快速看到实际产出又能顺着代码一层层往下挖天然就适合“刨根问底”型的学习方式。下面我把这套“不啃龙书、以Triton为蓝本”的十周路线完整拆开包含每一周做什么、为什么这么做、实际会遇到哪些坑以及我在带人过程中积累的一些独家心得。1. 为什么选Triton当编译器教材而不是先啃龙书1.1 龙书的困境学术正确但离“跑起来”太远龙书《Compilers: Principles, Techniques, and Tools》在编译器领域确实是经典中的经典词法分析、语法分析、语义分析、中间代码生成、优化、目标代码生成体系完整理论密度极高。但问题也恰恰出在“完整”上它是在教你建一座可以住几十年的宫殿而大多数人的真实需求是今晚先有个能做饭的厨房。我见过太多人被卡死在第2章。词法分析里正则表达式、NFA、DFA、最小化DFA一套组合拳下来还没见到真正的编译器长什么样人已经跑了。不是说这些理论没用而是对绝大多数做应用层开发、深度学习框架、性能优化的人来说理论视角和工程实践之间隔着一道巨大的鸿沟。等你按龙书的节奏把词法分析搞完一年过去了你依然不知道自己每天用的PyTorch里那个torch.compile到底在干什么。龙书适合什么人适合准备做编译器学术研究、或者要在一家大公司里从零开始造一门语言的人。它的知识颗粒度是按“你未来二十年都要靠这行吃饭”的标准设计的。但对于只想搞懂“IR长什么样、优化Pass怎么运作、我的算子为什么快/慢”的普通工程师来说投入产出比太低了。1.2 Triton的独特价值现代工业级编译器的“解剖样本”Triton和龙书面对的问题完全不同。龙书教你解决“通用编程语言的编译”Triton面对的则是“领域特定语言DSL的编译”——只服务深度学习算子语言范围小、约束明确、优化目标单一GPU吞吐最大化。这种简化恰恰是学习者的福音因为复杂度被控制在了一个可操作的范围里。打个比方龙书是让你学怎么造一整架波音747Triton是让你拆一台摩托车发动机。747拆错了是要出人命的而摩托车的发动机你拆完了装回去大概率还能跑就算装不回去也能很直观地看到每个零件原本是干什么用的。Triton正是这样一台“能跑起来的摩托车发动机”。具体来说Triton的编译流程包含几个关键层级Triton Python源程序 ↓ [前端AST解析 类型推断] Triton IRTTIR ↓ [中端优化布局分配、缓存修改、循环展开等Pass] Triton GPU IRTTGIR ↓ [后端LLVM IR生成] LLVM IR ↓ [LLVM Pass优化生成目标代码] PTXGPU汇编 ↓ [由显卡驱动编译为SASS最终执行]这一整套流程在代码里全部真实存在并且所有中间产物都能通过环境变量轻松导出查看。学习者完全可以“看着每一层的变化”来理解编译器的工作过程。这体验比龙书里那些抽象的图生动太多了。1.3 两条学习路径的对比到底谁更适合谁对比项龙书路线Triton实战路线入门门槛需要大量理论积累第2章就能劝退大半人会写Python能看懂简单的PyTorch代码即可见效速度几周甚至几个月后才见到“编译器”的影子第一周就能写出并运行第一个Triton kernel亲眼看到IR覆盖范围完整覆盖前端、中端、后端理论前端弱化Triton已封装好中端Pass和后端优化直接可见与实际工业界的关系理论模型和LLVM、GCC等实际实现差别较大本身就是生产级编译器学的就是天天在跑的工业化实现适合人群准备长期深耕编译器研究的开发者深度学习工程师、高性能计算开发者、系统方向的实践者我个人的观点非常直白如果你只是想“搞懂现代编译器是怎么工作的”“GPU算子为什么能并行”“编译器优化Pass到底在优化什么”那Triton这条实战路线效率是龙书的五倍以上。如果你未来真要去做一门新语言的编译器、要去改LLVM后端那龙书迟早要补但完全可以先通过Triton把全局轮廓建立起来回头再啃理论时效率反而更高。我甚至觉得先实战后理论才是更适合今多数人的路径。2. 10周学习路径总览不啃书但把编译器“拆”着学2.1 十周路线图每周一个主题产出可检验这套路线我在实际带人时用过两轮针对不同基础的学员做过微调。整体思路是一样的先跑起来再看中间产物再读源码再动手改最后做自动调优。周次学习主题核心任务可检验产出第1周Triton环境搭建与基本语法安装Triton跑通官方教程手写第一个向量加法kernel一个可运行的Triton kernel与对应PyTorch对比第2周Triton Kernel与GPU并行模型理解Grid、Block、Thread三层结构掌握tl.arange、tl.load/store、块内并行语义手写矩阵乘法kernel并验证正确性第3周编译器的“第一现场”TTIR与TTGIR通过环境变量导出IR逐行对比源程序与IR的对应关系导出矩阵乘法的TTIR和TTGIR并在源码注释对应第4周核心概念Tile、布局与块级并行搞懂Triton的核心抽象——Tile数据块理解布局分配和线程映射画出矩阵乘法Tile划分的示意图解释数据如何映射到线程第5周优化Pass初探以布局分配为例读源码中布局分配相关Pass理解编译器如何自动为数据分配布局能手动推导一个简单kernel的布局分配结果第6周自动调优机制Autotune与Config理解Triton的自动调优原理写一个带autotune的kernel观察不同配置的性能差异完成一个自动调优矩阵乘法并绘制性能对比曲线第7周更深的编译后端LLVM IR与PTX导出LLVM IR和PTX了解GPU汇编的基本结构能指出PTX中ld.global、fma.rn等指令对应源程序的哪一行第8周算子优化实战向量化与内存访问优化一个实际算子尝试调整Block大小、num_warps、num_stages理解它们对性能的影响性能达到手写CUDA的90%以上或至少明显超越朴素实现第9周阅读Triton源码Compiler.py到Pass管线精读triton/compiler.py和核心Pass代码画出Triton编译链条的完整流程图能对他人讲清楚一条Triton源程序从Python到PTX的完整旅程第10周动手改编译器加一个自定义Pass或扩展算子在Triton源码里新增一个小Pass或者给前端增加一个自定义内建函数完成一个可运行的自定义Pass demo或新算子这张表看着密集但实际操作下来每周投入大约10到15小时是足够的。如果是有CUDA基础的人第1周到第4周可以压缩到一个半星期整个周期甚至能缩到六到八周。2.2 为什么是这种节奏先“看到”再“读懂”最后“改”这套路线的设计逻辑和传统的“先学理论、再做练习”完全相反是一种自底向上的“逆向工程”式学习体验。前三周的核心目标只有一个让学习者对编译器不恐惧。很多人学不下去编译器最大的障碍不是难而是“不知道自己不知道什么”——根本想象不出编译器内部长什么样。第1周直接写出并运行一个kernel第2周理解GPU并行模型第3周导出IR到文件里看这三步走完学习者脑海里就有了一个具体的画面后续的所有学习都是在往这个画面上补充细节而不是面对一团迷雾去猜测。中间四周第4到第7周是对编译器核心内容的系统拆解。这一段是整个路径里信息密度最高的部分。布局分配、优化Pass、LLVM后端这一套东西在龙书里对应的是“存储布局”“代码优化”和“目标代码生成”三大篇章但龙书用抽象的数学语言讲Triton直接给你看活生生的实现。这个时候学习者已经能把“编译器内部的某次具体优化”对应到“代码的哪一个Pass改了哪些IR”。最后三周第8到第10周是化被动为主动的阶段。前面的学习是在读别人的代码从第8周开始是在读“我的算子在编译器眼里长什么样”。性能调优、读源码、改Pass一环扣一环本质上是让学习者从“用户”变成“开发者”。到这里我那个新人的变化特别明显他不再说“编译器好难”而是开始说“这个Pass如果改成这样会不会更优”——这就是质变的信号。2.3 一个容易被忽视的准备工作选对参考资料这十周我已经说了不啃龙书但完全不看书也不太现实。我的核心建议是把这个阶段要用的资料控制在三份以内避免陷入资料的海洋。第一份是Official Triton Tutorial这是最基础的入门资料。第二份是Triton的源码本身GitHub上开源重点看python/triton/compiler.py和lib/Dialect目录下的文件。第三份是LLVM官方文档里关于MLIR的浅显介绍但只看概念部分就够了不用深入。我特别建议第二份资料要用好因为Triton源码的注释虽然不多但代码组织非常清晰命名也很见功力很多问题直接看代码比搜博客更准确。还有一个小建议把你日常使用的显卡信息记录下来。Triton针对不同GPU架构生成的PTX差异很大比如Ampere架构和Hopper架构在wgmma等新指令的使用上完全不同。你在学习时如果发现IR和教程不一致先别急着怀疑自己先确认显卡架构是不是不一样。这个坑我在后面专门讲。3. 从Triton内核到PTX指令一条龙看完编译器的“五脏六腑”3.1 写一个最简Kernel打开“上帝视角”纸上谈兵再多不如亲手打开一次“上帝视角”。我建议第3周的时候每个人都必须完成一次完整的IR导出实验。动手第一步准备一个最简单的Triton内核。我用的例子是经典的向量加法import torch import triton import triton.language as tl triton.jit def vector_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) z x y tl.store(Z_ptr offsets, z, maskmask)然后设置环境变量把三样东西导出来TTIR、TTGIR和PTX。export TRITON_KERNEL_DUMP1 export MLIR_DUMP_TO_STDOUT1在Python里调用一次x torch.randn(1024, devicecuda) y torch.randn(1024, devicecuda) z torch.zeros(1024, devicecuda) vector_add_kernel[(1,)](x, y, z, 1024, BLOCK_SIZE1024) torch.cuda.synchronize()Triton会在当前目录生成一个__triton_*目录里面按顺序存放了.ttir、.ttgir、.llir和.ptx文件。逐个打开它们你会亲眼看到源码是如何一层层“变形”的。3.2 对比TTIR与TTGIR看编译器怎样“做规划”打开.ttir文件你看到的是一个相当接近原始逻辑的中间表示经过整理后大致长这样#vector_add_kernel tt.func vector_add_kernel(...) { %0 tt.program_id x {axis 0 : i32} %1 tt.make_block_ptr ... %2 tt.load %1 {boundaryCheck arrayi32: 1, padding 0 : i32} %3 arith.addi %2, %2 : tensor1024xi32 tt.store %4, %3 {boundaryCheck arrayi32: 1} tt.return }这里面tt.program_id对应Python里的tl.program_idarith.addi对应Python里的非常直观看不出太多不可名状的魔法。这是Triton IR作为“领域专用IR”的优势——它离领域语义足够近任何有基本编程经验的人都能轻松读懂。但重点戏在TTGIR。TTGIR里会比TTIR多出非常多的信息量ttg.local_load、ttg.local_store、scf.for、分布到thread维度的ttg.convert_layout等指令开始出现。你第一次看到它时大概率是懵的这是正常的。我建议不要急着一行一行读而是直接观察“代码行数增加了多少倍”这个事实就够了。为什么一个简单的向量加法会膨胀出这么多代码因为TTGIR阶段是Triton编译器做“并行规划”的阶段——它要把“一块数据上的操作”Tile语义映射到“大量线程上的操作”SIMT语义。这个映射过程涉及数据划分、线程分配、布局转换、共享内存使用等众多决策而所有这些决策都是编译器自动完成的。你写kernel时从来没有手动指定过“线程0算哪几个元素线程1算哪几个元素”编译器帮你安排了TTGIR里记录的就是它的安排结果。3.3 深入一个核心Pass看编译器如何“分配布局”第5周的重点是布局分配Layout Assignment。这个概念值得每一个想理解编译器的人花时间去啃因为它是Triton能高效生成GPU代码的核心机制之一。用一个简单比喻来说Triton的Tile抽象就像是给了一张“100个元素的大桌子”而GPU上的线程是一群“只能伸手拿自己面前格子的人”。布局分配要解决的问题就是——这张桌子怎么摆才能让每个人拿东西的时候都顺手、不打架、传输量最少。Triton里有一个很重要的概念叫“分布式布局”Distributed Layout。它描述了Tile中的每个元素分别被映射到哪个线程的哪个寄存器上。比如一个tensor128xf32可以被布局为每个线程持有4个元素也可以被布局为每个线程持有1个元素前者所需线程数为32后者需要128个线程。布局的选择直接影响后续指令的执行效率和内存访问的合并程度。在Triton源码里和布局相关的Pass主要路径是triton/lib/Dialect/TritonGPU/Transforms/目录其中核心Pass之一叫RemoveLayoutConversions和OptimizeThreadLocality。你可以给编译器加上--print-after-all之类的MLIR参数看到每个Pass执行前后的IR快照。这一步建议选择一个最简kernel来做因为IR会爆炸式增长选大kernel只会把自己绕晕。我在带新人时要求他们必须完整推导一次“一个128元素向量加法在4个线程上执行”的布局转换全过程。手动推导完这一次后续读任何Pass代码都轻松得多。3.4 一路看到PTXGPU汇编并不神秘第7周是导出LLVM IR和PTX。这个过程的核心目的是建立“高级语言如何变成机器码”的具体认知。前面TTIR、TTGIR你还能看出点“代码的样子”到了LLVM IR结构语法会变得更琐碎getelementptr、load、store、call开始有一种“底层感”。再往下到PTX画风彻底变了ld.global.b32 %r1, [%rd2]; add.s32 %r3, %r1, %r4; st.global.b32 [%rd6], %r3;ld.global.b32是从全局加载32位数据.s32是有符号32位整数运算fma.rn.f32是带舍入的单精度乘加指令。如果你之前完全不懂汇编看到这里也不要慌几条常见指令查一次就记住了。更值得关注的是PTX里的指令顺序反映了编译器在指令调度层面的决策比如它把哪些计算放在了一起、插入了哪些bar.sync线程同步指令。学到这一步我已经觉得“编译器”这个词对你不再是黑盒了。你知道了它中间有几个过程、每个过程输入什么输出什么、关键决策在哪里做的。这和啃完龙书第三章后“看见正则表达式就想吐”的状态相比进步的是真实世界的工程认知。4. 实操过程与关键工具环境搭建、调试、问题排查4.1 环境准备半小时跑通第一个Triton Kernel环境是很多人迈不过去的第一道坎我先把我实测最顺畅的方案写出来。硬件是NVIDIA GPU驱动版本随便一个新的都行。Torch版本建议2.0以上Triton版本建议跟上PyTorch官方内置的版本大多数情况两者是绑定的。# 创建虚拟环境以conda为例 conda create -n triton-dev python3.10 -y conda activate triton-dev # 安装PyTorch根据自己的CUDA版本选命令这里以cu121为例 pip install torch torchvision --index-url https://download.pytorch.org/whl/cu121 # 安装Triton pip install triton安装完成后跑一个最简验证import triton print(triton.__version__)如果能正常输出版本号环境就算通了。如果版本号后面带githash也不影响使用。有个细节要提一下一部分人的安装过程会出现LLVM相关报错这通常是因为Triton的llvm依赖包版本与当前平台不匹配。这时候最简单的解法是换用官方预编译wheel包或者直接用pip install triton2.x.x指定一个较新且稳定的版本。不要自己去编译LLVM那是在给自己找事。4.2 导出IR的正确姿势三个环境变量一个不能少我用的是Triton 2.x版本导出IR最常用的方式是设置环境变量。网上很多教程只提了TRITON_KERNEL_DUMP但实战中我还会用另外两个。export TRITON_KERNEL_DUMP1 # 在cwd生成 __triton_* 目录包含全部机器码与IR export MLIR_DUMP_TO_STDOUT1 # 把MLIR中间结果打印到stdout方便即时对比 export TRITON_PRINT_AUTOTUNING1 # 如果用了autotune打印每种config的耗时TRITON_KERNEL_DUMP生成的目录里文件命名通常类似__triton_launcher/vector_add_kernel/0/vector_add_kernel.ttir __triton_launcher/vector_add_kernel/0/vector_add_kernel.ttgir __triton_launcher/vector_add_kernel/0/vector_add_kernel.llir __triton_launcher/vector_add_kernel/0/vector_add_kernel.ptx建议每次跑完把目录清空一次再重新跑不然新旧文件混在一起容易看岔。4.3 自己加一个Pass给编译器动一次“小手术”第10周的实战是个重头戏。我推荐做一个“不加功能、只做观察”的小Pass先熟悉改代码-重新编译-验证的完整链路再做实际优化。一个简单可行的实验是在triton/lib/Dialect/TritonGPU/Transforms/下新增一个打印当前IR的Pass。这个Pass的作用仅仅是遍历一遍IR把结果输出到日志里但它能帮你把整套工具链跑通// 在 TritonGPU/Transforms 目录下新增一个自定义 Pass 文件 struct PrintIRPass : public PassWrapperPrintIRPass, OperationPassModuleOp { MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(PrintIRPass) void runOnOperation() override { Operation* op getOperation(); op-dump(); } };然后在已有的Pass Pipeline注册入口里加一行注册代码。重新编译Triton后你的这个Pass就跑起来了。听上去平平无奇但这个过程中你会理解几件特别重要的事情Triton的编译管线是怎么组织Pass的、如何在MLIR框架里加自定义操作、重新编译的工作量和耗时。相当于一次“编译器开发者的岗前培训”。做完这些一个更进阶的玩法是做一个真正修改代码的Pass。我周围有同学尝试过加一个实验性的“循环展开因子调整”Pass最终成功让某个算子的性能提升了约15%。虽然这只是实验性修改但那种“我改了编译器的行为并且真的跑出了效果”的成就感远比读完一整本龙书来得踏实。4.4 常见问题与排查技巧实录这一节是我实际带人过程中高频踩坑的集合按频率排序直接给结论。问题一IR导出为空或目录不存在。多数情况是Triton的kernel在JIT缓存里直接命中了没有重新编译。解法很简单换一个BLOCK_SIZE比如1024改成2048或者清除Triton缓存目录~/.triton/cache。问题二PTX里看不到预期的wgmma或mma指令。大概率是GPU架构不支持Triton对老架构会自动降级。可以用torch.cuda.get_device_capability()查看算力如果小于8.0就不要期望看到Hopper相关的张量核心指令了。问题三Autotune结果和第一次手动调优时差异巨大。这个现象多半是受到了首次运行的CUDA context初始化、显存频率动态调整的影响。排查思路是先跑几轮“热身”多次重复同一实验取中位数并且固定GPU时钟频率我这边用的是nvidia-smi -lgc 1500来锁定频率。问题四自定义Pass改了源码后不生效。90%的情况是Triton的Python绑定时没有重编译。我自己的习惯是改C源码后执行一条龙cd triton python setup.py develop然后重启Python进程。如果还不生效就查一下环境变量TRITON_ALWAYS_COMPILE1确保每次都触发JIT重新编译。问题五想看的Pass根本不在默认Pass管线里。Triton 2.x默认的编译管线在triton/tools/compile.py和compiler.py里有一份完整Pass列表很多Pass默认没开需要加参数。比如布局冲突消除相关的一些Pass默认并不参与。所以如果要研究某个特定Pass最好先把Pass列表打印出来或者从python/triton/compiler.py里把当前编译管线复制出来手动增删Pass再用自定义版本的编译管线去编译。我把这些常见问题整理成了一张速查表方便大家查阅问题现象可能原因快速解法IR导出目录为空JIT缓存命中清空~/.triton/cache或修改kernel参数编译时LLVM相关报错Triton内置LLVM包与平台不匹配更新Triton到较新版本或换预编译wheelPTX指令与教程不一致GPU架构不同确认get_device_capability匹配对应架构自定义Pass不生效Python侧缓存了旧编译结果python setup.py develop后重启进程自动调优结果波动大CUDA上下文初始化、时钟频率多次运行取中位数锁定GPU时钟4.5 学习Triton时的三个“黄金法则”第一所有结论都必须形成闭环。很多人看完IR导出觉得“哦原来如此”关掉文件就完了。我的建议是每学一个概念都要“修改源程序→导出IR→观察差异→写出总结”四步走。比如想理解num_warps的影响就分别设num_warps1、4、8看TTGIR里的线程数如何变化。这一步做扎实知识就不是“听说”而是“验证过”的。第二把编译器当代码库来读而不是当文档来读。Triton的源码并不难读它比大多数工业级C项目要干净得多。如果你在读某个Pass时卡住了我的习惯是从triton/compiler.py里找到这个Pass被加入管线的位置然后看它的输入输出类型再看具体实现。自顶向下先骨架后肌肉比一头扎进一个几百行的pass文件里高效太多。第三和大模型配合着学但别全信。现在用ChatGPT或Claude来辅助读懂IR、解释某段PTX的语义非常高效。但有一点务必注意——让AI解释代码时一定要把真实的IR片段贴给它而不是让它凭记忆编造。AI容易一本正经地编出根本不是你这份IR的内容这种“幻觉”在编译器领域的蒙蔽性极强会硬生生让你绕几个小时弯路。5. 由Triton延伸出去懂了一个编译器其他编译器还远吗学完十周内容后再回头看那些之前觉得完全不在一个维度的概念比如GCC、Clang、MSVC、Keil的AC5/AC6编译器还有Python解释器之类的都会有一个全新的坐标。我从入门到玩转Triton这段时间最值钱的收获不是记住了多少IR语法而是建立了“所有编译器本质上都在做同一件事”的认知框架——把人类写的高级语言逐步降级为机器能执行的指令每一步都涉及分析和转换。以前看到一个报错“编译器错误信息: CS1056: 意外的字符”会觉得不知所云现在第一反应是“这应该是词法分析阶段报的错”。以前看Keil里编译器未包含main类型这种提示会觉得编译器像是个脾气古怪的黑箱现在会意识到这大概率是链接阶段没找到入口函数是工程配置问题而非语言问题。这种“把黑箱变白箱”的认知升级才是这十周真正留给你的资产。你不需要记住Triton的每一个Pass叫什么名字但你需要牢记住那条从源码走向机器指令的路径是什么样子以及在哪里可以打断它、修改它、优化它。从这个意义上说十周时间学到的不仅是Triton编译器本身更是现代编译器设计的“通用语法”。换到其他领域编译器无非是词汇表和方言不一样语法规则是相通的。以后再面对Keil的AC5/AC6切换、GCC的-O2优化开关、MSVC的编译报错时你都会有底气支棱起来。我个人在实际操作中的体会是十周“刨根问底”Triton这条路最大的价值不在于你最后对Triton本身有多了解而在于你建立起了“编译器可以被理解、可以被修改、可以被掌控”的根深蒂固的自信。有了这种自信以后不管是做算子优化还是做推理加速你都不会再觉得“底层的事情交给框架就好”。这种掌控感是啃一年龙书给不了的。最后再分享一个小技巧学完这套内容后如果你还想更上一层楼可以挑一个自己最常用的PyTorch算子里性能不够好的那个用Triton重写一遍然后用torch.compile的默认后端跑一遍比较一下两者的性能差异。如果Triton版本赢了试着通过改编译器Pass再追一截性能。这个真实世界的“擂台赛”比任何教程都更有吸引力也是从“会读编译器”到“能驾驭编译器”之间最短的桥。
