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();

}

#define GM_ADDR __gm__ uint8_t* __restrict__

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:

TPipe pipe; //Pipe内存管理对象

TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY; //输入数据Queue队列管理对象，QuePosition为VECIN

TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueZ; //输出数据Queue队列管理对象，QuePosition为VECOUT

GlobalTensor<half> xGm, yGm, zGm; //管理输入输出Global Memory内存地址的对象，其中xGm, yGm为输入，zGm为输出

};

__aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z)

{

// 多核并行，设定当前核上Global Memory的起始地址以及长度 xGm.SetGlobalBuffer((__gm__ half*)x + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);

yGm.SetGlobalBuffer((__gm__ half*)y + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);

zGm.SetGlobalBuffer((__gm__ half*)z + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);

// 通过pipe为queue分配内存，单位为Bytes

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));

}

__aicore__ inline void Process()

{

// 开启double buffer后循环次数需要乘以2

constexpr int32_t loopCount = TILE_NUM * BUFFER_NUM;

// 多个任务实现流水并行

for (int32_t i = 0; i < loopCount; i++) {

CopyIn(i);

Compute(i);

CopyOut(i);

}

}

__aicore__ inline void CopyIn(int32_t progress)

{

// 1、从队列中分配Tensor

LocalTensor<half> xLocal = inQueueX.AllocTensor<half>();

LocalTensor<half> yLocal = inQueueY.AllocTensor<half>();

// 2、使用DataCopy接口将GlobalTensor数据拷贝到LocalTensor

DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH);

DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH);

// 3、将LocalTensor放入搬入数据的存放位置VecIn的Queue中

inQueueX.EnQue(xLocal);

inQueueY.EnQue(yLocal);

}

__aicore__ inline void Compute(int32_t progress)

{

// 1、使用DeQue从VecIn中取出LocalTensor

LocalTensor<half> xLocal = inQueueX.DeQue<half>();

LocalTensor<half> yLocal = inQueueY.DeQue<half>();

LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>();

// 2、调用Add指令完成双目矢量计算

Add(zLocal, xLocal, yLocal, TILE_LENGTH);

// 3、使用EnQue将计算结果LocalTensor放入到搬出数据的存放位置VECOUT的Queue中

outQueueZ.EnQue<half>(zLocal);

// 4、使用FreeTensor将释放不再使用的LocalTensor

inQueueX.FreeTensor(xLocal);

inQueueY.FreeTensor(yLocal);

}

__aicore__ inline void CopyOut(int32_t progress)

{

// 1、使用DeQue接口从VecOut的Queue中取出LocalTensor

LocalTensor<half> zLocal = outQueueZ.DeQue<half>();

// 2、使用DataCopy接口将LocalTensor拷贝到GlobalTensor上

DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH);

// 3、使用FreeTensor将不再使用的LocalTensor进行回收

outQueueZ.FreeTensor(zLocal);

}

int32_t main(int32_t argc, char* argv[])

{

size_t inputByteSize = 8 * 2048 * sizeof(uint16_t); // uint16_t represent half

size_t outputByteSize = 8 * 2048 * sizeof(uint16_t); // uint16_t represent half

uint32_t blockDim = 8;

#ifdef __CCE_KT_TEST__

// 用于CPU调试的调用程序



#else

// NPU侧运行算子的调用程序

#endif

return 0;

}

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);

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);

AscendC::GmFree((void *)x);

AscendC::GmFree((void *)y);

AscendC::GmFree((void *)z);

// AscendCL初始化

CHECK_ACL(aclInit(nullptr));

// 创建Context绑定设备

aclrtContext context;

int32_t deviceId = 0;

CHECK_ACL(aclrtSetDevice(deviceId));

CHECK_ACL(aclrtCreateContext(&context, deviceId));

aclrtStream stream = nullptr;

CHECK_ACL(aclrtCreateStream(&stream));

// 分配Host内存，并进行数据初始化

uint8_t *xHost, *yHost, *zHost;

uint8_t *xDevice, *yDevice, *zDevice;

CHECK_ACL(aclrtMallocHost((void**)(&xHost), inputByteSize));

CHECK_ACL(aclrtMallocHost((void**)(&yHost), inputByteSize));

CHECK_ACL(aclrtMallocHost((void**)(&zHost), outputByteSize));

// 分配Device内存，并将数据从Host上拷贝到Device上

CHECK_ACL(aclrtMalloc((void**)&xDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));

CHECK_ACL(aclrtMalloc((void**)&yDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));

CHECK_ACL(aclrtMalloc((void**)&zDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));

ReadFile("./input/input_x.bin", inputByteSize, xHost, inputByteSize);

ReadFile("./input/input_y.bin", inputByteSize, yHost, inputByteSize);

CHECK_ACL(aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE));

CHECK_ACL(aclrtMemcpy(yDevice, inputByteSize, yHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE));

// 用内核调用符<<<>>>调用核函数完成指定的运算

add_custom_do(blockDim, nullptr, stream, xDevice, yDevice, zDevice);

CHECK_ACL(aclrtSynchronizeStream(stream));

// 将Device上的运算结果拷贝回Host

CHECK_ACL(aclrtMemcpy(zHost, outputByteSize, zDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST));

WriteFile("./output/output_z.bin", zHost, outputByteSize);

// 释放申请的资源

CHECK_ACL(aclrtFree(xDevice));

CHECK_ACL(aclrtFree(yDevice));

CHECK_ACL(aclrtFree(zDevice));

CHECK_ACL(aclrtFreeHost(xHost));

CHECK_ACL(aclrtFreeHost(yHost));

CHECK_ACL(aclrtFreeHost(zHost));

CHECK_ACL(aclrtDestroyStream(stream));

CHECK_ACL(aclrtDestroyContext(context));

CHECK_ACL(aclrtResetDevice(deviceId));

CHECK_ACL(aclFinalize());