扶稳坐好,三分钟带你了解Ascend C
——算子开发与模型调优部署的利器
目录:
- 基本知识
- Asend C
- 算子基本概念
- 自定义算子
- 核函数
- Kernel侧算子
- Host侧算子
- 算子优化
- 基本知识
昇腾AI处理器:
采用自研达芬奇架构+高集成的SoC设计
实现性能、功耗、成本的三维均衡

昇腾AI异构计算架构CANN:
NPU负责调度;CPU负责同步信息。
特点/优势:
支持融合算子库,支持cpu生态向npu高效迁移,Ascend C支持算子极简开发,图编译加速技术(串行->并行;算子融合;计算图下沉:计算图分割子图,交付相应的芯片)
概述:
Ascend C开发的算子运行在AI Core上,AI Core是昇腾NPU硬件平台的计算核心,NPU内部有多个AI Core。多个AI Core共享相同的指令代码,每个核上的运行实例唯一的区别是block_idx不同,开发者只需要关注单核上的处理程序——核函数
Ascend C提供流水线式的编程范式,基于编程范式可以快速搭建算子实现的代码框架,实现流水并行。Ascend C把算子核内的处理程序,分成多个流水任务:“搬入、计算、搬出”,通过队列完成任务间通信和同步,并通过统一的内存管理模块管理任务间通信内存。
关键部分:
AI Core(升腾AI处理器计算核心):
计算单元:标量、向量、矩阵。
存储系统:多层级
控制单元:提供指令控制,并行控制时间同步。

什么是算子:
算子名称:一个算子的表示,用于定位算子出现的问题
算子类型:如卷积层(Convolution Layer)是一个算子;全连接层(Full-connected Layer, FC layer)中的权值求和过程,是一个算子
数据容器(张量Tensor):存储输入/输出数据的容器。名称、形状、数据类型、数据排布格式。

数据排布格式:怎么存决定怎么取,不同存储方式实现不同功能。
在昇腾AI处理器中,为了提高数据的访问效率,张量数据采用NC1HWC0的五维格式。其中C0与微架构强相关,等于AI Core中矩阵计算单元的大小,这部分数据需要连续存储;C1是将C维度按照C0进行拆分后的数目,即C1=C/C0。如果不整除,最后一份数据需要补齐以对齐C0。
优势:
基于C/C++编程
屏蔽硬件差异
多层级API封装
孪生调试:CPU可以模拟NPU实现
Host:服务器/主机。Host memary:主机内存。
Device模块:AI加速卡。Device memary:卡内内存。
Local memary核内内存。Globle memary核外内存。
调用核函数:
kernel_name<<<blockDim(几个核上进行),12ctrl(保留参数),stream(任务队列)>>>(argument list)
blockDim:几个核上进行
12ctrl:保留参数
stream任务队列
Init():数据偏移
process():执行命令
Tiling实现:计算数据切分的参数。
Shape推导:算子构图阶段可以提前静态申请内存。
算子原型注册。
MatMul 计算公式:C=A*B+Bias 、矩阵切分操作、原理
具体切分原理可见:MindSpore
考察两个指标:计算&搬运
算子构成:数学公式、计算机算法、编程
~Tiling:关键是如何切分
~Scalar优化:减少标量运算,如用Host中计算,Tiling搬运到Device
~0级API指令:十分灵活,但是参数多
~Shape 32B对齐.
扶稳坐好,三分钟带你了解Ascend C
——算子开发与模型调优部署的利器
目录:
昇腾AI处理器:
采用自研达芬奇架构+高集成的SoC设计
实现性能、功耗、成本的三维均衡
昇腾AI异构计算架构CANN:
NPU负责调度;CPU负责同步信息。
特点/优势:
支持融合算子库,支持cpu生态向npu高效迁移,Ascend C支持算子极简开发,图编译加速技术(串行->并行;算子融合;计算图下沉:计算图分割子图,交付相应的芯片)
概述:
Ascend C开发的算子运行在AI Core上,AI Core是昇腾NPU硬件平台的计算核心,NPU内部有多个AI Core。多个AI Core共享相同的指令代码,每个核上的运行实例唯一的区别是block_idx不同,开发者只需要关注单核上的处理程序——核函数
Ascend C提供流水线式的编程范式,基于编程范式可以快速搭建算子实现的代码框架,实现流水并行。Ascend C把算子核内的处理程序,分成多个流水任务:“搬入、计算、搬出”,通过队列完成任务间通信和同步,并通过统一的内存管理模块管理任务间通信内存。
关键部分:
AI Core(升腾AI处理器计算核心):
计算单元:标量、向量、矩阵。
存储系统:多层级
控制单元:提供指令控制,并行控制时间同步。
什么是算子:
算子名称:一个算子的表示,用于定位算子出现的问题
算子类型:如卷积层(Convolution Layer)是一个算子;全连接层(Full-connected Layer, FC layer)中的权值求和过程,是一个算子
数据容器(张量Tensor):存储输入/输出数据的容器。名称、形状、数据类型、数据排布格式。
数据排布格式:怎么存决定怎么取,不同存储方式实现不同功能。
在昇腾AI处理器中,为了提高数据的访问效率,张量数据采用NC1HWC0的五维格式。其中C0与微架构强相关,等于AI Core中矩阵计算单元的大小,这部分数据需要连续存储;C1是将C维度按照C0进行拆分后的数目,即C1=C/C0。如果不整除,最后一份数据需要补齐以对齐C0。
优势:
基于C/C++编程
屏蔽硬件差异
多层级API封装
孪生调试:CPU可以模拟NPU实现
Host:服务器/主机。Host memary:主机内存。
Device模块:AI加速卡。Device memary:卡内内存。
Local memary核内内存。Globle memary核外内存。
调用核函数:
kernel_name<<<blockDim(几个核上进行),12ctrl(保留参数),stream(任务队列)>>>(argument list)
blockDim:几个核上进行
12ctrl:保留参数
stream任务队列
Init():数据偏移
process():执行命令
Tiling实现:计算数据切分的参数。
Shape推导:算子构图阶段可以提前静态申请内存。
算子原型注册。
MatMul 计算公式:C=A*B+Bias 、矩阵切分操作、原理
具体切分原理可见:MindSpore
考察两个指标:计算&搬运
算子构成:数学公式、计算机算法、编程
~Tiling:关键是如何切分
~Scalar优化:减少标量运算,如用Host中计算,Tiling搬运到Device
~0级API指令:十分灵活,但是参数多
~Shape 32B对齐.