【2023 · CANN训练营第一季】TIK C++算子开发入门 第一章——TIK C++算子开发入门

1.TIK C++介绍

TIK C++是一种使用C/C++作为前端语言的算子开发工具，通过四层接口抽象、并行编程范式、孪生调试等技术，极大提高算子开发效率，助力AI开发者低成本完成算子开发和模型调优部署

使用TIK C++开发自定义算子的优势：
(1)C/C++原语编程
(2)编程模型屏蔽硬件差异，编程范式提高开发效率
(3)多层级API封装，从简单到灵活，兼顾易用与高效
(4)孪生调试，CPU侧模拟NPU侧的行为，可先在CPU侧调试

2.核函数

核函数（Kernel Function）是TIK C++算子设备侧的入口。TIK C++允许用户使用核函数这种C/C++函数的语法扩展来管理设备侧的运行代码，用户在核函数中实现算子逻辑的编写，例如自定义算子类及其成员函数以实现该算子的所有功能。核函数是主机侧和设备侧连接的桥梁

核函数是直接在设备侧执行的代码。在核函数中，需要为在一个核上执行的代码规定要进行的数据访问和计算操作，当核函数被调用时，多个核将并行执行同一个计算任务

2.1编写核函数

(1)使用函数类型限定符
除了需要按照C/C++函数声明的方式定义核函数之外，还要为核函数加上额外的函数类型限定符，包含__global__和__aicore__
使用__global__函数类型限定符来标识它是一个核函数，可以被<<<…>>>调用；使用__aicore__函数类型限定符来标识该核函数在设备侧AI Core上执行：
__global__ __aicore__ void kernel_name(argument list);
(2)使用变量类型限定符
为了方便：指针入参变量统一的类型定义为__gm__ uint8_t*
用户可统一使用uint8_t类型的指针，并在使用时转化为实际的指针类型；亦可直接传入实际的指针类型

其他规则
必须具有void返回类型
使用extern “C”
仅支持入参为指针类型或C/C++内置数据类型(Primitive Data Types)，如：half* s0、float* s1、int32_t c

核函数的调用语句是C/C++函数调用语句的一种扩展
常见的C/C++函数调用方式是如下的形式：
function_name(argument list);
核函数使用内核调用符<<<…>>>这种语法形式，来规定核函数的执行配置：
kernel_name<<<blockDim, l2ctrl, stream>>>(argument list);
blockDim，规定了核函数将会在几个核上执行，每个执行该核函数的核会被分配一个逻辑ID，表现为内置变量block_idx，编号从0开始，可为不同的逻辑核定义不同的行为，可以在算子实现中使用
l2ctrl，保留参数，暂时设置为固定值nullptr
stream，类型为aclrtStream，stream是一个任务队列，应用程序通过stream来管理任务的并行

3.接口介绍

常用数据定义：GlobalTensor
GlobalTensor用来存放Global Memory（外部存储）的全局数据
定义原型
template <typename T> class GlobalTensor { // 传入全局数据的指针，并手动设置一个buffer size，初始化GlobalTensor void SetGlobalBuffer(__gm__ T* buffer, uint32_t bufferSize);

LocalTensor
LocalTensor用于存放核上Local Memory（内部存储）的数据
定义原型
template <typename T> class LocalTensor { T GetValue(const uint32_t offset) const; // 获取 LocalTensor 中的某个值，返回 T 类型的立即数。 template <typename T1> void SetValue(const uint32_t offset, const T1 value) const; // 设置 LocalTensor 中的某个值。offset单位为element // 获取距原LocalTensor起始地址偏移量为offset的新LocalTensor，注意offset不能超过原有LocalTensor的size大小。offset单位为element LocalTensor operator[](const uint32_t offset) const; uint32_t GetSize() const; // 获取当前LocalTensor size大小

矢量计算指令接口