当前位置：首页 > news >正文

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

news 2025/7/7 23:22:03

本次主要讨论下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为VECINAscendC::TQue<AscendC::QuePosition::VECOUT, BUFFER_NUM> outQueueZ;  //输出数据Queue队列管理对象，QuePosition为VECOUTAscendC::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 * sizeof(half));}
Process函数中通过如下方式调用这三个：__aicore__ inline void Process(){// loop count need to be doubled, due to double bufferconstexpr int32_t loopCount = TILE_NUM * BUFFER_NUM;// tiling strategy, pipeline parallelfor (int32_t i = 0; i < loopCount; i++) {CopyIn(i);Compute(i);CopyOut(i);}}

CopyIn函数中通过如下方式调用这三个：
1、使用DataCopy接口将GlobalTensor数据拷贝到LocalTensor。
2、使用EnQue将LocalTensor放入VecIn的Queue中。

__aicore__ inline void CopyIn(int32_t progress){// alloc tensor from queue memoryAscendC::LocalTensor<half> xLocal = inQueueX.AllocTensor<half>();AscendC::LocalTensor<half> yLocal = inQueueY.AllocTensor<half>();// copy progress_th tile from global tensor to local tensorAscendC::DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH);AscendC::DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH);// enque input tensors to VECIN queueinQueueX.EnQue(xLocal);inQueueY.EnQue(yLocal);}

Compute函数实现。
1、使用DeQue从VecIn中取出LocalTensor。
2、使用Ascend C接口Add完成矢量计算。
3、使用EnQue将计算结果LocalTensor放入到VecOut的Queue中。
4、使用FreeTensor将释放不再使用的LocalTensor。

__aicore__ inline void Compute(int32_t progress)
{// deque input tensors from VECIN queueAscendC::LocalTensor<half> xLocal = inQueueX.DeQue<half>();AscendC::LocalTensor<half> yLocal = inQueueY.DeQue<half>();AscendC::LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>();// call Add instr for computationAscendC::Add(zLocal, xLocal, yLocal, TILE_LENGTH);// enque the output tensor to VECOUT queueoutQueueZ.EnQue<half>(zLocal);// free input tensors for reuseinQueueX.FreeTensor(xLocal);inQueueY.FreeTensor(yLocal);
}

CopyOut函数实现。
1、使用DeQue接口从VecOut的Queue中取出LocalTensor。
2、使用DataCopy接口将LocalTensor拷贝到GlobalTensor上。
3、使用FreeTensor将不再使用的LocalTensor进行回收。

 __aicore__ inline void CopyOut(int32_t progress)
{// deque output tensor from VECOUT queueAscendC::LocalTensor<half> zLocal = outQueueZ.DeQue<half>();// copy progress_th tile from local tensor to global tensorAscendC::DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH);// free output tensor for reuseoutQueueZ.FreeTensor(zLocal);
}

3 核函数的运行验证

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

3.1 编写CPU侧调用程序

添加图片注释，不超过 140 字（可选）

 // 使用GmAlloc分配共享内存，并进行数据初始化uint8_t* x = (uint8_t*)AscendC::GmAlloc(inputByteSize);uint8_t* y = (uint8_t*)AscendC::GmAlloc(inputByteSize);uint8_t* z = (uint8_t*)AscendC::GmAlloc(outputByteSize);ReadFile("./input/input_x.bin", inputByteSize, x, inputByteSize);ReadFile("./input/input_y.bin", inputByteSize, y, inputByteSize);// 调用ICPU_RUN_KF调测宏，完成核函数CPU侧的调用AscendC::SetKernelMode(KernelMode::AIV_MODE);ICPU_RUN_KF(add_custom, blockDim, x, y, z); // use this macro for cpu debug// 输出数据写出WriteFile("./output/output_z.bin", z, outputByteSize);// 调用GmFree释放申请的资源AscendC::GmFree((void *)x);AscendC::GmFree((void *)y);AscendC::GmFree((void *)z);

3.2 编写NPU侧运行算子的调用程序