Triton 2023 年 8 月开发者例会纪要解读:Hopper FP8 支持、3.0 发布路线与多后端生态布局
【免费下载链接】tritonDevelopment repository for the Triton language and compiler项目地址: https://gitcode.com/GitHub_Trending/tri/triton
本篇基于 Triton 项目 2023-08-22 开发者例会的会议纪要(docs/meetups/08-22-2023/notes.md),梳理本次会议的核心议题:H100(Hopper)初步支持与 FP8 性能数据、Triton 3.0 发布计划、Linalg 中间层讨论,以及 Intel GPU/CPU 后端与 AMD 后端的生态进展。结合当前仓库源码,读者可以了解这些规划在 Triton 编译流水线中的具体落点(如 CUDA 后端的 FP8 支持逻辑、Hopper 专属 pass 序列、block pointer 语言原语等),并掌握如何通过环境变量与编译选项验证和使用这些能力。
会议概况与议程
本次会议是 Triton 项目 2023 年 8 月的例行开发者例会,会议纪要记录于docs/meetups/08-22-2023/notes.md。议程(Agenda)围绕六项议题展开:
- H100(Hopper)支持更新;
- Triton 发布计划更新;
- Linalg 中间层讨论;
- Intel GPU 后端状态更新;
- Intel 正在进行的 CPU 后端工作;
- AMD 后端更新;
- 开放式讨论。
同时,例会预告了首届 Triton 开发者大会:大会定于 2023 年 9 月 20 日在微软硅谷园区(Microsoft Silicon Valley Campus)举行,注册即将开放。本次会议的全部录音可查阅会议纪要中提供的录制链接。
从仓库结构看,docs/meetups/目录按日期组织例会笔记(如docs/meetups/08-22-2023/、docs/meetups/10-25-2023/等),是了解 Triton 社区演进与路线图的第一手资料;本次会议还配套了两份幻灯片:Intel XPU 后端更新 与 AMD 更新,分别存放于同一目录下。
H100 支持更新:Hopper 与 FP8 的关键进展
本次会议的第一项议题是 H100 支持状态。纪要归纳了以下几点:
- 初步支持已合入主分支,但默认关闭,需要通过环境变量启用;
- 支持最新的张量核心(Tensor Cores)与 FP8 数据类型;Flash Attention 支持即将合入主分支;
- 矩阵乘法(Matmul)性能表现优异:在大规模 Matmul 上已达到 cuBLAS 的 80%~90%,目标是与 cuBLAS 对齐;在 xxm 卡上 fp16 性能超过 600 TFLOPS(cuBLAS 在随机输入数据上约为 670 TFLOPS),FP8 则约为其两倍,达到 1.2 PFLOPS 量级;
- Hopper 支持包含完整的 FP8 计算支持。
源码印证:CUDA 后端的 FP8 与 Hopper 支持
从当前仓库的 NVIDIA 后端实现(third_party/nvidia/backend/compiler.py)可以看到与纪要对应的工程落点:
- FP8 数据类型开关:
CUDAOptions中定义了supported_fp8_dtypes: Tuple[str] = ("fp8e5", "fp8e4b15"),并在parse_options中根据 GPU 算力动态扩展——当capability >= 89(即 Ada Lovelace 及之后的 Hopper/Blackwell)时,会追加"fp8e4nv"(NVIDIA 专用 FP8 E4M3 变体)。这印证了纪要中"完整的 FP8 支持"并非静态配置,而是随目标架构能力动态启用的。 - 架构级代码生成适配:
get_codegen_implementation根据算力选择 FP8 类型转换函数——capability >= 80使用convert_custom_float8_sm80,否则回退到convert_custom_float8_sm70,说明 FP8 代码生成在 Ampere 与 Hopper 两代架构上有不同的指令路径。 - Hopper 专属编译 pass 序列:在
make_ttgir中,当capability // 10 in [8, 9](Ampere/Hopper)时,会执行nvidia.passes.hopper.add_hopper_warpspec(Hopper warp 特化)、assign_latencies、schedule_loops、pipeline等一系列针对 Hopper 的优化与流水线 pass;make_ttir中capability // 10 < 9时还会执行rewrite_tensor_descriptor_to_pointer将张量描述符改写为指针形式。这些正是支撑 Hopper 上 FP8 Matmul 高性能的编译期基础。
需要说明的是,纪要中的具体性能数字(600+ TFLOPS、1.2 PFLOPS 等)属于当时开发中状态的基准测试结果,不代表当前仓库版本的实测性能;如需复现,应以本地硬件与当前版本的基准测试(如 python/triton_kernels/bench/bench_dense_matmul.py)为准。
Triton 3.0 发布计划:版本策略与生态目标
会议第二项议题讨论了 Triton 的版本发布规划:
- 时间窗口:尚无具体日期,但计划在 2023 年底前发布;
- 主版本升级至 3.0:由于存在少量向后兼容性破坏的改动(例如将索引算子中的编译器选项改为内核中的硬编码算子),将提升主版本号;
- 功能目标:核心目标是让 Intel 与 AMD GPU 拥有第三方插件(3rd party plugins)支持;
- 与 PyTorch 的同步策略:可能配合一次 PyTorch 发布,使 PyTorch 受益于最新特性;但持续集成(CI)工作流仍是默认的发布节奏预期;
- 默认优化模式:发布时将默认切换为优化模式,该决策需要与 NVIDIA 进一步讨论;
- 开放内核选择:将向用户暴露标志(flags),允许用户自行启用内核选择(kernel selection);
- 待解问题:PyTorch 尚未 rebase 到最新 Triton,且临近 PyTorch 代码冻结——PyTorch 是否会同步 Triton 2.0?是否需要为支持 Triton 2.0 再做一次发布?
- 社区协作方式:社区可以从最新稳定分支出发,在其上 rebase 第三方插件;OAI(OpenAI)没有承诺额外资源,但社区可以贡献。
源码印证:插件机制与架构覆盖
上述"第三方插件"路线在当前仓库中已有具体实现载体:
- 插件目录约定:
python/triton侧通过TRITON_PLUGIN_DIRS环境变量指定插件目录,仓库内置的插件示例位于 examples/plugins/DialectPlugins/DialectPlugin(包含 CMake 构建脚本与自定义 dialect/pass 实现),并在 examples/plugins/README.md 中给出说明; - 后端驱动架构:NVIDIA 后端实现了
BaseBackend接口(CUDABackend,见 third_party/nvidia/backend/compiler.py),AMD 后端位于third_party/amd/backend/,Intel XPU 后端则通过插件形式接入——这与纪要中"Intel/AMD 以第三方插件形态参与"的路线图一致; - 内核选择相关选项:
CUDAOptions中与编译/优化相关的开关(如enable_fp_fusion、maxnreg、num_warps、ptx_options等)均可在运行时通过编译选项传入,为用户侧精细控制内核生成提供了入口。
Linalg 中间层讨论:语言与目标硬件之间的桥梁
会议第三项议题讨论了在 GitHub 上进行的 Linalg 方案讨论:将 Linalg 作为语言(Triton DSL)与目标硬件之间的中间层,其内容包括对 block pointers(块指针)与取模运算符(modulo operators)的支持。讨论对应的分支当时落后于主干 tip,后续计划将其同步到最新。
源码印证:block pointer 已是语言一级原语
纪要中提到的 "block pointers" 在当前仓库中已成为 Triton 语言的核心特性:
- 语言前端提供
triton.language.make_block_ptr原语(定义于 python/triton/language/core.py),用于构造块指针,可指定base、shape、strides、offsets、block_shape与order等参数,从而表达按块访问张量数据的语义; - 张量描述符(tensor descriptor)相关改写逻辑分散在 TTIR 与 TTGIR 转换 pass 中(如 lib/Conversion/TritonToTritonGPU/TritonGPUConversion.cpp),并配套了 test/Triton/rewrite-tensor-descriptor-to-pointer.mlir 等测试;
- 关于 Linalg 方案的完整讨论过程记录于当时的 GitHub 讨论线程(讨论编号 1842),仓库内不包含该讨论的正文内容,此处不再展开。
Intel 后端进展:GPU 与 CPU 双线并行
会议第四、五项议题聚焦 Intel 的贡献:
- GPU 后端状态更新:Intel 团队分享了 XPU(Intel GPU)后端的进展,详见会议配套幻灯片 Intel XPU 后端更新;
- CPU 后端:Intel 正在着手为 Triton 开发 CPU 后端,相关细节同样见于上述幻灯片。
从仓库现状看,Intel XPU 后端的实现以third_party/intel形态演进,可通过TRITON_PLUGIN_DIRS以插件方式加载;CPU 后端则属于当时公布的在研方向,仓库内尚无对应的合入实现。两者的详细设计与进度以官方幻灯片和后续例会笔记为准。
AMD 后端更新
会议第六项议题为 AMD 后端更新,具体细节参见配套幻灯片 AMD 更新。从当前仓库看,AMD 相关支持已相当完整:
third_party/amd/下包含完整的后端实现(backend/、lib/、include/、python/、test/等子目录),其中third_party/amd/lib/下约 89 个源文件覆盖了从 dialect 到 LLVM 转换的编译路径;- 测试方面,test/Conversion/amd/ 与 test/TritonGPU/amd/ 目录包含大量 AMD 专属的 MLIR 测试用例(如 mfma/wmma dot 加速、异步拷贝、buffer 操作等),与本次会议纪要中"AMD 后端持续推进"的信息相互印证。
从例会到代码:本次纪要的价值与局限
综合来看,这份例会纪要的价值在于它记录了 Triton 生态在 2023 年下半年的关键节点:Hopper/FP8 支持进入主分支、3.0 主版本发布规划、Linalg 中间层讨论的发起,以及 Intel/AMD 多后端生态的布局。这些议题大多在当前仓库中能找到对应的实现痕迹:
| 纪要议题 | 仓库中的对应落点 |
|---|---|
| H100/FP8 支持 | CUDAOptions.supported_fp8_dtypes与算力相关的 FP8 动态启用(third_party/nvidia/backend/compiler.py) |
| 3.0 发布 / 第三方插件 | TRITON_PLUGIN_DIRS插件机制与 examples/plugins/ 示例 |
| Linalg / block pointers | tl.make_block_ptr语言原语与张量描述符改写 pass |
| Intel GPU/CPU 后端 | 会议幻灯片 intel-xpu-update.pptx |
| AMD 后端 | third_party/amd/ 全量实现与 test/Conversion/amd/ 测试 |
需要提醒读者的是:纪要中的性能数字与时间规划均反映 2023 年 8 月时点的状态,会议中提到的"Flash Attention 即将合入""年底发布 3.0"等属于当时的预期,不应作为当前版本能力的断言;若要验证最新能力,应直接参考仓库中RELEASE.md、README.md与各后端测试用例。对于想要深度参与多后端生态的开发者,从docs/meetups/的历次例会笔记入手跟踪路线图,再结合third_party/下各后端源码与test/下的 MLIR 测试进行实践,是一条高效的学习路径。
【免费下载链接】tritonDevelopment repository for the Triton language and compiler项目地址: https://gitcode.com/GitHub_Trending/tri/triton
创作声明:本文部分内容由AI辅助生成(AIGC),仅供参考