Ascand C 学习笔记

让我们一起来探索这个充满可能性的领域吧！

待办事项

回答问题链接：点击这里
提交笔记链接：点击这里
香橙派实验手册：点击这里
AscendC算子中级认证考试链接：点击这里

异构计算架构CANN

背景

AI创新正加速发展：从“预测推断”到“内容生成”，我们正经历一场从感知理解世界到生成创造世界的变革。

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

昇腾基础架构介绍（CANN）

1. 华为在AI行业的投入与发展

2018.10：华为正式进军AI行业
2022.11：昇腾基础软硬件全面升级

华为致力于打造性能卓越、易于使用的全场景人工智能平台，推动AI行业的全面发展。

2. 昇腾AI基础软硬件架构的创新与优势

硬件与计算框架：昇腾与英伟达对比，支持主流AI框架（如PyTorch、TensorFlow），以及昇思MindSpore，CANN对标英伟达的CUDA。
自研达芬奇架构与高集成SoC设计：专为AI计算设计，具备最优配比的算力、内存和带宽，匹配AI应用场景，实现极致能效和面效比。

3. 昇腾AI异构计算架构CANN

通用计算：操作系统使能CPU硬件指令抽象编码及执行。
AI计算：CANN使能CPU与NPU协同编码，充分发挥NPU计算能力，通过CPU分配计算指令，实现NPU并行计算。

4. 昇腾CANN的功能与特点

处理器并行加速：支持大模型并行计算加速，支持原生开发（Ascend C）和生态迁移（GPU到NPU）。
算子设计与实现能力：从“追赶”到“自我超越”，持续突破。
高性能基础算子库：实现网络瞬时加速，开发核心融合算子库，涵盖多种计算需求。

5. 软硬协同优化

二进制算子库：提升动态shape能力，解决Python运行速度慢的问题，使用最优调度模板自动调用。

6. 昇腾图优化引擎

图优化引擎：优化串行指令为并行执行，提升计算效率；通过“工序合并”减少计算节点和数据交互，降低内存开销。
计算图执行下沉技术：将计算任务下放到硬件设备，实现算力均衡和高效调度。

7. 自动调优引擎AOE

持续升级：优化算子默认优化策略，实现零门槛低成本、零代码模型性能自动优化。

8. 端到端开发部署

快速构建AI应用：从算子到模型的开发部署，基于昇腾平台。

9. Ascend C算子编程语言

使能算子极简开发：遵循C/C++标准规范，提供自动化流水并行调度、结构化核函数编程、CPU/NPU孪生调试。

10. 开放的CANN开发体系

支持开发者：开放对接北向生态，提供丰富的融合算子和Runtime库，共同推动大模型创新生态。

11. 昇腾AI开发者工具

Atlas 200 DK A2开发者套件：提供开箱即用的开发者套件，支持AI应用创新和高校教学，售价1999元。
Orange Pi AIPro开发板：昇腾AI与香橙派联合发布，适合开发者使用。

Ascend C算子开发

架构说明

AI Core架构

逻辑架构

计算单元：包含三种基础计算资源：矩阵计算单元、向量计算单元、标量计算单元。
存储系统：包括张量缓冲器、统一缓冲区、标量缓冲区、L1缓冲区等。
控制单元：提供指令控制，相当于AI Core的司令部。

并行计算架构抽象

AI Core：NPU的内部单元。
Global Memory：位于NPU内部，AI Core外部。
Local Memory：位于AI Core内部。

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

Scalar计算单元
Cube计算单元
Vector计算单元

此外，还有存储单元和搬运单元：

异步指令流：Scalar计算单元读取指令序列并分发指令。
同步信号流：管理指令间的依赖关系。
计算数据流：DMA将数据搬运到Local Memory，完成计算后再次搬出。

算子

算子对应网络中层或节点的计算逻辑，广义地讲，对任何函数进行某一项操作都可以算一个算子。

算子基本概念

算子名称（Name）
算子类型（Type）
数据容器（Tensor）

Tensor

张量是存储算子输入和输出数据的容器，张量描述符（TensorDesc）描述张量的属性，包括名称、形状、数据类型和数据排布格式。在深度学习中，多维度数据通常用多维数据存储。

轴（Axis）

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

Ascend C算子

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

自定义算子的场景

将第三方框架转化为适配昇腾平台的离线模型时遇到不支持的算子。
在应用中设计某些数学运算，通过自定义算子实现。
迁移第三方框架的网络训练脚本到昇腾平台时遇到未定义行为。

Device模块

Device模块负责指定计算运行的设备，包含多个接口用于Device管理。

通过aclrtSetDevice接口指定计算使用的设备，分配计算资源。
在分配的Device上进行操作。
通过aclrtResetDevice接口释放设备。

核函数

核函数是算子设备侧的入口，直接在设备侧执行。指针入参变量统一为`__gm

__ uint8_t*`，指向Global Memory上的内存地址。

核函数调用语句格式：

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

blockDim：规定核函数在几个核上执行。
l2ctrl：保留参数，设置为nullptr。
stream：任务队列，管理任务并行。

SPMD模型

Ascend C算子编程采用SPMD模型，将数据拆分并分布在多个计算核心上运行。

编程范式

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

流程：搬入——计算——搬出。

流水任务

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

矢量编程流水线设计

逻辑位置：搬入（VECIN）、搬出（VECOUT）。
每次需要先EnQue，再DeQue。
VECIN和VECOUT采用双缓冲区设计，并行存取。

资源管理模块

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

临时变量

临时变量通过Pipe管理，使用TBuf数据结构申请空间。

Ascend C 进阶

Host侧开发

在Host侧定义切分策略
Host侧算子实现
- Tiling实现：计算数据切分过程的参数。
- Shape推导：推理算子的输出张量描述，避免动态内存分配。
- 算子原型注册：描述算子的输入、输出及算子在AI处理器上的实现信息。

矩阵编程

矩阵乘基础

MatMul：$C = A*B+Bias$

基础知识

分形格式：ND、NZ
TPosition

Matmul矩阵乘

步骤
- 创建Matmul对象
- 初始化
- 设置左矩阵A，右矩阵B，Bias
- 完成矩阵乘操作
- 结束矩阵乘操作

Matmul Api算法

尽量不要切k轴，避免多个核之间的数据依赖。
m k n
singleCorem singleCorek singleCore n
depth A1, depth B1——基本块的数量
stepM stepN

性能优化

理论性能评估

达成某个执行单元的bound并不代表算子性能最优。
如果算法没有重复计算过程，并且达到计算单元的bound，则认为算子性能最优。
如果搬运单元达到bound，并且算法达到搬运量最小，则认为算子性能最优。