
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梳理本次会议的核心议题H100Hopper初步支持与 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围绕六项议题展开H100Hopper支持更新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 TFLOPScuBLAS 在随机输入数据上约为 670 TFLOPSFP8 则约为其两倍达到 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时会追加fp8e4nvNVIDIA 专用 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_warpspecHopper warp 特化、assign_latencies、schedule_loops、pipeline等一系列针对 Hopper 的优化与流水线 passmake_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 第三方插件OAIOpenAI没有承诺额外资源但社区可以贡献。源码印证插件机制与架构覆盖上述第三方插件路线在当前仓库中已有具体实现载体插件目录约定python/triton侧通过TRITON_PLUGIN_DIRS环境变量指定插件目录仓库内置的插件示例位于 examples/plugins/DialectPlugins/DialectPlugin包含 CMake 构建脚本与自定义 dialect/pass 实现并在 examples/plugins/README.md 中给出说明后端驱动架构NVIDIA 后端实现了BaseBackend接口CUDABackend见 third_party/nvidia/backend/compiler.pyAMD 后端位于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 团队分享了 XPUIntel 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.py3.0 发布 / 第三方插件TRITON_PLUGIN_DIRS插件机制与 examples/plugins/ 示例Linalg / block pointerstl.make_block_ptr语言原语与张量描述符改写 passIntel GPU/CPU 后端会议幻灯片 intel-xpu-update.pptxAMD 后端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),仅供参考