初始AscendC

本文记录笔者在杭研学习Ascend C的所得和体会。本次学习AscendC主要涵盖从基础概念、产品定位到AscendC算子开发。

AI Core架构

物理结构

昇腾AI处理器架构图如下：

  • AI Core是昇腾AI处理器的核心，采用达芬奇架构。
  • 达芬奇架构的主要部分有：
    • 计算单元：包含矩阵计算单元Cube、向量计算单元Vector、标量计算单元Scalar
    • 存储系统
    • 控制单元

逻辑结构

逻辑架构抽象：SIMD架构（单指令多数据计算 Single in struction, Multiple Data）

  • 计算单元：
    • Scalar：地址计算；循环控制等；把向量计算、矩阵计算、数据搬运、同步指令发射给对应单元。
    • Vector：向量计算。
    • Cube：矩阵计算。
  • 存储抽象：Local Memory。PS，外部内存统称为Global Memory。
  • 搬运单元：负责再Local和Global之间搬运。包括MTE2（Mem Transfer Engine）搬入、MTE3搬出

异步计算过程：（指令流）

  • 读取指令序列
  • 发射指令到对应单元
  • 并行执行指令
    内存搬运过程：（数据流）
  • 数据输入到Local Mem
  • 完成计算，回写Local Mem
  • 数据搬出到Global Mem

AscendC编程对象

外部存储：GlobalTensor，用于存放全局数据

  • 原型：
  • 案例：

内部存储：LocalTensor，用于存放内部存储数据。

  • 原型：
  • 案例：

逻辑位置

  • AscendC使用Queue队列完成任务之间的数据通信和同步。
  • Queue使用逻辑位置来表达各个存储级别，TPosition包括：详见下表
    • 用于向量编程：VECIN、VECCALC、VECOUT
    • 用于矩阵编程：A1、A2、B1、B2、CO1、CO2

Vector算子的开发流程（静态形状）

算子开发流程

算子的开发遵循一定的开发流程，一般来说算子的开发流程如下：

  • 算子分析：根据表达式、输入、输出等明确函数接口。
    • 明确算子数学表达式和计算逻辑。对于Add算子而言，表达式就是z=x+y。计算逻辑就是先搬到片上存储，然后使用计算接口完成两个加法运算，再搬到外部存储。
    • 输入输出。Add而言，输入输出类型时half，排布是ND，shape是固定的[8,2048]
    • 函数名和参数。add_custom为函数名，参数是x,y,z，在global mem上。
  • 核函数定义：定义AscendC入口函数。包括输入数据的shape、数据类型、数据排布的format、核函数名字等。
    • 即extern "C" __global__ __aicore__ void add_custom()这部分的代码。
    • 先实例化KernelAdd op算子类。
    • 调用Init方法完成内部初始化。
    • 调用Process方法完成核心逻辑。
  • 根据编程范式实现算子类：完成内部实现。包括：
    • CopyIn
    • Compute
    • CopyOut

值得注意，算子的所需接口可以参考官方文档。由于没接触过算子开发，刚开始我还以为要用cmath这种库。

例如，ADD算子，就需要用到：

  • 数据搬移接口：DataCopy
  • 矢量双目加法：Add
  • 队列：EnQue、DeQue

算子开发编程范式

Ascend C的编程范式：

  • 把算子内部处理程序分成多个流水任务（Stage）：指的是单核内部并行。
    • 示意图：拆分为三个流水任务
    • 矢量编程范式的3个基本任务：CopyIn、CopyOut、Compute，对应三个Stage
  • 数据载体：张量Tensor
  • 任务间的通信和同步：队列Queue
  • 管理任务间的通信内存：内存管理模块Pipe

SPMD模型：将数据拆分到多个核上运行。多个AI Core共享代码，用block_idx区分。

Add算子实现

三个任务的设计：

  • 第一步，CopyIn：就是将全局内存的张量xGm和yGm搬运为xLocal，yLocal。
  • 中间：通过VECIN队列inQeueX和inQueueY通信和同步。
  • 第二步，Compute：zLocal = xLocal+yLocal
  • 中间：通过VECOUT队列outQueueZ进行通信和同步。
  • 第三步，Copyout：将zLocal搬运到zGm

算子类的实现：

kernelAdd算子类包括：

  • public：Init函数和Process函数。
  • private：CopyIn、CopyOut、Compute三个私有函数。
  • 私有变量：TPipe用于管理内存、TQue<QuePosition::VECIN,BUFFER_NUM>队列，还有一个VECOUT的队列，以及GlobalTensor<half>xGm等。

Init函数实现：

主要是实现多核并行计算的数据切片，以及单核处理函数的数据切块。

Process函数的实现

迭代，依次完成CopyIn、Compute和Copyout。

本次实验的工程目录介绍

本次ADD算子的sample代码工程如下：

  • script：用于生成真值的脚本和校验结果的脚本。
  • cmake：包含CPU和NPU模式下的cmake脚本。
  • CMakeLists.txt
  • add_custom.cpp：算子实现文件
  • main.cpp：调用程序源码。
  • run.sh：编译运行算子的脚本。

后续又完成了sinh算子的开发，就是再此文件目录的基础上，修改script中的校验脚本并修改算子实现文件的compute部分代码即可实现。

AsendC Host侧实现

本部分内容有利于理解和完成昇腾算子开发中级微认证。

Host设备

Host与Device的区别：

  • Host设备：X86服务器、ARM服务器等。利用Device计算能力完成任务。
  • Device设备：昇腾AI处理器的硬件设备。通过PCle连接Host。提供计算。

Host算子实现：

  • Tiling实现：计算数据切分过程的相关参数。
  • Shape推导：根据输入张量的shape、类型和排布格式等推算输出张量的描述。推算好计算图的所有数据格式，避免动态内存分配的开销。
  • 算子原型注册
  • 算子类注册

工程上：

  • add_custom.cpp：Host侧实现
  • add_custom_tiling.h：Tiling结构体声明

Tiling

  • 背景：Local Mem无法完全容纳算子输入输出的所有数据，需要每次搬运一块数据进行计算，最后聚合。
  • 概念：这个数据切分、分块计算的过程叫做Tiling实现。
    • Tiling块：每次搬运的数据块。
    • Tiling算法：确定基本块大小的算法。
    • Tiling结构体：承载Tiling算法信息的数据结构。
    • Tiling函数：实现Tiling算法并将Tiling结构体下发给Kernel。

值得注意的是，Tiling实现中完成的均为标量计算，AI Core并不擅长，一般在Host侧执行。这也是调优关注的点，尽量避免标量运算在AI Core上进行。

在核函数内使用Tiling的信息：

extern "C" __global__ __aicore__ void add_custom(GM_AddR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling){
	GET_TILING_DATA(tiling_Data,tiling); //拆包
	KernelAdd op;
	op.Init(x,y,z,tailing_data.totalLength,tiling_data.tileNum);
	op.Process();
}

固定Shape和动态Shape的代码文件区别，在微认证中需要使用动态Shape。

Shape

计算图不可能一步一步调用，而是需要预先推导出每个Tensor Shape和 data type，避免动态内存分配的开销。

算子原型注册

算子原型注册描述了算子的输入输出等信息，关联Tiling实现和Shape推导函数。