昇腾算子开发体系：AOL、Ascend C、CATLASS、Triton-Ascend、TileLa

昇腾算子开发并不是只有一条路径。不同开发接口本质上是在“开发效率、可移植性、可控性、极致性能”之间做不同取舍。

整理日期：2026-09-10。

1. 先区分“调用算子”和“开发算子”

这是理解昇腾算子生态最重要的第一步。

调用现有算子

如果框架需要的算子已经存在，优先直接调用 CANN 提供的高性能实现。例如 AOL/ops 中已有 MatMul、Normalization、Attention 相关基础或融合算子。

开发自定义算子

只有当以下情况出现时，才需要进入 Kernel 开发：

模型出现新的 Operator；
现有 Operator Shape/数据类型不匹配；
希望做算子融合；
希望针对特定模型/硬件做性能极致优化；
上游框架缺失 NPU Kernel。

因此优化顺序通常不是“所有东西都自己写 Kernel”，而是：

已有高性能算子
→ 组合/融合
→ 模板库
→ Tile DSL
→ 底层 Kernel 手写

2. AOL：昇腾的高性能算子库

AOL（Ascend Operator Library）是 CANN 提供的高性能算子库。官方描述强调其 Operator 针对 Ascend AI Processor 做硬件亲和优化，并提供可直接调用的 API。

从开发者角度，可以把 AOL 看成“已经优化好的常见算子资产”。

但不要把它机械等同于 cuDNN。NVIDIA 生态中算子能力分散在 cuDNN、cuBLAS、cuBLASLt、CUTLASS、FlashAttention 等不同层次；昇腾同样存在 ops、NNAL、ATB、CATLASS 等多个组件。二者只能做功能角色类比，不能一一对应。

3. Ascend C：底层 Kernel 编程主路径

Ascend C 是面向昇腾 AI Core 的 Kernel 编程语言/接口体系。它允许开发者显式控制：

数据从 Global Memory 到片上存储的搬运；
Tile 划分；
Vector/Matrix 计算；
Pipeline；
多核任务切分；
同步和事件；
数据布局与 Double Buffer 等优化。

它的学习价值非常高，因为它直接暴露出 NPU 的数据流思维。

对 AI Infra 工程师来说，学习 Ascend C 不只是为了“写算子”，更重要的是理解：

为什么某个算子在 NPU 上快或慢。

这与学习 CUDA Kernel 的价值类似。

4. BiSheng：Ascend C 如何变成可执行代码

BiSheng 编译器面向 Ascend AI Processor，可将 Ascend Operator/Device 代码编译成二进制可执行文件和动态库，并支持 AI Core Device 指令集编译。

典型链路可以抽象为：

Ascend C / Kernel Source
↓
BiSheng Frontend + Optimization + Device Compilation
↓
NPU Target Code
↓
CANN Runtime Launch

因此 BiSheng 不是一个“算子库”，而是代码生成链路中的编译器。

5. CATLASS：昇腾的高性能线性代数算子模板库

CATLASS 全称 CANN Templates for Linear Algebra Subroutines。

它重点面向：

MatMul / GEMM；
Grouped MatMul；
量化 MatMul；
Flash Attention 相关计算；
其他以矩阵乘为核心的融合算子。

CATLASS 的核心思想与 NVIDIA CUTLASS 有相似之处：通过多层抽象和模板组件，把高性能 Kernel 的常见结构沉淀下来，使开发者不需要每次从裸底层重新实现。

CATLASS 典型层次包括 Kernel、Block、Tile、Basic 等抽象。

2026 年值得注意的新变化

CATLASS 2.0 在 2026 年 8 月加入了 CATLASS DSL：使用 Pythonic Kernel 表达，原生支持 Tile 编程，并通过 MLIR/AscendNPU-IR 工具链生成 Kernel。

这说明昇腾算子开发正在同时向两个方向发展：

底层模板化：保留性能控制能力；
高层 DSL 化：减少手写 C++ Kernel 的复杂度。

6. Triton-Ascend：把 GPU 社区成熟的 Tile 编程范式带到 NPU

Triton 的核心抽象不是“一个线程写一份标量代码”，而是让程序员以一个数据块（Tile）为中心描述计算。

例如矩阵乘可以抽象成：

Load Tile A
Load Tile B
→ Tile MatMul
→ Accumulate
→ Store Tile C

Triton-Ascend 是 Triton 面向 Ascend 的硬件适配实现。对开发者而言，其价值主要是：

Python 开发体验；
Tile 级抽象；
编译器自动完成较多底层映射；
更容易迁移已有 Triton Kernel 设计思想。

但“同一段 Triton Kernel 在 CUDA 和 Ascend 上完全无差异”并不成立。硬件 Memory Hierarchy、执行单元、Intrinsic、编译后端都不同，极致性能仍需要硬件特化。

7. TileLang：更高层的 Tile DSL 路线

TileLang 同样属于 Tile 编程思想：开发者描述块级计算、数据搬运和调度，编译系统负责向底层映射。

从学习角度，你可以把 Triton、TileLang、CATLASS DSL、PyPTO 放在同一张“抽象层级地图”里理解，而不是孤立背工具名。

8. PyPTO / PTO：为什么它不是简单的 PTX 替代品

你之前问过 PTO 与 PTX 的关系，最重要的结论是：不要简单写 PTO = PTX。

PTO 更强调 Tile 级操作与编程语义，抽象层通常高于硬件机器指令。PTO 项目的公开工具链可以进一步经过编译器 Lowering、Bufferization 或 C++ Library 实现路径，最终映射到底层 Ascend Device Code。

更合理的心智模型是：

Python / DSL
↓
Tile-level Program / PTO
↓
Compiler Lowering
↓
Ascend Device Representation
↓
NPU Executable Code

而 NVIDIA PTX 通常被看作 GPU 的虚拟 ISA，距离机器 ISA 更近。

9. HCCL 为什么也出现“通信算子开发”

原始笔记写“基于 CANN 的 Ascend C 和 HCCL 开发算子”，容易产生误解。

HCCL 的核心定位是 集合通信库，典型原语包括：

AllReduce；
AllGather；
ReduceScatter；
AllToAll；
Broadcast；
Send / Recv。

现代大模型中通信本身也会被实现成高度优化的 Device Kernel，因此 HCCL 文档里确实存在“通信算子开发”。但它解决的是 Collective Communication Kernel，而不是普通 MatMul/Softmax 等计算算子。

10. 不同算子开发路径应该怎么选

路径	开发效率	性能控制	适用场景
AOL / 已有 ops	最高	低	已有成熟算子，直接调用
Triton-Ascend / TileLang	高	中高	快速开发 Tile Kernel、模型适配
CATLASS DSL	高	中高	矩阵乘/Attention 类算子，Pythonic Tile 开发
CATLASS C++	中	高	GEMM 类极致优化与模板复用
Ascend C	较低	最高	底层特化、融合、新算子、性能极致优化
PyPTO / PTO	高	取决于编译链	探索 Tile ISA/编译器自动优化与更高层编程模型

11. AI Infra 学习者应该怎么学

建议不要同时学五套 DSL。按以下顺序效果最好：

先学硬件与 Roofline：知道算力、带宽、容量和通信的关系。
学 Ascend C：理解 NPU 的真实数据流和 Kernel 执行模型。
学 CATLASS：理解高性能 GEMM 如何模板化。
学 Triton-Ascend：理解编译器如何把 Tile DSL 自动映射到底层。
再看 PyPTO/PTO：研究下一代 Tile ISA/编译抽象。

最终目标不是“工具都会用”，而是能够回答：

同一个 MatMul / Attention Kernel，为什么 Ascend C、CATLASS、Triton 和 PTO 的写法不同？这些差异最终如何映射到同一套硬件资源？

参考资料

2026-09-18 增量：Mmad、GEMM与融合优化

问题起点：Mmad是什么矩阵运算，与GEMM/MMA区别在哪？优化是否只盯最耗时kernel，GM/UB分别是什么？

数学层面，GEMM通常表示C←αAB+βC；矩阵乘累加描述的是累加器加上矩阵块乘积。MMA/Mmad属于更靠近硬件或设备API的块级运算语境，完整GEMM还涉及分块、搬运、同步、尾块和写回；不能将单条矩阵运算调用等同于完整高性能算子。不同平台的名字、形状和布局契约不同，具体Mmad签名须查目标CANN版本API。本次不补造支持类型/块尺寸。

GM是设备全局存储语境；UB是片上Unified Buffer语境，常用于向量侧局部数据处理。完整矩阵通路还可能涉及其他片上缓冲，不能画成“所有矩阵都直接GM→UB→Mmad”的统一流程。Ascend C官方开发概述、CUTLASS官方仓库（2026-09-20核查）

优化优先级看端到端关键路径和可消除成本，而不是只看累计耗时第一名。多个小kernel间反复写回GM、启动和同步的成本可能可通过融合消除；但融合也可能增加寄存器/UB占用，降低并行度。通信重叠区内的热点和关键路径上的热点价值不同。

教学上界：若某段占总时长30%，该段加速2倍，则整体最多约1/(0.7+0.3/2)=1.176倍；这要求其他部分不变。对有重叠的时间线不能直接将所有kernel时间求和作为占比。

实践保留未优化基线、真实shape分布、数值误差和端到端延迟；先证明中间张量搬运或同步是瓶颈，再写融合算子。

2026-09-21—23 深度整理入口

本轮从工具名进一步下钻到数据流与工程验收：

TileLang的上游与Ascend适配仓是不同核查对象。TileLang、tilelang-ascend（2026-09-23）。有适配仓不等于任意上游kernel可原样运行；后端、目标芯片、dtype、intrinsic和编译版本需匹配。

PyPTO、TileLang等更高层表达可以减少样板，但不能免去数据依赖和存储预算。不要由“Python写kernel”推导“不需要理解硬件”。

仿真应区分功能语义验证、指令/周期性能模型与真实板卡。CPU执行仿真程序可以验证某些设备语义，不意味着CPU内部真的拥有对应NPU通路；没有误差标定的仿真时间不能预测真机吞吐。本次未取得950仿真产品版本与精度说明，未宣称能逼真模拟全部硬件。

FlashAttention属于减少标准注意力IO的精确算法实现路线，不是线性注意力的同义词；不同顺序仍有浮点差异。官方项目。FA、KV、MoE、通信、量化是相互连接的优化维度，学习项目可以贯通它们，而不是把五项当作一个固定算子名称。

2026-10-07 增量核查：新算子封装、样例拆分与DeepSeek昇腾组件

来源：9月23日新自定义算子/CCE讨论，9月30日样例与950/910B/C支持追问。以下状态按2026-10-07官方仓库核查。

新算子必须同时解释Kernel和集成

CUDA扩展的host入口、动态库、设备kernel与框架注册是不同对象；Ascend同样需区分host tiling、device代码、编译产物及调用/注册路径。Ascend C直调与框架集成、旧TBE/OPP路径不必生成相同目录/包。比较“.so对一大堆文件”前，先确认比较的是同一种集成方式和开发/运行产物。

你问的是新自定义算子，不能只回答“已有算子怎么调用”。实践以目标版本一个Add或Matmul融合样例为起点，画出编译—链接/注册—运行时launch链，逐个标明产物消费者，再判断哪些文件真正需要分发。

TBE/CCE的“非主流入口”“产品弃用”“底层编译目标不再支持”是不同命题。本轮没有匹配未来产品的正式弃用公告，不采纳旧回答“TBE/CCE都死了”或未来必然不能运行的断言。样例旧路径失效也不是ISA弃用证据。

官方样例入口分别核对Ascend/samples与CANN asc-devkit。9月30日你提到operator_contrib、MatmulLeakyReluCustom路径；尚未锁定当时commit与本地环境，本文不补造一条通用cd/编译命令。

开源不等于跨代立即兼容

DeepGEMM-Ascend标注2026-09-30初始发布，开发/验证目标是Ascend 950，要求CANN 9.20等配套；scale布局与NVIDIA路径有差异，不可原样互换packed数据。

DeepEP-Ascend当前验证950DT及指定PoC软件栈，对其他代际支持未由这些测量建立。仓库中的极限带宽/利用率是作者特定配置结果，不是你的机器实测。

FlashMLA在2026-09-30加入950稀疏prefill/decode；当前融合norm/RoPE/attention/cast测试路径仍标CUDA-only。该发布还移除Hopper及部分早期模型支持并改变FP8/FP4 KV格式，因此升级不能只替换包而继续假设旧缓存/模型兼容。需要旧组合时，按官方链接选择相应历史commit。

tilelang-ascend当前README明确A2/A3已测试；这与上面950特化库是独立支持范围。不能概括为“昇腾的新开源都不支持910”，也不能由TileLang支持A2/A3推出950特化kernel兼容910。

把支持状态做成自己的验证矩阵

分别记录芯片、CANN/driver/HDK、torch_npu、仓库commit、算子类型、dtype/布局、模型、正确性与性能。区分官方声明、可编译、可运行、质量通过、性能达标五种状态；代码量/star数都不能代替完成度。

尚未取得你设备的软件栈及这些新kernel的运行记录，未宣称已完成部署。高层编程与tile成本的实践入口见GEMM、Triton 与 TileLang：从分块代码到 Roofline 性能账本。