2023年5月6日，在昇騰AI開發者峰會上，華為正式發布了面向算子開發場景的昇騰Ascend C編程語言。Ascend C原生支持C/C++編程規范，通過多層接口抽象、并行編程范式、孿生調試等技術，極大提高了算子的開發效率，幫助AI開發者低成本完成算子開發和模型調優部署。

昇騰AI軟硬件基礎

和CUDA開發的算子運行在GPU上一樣，基于Ascend C開發的算子，可以通過異構計算架構CANN（Compute Architecture for Neural Networks）運行在昇騰AI處理器（可簡稱NPU）上。CANN是使能昇騰AI處理器的一個軟件棧，通過軟硬件協同優化，能夠充分發揮昇騰AI處理器的強大算力。從下面的架構圖可以清楚的看到，使用Ascend C編程語言開發的算子通過編譯器編譯和運行時調度，最終運行在昇騰AI處理器上。

我們知道，通用計算就是我們常寫的一些在CPU上運行的計算，它擅長邏輯控制和串行計算，而AI計算相對通用計算來說，更擅長并行計算，可支持大規模的計算密集型任務。如下面左圖所示，做一個矩陣乘，使用CPU計算需要三層for循環，而右圖在昇騰AI處理器上使用vector計算單元，只需要兩層for循環，最小計算代碼能同時計算多個數據的乘加，更近一步，如果使用Cube計算單元，只需要一條語句就能完成一個矩陣乘的計算，這就是我們所說的SIMD（單指令多數據）。因此，我們通常使用AI處理器來進行大量的并行計算。

NPU不能獨立運行，需要與CPU協同工作，可以看成是CPU的協處理器，CPU負責整個操作系統運行，管理各類資源并進行復雜的邏輯控制，而NPU主要負責并行計算任務。在基于CPU+NPU的異構計算架構中，NPU與CPU通過PCIe總線連接在一起來協同工作，CPU所在位置稱為主機端（host），而NPU所在位置稱為設備端（device），示意圖如下：

這里再詳細介紹一下昇騰AI處理器。昇騰AI處理器有不同的型號和產品形態，小到模塊、加速卡，大到服務器、集群。昇騰AI處理器里面最核心的部件是AI Core，有多個，是神經網絡加速的計算核心，每一個AI Core就相當于我們大家平時理解的多核cpu里的每個核，使用Ascend C編程語言開發的算子就運行在AI Core上，因為核心的神經網絡計算的加速都來源于AI Core的算力。

AI Core內部的并行計算架構抽象如下圖所示：

這個并行計算架構抽象核心包含了幾個大的部件，AI Core外面有一個Gobal Memory，是多個AI Core共享的，在AI Core內部有一塊本地內存Local Memory，因為靠近計算單元，所以它的帶寬會非常高，相對的容量就會很小，比如一般是幾百K到1M。AI Core內部的核心組件有三個計算單元，標量計算單元、向量計算單元，矩陣計算單元。另外還有一個DMA搬運單元，DMA搬運單元負責在Global Memory和Local Memory之間搬運數據。

AI Core內部的異步并行計算過程：Scalar計算單元讀取指令序列，并把向量計算、矩陣計算、數據搬運指令發射給對應單元的指令隊列，向量計算單元、矩陣計算單元、數據搬運單元異步并行執行接收到的指令。該過程可以參考上圖中藍色箭頭所示的指令流。不同的指令間有可能存在依賴關系，為了保證不同指令隊列間的指令按照正確的邏輯關系執行，Scalar計算單元也會給對應單元下發同步指令。各單元之間的同步過程可以參考上圖中的橙色箭頭所示的同步信號流。

AI Core內部數據處理的基本過程：DMA搬入單元把數據搬運到Local Memory，Vector/Cube計算單元完成數據，并把計算結果寫回Local Memory，DMA搬出單元把處理好的數據搬運回Global Memory。該過程可以參考上圖中的紅色箭頭所示的數據流。

Ascend C編程模型基礎

Ascend C編程范式

Ascend C編程范式是一種流水線式的編程范式，把算子核內的處理程序，分成多個流水任務，通過隊列（Queue）完成任務間通信和同步，并通過統一的內存管理模塊（Pipe）管理任務間通信內存。流水編程范式應用了流水線并行計算方法。

若n=3，即待處理的數據被切分成3片，則上圖中的流水任務運行起來的示意圖如下，從運行圖中可以看出，對于同一片數據，Stage1、Stage2、Stage3之間的處理具有依賴關系，需要串行處理；不同的數據切片，同一時間點，可以有多個任務在并行處理，由此達到任務并行、提升性能的目的。

Ascend C分別針對Vector、Cube編程設計了不同的流水任務。開發者只需要完成基本任務的代碼實現即可，底層的指令同步和并行調度由Ascend C框架實現，開發者無需關注。

矢量編程范式

矢量編程范式把算子的實現流程分為3個基本任務：CopyIn，Compute，CopyOut。CopyIn負責搬入操作，Compute負責矢量計算操作，CopyOut負責搬出操作。

我們只需要根據編程范式完成基本任務的代碼實現就可以了，底層的指令同步和并行調度由Ascend C框架來實現。

那Ascend C是怎么完成不同任務之間的數據通信和同步的呢？這里Ascend C提供了Queue隊列管理的API，主要就是兩個隊列操作API EnQue、DeQue以及內存的邏輯抽象。

矢量編程中使用到的邏輯位置（QuePosition）定義如下：

• 搬入數據的存放位置：VECIN；
• 計算中間變量的位置：VECCALC；
• 搬出數據的存放位置：VECOUT。

從前面可以看到，矢量編程主要分為CopyIn、Compute、CopyOut三個任務。CopyIn任務中將輸入數據從Global內存搬運至Local內存后，需要使用EnQue將LocalTensor放入VECIN的Queue中；Compute任務等待VECIN的Queue中LocalTensor出隊之后才可以完成矢量計算，計算完成后使用EnQue將計算結果LocalTensor放入到VECOUT的Queue中；CopyOut任務等待VECOUT的Queue中LocalTensor出隊，再將其拷貝到Global內存。這樣 ，Queue隊列就完成了三個任務間的數據通信和同步。具體流程和流程圖如下：

1. Stage1：CopyIn任務。使用DataCopy接口將GlobalTensor數據拷貝到LocalTensor。
   使用EnQue接口將LocalTensor放入VECIN的Queue中。
2. Stage2：Compute任務。使用DeQue接口從VECIN中取出LocalTensor。
   使用Ascend C接口完成矢量計算。
   使用EnQue接口將計算結果LocalTensor放入到VECOUT的Queue中。
3. Stage3：CopyOut任務。

使用DeQue接口從VECOUT的Queue中去除LocalTensor。

使用DataCopy接口將LocalTensor拷貝到GlobalTensor上。

這樣我們的kernel實現代碼就很清晰了。先初始化內存和隊列，然后通過編程范式實現CopyIn、Compute、CopyOut三個Stage就可以了。

SPMD并行編程-多核

最前面介紹昇騰AI處理器的時候，有介紹過AI Core是有多個的，那我們怎么把多個AI Core充分利用起來呢？常用的并行計算方法中，有一種SPMD（Single-Program Multiple-Data）數據并行的方法，簡單說就是將數據分片，每片數據經過完整的一個數據處理流程。這個就能和昇騰AI處理器的多核匹配上了，我們將數據分成多份，每份數據的處理運行在一個核上，這樣每份數據并行處理完成，整個數據也就處理完了。Ascend C是SPMD（Single-Program Multiple-Data）編程，多個AI Core共享相同的指令代碼，每個核上的運行實例唯一的區別是就是block_idx（內置變量）不同，這樣我們就可以通過block_idx來區分不同的核，只要對Global Memory上的數據地址進行切分偏移，就可以讓每個核處理自己對應的那部分數據了。

算子被調用時，所有的計算核心都執行相同的實現代碼，入口函數的入參也是相同的。每個核上處理的數據地址需要在起始地址上增加block_idx*BLOCK_LENGTH（每個block處理的數據長度）的偏移來獲取。這樣也就實現了多核并行計算的數據切分。

class KernelAdd {

public:

__aicore__ inline KernelAdd() {}

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

{

// get start index for current core, core parallel

GM_ADDR xGmOffset = x + BLOCK_LENGTH * GetBlockIdx();

GM_ADDR yGmOffset = y + BLOCK_LENGTH * GetBlockIdx();

GM_ADDR zGmOffset = z + BLOCK_LENGTH * GetBlockIdx();

xGm.SetGlobalBuffer((__gm__ half*)xGmOffset, BLOCK_LENGTH);

yGm.SetGlobalBuffer((__gm__ half*)yGmOffset, BLOCK_LENGTH);

zGm.SetGlobalBuffer((__gm__ half*)zGmOffset, BLOCK_LENGTH);

……

}

……

}

Ascend C API介紹

在整個kernel實現中，最最核心的代碼就是Add(zLocal, xLocal, yLocal, TILE_LENGTH);通過一個Ascend C提供的API接口完成了所有數據的加法計算，對，沒看錯，就是這個接口完成了計算。

接下來就介紹下Ascend C提供的API。Ascend C算子采用標準C++語法和一組類庫API進行編程，類庫API主要包含以下幾種，大家可以在核函數的實現中根據自己的需求選擇合適的API：

• 計算類API，包括標量計算API、向量計算API、矩陣計算API，分別實現調用Scalar計算單元、Vector計算單元、Cube計算單元執行計算的功能。
• 數據搬運API，上述計算API基于Local Memory數據進行計算，所以數據需要先從Global Memory搬運至Local Memory，再使用計算接口完成計算，最后從Local Memory搬出至Global Memory。執行搬運過程的接口稱之為數據搬移接口，比如DataCopy接口。
• 內存管理API，用于分配管理內存，比如AllocTensor、FreeTensor接口。
• 任務同步API，完成任務間的通信和同步，比如EnQue、DeQue接口。

Ascend C API的計算操作數都是Tensor類型：GlobalTensor和LocalTensor。

介紹完Ascend C API種類后，下面來解釋下為什么一個Add接口就可以計算所有的數。原來Ascend C編程模型是基于SIMD（單指令多數據）架構的，單條指令可以完成多個數據操作，同時在API內部封裝了一些指令的高級功能。

算子執行基本流程

前面有提到，在異構計算架構中，NPU與CPU是協同工作的，在Ascend C編程模型中，我們需要實現NPU側的代碼和CPU側的代碼。在NPU側的代碼我們通常叫做Kernel實現代碼，CPU側的代碼我們一般叫做Host實現代碼，一份完整的Ascend C代碼，通常包括Host側實現代碼和Kernel側實現代碼。Ascend C算子執行的基本流程如下：

1. 初始化Device設備；
2. 創建Context綁定設備；
3. 分配Host內存，并進行數據初始化；
4. 分配Device內存，并將數據從Host上拷貝到Device上；
5. 用內核調用符<<<>>>調用核函數完成指定的運算；
6. 將Device上的運算結果拷貝回Host；
7. 釋放申請的資源。

核函數介紹

上面的流程中，最重要的一步就是調用核函數來進行并行計算任務。核函數（Kernel Function）是Ascend C算子Device側實現的入口。在核函數中，需要為在AI核上執行的代碼規定要進行的數據訪問和計算操作。

extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z);

上面這個是一個核函數聲明的示例，extern "C"表示核函數按照類C的編譯和連接規約來編譯和連接，__global__函數類型限定符表示它是一個核函數， __aicore__函數類型限定符表示該核函數在device側的AI Core上執行。參數列表中的變量類型限定符__gm__，表明該指針變量指向Global Memory上某處內存地址，注意這里的入參只能支持指針或C/C++內置數據類型，樣例里指針使用的類型為uint8_t，在后續的使用中需要將其轉化為實際的指針類型。

Ascend C編程模型中的核函數采用內核調用符<<<...>>>來調用，樣例如下：

kernel_name<<<blockDim, l2ctrl, stream>>>(argument list);

kernel_name即為上面講的核函數名稱，argument list是核函數的函數入參，在<<<>>>中間，有3個參數：

• blockDim，規定了核函數將會在幾個核上執行，我們可以先設置為1；
• l2ctrl，保留參數，暫時設置為固定值nullptr，我們不用關注；
• stream，使用aclrtCreateStream創建，用于多線程調度。

樣例開發講解

樣例代碼結構

|-- CMakeLists.txt //編譯工程文件

|-- cmake //編譯工程文件

|-- data_utils.h //數據讀入寫出函數

|-- input //存放腳本生成的輸入數據目錄

|-- leakyrelu_custom.cpp //算子kernel實現

|-- leakyrelu_custom.py //輸入數據和真值數據生成腳本文件

|-- leakyrelu_custom_tiling.h //host側tiling函數

|-- main.cpp //主函數，host側調用代碼，含cpu域及npu域調用

|-- output //存放算子運行輸出數據和標桿數據的目錄

|-- readme.md //執行命令說明

|-- run.sh //運行腳本

主要文件

輸入數據和真值數據生成腳本文件：KERNEL_NAME.py。

根據算子的輸入輸出編寫生成輸入數據和真值數據的腳本。

本例子生成8 * 200 * 1024大小的fp16數據：

……

def gen_golden_data_simple():

total_length_imm = 8 * 200 * 1024

tile_num_imm = 8

//生成tilling的bin文件

total_length = np.array(total_length_imm, dtype=np.uint32)

tile_num = np.array(tile_num_imm, dtype=np.uint32)

scalar = np.array(0.1, dtype=np.float32)

tiling = (total_length, tile_num, scalar)

tiling_data = b''.join(x.tobytes() for x in tiling)

with os.fdopen(os.open('./input/tiling.bin', WRITE_FILE_FLAGS, PEN_FILE_MODES_640), 'wb') as f:

f.write(tiling_data)

//生成輸入數據

input_x = np.random.uniform(-100, 100, [8, 200, 1024]).astype(np.float16)

//生成golden數據，功能和LeakyRelu相同

golden = np.where(input_x > 0, input_x, input_x * scalar).astype(np.float16)

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

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

編譯工程文件：CMakeLists.txt

用于編譯cpu側或npu側運行的Ascend C算子。主要關注CMakeLists.txt中源文件是否全部列全。

調用算子的應用程序：main.cpp

主要是內存申請，數據拷貝和文件讀寫等操作，并最終調用算子，相關API的介紹如下：

1. AscendCL初始化接口aclInit，用于運行時接口AscendCL的初始化，是程序最先調用的接口；aclrtCreateContext和aclrtCreateStream用于創建Context和Stream，主要用于線程相關的資源管理。
2. aclrtMallocHost接口，用于在Host上申請內存：aclError aclrtMallocHost(void **hostPtr, size_t size)
   這個函數和C語言中的malloc類似，用于在Host上申請一定字節大小的內存，其中hostPtr是指向所分配內存的指針，size是申請的內存大小，如果需要釋放這塊內存的話，使用aclrtFreeHost接口釋放，這和C語言中的free函數對應。
3. aclrtMalloc接口，用于在Device上申請內存：aclError aclrtMalloc(void **devPtr, size_t size, aclrtMemMallocPolicy policy)
   和Host上的內存申請接口相比，多了一個policy參數，用于設置內存分配規則，一般設置成ACL_MEM_MALLOC_HUGE_FIRST就可以了。使用完畢后可以用對應的aclrtFree接口釋放內存。
4. aclrtMemcpy接口，用于Host和Device之間數據拷貝：前面申請的內存區分了Host內存和Device內存，那就會涉及到數據同步的問題，aclrtMemcpy就是用于Host和Device之間數據通信的接口：
   aclError aclrtMemcpy(void *dst, size_t destMax, const void *src, size_t count, aclrtMemcpyKind kind)
   其中src指向數據源，而dst是目標內存地址，destMax 是目的內存地址的最大內存長度，count是拷貝的字節數，其中aclrtMemcpyKind控制復制的方向：ACL_MEMCPY_HOST_TO_HOST、ACL_MEMCPY_HOST_TO_DEVICE、ACL_MEMCPY_DEVICE_TO_HOST和ACL_MEMCPY_DEVICE_TO_DEVICE，像ACL_MEMCPY_HOST_TO_DEVICE就是將Host上數據拷貝到Device上。
5. 核心函數為CPU側的調用kernel函數

ICPU_RUN_KF(leakyrelu_custom, blockDim, x, y, usrWorkSpace, tiling);

和NPU側調用的

leakyrelu_custom_do(blockDim, nullptr, stream, xDevice, yDevice, workspaceDevice, tilingDevice);

完整代碼如下：

//This file constains code of cpu debug and npu code.We read data from bin file and write result to file.

#include "data_utils.h"

#include "leakyrelu_custom_tiling.h"

#ifndef __CCE_KT_TEST__

#include "acl/acl.h"

extern void leakyrelu_custom_do(uint32_t coreDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y,

uint8_t* workspace, uint8_t* tiling);

#else

#include "tikicpulib.h"

extern "C" __global__ __aicore__ void leakyrelu_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling);

#endif

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

{

size_t tilingSize = sizeof(LeakyReluCustomTilingData);

size_t usrWorkspaceSize = 4096;

size_t sysWorkspaceSize = 16 * 1024 * 1024;

uint32_t blockDim = 8;

#ifdef __CCE_KT_TEST__ //CPU側調用

//申請內存用于存放workspace和tilling數據

uint8_t* usrWorkSpace = (uint8_t*)AscendC::GmAlloc(usrWorkspaceSize);

uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tilingSize);

ReadFile("./input/tiling.bin", tilingSize, tiling, tilingSize);

size_t inputByteSize = blockDim * 200 * 1024 * sizeof(uint16_t); // uint16_t represent half

size_t outputByteSize = blockDim * 200 * 1024 * sizeof(uint16_t); // uint16_t represent half

//申請內存用于存放輸入和輸出數據

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

uint8_t* y = (uint8_t*)AscendC::GmAlloc(inputByteSize);

//獲取輸入數據

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

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

//在AIV上執行

AscendC::SetKernelMode(KernelMode::AIV_MODE);

//調用kernel函數

ICPU_RUN_KF(leakyrelu_custom, blockDim, x, y, usrWorkSpace, tiling); // use this macro for cpu debug

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

WriteFile("./output/output_y.bin", y, outputByteSize);

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

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

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

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

#else //NPU側調用

CHECK_ACL(aclInit(nullptr));

aclrtContext context;

int32_t deviceId = 0;

CHECK_ACL(aclrtSetDevice(deviceId));

CHECK_ACL(aclrtCreateContext(&context, deviceId));

aclrtStream stream = nullptr;

CHECK_ACL(aclrtCreateStream(&stream));

uint8_t *xHost, *yHost, *tilingHost, *workspaceHost;

uint8_t *xDevice, *yDevice, *tilingDevice, *workspaceDevice;

//申請host上tilling內存并讀入tilling數據

CHECK_ACL(aclrtMallocHost((void**)(&tilingHost), tilingSize));

ReadFile("./input/tiling.bin", tilingSize, tilingHost, tilingSize);

//申請host上workspace內存

CHECK_ACL(aclrtMallocHost((void**)(&workspaceHost), tilingSize));

size_t inputByteSize = blockDim * 200 * 1024 * sizeof(uint16_t); // uint16_t represent half

size_t outputByteSize = blockDim * 200 * 1024 * sizeof(uint16_t); // uint16_t represent half

size_t workspaceByteSize = sysWorkspaceSize + usrWorkspaceSize;

//申請host和device上的輸入輸出內存和device上的workspace和tilling內存

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

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

CHECK_ACL(aclrtMallocHost((void**)(&workspaceHost), workspaceByteSize));

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**)&tilingDevice, tilingSize, ACL_MEM_MALLOC_HUGE_FIRST));

CHECK_ACL(aclrtMalloc((void**)&workspaceDevice, workspaceByteSize, ACL_MEM_MALLOC_HUGE_FIRST));

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

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

//從host上拷貝輸入數據和tilling數據到device

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

CHECK_ACL(aclrtMemcpy(tilingDevice, tilingSize, tilingHost, tilingSize, ACL_MEMCPY_HOST_TO_DEVICE));

//調用核函數

leakyrelu_custom_do(blockDim, nullptr, stream, xDevice, yDevice, workspaceDevice, tilingDevice);

//等待核函數運行完成

CHECK_ACL(aclrtSynchronizeStream(stream));

//拷回運行結果到host

CHECK_ACL(aclrtMemcpy(yHost, outputByteSize, yDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST));

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

WriteFile("./output/output_y.bin", yHost, outputByteSize);

//釋放資源

CHECK_ACL(aclrtFree(xDevice));

CHECK_ACL(aclrtFree(yDevice));

CHECK_ACL(aclrtFree(workspaceDevice));

CHECK_ACL(aclrtFree(tilingDevice));

CHECK_ACL(aclrtFreeHost(xHost));

CHECK_ACL(aclrtFreeHost(yHost));

CHECK_ACL(aclrtFreeHost(workspaceHost));

CHECK_ACL(aclrtFreeHost(tilingHost));

CHECK_ACL(aclrtDestroyStream(stream));

CHECK_ACL(aclrtDestroyContext(context));

CHECK_ACL(aclrtResetDevice(deviceId));

CHECK_ACL(aclFinalize());

#endif

return 0;

}

一鍵式編譯運行腳本run.sh

編譯和運行應用程序。

cpu側運行命令：

bash run.sh leakyrelu_custom ascend910B1 VectorCore cpu

npu側運行命令：

bash run.sh leakyrelu_custom ascend910B1 VectorCore npu

參數含義如下：

bash run.sh <kernel_name> <soc_version> <core_type> <run_mode>

<kernel_name>表示需要運行的算子。
<soc_version>表示算子運行的AI處理器型號。
<core_type>表示在AI Core上或者Vector Core上運行，參數取值為AiCore/VectorCore。
<run_mode>表示算子以cpu模式或npu模式運行，參數取值為cpu/npu。

kernel實現

函數原型定義

本樣例中，函數名為leakyrelu_custom，根據對算子輸入輸出的分析，確定有2個參數x，y，其中x為輸入內存，y為輸出內存。核函數原型定義如下所示：

extern "C" __global__ __aicore__ void leakyrelu_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling){ }

使用__global__函數類型限定符來標識它是一個核函數，可以被<<<...>>>調用；使用__aicore__函數類型限定符來標識該核函數在設備端AI Core上執行；為方便起見，統一使用GM_ADDR宏修飾入參，GM_ADDR宏定義：

#define GM_ADDR __gm__ uint8_t* __restrict__

獲取tilling數據，并調用算子類的Init和Process函數。

算子類的Init函數，完成內存初始化相關工作，Process函數完成算子實現的核心邏輯。

extern "C" __global__ __aicore__ void leakyrelu_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)

{

GET_TILING_DATA(tilingData, tiling);

KernelLeakyRelu op;

op.Init(x, y, tilingData.totalLength, tilingData.tileNum, tilingData.scalar);

op.Process();

}

對核函數的調用進行封裝

封裝后得到leakyrelu_custom_do函數，便于主程序調用。#ifndef __CCE_KT_TEST__表示該封裝函數僅在編譯運行NPU側的算子時會用到，編譯運行CPU側的算子時，可以直接調用add_custom函數。調用核函數時，除了需要傳入輸入輸出參數x，y，切分相關參數tiling，還需要傳入blockDim（核函數執行的核數）, l2ctrl（保留參數，設置為nullptr）, stream（應用程序中維護異步操作執行順序的stream）來規定核函數的執行配置。

#ifndef __CCE_KT_TEST__

// call of kernel function

void leakyrelu_custom_do(uint32_t blockDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y,

uint8_t* workspace, uint8_t* tiling)

{

leakyrelu_custom<<<blockDim, l2ctrl, stream>>>(x, y, workspace, tiling);

}

#endif

獲取tiling參數

主要從tilingPointer中獲取tiling的參數totalLength(總長度)、tileNum（切分個數，單核循環處理數據次數）和scalar（LeakyRelu計算標量）。

#define GET_TILING_DATA(tilingData, tilingPointer) \

LeakyReluCustomTilingData tilingData; \

INIT_TILING_DATA(LeakyReluCustomTilingData, tilingDataPointer, tilingPointer); \

(tilingData).totalLength = tilingDataPointer->totalLength; \

(tilingData).tileNum = tilingDataPointer->tileNum; \

(tilingData).scalar = tilingDataPointer->scalar;

#endif // LEAKYRELU_CUSTOM_TILING_H

Init函數

主要獲取tiling數據后，設置單核上gm的地址和Buffer的初始化。

__aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t totalLength, uint32_t tileNum, float scalar)

{

ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");

this->blockLength = totalLength / GetBlockNum();

this->tileNum = tileNum;

this->scalar = static_cast<half>(scalar);

ASSERT(tileNum != 0 && "tile num can not be zero!");

this->tileLength = this->blockLength / tileNum / BUFFER_NUM;

// get start index for current core, core parallel

xGm.SetGlobalBuffer((__gm__ half*)x + this->blockLength * get_block_idx(), this->blockLength);

yGm.SetGlobalBuffer((__gm__ half*)y + this->blockLength * get_block_idx(), this->blockLength);

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

pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(half));

pipe.InitBuffer(outQueueY, BUFFER_NUM, this->tileLength * sizeof(half));

}

Process函數

主要實現三個CopyIn、Compute、CopyOut這三stage。

__aicore__ inline void Process()

{

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

int32_t loopCount = this->tileNum * BUFFER_NUM;

// tiling strategy, pipeline parallel

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

CopyIn(i);

Compute(i);

CopyOut(i);

}

}

CopyIn函數

負責從Global Memory拷貝數據到Local Memory，并將數據加入Queue

__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 * tileLength], tileLength);

// enque input tensors to VECIN queue

inQueueX.EnQue(xLocal);

}

Compute函數

負責從Queue中取出數據，進行計算，并將結果放入Queue

__aicore__ inline void Compute(int32_t progress)

{

// deque input tensors from VECIN queue

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

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

// call LeakyRelu instr for computation

LeakyRelu(yLocal, xLocal, scalar, tileLength);

// enque the output tensor to VECOUT queue

outQueueY.EnQue<half>(yLocal);

// free input tensors for reuse

inQueueX.FreeTensor(xLocal);

}

CopyOut函數

負責從Queue中將數據取出，并將數據從Local Memory拷貝到Global Memory。

__aicore__ inline void CopyOut(int32_t progress)

{

// deque output tensor from VECOUT queue

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

// copy progress_th tile from local tensor to global tensor

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

// free output tensor for reuse

outQueueY.FreeTensor(yLocal);

}

編譯和執行

在CPU側執行

執行結果如下：

可以看到最后的輸出結果output_y.bin和標桿數據golden.bin的MD5值相同，說明計算結果相同。

執行完成后，在input下存放輸入數據和tiling數據，在output下面存放了輸出數據和標桿數據，npuchk目錄下是每個核的npu_check執行結果

在當前目錄還有一個可執行二進制文件leakyrelu_custom_cpu，如果執行報錯，可以通過gdb調試這個可執行文件，具體調試可參考文末官方教程。

在NPU側執行

在NPU側執行有兩種方式：仿真執行和上板運行，命令都相同，只是編譯選項不同，我們可以通過修改編譯選項-DASCEND_RUN_MODE為SIMULATOR運行CAModel仿真，設置為 ONBOARD是上板運行。

function compile_and_execute() {

# 使用cmake編譯cpu側或者npu側算子, SIMULATOR or ONBOARD

mkdir -p build; cd build; \

cmake .. \

-Dsmoke_testcase=$1 \

-DASCEND_PRODUCT_TYPE=$2 \

-DASCEND_CORE_TYPE=$3 \

-DASCEND_RUN_MODE="SIMULATOR" \

-DASCEND_INSTALL_PATH=$ASCEND_HOME_DIR

VERBOSE=1 cmake --build . --target ${1}_${4}

……

}

參考資料

總之，學習Ascend C，僅需了解C++編程、理解對列通信與內存申請釋放機制、通過調用相應的計算接口與搬運接口，就可以寫出運行在昇騰AI處理器上的高性能算子。

了解更多Ascend C學習資源，請訪問官方教程：Ascend C編程指南（官方教程）