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

TIK C++简介

TIK C++是一种使用C/C++作为前端语言的算子开发工具，通过四层接口抽象、并行编程范式、孪生调试等技术，极大提高算子开发效率。

TIK C++的优势有：

C/C++原语编程
编程模型屏蔽硬件差异，编程范式提高开发效率
多层级API封装，从简单到灵活，兼顾易用与高效
孪生调试，CPU侧使用开发者熟悉的gdb工具即可进行功能调试，调试后可无缝移植到AI处理器运行。

核函数

核函数（Kernel Function）是TIK C++算子的入口。TIK C++允许用户使用核函数这种C/C++函数的语法扩展来管理设备端的运行代码，用户在核函数中进行算子类的声明和其成员函数的调用，由此实现该算子的所有功能。核函数是主机端和设备端连接的桥梁。

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

使用函数类型限定符

除了需要按照C/C++函数声明的方式定义核函数之外，还要为核函数加上额外的函数类型限定符，包含__global__和__aicore__。
使用__global__函数类型限定符来标识它是一个核函数在设备端执行，可以被<<<...>>>调用，必须有一个void返回值类型；使用__aicore__函数类型限定符来标识该核函数在设备端aicore上执行，仅从设备端调用。
使用变量类型限定符
指针入参变量统一的类型定义为__gm__ uint8_t*，这里统一使用uint8_t类型的指针，在后续的使用中需要将其转化为实际的指针类型；用户亦可直接传入实际的指针类型。
变量类型限定符__gm__表明该指针变量指向Global Memory上某处内存地址
其他规则
1. 必须具有void返回类型。
2. 使用extern “C”。
3. 仅支持入参为指针或C/C++内置数据类型（Primitive data types），如：halfs0、float s1、int32_t c。

核函数的调用

核函数的调用语句是C/C++函数调用语句的一种扩展，使用内核调用符<<<...>>>这种语法形式，来规定核函数的执行配置：

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

内核调用符仅可在NPU侧编译时调用，CPU侧编译无法识别该符号，参数：

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

核函数的调用和主机线程是异步的，核函数的调用结束后，控制权立刻返回给主机端，可以强制主机测程序等待所有核函数执行完毕。

1	aclError aclrtSynchronizeStream(aclrtStream stream);

示例：

// 核函数的实现
extern "C" __global__ __aicore__ void add_tik2(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z)
{
    // 实现算子的具体逻辑，具体见Add算子示例
    // 初始化算子类，算子类中定义了数据搬运和计算逻辑
    KernelAdd op;
    // 接收外界传入的内存地址指针，做内存初始化和分配操作
    op.Init(x, y, z);
    // 完成算子的数据搬运与数据计算等核心逻辑
    op.Process();
}

// 核函数的调用
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);
}

执行模式

TIK C++算子可以在CPU模式或NPU模式执行：

CPU模式：算子功能调试用，可以模拟在NPU上计算行为，不需要依赖昇腾设备；
NPU模型：算子功能/性能调试，使用NPU算子加速。

内置宏__CCE_KT_TEXT__标识被宏包括的代码在特定的模式下编译。

#ifdef __CCE_KT_TEXT__表示CPU模型下编译该段代码
#ifndef __CCE_KT_TEXT__表示NPU模型下编译该段代码

主机侧

对于CPU模式，主机侧使用封装的执行宏ICPU_RNN_KF执行，ICPU_RNN_KF(kernel_name, bockDim, foo)，第一个参数为核函数名称，第二个为核数量，第三个为使用GmAlloc申请的内存空间。

#ifdef __CCE_KT_TEST__
#include "tikicpulib.h"
#define __aicore__
#endif

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

int32_t main(int32_t argc, char* argv[])
{
    // 定义申请数据内存大小核执行核函数的逻辑核个数
    size_t fooSize = 256;
    uint32_t blockDim = 8;
// CPU模式
#ifdef __CCE_KT_TEST__
    // 申请CPU模式下的内存空间
    uint8_t *foo = (uint8_t *)tik2::GmAlloc(fooSize);
    // 封装的执行宏
    ICPU_RNN_KF(HelloWorld, blockDim, foo);
    // 释放CPU模式下的内存空间
    tik2::GmFree((void *) foo);
#endif
    return 0;
}

对于NPU模式，主机侧负责数据在主机侧内存的申请，主机到设备的拷贝，核函数执行同步核回收资源的工作。

// NPU模式头文件
#ifndef __CCE_KT_TEST__
#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[])
{
    // 定义申请数据内存大小核执行核函数的逻辑核个数
    size_t fooSize = 256;
    uint32_t blockDim = 8;
// NPU模式
#ifndef __CCE_KT_TEST__
    // 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);

    // 去初始化
    aclFinalize();
#endif
    return 0;
}

编译

对于CPU模式，使用gcc进行编译；

对于NPU模式，使用ccec进行编译。

接口

常用数据定义

GlobalTensor
用来存在Global Memory的全局数据。

定义

template <typename T> class GlobalTensor {
void SetGlobalBuffer(__gm__ T* buffer, int32_t bufferSize);
void SetGlobalBuffer(__gm__ T* buffer);
}
// buffer：host 侧传入的全局数据指针
// bufferSize：设置 GlobalTensor size

调用

void Init(__gm__ uint8_t *__restrict__ src_gm, __gm__ uint8_t*__restrict__ dst_gm)
{
    int data_size = 256;

    GlobalTensor<int32_t> input_global;
    input_global.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t *>(src_gm), data_size);
    ...
}

LocalTensor
用于存放核上内部储存的数据

定义

template <typename T> class LocalTensor {
T GetValue(const int32_t offset) const;
template <typename T1> void SetValue(const int32_t offset, const T1 value) const;
LocalTensor operator[](const int32_t offset) const;
int32_t GetSize() const;
}
// GetValue: 获取 LocalTensor 中的某个值，返回 T 类型的立即数。
    // offset：偏移量，单位为 element
// SetValue: 设置 LocalTensor 中的某个值。
    // index：偏移量，单位为 element; value：设置值，单位为任意类型
// operator[]: 获取距原 LocalTensor 起始地址偏移量为 offset 的新 LocalTensor，注意 offset 不能超过原有 LocalTensor 的 size 大小。
    // offset：偏移量
// GetSize: 获取当前 LocalTensor size 大小。

调用

if(HasTensorInQue(qIDVecIn0)){
auto input0_local = DeQue<half>(qIDVecIn0);   //从Que队列中获取放入的 input0_local
}
if(HasTensorInQue(qIDVecIn1)){
auto input1_local = DeQue<half>(qIDVecIn1);   //从Que队列中获取放入的 input1_local
}

Add(output_local, input0_local, input1_local, mask, 1, {2, 1, 1, 8, 8, 8});
// 偏移一个block，两次矢量计算可穿插写入dst
Add(output_local[8], input2_local, input3_local, mask, 1, {2, 1, 1, 8, 8, 8});

矢量计算指令接口

矢量计算指令接口能够启动AI Core中的Vector单元执行计算，指令按照由易到难分为3级至0级接口，3级最为简单，0级接口最为复杂。多层级API有助于降低负责指令的使用难度、有助于提高跨代兼容性保障、有助于保留最大灵活度可能。

三级接口运算符重载，支持+,-,*,/,|,&,^,<,>,<=,>=,!=,==等2级接口的简化表达；
二级接口针对源操作数的连续数据进行计算，并连续写入目的操作数；
零级接口功能灵活，可以进行非连续计算。

零级接口的通用参数包括：

Repeat times（迭代次数）
Block stride（单次迭代内不同block间地址步长）
Repeat stride（相邻迭代间相同block的地址步长）
Mask（控制参与运算的计算单元）

Repeat times
矢量计算单元，一次最多计算256Bytes的数据，每次读取连续的8个block数据进行计算，为了完成对输入数据的处理，必须通过多次迭代才能完成所有数据的读取与计算。
Repeat stride
相邻迭代间相同block的地址步长，对于连续计算场景，Repeat stride取值为8；对于非连续计算场景Repeat stride取值大于8时便会出现间隔（即有数据不参与计算）；对于反复计算场景，Repeat stride取值为0时对当前的8个block反复读取和计算；对于部分重复计算场景，Repeat stride取值为大于0且小于8。
Block stride
单次迭代内不同block间的步长，对于连续计算，Block stride取值为1，大于1时不同block间会出现间隔。
Mask
控制每次迭代内参与计算的元素。
连续模式：表示前面连续的多少个元素参与计算，取值范围与操作数的数据类型有关，数据类型不同，每次迭代内能够处理的元素个数最大值不同（最大值为256/sizeof(数据类型)）。
逐比特模式：可以按位控制哪些元素参与计算。

结束语

本文是对TIK C++算子开发的入门，介绍了TIK C++算子开发的基础和核函数，而对于四级接口的API最灵活的零级接口其实可以类比卷积计算。