【2023 · CANN训练营第一季】- TIK C++算子开发入门 第一章 学习笔记

一、学习目标

1. 了解TIK C++基本概念

2. 掌握核函数的使用

3. 编写TIK C++的helloworld样例

二、TIK C++介绍

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

使用TIKC++开发自定义算子的优势:

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

三、核函数

1.介绍

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

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

2.编写核函数

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

使用__global__函数类型限定符来标识它是一个核函数，可以被<<<...>>>调用；使用__aicore__函数类型限定符来标识该核函数在设备端aicore上执行：

 （2）指针入参变量统一的类型定义为__gm__ uint8_t*，这里统一使用uint8_t类型的指针，在后续的使用中需要将其转化为实际的指针类型；用户亦可直接传入实际的指针类型。

（3）

• 必须具有void返回类型。
• 使用extern "C"。
• 仅支持入参为指针或C/C++内置数据类型（Primitive data types），如：half* s0、float* s1、int32_t c。

4. 调用核函数

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

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

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

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

aclError aclrtSynchronizeStream(aclrtStream stream);

四、样例演示

1.算子执行的不同模式

TIK C++算子可用CPU模式或NPU模式执行

• CPU模式:算子功能调试用，可以模拟在NPU上的计算行为，不需要依赖异腾设备
• NPU模式:算子功能/性能调试，可以使用NPU的强大算力进行运算加速

使用内置宏CCE_KTTEST 标识被宏包括的代码在特定的模式下编译

• #ifdef_CCE_KT_TEST 表示在CPU模式下会编译该段代码
• #ifndef_CCE_KT_TEST表示在NPU模式下会编译该段代码
• 

 2.样例演示