Skip to content

Triton Version Upgrade Report

Jump to bottom
zhzhcookie edited this page Jul 15, 2026 · 1 revision

Triton 3.7 版本变更报告

1. 对比范围与结论摘要

本报告比较 Triton 的两个 Release 分支：

  • Triton 3.6：标签 v3.6.0，提交 7c56a5e40f7fd928dfd5c72902d5def0097db73a，发布日期 2026-01-20。
  • Triton 3.7：标签 v3.7.1，提交 f797708c0626e5f9840ca5b0a98790e2c7cb09ad，发布日期 2026-06-16。

两个发布分支从提交 85400f80bf859a34ad7a746ffda877faf80312ab 分叉；3.6 分支在共同基线后还有 6 个提交，3.7 分支有 408 个提交。因此本报告以两个最终源码快照的文件差异为准，以 Git 提交历史、测试和文档为交叉证据，而不是简单把 v3.6.0..v3.7.1 当作线性提交区间。

快照对比约涉及 706 个文件、55,540 行新增和 18,530 行删除。总体上，3.7 是一次较大的功能版本：

  1. Python 稳定 API 以增量扩展为主：新增 tl.squeeze、tl.unsqueeze、公开 tl.to_tensor，完善 tl.cat，Tensor Descriptor 增加显式 TF32 舍入。
  2. Python 兼容性风险主要位于实验性 Gluon、AOT 工具、Proton 和独立的 triton_kernels 包。
  3. 新增完整的树外 TTIR/TTGIR dialect/pass/op 插件体系。
  4. NVIDIA 侧重点是 Blackwell/Thor、TMA im2col、tcgen05 multicast/ld.red 和 Warp Specialization。
  5. AMD 侧重点是 gfx1250、TDM scatter/gather、Warp Pipeline/Specialization、multi-CTA WMMA 和 MIR 调试。
  6. TTGIR 的多 CTA 布局接口从 CTA 统一重命名为 CGA，对直接处理 IR 或编写 pass 的下游代码有破坏性。
  7. LLVM 最终升级到 1f126a6，Blackwell CUDA 工具链升级到 13.1。

2. 需要优先处理的兼容性变化

升级前建议优先搜索并处理以下用法：

  • tl.make_block_ptr：3.7 开始每次调用都发出弃用警告；应迁移至 tl.make_tensor_descriptor 或 triton.tools.tensor_descriptor.TensorDescriptor。
  • Gluon thread_barrier()：删除并重命名为 barrier(*, cluster=False)。
  • triton.tools.ragged_tma：batch_offset/batch_size 形参改名为 slice_off/slice_size，关键字调用会失效。
  • AOT：3.7 改变 specialization suffix 和生成符号名，并引入 CUDA/HIP 通用链接类型；3.6 生成物不能与 3.7 混用，必须重新编译和链接。
  • TTGIR/原生 pass：CTAEncodingAttr、CTALayout、getCTALayout() 分别迁移到 CGAEncodingAttr、CGALayout、getCGALayout()。
  • triton_kernels：matmul_ogs 整套 API 重构为 matmul，路由、布局和分布式接口均有破坏性变化。
  • JIT preload：specialization 数据现在绑定编译 target；跨 target preload 会抛出 RuntimeError。
  • 自定义 JIT cache hook：hook 收到的 specialization JSON 新增 target，调用链内部也新增 target 参数。

3. Python API 完整差异

本节以 __all__、包级重导出、公开文档、测试和无下划线符号为准。以下模块的公开 API 未发生变化：triton.language.math、random、extra、target_info、triton.runtime.autotuner、runtime.errors、runtime.build、triton.backends 和 triton.errors。

3.1 顶层 triton

文件：python/triton/__init__.py

  • triton.__version__：3.6.0 → 3.7.1。

  • __all__ 新增 max_shared_mem。

  • 存在一个导出缺陷：顶层 __all__ 虽包含 max_shared_mem，但没有从 .compiler 导入该符号。因此可靠写法是：

    from triton.compiler import max_shared_mem

    不应依赖 from triton import max_shared_mem。

除上述项目外，jit、constexpr_function、compile、autotune、Config、heuristics、AsyncCompileMode、FutureKernel、JITFunction、set_allocator、cdiv、next_power_of_2 等顶层导出保持不变。

3.2 triton.language

文件：python/triton/language/__init__.py、core.py、standard.py、semantic.py

3.2.1 新增公开符号

  • tl.squeeze(x, dim: tl.constexpr)
    • 要求 x.shape[dim] == 1，删除指定维度后 reshape。
  • tl.unsqueeze(x, dim: tl.constexpr)
    • 在指定位置插入长度为 1 的维度。
  • tl.to_tensor(x, _semantic=None)
    • 3.6 已有内部实现，3.7 首次从 triton.language 正式导出。

tl.topk 不是 3.7 新 API：它在 3.6 已存在，3.7 只是把它加入 docs/python-api/triton.language.rst 并补充说明。

3.2.2 签名变化

tl.cat：

# 3.6
cat(input, other, can_reorder=False, _semantic=None)

# 3.7
cat(input, other, can_reorder=False, dim=0, _semantic=None)

3.7 在 can_reorder=False 时通过 join、permute、reshape 实现确定性拼接，并支持选择 dim；负维度会被规范化。输入 rank 必须相同，除拼接维外的 shape 必须一致。can_reorder=True 仍走原有编译器 hint 路径。

3.2.3 弃用与语义变化

  • tl.make_block_ptr(...)
    • 签名未变。
    • 3.7 新增运行时 warning，建议改用 Tensor Descriptor。
    • warning 未显式指定 DeprecationWarning 类别，实际是默认 warning。
  • tl.dot(...)
    • 签名和发布态默认精度未变。
    • 文档新增 TF32 输入可能被截断且产生偏差的警告，建议在精度敏感场景显式舍入，或在 Tensor Descriptor 上设置 round_f32_to_tf32=True。
    • 3.7 开发期间曾把默认 32 位 dot 精度改为 TF32x3，但发布前已 revert；不能把该中间态当作 3.7 最终行为。
  • tl.dot_scaled(...)
    • scale shape 校验扩展为允许前导 batch 维，从而支持 scaled BMM。
  • FP8 常量
    • semantic.full/scalar constant 路径开始支持 FP8 常量。
  • constexpr/JIT 返回值
    • JIT 函数可以返回 constexpr。
    • constexpr mangling 支持值对象自己的 mangle()。
  • tuple/namedtuple
    • constexpr 函数会保留 namedtuple 的类型和字段信息。
    • tuple_type.fields 现在可以为 None。
    • 对不存在的具名字段访问会明确抛出 AttributeError。
  • builtin min/max
    • 传入 propagate_nan 或对 tensor 使用 Python builtin 的 warning 改为默认可见 warning，而不是被默认过滤的 DeprecationWarning。
  • builtin 签名
    • builtin wrapper 预先缓存 inspect.signature，主要是前端编译性能优化，对调用签名无影响。

3.2.4 导出稳定性

tl.load、store、dot、dot_scaled、make_tensor_descriptor、tensor_descriptor、gather、histogram、associative_scan、range、static_range 等原有导出均未删除。

3.3 triton.compiler

文件：python/triton/compiler/__init__.py、compiler.py

新增公开 API：

max_shared_mem(device)

它返回 active driver 的 get_device_properties(device)["max_shared_mem"]。

原有 compile、make_backend、ASTSource、IRSource、CompiledKernel、CompilationError、LazyDict、get_cache_key 保持导出。

行为变化：

  • knobs.runtime.add_stages_inspection_hook 的标识和哈希会进入 JIT specialization/cache key，避免修改自定义 pass pipeline 后错误复用旧缓存。
  • 该 hook 存在必须注意的契约漂移：PipelineStagesHook Protocol 和 backend 的实际 stage inspection 仍按 hook(backend, stages, options, language, capability) 调用，但 compiler.compile() 与 JITFunction.run() 又会先按 hook() 零参数调用，并期望返回 (cache_key_suffix, hash)。 因此 3.6 风格的固定五参数 hook 在 3.7 会触发 TypeError；兼容实现必须同时处理零参数和五参数调用。
  • 支持通过插件向 TTIR/TTGIR pipeline 注入树外 pass，详见第 6 节。

3.4 triton.runtime

python/triton/runtime/__init__.py 的公开导出列表未变，但 JIT、driver、异步编译和缓存行为有重要变化。

3.4.1 Driver 选择

新增环境变量：

TRITON_DEFAULT_BACKEND=<backend-name>

当系统存在多个 active driver/backend 时，可以显式选择默认 backend；未设置时仍使用自动发现逻辑。

3.4.2 JIT specialization 与 preload

文件：python/triton/runtime/jit.py

  • specialization JSON 新增 target 字段。
  • JITFunction.preload(...) 会检查保存的 target 与当前 target 完全一致；不一致时抛出 RuntimeError。
  • constexpr 可以包含 JITFunction，序列化形式为 {"jit_function": "module:qualname"}。
  • 3.7 新增进程内 JITFunction 注册表以恢复上述引用；未注册的函数会导致 preload 失败。
  • JITCallable.get_capture_scope() 不再调用 inspect.getclosurevars，改为直接读取 closure cells，降低前端开销。
  • 自定义 pipeline hook 的哈希加入 specialization。

@triton.jit、@constexpr_function、JITFunction、KernelInterface 的公开调用签名不变。

3.4.3 异步编译

文件：python/triton/runtime/_async_compile.py

  • FutureKernel 新增 __getattr__，首次访问底层 kernel 属性时透明等待编译完成。
  • Future 会更早放入 kernel cache，避免重复提交相同的异步编译。
  • 异步编译增加 sentinel 机制，改善进程池故障检测。

3.4.4 缓存

文件：python/triton/runtime/cache.py

  • 读取损坏或不完整的 group JSON 时返回 cache miss 并重新编译，不再让损坏缓存直接终止程序。

3.5 triton.knobs 与环境变量

文件：python/triton/knobs.py

新增或改变的公开配置：

  • AMD：
    • knobs.amd.use_async_copy 从二值 env_bool 改为可区分 unset/true/false 的 env_opt_bool。
    • 新增 knobs.amd.dump_mir / TRITON_DUMP_MIR。
    • 新增 knobs.amd.swap_mir / TRITON_SWAP_MIR。
  • Proton：
    • 新增 knobs.proton.profile_buffer_size / TRITON_PROFILE_BUFFER_SIZE，默认 64 MiB。
    • 新增 knobs.proton.enable_hw_trace / TRITON_ENABLE_HW_TRACE，用于 Blackwell 低开销硬件 trace。
  • Runtime：
    • add_stages_inspection_hook 的结果现在影响 cache key。
  • Backend：
    • 新增 TRITON_DEFAULT_BACKEND。
  • 插件：
    • 新增 TRITON_PLUGIN_PATHS、TRITON_PLUGIN_VERSION_CHECK 和构建期开关 TRITON_EXT_ENABLED。
  • Warp Specialization 调试：
    • 新增 TRITON_PARTITION_SCHEDULING_ENABLE_DUMP_DOT。
    • 新增 TRITON_PARTITION_SCHEDULING_DUMP_DATA_ONLY。
    • 新增 TRITON_PARTITION_SCHEDULING_DUMP_LOOP_ONLY。

注意：3.7 发布态没有默认开启 AMD async copy；相关默认值变更已在发布前 revert，若需要必须显式配置。

3.6 triton.testing

文件：python/triton/testing.py

  • assert_close(x, y, ..., err_msg="") 在数组分支也会把 err_msg 传给 numpy.testing.assert_allclose。
  • 其余主要 API，如 do_bench、Benchmark、perf_report、get_dram_gbps 等，签名未变。

3.7 triton.tools

3.7.1 Tensor Descriptor

文件：python/triton/tools/tensor_descriptor.py

TensorDescriptor dataclass 新增：

round_f32_to_tf32: bool = False

工厂方法变为：

TensorDescriptor.from_tensor(
    tensor,
    block_shape,
    padding="zero",
    round_f32_to_tf32=False,
)

该选项只允许用于 FP32 tensor；其他 dtype 会触发断言。它用于在 descriptor load 路径显式把 FP32 舍入为 TF32。

3.7.2 Ragged TMA

文件：python/triton/tools/ragged_tma.py

以下公开函数的参数名从 batch_offset, batch_size 改为 slice_off, slice_size：

  • to_ragged_indices
  • load_ragged
  • store_ragged
  • atomic_add_ragged

位置参数调用语义基本不变，关键字调用不兼容。

3.7.3 AOT compile/link

文件：python/triton/tools/compile.py、link.py

CLI 参数基本未变，但生成物 ABI/命名规则改变：

  • 3.6 会把每个参数 index 都写进 specialization suffix；3.7 只为有 c/d hint 的参数写入 <index>c 或 <index>d。
  • C 符号名和文件名随之改变。
  • 生成头从 CUDA 专用 CUresult/CUstream 抽象为 TT_ResultTy/TT_StreamTy。
  • 模板增加 tt-linker-backend 标记和 backend_name。
  • 新增 HIP AOT link 模板。
  • Header parser 允许空 suffix，并重写 suffix 匹配算法。

升级后必须重新运行 triton.tools.compile 和 triton.tools.link，不能复用 3.6 产物。

triton.tools.__init__ 仍只重导出 LinearLayout；build_extern、disasm、mxfp 的公开签名未变。

3.8 triton._internal_testing

该模块不是稳定 API，但仓库测试和下游测试可能直接使用：

  • is_hip_gfx11() → is_hip_rdna3()
  • is_hip_gfx12() → is_hip_rdna4()
  • 新增 is_hip_rdna()
  • 新增 is_blackwell_ultra()

4. 实验性 Gluon Python API

路径：python/triton/experimental/gluon

Gluon 明确属于 experimental，但本节完整列出本次快照中可见的接口变化。

4.1 通用语言层

破坏性重命名：

# 3.6
thread_barrier()

# 3.7
barrier(*, cluster: bool = False)
  • barrier() 是 thread-block barrier。
  • barrier(cluster=True) 是 cluster barrier。

其他变化：

  • warp_specialize(functions_and_args, worker_num_warps, worker_num_regs=None, ...)
    • worker_num_regs 从必填变为可选。
  • tensor.gather(indices, axis)：新增 shared-memory/local gather。
  • tensor.scatter(values, indices, axis)：新增 shared-memory/local scatter。
  • 布局类新增 format_tensor_view()、format_hardware_view() 一类诊断格式化能力。
  • 多个 verifier 和错误信息收紧，原先可能进入 lowering 的非法布局会更早失败。

4.2 AMD Gluon

新增 triton.experimental.gluon.language.amd.warp_pipeline_stage 上下文管理器，用于给 gfx1250 Warp Pipeline 标记 stage 和优先级。

新增 gfx1250 cluster 模块：

triton.experimental.gluon.language.amd.gfx1250.cluster.arrive()
triton.experimental.gluon.language.amd.gfx1250.cluster.wait()

新增或扩展的 gfx1250 TDM API：

  • async_scatter(...)
  • async_gather(...)
  • prefetch(...)
  • shared_to_global(...)
  • 更完整的 1D–5D descriptor 支持。

get_wmma_scale_layout 的参数从 tiles_per_warp, warps_per_cta 泛化为 reg_bases, warp_bases。直接构造实验性 AMD 布局的代码需要迁移。

4.3 NVIDIA Gluon

新增 Hopper cluster 模块：

cluster.arrive(relaxed=False)
cluster.wait()

TMA/cluster API 变化：

  • mbarrier.allocate_mbarrier(...) 新增并公开。
  • mbarrier.expect(mbarrier, bytes, ...) 改为 expect(mbarrier, bytes_per_cta=None, ...)；按关键字传 bytes= 的代码不兼容。
  • async_copy_global_to_shared(..., multicast=False, ...) 新增 multicast。
  • fence_init_release_cluster() 新增。
  • sync_cluster_init() 新增。
  • Tensor Descriptor 增加 round_f32_to_tf32。
  • 对 descriptor、coordinate 的对齐和非负约束增加运行时/编译时检查。

Blackwell/tcgen05 API 变化：

  • tcgen05_mma(..., multicast=False, ...) 新增 multicast。
  • tcgen05_commit(barrier, pred=True, descs=(), ...) 用 descriptor 列表表达 multicast commit。
  • 新增 tcgen05_mma_barrier_count(...)。
  • tensor memory descriptor 增加 load_min / load_max，用于 tcgen05.ld.red。
  • 新增 TensorMemoryScalesLayout。

部分内部/实验符号虽然进入 __all__，但只用于 IR 打印，例如 _TensorMemoryLinearLayout；它不应作为可物化布局使用。

5. Proton / triton.profiler

third_party/proton/proton 通过 setup.py 安装为 triton.profiler。3.7 对其 Python 和 C++ API 做了较大扩展。

5.1 新增包级 API

  • metadata_state
  • data 子模块

全新的 triton.profiler.data 提供：

get(session=0, phase=0)
get_msgpack(session=0, phase=0)
advance_phase(session=0)
is_phase_complete(session=0, phase=0)
clear(session=0, phase=0, clear_up_to_phase=False)

这些 API 允许不经文件直接读取、分 phase 推进和清理 profiling 数据。

5.2 新增 metrics 能力

  • scope metrics 从标量扩展到 sequence 和设备 tensor。
  • 新增 metric 模块及 GPU metric 写入 kernel。
  • LaunchHook.configure(include=..., exclude=...) 可用正则筛选 kernel。
  • launch metadata 从固定 flops/bytes 扩展为任意非保留数值或 tensor metric。

5.3 签名和行为变化

  • start(..., hook=...)
    • hook 可传字符串或 Hook 实例。
    • 未知 hook 从静默忽略改为抛出 ValueError。
    • 返回类型明确为 Optional[int]。
  • deactivate(session=None, flushing=False)
    • 新增 flushing。
    • 当前 docstring 写默认 True，但实现默认是 False，应以签名为准。
  • finalize(..., output_format=...)
    • 新增 hatchet_msgpack。
  • 新增 profiler mode：
    • periodic_flushing
    • 支持 hatchet、hatchet_msgpack、chrome_trace 格式。
  • proton CLI 改用 runpy.run_path(..., run_name="__main__") 执行目标脚本，并正确传播退出码。
  • proton-viewer 对缺失 device_info 更健壮。

5.4 新增配置

  • TRITON_PROFILE_BUFFER_SIZE
  • TRITON_ENABLE_HW_TRACE

Blackwell 硬件 trace 必须在 CUDA context 创建前启动 Proton。

5.5 移除/内部迁移

  • pybind 类型 MetricMap 删除。
  • 废弃的 instrumentation runtime 目录重构到统一 Runtime 层。
  • Proton 自有 GlobalScratchAllocOp 迁移到 TritonGPU 对应能力；直接依赖该 dialect op 的树外代码需要迁移。

6. 原生 C++、MLIR 与插件 API

6.1 树外插件体系

3.7 新增 include/triton/Tools/PluginUtils.h 和 lib/Tools/PluginUtils.cpp，提供 pass、dialect、custom op 的动态插件接口。

关键公开定义：

  • TRITON_PLUGIN_API_VERSION：当前为 2。
  • TRITON_PLUGIN_API：插件入口导出宏。
  • PassInfo
  • DialectInfo
  • OpInfo
  • PluginInfo
  • TritonPlugin
  • 公共入口 tritonGetPluginInfo()

PluginInfo 包含：

  • API version
  • 插件名和插件版本
  • pass/dialect/op 数组及数量
  • 编译所针对的 Triton 版本

运行与构建接口：

TRITON_EXT_ENABLED=1
TRITON_PLUGIN_PATHS=/path/to/plugin1.so:/path/to/plugin2.so
TRITON_PLUGIN_VERSION_CHECK=<policy>

完整示例位于 examples/plugins/，测试位于 test/Plugins/ 和 python/test/unit/plugins/。

6.1.1 示例具体扩展了什么

示例包含两个共享库，分别演示“只增加 pass”和“同时增加 dialect、类型、op、lowering 与前端 op”。

TritonPluginsTestLib：增加一个树外优化 pass

  • examples/plugins/Passes.td 定义 tritongpu-plugin pass，并提供 num-warps 选项。
  • examples/plugins/TritonPlugin.cpp 实现该 pass：遍历模块中的函数， 默认把函数名改为 foo；当 num-warps != 4 时改为 foo_num_warps_<N>。该变换本身只是便于观察插件是否执行的最小示例， 实际插件可在相同位置执行优化、分析、插桩或 kernel 专用变换。
  • 同一文件提供 addTritonPluginPass 和 registerTritonPluginPass： 前者把 pass 实例加入指定 PassManager，后者使该 pass 可被 triton-opt -tritongpu-plugin 按命令行名称调用。

MLIRDialectPlugin：扩展 Triton 可识别和 lower 的 IR

  • DialectPluginDialect.td 注册名为 plugin 的新 MLIR dialect。
  • DialectPluginTypes.td 增加参数化类型 !plugin.custom<"...">，展示树外类型的 parser/printer 和注册方式。
  • DialectPluginOps.td 增加纯操作 plugin.magic %input : i32；它接收并返回 i32。
  • DialectPluginPasses.cpp 实现 convert-plugin-gpu-to-llvm lowering： 将 plugin.magic(x) 改写为 x + gpu.thread_id x。转换目标把整个 plugin dialect 标记为 illegal，因而残留未 lower 的插件 op 会使 pass 失败，而不会悄悄进入后续 LLVM pipeline。
  • DialectPluginDialect.cpp 负责注册 dialect、type 和 op，并额外注册 pass 与一个名为 create_custom_op 的前端 builder callback。
  • python/test/unit/plugins/custom_ops.py 展示如何用 @builtin 包装 builder.create_custom_op(...)，从 Triton Python DSL 调用插件回调； 示例回调生成 arith.addf(src, src)。这条 callback 路径与 plugin.magic dialect op 是两种独立扩展方式：前者直接生成已有 MLIR op，后者引入新的 dialect op 并在稍后 lower。

因此，该示例覆盖了四层扩展：

  1. 在已有 TTIR/TTGIR 上运行自定义 pass。
  2. 引入 Triton 核心源码之外的 MLIR dialect、type 和 op。
  3. 为新 op 提供到 LLVM/NVVM 可接受 IR 的 lowering。
  4. 给 Triton Python 前端增加可在 JIT kernel 中调用的 builtin。

6.1.2 扩展是如何接入 Triton 的

  1. 编译共享库：examples/plugins/CMakeLists.txt 使用 TableGen 生成 pass/dialect/op/type 的 C++ 声明与定义，构建 libTritonPluginsTestLib.so 和 libMLIRDialectPlugin.so，并链接 libtriton。宿主 Triton 必须以 TRITON_EXT_ENABLED=1 构建， 才会导出树外插件所需符号。
  2. 暴露统一入口：每个插件实现 TRITON_PLUGIN_API tritonGetPluginInfo()，返回静态 PluginInfo。 其中列出 API 版本、插件/Triton 版本，以及 pass、dialect 和 op callback。加载器据此完成 ABI 与版本检查。
  3. 运行时发现：进程启动后，PluginUtils 从冒号分隔的 TRITON_PLUGIN_PATHS 加载动态库，调用 tritonGetPluginInfo()， 再把 dialect、pass 和 builder callback 注册到 Triton/MLIR。 仅加载共享库不会自动执行 pass。
  4. 选择插入位置：Python 设置 knobs.runtime.add_stages_inspection_hook，通过修改 stages["ttir"] 或 stages["ttgir"] 包装/替换标准 stage，并在目标位置调用 passes.plugin.add_plugin(pm)。示例既演示在 TTIR stage 末尾追加， 也演示根据标准 make_ttir 源码把 pass 插到 add_loop_unroll 之后，还演示导出并完全替换 TTIR/TTGIR stage。
  5. 保证缓存正确：hook 的零参数调用返回 (key, hash)，使插件代码 或 pipeline 配置成为 JIT cache key 的一部分；五参数调用才实际修改 stages。若忽略前者，插件变化可能错误复用旧二进制。
  6. 调试与验证：pass 可直接由 TRITON_PLUGIN_PATHS=... triton-opt -tritongpu-plugin input.mlir 调试；单元测试则检查函数重命名、plugin.magic 被 lower 为 gpu.thread_id，以及 Python custom op 生成 arith.addf。

当前 dialect lowering 示例依赖 NVIDIA TargetInfo，并固定以 compute capability/PTX 80 构造转换；测试也明确跳过 HIP。这说明统一 插件 ABI 本身支持通用扩展，但示例中的具体 lowering 不是可直接用于 AMD 的生产实现。真实插件还应为目标架构传递正确的 capability、完善 错误处理，并在 TRITON_PLUGIN_API_VERSION 或 Triton 版本变化后重新构建。

新插件系统目前在 NVIDIA/AMD 中是“已预留可用接入点”，但仓库内生产后端没有用它 实现自身功能；实际用例主要是 examples/plugins 和对应测试。示例 lowering 还依赖 NVIDIA TargetInfo，但这只是演示，并非 NVIDIA 后端自身采用了插件系统。

6.2 版本头

新增 include/triton/Version.h.in，构建后提供：

#define TRITON_VERSION "<configured-version>"

插件可据此校验宿主版本。

6.3 布局 API：CTA → CGA

3.7 删除：

  • include/triton/Dialect/TritonGPU/IR/CTAEncodingAttr.h
  • CTAEncodingAttr.td

新增：

  • CGAEncodingAttr.h
  • CGAEncodingAttr.td

主要迁移：

CTAEncodingAttr  -> CGAEncodingAttr
CTALayout        -> CGALayout
getCTALayout()   -> getCGALayout()
cta_encoding     -> cga_encoding

该变化对普通 Python kernel 基本透明，但会破坏：

  • 解析 TTGIR 文本的工具。
  • 直接构造布局属性的 C++/Python 扩展。
  • 树外 TritonGPU pass。
  • 依赖旧 attribute 打印格式的 golden tests。

6.4 新增/改变的 IR 接口

TritonGPU：

  • 新增 ttg.local_gather。
  • 新增 ttg.local_scatter。
  • 新增 ttg.warp_id。
  • ttg.local_barrier 重命名为 ttg.barrier；手写或 override TTGIR 必须同步修改。
  • ttg.async_wait 增加 memory-wait trait，语义与 local barrier 分工更明确。

NVIDIA：

  • 新增 TMA im2col descriptor type。
  • AsyncTMACopyGlobalToLocalOp 支持 tiled/im2col、offset 和 multicast。
  • TCGen5CommitOp 从 two_ctas 迁移到 descriptor 列表。
  • TMEMLoadOp 支持 reduction op、abs 和 NaN 处理。
  • 新增 cluster mbarrier init-release fence op。

AMD：

  • 新增 TDM scatter/gather/prefetch 相关 op。
  • 新增 Warp Pipeline 和 Warp Specialization lowering。
  • multi-CTA/cluster metadata 扩展。

Instrumentation：

  • 新增 BufferRegionAnalysis。
  • buffer pointer 描述迁移为含 length 信息的 buffer descriptor。

7. 编译器和通用优化

7.1 前端

  • 支持 scaled BMM。
  • 支持 FP8 常量。
  • JIT 函数可返回 constexpr。
  • namedtuple 可穿过 constexpr function 并保留字段。
  • 修复字符串 constexpr 被误识别为运行时参数属性的问题：
    • JIT specialization 中每个参数都记录 (类型, 特化值)；动态参数的第二项可能是 D 等对齐/整除属性，而 constexpr 的第二项是编译期常量本身。
    • 当 constexpr 恰好是字符串（例如 "SD"）时，3.6 会把它误当作动态参数的属性， 并尝试设置到 MLIR func 参数上；但 constexpr 已在编译期内联，不对应任何 MLIR 运行时参数，因此可能造成属性与参数错位或编译失败。
    • 3.7 在收集参数属性前过滤 constexpr 项，只对动态参数调用 backend 的 parse_attr。新增测试以 kernel.warmup("SD") 验证字符串 constexpr 可以正常编译。
  • 多项前端性能优化：预计算 builtin signature、避免 inspect.getclosurevars、减少 scope copy 和路径搜索。

7.2 Pass 与分析

  • 新增 BufferRegion 分析，供 ConSan/iisan 使用。
  • ProxyFenceInsertion 开始追踪 async proxy read 依赖；3.7.1 将该正确性修复回合并到发布分支。
  • RemoveLayoutConversions 在 scf.if cleanup 不收敛时改为非致命。
  • AxisInfo 的加减法 divisibility 推导增强。
  • Warp Specialization 分区调度和嵌套循环支持增强。
  • 提取 NVIDIA/AMD 共用 Warp Specialization lowering 工具。

7.3 正确性修复

3.7.1 尤其包含：

  • async read 的 proxy fence。
  • AMD mixed FP8 WMMA/MFMA promotion 和指令选择修复。
  • AMD CanonicalizePointers 对不同 base/offset 的处理。
  • 非 MFMA dot 的 BlockPingpong 修复。
  • GFX950 padded layout pipelining OOM 修复。
  • RangeAnalysis trip count 修复。

8. NVIDIA 后端

8.1 架构与工具链

  • Blackwell ptxas 从 12.9.86 升级到 13.1.80。
  • cuobjdump、nvdisasm、cudacrt、cudart 升级到 13.1.80。
  • LLVM 可接受的 PTX 上限从 min(86, ptx_version) 提升为 min(90, ptx_version)。
  • 为 ptxas 13.x 增加 --regAllocOptLevel=2 workaround。
  • 新增/增强 sm_103（Blackwell Ultra）和 sm_110（Jetson Thor）路径。

8.2 新能力

  • TMA im2col：
    • 新 descriptor type。
    • driver 侧 cuTensorMapEncodeIm2col。
    • async TMA copy 支持 im2col。
  • tcgen05：
    • MMA multicast。
    • M=64 2CTA。
    • tcgen05.ld.red。
    • 更完整的 commit/synchronization 语义。
  • TMA multicast 和 descriptor driver 支持。
  • Variadic 预编译 CUDA launcher。
  • Blackwell 低开销 Proton hardware trace。

8.3 性能与正确性

  • 改进 Warp Specialization partition scheduling。
  • 提升 nested loop 中的 TMEM store。
  • WGMMA pipeline 补充 wait 并推迟不必要的等待。
  • TMA index translation 后移到 lowering。
  • 恢复/调整 membar 的跳过同步分析。
  • FenceAsync 覆盖 async read source。

9. AMD 后端

9.1 gfx1250/RDNA4

3.7 将 gfx1250 从基础支持推进到较完整的计算栈：

  • Warp Specialization。
  • Warp Pipeline。
  • TDM 1D–5D descriptor。
  • TDM scatter/gather/prefetch。
  • multi-CTA WMMA。
  • scaled dot / mixed precision。
  • 8-warp pingpong。
  • Stream-K 示例和内核。
  • cluster barrier。

内部测试辅助命名从 gfx11/gfx12 改为 RDNA3/RDNA4，但目标 arch 字符串没有因此改名。

9.2 编译管线重构

新增或接入：

  • SinkLayoutConversions
  • PrepareIfCombining
  • MoveUpPrologueLoads
  • ConvertToTensorOps
  • WarpPipeline
  • AllocateWarpGroups
  • WarpPipelineConversion
  • WarpSpecializeToLLVM

替换或移除：

  • ReorderInstructions 被 MoveUpPrologueLoads 取代。
  • OptimizeLDSUsage 从管线移除，由新的 pipeline/layout 策略替代。

其他行为：

  • TDM 架构跳过 descriptor-to-pointer rewrite。
  • Warp Specialization 下 metadata 的 warp 数可来自 ttg.total-num-warps。
  • gfx1250 或 MIR dump/swap 模式下禁用部分 inreg 设置。
  • gfx11 汇编显式禁用 True16。

9.3 调试与性能

  • 新增 TRITON_DUMP_MIR / TRITON_SWAP_MIR。
  • MIR swap pipeline 增加 NoAlias AA wrapper。
  • v_perm 加速 convert-layout。
  • contiguity 信息用于 buffer/direct-to-LDS。
  • mixed FP8、RangeAnalysis、CanonicalizePointers 等获得发布分支修复。

10. triton_kernels 独立包

python/triton_kernels 是仓库内的独立/附赠 Python 包，不是 triton.language 稳定 API，但 3.7 的改动非常大。

10.1 Matmul 破坏性重构

主要重命名：

triton_kernels.matmul_ogs             -> triton_kernels.matmul
matmul_ogs_set_idle_sms               -> matmul_set_idle_sms
matmul_ogs(...)                       -> matmul(...)
matmul_ogs_torch(...)                 -> matmul_torch(...)
matmul_ogs_details/                   -> matmul_details/
_matmul_ogs / _p_matmul_ogs           -> _matmul / _p_matmul
initialize_matmul_ogs                 -> initialize_matmul

这不只是重命名：

  • 主输入 x, w 改为 a, b。
  • RoutingData/InnerRoutingData 迁移为 ragged metadata。
  • GatherIndx/ScatterIndx dataclass 改为 tensor index。
  • act_scale/weight_scale/out_scale 改为 a_mx_scale/b_mx_scale/c_mx_scale。
  • 输出和累加参数从 y 系列改为 c 系列。

删除或不再直接提供的类型包括 GatherIndx、ScatterIndx、RoutingData、InnerRoutingData、Bitmatrix 和旧 tensor 模块中的 bitwidth/get_layout。

10.2 Tensor 与 layout API

  • Tensor(storage=...) 现在要求 Storage，不再自动包装裸 torch.Tensor。
  • dtype 迁移到独立的 IntegerType/FloatType 体系。
  • wrap_torch_tensor 新增 shape、shape_max、layout。
  • convert_layout 从接收 layout class 改为接收 layout instance。
  • 默认 matmul layout 工厂从返回 (LayoutClass, kwargs) 改为直接返回 Layout 实例。
  • 新增 empty、is_tma_compliant、make_dense_tma、make_tma 和 dtype/torch dtype 转换函数。
  • 新增 Blackwell activation MX scale layout 和多种 LayoutTransformation。

10.3 分布式 API

  • 新增 Mesh(process_group)。
  • SymmetricMemoryPool 从无参/全局单例改为显式接收 Mesh。
  • 删除模块级 distributed.symm_mem_pool。
  • convert_dp_to_ep、convert_ep_to_dp 需要显式传入 pool。
  • topk_forward(..., symm_mem_pool=None) 新增 pool；all_gather=True 时必须提供。

10.4 Reduce、numerics 和 roofline

  • reduce 拆分为 reduce_forward / reduce_backward，新增 autograd 支持。
  • reduce 新增 global scale 和 unpadded batch size 等能力。
  • downcast_to_mxfp 参数名和 dtype 体系调整。
  • get_cublas_tflops(dtype) → get_blas_tflops(dtype, workspace_size=..., device="cuda")。
  • 新增 get_rdna_version()。

11. 构建、打包和依赖

11.1 LLVM 和 CUDA 工具链

  • LLVM：
    • 3.6：f6ded0be897e2878612dd903f7e8bb85448269e5
    • 3.7.1：1f126a6dea50d185c0781743a667390037ae88bd
  • LLVM 预编译包名新增 build number 后缀。
  • Blackwell CUDA 工具链升级到 13.1.80。

11.2 构建开关

  • 新增 TRITON_EXT_ENABLED，用于导出树外插件需要的符号。
  • CMake 新增 TRITON_OFFLINE_BUILD option；setup.py 在 3.6 已支持同名环境变量，3.7 将离线构建控制补齐到 CMake 层。
  • CMake 接收 TRITON_VERSION 并生成版本头。
  • 插件构建时调整 libtriton 的符号可见性和接口链接属性。

11.3 Python 与 wheel

  • 支持的 Python 范围不变：3.10–3.14。
  • 主要 Python build dependencies 未变。
  • 3.7.1 source distribution 新增 examples/，从而包含插件示例。
  • wheel 构建启用扩展支持并减小产物体积。
  • CI pin pandas < 3.0。

12. 文档、示例和测试

新增或显著扩展：

  • examples/plugins/：完整 pass、dialect、custom op 插件示例。
  • Gluon 教程：
    • TMA gather/scatter。
    • tcgen05 copy。
    • tcgen05 scaled MMA。
  • AMD gfx1250 示例：
    • Warp Pipeline。
    • Stream-K。
    • multi-CTA GEMM。
  • docs/python-api/triton.language.rst：补充 topk。
  • test/Plugins/、python/test/unit/plugins/：插件测试。
  • test/Analysis/test-buffer-region.mlir：BufferRegion。
  • NVIDIA/AMD Warp Specialization、TMA/TDM、cluster、mixed FP8 和 fence 测试。

13. 从 3.6 升级到 3.7 的建议清单

普通 Triton Python 用户

  • 搜索 make_block_ptr，规划迁移到 Tensor Descriptor。
  • 检查 tl.cat 测试，尤其是过去依赖未定义/可重排顺序的代码。
  • 精度敏感的 TF32 路径考虑 round_f32_to_tf32=True。
  • 若使用 preload，确保 specialization 与当前 backend target 一致。
  • 升级后删除或隔离旧 JIT cache。

Gluon 用户

  • thread_barrier 改为 barrier。
  • 检查 barrier(cluster=True)、cluster arrive/wait 语义。
  • 检查直接构造的 AMD/NVIDIA experimental layout。
  • 对新 verifier 重新运行全部 kernel 编译测试。

AOT 用户

  • 使用 3.7 的 compile/link 工具全量重建。
  • 不复用 3.6 的头、C 文件、对象或共享库。
  • 若支持 HIP，改用 TT_ResultTy/TT_StreamTy 抽象。

插件、后端和 IR 工具作者

  • CTA/CTALayout API 迁移到 CGA/CGALayout。
  • 手写 TTGIR 中的 ttg.local_barrier 改为 ttg.barrier。
  • 自定义 add_stages_inspection_hook 同时兼容零参数 cache-key 调用和五参数 stage-inspection 调用。
  • 使用 TRITON_PLUGIN_API_VERSION == 2 的 PluginInfo。
  • 构建宿主时启用 TRITON_EXT_ENABLED=1。
  • 设置 TRITON_PLUGIN_PATHS 并处理版本校验。
  • 用新 LLVM revision 重建全部 native 扩展。

AMD 用户

  • 对 gfx1250/TDM/WS 内核做性能和数值回归。
  • 不假设 async copy 默认开启；需要时显式设置。
  • 若使用 MIR 调试，检查 TRITON_DUMP_MIR/TRITON_SWAP_MIR。

NVIDIA/Blackwell 用户

  • 验证 ptxas 13.1 对现有 kernel 的寄存器分配和性能影响。
  • 检查 TMA descriptor 所需 driver 能力。
  • 对 tcgen05 multicast、2CTA、ld.red 路径做架构定向测试。

Proton 用户

  • 评估迁移到 triton.profiler.data.* 的内存内/phase API。
  • 若用 deactivate(flushing=...)，显式传值，不依赖有冲突的 docstring。
  • Blackwell hardware trace 在创建 CUDA context 前启动。

triton_kernels 用户

  • 按新 matmul API 重写 import、路由和 scale 参数。
  • 迁移 Tensor、layout instance 和 dtype API。
  • 使用显式 Mesh/SymmetricMemoryPool。
  • 检查 reduce autograd 和 MX scale 路径。

14. 关键证据索引

  • 版本：python/triton/__init__.py、setup.py
  • Python 语言 API：python/triton/language/
  • JIT/runtime：python/triton/runtime/jit.py、_async_compile.py、cache.py、driver.py
  • 工具：python/triton/tools/
  • 配置：python/triton/knobs.py
  • Gluon：python/triton/experimental/gluon/
  • Proton：third_party/proton/proton/、third_party/proton/csrc/
  • 插件：include/triton/Tools/PluginUtils.h、lib/Tools/PluginUtils.cpp、examples/plugins/
  • Triton/TritonGPU IR：include/triton/Dialect/
  • NVIDIA：third_party/nvidia/
  • AMD：third_party/amd/
  • triton_kernels：python/triton_kernels/
  • LLVM revision：cmake/llvm-hash.txt
  • NVIDIA 工具链：cmake/nvidia-toolchain-version.json

Clone this wiki locally