TIK C++算子开发入门
一、TIK C++基本概念
TIK C++是一种使用**C/C++**作为前端语言的算子开发工具，通过**四层接口抽象**、**并行编程范式**、**孪生调试**等技术，极大提高算子开发效率，助力AI开发者低成本完成算子开发和模型调优部署。
二、核函数
1.基本定义
核函数（Kernel Function）是TIK C++算子**设备侧**的入口。TIK C++允许用户使用核函数这种C/C++函数的语法扩展来管理设备侧的运行代码，用户在核函数中实现算子逻辑的编写，例如自定义算子类及其成员函数以实现该算子的所有功能。
核函数是直接在设备侧执行的代码。在核函数中，需要为在**一个核**上执行的代码规定要进行的**数据访问**和**计算操作**，当核函数被调用时，**多个核**将并行执行同一个计算任务。
2.编写规范
```c++
__global__ __aicore__ void kernel_name(argument list);
```
a.使用函数限定符
使用**`__global__`**函数类型限定符来标识它是一个核函数，可以被**`<<<...>>>`**调用；使用**`__aicore__`**函数类型限定符来标识该核函数在设备侧**`AI Core`**上执行。
b.使用变量类型限定符
指针入参变量统一的类型定义为**`__gm__ uint8_t*`**。用户可统一使用**`uint8_t`**类型的指针，并在使用时转化为实际的指针类型；亦可直接传入实际的指针类型。
c.其他规则
必须具有**void**返回类型；使用**extern "C"**；仅支持入参为指针类型或C/C++内置数据类型(Primitive Data Types)，如：`half* s0`、`float* s1`、`int32_t c`。
3.调用方式
a.核函数使用内核调用符**`<<<...>>>`**这种语法形式，来规定核函数的执行配置。
```c++
kernel_name<<<blockDim, l2ctrl, stream>>>(argument list);
```
- **blockDim**，规定了核函数将会在几个核上执行，每个执行该核函数的核会被分配一个逻辑`ID`，表现为内置变量`block_idx`，编号从0开始，可为不同的逻辑核定义不同的行为，可以在算子实现中使用。
- **`l2ctrl`**，保留参数，暂时设置为固定值`nullptr`。
- **`stream`**，类型为`aclrtStream`，`stream`是一个任务队列，应用程序通过`stream`来管理任务的并行。
b.核函数的调用是**异步**的，核函数的调用结束后，控制权立刻返回给主机侧。
强制主机侧程序等待所有核函数执行完毕的API（阻塞应用程序运行，直到指定Stream中的所有任务都完成，同步接口）为**`aclrtSynchronizeStream`**。
```
aclError aclrtSynchronizeStream(aclrtStream stream);
```
三、TIK C++`演示样例`
```c++
/*
* Copyright (c) Huawei Technologies Co., Ltd. 2022-2023. All rights reserved.
*/
#ifdef __CCE_KT_TEST__
// 运行CPU模式包含的头文件
#include "tikicpulib.h"
#define __aicore__
#else
// 运行NPU模式包含的头文件
#include "acl/acl.h"
#define __aicore__ [aicore]
#endif
// 核函数的定义
extern "C" __global__ __aicore__ void HelloWorld(__gm__ uint8_t* foo) {}
int32_t main(int32_t argc, char* argv[])
{
// 申请的数据内存大小Bytes
size_t fooSize = 256;
// 执行核函数的逻辑核个数
uint32_t blockDim = 8;
// 内置宏（主机侧执行CPU模式逻辑）
#ifdef __CCE_KT_TEST__
uint8_t *foo = (uint8_t *)tik2::GmAlloc(fooSize);
// 执行宏（主机侧执行CPU模式逻辑）
ICPU_RUN_KF(HelloWorld, blockDim, foo);
// 释放CPU模式下的内存空间
tik2::GmFree((void *)foo);
// 内置宏（主机侧执行NPU模式逻辑）
#else
// AscendCL初始化
aclInit(nullptr);
// 运行管理资源申请
aclrtStream stream = nullptr;
aclrtCreateStream(&stream);
uint8_t *fooDevice;
aclrtMalloc((void**)&fooDevice, fooSize, ACL_MEM_MALLOC_HUGE_FIRST);
// 执行任务并等待
HelloWorld<<<blockDim, nullptr, stream>>>(fooDevice);
aclrtSynchronizeStream(stream);
// 运行管理资源释放
aclrtFree(fooDevice);
aclrtDestroyStream(stream);
// AscendCL去初始化
aclFinalize();
#endif
return 0;
}
```
`四`、TIK C++数据结构
**`GlobalTensor`**用来存放**`Global Memory`**（外部存储）的全局数据；**`LocalTensor`**用于存放核上**`Local Memory`**（内部存储）的数据。
五、多层级API封装
1.基本概念
矢量计算指令接口，能够启动AI Core中的`Vector`单元执行计算。为了降低开发者的使用门槛，指令按照由易到难，分成了3级到0级接口。其中3级接口最为简单，0级接口最为复杂。
a. **3级接口**，运算符重载，支持`+, -, *, /, |, &, ^, >, < , >=, <=，!=，==`实现2级接口的简化表达。
b. **2级接口**，针对源操作数`srcLocal`的连续`COUNT`个数据进行计算，并连续写入目的操作数`dstLocal`，提供了一维`Tensor`的连续`COUNT`个数据的计算支持。
c. **1级接口**还未发布。
d. **0级接口**，是最底层的开发接口，可以完整发挥硬件优势的计算API，可以进行非连续的计算该功能可以充分发挥CANN系列芯片的强大功能指令，支持对每个操作数的`Block stride`，`Repeat stride`，`MASK`的操作，允许用户使用诸多的通用参数来定制化所需要的操作。
2.使用方法
```
// int16_t数据类型， dstLocal长度为512个int16_t
// 0级接口样例-mask连续模式
uint64_t mask = 128;
// repeatTimes = 4, 一次迭代计算128个数, 共计算512个数
// dstBlkStride, src0BlkStride, src1BlkStride = 1, 单次迭代内数据连续读取和写入
// dstRepStride, src0RepStride, src1RepStride = 8, 相邻迭代间数据连续读取和写入
Add(dstLocal, src0Local, src1Local, mask, 4, { 1, 1, 1, 8, 8, 8 });
// 0级接口样例-mask逐bit模式
uint64_t mask[2] = { UINT64_MAX, UINT64_MAX };
// repeatTimes = 4, 一次迭代计算128个数, 共计算512个数
// dstBlkStride, src0BlkStride, src1BlkStride = 1, 单次迭代内数据连续读取和写入
// dstRepStride, src0RepStride, src1RepStride = 8, 相邻迭代间数据连续读取和写入
Add(dstLocal, src0Local, src1Local, mask, 4, { 1, 1, 1, 8, 8, 8 });
// 2级接口样例
Add(dstLocal, src0Local, src1Local, 512);
// 3级接口样例
dstLocal = src0Local + src1Local;
```
