Triton-MUSA

Triton-MUSA 是一套面向 MUSA 后端适配的 Triton 语言与编译器栈，用于使用 Python 编写并运行面向 MUSA 设备的 GPU Kernel。当前 MUSA 已完成 Triton 3.1、Triton 3.2 和 Triton 3.6 三个版本的适配，其中 Triton 3.6 适配已在 MUSA SDK 5.2.0 中正式推出。

概述

什么是 Triton-MUSA

Triton-MUSA（MUSA 后端适配的 Triton 语言与编译器栈）是 Triton 语言和编译器栈的 MUSA 后端适配，允许用户使用 Triton DSL 以 Python 方式编写 GPU Kernel，并将其编译为能够在 MUSA 设备上运行的后端产物。它是一个面向 MUSA 硬件的、以 Python 为中心的 Kernel 编程环境和编译工具链。

现代工作负载经常需要自定义 Kernel 来实现算子融合、带宽优化、特殊布局支持或硬件特性利用。Triton-MUSA 缩小了开发效率与运行性能之间的差距：用户可以使用 Triton 风格的 Python Kernel 表达计算逻辑，同时由 MUSA 后端完成目标相关的编译、运行和调优支持。

当前 Triton-MUSA 的版本适配可以理解为三个阶段：

适配版本	发布状态	后端形态	主要定位
Triton 3.1	已适配	MTGPU 后端	支持 MUSA 设备上的基础 Triton Kernel 编写、编译、运行、benchmark 和验证流程
Triton 3.2	已适配	MTGPU 后端	在 Triton 3.1 适配基础上增强 MUSA-only 构建、AOT、benchmark 接口、WMMA/SQMMA/TME 测试和后端稳定性
Triton 3.6	已随 MUSA SDK 5.2.0 推出	MUSA 后端	面向 Triton 3.6 的 MUSA 后端重构，补充 `TensorDescriptor`、target info、descriptor/TME、WMMA/SQMMA、AOT/JIT、更多 dtype 和更完整测试覆盖

本文档中已有的基础使用方式、安装验证、最小 Kernel、benchmark 和调试排查内容，可以视为 Triton 3.1 与 Triton 3.2 适配现状。后续涉及 TensorDescriptor、TME、WMMA/SQMMA、MUSA-only 构建、更多 dtype 和更完整测试覆盖的内容，主要对应 Triton 3.6 适配方向。

当前 Triton-MUSA 3.1、3.2 和 3.6 适配版本都会调用 MUSA SDK 中的后端编译器完成目标代码生成，因此对应 Python 发布包需要与 MUSA SDK 版本匹配使用。MUSA SDK 5.3.0 起，Triton-MUSA 计划不再与 MUSA SDK 版本强绑定，届时请以对应版本发布说明为准。

版本适配状态

对比项	Triton 3.1 适配	Triton 3.2 适配	Triton 3.6 适配
版本基础	基于 Triton 3.1.0，Wheel 版本带 MUSA 后缀	基于 Triton 3.2.0，Wheel 版本带 MUSA 后缀	基于 Triton 3.6.0，面向 MUSA SDK 5.2.0 发布
SDK 依赖	需要与对应 MUSA SDK 版本匹配	需要与对应 MUSA SDK 版本匹配	需要与 MUSA SDK 5.2.0 版本组合匹配；MUSA SDK 5.3.0 起计划解除 SDK 版本强绑定
构建与发布	完成 MUSA 后端基础接入	强化 MUSA-only 构建路径，减少非 MUSA 工具链依赖	面向 Triton 3.6 工程结构升级，强化 MUSA-only 构建和发布流程
后端组织	以 MTGPU 后端形态承载 MUSA 编译、运行和测试能力	延续 MTGPU 后端形态，同时增强编译选项、运行时 benchmark、AOT 和缓存稳定性	升级为更清晰的 MUSA 后端形态，系统性增强编译、运行和测试覆盖
支持架构	支持 PH1（MP_31）和 QY2（MP_22）	支持 PH1（MP_31）和 QY2（MP_22）	支持 PH1（MP_31）和 QY2（MP_22），其中 QY2 支持相对 PH1 较弱
编程模型	支持 `@triton.jit`、`tl.load`、`tl.store`、`tl.arange`、`tl.dot`、autotune 等基础 Triton 写法	在 3.1 基础上补充更多语言语义、模块映射、类型转换和矩阵路径修复	增加 `tl.target_info.is_musa()`、`tl.target_info.musa_capability_geq()`、`TensorDescriptor.from_tensor`、`tl.load_tensor_descriptor` / `tl.store_tensor_descriptor` 等接口
矩阵与数据搬运	已具备 WMMA、SQMMA、TME 相关基础验证；其中 WMMA 需要通过 `ENABLE_MUSA_MMA=1` 显式开启，SQMMA 需要通过 `MUSA_ENABLE_SQMMA=1` 显式开启	扩展 PH1 WMMA、SQMMA K16、TME load/store、inplace、transpose、FP8/BF16/TF32 等场景，满足条件时由后端自动选择矩阵加速路径	对 descriptor/TME、WMMA/SQMMA、rank-3 dot、int8、FP8/BF16/TF32、转置 operand、accumulator 布局等路径做更完整覆盖，满足条件时由后端自动选择矩阵加速路径
AOT/JIT	以 JIT 为主，具备 MUSA 后端编译与运行基础	增加面向 `musa` target 的 AOT 编译测试和工具链验证	AOT 与 JIT 适配到 Triton 3.6 工具链，支持显式 MUSA target 编译、运行时 target 检测和后端产物生成
测试覆盖	覆盖基础语言、runtime、operator、SQMMA/TME 和 WMMA 场景	新增 AOT、PH1 WMMA、更多 SQMMA benchmark、inplace 和回归测试	补充 target info、TensorDescriptor、dot path coverage、MUSA 后端工具链和 MLIR 级测试

关键特性

Triton-MUSA 提供以下能力：

基于 @triton.jit 的 Python Kernel 编程方式
MUSA 后端支持，运行时可自动检测当前可用的 MUSA target
从 Triton Kernel 到后端二进制产物的 JIT 编译流程
对 tl.load、tl.store、tl.arange、tl.dot 以及 launch grid 等标准 Triton 语言构造的支持
通过 triton.autotune、triton.Config、num_warps 和 num_stages 进行性能调优
通过 triton.musa_testing.do_bench 进行 MUSA 性能测试
面向 MUSA 的 AOT 编译路径，可生成 C 启动代码和嵌入式 mubin
面向 MUSA 架构特性的矩阵加速和数据搬运路径，包括 TME、WMMA、WGMMA 风格的 warp-group 矩阵计算模型以及 SQMMA
Triton 3.6 适配中补充 MUSA-only 构建、TensorDescriptor、MUSA target info、更多矩阵 Kernel 验证和更完整的后端测试覆盖

适用读者

本文档适用于：

希望在 MUSA 设备上编写和运行 Triton Kernel 的用户
需要调试、压测或集成 MUSA 定向 Kernel 的开发者

工作原理

工作流程

主要组成部分

前端用户使用 Triton 装饰器和语言 API 在 Python 中编写 Kernel。
编译器 Triton 将 Kernel Lower 到多个中间表示，执行优化 Pass，最终为 MUSA 生成后端代码。
运行时 Runtime 负责选择当前激活的后端驱动并启动已编译的 Kernel。
后端 MUSA 后端提供目标相关的代码生成、启动桩代码以及 AOT 集成能力。

驱动选择

通常情况下，用户不需要手动选择后端驱动。运行时会根据当前环境检测可用的 MUSA target；用户只需要确认 MUSA Runtime、Driver 和 torch_musa 环境可用。

架构

从用户视角看，整个栈可以理解为：

使用 Python 编写 Kernel
通过 JIT 或 AOT 完成编译
在 device="musa" 上分配 Tensor 并发起运行
与 PyTorch 参考实现进行正确性校验
对性能进行测试和调优

环境与前置条件

在使用 Triton-MUSA 之前，需要准备以下环境：

Linux 环境
Python 3.10
正确安装并可用的 MUSA Runtime 和 Driver
torch_musa
与当前 Python 版本、操作系统和硬件环境匹配的 Triton-MUSA Python 包
与 Triton-MUSA Python 包匹配的 MUSA SDK 版本

Triton-MUSA 3.1、3.2 和 3.6 适配版本会调用 MUSA SDK 中的后端编译器，例如 /usr/local/musa/bin/llc，因此不同于部分 NV 后端适配方式，当前 Triton-MUSA 版本需要与 MUSA SDK 版本匹配。如果使用 MUSA SDK 5.2.0 中推出的 Triton 3.6 适配版本，请以 MUSA SDK 5.2.0 随包说明为准确认 Triton、torch_musa、MUSA Runtime 和 Driver 的版本组合。MUSA SDK 5.3.0 起，Triton-MUSA 计划不再与 MUSA SDK 版本强绑定。

推荐先执行以下检查：

python - <<'PY'
import torch
import torch_musa
import triton
print("triton version:", triton.__version__)
print("has torch.musa:", hasattr(torch, "musa"))
print("musa available:", hasattr(torch, "musa") and torch.musa.is_available())
PY

如果还想进一步查看当前 Triton 目标信息：

python - <<'PY'
import torch_musa
from triton.runtime.driver import driver
target = driver.active.get_current_target()
print("backend:", target.backend)
print("arch:", target.arch)
print("warp_size:", target.warp_size)
PY

快速开始

配置 pip 软件源

请根据使用场景，选择以下任一种方式配置 pip 软件源：

为当前 shell 会话或 CI 任务配置软件源：

export PIP_INDEX_URL=https://dl.mthreads.com/repo/api/pypi/pypi/simple

将软件源写入 pip 配置文件，并查看当前配置：

python -m pip config set global.index-url \
  https://dl.mthreads.com/repo/api/pypi/pypi/simple
python -m pip config list

不修改 pip 配置，在安装时指定软件源：

python -m pip uninstall -y triton
python -m pip install triton \
  --index-url https://dl.mthreads.com/repo/api/pypi/pypi/simple

备注

请勿在同一次安装中同时使用摩尔线程 Python 包源和公开 PyPI 镜像源。安装普通第三方包时，请切换软件源后单独安装。

安装 Python 包

使用环境变量或 pip 配置文件完成配置后，可通过 pip 直接安装 Triton-MUSA：

python -m pip uninstall -y triton
python -m pip install triton

triton_musa 仓库发布到 PyPI 后的包名是 triton。

如需指定 triton 版本，可选择：

python -m pip install triton==3.2.0
python -m pip install triton==3.6.0

安装完成后仍通过 import triton 使用。

首次功能验证

安装完成后，先验证 import 与 target 检测是否正常。

如需先确认当前安装的分发包信息，可执行：

python -m pip show triton

python - <<'PY'
import triton
print("triton:", triton.__version__)
PY

python - <<'PY'
import triton
import torch
import torch_musa
from triton.runtime.driver import driver

print("triton:", triton.__version__)
print("musa available:", torch.musa.is_available())
print("target:", driver.active.get_current_target())
PY

主要操作与 API

面向用户最常用的 API 包括：

@triton.jit 将 Python 函数标记为 Triton Kernel。
triton.cdiv 用于 launch grid 计算的向上取整除法辅助函数。
triton.autotune 与 triton.Config 用于配置搜索与运行时调优。
triton.testing.do_bench 通用性能测试辅助函数。
triton.musa_testing.do_bench Triton 3.1 / 3.2 适配中常用的 MUSA 专用性能测试辅助函数。
triton.compile 编程方式调用的编译入口。
python -m triton.tools.compile 支持 musa 目标的 AOT 编译工具。
tl.target_info.is_musa Triton 3.6 适配中可用于在 JIT Kernel 内判断当前 target 是否为 MUSA。
tl.target_info.musa_capability_geq Triton 3.6 适配中可用于在 JIT Kernel 内按 MUSA capability 选择代码路径。
triton.tools.tensor_descriptor.TensorDescriptor Triton 3.6 适配中用于创建 descriptor-based load/store 所需的张量描述对象。
tl.load_tensor_descriptor / tl.store_tensor_descriptor Triton 3.6 适配中用于通过 descriptor 读写张量 tile。
triton.tools.experimental_descriptor.create_1d_tma_descriptor 在部分既有适配版本中用于创建一维 TME/TMA descriptor。
triton.tools.experimental_descriptor.create_2d_tma_descriptor 在部分既有适配版本中用于创建二维 TME/TMA descriptor，常用于 GEMM 的矩阵 tile 搬运。
triton.tools.experimental_descriptor.create_3d_tma_descriptor 在部分既有适配版本中用于创建三维 TME/TMA descriptor。

常用 Triton 语言 API 包括：

tl.program_id
tl.arange
tl.load
tl.store
tl.dot
tl.zeros
tl.constexpr
tl.load_tensor_descriptor
tl.store_tensor_descriptor

备注

不同 Triton-MUSA 适配版本的高级 descriptor API 可能存在差异。Triton 3.1 / 3.2 适配以随包示例和既有接口为准；Triton 3.6 适配中推荐关注 TensorDescriptor.from_tensor、tl.load_tensor_descriptor 和 tl.store_tensor_descriptor 这一类接口。

示例代码

最小 Vector Add 示例

下面是一个可直接保存并运行的最小示例。请将示例保存为 .py 文件后执行，@triton.jit 需要读取 Kernel 函数源码，不建议从标准输入或临时交互片段中定义 Kernel。

import torch
import torch_musa
import triton
import triton.language as tl


@triton.jit
def add_kernel(x_ptr, y_ptr, output_ptr, n_elements, BLOCK_SIZE: tl.constexpr):
    pid = tl.program_id(axis=0)
    block_start = pid * BLOCK_SIZE
    offsets = block_start + tl.arange(0, BLOCK_SIZE)
    mask = offsets < n_elements
    x = tl.load(x_ptr + offsets, mask=mask)
    y = tl.load(y_ptr + offsets, mask=mask)
    tl.store(output_ptr + offsets, x + y, mask=mask)


def add(x: torch.Tensor, y: torch.Tensor):
    output = torch.empty_like(x)
    n_elements = output.numel()
    grid = lambda meta: (triton.cdiv(n_elements, meta["BLOCK_SIZE"]),)
    add_kernel[grid](x, y, output, n_elements, BLOCK_SIZE=1024)
    return output


x = torch.rand(98432, device="musa", dtype=torch.float32)
y = torch.rand(98432, device="musa", dtype=torch.float32)
out = add(x, y)
ref = x + y
torch.testing.assert_close(out, ref)
print("vector add passed")

Benchmark 示例

ms = triton.musa_testing.do_bench(lambda: add(x, y))
print(f"triton add latency: {ms:.4f} ms")

在 Triton 3.6 适配版本中，也可以使用 triton.testing.do_bench 并显式指定 MUSA 设备类型：

ms = triton.testing.do_bench(lambda: add(x, y), device_type="musa")
print(f"triton add latency: {ms:.4f} ms")

面向 MUSA 的高级示例入口

如果发布包或配套材料中随附高级示例，可以优先参考其中的 descriptor-based load/store、GEMM benchmark、WMMA、SQMMA 和 AOT 代码生成示例。Triton 3.6 适配中已覆盖 MM、BMM、AddMM、TME load/store、WMMA/SQMMA、inplace、target info 等测试和示例方向。未随包提供示例时，可先基于本文档的最小 Vector Add 示例确认安装和运行环境。

测试与验证

用户侧推荐以安装验证、最小 Kernel 正确性验证和性能基准验证为主。

安装验证

python - <<'PY'
import triton
import torch
import torch_musa
from triton.runtime.driver import driver

print("triton:", triton.__version__)
print("musa available:", torch.musa.is_available())
print("target:", driver.active.get_current_target())
PY

Kernel 正确性验证

运行本文档中的最小 Vector Add 示例，并确认输出包含：

vector add passed

性能验证

在正确性验证通过后，再运行 Benchmark 示例，确认 triton.musa_testing.do_bench 能返回稳定的耗时数据。

验证检查项：

Import 成功
torch.musa.is_available() 返回 True
当前激活目标后端为 musa
Kernel 输出与 PyTorch 参考实现一致
Benchmark 可以正常运行
如果使用 Triton 3.6 适配版本，可结合随包测试或示例进一步验证 descriptor/TME、MM、BMM、AddMM、WMMA/SQMMA、FP8/BF16/TF32/int8 等场景

调试与故障排查

常用环境变量

这些开关只在需要时启用：

# 始终强制重新编译 Kernel
TRITON_ALWAYS_COMPILE=1

# 打印 autotune 过程信息
TRITON_PRINT_AUTOTUNING=1

# 查看前端解释执行行为，适合定位基础语义问题
TRITON_INTERPRET=1

# 统计编译阶段耗时
MLIR_ENABLE_TIMING=1
LLVM_ENABLE_TIMING=1

# Triton 3.1 适配版本中显式开启 WMMA / SQMMA 矩阵加速路径
ENABLE_MUSA_MMA=1
MUSA_ENABLE_SQMMA=1

常见问题

现象	可能原因	建议处理方式
`torch.musa.is_available()` 为 `False`	Runtime 或 Driver 未准备好	检查 MUSA 安装和 `torch_musa` 环境
当前后端不是 `musa`	Runtime 选择了其他后端	检查环境并打印 `driver.active.get_current_target()`
安装 Python 包失败	当前 Python 版本、系统平台或安装配置不匹配	确认当前环境与发布包要求匹配，必要时重新检查安装指南中的 Python 包安装说明
`import triton` 后版本不符合预期	环境中已有其他 `triton` 包	使用 `python -m pip show triton` 确认安装位置，必要时卸载 `triton` 后重新安装
上游 Triton 示例里仍然使用 `cuda`	上游示例面向 CUDA 后端	将示例中的 device 字符串和环境依赖改成 MUSA 对应写法
Triton 3.6 示例中的 descriptor API 不存在	当前安装的是 Triton 3.1 / 3.2 适配包	以当前发布包随附示例为准，或切换到 MUSA SDK 5.2.0 中的 Triton 3.6 适配版本

建议的调试流程

推荐顺序：

先验证设备是否可用
再运行一个最小 Kernel
与 PyTorch 参考结果做对比
在正确版本上进行性能测试
如果使用 Triton 3.6 适配版本，再逐步验证 descriptor/TME、矩阵 Kernel 和目标 dtype
如问题仍然存在，记录 Triton-MUSA 发布包版本、MUSA Driver/Runtime 版本、torch_musa 版本和最小复现代码后反馈给支持人员

高级主题

三个适配版本的演进关系

Triton 3.1、Triton 3.2 和 Triton 3.6 的 MUSA 适配是逐步演进关系：

Triton 3.1 适配完成了 MUSA 后端的基础闭环，包括 Python JIT Kernel、MUSA runtime/driver 接入、基础语言算子、tl.dot 矩阵路径、benchmark 辅助接口以及基础验证。
Triton 3.2 适配在 3.1 基础上提升工程化和可用性，包括 MUSA-only 构建、AOT 编译、autotune benchmark 入口、PH1 WMMA、更多 SQMMA/TME 矩阵测试、inplace load 以及回归测试。
Triton 3.6 适配同步到 Triton 3.6 工程结构，并升级为更清晰的 MUSA 后端形态，重点增强 TensorDescriptor、target info、descriptor/TME、WMMA/SQMMA、AOT/JIT、dtype 覆盖和 MLIR 级测试。

从用户视角看，Triton 3.1 / 3.2 更适合理解当前已发布版本的基础能力；Triton 3.6 则是面向 MUSA SDK 5.2.0 的新一代适配版本，推荐新项目优先关注其发布包和随包示例。

Triton 3.6 MUSA 适配能力

Triton 3.6 MUSA 适配已在 MUSA SDK 5.2.0 中正式推出。相较 Triton 3.1 / 3.2 适配版本，Triton 3.6 适配重点补充以下能力：

MUSA-only 构建与发布流程：支持只构建 MUSA backend，减少与 CUDA / AMD 后端工具链的耦合。
更完整的 MUSA runtime 集成：通过 torch_musa 检测当前 MUSA 设备、stream、target arch 和 warp size。
后端编译链路升级：支持从 Triton Kernel 经 TTIR、TTGIR、LLVM IR 到 MUSA 后端产物的 JIT 编译流程，并系统性增强 MUSA 后端编译与运行能力。
AOT 编译支持：支持通过 Triton AOT 工具生成面向 musa target 的启动代码和后端产物。
MUSA target info：支持在 JIT Kernel 中使用 tl.target_info.is_musa() 和 tl.target_info.musa_capability_geq() 做 target 判断。
Descriptor/TME 路径：支持 TensorDescriptor.from_tensor、tl.load_tensor_descriptor 和 tl.store_tensor_descriptor 等 descriptor-based load/store。
矩阵 Kernel 路径：围绕 tl.dot 补充 WMMA、SQMMA、TME 与矩阵 tile 相关 lowering 和优化。
数据类型覆盖：补充 FP16、BF16、FP8、TF32、int8 等常见矩阵计算场景的编译和验证能力。
测试覆盖增强：新增 MM、BMM、AddMM、TME load/store、WMMA/SQMMA、inplace load、target info、dot path coverage、TensorDescriptor、MUSA 后端工具链和 MLIR 级测试方向。

AOT 编译

如果安装包包含 triton.tools.compile 模块，可以通过 python -m triton.tools.compile 使用面向 musa 的 AOT 编译能力。它可以读取用户自己的 Triton Kernel Python 文件，并生成自包含的 C 启动代码以及嵌入式后端二进制数据。

用于 AOT 编译的 Python 文件建议只包含 Kernel 定义和必要的 import，避免在文件顶层执行设备分配或 Kernel launch。--signature 需要按照 Kernel 参数顺序填写；如果 Kernel 中包含 tl.constexpr meta 参数，也需要在对应位置提供常量值。

典型流程：

查询当前激活目标：

python - <<'PY'
import torch_musa
from triton.runtime.driver import driver
t = driver.active.get_current_target()
print(t.backend, t.arch, t.warp_size)
PY

使用显式目标参数编译 Kernel：

python -m triton.tools.compile \
  --kernel-name add_kernel \
  --signature "*fp32,*fp32,*fp32,i32,1024" \
  --grid 1,1,1 \
  --target-backend musa \
  --target-arch <arch> \
  --target-warp-size <warp_size> \
  path/to/kernel.py

MUSA 架构特性概览

Triton-MUSA 的用户层编程模型尽量保持 Triton 风格：普通访存使用 tl.load / tl.store，矩阵计算使用 tl.dot，性能调优通过 block size、num_warps、num_stages 和 autotune 配置完成。面向具体 MUSA 架构的 TME、WMMA、WGMMA 风格矩阵计算模型和 SQMMA 由后端在编译阶段选择和降级，用户通常不需要直接写底层指令。

当前 Triton 3.1、Triton 3.2 和 Triton 3.6 适配均覆盖 PH1（MP_31）和 QY2（MP_22）两类 MUSA 架构。需要注意的是，Triton 3.6 适配中 PH1 相关的 descriptor/TME、WMMA/SQMMA、矩阵 Kernel 和测试覆盖更完整；QY2 仍可作为支持架构使用，但高级矩阵加速与验证覆盖相对 PH1 较弱，建议以当前发布包说明和随包示例确认具体可用能力。

这些能力是否生效取决于当前硬件架构、数据类型、矩阵 tile shape、访存布局以及发布包版本。建议先通过下面的方式确认当前目标架构：

python - <<'PY'
import torch_musa
from triton.runtime.driver import driver
t = driver.active.get_current_target()
print("backend:", t.backend)
print("arch:", t.arch)
print("warp_size:", t.warp_size)
PY

TME 与 Descriptor Load/Store

TME（Tensor Memory Engine）用于以 descriptor 描述张量 tile 的形状、跨度、元素大小和基地址，再由后端选择更适合矩阵 tile 的数据搬运路径。Triton-MUSA 中，descriptor-based 工作流通常由两部分组成：

Host 侧创建 descriptor：Triton 3.6 适配中推荐使用 TensorDescriptor.from_tensor；Triton 3.1 / 3.2 适配中请以随包示例提供的 descriptor 接口为准。
Kernel 侧读写 descriptor：Triton 3.6 适配中使用 tl.load_tensor_descriptor / tl.store_tensor_descriptor 读写 descriptor 描述的 tile。

这种方式常用于 GEMM、batched GEMM、带特殊布局的矩阵读写，尤其适合与 tl.dot 组成矩阵 Kernel。descriptor 接口属于高级接口，不同 Triton-MUSA 适配版本可能存在 API 差异；生产代码建议以当前发布说明和随包示例为准。

WMMA、WGMMA 与 SQMMA

WMMA、WGMMA 和 SQMMA 都属于矩阵乘加相关的硬件加速路径，但关注层级不同：

WMMA 通常表示 warp 级矩阵计算路径，适合较传统的矩阵 tile 计算。
WGMMA 表示 warp-group 级矩阵计算模型，常见于上游 Triton/Hopper 相关材料。对 Triton-MUSA 用户来说，它更适合作为理解“多 warp 协作完成矩阵乘加”的概念；实际 MUSA 后端会根据目标架构选择对应的矩阵加速 lowering。
SQMMA 是 MUSA/MTGPU 后端面向特定架构的矩阵加速路径。对满足条件的 tl.dot，后端可以在编译阶段选择 SQMMA 相关实现，从而利用硬件矩阵乘加能力。

实际开发时，推荐从标准 tl.dot 写法开始，再结合 TME descriptor、合适的 block size 和 autotune 搜索优化性能。不是所有 dtype、tile shape 和 num_warps 组合都会触发同一种矩阵加速路径，具体行为以当前 Triton-MUSA 发布包、目标硬件架构和随包示例为准。

Triton-MUSA

概述

什么是 Triton-MUSA

版本适配状态

关键特性

适用读者

工作原理

工作流程

主要组成部分

驱动选择

架构

环境与前置条件

快速开始

配置 pip 软件源

安装 Python 包

首次功能验证

主要操作与 API

示例代码

最小 Vector Add 示例

Benchmark 示例

面向 MUSA 的高级示例入口

测试与验证

安装验证

Kernel 正确性验证

性能验证

调试与故障排查

常用环境变量

常见问题

建议的调试流程

高级主题

三个适配版本的演进关系

Triton 3.6 MUSA 适配能力

AOT 编译

MUSA 架构特性概览

TME 与 Descriptor Load/Store

WMMA、WGMMA 与 SQMMA

相关文档

这篇文章对您有帮助吗？

概述​

什么是 Triton-MUSA​

版本适配状态​

关键特性​

适用读者​

工作原理​

工作流程​

主要组成部分​

驱动选择​

架构​

环境与前置条件​

快速开始​

配置 pip 软件源​

安装 Python 包​

首次功能验证​

主要操作与 API​

示例代码​

最小 Vector Add 示例​

Benchmark 示例​

面向 MUSA 的高级示例入口​

测试与验证​

安装验证​

Kernel 正确性验证​

性能验证​

调试与故障排查​

常用环境变量​

常见问题​

建议的调试流程​

高级主题​

三个适配版本的演进关系​

Triton 3.6 MUSA 适配能力​

AOT 编译​

MUSA 架构特性概览​

TME 与 Descriptor Load/Store​

WMMA、WGMMA 与 SQMMA​

相关文档​

这篇文章对您有帮助吗？

概述

什么是 Triton-MUSA

版本适配状态

关键特性

适用读者

工作原理

工作流程

主要组成部分

驱动选择

架构

环境与前置条件

快速开始

配置 pip 软件源

安装 Python 包

首次功能验证

主要操作与 API

示例代码

最小 Vector Add 示例

Benchmark 示例

面向 MUSA 的高级示例入口

测试与验证

安装验证

Kernel 正确性验证

性能验证

调试与故障排查

常用环境变量

常见问题

建议的调试流程

高级主题

三个适配版本的演进关系

Triton 3.6 MUSA 适配能力

AOT 编译

MUSA 架构特性概览

TME 与 Descriptor Load/Store

WMMA、WGMMA 与 SQMMA

相关文档