Triton 2023 开发者大会全回顾编译器路线图、多后端生态与关键议题解析【免费下载链接】tritonDevelopment repository for the Triton language and compiler项目地址: https://gitcode.com/GitHub_Trending/tri/triton2023 年 9 月 20 日首届 Triton Developer Conference 在美国加州山景城微软硅谷园区以混合模式举行来自 OpenAI、NVIDIA、AMD、Intel、Qualcomm、Meta、Google、Microsoft 的十余位演讲者齐聚一堂围绕 Triton 语言与编译器的过去、现在与未来展开了全景式分享。本文以 docs/meetups/dev-meetup-2023.md 记录的官方议程为骨架逐一解读每个议题的技术内涵并结合当前 Triton 仓库GitHub_Trending/tri/triton中的源码、教程与文档为读者呈现一场会议议程 仓库现状的双重视角回顾帮助理解 Triton 生态的多后端架构、设备无关设计目标与高性能内核编程模型。会议概况一场汇聚硬件厂商与编译器团队的行业峰会根据 docs/meetups/dev-meetup-2023.md本次大会于2023 年 9 月 20 日在加州山景城的 Microsoft Silicon Valley Campus举行采用线上 线下混合模式。正式议程从上午10:00持续到下午4:00随后是持续到5:30 PM的社交招待环节Reception。从参会阵容可以清晰看出 Triton 生态的横向扩张特征除了 Triton 语言的缔造者 OpenAI还包括三大 GPU 厂商NVIDIA、AMD、Intel、移动端 DSP 厂商 Qualcomm、以及正在自研 AI 芯片和编译栈的 MetaMTIA与 GoogleXLA / JAX。这场会议事实上勾勒出了 Triton 从面向 NVIDIA GPU 的 CUDA 替代品向硬件无关的深度学习编译器中间层演进的早期蓝图。仓库现状也印证了这一多后端格局当前 third_party 目录下同时维护着 amd、nvidia 与 proton 三个第三方后端README.md 中列出的受支持硬件包括 NVIDIA GPUsCompute Capability 8.0与 AMD GPUsROCm 6.2CPU 后端处于开发中。大会完整议程一览以下是大会官方议程来源docs/meetups/dev-meetup-2023.md时间、标题与演讲者均为原文内容时间议题演讲者10:00 AMWelcomeKevin ScottMicrosoft10:20 AMThe Triton Compiler: Past, Present and FuturePhil TilletOpenAI11:00 AMBreak11:20 AMHopper support in TritonGustav ZhuNvidia11:40 AMBringing Triton to AMD GPUsJason Furmanek, Lixun ZhangAMD12:00 PMIntel XPU Backend for TritonEikan WangIntel12:20 PMVectorization of Triton Kernels for Qualcomm Hexagon BackendJaved AbsarQualcomm12:30 PMLunch1:40 PMTriton for MTIARoman Levenstein et al.Meta2:00 PMUsing Triton IR for high-performance fusions in XLAGeorge KarpenkovGoogle2:20 PMTriton for All: Triton as a device-independent languageIan BearmanMicrosoft2:40 PMBreak3:00 PMPyTorch 2.0 and TorchInductorJason Ansel, Horace HeMeta3:20 PMPallas: A JAX Kernel LanguageSharad VikramGoogle3:40 PMWriting Grouped GEMMs in TritonVinod GroverNvidia4:00 PMReception下半场议程尤其值得注意1:40 PM 之后连续四个议题MTIA、XLA、设备无关语言、TorchInductor、Pallas全部来自硬件厂商之外的大型框架/编译器团队说明 Triton 在当时已经不再只是GPU 厂商的方言而是被 PyTorch、JAX、XLA 等主流框架视为内核生成与算子融合的公共技术基座。下文按议题逐一展开。主题演讲The Triton Compiler——过去、现在与未来大会开场技术演讲由 Triton 作者Phil TilletOpenAI带来主题为编译器本身的演进史。虽然会议材料未在仓库内留存详细讲义但当前仓库中的大量事实可以与这场演讲的三段式框架互证过去PastTriton 的学术根基可追溯至 2019 年 MAPL 论文Triton: An Intermediate Language and Compiler for Tiled Neural Network ComputationsREADME.md 明确要求引用该论文其核心思想是以分块tiled为单位编写神经网络计算从而在保持接近 CUDA 性能的同时大幅提升编程生产力。现在PresentREADME 的 Changelog 记载了Triton 2.0的重大转折——Backend rewritten to use MLIR后端使用 MLIR 重写并支持包含连续多个 matmul 的内核如 Flash Attention。这一事实直接对应仓库的目录结构include/triton与lib下按Dialect/Triton、TritonGPU、TritonNvidiaGPU、TritonInstrument、Gluon、Conversion/如 lib/Conversion/TritonGPUToLLVM组织是典型的 MLIR 多级方言分层设计。未来Future会议后 Triton 的演进方向——设备无关见下文 Microsoft 议题与 Gluon 显式布局编程模型见 docs/gluon/api/index.rst——如今都能在仓库中直接找到对应实现。硬件后端议题一NVIDIA Hopper 支持Gustav ZhuNVIDIA介绍了 Triton 对 Hopper 架构Compute Capability 9.0的支持。Hopper 引入了新的硬件能力wgmmaWarpgroup MMA指令、Tensor Memory AcceleratorTMA、张量内存Tensor Memory以及分布式共享内存DSMEM。当前仓库中Hopper 支持有大量直接证据后端实现在 third_party/nvidia/hopper包含 12 个.cpp与 4 个.h文件方言层面include/triton/Dialect/TritonNvidiaGPU 及其转换 passlib/Dialect/TritonNvidiaGPU/Transforms承载了wgmma、TMA、Tensor Memory 等 Hopper 特性的 IR 表示测试覆盖方面test/Conversion下存在 tritongpu_to_llvm_hopper.mlir 与 tritongpu_to_llvm_hopper_ptx80.mlirtest/TritonNvidiaGPU下有 tma_lowering.mlir 等用例Python 侧docs/gluon/api/nvidia.hopper.rst 提供了 Hopper 特有 intrinsics 的 API 文档其中可看到mma_v2、warpgroup_mma等条目python/tutorials/05-wgmma.py 则是针对 Hopper 的完整上手教程。对开发者而言这意味着自 2023 年起 Triton 便已为 Hopper 的 TMA 与 wgmma 提供了从方言到 PTX 的完整编译路径这也是 Flash Attention 等访存密集型内核能在 Triton 中获得高性能的关键前提。硬件后端议题二Bringing Triton to AMD GPUsJason Furmanek、Lixun ZhangAMD介绍了将 Triton 移植到 AMD GPU 的工作。这一议题如今已完全落地仓库中的 third_party/amd 是一个体量庞大的独立后端79 个.h、89 个.cpp、35 个 Python 文件包含后端驱动与运行时third_party/amd/backend语言扩展third_party/amd/language4 个.py库实现third_party/amd/lib与 Python 侧支持third_party/amd/python专项测试third_party/amd/test。在方言与转换层面AMD 支持同样深度集成test/Conversion/amd下拥有约 40 个 MLIR 测试文件覆盖 MFMA/WMMA 矩阵指令、buffer 原子操作、TMAtritongpu_tdm等test/TritonGPU/amd下则包含 accelerate-amd-matmul-mfma.mlir 等加速 pass 测试。README 将 AMD GPUROCm 6.2列为正式受支持硬件且TRITON_ENABLE_ASAN1地址消毒能力目前仅在 AMD 后端可用见 README.md。硬件后端议题三Intel XPU Backend for TritonEikan WangIntel介绍了 Intel XPU含 GPU 等加速硬件上的 Triton 后端移植。该议题在仓库 meetups 目录中留有后续资料docs/meetups/08-22-2023/intel-xpu-update.pptx 与 docs/meetups/10-25-2023/intel-xpu-update.pdf 分别对应 2023 年 8 月与 10 月的月度社区例会更新说明 Intel 后端的工作自大会起便持续在 Triton 月度社区会议中跟进汇报。需要注意的是当前仓库主树的 third_party 下尚无独立的 Intel 后端目录Intel 支持主要通过独立仓库/插件机制推进这与仓库中 examples/plugins 提供的插件化后端扩展能力Dialect 插件示例见 examples/plugins/README.md相呼应。硬件后端议题四面向 Qualcomm Hexagon 的 Triton 内核向量化Javed AbsarQualcomm讨论了 Triton 内核在 QualcommHexagonDSP 后端上的向量化。Hexagon 是 Qualcomm 移动 SoC 中的数字信号处理器具有独特的宽向量 SIMD 架构将 Triton 的分块张量语义映射到 Hexagon 向量单元需要专门的向量化vectorization策略——例如将tl.load/tl.store与逐元素运算转换成 Hexagon 的 HVX 向量指令。该议题在仓库主树中没有对应的后端实现目录Hexagon 后端以外部插件形式存在但可以从 Triton 的通用编译管线理解其难点Triton 的 python/triton/compiler/compiler.py 定义了从 TTGIR 到 LLVM IR 再到目标汇编的标准流水线向量化质量直接决定最终指令选择ISA selection的优劣。这也解释了为何 Triton 团队在 docs/gluon/api/index.rst 中强调 Gluon 布局layout的显式控制——布局决定线程/向量如何映射到硬件资源是跨硬件向量化的核心抽象。硬件后端议题五Triton for MTIARoman Levenstein 等人Meta介绍了 Triton 在 Meta 自研 AI 芯片MTIAMeta Training and Inference Accelerator上的应用。作为 Meta 推理基础设施的一部分MTIA 使用 Triton 编写其自定义硬件上的算子。这一议题反映了 Triton 编译器中间表示IR的可移植性设计内核可以用 Triton DSL 编写编译期将硬件无关的 Triton IR 逐步降低lower到目标架构相关的方言与指令。需要强调的是MTIA 后端同样不在当前公开仓库主树中属 Meta 内部/专有工作。但其技术路径——复用 Triton 前端与 IR替换后端 lowering 与代码生成——与仓库中third_party的目录组织每个后端一个独立子树以及TRITON_DEFAULT_BACKEND运行时选择机制见 README.md 的环境变量说明完全一致。中间表示与框架生态XLA、设备无关语言与 Pallas下午时段围绕Triton 作为公共中间层展开了三个相互关联的议题。Using Triton IR for high-performance fusions in XLAGeorge KarpenkovGoogle介绍了 XLA 如何借用Triton IR实现高性能算子融合。XLA 的 GPU 后端长期面临融合粒度问题手写 fusion 模式难以覆盖所有组合。Triton 提供的分块式内核生成能力tiled kernel generation恰好适合作为 XLA fusion 的后端——将 XLA 的融合子图翻译为 Triton IR再经由 Triton 的编译管线生成 GPU 代码。这一思路后来演化成 XLA 中的TritonGemm/triton fusion生态。对于读者而言这说明 Triton IR 在设计之初就具备被其他编译器嵌入的中间表示属性而不只是面向最终用户的前端语言。Triton for All: Triton as a device-independent languageIan BearmanMicrosoft提出了本次大会最具方向性的命题把 Triton 打造为设备无关device-independent语言。其核心主张是内核源代码应当只描述计算语义与分块策略而将布局layout、线程映射、访存模式等硬件相关决策交给编译器在不同目标上分别解决从而实现一套内核多硬件运行。当前仓库中这一愿景最具体的落地是Gluon编程模型设计文档入口见 docs/gluon/index.rstAPI 参考见 docs/gluon/api/index.rst其中语言 API 明确覆盖layouts、memory operations、math operations运行时与语言层实现在 python/triton/experimental/gluon并细分为language/、nvidia/、amd/三部分——语言核心与目标 intrinsics 分离正是设备无关设计的直接体现配套教程 python/tutorials/gluon 含 14 个示例01-intro.py 至 14-multicta.py覆盖布局、异步拷贝、TMA、wgmma、tcgen05、warp specialization 等主题测试方面python/test/gluon 提供前端、布局、lowerings 等共 7 个测试模块。此外方言层面还存在TritonToTritonGPU转换lib/Conversion/TritonToTritonGPU与 Gluon 方言include/triton/Dialect/Gluon它们共同支撑设备无关写法 → 设备相关布局的编译期分化。Pallas: A JAX Kernel LanguageSharad VikramGoogle介绍了 JAX 生态中的内核语言Pallas。Pallas 与 Triton 同属从 Python DSL 生成 GPU 内核的路线区别在于 Pallas 深度绑定 JAX 的抽象如jax.Array、jax.sharding用于在 JAX 计算图中插入自定义内核。该议题与本仓库的直接关联在于两者都面对如何把高层张量语义高效映射到硬件这一共同难题且 JAX/XLA 与 Triton 生态后续在多处产生了协作例如 Triton 被嵌入 XLA 作为 gemm 生成后端。框架集成议题PyTorch 2.0 与 TorchInductorJason Ansel、Horace HeMeta介绍了PyTorch 2.0与TorchInductor。PyTorch 2.0 引入了torch.compile其 GPU 后端 TorchInductor 将 PyTorch 算子图编译为 Triton 内核——这是 Triton 迄今为止用户量最大的集成场景之一开发者无需手写 Triton只需调用torch.compileInductor 即可自动生成并调优 Triton 内核。这一集成关系在当前仓库中有迹可循Triton 的 Python API 面向以 Python 函数为内核、triton.jit装饰、autotune 自动搜索最优配置的编程模型见 python/triton/language/core.py、python/triton/runtime/autotuner.py这正是 TorchInductor 代码生成器所消费的接口形态python/triton/compiler 下的编译驱动compiler、code_generator、make_launcher构成了内核从 Python AST 到 PTX/AMDGPU 二进制 的完整通路。大会压轴Writing Grouped GEMMs in TritonVinod GroverNVIDIA的收尾演讲聚焦Grouped GEMM的 Triton 实现这是 MoEMixture of Experts推理等场景的核心算子一批形状各异的矩阵乘法共享同一个内核实例需要通过分组调度group scheduling在设备端动态/静态地分配计算资源。仓库中现成的实战范本是 python/tutorials/08-grouped-gemm.py其头部注释点明了实现思路This group gemm kernel launches a fixed number of CTA to compute a group of gemms. The scheduling is static and we do it on device.固定数量 CTA、设备端静态调度。该教程的关键实现要素包括使用triton.autotune搜索BLOCK_SIZE_M/N/K与NUM_SM等配置见 python/tutorials/08-grouped-gemm.py 第 52 行起的triton.Config列表通过tl.dot调用硬件矩阵指令。tl.dot的签名与精度控制见 python/triton/language/core.py 第 2226 行支持input_precision参数NVIDIA 上可选tf32、tf32x3、ieee、bf16x3、bf16x6默认tf32AMD 上可选ieee、bf16x3、bf16x6CDNA3 另有tf32相关的持续回归测试位于 python/test/unit/language/test_matmul.py而test/Conversion下的 tritongpu_to_llvm_block_dot_shortcut.mlir 等文件则展示了 dot 操作在编译后端中的 lowering 路径。此外仓库中还有多组可交叉参考的分组 GEMM 素材Gluon 侧的多 CTA 示例 python/examples/gluon/03-matmul-multicta.py 与 python/examples/gluon/05-moe-bmm1-fused-gather.pyMoE 场景的 fused gather以及 Triton 内核库中的 python/triton_kernels/triton_kernels/matmul.py。结语从 2023 年议程看 Triton 的生态坐标回看 2023 年这场开发者大会的完整议程可以提炼出 Triton 生态的三个清晰脉络硬件横向扩张从 NVIDIAHopper起步逐步覆盖 AMD ROCm、Intel XPU、Qualcomm Hexagon、Meta MTIA 等目标——对应到今天仓库中的 third_party/amd 完整后端与插件化扩展机制examples/plugins成为框架的公共编译器基座PyTorch 的 TorchInductor、Google 的 XLA 融合与 Pallas 均围绕以 Triton 生成高性能内核展开——对应仓库中完备的 Python 编译驱动python/triton/compiler与tl.dot等丰富的语言原语设备无关的长期愿景Microsoft 提出的 Triton for All 在今日的 Gluon 编程模型中落地——显式布局、目标 intrinsics 分层、多后端测试矩阵python/test都是这一愿景的工程化成果。对希望深入 Triton 的读者建议结合本仓库按三条路径继续探索阅读 python/tutorials 中的 14 个基础教程上手内核编写通过 test/Conversion/tritongpu_to_llvm.mlir 与 test/Conversion/tritongpu_to_llvm_hopper.mlir 理解编译后端再通过 docs/gluon/index.rst 与 python/tutorials/gluon 体验下一代设备无关编程模型。大会后的后续社区活动资料月度例会记录与 2024 年大会材料分别存档于 docs/meetups 目录与 docs/meetups/dev_conference_2024.md。【免费下载链接】tritonDevelopment repository for the Triton language and compiler项目地址: https://gitcode.com/GitHub_Trending/tri/triton创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考