Triton 开发者会议纪要解读Block Pointer 迁移路线、性能测试基础设施与 2025 开发者峰会规划【免费下载链接】tritonDevelopment repository for the Triton language and compiler项目地址: https://gitcode.com/GitHub_Trending/tri/triton2025 年 5 月 1 日的 Triton 月度开发者会议聚焦三个核心议题Block Pointer 编程模型的迁移与弃用时间表、Triton 性能测试基础设施的建设以及 2025 Triton Developers Summit 的内容规划。本文基于 docs/meetups/05-01-2025/notes.md 的会议记录结合当前仓库源码梳理这三大议题的技术背景、社区共识与实际落地状态帮助开发者理解 Triton 内存访问编程模型的演进方向以及上游社区在性能回归测试与开发者生态建设上的协作路径。议题一Block Pointer 编程模型的去留与 Tensor Descriptor 迁移背景从厂商专属 TMA 走向通用 Tensor Descriptor会议首先由 Intel 的 Jianhui Li 提出一个关键问题现有 Block Pointer 编程模型将何去何从Intel GPU 后端目前严重依赖 Block Pointer 模型全面迁移到 Tensor Descriptor 模型需要时间。与会者普遍认同 Triton 正在从厂商专属的 TMATensor Memory Accelerator抽象走向通用 Tensor Descriptor这一方向但迁移过程涉及 API 设计、调用约定、硬件兼容性等多重问题。从当前仓库源码可以确认这一迁移方向最终落到了实处。在 python/triton/language/core.py 中make_block_ptr的实现已经变为def make_block_ptr(base: tensor, shape, strides, offsets, block_shape, order, _semanticNone): Block pointers have been removed. Use a tensor descriptor instead. raise NotImplementedError(Block pointers have been removed in favor of the tensor descriptor API)同样tl.advance也以相同的方式被移除。这印证了会议中移除 block pointer 支持的表态——Triton 官方认为 Tensor Descriptor 已能覆盖 Block Pointer 的全部使用场景。Host 侧与 Device 侧两种 Tensor Descriptor API 之争会议中一个核心争论点是仓库中存在两种 Tensor Descriptor 风格——CreateTensorDescriptorFromHostHost 侧创建与 Device 侧创建。与会者提出希望只保留一种 API倾向于保留 Host 侧版本理由包括减少程序员的选择面限制编程模型数量有助于 Triton 在更多平台上落地简化支持成本只维护一种 descriptor 更新路径对其它硬件后端更友好。OpenAI 的回答明确了设计动机Device 侧 descriptor API 之所以存在是因为需要支持在 kernel 内部更新 descriptor的场景——当 descriptor 的更新发生在设备端而非主机端时仅靠 Host 侧 API 无法满足。最终结论是两种 API 大概率都需要保留但建议开发者参考官方教程如 0.9 版本的教程练习不同 descriptor API 会表现出不同的性能特征。在仓库中两种风格都有对应实现痕迹。例如 python/triton/tools/triton_to_gluon_translator/nvidia_helpers.py 中的convert_host_descriptor以及 Gluon 方言对共享内存 descriptor、张量内存 descriptor 的完整类型体系见 python/triton/experimental/gluon/language/_core.py 中的shared_memory_descriptor_type及其slice、index、permute、reshape等操作。Block Pointer 无法直接 Lower 到 TMA 的工程现实会议中一个技术细节值得注意某位与会者指出自己用 Block Pointer 编写的 GEMM kernel 改写成 Device 侧 Tensor Descriptor 后可以正常工作但Tensor Descriptor 的 load 本身不携带偏移信息需要同时查看 load 指令与 descriptor 才能物化出 Block Pointer过程内Intra-procedural场景同一函数内可以重建 Block Pointer工作正常过程间Inter-procedural场景如果调用非内联函数descriptor 只存在于调用方caller被调用方callee拿不到重建所需的信息因而无法在 callee 内完成重建——调用约定calling convention在这种情况下会变得令人困惑。对此Triton 方的回应是Block Pointer 不能简单直接 lower 到 TMA需要设计中间 pass 将其翻译为与 Block Pointer 相似的东西如果与 CTACluster不兼容才会 lower 到 TMA。但目前这项工作尚无排期和 Peter 讨论过但没有时间投入。迁移时间表与约束条件会议明确了两点现状没有明确的时间表Block Pointer 的弃用时间线尚未确定与会者强调需要一个宽限期grace period存在功能覆盖缺口有重度的 Block Pointer 用户同时封装了两种 API 的 wrapper指出Block Pointer 是 Tensor Descriptor 的超集——他们用 Block Pointer 处理一些小规模写入而这类写入不被 TMA 支持。此外 Block Pointer 的 load 参数还有一个限制最小 stride 必须为 1其余所有 stride 必须是 16 的倍数。当前 Tensor Descriptor API 的实际约束源码级佐证会议讨论的诸多限制可以从当前仓库的 make_tensor_descriptor 实现中得到印证builtin def make_tensor_descriptor( base: tensor, shape: List[tensor], strides: List[tensor], block_shape: List[constexpr], padding_optionzero, _semanticNone, ) - tensor_descriptor: Make a tensor descriptor object :param base: the base pointer of the tensor, must be 16-byte aligned :param shape: A list of non-negative integers representing the tensor shape :param strides: A list of tensor strides. Leading dimensions must be multiples of 16-byte strides and the last dimension must be contiguous. :param block_shape: The shape of block to be loaded/stored from global memory 配套的语义检查位于 python/triton/language/semantic.py维度限制当前仅支持 15 维张量1 ndim 5末维连续性最后维度的 stride 必须为 1last_stride ! 1时报错块大小下限末维的block_shape[-1] * 元素字节数必须 ≥ 16 字节对齐要求base必须是 16 字节对齐的指针硬件映射在支持 TMA 的 NVIDIA GPUHopper 及更新架构上该 API 会生成 TMA descriptor 对象load/store 由 TMA 硬件承载。此外make_tensor_descriptor的文档示例还演示了完整用法由于 TMA descriptor 需要全局内存分配必须通过triton.set_allocator(alloc_fn)提供分配器见 python/triton/language/core.py。Gluon 中的 Tensor Descriptor 与 TMA 实践Tensor Descriptor 与 TMA 的配合使用在 Gluon 教程中有完整演示。python/tutorials/gluon/04-tma.py 的说明非常清晰Tensor Descriptor驻留在全局内存中包含 shape、strides、base pointer、layout 等信息这与会议中让 tensor descriptor 驻留在设备端的诉求一致TMA 读写本质上是异步的需要 mbarrier 对象进行同步每个 TMA 请求拷贝一个 block坐标通过偏移量指定到 block 起点越界读写由 TMA 依据 descriptor 的 shape 自动处理完成 TMA 读后通过mbarrier.wait(bar, phase0)等待TMA store 的完成由 commit group 跟踪可用tma.store_wait(pendings0)等待。教程同时提醒了两个易错点TMA descriptor 的 layout 必须是NVMMASharedLayout且stride 必须 16 字节对齐python/tutorials/gluon/04-tma.py这与上述make_tensor_descriptor的约束一脉相承。议题二Triton 性能测试基础设施的共建在最新 NVIDIA 硬件上跑通 TritonBench第二个议题由 Google 的 Sayce Falk 提出如何为 NVIDIA 最新硬件建立公开基准测试。讨论要点如下MetaCicie WangMeta 曾与 NVIDIA 讨论在 B200 上运行 TritonBenchNVIDIA 建议与 OpenAI 合作OpenAI 拥有硬件。现在 Meta 已获得硬件NVIDIA 的 Jason 正在搭建 CI第一步是在该硬件上跑通 TritonBench基础设施难题搭建 devrunner 机器的复杂度和安全性都很高社区询问能否复用 triton CI 中现有的 GB200 runner代码仓库归属可能放在facebookexperimental/triton仓库也可能另建仓库甚至考虑 PyTorch 仓库硬件来源Triton 的 Blackwell 硬件在 OpenAI 内部每个类型只有一台测试节点希望获得更好的资源PyTorch Foundation 的云实例由 Top 云厂商的 credit 资助AMD 也在视线内同时关注 AMD MI300x 与 MI350x开放诉求Xu ZhaoMeta目前已在内部运行 TritonBench但仅供 Meta 内部消费目标是对外开放。回归测试与基准测试的社区协作讨论进一步延伸到回归测试层面Intel 通常看不到来自 OpenAI 的回归因为只有一周的滞后而 Google 有 XLA 的既有经验可借鉴社区希望先从部分性能测试 部分回归测试开始包括 Llama 4 和 MoE 算子需要明确哪些 kernel 与算子应当阻塞 releaseIntel 表达了共建通用基准测试基础设施与回归测试基础设施的意愿一个有趣的提议是不要只做 lit 式测试还要观察pass 的变化如何影响生成的代码——第一步可以先量化一个 PR 影响了多少生成代码从而给出变更爆炸半径blast radius的信号Intel 曾有实习生研究过此方向感兴趣者可联系 Intel 的 Alexander。这一诉求在仓库中已有初步呼应python/triton_kernels/bench/下已存在bench_dense_matmul.py、bench_matmul_metadata.py、bench_mlp.py、bench_reduce.py等基准脚本见 python/triton_kernels/bench/bench_dense_matmul.pypython/triton_kernels/tests/下则有配套的 matmul、reduce、topk、roofline 等测试可视为社区性能/回归测试基础设施的早期形态。议题三2025 Triton Developers Summit 的内容规划往届回顾与本届安排Adnan AzizMeta回顾了去年峰会的组织情况由 Phil、Elena、Mithra 与 Adnan 共同筹备包含 MLIR 教程、主题演讲、closed-end backends、OSS 项目、Intel Triton 工作等主题覆盖异构硬件话题参会人数超过 500 人。2025 年峰会由 Microsoft 主办计划与 PyTorch 大会共址举办地点可能在旧金山 Moscone Center尚待 Microsoft 确认。社区提出的内容建议与会者提出的峰会内容方向包括用户向教程编写 Triton 代码、kernel 性能剖析器profiler使用用户座谈小组覆盖资深用户与新手用户两个群体学术与科学计算Keren 提议物理学家正用 Triton 做仿真应扩展到更广泛的 HPC 领域学习外部 DSLJason 提议EVO 与 mosaic 的演讲值得借鉴例如 Cutlass DSLTriton 应从这些项目中学习提案流程反思Cicie 提问目前没有正式的提案提交流程去年筹备时间被压缩到 10 周部分提案因时间原因未能入选。社区承诺将整理反馈转交给 Microsoft以改进 2025 年的组织流程。结语三个议题背后的共同主线纵览本次会议三条线索指向同一个方向Triton 正在从以厂商硬件为中心走向以通用、可移植的编程模型为中心。Block Pointer 的移除与 Tensor Descriptor 的普及本质上是将内存访问抽象统一到跨厂商可移植的 descriptor 模型上让 NVIDIA TMA、AMD 相关硬件与 Intel 后端共享同一套上层语义性能测试基础设施的共建则是让这一统一模型在不同硬件上都能被持续验证、防止回归开发者峰会的规划则为这些技术演进提供了社区协作与反馈的机制。对开发者而言最重要的实操建议是新写的 kernel 应直接基于 Tensor Descriptor APItl.make_tensor_descriptor注意 16 字节对齐、末维 stride 为 1、块末维 ≥ 16 字节等约束对仍在使用 Block Pointer 的存量代码应规划迁移路径并关注上游关于弃用时间表的后续公告。可参考的实践素材包括 python/triton/language/core.py 中的完整示例、python/tutorials/gluon/04-tma.py 的 TMA 教程以及 python/triton_kernels/bench/ 下的性能基准脚本。【免费下载链接】tritonDevelopment repository for the Triton language and compiler项目地址: https://gitcode.com/GitHub_Trending/tri/triton创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 Triton 开发者会议纪要解读：Block Pointer 迁移路线、性能测试基础设施与 2025 开发者峰会规划
Triton 开发者会议纪要解读Block Pointer 迁移路线、性能测试基础设施与 2025 开发者峰会规划【免费下载链接】tritonDevelopment repository for the Triton language and compiler项目地址: https://gitcode.com/GitHub_Trending/tri/triton2025 年 5 月 1 日的 Triton 月度开发者会议聚焦三个核心议题Block Pointer 编程模型的迁移与弃用时间表、Triton 性能测试基础设施的建设以及 2025 Triton Developers Summit 的内容规划。本文基于 docs/meetups/05-01-2025/notes.md 的会议记录结合当前仓库源码梳理这三大议题的技术背景、社区共识与实际落地状态帮助开发者理解 Triton 内存访问编程模型的演进方向以及上游社区在性能回归测试与开发者生态建设上的协作路径。议题一Block Pointer 编程模型的去留与 Tensor Descriptor 迁移背景从厂商专属 TMA 走向通用 Tensor Descriptor会议首先由 Intel 的 Jianhui Li 提出一个关键问题现有 Block Pointer 编程模型将何去何从Intel GPU 后端目前严重依赖 Block Pointer 模型全面迁移到 Tensor Descriptor 模型需要时间。与会者普遍认同 Triton 正在从厂商专属的 TMATensor Memory Accelerator抽象走向通用 Tensor Descriptor这一方向但迁移过程涉及 API 设计、调用约定、硬件兼容性等多重问题。从当前仓库源码可以确认这一迁移方向最终落到了实处。在 python/triton/language/core.py 中make_block_ptr的实现已经变为def make_block_ptr(base: tensor, shape, strides, offsets, block_shape, order, _semanticNone): Block pointers have been removed. Use a tensor descriptor instead. raise NotImplementedError(Block pointers have been removed in favor of the tensor descriptor API)同样tl.advance也以相同的方式被移除。这印证了会议中移除 block pointer 支持的表态——Triton 官方认为 Tensor Descriptor 已能覆盖 Block Pointer 的全部使用场景。Host 侧与 Device 侧两种 Tensor Descriptor API 之争会议中一个核心争论点是仓库中存在两种 Tensor Descriptor 风格——CreateTensorDescriptorFromHostHost 侧创建与 Device 侧创建。与会者提出希望只保留一种 API倾向于保留 Host 侧版本理由包括减少程序员的选择面限制编程模型数量有助于 Triton 在更多平台上落地简化支持成本只维护一种 descriptor 更新路径对其它硬件后端更友好。OpenAI 的回答明确了设计动机Device 侧 descriptor API 之所以存在是因为需要支持在 kernel 内部更新 descriptor的场景——当 descriptor 的更新发生在设备端而非主机端时仅靠 Host 侧 API 无法满足。最终结论是两种 API 大概率都需要保留但建议开发者参考官方教程如 0.9 版本的教程练习不同 descriptor API 会表现出不同的性能特征。在仓库中两种风格都有对应实现痕迹。例如 python/triton/tools/triton_to_gluon_translator/nvidia_helpers.py 中的convert_host_descriptor以及 Gluon 方言对共享内存 descriptor、张量内存 descriptor 的完整类型体系见 python/triton/experimental/gluon/language/_core.py 中的shared_memory_descriptor_type及其slice、index、permute、reshape等操作。Block Pointer 无法直接 Lower 到 TMA 的工程现实会议中一个技术细节值得注意某位与会者指出自己用 Block Pointer 编写的 GEMM kernel 改写成 Device 侧 Tensor Descriptor 后可以正常工作但Tensor Descriptor 的 load 本身不携带偏移信息需要同时查看 load 指令与 descriptor 才能物化出 Block Pointer过程内Intra-procedural场景同一函数内可以重建 Block Pointer工作正常过程间Inter-procedural场景如果调用非内联函数descriptor 只存在于调用方caller被调用方callee拿不到重建所需的信息因而无法在 callee 内完成重建——调用约定calling convention在这种情况下会变得令人困惑。对此Triton 方的回应是Block Pointer 不能简单直接 lower 到 TMA需要设计中间 pass 将其翻译为与 Block Pointer 相似的东西如果与 CTACluster不兼容才会 lower 到 TMA。但目前这项工作尚无排期和 Peter 讨论过但没有时间投入。迁移时间表与约束条件会议明确了两点现状没有明确的时间表Block Pointer 的弃用时间线尚未确定与会者强调需要一个宽限期grace period存在功能覆盖缺口有重度的 Block Pointer 用户同时封装了两种 API 的 wrapper指出Block Pointer 是 Tensor Descriptor 的超集——他们用 Block Pointer 处理一些小规模写入而这类写入不被 TMA 支持。此外 Block Pointer 的 load 参数还有一个限制最小 stride 必须为 1其余所有 stride 必须是 16 的倍数。当前 Tensor Descriptor API 的实际约束源码级佐证会议讨论的诸多限制可以从当前仓库的 make_tensor_descriptor 实现中得到印证builtin def make_tensor_descriptor( base: tensor, shape: List[tensor], strides: List[tensor], block_shape: List[constexpr], padding_optionzero, _semanticNone, ) - tensor_descriptor: Make a tensor descriptor object :param base: the base pointer of the tensor, must be 16-byte aligned :param shape: A list of non-negative integers representing the tensor shape :param strides: A list of tensor strides. Leading dimensions must be multiples of 16-byte strides and the last dimension must be contiguous. :param block_shape: The shape of block to be loaded/stored from global memory 配套的语义检查位于 python/triton/language/semantic.py维度限制当前仅支持 15 维张量1 ndim 5末维连续性最后维度的 stride 必须为 1last_stride ! 1时报错块大小下限末维的block_shape[-1] * 元素字节数必须 ≥ 16 字节对齐要求base必须是 16 字节对齐的指针硬件映射在支持 TMA 的 NVIDIA GPUHopper 及更新架构上该 API 会生成 TMA descriptor 对象load/store 由 TMA 硬件承载。此外make_tensor_descriptor的文档示例还演示了完整用法由于 TMA descriptor 需要全局内存分配必须通过triton.set_allocator(alloc_fn)提供分配器见 python/triton/language/core.py。Gluon 中的 Tensor Descriptor 与 TMA 实践Tensor Descriptor 与 TMA 的配合使用在 Gluon 教程中有完整演示。python/tutorials/gluon/04-tma.py 的说明非常清晰Tensor Descriptor驻留在全局内存中包含 shape、strides、base pointer、layout 等信息这与会议中让 tensor descriptor 驻留在设备端的诉求一致TMA 读写本质上是异步的需要 mbarrier 对象进行同步每个 TMA 请求拷贝一个 block坐标通过偏移量指定到 block 起点越界读写由 TMA 依据 descriptor 的 shape 自动处理完成 TMA 读后通过mbarrier.wait(bar, phase0)等待TMA store 的完成由 commit group 跟踪可用tma.store_wait(pendings0)等待。教程同时提醒了两个易错点TMA descriptor 的 layout 必须是NVMMASharedLayout且stride 必须 16 字节对齐python/tutorials/gluon/04-tma.py这与上述make_tensor_descriptor的约束一脉相承。议题二Triton 性能测试基础设施的共建在最新 NVIDIA 硬件上跑通 TritonBench第二个议题由 Google 的 Sayce Falk 提出如何为 NVIDIA 最新硬件建立公开基准测试。讨论要点如下MetaCicie WangMeta 曾与 NVIDIA 讨论在 B200 上运行 TritonBenchNVIDIA 建议与 OpenAI 合作OpenAI 拥有硬件。现在 Meta 已获得硬件NVIDIA 的 Jason 正在搭建 CI第一步是在该硬件上跑通 TritonBench基础设施难题搭建 devrunner 机器的复杂度和安全性都很高社区询问能否复用 triton CI 中现有的 GB200 runner代码仓库归属可能放在facebookexperimental/triton仓库也可能另建仓库甚至考虑 PyTorch 仓库硬件来源Triton 的 Blackwell 硬件在 OpenAI 内部每个类型只有一台测试节点希望获得更好的资源PyTorch Foundation 的云实例由 Top 云厂商的 credit 资助AMD 也在视线内同时关注 AMD MI300x 与 MI350x开放诉求Xu ZhaoMeta目前已在内部运行 TritonBench但仅供 Meta 内部消费目标是对外开放。回归测试与基准测试的社区协作讨论进一步延伸到回归测试层面Intel 通常看不到来自 OpenAI 的回归因为只有一周的滞后而 Google 有 XLA 的既有经验可借鉴社区希望先从部分性能测试 部分回归测试开始包括 Llama 4 和 MoE 算子需要明确哪些 kernel 与算子应当阻塞 releaseIntel 表达了共建通用基准测试基础设施与回归测试基础设施的意愿一个有趣的提议是不要只做 lit 式测试还要观察pass 的变化如何影响生成的代码——第一步可以先量化一个 PR 影响了多少生成代码从而给出变更爆炸半径blast radius的信号Intel 曾有实习生研究过此方向感兴趣者可联系 Intel 的 Alexander。这一诉求在仓库中已有初步呼应python/triton_kernels/bench/下已存在bench_dense_matmul.py、bench_matmul_metadata.py、bench_mlp.py、bench_reduce.py等基准脚本见 python/triton_kernels/bench/bench_dense_matmul.pypython/triton_kernels/tests/下则有配套的 matmul、reduce、topk、roofline 等测试可视为社区性能/回归测试基础设施的早期形态。议题三2025 Triton Developers Summit 的内容规划往届回顾与本届安排Adnan AzizMeta回顾了去年峰会的组织情况由 Phil、Elena、Mithra 与 Adnan 共同筹备包含 MLIR 教程、主题演讲、closed-end backends、OSS 项目、Intel Triton 工作等主题覆盖异构硬件话题参会人数超过 500 人。2025 年峰会由 Microsoft 主办计划与 PyTorch 大会共址举办地点可能在旧金山 Moscone Center尚待 Microsoft 确认。社区提出的内容建议与会者提出的峰会内容方向包括用户向教程编写 Triton 代码、kernel 性能剖析器profiler使用用户座谈小组覆盖资深用户与新手用户两个群体学术与科学计算Keren 提议物理学家正用 Triton 做仿真应扩展到更广泛的 HPC 领域学习外部 DSLJason 提议EVO 与 mosaic 的演讲值得借鉴例如 Cutlass DSLTriton 应从这些项目中学习提案流程反思Cicie 提问目前没有正式的提案提交流程去年筹备时间被压缩到 10 周部分提案因时间原因未能入选。社区承诺将整理反馈转交给 Microsoft以改进 2025 年的组织流程。结语三个议题背后的共同主线纵览本次会议三条线索指向同一个方向Triton 正在从以厂商硬件为中心走向以通用、可移植的编程模型为中心。Block Pointer 的移除与 Tensor Descriptor 的普及本质上是将内存访问抽象统一到跨厂商可移植的 descriptor 模型上让 NVIDIA TMA、AMD 相关硬件与 Intel 后端共享同一套上层语义性能测试基础设施的共建则是让这一统一模型在不同硬件上都能被持续验证、防止回归开发者峰会的规划则为这些技术演进提供了社区协作与反馈的机制。对开发者而言最重要的实操建议是新写的 kernel 应直接基于 Tensor Descriptor APItl.make_tensor_descriptor注意 16 字节对齐、末维 stride 为 1、块末维 ≥ 16 字节等约束对仍在使用 Block Pointer 的存量代码应规划迁移路径并关注上游关于弃用时间表的后续公告。可参考的实践素材包括 python/triton/language/core.py 中的完整示例、python/tutorials/gluon/04-tma.py 的 TMA 教程以及 python/triton_kernels/bench/ 下的性能基准脚本。【免费下载链接】tritonDevelopment repository for the Triton language and compiler项目地址: https://gitcode.com/GitHub_Trending/tri/triton创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考