Sinh算子kernel直调实现

这次的CANN训练营讲述的是AscendC语言开发算子，先来简单的介绍下Ascend C语言：

面向算子开发场景的编程语言Ascend C，原生支持C和C++标准规范，最大化匹配用户开发习惯；通过多层接口抽象、自动并行计算、孪生调试等关键技术，极大提高算子开发效率，助力AI开发者低成本完成算子开发和模型调优部署。

使用Ascend C进行自定义算子开发的突出优势有：

C/C++原语编程

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

类库API封装，从简单到灵活，兼顾易用与高效

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

在开发一个算子之前，首先要分析一个算子的逻辑，以下就用sinh算子为例：

y = (e^x + e^(-x)) / 2

上面是sinh的算子公式，通过公式我们可以看出总共有一个输入，一个输出，为了方便计算还引用了两个tembuf临时变量。需要调用到的计算接口为幂函数Exp、取倒数Reciprocal、加法Add、标量乘法接口Muls；除此之外，还有数据搬运接口Datacopy，内存管理接口AllocTensor、FreeTensor，队列管理接口Enque和Deque。对于输入和输出数据的类型和shape，在这里我们先使用固定shape，类型的话是fp16，后面我们会不断的优化这个算子，并且也会尝试其他的算子实现，尽情期待吧。

核函数实现

核函数创建：

extern "C" __global__ __aicore__ void sinh_custom(GM_ADDR x, GM_ADDR y)

{

    KernelSinh op;

    op.Init(x, y);

    op.Process();

}

使用__global__函数类型限定符来标识它是一个核函数，可以被<<<...>>>调用；使用__aicore__函数类型限定符来标识该核函数在设备端AI Core上执行。指针入参变量需要增加变量类型限定符__gm__，表明该指针变量指向Global Memory上某处内存地址。为了统一表达，使用GM_ADDR宏来修饰入参，GM_ADDR宏定义如下：

#define GM_ADDR __gm__ uint8_t*

Kernel算子类的实现概览如下：

其中包含了有关的通道，需要用到的各种的变量，队列等要在这里声明。

基于矢量编程范式，将核函数的实现分为3个基本任务：CopyIn，Compute，CopyOut。Process函数中通过如下方式调用这三个函数。

CopyIn函数实现。

使用DataCopy接口将GlobalTensor数据拷贝到LocalTensor。

使用EnQue将LocalTensor放入VecIn的Queue中。

Compute函数实现。

使用DeQue从VecIn中取出LocalTensor。

使用Ascend C接口Add完成矢量计算。

使用EnQue将计算结果LocalTensor放入到VecOut的Queue中。

使用FreeTensor将释放不再使用的LocalTensor。

核函数运行验证

异构计算架构中，NPU（kernel侧）与CPU（host侧）是协同工作的，完成了kernel侧核函数开发后，即可编写host侧的核函数调用程序，实现从host侧的APP程序调用算子，执行计算过程。

除了上文核函数实现文件sinh_custom.cpp外，核函数的调用与验证还需要准备以下文件：

调用算子的应用程序：main.cpp。

输入数据和真值数据生成脚本文件：gen_data.py；验证输出数据和真值数据是否一致的验证脚本：verify_result.py。

编译cpu侧或npu侧运行的算子的编译工程文件：CMakeLists.txt。

编译运行算子的脚本：run.sh。

核心代码如下：

CPU侧：

NPU侧：

其他的代码需要修改的是：

还需要修改的是run.sh文件

修改完成后，就可以执行看看了。各位开发者可以参考上面的编写自己的代码。