【CANN】-Ascend C 算子开发

适用: 有 CUDA / HIP 算子开发经验, 或 PyTorch 深度用户想深入 AI Core 优化, 第一次在昇腾 NPU 上写自定义算子的工程师
配套仓: Ascend/samples (cplusplus/level1_single_api/4_op_dev/6_ascendc_custom_op)
配套版本: CANN 6.0.RC1.alpha005+ (主) / 5.0.4+ (历史)
关联博客: 《【昇腾训练】-MindSpeed 开发入门》 (训练) / 《【昇腾推理】-MindIE 极速入门》(推理)

---

一、AscendC 不是"用 C++ 重写 Python 算子", 是面向昇腾 AI Core 的专用算子开发语言

关键洞察: 多数新人把 AscendC 当作"用 C++ 写 PyTorch 算子"或"昇腾版的 CUDA", 实际它是 面向昇腾 AI Core 硬件的专用算子开发语言, 范式是 CopyIn → Compute → CopyOut 三段式, 配合 Tiling 设计实现多核并行 + 内存分块. 这与 CUDA 写 kernel 的思路完全不同.

AscendC 定位 = 昇腾 AI Core 专用算子开发语言

维度	AscendC	CUDA C	TBE (老方案)
出身	华为官方 (CANN 6.0+)	NVIDIA 官方	华为官方 (CANN 5.x 主流)
目标硬件	昇腾 AI Core / Vector Core	NVIDIA GPU SM	昇腾 AI Core (DSL 风格)
范式	CopyIn → Compute → CopyOut	自由风格	TIK DSL (Python 风格)
内存管理	TBuf / TQue 队列	shared memory / registers	自动
多核切分	Tiling 函数 (用户写)	block / grid	自动
推荐用法	2026 新项目首选	NVIDIA 项目	旧项目维护

新人最常踩的概念坑: 把 AscendC 当 “CUDA for 昇腾” — 实际思路完全不同:

• CUDA: 一个 thread 写一个元素, block 内同步
• AscendC: 全局内存 → Local Memory → 算子计算 → 写回, 强调显式数据搬运 + 多核切分

AscendC 在昇腾生态中的位置

昇腾技术栈 (从底向上)
├── 硬件层: NPU 芯片 (910B3 / 310P)
├── 驱动层: npu driver
├── 算子层: AscendC 算子 / TBE 算子 / 内置算子库     ← 本博客主角
├── 框架层: torch_npu / MindSpore
├── 模型层: PyTorch / MindSpore / Megatron-LM / DeepSpeed
└── 应用层: vLLM-Ascend / MindIE / MindSpeed

关键洞察: AscendC 是算子层语言, 不是模型层. 多数应用开发者 (训练 / 推理) 不需要直接写 AscendC, 但需要理解它的能力 (知道什么算子能优化, 什么瓶颈在算子层). 模型层的人通过 torch_npu 间接使用 AscendC 写好的算子.

---

二、AscendC 范式: CopyIn → Compute → CopyOut

关键洞察: AscendC 的"三段式"不是风格选择, 是硬件决定 —— 昇腾 AI Core 的存储层次 (Global Memory → Local Memory → Register) 决定了数据必须显式搬运 + 多核切分, 范式直接对应这个层次.

三段式范式详解

输入 GlobalTensor
    ↓ CopyIn (从 Global → Local)
LocalTensor (在 Local Memory / UB)
    ↓ Compute (在 AI Core / Vector Core 上计算)
LocalTensor (结果)
    ↓ CopyOut (从 Local → Global)
输出 GlobalTensor

3 段对应 3 个核心 API

阶段	核心 API	作用
CopyIn	DataCopy / DataCopyPad	全局 → 本地 (含 32B 对齐)
Compute	Add / Mul / ReduceSum 等内置指令	本地计算 (在 AI Core 上)
CopyOut	DataCopy / EnQue / DeQue	本地 → 全局

对比 CUDA 写 kernel 的差异

步骤	CUDA	AscendC
数据读取	array[i] 直接读	必须显式 DataCopy 搬到 Local
计算	result[i] = a[i] + b[i]	Add(local_a, local_b, local_c)
结果写回	output[i] = result[i]	必须显式 DataCopy 写回 Global
同步	__syncthreads()	EnQue / DeQue (队列同步)

关键洞察: AscendC 的"显式数据搬运"是性能来源 —— AI Core 不知道你要算什么, 你必须告诉它"从哪里读 → 算 → 写到哪里". 显式意味着可控, 也意味着初学者会觉得繁琐.

---

三、核心数据结构 (GlobalTensor / LocalTensor / TBuf / TQue)

关键洞察: AscendC 的 4 类数据结构对应 AI Core 存储层次 + 计算流程, 不是"随便起的类名". 每个都有明确语义.

4 类数据结构对照

数据结构	存储位置	生命周期	何时用
GlobalTensor	Global Memory (HBM)	整个算子	输入 / 输出张量
LocalTensor	Local Memory (UB / L1)	一个核内	计算的中间结果
TBuf	Unified Buffer 分配器	静态分配	申请 Local Memory 空间
TQue	队列 (双缓冲)	动态队列	CopyIn / Compute / CopyOut 之间的同步

典型用法示例 (Add 算子)

// 1. 准备 TBuf (申请 UB 空间)
TBuf<QuePosition::VECIN, TPosition::LCM> buf1, buf2, buf3;
buf1.InitBuffer(que1, totalLength * sizeof(TYPE));
buf2.InitBuffer(que2, totalLength * sizeof(TYPE));
buf3.InitBuffer(que3, totalLength * sizeof(TYPE));

// 2. CopyIn: Global → Local
LocalTensor<TYPE> aLocal = que1.AllocTensor<TYPE>();
LocalTensor<TYPE> bLocal = que2.AllocTensor<TYPE>();
DataCopy(aLocal, aGlobal, totalLength);
DataCopy(bLocal, bGlobal, totalLength);
que1.EnQue(aLocal);
que2.EnQue(bLocal);

// 3. Compute: Local 上做加法
LocalTensor<TYPE> aLocal2 = que1.DeQue<TYPE>();
LocalTensor<TYPE> bLocal2 = que2.DeQue<TYPE>();
LocalTensor<TYPE> cLocal = que3.AllocTensor<TYPE>();
Add(cLocal, aLocal2, bLocal2, totalLength);
que1.FreeTensor(aLocal2);
que2.FreeTensor(bLocal2);

// 4. CopyOut: Local → Global
DataCopy(cGlobal, cLocal, totalLength);
que3.EnQue(cLocal);
LocalTensor<TYPE> cLocal2 = que3.DeQue<TYPE>();
que3.FreeTensor(cLocal2);

关键洞察: 这段代码体现了 AscendC 的 3 大特点 — (1) 显式内存管理 (TBuf / AllocTensor / FreeTensor) (2) 队列同步 (EnQue / DeQue 实现 CopyIn 与 Compute 的解耦) (3) 多核可并行 (每核独立 TBuf, 全局汇总).

---

四、Tiling 设计 (多核切分的核心)

关键洞察: Tiling 是 AscendC 最难也最关键的优化. 没 Tiling 就没多核, 没多核就没性能. 8 卡 910B3 的 80%+ 算力 = Tiling 表 + 多核并行 + UB 复用.

Tiling 的作用

把一个大 tensor (例如 1024×1024) 切成多个小 tile (例如 64×64), 每个 tile 分配给一个核 (AI Core), 核内完成该 tile 的计算. 多核并行是 AI Core 算力的来源.

Tiling 设计的 3 个核心参数

参数	含义	典型值
totalLength	输入张量总元素数	1024*1024 = 1048576
tileNum	切多少个 tile (= 核数)	32 - 64 (8 卡共 32*8 = 256 核)
tileLength	每个 tile 的元素数	16384 (1024*1024 / 64)
bufferCoefficient	UB 复用系数	2-3 (双缓冲 / 三缓冲)

Tiling 函数示例

// tiling 函数 (host 端调用, 算出每核处理多少数据)
void Tiling(TCubeTiling& tiling, int totalLength, int coreNum) {
    tiling.totalLength = totalLength;
    tiling.blockLength = 16384;  // 每个 tile 16384 元素
    tiling.tileNum = (totalLength + tiling.blockLength - 1) / tiling.blockLength;
    tiling.bufferCoefficient = 2;  // 双缓冲
    // ... 其他参数
}

关键洞察: Tiling 表写得好, 算子跑得快; 写得差, 算子比 torch 还慢. 经验值: 32 核以上用 16K-64K 的 tile 大小, bufferCoefficient = 2 (双缓冲) 通常最优, 三缓冲适合算 Compute 阶段的算子.

---

五、3 大官方样例 (Add / MatMul / TopK)

关键洞察: samples 仓的 6_ascendc_custom_op 提供了 Add / Add_tile / MatMul / TopK / kernel_template 5 个真实算子, 从最简单到复杂递进. 入门建议按 Add → Add_tile → TopK → MatMul 顺序.

样例对照

样例	难度	核心学到的
Add	入门	CopyIn/Compute/CopyOut 三段式完整流程, 单核版
Add_tile	入门	引入 Tiling, 多核并行版
TopK	中等	Reduce 算子, 非线性访存
MatMul	进阶	Cube 算子, 大矩阵优化
kernel_template	模板	算子工程脚手架 (CMakeLists.txt / build.sh)

运行样例的命令 (samples 仓提供)

# 1. 克隆 samples 仓
git clone https://gitcode.com/Ascend/samples.git
cd samples/cplusplus/level1_single_api/4_op_dev/6_ascendc_custom_op

# 2. 跑 Add 算子 (910 + AiCore + CPU 模拟)
(cd Add; bash run.sh add_custom ascend910 AiCore cpu)

# 3. 跑 Add 算子 (910 + AiCore + NPU 真机)
(cd Add; bash run.sh add_custom ascend910 AiCore npu)

# 4. 跑 TopK
(cd TopK; bash run.sh topk_custom ascend910 AiCore npu)

# 5. 跑 MatMul
(cd MatMul; bash run.sh matmul_custom ascend910 AiCore npu)

关键洞察: 命令格式 [KERNEL_NAME] [SOC_VERSION] [CORE_TYPE] [RUN_MODE]. SOC_VERSION 是 ascend910 / ascend310p; CORE_TYPE 是 AiCore / VectorCore; RUN_MODE 是 cpu (模拟) / npu (真机). 入门先 cpu 跑通逻辑, 再上 npu 测性能.

---

六、性能分析: msprof (算子级的 profiler)

关键洞察: AscendC 算子写完后, 性能分析是必经一步. msprof 是昇腾算子级 profiler, 关键指标是 PipeUtilization (流水线利用率) + Memory (内存使用).

msprof 常用命令 (在算子目录下)

# 1. 流水线利用率 (关键指标, 反映多核并行效率)
msprof --application="./add_custom_npu" \
       --output="./out" \
       --aic-metrics="PipeUtilization" \
       --sys-hardware-mem=on \
       --sys-io-profiling=on \
       --sys-interconnection-profiling=on

# 2. 内存使用 (UB / L1 占用)
msprof --application="./add_custom_npu" \
       --output="./out" \
       --aic-metrics="Memory"

# 3. UB 详细使用 (Local Memory 视角)
msprof --application="./add_custom_npu" \
       --output="./out" \
       --aic-metrics="MemoryUB"

3 大性能指标对照

指标	含义	健康值
PipeUtilization	流水线利用率	60-80% 算好, 80%+ 优秀, <40% 有瓶颈
Memory (L1/UB)	内存占用	接近 limit 时考虑 Tiling 调整
MemoryUB	UB 详细分配	用于找 UB 碎片

关键洞察: msprof 输出有 .csv 和可视化工具, 跑完看 PipeUtilization 是首要动作. 如果 < 40%, 多半是 Tiling 没切好, 或者 Compute 阶段太短, 或者 CopyIn/CopyOut 等待时间太长.

---

七、AscendC vs TBE 选型 (2026 现状)

关键洞察: 2026 年的新项目应该直接用 AscendC, 不用 TBE. TBE 是 CANN 5.x 主流, AscendC 是 CANN 6.0+ 主流. 老项目维护用 TBE 没问题, 新项目上 AscendC 收益更大.

选型对照

维度	TBE (老)	AscendC (新)
风格	DSL (Python 风格)	C++ 风格
适配 CANN	5.x 主流	6.0+ 主流
维护状态	维护中 (旧项目)	推荐 (新项目)
性能上限	中	高 (更细的控制)
上手难度	中 (DSL 学习)	高 (C++ + 硬件理解)
调试	TIK 调试器	msprof / 工具链成熟

立场: 2026 年新写算子选 AscendC, 2024 年前写的 TBE 算子继续维护. 不要混用 (同一个算子内不要 TBE + AscendC 拼接).

---

八、4 个常见踩坑 (AscendC 专属)

#	现象	根因	解决
1	“DataCopy 对齐错误”	LocalTensor 32B 没对齐	用 DataCopyPad 替代, 或加 padding
2	“UB 内存不足”	单 tile 太大, 超 UB 容量	缩小 tileLength 或加 bufferCoefficient
3	“npu 跑结果不对”	cpu 模拟通过但 npu 算错	多半是浮点精度问题, 加 FP32→FP16 转换保护
4	“Tiling 表算错, 算子卡死”	tileNum 算错, 部分核空跑	用 tiling.blockLength 严格向上取整

避坑心法: AscendC 的 80% 错误集中在 Tiling 表 和 数据对齐. 写完算子先用 cpu 跑通 (验证逻辑), 再上 npu 跑 (验证性能). 永远不要在 npu 上第一次跑就调 Tiling.

---

九、关联资源

资源	链接
samples 仓 (AscendC 6_ascendc_custom_op)	https://gitcode.com/Ascend/samples/tree/master/cplusplus/level1_single_api/4_op_dev/6_ascendc_custom_op
算子开发总览文档	https://www.hiascend.com/document (CANN 文档站 → 算子开发)
TBE → AscendC 迁移指南	samples 仓内 1_custom_op 对比 6_ascendc_custom_op
msprof 用户指南	https://www.hiascend.com/document (CANN 文档站 → 性能调优 → profiler)
社区会议	https://meeting.ascend.osinfra.cn/
Issue 反馈	https://gitcode.com/Ascend/samples/issues
关联博客	《【昇腾训练】-MindSpeed 开发入门》/ 《【昇腾推理】-MindIE 极速入门》/ 《【昇腾推理】-MindIE-SD 极速入门》

---

总结

AscendC = 昇腾 AI Core 专用算子开发语言, 不是 CUDA for 昇腾. 它用 CopyIn → Compute → CopyOut 三段式 + 显式 TBuf/TQue 内存管理 + 用户写 Tiling 函数实现多核并行, 是 2026 年新算子项目的推荐方案.

三个关键事实

1. 范式 = 三段式 + 显式数据搬运: 与 CUDA 完全不同, 数据从 Global 到 Local 必须显式 DataCopy / EnQue, 不是隐式访存
2. 性能 = Tiling 表写得好: 8 卡 910B3 的 80%+ 算力 = Tiling 设计 (切多少 tile / 多大 tile / 双缓冲系数) + 多核并行
3. 上手 = 5 步走通: 跑 Add (cpu) → 跑 Add_tile (cpu) → 跑 Add (npu) → msprof 看 PipeUtilization → 调 Tiling

一句话给新人: 别再把 AscendC 当 “用 C++ 写 PyTorch 算子” 了, 它是面向 AI Core 硬件的专用算子开发语言, 范式是显式数据搬运 + 三段式 + 多核切分. 入门从 samples 仓的 Add 算子跑起, 跑通 cpu 模拟再上 npu 真机.

---

参考

1. Ascend/samples 仓 - 2026-06-24
2. CANN 官方文档 - 算子开发章节
3. AscendC 编程指南 (CANN 文档站 → 算子开发)
4. TBE → AscendC 迁移说明 - 1_custom_op vs 6_ascendc_custom_op
5. 《【昇腾训练】-MindSpeed 开发入门》 - 算子之上是框架
6. 《【昇腾推理】-MindIE 极速入门》 - 算子之上是推理