【2023 · CANN訓練營第一季】——Ascend C sqrt算子實戰(zhàn)

前言：編寫一個Ascend C的sqrt算子，并通過內(nèi)核調(diào)用方式在cpu和npu模式下進行驗證。在訓練營沙箱環(huán)境下，cpu模式工作正常結果正確，npu模式下編譯報錯，以后有機會再研究。

一、概述

先簡單回顧下TIK C++算子矢量編程的流程和實現(xiàn)。

矢量算子開發(fā)流程如下：

主要工作內(nèi)容有：

1、算子分析：確定輸入輸出，確定數(shù)學表達式以及底層實現(xiàn)接口，確定核函數(shù)定義。

2、算子類的實現(xiàn)：實現(xiàn)init()和process()。init()完成內(nèi)存初始化，實質(zhì)上體現(xiàn)的是多核運行，和單核數(shù)據(jù)切分以及是否開啟double buffer優(yōu)化；Process()實現(xiàn)的是CopyIn，Compute、CopyOut三個流水任務。

3、算子驗證：通過核函數(shù)的內(nèi)核調(diào)用符的方式調(diào)用算子，計算出結果，并于使用相同輸入用numpy計算結果進行比對，誤差在一定范圍內(nèi)即可。實際應用中，需要使用原有框架的算子進行計算精度比對。

二、算子分析

算子定義如下：假定仍是8個邏輯核。


?????????查詢TIK C++的API可知，可以使用（TIK C++ API/矢量計算/單目/Sqrt，采用2級接口）完成運算，得到最終結果。

三、代碼分析

?????????直接在訓練營課程提供的add_tik2算子工程上修改。代碼地址：https://gitee.com/zgx950813/samples/tree/master/tik2_demo/kernel_samples/kernel_add_sample

????????修改代碼目錄結構如下：CMakeLists.txt和data_utils.h未作修改，編譯和執(zhí)行腳本run.sh只改了計算結果與真值比對部分。

一）、核函數(shù)定義

與例程相比，輸入?yún)?shù)只有x。

extern"C"__global__ __aicore__voidsqrt_tik2(__gm__ uint8_t*x,__gm__ uint8_t*z)

{

KernelSqrt op;

op.Init(x,z);

op.Process();

}

二）、算子類

????????實現(xiàn)方式與add例程類似。init（）函數(shù)里初始化內(nèi)存：x，y的Global Memory ；流水線任務通訊內(nèi)存；Process（）實現(xiàn)流水線任務；按范式編寫CopyIn、Compute、CopyOut。與add例程最大差異是，在compute函數(shù)中，調(diào)用sqrt的2類接口API實現(xiàn)計算。

classKernelSqrt{

public:

__aicore__ inlineKernelSqrt(){}

__aicore__ inlinevoidInit(__gm__ uint8_t*x,__gm__ uint8_t*z)

{

// get start index for current core, core parallel

xGm.SetGlobalBuffer((__gm__ half*)x+block_idx*BLOCK_LENGTH,BLOCK_LENGTH);

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

// pipe alloc memory to queue, the unit is Bytes

pipe.InitBuffer(inQueueX,BUFFER_NUM,TILE_LENGTH*sizeof(half));

pipe.InitBuffer(outQueueZ,BUFFER_NUM,TILE_LENGTH*sizeof(half));

}

__aicore__ inlinevoidProcess()

{

// loop count need to be doubled, due to double buffer

constexpr int32_t loopCount=TILE_NUM*BUFFER_NUM;

// tiling strategy, pipeline parallel

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

CopyIn(i);

Compute(i);

CopyOut(i);

}

}

private:

__aicore__ inlinevoidCopyIn(int32_t progress)

{

// alloc tensor from queue memory

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

// copy progress_th tile from global tensor to local tensor

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

// enque input tensors to VECIN queue

inQueueX.EnQue(xLocal);

}

__aicore__ inlinevoidCompute(int32_t progress)

{

// deque input tensors from VECIN queue

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

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

// call Sqrt instr for computation

Sqrt(zLocal,xLocal,TILE_LENGTH);

// enque the output tensor to VECOUT queue

outQueueZ.EnQue<half>(zLocal);

// free input tensors for reuse

inQueueX.FreeTensor(xLocal);

}

__aicore__ inlinevoidCopyOut(int32_t progress)

{

// deque output tensor from VECOUT queue

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

// copy progress_th tile from local tensor to global tensor

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

// free output tensor for reuse

outQueueZ.FreeTensor(zLocal);

}

private:

TPipe pipe;

// create queues for input, in this case depth is equal to buffer num

TQue<QuePosition::VECIN,BUFFER_NUM>inQueueX;

// create queue for output, in this case depth is equal to buffer num

TQue<QuePosition::VECOUT,BUFFER_NUM>outQueueZ;

GlobalTensor<half>xGm,zGm;

};

三）、核函數(shù)調(diào)用

1、在CPU模式下，通過ICPU_RUN_KF調(diào)用

ICPU_RUN_KF(sqrt_tik2,blockDim,x,z);// use this macro for cpu debug

2、在NPU模式下，通過<<<>>>調(diào)用

#ifndef __CCE_KT_TEST__

// call of kernel function

voidsqrt_tik2_do(uint32_t blockDim,void*l2ctrl,void*stream,uint8_t*x,uint8_t*z)

{

sqrt_tik2<<<blockDim,l2ctrl,stream>>>(x,z);

}

#endif

由于<<<>>>，只能在NPU模式下調(diào)用，所以需要用條件編譯，不在CPU調(diào)試模式下有效。在調(diào)用sqrt_tik2_do，需要按ascendcl應用編程的要求進行。

3、調(diào)用代碼

????????通過“__CCE_KT_TEST__”宏區(qū)分CPU和NPU模式。

int32_tmain(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__

uint8_t*x=(uint8_t*)tik2::GmAlloc(inputByteSize);

uint8_t*z=(uint8_t*)tik2::GmAlloc(outputByteSize);

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

// PrintData(x, 16, printDataType::HALF);

ICPU_RUN_KF(sqrt_tik2,blockDim,x,z);// use this macro for cpu debug

// PrintData(z, 16, printDataType::HALF);

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

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

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

#else

aclInit(nullptr);

aclrtContext context;

aclError error;

int32_t deviceId=0;

aclrtCreateContext(&context,deviceId);

aclrtStream stream=nullptr;

aclrtCreateStream(&stream);

uint8_t*xHost,*zHost;

uint8_t*xDevice,*zDevice;

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

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

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

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

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

// PrintData(xHost, 16, printDataType::HALF);

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

sqrt_tik2_do(blockDim,nullptr,stream,xDevice,zDevice);// call kernel in this function

aclrtSynchronizeStream(stream);

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

// PrintData(zHost, 16, printDataType::HALF);

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

aclrtFree(xDevice);

aclrtFree(zDevice);

aclrtFreeHost(xHost);

aclrtFreeHost(zHost);

aclrtDestroyStream(stream);

aclrtResetDevice(deviceId);

aclFinalize();

#endif

return0;

}

四）、基準數(shù)據(jù)生成——sqrt_tik2.py

使用numpy生成input_x和基準結果golden。

importnumpyasnp

defgen_golden_data_simple():

input_x=np.random.uniform(0,100,[8,2048]).astype(np.float16)

golden=np.sqrt(input_x).astype(np.float16)

input_x.tofile("./input/input_x.bin")

golden.tofile("./output/golden.bin")

if__name__=="__main__":

gen_golden_data_simple()

五）、計算結果比較

????????使用numpy的allclose()函數(shù)比較算子計算與基準數(shù)據(jù)的結果。實際上由于npu模式編譯出錯，實際未執(zhí)行改函數(shù)進行比較。CPU模式下，算子計算出的結果與基準golden數(shù)據(jù)完全一致，兩者的md5相同。

四、編譯運行

????????本次課程提供了沙箱運行環(huán)境，想個辦法把代碼搞進去。

一）、配置環(huán)境變量

二）、CPU模式

????????cpu模式順利編譯運行，結果與對比組完全一致。

二）、NPU模式

????????npu模式下編譯報錯，因為沙箱時間有限，以后有機會再研究。

?著作權歸作者所有,轉(zhuǎn)載或內(nèi)容合作請聯(lián)系作者
【社區(qū)內(nèi)容提示】社區(qū)部分內(nèi)容疑似由AI輔助生成，瀏覽時請結合常識與多方信息審慎甄別。
平臺聲明：文章內(nèi)容（如有圖片或視頻亦包括在內(nèi)）由作者上傳并發(fā)布，文章內(nèi)容僅代表作者本人觀點，簡書系信息發(fā)布平臺，僅提供信息存儲服務。

相關閱讀更多精彩內(nèi)容

友情鏈接更多精彩內(nèi)容