算子实现
基于Ascend C方式实现矢量算子的流程如下图图1所示。
- 算子分析:分析算子的数学表达式、输入、输出以及计算逻辑的实现,明确需要调用的Ascend C接口。
- 核函数定义:定义Ascend C算子入口函数。
- 根据矢量编程范式实现算子类:完成核函数的内部实现。
下文以ElemWise(Add)算子为例,对上述步骤进行详细介绍。本样例中介绍的算子完整代码参见add_custom.cpp。
算子分析
在开发算子代码之前需要分析算子的数学表达式、输入、输出以及计算逻辑的实现,明确需要调用的Ascend C接口。
- 明确算子的数学表达式及计算逻辑。
Add算子的数学表达式为:
z = x + y
计算逻辑是:Ascend C提供的矢量计算接口的操作元素都为LocalTensor,输入数据需要先搬运进片上存储,然后使用计算接口完成两个输入参数相加,得到最终结果,再搬出到外部存储上。Ascend C Add算子的计算逻辑如下图所示。
图2 算子计算逻辑
- 明确输入和输出。
- Add算子有两个输入:x与y,输出为z。
- 本样例中算子的输入支持的数据类型为half(float16),算子输出的数据类型与输入数据类型相同。
- 算子输入支持shape(8,2048),输出shape与输入shape相同。
- 算子输入支持的format为:ND。
- 确定核函数名称和参数。
- 您可以自定义核函数名称,本样例中核函数命名为add_custom。
- 根据对算子输入输出的分析,确定核函数有3个参数x,y,z;x,y为输入在Global Memory上的内存地址,z为输出在Global Memory上的内存地址。
- 确定算子实现所需接口。
通过以上分析,得到Ascend C Add算子的设计规格如下:
算子类型(OpType) |
Add |
|||
---|---|---|---|---|
算子输入输出 |
name |
shape |
data type |
format |
x(输入) |
(8, 2048) |
half |
ND |
|
y(输入) |
(8, 2048) |
half |
ND |
|
z(输出) |
(8, 2048) |
half |
ND |
|
核函数名称 |
add_custom |
|||
使用的主要接口 |
DataCopy:数据搬移接口 |
|||
Add:矢量双目指令接口 |
||||
EnQue、DeQue等接口:Queue队列管理接口 |
||||
算子实现文件名称 |
add_custom.cpp |
核函数定义
根据核函数定义中介绍的规则进行核函数的定义。
- 函数原型定义
本样例中,函数名为add_custom(核函数名称可自定义),根据算子分析中对算子输入输出的分析,确定有3个参数x,y,z,其中x,y为输入内存,z为输出内存。根据核函数定义核函数的规则介绍,函数原型定义如下所示:使用__global__函数类型限定符来标识它是一个核函数,可以被<<<>>>调用;使用__aicore__函数类型限定符来标识该核函数在设备端aicore上执行;为方便起见,统一使用GM_ADDR宏修饰入参,GM_ADDR宏定义请参考核函数定义。
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) { }
- 调用算子类的Init和Process函数。算子类的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(); }
- 对核函数的调用进行封装,得到add_custom_do函数,便于主程序调用。#ifndef ASCENDC_CPU_DEBUG表示该封装函数仅在编译运行NPU侧的算子时会用到,编译运行CPU侧的算子时,可以直接调用add_custom函数。根据核函数调用章节,调用核函数时,除了需要传入参数x,y,z,还需要传入blockDim(核函数执行的核数),l2ctrl(保留参数,设置为nullptr),stream(应用程序中维护异步操作执行顺序的stream)来规定核函数的执行配置。
#ifndef ASCENDC_CPU_DEBUG // call of kernel function void 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
算子类实现
根据上一节介绍,核函数中会调用算子类的Init和Process函数,本节具体讲解如何基于编程范式实现算子类。
根据矢量编程范式对Add算子的实现流程进行设计的思路如下,矢量编程范式请参考Vector编程范式,设计完成后得到的Add算子实现流程图参见图3:
- Add算子的实现流程分为3个基本任务:CopyIn,Compute,CopyOut。CopyIn任务负责将Global Memory上的输入Tensor xGm和yGm搬运至Local Memory,分别存储在xLocal,yLocal,Compute任务负责对xLocal,yLocal执行加法操作,计算结果存储在zLocal中,CopyOut任务负责将输出数据从zLocal搬运至Global Memory上的输出Tensor zGm中。
- CopyIn,Compute任务间通过VECIN队列inQueueX,inQueueY进行通信和同步,Compute,CopyOut任务间通过VECOUT队列outQueueZ进行通信和同步。
- 任务间交互使用到的内存、临时变量使用到的内存统一使用pipe内存管理对象进行管理。
算子类中主要实现上述流程,包含对外开放的初始化Init函数和核心处理函数Process,Process函数中会对上图中的三个基本任务进行调用;同时包括一些算子实现中会用到的私有成员,比如上图中的Global Tensor和VECIN、VECOUT队列等。KernelAdd算子类具体成员如下:
class KernelAdd { public: __aicore__ inline KernelAdd() {} // 初始化函数,完成内存初始化相关操作 __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z){} // 核心处理函数,实现算子逻辑,调用私有成员函数CopyIn、Compute、CopyOut完成矢量算子的三级流水操作 __aicore__ inline void Process(){} private: // 搬入函数,完成CopyIn阶段的处理,被核心Process函数调用 __aicore__ inline void CopyIn(int32_t progress){} // 计算函数,完成Compute阶段的处理,被核心Process函数调用 __aicore__ inline void Compute(int32_t progress){} // 搬出函数,完成CopyOut阶段的处理,被核心Process函数调用 __aicore__ inline void CopyOut(int32_t progress){} private: AscendC::TPipe pipe; //Pipe内存管理对象 AscendC::TQue<AscendC::QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY; //输入数据Queue队列管理对象,QuePosition为VECIN AscendC::TQue<AscendC::QuePosition::VECOUT, BUFFER_NUM> outQueueZ; //输出数据Queue队列管理对象,QuePosition为VECOUT AscendC::GlobalTensor<half> xGm; //管理输入输出Global Memory内存地址的对象,其中xGm, yGm为输入,zGm为输出 AscendC::GlobalTensor<half> yGm; AscendC::GlobalTensor<half> zGm; };
初始化函数主要完成以下内容:
- 设置输入输出Global Tensor的Global Memory内存地址。
本样例中使用多核并行计算,即把数据进行分片,分配到多个核上进行处理。Ascend C核函数是在一个核上的处理函数,所以只处理部分数据,需要在初始化函数中获取该核函数需要处理的输入输出在Global Memory上的内存偏移地址,并将该偏移地址设置在Global Tensor中。
以获取输入x在Global Memory上的内存偏移地址为例:
xGm.SetGlobalBuffer((__gm__ half*)x + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);
本样例中的分配方案是:数据整体长度TOTAL_LENGTH为8* 2048,平均分配到8个核上运行,每个核上处理的数据大小BLOCK_LENGTH为2048。x + BLOCK_LENGTH * GetBlockIdx()即为单核处理程序中x在Global Memory上的内存偏移地址,获取偏移地址后,使用GlobalTensor类的SetGlobalBuffer接口设定该核上Global Memory的起始地址以及长度。具体示意图请参考图4。
- 通过Pipe内存管理对象为输入输出Queue分配内存。
比如,为输入x的Queue分配内存,可以通过如下代码段实现:
pipe.InitBuffer(inQueueX, BUFFER_NUM, TILE_LENGTH * sizeof(half))
对于单核上的处理数据,可以进行数据切块(Tiling),在本示例中,仅作为参考,将数据切分成8块(并不意味着8块就是性能最优)。切分后的每个数据块再次切分成2块,即可开启double buffer,实现流水线之间的并行。
这样单核上的数据(2048个数)被切分成16块,每块TILE_LENGTH(128)个数据。上文代码表示Pipe为inQueueX分配了两块大小为TILE_LENGTH * sizeof(half)个字节的内存块,每个内存块能容纳TILE_LENGTH(128)个half类型数据。数据切分示意图如图5所示。
具体的初始化函数代码如下:
constexpr int32_t TOTAL_LENGTH = 8 * 2048; // total length of data constexpr int32_t USE_CORE_NUM = 8; // num of core used constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM; // length computed of each core constexpr int32_t TILE_NUM = 8; // split data into 8 tiles for each core constexpr int32_t BUFFER_NUM = 2; // tensor num for each queue constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; // separate to 2 parts, due to double buffer __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z) { // get start index for current core, core parallel xGm.SetGlobalBuffer((__gm__ half*)x + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH); yGm.SetGlobalBuffer((__gm__ half*)y + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH); zGm.SetGlobalBuffer((__gm__ half*)z + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH); // pipe alloc memory to queue, the unit is Bytes pipe.InitBuffer(inQueueX, BUFFER_NUM, TILE_LENGTH * sizeof(half)); pipe.InitBuffer(inQueueY, BUFFER_NUM, TILE_LENGTH * sizeof(half)); pipe.InitBuffer(outQueueZ, BUFFER_NUM, TILE_LENGTH * sizeof(half)); }
基于矢量编程范式,将核函数的实现分为3个基本任务:CopyIn,Compute,CopyOut。Process函数中通过如下方式调用这三个函数。
__aicore__ inline void Process() { // loop count need to be doubled, due to double buffer constexpr int32_t loopCount = TILE_NUM * BUFFER_NUM; // tiling strategy, pipeline parallel for (int32_t i = 0; i < loopCount; i++) { CopyIn(i); Compute(i); CopyOut(i); } }
根据编程范式上面的算法分析,将整个计算拆分成三个Stage,用户单独编写每个Stage的代码,三阶段流程示意图参见图3,具体流程如下:
- Stage1:CopyIn函数实现。
__aicore__ inline void CopyIn(int32_t progress) { // alloc tensor from queue memory AscendC::LocalTensor<half> xLocal = inQueueX.AllocTensor<half>(); AscendC::LocalTensor<half> yLocal = inQueueY.AllocTensor<half>(); // copy progress_th tile from global tensor to local tensor AscendC::DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH); AscendC::DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH); // enque input tensors to VECIN queue inQueueX.EnQue(xLocal); inQueueY.EnQue(yLocal); }
- Stage2:Compute函数实现。
- 使用DeQue从VecIn中取出LocalTensor。
- 使用Ascend C接口Add完成矢量计算。
- 使用EnQue将计算结果LocalTensor放入到VecOut的Queue中。
- 使用FreeTensor释放不再使用的LocalTensor。
__aicore__ inline void Compute(int32_t progress) { // deque input tensors from VECIN queue AscendC::LocalTensor<half> xLocal = inQueueX.DeQue<half>(); AscendC::LocalTensor<half> yLocal = inQueueY.DeQue<half>(); AscendC::LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>(); // call Add instr for computation AscendC::Add(zLocal, xLocal, yLocal, TILE_LENGTH); // enque the output tensor to VECOUT queue outQueueZ.EnQue<half>(zLocal); // free input tensors for reuse inQueueX.FreeTensor(xLocal); inQueueY.FreeTensor(yLocal); }
- Stage3:CopyOut函数实现。
- 使用DeQue接口从VecOut的Queue中取出LocalTensor。
- 使用DataCopy接口将LocalTensor拷贝到GlobalTensor上。
- 使用FreeTensor将不再使用的LocalTensor进行回收。
__aicore__ inline void CopyOut(int32_t progress) { // deque output tensor from VECOUT queue AscendC::LocalTensor<half> zLocal = outQueueZ.DeQue<half>(); // copy progress_th tile from local tensor to global tensor AscendC::DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH); // free output tensor for reuse outQueueZ.FreeTensor(zLocal); }
运行验证
核函数即算子kernel程序开发完成后,即可编写host侧的核函数调用程序,实现从host侧的APP程序调用算子,进行运行验证。包括CPU侧和NPU侧两种运行验证方法:
- CPU侧运行验证主要通过ICPU_RUN_KF CPU调测宏等CPU调测库提供的接口来完成;
- NPU侧运行验证主要通过使用Kernel Launch接口或者<<<>>>内核调用符,以及AscendCL API提供的运行时接口来完成。
具体步骤请参考核函数运行验证简介。