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