AscendNPU IR 语言生态对接指南：Triton / TileLang / DLCompiler / FlagTree 如何使能昇腾

• 编译器
• 人工智能
• Ascend
• CANN

AscendNPU IR（AscendNPU Intermediate Representation）是基于 MLIR 构建、面向昇腾亲和算子编译的中间表示。本文以官方"用户生态"文档为主体，系统介绍已对接或正在使用 AscendNPU IR 的编程语言与框架（Triton-Ascend、TileLang-Ascend、DLCompiler、FlagTree），并结合仓库内集成文档与源码，讲解环境搭建、算子编译流程、IR 转换路径与编译选项，帮助读者理解"生态框架如何通过 AscendNPU IR 使能昇腾 AI 处理器"。

生态对接总览

AscendNPU IR 的目标之一是通过开源社区开放接口，支持生态框架灵活对接，高效使能昇腾 AI 处理器。当前项目已使能多个算子编程框架接入昇腾后端，并提供面向昇腾的编译优化能力。从实现角度看，生态接入主要发生在 AscendNPU IR 的Conversion（转换）层：仓库 bishengir/include/bishengir/Conversion 目录下可以看到TorchToHFusion、HFusionToHIVM、HIVMToTritonGPU、TritonAscendGPUToLLVM等一系列转换目录，它们分别承担"生态框架 IR → HFusion"与"AscendNPU IR 内部方言 → 底层指令"的桥接职责。

对接方式分为两类（详见 框架集成）：

• DSL 集成：通过 Triton、TileLang 等领域特定语言编写算子，编译为 AscendNPU IR 执行；
• IR 集成：通过 Torch IR、Linalg/HFusion IR、HIVM IR 等多级表示接入，支持自动融合与拆分，生成昇腾亲和的高性能算子。

语言生态总览

官方"用户生态"文档（英文版、中文版）列出的已对接或使用 AscendNPU IR 的语言与框架如下：

其中，Triton-Ascend 与 TileLang-Ascend 在仓库中拥有完整的集成文档与代码示例，是理解"生态框架对接 AscendNPU IR"的最佳入口；DLCompiler 与 FlagTree 则代表了基于 Triton 语言的跨架构编译生态对昇腾后端的支持。

多层次 IR 抽象：生态对接的底层基础

生态框架之所以能灵活对接，源于 AscendNPU IR 提供的多级 IR 抽象（详见 IR 集成总览）：

Torch-MLIR / Triton (框架 / DSL 层) | v Linalg / Tensor (通用张量代数层) | v HFusion (硬件感知的融合调度层) | v HIVM (NPU 指令层) | v LIR -> Binary (机器码生成)

• Linalg / Tensor 层：使用标准 MLIR 方言表达算子语义，支持Elemwise、Broadcast、Reduce、Transpose、Concat等操作；
• HFusion 层：提供昇腾感知的 Named Op（如hfusion.elemwise_unary、hfusion.cast、hfusion.select、hfusion.reduce_with_index），自动完成 bufferization、切分与调度；
• HIVM 层：直接映射 NPU 硬件指令，显式控制内存层次（GM/UB/L1/L0）、计算流水（Vector/Cube/MTE）与同步原语，支持细粒度性能调优。

这套分层设计正是 架构设计 中描述的 HFusion / HIVM / HACC 三大自研方言的分工：HFusion 负责硬件无关优化与生态对接，HIVM 负责细粒度感知 NPU 硬件细节并将高层语言映射为底层 NPU 指令，HACC 负责异构硬件抽象（Host/Device 编程模型与启动语义）。

Triton-Ascend：Triton 生态接入昇腾

Triton-Ascend 集成指南 是仓库中最完整的生态对接文档。Triton 是主流的算子高性能开发语言，Triton-Ascend 将 Triton 算子转换为 MLIR 并接入 AscendNPU IR 生态，使 Triton Kernel 能在昇腾 NPU 上运行。

环境准备与安装

Triton-Ascend 要求 Python 3.9～3.11（含），运行时依赖昇腾 CANN 环境、torch_npu包与triton-ascend包。以 Atlas A3 系列产品的 x86 环境为例：

# 1. 安装 CANN Toolkit 与对应 ops 包（{version} 替换为实际版本号，如 9.0.0） chmod +x Ascend-cann_{version}_linux-x86_64.run chmod +x Ascend-cann-A3-ops_{version}_linux-x86_64.run ./Ascend-cann_{version}_linux-x86_64.run --full [--install-path=${PATH-TO-CANN}] ./Ascend-cann-A3-ops_{version}_linux-x86_64.run --install [--install-path=${PATH-TO-CANN}] pip install attrs==24.2.0 numpy==1.26.4 scipy==1.13.1 decorator==5.1.1 psutil==6.0.0 pyyaml # 2. 设置环境变量（8.5.0 及更早版本路径为 ${PATH-TO-CANN}/ascend-toolkit/set_env.sh） source ${PATH-TO-CANN}/cann/set_env.sh # 3. 安装 torch_npu（固定匹配版本 2.7.1）与 triton-ascend pip install torch_npu==2.7.1 pip install triton-ascend

Triton Kernel 调用与验证

安装完成后，可通过pytest -sv <file>.py运行 Triton Kernel 并验证功能，正确时输出PASS。以tl.load/tl.store实现的逐元素比较 Kernel 为例：

import triton import triton.language as tl import torch import torch_npu @triton.jit def triton_lt(in_ptr0, in_ptr1, out_ptr0, XBLOCK: tl.constexpr, XBLOCK_SUB: tl.constexpr): offset = tl.program_id(0) * XBLOCK base1 = tl.arange(0, XBLOCK_SUB) loops1: tl.constexpr = XBLOCK // XBLOCK_SUB for loop1 in range(loops1): x_index = offset + (loop1 * XBLOCK_SUB) + base1 tmp0 = tl.load(in_ptr0 + x_index, None) tmp1 = tl.load(in_ptr1 + x_index, None) tmp2 = tmp0 < tmp1 tl.store(out_ptr0 + x_index, tmp2, None) # 调用：grid 中配置并行粒度（ncore），XBLOCK / XBLOCK_SUB 控制切分大小 triton_ltncore, 1, 1

该示例同时展示了 Triton-Ascend 的动态切分与动态 shape 能力：并行粒度通过[]中的 grid 参数配置，Kernel 可自动适配任意长度的一维张量。

Triton 算子到 AscendNPU IR 的转换

Triton-Ascend 逐步将 Triton 方言的高层 GPU 抽象算子降级为目标方言（Linalg、HFusion、HIVM 等）。以下是文档中给出的部分映射关系：

从仓库实现看，这一降级链路对应 HIVMToTritonGPU 与 TritonAscendGPUToLLVM 等转换目录，后者在昇腾 950PR/950DT 上还实现了 SIMT 编译支持，将社区 TritonGPU 方言连接到 HIVM。

Triton Ascend 扩展能力

AscendNPU IR 为 Triton-Ascend 提供了一套增量语言特性，需要导入如下模块使用：

import triton.language.extra.cann.extension as al

主要包括：

• 同步与调试操作：debug_barrier（向量流水内同步模式，如al.SYNC_IN_VF.VV_ALL）、sync_block_set/sync_block_wait（计算单元与向量单元间的同步事件，须成对使用）、sync_block_all（全局同步，支持"all_cube"、"all_vector"、"all"、"all_sub_vector"模式，event_id取值 0～15）；
• 硬件查询与控制操作：sub_vec_id/sub_vec_num（查询当前 AI Core 上 Vector core 索引与数量）、parallel（带并行执行语义的迭代器，支持num_stages流水级数与loop_unroll_factor展开因子）；
• 编译优化提示：compile_hint（向编译器传递优化提示，如"bitwise_mask"）、multibuffer（多缓冲/双缓冲）、scope（指定块内计算核心类型"cube"或"vector"）；
• 张量切片与计算操作：insert_slice/extract_slice、get_element、sort、flip、cast；
• 索引与集合操作：_index_select、index_put、gather_out_to_ub、scatter_ub_to_out、index_select_simd等，均支持 2D～5D 张量与边界检查。

此外，Triton-Ascend 还支持通过register_custom_op注册自定义算子，通过al.custom('my_custom_op', src, index, out=dst)调用，并可在注册类中通过构造函数做参数合法性校验、动态设置符号名与编译命令（source/compile属性，%<与%@分别表示源文件与目标文件，类似 Makefile 语法）。

TileLang-Ascend：Tile 级编程框架接入

TileLang 集成指南 介绍了 tilelang-ascend（npuir分支）：一个面向昇腾 NPU 架构的 tile-lang 方言变体，基于 tile-lang 的 Python 风格语法与 TVM 编译器基础设施，支持 GEMM、向量运算、注意力机制等算子的开发。其编译后端提供 AscendNPU IR 与 Ascend C & PTO 两条技术路线，AscendNPU IR 路线使 tilelang 与基于 MLIR 的开源 AI 编译器生态无缝集成。

环境与构建

基础环境需安装 CANN Toolkit 并设置环境变量，Python 版本要求 3.7.0～3.11.4：

source /path/to/install/Ascend/ascend-toolkit/set_env.sh pip3 install attrs cython 'numpy>=1.19.2,<=1.24.0' decorator sympy cffi pyyaml pathlib2 psutil protobuf==3.20.0 scipy requests absl-py export ACL_OP_INIT_MODE=1

源码构建有两种方式：一是直接拉取代码后由脚本构建内置的 AscendNPU IR，二是复用本地已构建的 AscendNPU IR：

git clone https://github.com/tile-ai/tilelang-ascend.git --recursive -b npuir cd tilelang-ascend bash install_npuir.sh # 或者使用本地构建的 AscendNPU IR bash install_npuir.sh --bishengir-path=/path/to/AscendNPU-IR/build/install

构建完成后还需安装pybind11与torch_npu。

快速开始：向量加法 Kernel

以下示例用 TileLang 实现长度为 4096 的float32向量加法：将数据从全局内存拷入片上统一缓冲区（UB），调用底层 NPU 指令npuir_add完成逐元素相加，再写回全局内存，最后与 PyTorch 原生加法对比验证正确性：

import tilelang import tilelang.language as T import torch import torch_npu dtype = "float32" seq_len = 4096 def vec_add(N, block_N, dtype="float32"): n_num = N // block_N @T.prim_func def main(A: T.Tensor((N), dtype), B: T.Tensor((N), dtype), C: T.Tensor((N), dtype), shape: T.int32): with T.Kernel(n_num, is_npu=True) as (cid, _): A_VEC = T.alloc_ub((block_N), dtype) B_VEC = T.alloc_ub((block_N), dtype) C_VEC = T.alloc_ub((block_N), dtype) start_idx = cid * block_N remaining = shape - start_idx tail_size = T.min(block_N, remaining) T.copy(A[start_idx], A_VEC[:tail_size]) T.copy(B[start_idx], B_VEC[:tail_size]) T.npuir_add(A_VEC, B_VEC, C_VEC) T.copy(C_VEC[:tail_size], C[start_idx]) return main func = vec_add(seq_len, seq_len) compiled_kernel = tilelang.compile(func, target="npuir") compiled_kernel(v1, v2, v3, seq_len)

运行python3 test_tilelang.py后，TileLang Kernel 输出与 PyTorch 参考结果一致。

生成的 AscendNPU IR

开启export TILELANG_DUMP_IR=1可同时打印 TVM IR 与 AscendNPU IR。生成的 IR 中可以看到：hacc.entry/hacc.function_kind = #hacc.function_kind<DEVICE>标记 kernel 入口与 DEVICE 侧属性，#hivm.address_space<gm>/<ub>标注内存层次，memref.copy完成 GM→UB 搬运，hivm.hir.vadd执行片上向量加法。这正是"TileLang 算子 → AscendNPU IR(HIVM) → 昇腾 NPU 执行"的完整证据链。

DLCompiler 与 FlagTree

除上述两个在仓库中有完整集成文档的项目外，官方"用户生态"文档还列出了两个基于 Triton 语言的编译器项目：

• DLCompiler：扩展 Triton 的深度学习编译器，支持跨架构 DSL 扩展与智能自动优化；
• FlagTree：基于 Triton 语言的开源 AI 编译器，提供跨多后端的统一编译能力。

它们与 Triton-Ascend 一样，均以 Triton 语言为编程入口。从源码结构看，仓库 triton 目录完整引入了 Triton 社区代码（含lib/Conversion、lib/Dialect等），并提供了ProtonAscendGPUToLLVM、TritonAscendGPUToLLVM等昇腾后端转换，这为"跨架构 DSL 扩展 / 多后端统一编译"类项目复用 Triton 前端、借助 AscendNPU IR 后端使能昇腾提供了基础设施支撑。

框架级 IR 集成路径

除了 DSL 接入，AscendNPU IR 还支持 PyTorch / TensorFlow / MindSpore 等框架通过 IR 形式集成（详见 框架集成）：

• Torch IR 集成：使用 Torch 方言的 ATen 算子，经convert-torch-to-hfusion等 pass 自动转换为 Linalg/HFusion Named Op。文档中给出了覆盖 55+ 个 ATen 算子的映射表，例如aten.mul.Tensor→linalg.binary_fn<mul>、aten.relu→hfusion.unary_fn<relu>、aten.sum→linalg.reduce+arith.addf/addi。命令行调用示例：bishengir-opt -torch-backend-to-named-op-backend-pipeline torch.mlir -o torch_to_hfusion.mlir；
• Linalg/HFusion IR 集成：直接使用标准 MLIR 方言表达算子语义，进入自动融合与调度流程，支持Elemwise、Broadcast、Reduce、Transpose、Concat等算子类型的融合切分；
• HIVM IR 集成：面向需要细粒度控制硬件行为的场景，直接用 HIVM 方言编写 Kernel，显式管理内存层次与计算流水，如hivm.hir.load/hivm.hir.store执行显式 DMA 搬运、hivm.hir.vadd完成片上计算。相关实现位于仓库 bishengir/include/bishengir/Dialect/HIVM。

编译选项与函数属性

bishengir-compile是 AscendNPU IR 编译器的命令行驱动（见 bishengir/tools/bishengir-compile），它提供的公共编译选项同样适用于上述生态集成路径：

支持的目标设备包括：Ascend 950PR/Ascend 950DT、Atlas A3 训练/推理产品、Atlas A2 训练/推理产品。所有集成路径共享以下 kernel 入口函数属性：hacc.entry（标记 kernel 入口）、hacc.function_kind = #hacc.function_kind<DEVICE>（DEVICE 侧函数）与hacc.function_kind = #hacc.function_kind<HOST>（HOST 侧函数，HFusion 自动 outline 设备 kernel）。

以 Torch IR 端到端编译为例：

bishengir-compile -enable-torch-compile=true -enable-hfusion-compile=true -enable-hivm-compile=true -target=Ascend910B1 torch.mlir -o torch_kernel.o

更多资料

• Triton 集成指南：安装、环境、算子映射与 Ascend 扩展接口；
• TileLang 集成指南：环境、构建与快速上手；
• 框架集成指南：Torch / Linalg-HFusion / HIVM 三级 IR 接入；
• IR 集成总览：多级 IR 抽象架构与编译选项；
• 架构设计：HFusion / HIVM / HACC 方言分工与编译流程；
• 性能调优最佳实践：Triton 算子迁移昇腾的切分策略、bitwise_mask与对齐优化案例。

• 编译器
• 人工智能
• Ascend
• CANN