目录
Kafka-Hatsune

[Triton] Fea/warp layout broadcast from upstream triton (#1150)

  • [Analysis] Backport layout prerequisites for Triton #11646

Add the input/output basis-mask helpers and update packed transposition selectors to the upstream algorithm preceding #11646. Keep selector ordering correct when low register bits participate in multiple factors, and fold remaining low-bit register permutations into pre-shuffle selectors.

The selector implementation comes from upstream af85fc304db5, the parent of 2074a1b8a5cd (Triton #11646). The basis-mask helpers are backported from the same upstream API with compatibility for FlagTree’s LinearLayout. This commit keeps the existing conversion structure and lowering interfaces.

Validation on H20: NVIDIA/AMD build, pre-commit, five C++ test groups, focused conversion IR tests, and existing warp-local GPU conversion tests.

  • [BACKEND] Backport broadcast-local warp shuffles from Triton #11646

Allow layout conversions to read broadcast copies in the current warp or CTA, carrying warp/block-dependent owner lane bits into the shuffle map. Extend mixed transpositions to support one-sided exchanges and retain the existing mixedTranspositions.size() < 2 cost check.

Backport upstream 2074a1b8a5cd283ce584f8ebca825c4f0b215fc1: https://github.com/triton-lang/triton/pull/11646

Keep the required FlagTree compatibility changes with the feature so this commit builds independently: retain tensor-type analysis interfaces, explicit zero padding for the older resizeInDim API, existing register packing/broadcast helpers, and the existing shared-memory fallback. Adapt both common lowering and AMD permlane selection to the new fields. Backport permutation/identity checks as local analysis helpers instead of expanding FlagTree’s public layout APIs. Upstream force_warp_shuffle and GenericLinearEncoding are outside this backport.

Include applicable upstream IR coverage and updated duplication-reduction expectations. Validation: NVIDIA/AMD build, pre-commit, five C++ test groups, and focused NVIDIA/AMD conversion IR tests.

  • [Test] Cover broadcast-local layout shuffles in FlagTree

Add the 16x1 pointer layout regression from FlagTree #1047 for NVIDIA and AMD wave64. Check that conversion emits two 32-bit shuffles without shared memory or a barrier. Add layout-unit coverage for local broadcast copies, including partially broadcast warp bits and CTA locality.

Adapt and extend the upstream layout pairs into GPU tests covering forward and reverse conversions, multiple element widths, runtime pointers, owner lanes, packed exchanges, and register selection.

Validation on H20: pre-commit, five C++ test groups, five focused IR files, and 236 GPU cases passed (244 existing cases skipped). All 88 added GPU cases passed. AMD wave64 is validated by compilation, not AMD hardware.

  • [AMD][Metax] Preserve legacy layout decomposition in mixed builds

Metax builds the shared AMD permlane lowering against its own pre-#11646 analysis header. Reading regBit and dstLane from that older decomposition breaks compilation in FlagTree #1150.

Use the existing USE_MACA condition to select Metax’s pLane/transposition representation; other builds continue using shuffleMap/regBit/dstLane. Keep the original Metax shuffle gate and analysis implementation unchanged. Annotate each compatibility change with Triton #11646 and FlagTree #1150.

Validation on H20:

  • Reproduced the missing-member errors using the actual Metax analysis header; the same AMD translation unit passes syntax checking after the fix.
  • NVIDIA/AMD build and pre-commit passed.
  • All 24 existing gfx950 permlane cases produce identical LLVM IR before and after the fix.
  • Five focused conversion IR files and five C++ test groups passed.
  • All 88 broadcast-layout GPU regression cases passed.

The complete Metax SDK build and device tests remain for platform CI.

  • [CORE] Isolate layout shuffle backport with the FlagTree feature macro

Define FLAGTREE_SAME_WARP_LAYOUT_SHUFFLE directly in the default NVIDIA/AMD configuration, following the existing FlagTree feature macros.

Wrap all native source differences in the layout backport, including its prerequisites and C++ regression coverage, in matching ifdef/else/endif blocks. Preserve the original declarations, analysis, selectors, and lowering in the disabled branches. Select the legacy AMD interface when the feature is absent, so Metax does not depend on the new decomposition. Keep this feature independent of the TLE macro.

Validation on H20:

  • NVIDIA/AMD build, pre-commit, five C++ groups, five IR files, and all 88 broadcast-layout GPU cases passed.
  • All 24 gfx950 permlane LLVM IR outputs match the original backport.
  • Disabled branches restore every pre-backport nonblank source line.
  • Eight CMake backend configurations passed.
  • All five affected translation units pass syntax checking with the feature disabled and with it enabled without TLE.
  • Shared AMD lowering passes syntax checking with the actual Metax analysis header and the feature disabled.
  • [CORE] Match feature guard spacing to the original source

Remove blank lines next to conditional branch boundaries and preserve the original source spacing outside each guard. Drop empty else branches from the three blocks containing only new code.

Validation: pre-commit and git diff –check passed. All 11 feature guards match the original boundary spacing; enabled and disabled code tokens are unchanged across all seven affected native source files.


Co-authored-by: zhzhcookie zhengyang@baai.ac.cn

2天前5545次提交

github+banner-20260130 [中文版|English]

FlagTree

FlagTree 是 FlagOS 的一部分。 FlagOS 是一个面向多元AI芯片的开源、统一系统软件栈,旨在打通模型、系统与芯片层,培育开放协作的生态系统。 它支持 “一次开发,多芯运行” 的工作流,兼容多样化的 AI 加速芯片。 它释放硬件性能潜力,消除各类 AI 芯片专用软件栈之间的碎片化问题,并大幅降低大模型在多种 AI 硬件移植与维护的成本。

FlagTree 是面向多种 AI 芯片的开源、统一编译器。 FlagTree 致力于打造多元 AI 芯片编译器及相关工具平台,发展和壮大 Triton 上下游生态。 项目当前处于初期,目标是兼容现有适配方案,统一代码仓库,快速实现单仓库多后端支持。 对于上游模型用户,提供多后端的统一编译能力; 对于下游芯片厂商,提供 Triton 生态接入范例。

多后端支持

各后端基于不同版本的 Triton 适配,因此位于不同的主干分支。 各主干分支均为保护分支且地位相等,表格中所有后端均搭建了 CI/CD Runner。 有些后端适配了多个 Triton 版本,表格中仅展示最新版本。

主干分支 厂商 后端 Triton 版本 用户指引
main NVIDIA
NVIDIA TileIR
AMD
Enflame(燧原)
ILUVATAR(天数智芯)
HYGON(海光信息)
Moore Threads(摩尔线程)
DAMO ACADEMY(阿里达摩院)
Huixi(辉羲智能)
MetaX(沐曦股份)
Sunrise(曦望芯科)
KLX
T-Head(平头哥)
SpacemiT(进迭时空)
Tsingmicro(清微智能)
nvidia
tileir
amd
enflame
iluvatar
hcu
mthreads
damoacademy
rpu
metax
sunrise
xpu
ppu
spacemit
tsingmicro
3.6 nvidia guide
tileir guide
amd guide
enflame guide
iluvatar guide
hcu guide
mthreads guide
-
rpu guide
metax guide
sunrise guide
xpu guide
ppu guide
spacemit guide
tsingmicro guide
triton_v3.5.x Huawei Ascend(华为昇腾) ascend 3.5 ascend guide
triton_v3.3.x ARM China(安谋科技)
ARM64 cpu
x86_64 cpu
aipu
cpu
triton-shared
3.3 aipu guide
cpu guide
-
triton_v3.2.x Cambricon(寒武纪) cambricon 3.2 -

FlagTree 的扩展组件当前在部分后端可用:

主干分支 后端 Triton 版本 扩展组件
main nvidia
enflame
3.6 TLE-Lite
TLE-Struct GPU
TLE-Raw
HINTS
main amd
mthreads
sunrise
hcu
iluvatar
ppu
3.6 TLE-Lite
TLE-Struct GPU
main metax 3.6 TLE-Lite
main tsingmicro 3.3 TLE-Lite
TLE-Struct DSA
FLIR
triton_v3.5.x ascend 3.5 TLE-Struct DSA
FLIR
HINTS
triton_v3.3.x aipu 3.3 FLIR
HINTS

TLE(Triton Language Extensions)简介

如果要在 nvidia 后端使用 TLE 语言扩展,请使用 main 分支。 其他后端的 TLE 支持分支详见上面的表格。

Triton 在算子开发效率方面表现突出,但在多元 AI 芯片适配和更深层性能调优场景下,往往需要对分布式执行、内存访问模式和硬件相关原语提供更显式的控制。 TLE 以分层方式扩展 Triton,在保持现有 Triton 工作流兼容性的同时补齐这部分能力。

tle-speedup-20260626

TLE 的主要优势包括:

  • 从可移植到硬件导向调优的渐进式抽象(Lite / Struct / Raw)。
  • 更好覆盖多设备、架构特化与后端 lowering 场景。
  • 在保留优化空间的同时,降低现有 Triton kernel 的迁移改造成本。

详细设计、API 与示例请参考 TLE WikiTLE-Raw Wiki

性能改进

无需修改任何 Triton 算子代码,FlagTree 可在实际模型中的某些形状上获得性能增益。 下面以 Qwen 模型中调用的一些形状下的 mm 算子为例,展示 FlagTree 在不同芯片上的性能增益。

nv_h100_bf16_mm_1a nv_h100_fp32_mm_1a hcu_bf16_mm_1a hcu_fp32_mm_1a hcu_bf16_mm_3d hcu_fp32_mm_3d mthreads_bf16_mm_1a_3c mthreads_fp32_mm_1a mthreads_bf16_mm_1c mthreads_fp32_mm_1c enflame_bf16_mm_3d enflame_fp32_mm_3d

新特性

  • 2026/09/10 tsingmicro 后端升级到 Triton 3.6,加入 CI/CD。
  • 2026/08/21 amd 后端支持 TLE,加入 CI/CD。
  • 2026/08/18 新增接入 spacemit 后端(对应 Triton 3.6),加入 CI/CD。
  • 2026/08/04 新增接入 ppu 后端(对应 Triton 3.6),加入 CI/CD。
  • 2026/07/07 新增接入 NVIDIA tileir 后端(对应 Triton 3.6),加入 CI/CD。
  • 2026/07/03 iluvatar 后端升级到 Triton 3.6,加入 CI/CD。
  • 2026/07/02 xpu 后端升级到 Triton 3.6,加入 CI/CD。
  • 2026/06/30 sunrise 后端升级到 Triton 3.6,加入 CI/CD。
  • 2026/06/26 metax 后端升级到 Triton 3.6,加入 CI/CD。
  • 2026/06/10 新增接入 rpu 后端(对应 Triton 3.6),加入 CI/CD。
  • 2026/06/08 ascend 后端升级到 Triton 3.5,加入 CI/CD。
  • 2026/06/03 新增接入 ARM64 cpu 后端(对应 Triton 3.3)。
  • 2026/06/01 新增接入 damoacademy 后端(对应 Triton 3.6),加入 CI/CD。
  • 2026/05/12 mthreads 后端升级到 Triton 3.6,加入 CI/CD。
  • 2026/05/07 hcu 后端升级到 Triton 3.6,加入 CI/CD。
  • 2026/04/24 mthreads 后端升级到 Triton 3.2,加入 CI/CD。
  • 2026/04/17 enflame 后端升级到 Triton 3.6,加入 CI/CD。
  • 2026/03/13 enflame 后端升级到 Triton 3.5,加入 CI/CD。
  • 2026/01/23 新增接入 sunrise 后端(对应 Triton 3.4),加入 CI/CD。
  • 2026/01/08 添加 HINTSTLETLE-Raw 等新功能 WIKI。
  • 2025/12/08 新增接入 enflame 后端(对应 Triton 3.3),加入 CI/CD。
  • 2025/11/26 添加 FlagTree 后端特化统一设计文档 FlagTree-Backend-Specialization
  • 2025/10/28 提供离线构建支持(预下载依赖包),改善网络环境受限时的构建体验,使用方法见后文。
  • 2025/09/30 在 GPGPU 上支持编译指导 shared memory。
  • 2025/09/29 SDK 存储迁移至金山云,大幅提升下载稳定性。
  • 2025/09/25 支持编译指导 ascend 的后端编译能力。
  • 2025/09/16 新增接入 hcu 后端(对应 Triton 3.0),加入 CI/CD。
  • 2025/09/09 Fork 并修改 llvm-project,承接 FLIR 的功能。
  • 2025/09/01 新增适配 Paddle 框架,加入 CI/CD。
  • 2025/08/16 新增适配北京超级云计算中心 AI 智算云。
  • 2025/08/04 新增接入 T*** 后端(对应 Triton 3.1)。
  • 2025/08/01 FLIR 支持编译指导 shared memory loading。
  • 2025/07/30 更新 cambricon 后端(对应 Triton 3.2)。
  • 2025/07/25 浪潮团队新增适配 OpenAnolis 龙蜥操作系统。
  • 2025/07/09 FLIR 支持编译指导 Async DMA。
  • 2025/07/08 新增多后端编译统一管理模块。
  • 2025/07/02 新增接入 S*** 后端(对应 Triton 3.3)。
  • 2025/06/20 FLIR 开始承接 MLIR 扩展功能。
  • 2025/06/06 新增接入 tsingmicro 后端(对应 Triton 3.3),加入 CI/CD。
  • 2025/06/04 新增接入 ascend 后端(对应 Triton 3.2),加入 CI/CD。
  • 2025/06/03 新增接入 metax 后端(对应 Triton 3.1),加入 CI/CD。
  • 2025/05/21 FLIR 开始承接到中间层的转换功能。
  • 2025/04/09 新增接入 aipu 后端(对应 Triton 3.3),提供 torch 标准扩展范例,加入 CI/CD。
  • 2025/03/26 接入安全合规扫描。
  • 2025/03/19 新增接入 xpu 后端(对应 Triton 3.0),加入 CI/CD。
  • 2025/03/19 新增接入 mthreads 后端(对应 Triton 3.1),加入 CI/CD。
  • 2025/03/12 新增接入 iluvatar 后端(对应 Triton 3.1),加入 CI/CD。

环境准备

避免环境匹配问题的最佳实践是使用 用户手册 中推荐的镜像。

从源码安装

安装依赖(注意使用正确的 python3.x 执行):

apt update; apt install zlib1g zlib1g-dev libxml2 libxml2-dev nlohmann-json3-dev
python3 -m pip install -r python/requirements.txt

通用的构建安装方式(网络畅通环境下推荐使用):

# Set FLAGTREE_BACKEND using the backend name from the table above
export FLAGTREE_BACKEND=${backend_name}  # Do not set it on nvidia/amd/triton-shared

# For Triton 3.1/3.2/3.3 (branch: triton_v3.1.x, triton_v3.2.x, triton_v3.3.x)
cd python; python3 -m pip install . --no-build-isolation -v  # Install flagtree and uninstall triton

# For Triton 3.4/3.5/3.6 (branch: triton_v3.4.x, triton_v3.5.x, main)
python3 -m pip install . --no-build-isolation -v             # Install flagtree and uninstall triton

安装 flagtree 后,可通过下列命令查看:

python3 -m pip show flagtree
cd ${ANY_DIR_OTHER_THAN_FLAGTREE_PYTHON}; python3 -c 'import triton; print(triton.__path__)'

免源码安装

参见 用户手册

关于贡献

欢迎参与 FlagTree 的开发并贡献代码,详情参见 CONTRIBUTING.md

许可证

FlagTree 使用 MIT license

邀请码
    Gitlink(确实开源)
  • 加入我们
  • 官网邮箱:gitlink@ccf.org.cn
  • QQ群
  • QQ群
  • 公众号
  • 公众号

版权所有:中国计算机学会技术支持:开源发展技术委员会
京ICP备13000930号-9 京公网安备 11010802047560号