Ascend C sqrt算子實戰-有解無憂

摘要：撰寫一個Ascend C的sqrt算子，并通過內核呼叫方式在cpu和npu模式下進行驗證，

本文分享自華為云社區《【2023 · CANN訓練營第一季】——Ascend C sqrt算子實戰》，作者：dayao，

前言

撰寫一個Ascend C的sqrt算子，并通過內核呼叫方式在cpu和npu模式下進行驗證，在訓練營沙箱環境下，cpu模式作業正常結果正確，

一、概述

先簡單回顧下TIK C++算子矢量編程的流程和實作，

矢量算子開發流程如下：

主要作業內容有：

1、算子分析：確定輸入輸出，確定數學運算式以及底層實作介面，確定核函式定義，

2、算子類的實作：實作init()和process()，init()完成記憶體初始化，實質上體現的是多核運行，和單核資料切分以及是否開啟double buffer優化；Process()實作的是CopyIn，Compute、CopyOut三個流水任務，

3、算子驗證：通過核函式的內核呼叫符的方式呼叫算子，計算出結果，并于使用相同輸入用numpy計算結果進行比對，誤差在一定范圍內即可，實際應用中，需要使用原有框架的算子進行計算精度比對，

二、算子分析

算子定義如下：假定仍是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未作修改，編譯和執行腳本run.sh只改了計算結果與真值比對部分，

一）、核函式定義

與例程相比，輸入引數只有x，

extern "C" __global__ __aicore__ void sqrt_tik2(__gm__ uint8_t* x, __gm__ uint8_t* z)
{
 KernelSqrt op;
 op.Init(x, z);
 op.Process();
}

二）、算子類

實作方式與add例程類似，init（）函式里初始化記憶體：x，y的Global Memory ；流水線任務通訊記憶體；Process（）實作流水線任務；按范式撰寫CopyIn、Compute、CopyOut，與add例程最大差異是，在compute函式中，呼叫sqrt的2類介面API實作計算，

class KernelSqrt {
public:
    __aicore__ inline KernelSqrt() {}
    __aicore__ inline void Init(__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__ inline void Process()
    {
 // 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__ inline void CopyIn(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__ inline void Compute(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__ inline void CopyOut(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;
};

三）、核函式呼叫

1、在CPU模式下，通過ICPU_RUN_KF呼叫

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

2、在NPU模式下，通過<<<>>>呼叫

#ifndef __CCE_KT_TEST__
// call of kernel function
void sqrt_tik2_do(uint32_t blockDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* z)
{
    sqrt_tik2<<<blockDim, l2ctrl, stream>>>(x, z);
}
#endif

由于<<<>>>，只能在NPU模式下呼叫，所以需要用條件編譯，不在CPU除錯模式下有效，在呼叫sqrt_tik2_do，需要按ascendcl應用編程的要求進行，

3、呼叫代碼

通過“__CCE_KT_TEST__”宏區分CPU和NPU模式，

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__
    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
 return 0;
}

四）、基準資料生成——sqrt_tik2.py

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

import numpy as np
def gen_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()函式比較算子計算與基準資料的結果，實際上由于npu模式編譯出錯，實際未執行改函式進行比較，CPU模式下，算子計算出的結果與基準golden資料完全一致，兩者的md5相同，

四、編譯運行

本次課程提供了沙箱運行環境，想個辦法把代碼搞進去，

一）、配置環境變數

二）、CPU模式

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

三）、NPU模式

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

點擊關注，第一時間了解華為云新鮮技術~

轉載請註明出處，本文鏈接：https://www.uj5u.com/qita/554453.html

標籤：其他

上一篇：讀改變未來的九大演算法筆記05_數字簽名

下一篇：返回列表