下面把Triton 编程范式作为一个研究方向来梳理它不只是“用 Python 写 GPU kernel”而是把 GPU 编程抽象从thread-level 提升到tile / block-program level 的一整套模型、编译器和工程方法论。1. 一句话定位CUDAScalar Program, Blocked Threads你写一条线程硬件复制成成千上万线程TritonBlocked Program, Scalar Threads你写一个 program 如何处理一整块数据编译器再把块拆给线程/warp/Tensor CoreTriton Python DSL MLIR 编译器 tile 抽象 autotune 工作流目标是用接近 PyTorch 的生产力写出接近手调 CUDA 的深度学习/GPU 算子。2. 核心编程范式Tile-Centric SPMD2.1 基本心智模型CUDA 里你写int i blockIdx.x * blockDim.x threadIdx.x; y[i] a[i] b[i];Triton 里你写triton.jit def add_kernel(x_ptr, y_ptr, out_ptr, n, BLOCK_SIZE: tl.constexpr): pid tl.program_id(0) offs pid * BLOCK_SIZE tl.arange(0, BLOCK_SIZE) mask offs n x tl.load(x_ptr offs, maskmask) y tl.load(y_ptr offs, maskmask) tl.store(out_ptr offs, x y, maskmask)关键点没有threadIdx/warpIdx没有__shared__声明没有__syncthreads()操作对象是一整块 tensor / tile边界用mask表达而不是 if-break2.2 三个核心概念概念含义Program一个并行执行实例大致对应 CUDA 的一个 thread block / CTATileprogram 一次处理的一块数据vector / sub-matrixconstexpr meta-paramBLOCK_M/N/K、num_warps、num_stages等编译期超参程序员决定数据怎么分块哪个 program 处理哪个 tile算法是否融合softmax residual、GEMM epilogue编译器决定tile → warp → thread 的映射访存合并shared memory 分配与 pipelineTensor Core 指令选择mma.sync/wgmma/mfma软件流水线、register/spill 管理3. 编译器范式MLIR 多层 IRTriton 不是“Python 转 CUDA 字符串”而是正经编译器Python triton.jit ↓ trace / AST Triton IR (TTIR) # tile-level, hardware agnostic ↓ Triton GPU IR (TTGIR) # 加入 layout / warp / shared / MMA ↓ LLVM IR ↓ PTX / cubin (NVIDIA) AMDGCN / hsaco (AMD)TTGIR 是灵魂#blocked、#dot_op、#shared、#nvidia_mma、#amd_mfmalayout propagationdistribute-to-warpsasync copy / software pipeliningMMA intrinsic matching这也是为什么同一份 Triton 源码可以跨架构生成不同机器码。4. Autotune 是范式的一部分不是附属品在 CUDA 里改 launch config → 改 template → 重编译 → 手测在 Triton 里triton.autotune( configs[ triton.Config({BLOCK_M:128,BLOCK_N:128,BLOCK_K:32}, num_warps4), triton.Config({BLOCK_M:64,BLOCK_N:256,BLOCK_K:64}, num_warps8), ], key[M,N,K] )Autotune 把“架构知识”外置成配置空间tile shapewarp 数pipeline stage是否用 dot / atomic / block ptr这使得 Triton 工程范式为写对算法 → 给配置空间 → 首次运行 benchmark → 缓存最优 kernel5. Triton 最适合什么问题✅ 强项融合算子LayerNorm Residual、Softmax、SwiGLU、Rotary内存瓶颈算子elementwise / reduction / softmax自定义 GEMM epilogueFlashAttention / fused attention非标准形状、非标准 dtypefp8 / bf16 / nvfp4torch.compile / Inductor 后端生成代码❌ 弱项强不规则访存hash table、spMV、graph scatterall-to-all particle interactionN-body需要精细 warp 级原语、custom TMA、warp specialization极端性能天花板场景专家级 CUTLASS / CUDA 仍更强有论文实测规则 compute-bound GEMMTriton 可接近甚至超过 cuBLASSpMV / N-bodyTriton 可能明显落后于手调 CUDA6. 与 CUDA / CUTLASS / TVM 的范式对比维度CUDA CCUTLASSTritonTVM抽象层级threadtileC templatetilePythontensor program共享内存手写template/layout编译器管schedule 管Tensor Core手调 PTXCuTe layouttl.dot自动映射intrinsic生产力低中低高中性能天花板最高很高高90–98%高但调参重适合人群底层架构师算子库作者算法/训练/推理工程师编译系统研究者一句话CUTLASS 是“用 C 把 Triton 想做的事手动做一遍”Triton 是“把 CUTLASS 的智力塞进编译器”。7. 研究层面的开放问题如果你要做“Triton 编程范式研究”值得盯的方向Tile abstraction 的理论边界什么算子天然不适合 tile-SPMDirregular / dynamic-shape / control-flow-heavy 如何建模Layout-aware IR 设计TritonGPU IR 的 layout propagation 是否足够Gluon / ML-Triton 的 workgroup-warp-intrinsic 多级 IR 是否是未来Autotune 自动化用代价模型替代暴力 bench结合 roofline / occupancy / register pressure 预测最优 config跨架构公平性问题Hopper / Blackwell / AMD / Intel XPU 上同一源码性能方差很大Triton 的“可移植性”到底是源码可移植还是性能不可移植与 LLM 代码生成结合LLM 生成 Triton kernelTriton 作为“GPU 算子中间表示”比 CUDA 更适合模型生成和 PyTorch 编译栈的关系Inductor → Triton未来会不会出现“Triton 作为 AI 系统通用后端”8. 典型研究题目可直接开做《Triton tile abstraction 对不规则访存算子的表达力边界》《TritonGPU IR 中 layout propagation 对 GEMM pipeline 的影响》《基于代价模型的 Triton autotune 替代方案》《Triton vs CUTLASS vs CuTetile 编程范式系统性评测》《LLM Triton自动生成 fused attention kernel 的可靠性研究》《Hopper/Blackwell 上 Triton 的 TMA / wgmma 利用率分析》9. 最小研究结论Triton 的范式创新不是“Python 写 GPU”而是把线程调度、访存合并、shared memory、Tensor Core 映射这些专家知识从“人脑”移进“编译器 autotune 配置空间”。它牺牲一点极限性能换来了更低认知负担更快算子迭代更接近算法本身的表达成为 PyTorch / LLM 基础设施的“GPU 汇编层” 
长河君 · 2026/10/14 15:04:14 · 阅读 1.2k 学习: Triton
下面把Triton 编程范式作为一个研究方向来梳理它不只是“用 Python 写 GPU kernel”而是把 GPU 编程抽象从thread-level 提升到tile / block-program level 的一整套模型、编译器和工程方法论。1. 一句话定位CUDAScalar Program, Blocked Threads你写一条线程硬件复制成成千上万线程TritonBlocked Program, Scalar Threads你写一个 program 如何处理一整块数据编译器再把块拆给线程/warp/Tensor CoreTriton Python DSL MLIR 编译器 tile 抽象 autotune 工作流目标是用接近 PyTorch 的生产力写出接近手调 CUDA 的深度学习/GPU 算子。2. 核心编程范式Tile-Centric SPMD2.1 基本心智模型CUDA 里你写int i blockIdx.x * blockDim.x threadIdx.x; y[i] a[i] b[i];Triton 里你写triton.jit def add_kernel(x_ptr, y_ptr, out_ptr, n, BLOCK_SIZE: tl.constexpr): pid tl.program_id(0) offs pid * BLOCK_SIZE tl.arange(0, BLOCK_SIZE) mask offs n x tl.load(x_ptr offs, maskmask) y tl.load(y_ptr offs, maskmask) tl.store(out_ptr offs, x y, maskmask)关键点没有threadIdx/warpIdx没有__shared__声明没有__syncthreads()操作对象是一整块 tensor / tile边界用mask表达而不是 if-break2.2 三个核心概念概念含义Program一个并行执行实例大致对应 CUDA 的一个 thread block / CTATileprogram 一次处理的一块数据vector / sub-matrixconstexpr meta-paramBLOCK_M/N/K、num_warps、num_stages等编译期超参程序员决定数据怎么分块哪个 program 处理哪个 tile算法是否融合softmax residual、GEMM epilogue编译器决定tile → warp → thread 的映射访存合并shared memory 分配与 pipelineTensor Core 指令选择mma.sync/wgmma/mfma软件流水线、register/spill 管理3. 编译器范式MLIR 多层 IRTriton 不是“Python 转 CUDA 字符串”而是正经编译器Python triton.jit ↓ trace / AST Triton IR (TTIR) # tile-level, hardware agnostic ↓ Triton GPU IR (TTGIR) # 加入 layout / warp / shared / MMA ↓ LLVM IR ↓ PTX / cubin (NVIDIA) AMDGCN / hsaco (AMD)TTGIR 是灵魂#blocked、#dot_op、#shared、#nvidia_mma、#amd_mfmalayout propagationdistribute-to-warpsasync copy / software pipeliningMMA intrinsic matching这也是为什么同一份 Triton 源码可以跨架构生成不同机器码。4. Autotune 是范式的一部分不是附属品在 CUDA 里改 launch config → 改 template → 重编译 → 手测在 Triton 里triton.autotune( configs[ triton.Config({BLOCK_M:128,BLOCK_N:128,BLOCK_K:32}, num_warps4), triton.Config({BLOCK_M:64,BLOCK_N:256,BLOCK_K:64}, num_warps8), ], key[M,N,K] )Autotune 把“架构知识”外置成配置空间tile shapewarp 数pipeline stage是否用 dot / atomic / block ptr这使得 Triton 工程范式为写对算法 → 给配置空间 → 首次运行 benchmark → 缓存最优 kernel5. Triton 最适合什么问题✅ 强项融合算子LayerNorm Residual、Softmax、SwiGLU、Rotary内存瓶颈算子elementwise / reduction / softmax自定义 GEMM epilogueFlashAttention / fused attention非标准形状、非标准 dtypefp8 / bf16 / nvfp4torch.compile / Inductor 后端生成代码❌ 弱项强不规则访存hash table、spMV、graph scatterall-to-all particle interactionN-body需要精细 warp 级原语、custom TMA、warp specialization极端性能天花板场景专家级 CUTLASS / CUDA 仍更强有论文实测规则 compute-bound GEMMTriton 可接近甚至超过 cuBLASSpMV / N-bodyTriton 可能明显落后于手调 CUDA6. 与 CUDA / CUTLASS / TVM 的范式对比维度CUDA CCUTLASSTritonTVM抽象层级threadtileC templatetilePythontensor program共享内存手写template/layout编译器管schedule 管Tensor Core手调 PTXCuTe layouttl.dot自动映射intrinsic生产力低中低高中性能天花板最高很高高90–98%高但调参重适合人群底层架构师算子库作者算法/训练/推理工程师编译系统研究者一句话CUTLASS 是“用 C 把 Triton 想做的事手动做一遍”Triton 是“把 CUTLASS 的智力塞进编译器”。7. 研究层面的开放问题如果你要做“Triton 编程范式研究”值得盯的方向Tile abstraction 的理论边界什么算子天然不适合 tile-SPMDirregular / dynamic-shape / control-flow-heavy 如何建模Layout-aware IR 设计TritonGPU IR 的 layout propagation 是否足够Gluon / ML-Triton 的 workgroup-warp-intrinsic 多级 IR 是否是未来Autotune 自动化用代价模型替代暴力 bench结合 roofline / occupancy / register pressure 预测最优 config跨架构公平性问题Hopper / Blackwell / AMD / Intel XPU 上同一源码性能方差很大Triton 的“可移植性”到底是源码可移植还是性能不可移植与 LLM 代码生成结合LLM 生成 Triton kernelTriton 作为“GPU 算子中间表示”比 CUDA 更适合模型生成和 PyTorch 编译栈的关系Inductor → Triton未来会不会出现“Triton 作为 AI 系统通用后端”8. 典型研究题目可直接开做《Triton tile abstraction 对不规则访存算子的表达力边界》《TritonGPU IR 中 layout propagation 对 GEMM pipeline 的影响》《基于代价模型的 Triton autotune 替代方案》《Triton vs CUTLASS vs CuTetile 编程范式系统性评测》《LLM Triton自动生成 fused attention kernel 的可靠性研究》《Hopper/Blackwell 上 Triton 的 TMA / wgmma 利用率分析》9. 最小研究结论Triton 的范式创新不是“Python 写 GPU”而是把线程调度、访存合并、shared memory、Tensor Core 映射这些专家知识从“人脑”移进“编译器 autotune 配置空间”。它牺牲一点极限性能换来了更低认知负担更快算子迭代更接近算法本身的表达成为 PyTorch / LLM 基础设施的“GPU 汇编层”