LLM 训练和推理性能优化走到今天已经绕不开 Triton。无论是 FlashAttention 一类的 fused attention还是 RMSNorm、RoPE、MoE 的 gate 分流几乎都能看到 Triton Kernel 的身影。但很多同学在调 Triton Kernel 时还是停留在“改 BLOCK_SIZE、试 num_warps、跑一下 benchmark”这个循环里。运气好能调快 10%运气不好反而更慢而且不知道为什么。这篇文章想给出的判断是Triton Kernel 优化的真正瓶颈往往不在 Python 层的参数而在编译器生成的中间表示和最终汇编里。与其凭经验调参不如建立一套“编译器视角的分层诊断”方法——从 TTIR、TTGIR、PTX 到 SASS一级一级往下看找到性能问题到底出在数据搬运、layout 转换、向量化宽度还是寄存器溢出。这套方法就是标题里说的 Compiler-Grounded Hierarchical Diagnosis编译器级分层诊断。文章会先讲清楚 Triton 的编译流水线然后给出一个可直接运行的 dump 示例最后用三个真实场景演示如何从 IR 层定位性能问题。读完这篇文章你应该能掌握一套可复用的 Triton Kernel 性能诊断流程而不是继续盲调参数。1. 为什么 LLM Kernel 优化不能只靠“经验调参”先说一个很常见的场景。你在优化一个 LLM 里的 softmax kernel发现它在 A100 上跑得比 PyTorch 原生实现还慢。于是你开始调 BLOCK_SIZE把 1024 改成 512发现快了一点点再试 num_warps8又慢回去了。继续调最后整个下午过去了性能提升不到 10%但代码却变得不可维护。这不是你一个人的问题。Triton 的好处是“把 CUDA 的复杂度藏起来”但坏处也是编译器自动做的 layout 推断、shared memory 分配、向量化决策对开发者来说像是一个黑盒。你在 Python 层看到的tl.load、tl.exp、tl.sum和 GPU 上真正执行的 SASS 指令之间隔了整整四层编译产物。盲调参数的核心问题在于BLOCK_SIZE、num_warps、num_stages 这些参数只是编译器优化策略的“开关”而不是性能问题的“原因”。比如一个 kernel 慢可能是因为 TTGIR 里出现了大量convert_layout导致数据在跨 warp 之间反复搬运也可能是因为 PTX 层的向量化宽度只有 32 bit内存带宽根本没有用满还可能是因为 SASS 层发生了严重的寄存器溢出大量的 local memory 访问把性能拖垮。这三种原因靠调 BLOCK_SIZE 是治不好的甚至可能越调越糟。这就是“Compiler-Grounded”的意义所在。它不是说不用 profiling 工具而是强调性能诊断的“事实依据”应该来自编译器的中间表示和最终汇编而不是来自猜测。你看到 TTGIR 里有一长串convert_layout就知道问题在 layout 策略看到 PTX 里是ld.global.b32而不是ld.global.b128就知道向量化没生效。每一层 IR 都在告诉你一件事你的代码让编译器做了什么。所以真正高效的优化流程应该是先跑通一个 baseline然后 dump 出编译中间产物按层级检查有没有异常特征最后再决定是改参数、改 kernel 结构还是改访存模式。这篇文章后面所有内容都是围绕这套流程展开的。2. Triton 编译器的本质从 Python DSL 到 GPU 机器码2.1 Triton 不是“写 kernel”而是“描述 kernel”很多初学者把 Triton 当成 CUDA 的 Python 替代品觉得它就是帮你在 Python 里写grid, block。这个理解不准确。Triton 本质上是一个面向 GPU 的 DSL 编译器你写的triton.jit函数并不是直接翻译成 CUDA C而是经过一条完整的编译流水线最终变成 GPU 可执行的机器码。这也解释了为什么 Triton Kernel 能跨 NVIDIA、AMD 甚至其他硬件平台运行因为它在编译期就把“逻辑计算”和“底层硬件”解耦了只要后端能处理 Triton IR就能生成对应平台的代码。2.2 Triton 的编译流水线五层中间产物一个 Triton Kernel 从 Python 代码到 GPU 机器码大致经过以下几步编译阶段产物主要作用前端解析Python AST把triton.jit的函数体解析成抽象语法树Triton IRTTIR与设备无关的中间表示描述核心计算逻辑TritonGPU IRTTGIR引入 block 布局、layout 转换、shared memory、barrier 等 GPU 概念LLVM IRLLIR经过 TritonGPU 优化后生成 LLVM 中间表示PTXPTXNVIDIA 虚拟指令集和具体 GPU 架构解耦SASS / CUBIN机器码真正在 GPU 上执行的硬件指令这里最关键的两层是 TTGIR 和 PTX / SASS。TTGIR 是定位“数据搬运问题”的核心。Triton 的编程单位是 block一个 block 里的数据如何分布到不同的线程、warp、子块上完全由 layout 决定。TTGIR 里会出现大量ttg.convert_layout、ttg.local_alloc、ttg.local_load这样的操作。如果你看到一段计算中间夹着很多 layout 转换意味着数据在跨线程搬运这会消耗 shared memory 带宽和同步开销。PTX 和 SASS 则是定位“底层执行效率问题”的核心。比如 load 指令的宽度是 32 bit 还是 128 bit访存是否 coalesced寄存器有没有保存到 local memory这些只能在 PTX / SASS 层看到。2.3 编译器自动做了哪些优化Triton 不是简单地把每个线程要做的计算展开而是会做一系列优化layout 推断决定一个 tensor 如何映射到线程块尽量避免跨线程搬运。自动向量化把多个 32 bit 访存合并成 128 bit 访存。shared memory 分配为tl.load的生命周期分配共享内存。循环优化对tl.range做 unrolling 和 pipelining。同步优化插入合理的 barrier 指令。这些优化给开发者带来了极大便利但也带来了新的问题当性能不达标时你很难确定是哪个优化环节没做好。Compiler-Grounded 的分层诊断就是要把这些问题暴露出来。3. 分层诊断模型五层定位法如果把 Triton Kernel 优化看成一个推理题那么每一层编译产物都是“证据”。我建议把诊断过程分成五个层级自顶向下逐层排查。3.1 L0源码与参数层这一层看的是你写的triton.jit函数本身BLOCK_SIZE、num_warps、num_stages、循环结构、mask 逻辑。它解决的问题是“我写的代码是否存在结构性低效”。比如在 softmax kernel 中如果 BLOCK_SIZE 小于一行的列数编译器会引入循环如果 BLOCK_SIZE 远大于实际数据量又会造成大量 mask 判断和空洞计算。这些在 L0 层就能发现。但 L0 层能发现的问题有限它看不到编译器真正做了什么。所以这一层只作为起点不作为终点。3.2 L1TTIR 层TTIR 是与设备无关的中间表示特点是还没有引入具体的 GPU block 布局。在这一层你可以看到是否存在冗余的tt.splat或tt.broadcast。是否有多余的 mask 计算。计算的“形状”是否符合预期。TTIR 在性能诊断中通常不是重点但它是理解后续 TTGIR 的前提。3.3 L2TTGIR 层这是整个诊断模型中最重要的层级。TTGIR 引入了#ttg.blocked、#ttg.slice、#ttg.shared等布局描述并且会明确写出ttg.convert_layout、ttg.local_alloc、ttg.local_load等操作。在 TTGIR 里你能看到tensor 的 layout 是否合理。是否存在大量 layout 转换导致跨 warp 数据搬运。shared memory 是否被过度使用。每个 block 内的处理流程是什么。如果 TTGIR 中convert_layout出现的次数明显偏多这就是一个强信号数据搬运开销过大。3.4 L3PTX / LLIR 层PTX 是 NVIDIA 的虚拟指令集。在这一层你可以看到ld.global.b32还是ld.global.b128这决定了访存宽度。ld.shared和st.shared的分布shared memory 访问频率。bar.sync的数量同步开销。fma.rn.f32、tex.exp.approx.f32等指令类型计算指令效率。PTX 层最能反映访存和计算的比例。3.5 L4SASS 层SASS 是最终在 GPU 上执行的机器码。通过 SASS你能看到寄存器分配数量。是否存在STL/LDLlocal memory 的 store / load即寄存器溢出。等待指令、空转指令的比例。最终的内存访问宽度。SASS 是“最终事实”但可读性不如 PTX。实际工程中我建议先看 TTGIR 和 PTX遇到寄存器溢出再深入 SASS。3.6 L5Profiler 层分层的 IR 诊断告诉你“为什么慢”而 NVIDIA Nsight Computencu这类 profiling 工具告诉你“慢在哪里、成本多高”。两者是交叉验证的关系。比如你在 TTGIR 里看到大量 convert_layout接下来可以用 ncu 看 shared memory 的 bank conflict 是不是特别高或者 barrier 等待时间是不是很长。层级诊断对象能发现的问题典型工具L0源码与参数结构性低效、mask 过多人工 reviewL1TTIR冗余计算、形状不匹配TF32 视图、IR dumpL2TTGIRlayout 转换、shared memory 开销TRITON_KERNEL_DUMPL3PTX / LLIR向量化宽度、同步频率反汇编、文本查看L4SASS寄存器溢出、最终机器指令cuobjdump、SASS dumpL5Profiler硬件利用率、stall 原因ncu、NVIDIA Nsight Systems这个五层定位法不是严格的顺序流程而是按照“从源码到硬件”的路径帮助你快速定位问题所在。4. 环境准备把 Triton Kernel 的“编译现场”打开要做编译器级诊断第一步是能够拿到 Triton 的编译中间产物。这里不需要自己写编译器插件Triton 本身就提供了 dump 机制。4.1 安装 Triton最直接的方式是 pip 安装pip install triton如果你用的是 vLLM、SGLang 等 LLM 推理框架它们会自带某个版本的 Triton通常位于框架的vllm/third_party或site-packages里。这时候要注意不要随意覆盖框架自带的 Triton不然很可能出现版本冲突。一个很常见的报错是module triton has no attribute language这个问题的本质是 Triton 版本异常比如你同时装了多个 Triton或者框架内置的 Triton 与系统层 Python 包发生了冲突。遇到这种问题先检查python -c import triton; print(triton.__version__)并确认当前解释器加载的 Triton 路径来自哪里。4.2 常用 dump 环境变量Triton 提供了多个环境变量来控制编译过程的可观测性下面是最常用的几个# 把编译中间产物 dump 到 /tmp/triton 目录通常包含 .ttir / .ttgir / .llir / .ptx / .cubin export TRITON_KERNEL_DUMP1 # 打印 autotune 过程便于观察不同配置的耗时 export TRITON_PRINT_AUTOTUNING1 # 强制每次都重新编译避免使用缓存中的旧编译结果 export TRITON_ALWAYS_COMPILE1TRITON_KERNEL_DUMP1是最常用的诊断开关。设置后每次运行 Triton Kernel编译器会把中间产物写到临时目录通常是/tmp/triton文件名会带上 kernel 名和 compile key。4.3 在代码中直接获取编译产物除了环境变量你也可以在 Python 代码里直接拿到编译后的 kernel 对象。Triton 的 JITFunction 在调用后会返回一个 CompiledKernel 对象它有一个asm属性里面包含了多层中间产物kernel my_kernel[grid](x, y, out, n, BLOCK_SIZE1024) print(kernel.asm[ttgir]) # 查看 TritonGPU IR print(kernel.asm[ptx]) # 查看 PTX需要说明的是不同版本的 Triton 在 API 细节上略有差异。如果kernel.asm在你的版本里不可用优先使用TRITON_KERNEL_DUMP1这是最通用、最稳定的方式。4.4 验证 dump 是否生效在正式诊断前可以先跑一个最小示例确认 dump 文件能正常生成。如果运行完没看到 dump 文件按如下顺序排查确认环境变量在当前进程里生效echo $TRITON_KERNEL_DUMP。确认 kernel 确实执行过且发生了编译而不是从缓存里加载的旧结果。检查/tmp/triton目录是否存在且有写权限。5. 完整示例编译并 dump 一个 Triton Kernel 的分层 IR下面用一个完整的 Python 脚本来演示如何在一个脚本里同时触发两个典型 kernel 的编译和 dump。5.1 示例脚本dump_triton_ir.py# dump_triton_ir.py import torch import triton import triton.language as tl triton.jit def add_kernel(x_ptr, y_ptr, out_ptr, n_elements, BLOCK_SIZE: tl.constexpr): pid tl.program_id(axis0) offsets pid * BLOCK_SIZE 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(out_ptr offsets, x y, maskmask) triton.jit def softmax_kernel(x_ptr, out_ptr, n_cols, BLOCK_SIZE: tl.constexpr): row tl.program_id(axis0) col_offsets tl.arange(0, BLOCK_SIZE) row_start row * n_cols x tl.load(x_ptr row_start col_offsets, maskcol_offsets n_cols) x x - tl.max(x, axis0) num tl.exp(x) den tl.sum(num, axis0) tl.store(out_ptr row_start col_offsets, num / den, maskcol_offsets n_cols) def main(): # elementwise add一个简单的访存密集型 kernel n 1024 x torch.randn(n, devicecuda) y torch.randn(n, devicecuda) out torch.empty(n, devicecuda) grid_add (triton.cdiv(n, 1024),) add_kernel[grid_add](x, y, out, n, BLOCK_SIZE1024) torch.cuda.synchronize() # block softmax一个包含 reduction 的 kernel更容易暴露 layout 问题 rows 8 cols 4096 x2 torch.randn(rows, cols, devicecuda) out2 torch.empty(rows, cols, devicecuda) grid_softmax (rows,) softmax_kernel[grid_softmax](x2, out2, cols, BLOCK_SIZEtriton.next_power_of_2(cols)) torch.cuda.synchronize() print(done) if __name__ __main__: main()这个脚本里包含两个 kerneladd_kernel简单的 elementwise 加法适合观察基础向量化和访存。softmax_kernel按行做 softmax包含跨线程的 reduction适合观察 layout 转换和 shared memory 使用。5.2 运行并触发 dumpTRITON_ALWAYS_COMPILE1 TRITON_KERNEL_DUMP1 python dump_triton_ir.py运行结束后查看/tmp/triton目录ls -lh /tmp/triton你会看到类似这样的文件add_kernel_...ttir add_kernel_...ttgir add_kernel_...llir add_kernel_...ptx softmax_kernel_...ttgir softmax_kernel_...ptx5.3 TTGIR 里能看到什么以softmax_kernel的 TTGIR 为例关键信息会包括每个 tensor 标注的 layout比如tensor4096xf32, #blocked或tensor4096xf32, #slice。ttg.local_alloc/ttg.local_load操作。是否有ttg.convert_layout操作。这里贴一段示意片段实际文件内容会因为你使用的 Triton 版本、GPU 型号而不同#blocked #ttg.blocked{sizePerThread [1], threadsPerWarp [32], warpsPerCTA [4], order [0]} module { func.func public softmax_kernel(...) { ... %x tt.load %arg0, %offsets : tensor4096xf32, #blocked %max tt.reduce.max %x : tensor4096xf32, #blocked %x_shifted arith.subi %x, %max %num tt.exp %x_shifted %den tt.reduce.sum %num : tensor4096xf32, #blocked %div arith.divf %num, %den tt.store %arg1, %offsets, %div : tensor4096xf32, #blocked } }在这段示意里注意 reduction 操作是否会触发额外的 layout 转换。真正的 TTGIR 会明确写出这些转换通常表现为ttg.convert_layout的多次调用。5.4 PTX 里能看到什么PTX 文件会包含具体的访存指令比如ld.global.b32 %r0, [%rd0]; ld.global.b32 %r1, [%rd04]; ...如果看到多条连续地址的b32指令说明编译器只做了 32 bit 访存。如果你在源码里确认数据是对齐的那么这通常意味着BLOCK_SIZE或 layout 设计导致向量化没有生效。更好的形式是ld.global.b128 %r0, [%rd0];一条 128 bit 指令完成 16 字节的读取效率会高很多。5.5 在代码中直接查看 asm如果你不想依赖环境变量也可以在 Python 里直接访问编译产物kernel softmax_kernel[(rows,)](x2, out2, cols, BLOCK_SIZEtriton.next_power_of_2(cols)) print(kernel.asm[ttgir]) print(kernel.asm[ptx])这种方式适合在 Jupyter Notebook 里快速验证。6. 诊断实战从 IR 层捕捉三个典型性能问题拿到 IR 不是目的能从 IR 里看出问题才是目的。下面用三个典型案例演示如何从 IR 中发现性能瓶颈。6.1 问题一TTGIR 中的 convert_layout 风暴诊断现象在softmax_kernel.ttgir中ttg.convert_layout出现次数明显偏多并且在它前后紧跟着ttg.local_alloc/ttg.local_load和bar.sync。问题本质Triton 的 block 计算模型里一个 tensor 的数据会按照某种 layout 分布在各个线程上。当两个操作的 layout 不一致时就需要通过 shared memory 做一次数据重排。假设你要对一行 4096 个元素做 softmax如果这行数据初始 layout 是sizePerThread [1]即每个线程负责一个元素那么在计算跨行 reduction 时数据需要被重排成适合 reduction 的布局。重排过程会产生convert_layout并且需要经过 shared memory伴随 barrier 同步。为什么性能差convert_layout 本质是数据搬运搬运次数多了kernel 的大部分时间都花在等数据上而不是花在计算指数和除法上。解决策略调整BLOCK_SIZE尽量让一个 block 覆盖的 shape 与目标 GPU 的 vectorized load 对齐减少 layout 切换的频率。如果 kernel 本身需要大量 reduction考虑把它拆成两阶段 kernel先做行内规约再做跨行规约避免让 Triton 在同一个 layout 里反复寻找最优解。适当提升num_warps让更多线程参与并发搬运降低单线程的序列化等待。6.2 问题二PTX 层向量化宽度不足诊断现象在 PTX 中看到大量ld.global.b32却很少看到ld.global.b128。同地址连续递增的 load 指令占多数。问题本质GPU 访存的理想形式是每个线程一次性加载 128 bit 数据。如果 Triton 编译器认为数据对齐信息不足或者 layout 让连续地址分配到不同线程向量化就会被破坏退化成多个 32 bit 访存。为什么性能差同样的数据量b32 指令数量是 b128 的 4 倍。这意味着需要更多的内存事务也占用更多的指令槽位和寄存器。解决策略确保输入 tensor 在内存中是连续且对齐的尽量使用torch.empty/torch.zeros这类对齐分配。尽量让BLOCK_SIZE是 16 字节的整数倍并且让每个线程处理的数据量足够大以支持 128 bit 访存。在 mask 判断上避免使用过于复杂的offsets计算因为复杂 index 计算会影响编译器的向量化推断。6.3 问题三SASS 层寄存器溢出诊断现象在 SASS 中观察到明显的STL和LDL指令或者 PTX 中出现st.local/ld.local。问题本质当 kernel 的寄存器需求超过硬件所能提供的数量时编译器会把部分频繁使用的数据保存到 local memory本质上是显存。这被称为 register spill。为什么性能差local memory 虽然名义上是“内存”但访问延迟远高于寄存器。一旦代码发生寄存器溢出性能会断崖式下跌。在 LLM kernel 里某些中间 tensor 过多、或者 BLOCK_SIZE 过大时很容易触发这个问题。解决策略减小BLOCK_SIZE减少每个线程需要同时存活的中间数据量。调整num_warps因为 register file 是有限的warp 数量会影响每个线程可用的寄存器配额。简化 kernel 中同时存活的中间变量。比如在 softmax 中不要过早保存x_shifted、num、den等多个大 tensor而是复用变量名减少编译器的 liveness 压力。6.4 三个问题放在一起怎么看问题类型观察层级特征典型动作convert_layout 风暴TTGIRconvert_layoutbar.sync频繁调整 block shape拆分 kernel向量化宽度不足PTXld.global.b32过多对齐数据调整 BLOCK_SIZE寄存器溢出SASS / PTXSTL/LDL/st.local减少中间变量调整 num_warps7. 用 Profiler 交叉验证让诊断结论落地IR 诊断能帮你定位“问题是什么”但要了解“这个问题造成了多大代价”最好用 NVIDIA Nsight Computencu做交叉验证。7.1 常用 ncu 命令# 对 softmax kernel 做基础的 memory 和 warp stall 分析 ncu --set full --section MemoryWorkloadAnalysis --section WarpStateStats python dump_triton_ir.py# 只看寄存器溢出等关键指标 ncu --metrics launch__registers_per_thread,sm__sass_inst_executed_op_shared_ldst_pred_on.sum \ python dump_triton_ir.py7.2 如何对应 IR 诊断如果 IR 里发现大量 convert_layoutncu 里通常能看到 shared memory 相关指标较高或者barrier相关的 stall 占比高。如果 PTX 里向量化宽度不足ncu 的 Memory Workload Analysis 会显示内存事务数偏多但实际吞吐利用率不高。如果发生寄存器溢出ncu 里可以看到 local memory 相关指标从 0 变为非 0。需要注意ncu 的每一项指标含义都比较深不建议一开始就无脑--set full跑全部指标那样输出会很庞大反而干扰判断。建议先跑--section或指定几个关键 metric围绕你从 IR 里看到的怀疑点去验证。7.3 一个推荐的验证闭环先跑TRITON_KERNEL_DUMP1观察 TTGIR 和 PTX。根据 IR 特征提出假设比如“这里 convert_layout 太多”。针对假设跑 ncu 的对应 section看硬件指标是否印证。修改 kernel 或参数重新 dump IR 和跑 benchmark。对比修改前后的 IR 差异和性能差异确认优化是否真的有效。这套闭环的好处是每一步都有据可查而不是“改个参数试试”。8. 常见问题与排查清单问题现象可能原因排查方式解决方案module triton has no attribute language多个 Triton 版本冲突或框架内置 Triton 被污染python -c import triton; print(triton.__file__, triton.__version__)确认加载路径在隔离环境安装 Triton不要手动覆盖 LLM 框架自带版本TRITON_KERNEL_DUMP1没有输出进程没读到环境变量kernel 走了缓存echo $TRITON_KERNEL_DUMP添加TRITON_ALWAYS_COMPILE1强制重编译设置环境变量后重新运行检查 dump 目录写权限PTX 里大量b32没有b128数据未对齐BLOCK_SIZE 不适合向量化mask 逻辑太复杂打开 PTX 确认连续地址对应的指令检查 tensor 对齐使用连续内存调整 BLOCK_SIZE简化 offsets 计算TTGIR 里convert_layout很多不同操作的 layout 不匹配reduction 频繁触发跨线程重排逐段查看 convert_layout 前后操作调整 block 结构两阶段 kernel调整 num_warpsSASS 里出现STL/LDL寄存器溢出ncu --metrics launch__registers_per_thread ...降低 BLOCK_SIZE减少中间变量调整 num_warps修改参数后性能没变化编译缓存未失效添加TRITON_ALWAYS_COMPILE1或清除/tmp/triton下对应文件强制重编译后重新跑 benchmarkdump 文件太多不知道看哪个缺少分析路径按 L0→L5 分层检查先看 TTGIR再看 PTX/SASS建立团队统一的 IR 检查清单按图索骥9. 最佳实践与工程建议9.1 建立可复现的 baseline做任何优化之前先把你当前 kernel 的编译产物和性能数据保存下来。这个 baseline 不只是“跑一次 benchmark”而是包括Triton 和 CUDA 版本。GPU 型号和 driver 版本。dump 出的 TTGIR / PTX。ncu 的关键指标。这样后续每次改动都能对照 baseline 判断效果。9.2 每次只改一个变量Triton Kernel 优化的参数空间很大BLOCK_SIZE、num_warps、num_stages、kernel 结构、数据布局、编译器选项。一次改多个变量你根本无法判断是哪个改动起了作用。推荐的做法是先只改 BLOCK_SIZE记录 IR 和 benchmark再只改 num_warps再改 kernel 结构。每次只改一个改动前后都 dump IR对比差异。9.3 把 IR 检查清单融入团队协作可以把前面提到的三层检查点整理成一份 checklist放在项目文档里让团队每次提交 kernel 优化时都按这个流程走一遍TTGIR 中是否出现大量convert_layoutPTX 中访存指令是否以b128为主SASS / PTX 中是否有STL/LDLncu 的 stall 原因是否集中在 shared memory 或 barrier有了这份 checklistkernal 优化的 review 就不再是“看几个 benchmark 数字”而是有中间产物可查。9.4 谨慎对待生产环境替换LLM 推理框架里的 kernel 往往和 attention、KV cache、quantization 绑定。替换一个 Triton Kernel 前一定要在测试环境做正确性验证比如对比 PyTorch 原生实现的输出误差并确认误差在可接受范围内。同时建议保留旧的 kernel 实现方便灰度回滚。尤其是涉及线上服务时不要在一个发布周期里同时替换多个 kernel否则出了问题很难定位。9.5 关注版本兼容Triton 的编译产物和 API 在不同版本之间变化较大。网上很多教程里的 IR 片段可能基于某个老版本。如果你使用的版本较新dump 出来的 IR 会有所差异。遇到对不上的情况优先查看你当前版本的官方文档和源码而不是硬套旧经验。10. 总结这篇文章的核心判断很简单LLM 场景下的 Triton Kernel 优化不应该停留在“改参数 看耗时”这个层面而应该把编译器的中间产物当成诊断对象。从 TTIR 到 TTGIR再到 PTX 和 SASS每一层都记录了编译器做了什么决策也暴露了性能问题的蛛丝马迹。你可以从两个最小实践开始跑通TRITON_KERNEL_DUMP1把一个真实 kernel 的 TTGIR 和 PTX dump 出来先熟悉每一层 IR 长什么样。打开 TTGIR数一下convert_layout出现的次数打开 PTX数一下ld.global.b32和ld.global.b128的比例。这两个动作就能帮你避免大量“无效调参”。以后再遇到 Triton Kernel 性能问题记住一句话不要猜编译器做了什么直接看它做了什么。优化之前先让编译器开口说话。