3 天上手 Ascend C 编程丨通过 Ascend C 编程范式实现一个算子实例

2023-09-11
广东
本文字数：3760 字
阅读完需：约 12 分钟

本文分享自华为云社区《3天上手Ascend C编程 | Day2 通过Ascend C编程范式实现一个算子实例》，作者：昇腾 CANN 。

一、Ascend C 编程范式

Ascend C 编程范式把算子内部的处理程序，分成多个流水任务( stage )，以张量( Tensor)为数据载体，以队列 ( Queue ) 进行任务之间的通信与同步，以内存管理模块( Pipe ) 管理任务间的通信内存。

1、流水任务

流水任务指的是单核处理程序中主程序调度的并行任务。在核函数内部，可以通过流水任务实现数据的并行处理，进一步提升性能。下面举例来说明，流水任务如何进行并行调度。以下面的流水任务示意图为例，单核处理程序的功能被拆分成 3 个流水任务：Stage1、Stage2、Stage3，每个任务专注于完成单一功能；需要处理的数据被切分成 n 片，使用 Progress1~n 表示，每个任务需要依次完成 n 个数据切片的处理。Stage 间的箭头表达数据间的依赖关系，比如 Stage1 处理完 Progress1 之后，Stage2 才能对 Progress1 进行处理。

若 n=3，即待处理的数据被切分成 3 片，则上图中的流水任务运行起来的示意图如下，从运行图中可以看出，对于同一片数据，Stage1、Stage2、Stage3 之间的处理具有依赖关系，需要串行处理；不同的数据切片，同一时间点，可以有多个任务在并行处理，由此达到任务并行、提升性能的目的。

矢量（Vector）编程范式把算子的实现流程分为 3 个基本任务：CopyIn，Compute，CopyOut。CopyIn 负责搬入操作，Compute 负责矢量计算操作，CopyOut 负责搬出操作。

2、任务间通信与同步

不同的流水任务之间存在数据依赖，需要进行数据传递。Ascend C 中使用 Queue 队列完成任务之间的数据通信和同步，提供 EnQue、DeQue 等基础 API。Queue 队列管理 NPU 上不同层级的物理内存时，用一种抽象的逻辑位置（QuePosition）来表达各级别的存储，代替了片上物理存储的概念，开发者无需感知硬件架构。

矢量编程中使用到的逻辑位置（QuePosition）定义如下：

搬入数据的存放位置：VECIN 搬出数据的存放位置：VECOUT 矢量编程主要分为 CopyIn、Compute、CopyOut 三个任务：

CopyIn 任务中将输入数据从 Global 内存搬运至 Local 内存后，需要使用 EnQue 将 LocalTensor 放入 VECIN 的 Queue 中；

Compute 任务等待 VECIN 的 Queue 中 LocalTensor 出队之后才可以完成矢量计算，计算完成后使用 EnQue 将计算结果 LocalTensor 放入到 VECOUT 的 Queue 中；

CopyOut 任务等待 VECOUT 的 Queue 中 LocalTensor 出队，再将其拷贝到 Global 内存。

Ascend C 使用 GlobalTensor 和 LocalTensor 作为数据的基本操作单元，它是各种指令 API 直接调用的对象，也是数据的载体。

3、内存管理机制

任务间数据传递使用到的内存统一由内存管理模块 Pipe 进行管理。Pipe 作为片上内存管理者，通过 InitBuffer 接口对外提供 Queue 内存初始化功能，开发者可以通过该接口为指定的 Queue 分配内存。

Queue 队列内存初始化完成后，需要使用内存时，通过调用 AllocTensor 来为 LocalTensor 分配内存，当创建的 LocalTensor 完成相关计算无需再使用时，再调用 FreeTensor 来回收 LocalTensor 的内存。

编程过程中使用到的临时变量内存同样通过 Pipe 进行管理。临时变量可以使用 TBuf 数据结构来申请指定 QuePosition 上的存储空间，并使用 Get()来将分配到的存储空间分配给新的 LocaLTensor 从 TBuf 上获取全部长度，或者获取指定长度的 LocalTensor。

使用 TBuf 申请的内存空间只能参与计算，无法执行 Queue 队列的入队出队操作。

二、使用 Ascend C 编程范式实现一个算子实例

矢量算子开发一般开发流程如下：

下面以 add 作为例子介绍 Ascend C 矢量算子的开发流程。完整样例大家可以参考代码样例。

1、算子分析

分析算子的数学表达式、输入、输出以及计算逻辑的实现，明确需要调用的 Ascend C 接口。

例子以 Add 算子为例，数学公式:z= x+y，为简单起见，设定输入张量 x, y,z 为固定 shape(8,2048)，数据类型 dtype 为 half 类型，数据排布类型 format 为 ND，核函数名称为 add_custom。

算子的数学表达式及计算逻辑。Add 算子的数学表达式为:z = x + y；计算逻辑：输入数据需要先搬入到片上存储，然后使用计算接口完成两个加法运算，得到最终结果，再搬出到外部存储。
输入和输出。Add 算子有两个输入:x 与 y，输出为 z。输入数据类型为 half，输出数据类型与输入数据类型相同。输入支持固定 shape(8，2048)输出 shape 与输入 shape 相同，输入数据排布类型为 ND。
确定核函数名称和参数。自定义核函数名，如 add_custom。根据输入输出，确定核函数有 3 个入参 x，y，z。x，y 为输入在 Global Memory 上的内存地址，z 为输出在 Global Memory 上的内存地址。
确定算子实现所需接口。涉及内外部存储间的数据搬运，使用数据搬移接口:Datacopy 实现；涉及矢量计算的加法操作，使用矢量双目指令:Add 实现；使用到 LocalTensor，使用 Queue 队列管理，会使用到 EnQue、DeQue 等接口。

2、核函数定义

在 add_custom 核函数的实现中实例化 kernelAdd 算子类，调用 Init()数完成内存初始化，调用 Process()函数完成核心逻辑。

// 实现核函数extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z){    // 初始化算子类，算子类提供算子初始化和核心处理等方法    KernelAdd op;    // 初始化函数，获取该核函数需要处理的输入输出地址，同时完成必要的内存初始化工作    op.Init(x, y, z);    // 核心处理函数，完成算子的数据搬运与计算等核心逻辑    op.Process();}

复制代码

3、根据矢量编程范式实现算子类

根据前面的知识，算子实现三个流水任务 CopyIn、Compute、CopyOut。任务间通过队列 VECIN、VECOUT 进行通信和同步，由 pipe 内存管理对象对任务间交互使用到的内存、临时变量使用到的内存统一进行管理。如下图所示：

CopyIn 任务:将 Global Memory 上的输入 Tensor xGm 和 yGm 搬运至 Local Memory，分别存储在 xLocal,yLocal；
Compute 任务:对 xLocal,yLocal 执行加法操作，计算结果存储在 zLocal 中；
CopyOut 任务:将输出数据从 zLocal 搬运至 Global Memory 上的输出 Tensor zGm 中

CopyIn，Compute 任务间通过 VECIN 队列 inQueuex，inQueuer 进行通信和同步；compute，copyout 任务间通过 VECOUT 队列 outQueuez 进行通信和同步。

第一步，我们进行算子类定义：

第二步，实现 Init()函数：多核并行+单核处理数据：

第三步，实现 Process()函数——CopyIn，Compute、CopyOut 三个流水任务：

第四步，通过 double buffer 机制进一步提升性能

double buffer 通过将数据搬运与矢量计算并行执行以隐藏数据搬运时间并降低矢量指令的等待时间，最终提高矢量计算单元的利用效率 1 个 Tensor 同一时间只能进行搬入、计算和搬出三个流水任务中的一个，其他两个流水任务涉及的硬件单元则处于 ldle 状态如果将待处理的数据一分为二，比如 Tensor1、Tensor2：

当矢量计算单元对 Tensor1 进行 Compute 时，Tensor2 可以执行 CopvIn 的任务

当矢量计算单元对 Tensor2 进行 Compute 时，Tensor1 可以执行 CopyOut 的任务

当矢量计算单元对 Tensor2 进行 CopyOut 时，Tensor1 可以执行 CopyIn 的任务。由此，数据的进出搬运和矢量计算之间实现并行，硬件单元闲置问题得以有效缓解

最后，基于内核调用符方式进行算子验证

先使用 python 脚本生成 x,y，并计算出 z（golden）并落盘。然后再用相同的 x，y，在 cpu 和 npu 模式下调用 add 算子，计算出结果 z，并与 python 脚本采用计算 md5sum 的方式进行对比，完全一样，则表示结果正确。

为了运行方便，我们使用一个 run.sh，写有 cpu 和 npu 模式的编译命令，通过输入参数进行选择 cpu 或 npu 模式进行编译，运行。

1）CPU 模式下：

使用 ICPU_RUN_KF 宏进行 CPU 调试。

ICPU_RUN_KF(add_custom, blockDim, x, y, z); // use this macro for cpu debug

复制代码

bash run.sh add_custom ascend910 AiCore cpu 运行结果：

2）NPU 模式下：

NPU 模式使用<<<>>>方式调用，由于 CPU 模式 g++没有<<<>>>的表达，需要使用内置宏 __CCE_KT_TEST。

#ifndef __CCE_KT_TEST__//call of kernel functionvoid add_custom_do(uint32_t blockDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y, uint8_t* z);{     add_custom<<<blockDim, l2ctrl, stream>>> (x, y, z);}#endif

复制代码

bash run.sh add_custom ascend910 AiCore npu 运行结果：