Ascend C算子学习笔记

2024年鲲鹏&昇腾创新大赛集训营

todo：

• 回答问题链接：https://www.hiascend.com/forum/thread-0265155791127017253-1-1.html?fid=0163125572293226003
• 提交笔记链接：昇腾论坛
• 香橙派实验手册：昇腾论坛
• AscendC算子中级认证考试链接：微认证

异构计算架构CANN

背景

AI创新在加速：从“预测推断”向“内容生成”延展：感知理解世界-->生成创造世界

第一次浪潮：符号主义+逻辑推理 第二次浪潮：专家系统+机器学习 第三次浪潮：深度学习 第四次浪潮：智能内容生成（chatGPT）

昇腾基础架构介绍（CANN）

1. 深耕基础，厚植技术，华为坚定人工智能产业投入 2018.10 华为进军ai行业 2022.11 昇腾基础软硬件全面升级

2. 持续打造极致性能、极简易用的全场景人工智能平台 硬件、计算框架（CANN）、AI框架、应用使能、行业应用...... 华为全流程推动ai行业发展

3. 昇腾AI基础软硬件架构全面创新，使能行业场景化应用 昇腾 vs 英伟达 昇腾除传统pyTorch/TF等，还支持昇思MS 昇腾CANN对标英伟达CUDA

4. 昇腾AI处理器：自研达芬奇架构+高集成SoC设计，实现极致能效与面效比

   • 达芬奇架构，为AI计算而生

   • 算力、内存、带宽 最优配比，匹配AI应用场景

   • CPU+NPU+网卡+编解码 高级程度SoC，极致能效

5. 昇腾AI异构计算架构CANN，使能NPU高性能计算

   • 通用计算：操作系统使能CPU硬件指令抽象编码及执行

   • AI计算：CANN使能CPU与NPU协同编码，发挥NPU计算能力 NPU并行计算，通过CPU为NPU分配计算指令

6. 昇腾CANN：向下使能处理器并行加速，向上使能高效开发

   1. 使能大模型并行计算加速

   2. 高效原生开发与生态迁移 原生开发：Ascend C 生态迁移：GPU-->NPU

   3. 全面开放，生态兼容

7. CANN：持续突破算子设计与实现能力，从“追赶”走向“自我超越”

8. 提供深度优化的高性能基础算子库，实现网络瞬时加速 大模型时代，开发核心融合算子库，例：FlashAttention类 算子库覆盖融合算子、神经网络算子、视频/图片处理算子、图像更改算子、线性代数算子、分布式训练算子等

9. 软硬协同优化技术 使能NPU高性能计算

10. 二进制算子库提升动态shape能力，夯实成熟场景，突破新场景 破除python运行速度慢特点 使用算子二进制Kernel包，包含最优调度模板，自动调用

11. 昇腾图优化引擎 极致图编译技术 充分释放澎湃算力

    • 可将原本的串行指令优化为并行执行，实现计算效率提升

    • 保持精度前提下，通过“工序合并”减少计算节点/数据交互等，有效减少中间数据搬入搬出的时间及内存开销

12. 计算图执行下沉技术实现处理器侧计算闭环，提升计算性能 下沉：将计算下放到硬件设备上 图编译过程中切分计算类型，实现算力均衡 在硬件层面进行计算任务调度，一次下发，无限次执行，消除调度开销

13. 自动调优引擎AOE持续升级 实现性能效率双效提升 优化算子默认优化策略

    • 零门槛低成本

    • 零代码实现模型性能能自动优化

    • 调优效果显著

14. 端到端使能从算子到模型的开发部署，快速构建基于昇腾平台的AI应用

15. Ascend C算子变成语言，使能算子极简开发

    • 遵循C/C++标准规范

    • 自动化流水并行调度

    • 结构化核函数变成

    • CPU/NPU孪生调试

16. CANN开发体系持续开放，使能开发者，共促大模型创新生态

    • 开放对接北向生态

    • 开放丰富的融合算子

    • 开放Runtime库

17. 面向开发者，提供开箱即用的昇腾AI开发者套件Atlas 2001 DK A2 Atlas 2001 DK A2开发者套件：AI应用创新、高校教学 1999￥一个（）

18. 昇腾AI + 香橙派 联合发布Orange Pi AIPro开发板 好东西 买就完了（真心话）

Ascend C算子开发

架构说明

AI Core架构

逻辑架构

• 计算单元：包含三种基础计算资源（矩阵计算单元、向量计算单元、标量计算单元）

• 存储系统：AI Core的片上存储单元和相应的数据通路构成了存储系统（其中缓冲区有张量缓冲器L0C、统一缓冲区、标量缓冲区、L1缓冲区/输入缓冲区、张量缓冲区等）

• 控制单元：整个计算过程提供了指令控制，相当于AI Core的司令部，负责整个AI Core的运行

并行计算架构抽象

AI Core是NPU的一个内部单元

Global Memory在NPU内部，AI Core外部

Local Memory在AI Core内部

计算单元包括了三种基础计算资源：

1. Scalar计算单元

2. Cube计算单元

3. Vector计算单元

此外还有存储单元、搬运单元

• 异步指令流 Scalar计算单元读取指令序列，将指令分发

• 同步信号流 指令间可能存在依赖关系，有Scalar计算单元给对应单元下发同步指令

• 计算数据流 DMA将数据搬运到Local Memory，数据计算好后(Vector/Cube)，再次搬出

算子

算子对应网络中层或者节点的计算逻辑

广义地讲：对任何函数进行某一项操作都可以算一个算子

算子基本概念

• 算子名称（Name）

• 算子类型（Type）

• 数据容器（Tensor）

Tensor

张量是存储算子输入数据与输出数据的容器，张量描述符（TensorDesc）是对输入数据与输出数据的描述，张量描述符的数据结构包含一下属性：

• 名称（name）

• 形状（shape）

• 数据类型（dtype）

• 数据排布格式（format） 在深度学习中，多维度数据通过多维数据存储，比如卷积神经网络的特征图通常用四维数据保存，即NHWC

轴（Axis）

axis代表张量中纬度的下标，可以理解为展开n次后的数据

Ascend C算子

Ascend C是CANN针对算子开发场景推出的编程语言

需要自定义算子的场景

• 推理场景下，将第三方框架转化为适配昇腾平台的离线模型时遇到了不支持的算子

• 推理场景下，应用中设计某些数学运算，开发者可以将这些操作通过自定义算子的方式进行实现

• 训练场景下，将第三方框架的网络训练脚本迁移到昇腾平台时遇到未定义行为

Device模块

Device模块负责指定计算运行的真是设备，包含多个aclrtSetDevice、aclrtResetDevice、aclrtGetDevice等众多运行时接口用于进行Device管理

流程主要为：

1. 通过aclrtSetDevice接口指定计算使用的设备，设置成功后分配Device的计算资源

2. 在分配的Device上使用计算资源进行相关操作

3. 计算任务结束后通过aclrtResetDevice接口在计算结束后释放设备

核函数

核函数（Kernel Function）是算子设备测侧的入口，是主机测和设备测链接的桥梁

核函数是直接在设备测执行的代码，在核函数中，需要为在一个核上执行的代码规定要进行的数据访问和计算操作

为了方便：指针入参变量统一的类型定义为__gm__ uint8_t*，参数列表中的变量类型限定符__gm__，表明该指针变量指向Global Memory上某处内存地址

用户可统一使用uint8_t 的类型的指针，并在使用时转化为实际的指针类型，也可以直接传入实际的指针类型

核函数的调用语句使用内核调用符<<<...>>>这种形式，如

 kernel_name<<<blockDim, l2ctrl, stream>>>(argument list);

• blockDim：规定了核函数会在几个核上执行，每个核会有一个逻辑ID，变现为内置变量block_idx，从0开始，可用GetBlockIdx()获得

• l2ctrl：保留参数，暂时设置为固定值nullptr

• stream：类型为aclrtStream，是一个任务队列，应用程序用通过stream来管理任务的并行

SPMD模型

Ascend C算子编程是SPMD的变成，将需要处理的数据拆分并分布在多个计算核心上运行

多个AI Core共享相同的指令代码，每个核上的运行实例的唯一区别是block_idx不同

block类似于进程，block_idx就是表示进程唯一性的进程ID

编程范式

抽象编程模型“TPIPE并行计算”

搬入——计算——搬出

流水任务

流水任务指的是单核处理程序中主程序调度的并行任务，可以通过流水任务实现数据的并行处理来提升性能

需要被处理的数据分成n份，使用Progress1~n表示，同一时间点，可以有多个流水任务Stage并行处理

矢量编程流水线设计

逻辑位置（QuePosition）搬入：VECIN、搬出：VECOUT

每次需要先EnQue，再DeQue

VECIN和VECOUT采用双缓冲区设计，并行存取

资源管理模块

任务间数据传递使用到的内存统一由内存管理模块Pipe进行管理

临时变量

临时变量同样通过Pipe进行管理，临时变量可以使用TBuf数据结构来申请空间，无法使用Que结构