TIKC++算子开发入门

一. 了解TIKC++基本概念

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

使用TIK C++开发自定义算子的优势∶

• C/C++原语编程
• 编程模型屏蔽硬件差异，编程范式提高开发效率·多层级API封装，从简单到灵活，兼顾易用与高效
• 孪生调试，CPU侧模拟NPU侧的行为，可先在CPU侧调试

二. 掌握核函数的使用

1 核函数定义

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

2 编写核函数

使用函数类型限定符

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

使用__global__函数类型限定符来标识它是一个核函数，可以被<<<…>>>调用;使用_aicore__函数类型限定符来标识该核函数在设备侧AI Core上执行:

__global__ __aicore__ void kernel_name( argument list);

使用变量类型限定符
为了方便:指针入参变量统一的类型定义为_gm_ uint8_t*
用户可统一使用uint8_t类型的指针，并在使用时转化为实际的指针类型;亦可直接传入实际的指针类型

2.3 其他规则
1 必须具有void返回类型
2 使用extern“c”
3 仅支持入参为指针类型或c/C++内置数据类型(原始数据类型)，如：半S0、浮点数S1、int 32_tc

3 调用核函数

核函数的调用语句是C/C++ 函数调用语句的一种扩展常见的C/C++ 函数调用方式是如下的形式：

function_name(argument list);

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

kernel_name<<<blockDim,12ctrl,stream>>>(argument list);

• blockDim，规定了核函数将会在几个核上执行，每个执行该核函数的核会被分配一个逻辑D，表现为内置变量block_idx，编号从o开始，可为不同的逻辑核定义不同的行为，可以在算子实现中使用.
• I2ctrl，保留参数，暂时设置为固定值nuilptr
• stream，类型为aclrtStream , stream是一个任务队列，应用程序通过stream来管理任务的并行

使用内核调用符<<<…>>>调用核函数︰

Hellowlorld<<<8,nullptr, stream>>>(fooDevice);

blockDim设置为8，表示在8个核上调用了Helloworld核函数，每个核都会独立且并行地执行核函数
Stream可以通过acLrtCreateStream来创建，它的作用是在当前进程或线程中显式创建一个aclrtStreamargument list设置为fooDevice这1个入参。
！！注意
核函数的调用是异步的，核函数的调用结束后，控制权立刻返回给主机侧强制主机侧程序等待所有核函数执行完毕的API

aclError aclrtSynchronizeStream(aclrtStream stream);

本小节介绍了核函数的概念、编写以及如何调用。