TIK C++算子编程范式

什么是编程范式？
TIK C++编程范式把算子内部的处理程序，分成多个流水任务（Stage），以张量（Tensor）为数据载体，以队列（Queue）进行任务之间的通信与同步，以内存管理模块（Pipe）管理任务间的通信内存。

• 快速开发编程的固定步骤
• 统一代码框架的开发捷径
• 使用者总结出的开发经验
• 面向特定场景的编程思想
• 定制化的方法论开发体验

1. 流水任务

1.1 流水任务定义

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

1.2 流水任务例子

举例来说，单核处理程序的功能可以被拆分成3个流水任务：Stage1、Stage2、Stage3，每个任务专注于完成单一功能；需要处理的数据被切分成n片，使用Progress1~n表示，每个任务需要依次完成n个数据切片的处理。Stage间的箭头表达数据间的依赖关系，比如Stage1处理完Progress1之后，Stage2才能对Progress1进行处理。

若Progress的n=3，待处理的数据被切分成3片，对于同一片数据，Stage1、Stage2、Stage3之间的处理具有依赖关系，需要串行处理；不同的数据切片，同一时间点，可以有多个流水任务Stage在并行处理，由此达到任务并行、提升性能的目的。

1.3 流水任务——矢量编程

矢量算子编程范式把算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut

• CopyIn负责数据搬入操作
• Compute负责矢量计算操作
• CopyOut负责数据搬出操作

2. 任务间通信与同步

2.1 数据通信与同步的管理者

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

2.2 数据的载体

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

2.3 任务间通信与同步——矢量编程

矢量编程中的逻辑位置（QuePosition）：搬入数据的存放位置：VECIN、搬出数据的存放位置：VECOUT
矢量编程主要分为CopyIn、Compute、CopyOut三个任务：
CopyIn任务中将输入数据从GlobalTensor搬运至LocalTensor后，需要使用EnQue将LocalTensor放入VECIN的Queue中
Compute任务等待VECIN的Queue中LocalTensor出队之后才可以进行矢量计算，计算完成后使用EnQue将计算结果LocalTensor放入到VECOUT的Queue中
CopyOut任务等待VECOUT的Queue中LocalTensor出队，再将其拷贝到GlobalTensor

Stage1：CopyIn任务
使用DataCopy接口将GlobalTensor拷贝到LocalTensor
使用EnQue将LocalTensor放入VECIN的Queue中

Stage2：Compute任务
使用DeQue从VECIN中取出LocalTensor
使用TIK C++指令API完成矢量计算：Add
使用EnQue将结果LocalTensor放入VECOUT的Queue中

Stage3：CopyOut任务
使用DeQue接口从VECOUT的Queue中取出LocalTensor
使用DataCopy接口将LocalTensor拷贝到GlobalTensor

任务通信与同步管理单元Queue
TQue
Queue接口
EnQue(…)
DeQue(…)
Queue的逻辑位置（QuePosition）
VECIN
VECOUT

3. 内存管理

3.1 内存管理模块Pipe

任务间数据传递使用到的内存统一由内存管理模块Pipe进行管理。
Pipe作为片上内存管理者，通过InitBuffer接口对外提供Queue内存初始化功能，开发者可以通过该接口为指定的Queue分配内存。
Queue队列内存初始化完成后，需要使用内存时，通过调用AllocTensor来为LocalTensor分配内存给Tensor，当创建的LocalTensor完成相关计算无需再使用时，再调用FreeTensor来回收LocalTensor的内存。

// 使用AllocTensor分配Tensor
TPipe pipe;
TQue<TPosition::VECOUT, 2> que;
int num = 4;
int len = 1024;
// InitBuffer分配内存块数为4，每块大小为1024Bytes
pipe.InitBuffer(que, num, len);
// AllocTensor分配Tensor长度为1024Bytes
LocalTensor<half> tensor1 = que.AllocTensor();
// 使用FreeTensor释放通过AllocTensor分配的Tensor，注意配对使用
que.FreeTensor<half>(tensor1);

3.2 临时变量内存管理

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

// 为TBuf初始化分配内存，分配内存长度为1024字节
TPipe pipe;
TBuf<TPosition::VECIN> calcBuf; // 模板参数为QuePosition中的VECIN类型
uint32_t byteLen = 1024;
pipe.InitBuffer(calcBuf, byteLen);
// 从calcBuf获取Tensor,Tensor为pipe分配的所有内存大小，为1024字节
LocalTensor<int32_t> tempTensor1 = calcBuf.Get<int32_t>();
// 从calcBuf获取Tensor,Tensor为128个int32_t类型元素的内存大小，为512字节
LocalTensor<int32_t> tempTensor1 = calcBuf.Get<int32_t>(128);

内存管理单元Pipe
TPipe
TBuf
Pipe接口
InitBuffer(…)
Queue接口
AllocTensor(…)
FreeTensor(…)
Buf接口
Get(…)

4. 算子开发流程

4.1 两种模式

4.1.1 快速TIK C++算子开发流程

完成算子核函数的开发
基于内核调用符方式进行算子运行验证

4.1.2 标准TIK C++算子开发流程

完成算子核函数的开发
完成单算子网络应用程序的开发
基于ACL单算子调用方式进行算子运行验证

4.2 TIK C++矢量算子的编程

**算子分析：**分析算子的数学表达式、输入、输出以及计算逻辑的实现，明确需要调用的TIK C++接口。
**核函数定义：**定义TIK C++算子入口函数。
**根据矢量编程范式实现算子类：**完成核函数的内部实现。

4.3 算子分析

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

4.4 核函数定义

在之前我们已经学过HelloWorld核函数的实现

extern "C" __global__ __aicore__ void HelloWorld(__gm__ uint8_t* foo) {}

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

注：算子类名和成员函数名无特殊要求，开发者可根据自身的C/C++编码习惯，决定核函数中的具体实现

// implementation of kernel function
extern "C" __global__ __aicore__ void add_tik2(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z)
{
    KernelAdd op;
    op.Init(x, y, z);
    op.Process();
}

对于核函数的调用，使用第一章中提到的内置宏__CCE_KT_TEST__来标识<<<…>>>仅在NPU模式下才会编译到（CPU模式g++没有<<<…>>>的表达），对核函数的调用进行封装，可以在封装函数中补充其他逻辑，这里仅展示对于核函数的调用

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

4.5 算子类实现

• 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内存管理对象对任务间交互使用到的内存、临时变量使用到的内存统一进行管理

实操部分：
算子类类名：KernelAdd
初始化函数Init()和核心处理函数Process()
三个流水任务：CopyIn(), Compute(), CopyOut()
Progress的含义：

TQue模板的BUFFER_NUM的含义：
该Queue的深度，double buffer优化技巧

class KernelAdd {
public:
    __aicore__ inline KernelAdd() {}
    // 初始化函数，完成内存初始化相关操作
    __aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* 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:
    // Pipe内存管理对象
    TPipe pipe;
    // 输入数据Queue队列管理对象，QuePosition为VECIN
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;
    // 输出数据Queue队列管理对象，QuePosition为VECOUT
    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueZ;
    // 管理输入输出Global Memory内存地址的对象，其中xGm, yGm为输入，zGm为输出
    GlobalTensor<half> xGm, yGm, zGm;
};

4.6 成员函数实现

4.6.1 Init()函数实现

• 多核并行运算
  使用多核并行计算，需要将数据切片，获取到每个核实际需要处理的在Global Memory上的内存偏移地址。
  数据整体长度TOTAL_LENGTH为8* 2048，平均分配到8个核上运行，每个核上处理的数据大小BLOCK_LENGTH为2048。block_idx为核的逻辑ID，(gm half*)x + block_idx * BLOCK_LENGTH即索引为block_idx的核的输入数据在Global Memory上的内存偏移地址。
  
• 单核处理数据
  对于单核处理数据，可以进行数据切块（Tiling），将数据切分成8块。切分后的每个数据块再次切分成BUFFER_NUM=2块，可开启double buffer，实现流水线之间的并行。
  单核需要处理的2048个数被切分成16块，每块TILE_LENGTH=128个数据。Pipe为inQueueX分配了BUFFER_NUM块大小为TILE_LENGTH * sizeof(half)个字节的内存块，每个内存块能容纳TILE_LENGTH=128个half类型数据。
  

// total length of data
constexpr int32_t TOTAL_LENGTH = 8 * 2048;
// num of core used
constexpr int32_t USE_CORE_NUM = 8;
// length computed of each core
constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM;
// split data into 8 tiles for each core
constexpr int32_t TILE_NUM = 8;
// tensor num for each queue
constexpr int32_t BUFFER_NUM = 2;
// each tile length is separated to 2 part, due to double buffer
constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM;
__aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z)
{
    // get start index for current core, core parallel
    xGm.SetGlobalBuffer((__gm__ half*)x + block_idx * BLOCK_LENGTH, BLOCK_LENGTH);
    yGm.SetGlobalBuffer((__gm__ half*)y + block_idx * BLOCK_LENGTH, BLOCK_LENGTH);
    zGm.SetGlobalBuffer((__gm__ half*)z + block_idx * BLOCK_LENGTH, 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));
}

4.6.2 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);
    }
}
__aicore__ inline void CopyIn(int32_t progress)
{
    // alloc tensor from queue memory
    LocalTensor<half> xLocal = inQueueX.AllocTensor<half>();
    LocalTensor<half> yLocal = inQueueY.AllocTensor<half>();
    // copy progress_th tile from global tensor to local tensor
    DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH);
    DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH);
    // enque input tensors to VECIN queue
    inQueueX.EnQue(xLocal);
    inQueueY.EnQue(yLocal);
}
__aicore__ inline void Compute(int32_t progress)
{
    // deque input tensors from VECIN queue
    LocalTensor<half> xLocal = inQueueX.DeQue<half>();
    LocalTensor<half> yLocal = inQueueY.DeQue<half>();
    LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>();
    // call Add instr for computation
    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);
}
__aicore__ inline void CopyOut(int32_t progress)
{
    // deque output tensor from VECOUT queue
    LocalTensor<half> zLocal = outQueueZ.DeQue<half>();
    // copy progress_th tile from local tensor to global tensor
    DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH);
    // free output tensor for reuse
    outQueueZ.FreeTensor(zLocal);
}

4.6.3 ouble buffer机制

double buffer通过将数据搬运与矢量计算并行执行以隐藏数据搬运时间并降低矢量指令的等待时间，最终提高矢量计算单元的利用效率
1个Tensor同一时间只能进行搬入、计算和搬出三个流水任务中的一个，其他两个流水任务涉及的硬件单元则处于Idle状态
如果将待处理的数据一分为二，比如Tensor1、Tensor2
当矢量计算单元对Tensor1进行Compute时，Tensor2可以执行CopyIn的任务
当矢量计算单元对Tensor2进行Compute时，Tensor1可以执行CopyOut的任务
当矢量计算单元对Tensor2进行CopyOut时，Tensor1可以执行CopyIn的任务
由此，数据的进出搬运和矢量计算之间实现并行，硬件单元闲置问题得以有效缓解

\

5 实际操作演示

5.1 核函数

5.2 Cmakelist编写

5.3 CPU模式

5.4 DoprintData用作数据打印

5.5 CPU框架

CPU框架：申请资源，读入，计算，写入，释放，NPU框架相似，但是语法会复杂一点

5.6 run.sh文件编写

5.7 CPU运行结果

执行完成，pid显示八个核

5.8 NPU运行结果