【2023 · CANN训练营第一季】TIK C++算子开发入门第一章知识总结（上篇）

一、什么是TIK C++？

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

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

C/C++原语编程

编程模型屏蔽硬件差异，编程范式提高开发效率

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

孪生调试，CPU侧模拟NPU侧的行为，可先在CPU侧调试

二、核函数的概念和使用

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

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

2.1 如何编写核函数？

1、使用函数类型限定符

        除了需要按照C/C++函数声明的方式定义核函数之外，还要为核函数加上额外的函数类型限定符，包含__global__和__aicore__ 使用__global__函数类型限定符来标识它是一个核函数，可以被<<<...>>>调用；使用__aicore__函数类型限定符来标识该核函数在设备侧AI Core上执行：

 其他规则：

__global__ __aicore__ void kernel_name(argument list)

2、使用变量类型限定符

         指针入参变量统一的类型定义为__gm__ uint8_t* 用户可统一使用uint8_t类型的指针，并在使用时转化为实际的指针类型；亦可直接传入实际的指针类型。

必须具有void返回类型

使用extern "C"

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

2.2 如何调用核函数？

常见的C/C++函数调用方式是如下的形式：

function_name(argument list)

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

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

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

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

blockDim设置为8，表示在8个核上调用了HelloWorld核函数，每个核都会独立且并行地执行该核函数 Stream可以通过aclrtCreateStream来创建，它的作用是在当前进程或线程中显式创建一个aclrtStream argument list设置为fooDevice这1个入参

其中：

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

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

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

核函数的调用是异步的，核函数的调用结束后，控制权立刻返回给主机侧 强制主机侧程序等待所有核函数执行完毕的API（阻塞应用程序运行，直到指定Stream中的所有任务都完成，同步接口）为aclrtSynchronizeStream

三、TIK C++的helloworld样例

3.1、算子执行的不同模式

CPU模式：算子功能调试用，可以模拟在NPU上的计算行为，不需要依赖昇腾设备

NPU模式：算子功能/性能调试，可以使用NPU的强大算力进行运算加速

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

#ifdef __CCE_KT_TEST__ 表示在CPU模式下会编译该段代码

#ifndef __CCE_KT_TEST__ 表示在NPU模式下会编译该段代码

 上述步骤：

3.2、样例展示及各模块功能

# 设置环境变量

install_path=/usr/local/Ascend/ascend-toolkit/latest
source ${install_path}/bin/setenv.bash
或 source ${install_path}/../set_env.sh

(根据不同的芯片版本而异)

CPU模式的流程命令（# 编译命令）

g++ hello_world.cpp                                             \
    -D__CCE_KT_TEST__=1 -D__CCE_AICORE__=100 -D__DAV_C100__     \
    -I${install_path}/tools/tikicpulib/lib/include              \
    -L${install_path}/toolkit/tools/simulator/Ascend910A/lib    \
    -L${install_path}/tools/tikicpulib/lib -ltikicpulib_stubreg \
    -L${install_path}/tools/tikicpulib/lib/Ascend910 -lstdc++   \
    -O2 -std=c++17 -o hello_world_cpu

CPU模式的流程命令（# 执行命令）

export LD_LIBRARY_PATH=${install_path}/tools/tikicpulib/lib:$LD_LIBRARY_PATH
./hello_world_cpu

CPU模式的流程命令执行结果：

 NPU模式的流程命令（# 编译命令）

ccec -x cce hello_world.cpp             \
     --cce-aicore-arch=dav-c100         \
     -I${install_path}/acllib/include   \
     -L${install_path}/atc/lib64        \
     -lruntime -lascendcl -lstdc++      \
     -O2 -std=c++17 -o hello_world_npu

NPU模式的流程命令（# 执行命令）

./hello_world_npu

NPU模式的流程命令执行结果