【2023 · CANN训练营第一季】第一节课程笔记-算子开发基础

一、算子开发优势简介

1.语言：C/C++

2.编程模式屏蔽硬件差异，编程范式提高编程效率

3.多层级API封装，从简单到灵活，兼顾易用与高效

4.孪生调试，CPU侧模仿NPU侧行为，可现在CPU侧调试

二、核函数简介

基本概念：

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

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

基本构成图如下：

使用函数类型限定符

上图的__global__和__aicore__是其两种类型限定符，用来标识它是一个核函数，使用<<<...>>>来进行调用,aicore是用来表示该核函数是用来在设备侧AI Core上执行 

还有一种入参变量的限定符:

为了方便，统一使用__gm__uint8_t*。用户可统一使用uint8_t,并在使用时转化为实际的指针类型;亦可直接传入实际的指针类型.

其他规则：

1.必须具有void返回类型

2.使用extern"C"

3.仅支持入参为指针类型或C/C++内置数据类型(Primitive Data Tvpes)，如: half* s0、float* s1、int32_t

函数调用方式：

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

• blockDim，规定了核函数将会在几个核上执行，每个执行该核函数的核会被分配一个逻辑ID，表现为内置变量block idx，编号从0开始，可为不同的逻辑核定义不同的行为，可以在算子实现中使用

• l2ctrl，保留参数，暂时设置为固定值nullptr

• stream，类型为acltStream，stream是一个任务队列，应用程序通过stream来管理任务的并行

eg:

HelloWorld<<<8，nullptr，stream>>>(fooDevice);

blockDim设置为8，表示在8个核上调用了HelLoworld核函数，每个核都会独立且并行地执行该核函数

aclrtstreamStream可以通过aclrtCreatestream来创建,它的作用是在当前进程或线程中显式创建一个arqument List设置为fooDevice这1个入参

三、算子执行的模式

算子的执行分为CPU模式和NPU模式

代码可以使用宏定义来进行分割CPU运行还是NPU运行

四、接口简介

常用数据定义：

GlobalTensor：

用来存放Global Memory(外部存储)的全局数据。

LocalTensor：

用于存放核上Local Memory (内部存储)的数据

矢量计算指令接口简介：

矢量计算指令接口，能够启动AI Core中的Vector单元执行计算为了降低开发者的使用门槛，指令按照由易到难，分成了3级到0级接口。其中3级接口最为简单，0级接口最为复杂，(1级接口还末发布)

多层级API封装的作用:

• 降低复杂指令的使用难度

• 跨代兼容性保障

• 保留最大灵活度的可能

3级接口

3级接口，运算符重载，支持+，-,*,/，|, &, ^,等C/C++运算符实现2级接口的简化表达

允许用户使用形如: dst = src*sr1，针对整个Tensor进行计算

2级接口

2级连续计算接口，针对源操作数srcLocal的连续COUNT个数据进行计算，并连续写入目的操作数dstLocal，提供了一维Tensor的连续COUNT个数据的计算支持。

0级接口

0级功能灵活计算接口，是最底层的开发接口，可以完整发挥硬件优势的计算API，可以进行非连续的计算该功能可以充分发挥CANN系列芯片的强大功能指令，支持对每个操作数的Block stride，Repeat stride，MASK的操作，允许用户使用诸多的通用参数来定制化所需要的操作

1.Repeat times：迭代的次数

矢量计算单元，一次最多可以计算256Bytes的数据，每次读取连续的8个block (每个block 32Bytes，共256Bytes)数据进行计算，为完成对输入数据的处理，必须通过多次迭代(repeat)才能完成所有数据的读取与计算。

2.Repeat stride：相邻迭代间相同block的地址步长

当Repeat times大于1，需要多次迭代完成矢量计算时，可以根据不同的使用场景合理设置相邻迭代间相同block的地址步长Repeat stride的值

• 连续计算场景:

假设定义一个Tensor供目的操作数和源操作数同时使用(即地址重叠)，Repeat stride取值为8.此时，矢量计算单元第一次迭代读取连续8个block，第二轮迭代读取下一个连续的8个block，通过多次选代即可完成所有输入数据的计算

• 非连续计算场景:

Repeat stride取值大8(如取10)时，则相邻选代间矢量计算单元读取的数据在地址上不连续，出现2个block的间隔

• 反复计算场景:

Repeat stride取值为0时，矢量计算单元会对首个连续的8个block进行反复读取和计算

• 部分重复计算:

Repeat stride取值大于且小于8时，相邻选代间部分数据会被矢量计算单元重复读取和计算，此种情形一般场景不涉及

3.Block stride：表示同一迭代内不同block的地址步长

如果需要控制单次迭代内，数据处理的步长，可以通过设置同一迭代内不同block的地址步长Block stride来实现连续计算，Block stride 设置为1，对同一迭代内的8个block数据连续进行处理非连续计算，Block stride值大于1(如取2)，同一选代内不同block之间在读取数据时出现一个block的间隔

4.Mask参数

Mask用于控制每次迭代内参与计算的元素。可通过连续模式和逐比特模式两种方式进行设置

• 连续模式:

表示前面连续的多少个元素参与计算。数据类型为uint64_t。取值范围和操作数的数据类型有关，数据类型不同，每次迭代内能够处理的元素个数最大值不同(当前数据类型单次迭代时能处理的元素个数最大值为: 256/sizeof(数据类型))。当操作数的数据类型占比特位16位时(如half，uint16_t)，mask[1,128];当操作数为32位时(如float, int32_t) ，maske取值范围[1, 64]。

• 逐比特模式:

可以按位控制哪些元素参与计算，比特位的值为1表示参与计算，0表示不参与。参数类型为长度为2的uint64_t类型数组 参数取值范围和操作数的数据类型有关，数据类型不同，每次迭代内能够处理的元素个数最大值不同。当操作数为16位时，mask[0]、mask[1]取值范围[0, 264-11; 当dst/src为32位时，mask[1]为0，maski[0]取值范围为[0,(0.2的64次方)-1]

三等级通用接口

ps:该文仅是为了记录CANN训练营的学习过程所用，不参与任何商业用途