【2023 · CANN训练营第一季】TIK2算子开发入门

TIK2简介

TIK2是一种使用C/C++作为前端语言的编程框架，开发者可以使用TIK2提供的API编写自定义算子，并通过CCEC编译器将自定义算子编译成为可运行在昇腾AI处理器上的应用程序。

TIK2继承了TIK数据操作灵活的优点，除此之外，更多优势如下：
1、C/C++原语编程
2、编程模型屏蔽硬件差异，编程范式提高开发效率
3、多层级API封装，从简单到灵活，兼顾易用与高效
4、调试方法简单

编程模型

1、使用函数类型限定符
除了需要按照C/C++函数声明的方式定义核函数之外，还要为核函数加上额外的函数类型限定符，包含__global__和__aicore__。

2、使用变量类型限定符
1)、指针入参变量统一的类型定义为__gm__ uint8_t*，这里统一使用uint8_t类型的指针，在后续的使用中需要将其转化为实际的指针类型；用户亦可直接传入实际的指针类型。
2)、使用__global__函数类型限定符来标识它是一个核函数，可以被<<<…>>>调用；使用__aicore__函数类型限定符来标识该核函数在设备端aicore上执行；参数列表中使用变量类型限定符__gm__来表明输入输出的指针变量指向Global Memory上的某处地址。

3、调用核函数
核函数使用内核调用符<<<…>>>这种语法形式，来规定核函数的执行配置：

kernel_name<<<blockDim, l2ctrl, stream>>>(argument list);

内核调用符仅可在NPU侧编译时调用，CPU侧编译无法识别该符号。4、

如：小王做算子调试时发现CPU域调试精度无问题，NPU域调试时编译报错，那么可能会出现这种场景的原因是（AC）

A.
小王调用了cmath函数实现算子功能

B.
小王在代码内添加了Add 计算API

C.
小王在代码内添加了std::cout打印

D.
小王在代码运算时使用了 + 号
解释：A：cmath函数不适用于NPU。
C: cpu可以使用打印，但是NPU不行。

执行配置由3个参数决定：
blockDim，规定了核函数将会在几个核上执行，blockDim的大小不能超过当前设备上核的配置个数。每个执行该核函数的核会被分配一个逻辑ID，表现为内置变量block_idx，可以在核函数的实现中直接使用；
l2ctrl，保留参数，暂时设置为固定值nullptr，开发者无需关注；
stream，类型为aclrtStream，stream是一个任务队列，应用程序通过stream来管理任务的并行。

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

5、算子分析Add
Add算子的数学表达式为: z = x + y

dstLocal–输出–目的操作数，类型为LocalTensor。
srcLocal–输入–源操作数，类型为LocalTensor
例如：小王想实现一个add算子，那么他可能会使用的代码为(AB)

A.
Add(zLocal, xLocal, yLocal, TILE_LENGTH);

B.
zLocal = xLocal + yLocal

C.
Adds(dstLocal, srcLocal, TILE_LENGTH);

D.
Sub(zLocal, xLocal, yLocal, TILE_LENGTH);
解释：B：根据Add计算规则 z=x+y。
A: 为调用。

注：其他规则
1.必须具有void返回类型。
2.使用extern “C”。
3.仅支持入参为指针或C/C++内置数据类型（Primitive data types），如：half* s0、float* s1、int32_t c。

注意事项

1、当前TIK2支持的AI处理器型号为昇腾310P AI处理器、昇腾910 AI处理器，其他型号暂不支持。
2、当前支持用户使用g++等C/C++编译器编译在cpu侧执行的TIK2算子，并使用gdb单步调试；
3、支持用户使用CCEC编译器编译在npu侧执行的TIK2算子，实现加速计算，暂不支持加载至网络模型中进行整网验证。
4、TIK2提供了多层级的0-3级API，随着级别增高，API使用的自由度降低，易用性增强。您可以根据需要选择合适的API，使用最通俗易懂的高级接口快速搭建算子逻辑，使用自由灵活的低级接口进行复杂的逻辑实现和性能调优。