【CANN】-Ascend C 算子开发

【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 专用算子开发语言维度AscendCCUDA CTBE (老方案)出身华为官方 (CANN 6.0)NVIDIA 官方华为官方 (CANN 5.x 主流)目标硬件昇腾 AI Core / Vector CoreNVIDIA 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) 输出 GlobalTensor3 段对应 3 个核心 API阶段核心 API作用CopyInDataCopy/DataCopyPad全局 → 本地 (含 32B 对齐)ComputeAdd/Mul/ReduceSum等内置指令本地计算 (在 AI Core 上)CopyOutDataCopy/EnQue/DeQue本地 → 全局对比 CUDA 写 kernel 的差异步骤CUDAAscendC数据读取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 类数据结构对照数据结构存储位置生命周期何时用GlobalTensorGlobal Memory (HBM)整个算子输入 / 输出张量LocalTensorLocal Memory (UB / L1)一个核内计算的中间结果TBufUnified Buffer 分配器静态分配申请 Local Memory 空间TQue队列 (双缓冲)动态队列CopyIn / Compute / CopyOut 之间的同步典型用法示例 (Add 算子)// 1. 准备 TBuf (申请 UB 空间)TBufQuePosition::VECIN,TPosition::LCMbuf1,buf2,buf3;buf1.InitBuffer(que1,totalLength*sizeof(TYPE));buf2.InitBuffer(que2,totalLength*sizeof(TYPE));buf3.InitBuffer(que3,totalLength*sizeof(TYPE));// 2. CopyIn: Global → LocalLocalTensorTYPEaLocalque1.AllocTensorTYPE();LocalTensorTYPEbLocalque2.AllocTensorTYPE();DataCopy(aLocal,aGlobal,totalLength);DataCopy(bLocal,bGlobal,totalLength);que1.EnQue(aLocal);que2.EnQue(bLocal);// 3. Compute: Local 上做加法LocalTensorTYPEaLocal2que1.DeQueTYPE();LocalTensorTYPEbLocal2que2.DeQueTYPE();LocalTensorTYPEcLocalque3.AllocTensorTYPE();Add(cLocal,aLocal2,bLocal2,totalLength);que1.FreeTensor(aLocal2);que2.FreeTensor(bLocal2);// 4. CopyOut: Local → GlobalDataCopy(cGlobal,cLocal,totalLength);que3.EnQue(cLocal);LocalTensorTYPEcLocal2que3.DeQueTYPE();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 1048576tileNum切多少个 tile ( 核数)32 - 64 (8 卡共 32*8 256 核)tileLength每个 tile 的元素数16384 (1024*1024 / 64)bufferCoefficientUB 复用系数2-3 (双缓冲 / 三缓冲)Tiling 函数示例// tiling 函数 (host 端调用, 算出每核处理多少数据)voidTiling(TCubeTilingtiling,inttotalLength,intcoreNum){tiling.totalLengthtotalLength;tiling.blockLength16384;// 每个 tile 16384 元素tiling.tileNum(totalLengthtiling.blockLength-1)/tiling.blockLength;tiling.bufferCoefficient2;// 双缓冲// ... 其他参数}关键洞察: 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 仓gitclone https://gitcode.com/Ascend/samples.gitcdsamples/cplusplus/level1_single_api/4_op_dev/6_ascendc_custom_op# 2. 跑 Add 算子 (910 AiCore CPU 模拟)(cd Add;bashrun.sh add_custom ascend910 AiCore cpu)# 3. 跑 Add 算子 (910 AiCore NPU 真机)(cd Add;bashrun.sh add_custom ascend910 AiCore npu)# 4. 跑 TopK(cd TopK;bashrun.sh topk_custom ascend910 AiCore npu)# 5. 跑 MatMul(cd MatMul;bashrun.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-metricsPipeUtilization\--sys-hardware-memon\--sys-io-profilingon\--sys-interconnection-profilingon# 2. 内存使用 (UB / L1 占用)msprof--application./add_custom_npu\--output./out\--aic-metricsMemory# 3. UB 详细使用 (Local Memory 视角)msprof--application./add_custom_npu\--output./out\--aic-metricsMemoryUB3 大性能指标对照指标含义健康值PipeUtilization流水线利用率60-80% 算好, 80% 优秀, 40% 有瓶颈Memory (L1/UB)内存占用接近 limit 时考虑 Tiling 调整MemoryUBUB 详细分配用于找 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 风格适配 CANN5.x 主流6.0 主流维护状态维护中 (旧项目)推荐 (新项目)性能上限中高 (更细的控制)上手难度中 (DSL 学习)高 (C 硬件理解)调试TIK 调试器msprof / 工具链成熟立场: 2026 年新写算子选 AscendC, 2024 年前写的 TBE 算子继续维护. 不要混用 (同一个算子内不要 TBE AscendC 拼接).八、4 个常见踩坑 (AscendC 专属)#现象根因解决1“DataCopy 对齐错误”LocalTensor 32B 没对齐用DataCopyPad替代, 或加 padding2“UB 内存不足”单 tile 太大, 超 UB 容量缩小 tileLength 或加 bufferCoefficient3“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_opmsprof 用户指南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 年新算子项目的推荐方案.三个关键事实范式 三段式 显式数据搬运: 与 CUDA 完全不同, 数据从 Global 到 Local 必须显式DataCopy/EnQue, 不是隐式访存性能 Tiling 表写得好: 8 卡 910B3 的 80% 算力 Tiling 设计 (切多少 tile / 多大 tile / 双缓冲系数) 多核并行上手 5 步走通: 跑 Add (cpu) → 跑 Add_tile (cpu) → 跑 Add (npu) → msprof 看 PipeUtilization → 调 Tiling一句话给新人: 别再把 AscendC 当 “用 C 写 PyTorch 算子” 了, 它是面向 AI Core 硬件的专用算子开发语言, 范式是显式数据搬运 三段式 多核切分. 入门从 samples 仓的 Add 算子跑起, 跑通 cpu 模拟再上 npu 真机.参考Ascend/samples 仓 - 2026-06-24CANN 官方文档 - 算子开发章节AscendC 编程指南 (CANN 文档站 → 算子开发)TBE → AscendC 迁移说明 - 1_custom_op vs 6_ascendc_custom_op《【昇腾训练】-MindSpeed 开发入门》 - 算子之上是框架《【昇腾推理】-MindIE 极速入门》 - 算子之上是推理