本次主要討論下AscendC算子的開發(fā)流程，基于Kernel直調工程的算子開發(fā)。

1 AscendC算子開發(fā)的基本流程

使用Ascend C完成Add算子核函數(shù)開發(fā)；
使用ICPU_RUN_KF CPU調測宏完成算子核函數(shù)CPU側運行驗證；
使用<<<>>>內核調用符完成算子核函數(shù)NPU側運行驗證。
在正式的開發(fā)之前，還需要先完成環(huán)境準備和算子分析工作，開發(fā)Ascend C算子的基本流程如下圖所示：

image.png

2 核函數(shù)開發(fā)

本次以add_custom.cpp作為參考用例。Gitee也有對應工程和完整代碼。
operator/AddCustomSample/KernelLaunch/AddKernelInvocationNeo · Ascend/samples - 碼云 - 開源中國 (gitee.com)

2.1 核函數(shù)定義

首先要根據(jù)核函數(shù)定義核函數(shù)-編程模型-Ascend C算子開發(fā)-算子開發(fā)-開發(fā)指南-CANN社區(qū)版8.0.RC3.alpha003開發(fā)文檔-昇騰社區(qū) (hiascend.com) 的規(guī)則進行核函數(shù)的定義，并在核函數(shù)中調用算子類的Init和Process函數(shù)。

// 給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 算子類定義

根據(jù)矢量編程范式實現(xiàn)算子類，本樣例中定義KernelAdd算子類，其具體成員如下：

class KernelAdd {
public:
    __aicore__ inline KernelAdd(){}
    // 初始化函數(shù)，完成內存初始化相關操作
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z){}
    // 核心處理函數(shù)，實現(xiàn)算子邏輯，調用私有成員函數(shù)CopyIn、Compute、CopyOut完成矢量算子的三級流水操作
    __aicore__ inline void Process(){}

private:
    // 搬入函數(shù)，完成CopyIn階段的處理，被核心Process函數(shù)調用
    __aicore__ inline void CopyIn(int32_t progress){}
    // 計算函數(shù)，完成Compute階段的處理，被核心Process函數(shù)調用
    __aicore__ inline void Compute(int32_t progress){}
    // 搬出函數(shù)，完成CopyOut階段的處理，被核心Process函數(shù)調用
    __aicore__ inline void CopyOut(int32_t progress){}

private:
    AscendC::TPipe pipe;  //Pipe內存管理對象
    AscendC::TQue<AscendC::QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;  //輸入數(shù)據(jù)Queue隊列管理對象，QuePosition為VECIN
    AscendC::TQue<AscendC::QuePosition::VECOUT, BUFFER_NUM> outQueueZ;  //輸出數(shù)據(jù)Queue隊列管理對象，QuePosition為VECOUT
    AscendC::GlobalTensor<half> xGm;  //管理輸入輸出Global Memory內存地址的對象，其中xGm, yGm為輸入，zGm為輸出
    AscendC::GlobalTensor<half> yGm;
    AscendC::GlobalTensor<half> zGm;
};

核函數(shù)調用關系圖

image.png

2.3 實現(xiàn)Init，CopyIn，Compute，CopyOut這個4個關鍵函數(shù)

Init函數(shù)初始化輸入資源

__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函數(shù)中通過如下方式調用這三個：
    __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);
        }
    }

CopyIn函數(shù)中通過如下方式調用這三個：
1、使用DataCopy接口將GlobalTensor數(shù)據(jù)拷貝到LocalTensor。
2、使用EnQue將LocalTensor放入VecIn的Queue中。

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

Compute函數(shù)實現(xiàn)。
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 queue
    AscendC::LocalTensor<half> xLocal = inQueueX.DeQue<half>();
    AscendC::LocalTensor<half> yLocal = inQueueY.DeQue<half>();
    AscendC::LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>();
    // call Add instr for computation
    AscendC::Add(zLocal, xLocal, yLocal, TILE_LENGTH);
    // enque the output tensor to VECOUT queue
    outQueueZ.EnQue<half>(zLocal);
    // free input tensors for reuse
    inQueueX.FreeTensor(xLocal);
    inQueueY.FreeTensor(yLocal);
}

CopyOut函數(shù)實現(xiàn)。
1、使用DeQue接口從VecOut的Queue中取出LocalTensor。
2、使用DataCopy接口將LocalTensor拷貝到GlobalTensor上。
3、使用FreeTensor將不再使用的LocalTensor進行回收。

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

3 核函數(shù)的運行驗證

異構計算架構中，NPU（kernel側）與CPU（host側）是協(xié)同工作的，完成了kernel側核函數(shù)開發(fā)后，即可編寫host側的核函數(shù)調用程序，實現(xiàn)從host側的APP程序調用算子，執(zhí)行計算過程。

3.1 編寫CPU側調用程序

image.png

 // 使用GmAlloc分配共享內存，并進行數(shù)據(jù)初始化
    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調測宏，完成核函數(shù)CPU側的調用
    AscendC::SetKernelMode(KernelMode::AIV_MODE);
    ICPU_RUN_KF(add_custom, blockDim, x, y, z); // use this macro for cpu debug
    // 輸出數(shù)據(jù)寫出
    WriteFile("./output/output_z.bin", z, outputByteSize);
    // 調用GmFree釋放申請的資源
    AscendC::GmFree((void *)x);
    AscendC::GmFree((void *)y);
    AscendC::GmFree((void *)z);

3.2 編寫NPU側運行算子的調用程序

image.png

  // AscendCL初始化
    CHECK_ACL(aclInit(nullptr));
    // 運行管理資源申請
    int32_t deviceId = 0;
    CHECK_ACL(aclrtSetDevice(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內存
    CHECK_ACL(aclrtMalloc((void**)&xDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
    CHECK_ACL(aclrtMalloc((void**)&yDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
    CHECK_ACL(aclrtMalloc((void**)&zDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
    // Host內存初始化
    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));
    // 用內核調用符<<<>>>調用核函數(shù)完成指定的運算,add_custom_do中封裝了<<<>>>調用
    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));
    // AscendCL去初始化
    CHECK_ACL(aclrtDestroyStream(stream));
    CHECK_ACL(aclrtResetDevice(deviceId));
    CHECK_ACL(aclFinalize());

3.3 完整main.cpp

/**
 * @file main.cpp
 *
 * Copyright (C) 2024. Huawei Technologies Co., Ltd. All rights reserved.
 *
 * This program is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
 */
#include "data_utils.h"
#ifndef ASCENDC_CPU_DEBUG
#include "acl/acl.h"
extern void add_custom_do(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, uint8_t *z);
#else
#include "tikicpulib.h"
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z);
#endif

int32_t main(int32_t argc, char *argv[])
{
    uint32_t blockDim = 8;
    size_t inputByteSize = 8 * 2048 * sizeof(uint16_t);
    size_t outputByteSize = 8 * 2048 * sizeof(uint16_t);

#ifdef ASCENDC_CPU_DEBUG
    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);
#else
    CHECK_ACL(aclInit(nullptr));
    int32_t deviceId = 0;
    CHECK_ACL(aclrtSetDevice(deviceId));
    aclrtStream stream = nullptr;
    CHECK_ACL(aclrtCreateStream(&stream));

    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));
    CHECK_ACL(aclrtMalloc((void **)&xDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
    CHECK_ACL(aclrtMalloc((void **)&yDevice, inputByteSize, 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, stream, xDevice, yDevice, zDevice);
    CHECK_ACL(aclrtSynchronizeStream(stream));

    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(aclrtResetDevice(deviceId));
    CHECK_ACL(aclFinalize());
#endif
    return 0;
}

整體運行起來，請參考operator/AddCustomSample/KernelLaunch/AddKernelInvocationNeo · Ascend/samples - 碼云 - 開源中國 (gitee.com)

色偷偷精品伊人,欧洲久久精品,欧美综合婷婷骚逼,国产AV主播,国产最新探花在线,九色在线视频一区,伊人大交九欧美,1769亚洲,黄色成人av

AscendC從入門到精通系列（二）基于Kernel直調開發(fā)AscendC算子

AscendC從入門到精通系列（二）基于Kernel直調開發(fā)AscendC算子

1 AscendC算子開發(fā)的基本流程

2 核函數(shù)開發(fā)

2.1 核函數(shù)定義

2.2 算子類定義

2.3 實現(xiàn)Init，CopyIn，Compute，CopyOut這個4個關鍵函數(shù)

3 核函數(shù)的運行驗證

3.1 編寫CPU側調用程序

3.2 編寫NPU側運行算子的調用程序

3.3 完整main.cpp

相關閱讀更多精彩內容

友情鏈接更多精彩內容

色偷偷精品伊人,欧洲久久精品,欧美综合婷婷骚逼,国产AV主播,国产最新探花在线,九色在线视频一区,伊人大交九 欧美,1769亚洲,黄色成人av

AscendC從入門到精通系列（二）基于Kernel直調開發(fā)AscendC算子

1 AscendC算子開發(fā)的基本流程

2 核函數(shù)開發(fā)

2.1 核函數(shù)定義

2.2 算子類定義

2.3 實現(xiàn)Init，CopyIn，Compute，CopyOut這個4個關鍵函數(shù)

3 核函數(shù)的運行驗證

3.1 編寫CPU側調用程序

3.2 編寫NPU側運行算子的調用程序

3.3 完整main.cpp

相關閱讀更多精彩內容

友情鏈接更多精彩內容

色偷偷精品伊人,欧洲久久精品,欧美综合婷婷骚逼,国产AV主播,国产最新探花在线,九色在线视频一区,伊人大交九欧美,1769亚洲,黄色成人av

2.3 實現(xiàn)Init，CopyIn，Compute，CopyOut這個4個關鍵函數(shù)