AscendC从入门到精通系列（二）基于Kernel直调开发AscendC算子

本次主要讨论下AscendC算子的开发流程，基于Kernel直调工程的算子开发。

1 AscendC算子开发的基本流程

使用Ascend C完成Add算子核函数开发；
使用ICPU_RUN_KF CPU调测宏完成算子核函数CPU侧运行验证；
使用<<<>>>内核调用符完成算子核函数NPU侧运行验证。
在正式的开发之前，还需要先完成环境准备和算子分析工作，开发Ascend C算子的基本流程如下图所示：

2 核函数开发

本次以add_custom.cpp作为参考用例。Gitee也有对应工程和完整代码。
operator/AddCustomSample/KernelLaunch/AddKernelInvocationNeo · Ascend/samples - 码云 - 开源中国 (gitee.com)

2.1 核函数定义

首先要根据核函数定义 核函数-编程模型-Ascend C算子开发-算子开发-开发指南-CANN社区版8.0.RC3.alpha003开发文档-昇腾社区 (hiascend.com) 的规则进行核函数的定义，并在核函数中调用算子类的Init和Process函数。

// 给CPU调用
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
   
    KernelAdd op;
    op.Init(x, y, z);
    op.Process();
}


// 给NPU调用
#ifndef ASCENDC_CPU_DEBUG
void add_custom_do(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, uint8_t *z)
{
   
    add_custom<<<blockDim, nullptr, stream>>>(x, y, z);
}
#endif

2.2 算子类定义

根据矢量编程范式实现算子类，本样例中定义KernelAdd算子类，其具体成员如下：

class KernelAdd {
   
public:
    __aicore__ inline KernelAdd(){
   }
    // 初始化函数，完成内存初始化相关操作
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z){
   }
    // 核心处理函数，实现算子逻辑，调用私有成员函数CopyIn、Compute、CopyOut完成矢量算子的三级流水操作
    __aicore__ inline void Process(){
   }

private:
    // 搬入函数，完成CopyIn阶段的处理，被核心Process函数调用
    __aicore__ inline void CopyIn(int32_t progress){
   }
    // 计算函数，完成Compute阶段的处理，被核心Process函数调用
    __aicore__ inline void Compute(int32_t progress){
   }
    // 搬出函数，完成CopyOut阶段的处理，被核心Process函数调用
    __aicore__ inline void CopyOut(int32_t progress){
   }

private:
    AscendC::TPipe pipe;  //Pipe内存管理对象
    AscendC::TQue<AscendC::QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;  //输入数据Queue队列管理对象，QuePosition为VECIN
    AscendC::TQue<AscendC::QuePosition::VECOUT, BUFFER_NUM> outQueueZ;  //输出数据Queue队列管理对象，QuePosition为VECOUT
    AscendC::GlobalTensor<half> xGm;  //管理输入输出Global Memory内存地址的对象，其中xGm, yGm为输入，zGm为输出
    AscendC::GlobalTensor<half> yGm;
    AscendC::GlobalTensor<half> zGm;
};

核函数调用关系图

2.3 实现Init，CopyIn，Compute，CopyOut这个4个关键函数

Init函数初始化输入资源

__aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z)
    {
   
        xGm.SetGlobalBuffer((__gm__ half *)x + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
        yGm.SetGlobalBuffer((__gm__ half *)y + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
        zGm.SetGlobalBuffer((__gm__ half *)z + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
        pipe.InitBuffer(inQueueX, BUFFER_NUM, TILE_LENGTH * sizeof(half));
        pipe.InitBuffer(inQueueY, BUFFER_NUM, TILE_LENGTH * sizeof(half));
        pipe.InitBuffer(outQueueZ, BUFFER_NUM, TILE_LENGTH