2023年5月6日，在昇騰AI開發(fā)者峰會上，華為正式發(fā)布了面向算子開發(fā)場景的昇騰Ascend C編程語言。Ascend C原生支持C/C++編程規(guī)范，通過多層接口抽象、并行編程范式、孿生調試等技術，極大提高了算子的開發(fā)效率，幫助AI開發(fā)者低成本完成算子開發(fā)和模型調優(yōu)部署。
1 昇騰AI軟硬件基礎
和CUDA開發(fā)的算子運行在GPU上一樣，基于Ascend C開發(fā)的算子，可以通過異構計算架構CANN（Compute Architecture for Neural Networks）運行在昇騰AI處理器（可簡稱NPU）上。CANN是使能昇騰AI處理器的一個軟件棧，通過軟硬件協(xié)同優(yōu)化，能夠充分發(fā)揮昇騰AI處理器的強大算力。從下面的架構圖可以清楚的看到，使用Ascend C編程語言開發(fā)的算子通過編譯器編譯和運行時調度，最終運行在昇騰AI處理器上。
我們知道，通用計算就是我們常寫的一些在CPU上運行的計算，它擅長邏輯控制和串行計算，而AI計算相對通用計算來說，更擅長并行計算，可支持大規(guī)模的計算密集型任務。如下面左圖所示，做一個矩陣乘，使用CPU計算需要三層for循環(huán)，而右圖在昇騰AI處理器上使用vector計算單元，只需要兩層for循環(huán)，最小計算代碼能同時計算多個數(shù)據(jù)的乘加，更近一步，如果使用Cube計算單元，只需要一條語句就能完成一個矩陣乘的計算，這就是我們所說的SIMD（單指令多數(shù)據(jù)）。因此，我們通常使用AI處理器來進行大量的并行計算。
NPU不能獨立運行，需要與CPU協(xié)同工作，可以看成是CPU的協(xié)處理器，CPU負責整個操作系統(tǒng)運行，管理各類資源并進行復雜的邏輯控制，而NPU主要負責并行計算任務。在基于CPU+NPU的異構計算架構中，NPU與CPU通過PCIe總線連接在一起來協(xié)同工作，CPU所在位置稱為主機端（host），而NPU所在位置稱為設備端（device），示意圖如下：
這里再詳細介紹一下昇騰AI處理器。昇騰AI處理器有不同的型號和產(chǎn)品形態(tài)，小到模塊、加速卡，大到服務器、集群。昇騰AI處理器里面最核心的部件是AI Core，有多個，是神經(jīng)網(wǎng)絡加速的計算核心，每一個AI Core就相當于我們大家平時理解的多核cpu里的每個核，使用Ascend C編程語言開發(fā)的算子就運行在AI Core上，因為核心的神經(jīng)網(wǎng)絡計算的加速都來源于AI Core的算力。
AI Core內部的并行計算架構抽象如下圖所示：
這個并行計算架構抽象核心包含了幾個大的部件，AI Core外面有一個Gobal Memory，是多個AI Core共享的，在AI Core內部有一塊本地內存Local Memory，因為靠近計算單元，所以它的帶寬會非常高，相對的容量就會很小，比如一般是幾百K到1M。AI Core內部的核心組件有三個計算單元，標量計算單元、向量計算單元，矩陣計算單元。另外還有一個DMA搬運單元，DMA搬運單元負責在Global Memory和Local Memory之間搬運數(shù)據(jù)。
AI Core內部的異步并行計算過程：Scalar計算單元讀取指令序列，并把向量計算、矩陣計算、數(shù)據(jù)搬運指令發(fā)射給對應單元的指令隊列，向量計算單元、矩陣計算單元、數(shù)據(jù)搬運單元異步并行執(zhí)行接收到的指令。該過程可以參考上圖中藍色箭頭所示的指令流。不同的指令間有可能存在依賴關系，為了保證不同指令隊列間的指令按照正確的邏輯關系執(zhí)行，Scalar計算單元也會給對應單元下發(fā)同步指令。各單元之間的同步過程可以參考上圖中的橙色箭頭所示的同步信號流。
AI Core內部數(shù)據(jù)處理的基本過程：DMA搬入單元把數(shù)據(jù)搬運到Local Memory，Vector/Cube計算單元完成數(shù)據(jù)，并把計算結果寫回Local Memory，DMA搬出單元把處理好的數(shù)據(jù)搬運回Global Memory。該過程可以參考上圖中的紅色箭頭所示的數(shù)據(jù)流。
2 Ascend C編程模型基礎
2.1 Ascend C編程范式
Ascend C編程范式是一種流水線式的編程范式，把算子核內的處理程序，分成多個流水任務，通過隊列（Queue）完成任務間通信和同步，并通過統(tǒng)一的內存管理模塊（Pipe）管理任務間通信內存。流水編程范式應用了流水線并行計算方法。
若n=3，即待處理的數(shù)據(jù)被切分成3片，則上圖中的流水任務運行起來的示意圖如下，從運行圖中可以看出，對于同一片數(shù)據(jù)，Stage1、Stage2、Stage3之間的處理具有依賴關系，需要串行處理；不同的數(shù)據(jù)切片，同一時間點，可以有多個任務在并行處理，由此達到任務并行、提升性能的目的。
Ascend C分別針對Vector、Cube編程設計了不同的流水任務。開發(fā)者只需要完成基本任務的代碼實現(xiàn)即可，底層的指令同步和并行調度由Ascend C框架實現(xiàn)，開發(fā)者無需關注。
2.2 矢量編程范式
矢量編程范式把算子的實現(xiàn)流程分為3個基本任務：CopyIn，Compute，CopyOut。CopyIn負責搬入操作，Compute負責矢量計算操作，CopyOut負責搬出操作。
我們只需要根據(jù)編程范式完成基本任務的代碼實現(xiàn)就可以了，底層的指令同步和并行調度由Ascend C框架來實現(xiàn)。
那Ascend C是怎么完成不同任務之間的數(shù)據(jù)通信和同步的呢？這里Ascend C提供了Queue隊列管理的API，主要就是兩個隊列操作API EnQue、DeQue以及內存的邏輯抽象。
矢量編程中使用到的邏輯位置（QuePosition）定義如下：
· 搬入數(shù)據(jù)的存放位置：VECIN；
· 計算中間變量的位置：VECCALC；
· 搬出數(shù)據(jù)的存放位置：VECOUT。
從前面可以看到，矢量編程主要分為CopyIn、Compute、CopyOut三個任務。CopyIn任務中將輸入數(shù)據(jù)從Global內存搬運至Local內存后，需要使用EnQue將LocalTensor放入VECIN的Queue中；Compute任務等待VECIN的Queue中LocalTensor出隊之后才可以完成矢量計算，計算完成后使用EnQue將計算結果LocalTensor放入到VECOUT的Queue中；CopyOut任務等待VECOUT的Queue中LocalTensor出隊，再將其拷貝到Global內存。這樣 ，Queue隊列就完成了三個任務間的數(shù)據(jù)通信和同步。具體流程和流程圖如下：
1. Stage1：CopyIn任務。
使用DataCopy接口將GlobalTensor數(shù)據(jù)拷貝到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實現(xiàn)代碼就很清晰了。先初始化內存和隊列，然后通過編程范式實現(xiàn)CopyIn、Compute、CopyOut三個Stage就可以了。
2.3 SPMD并行編程-多核
最前面介紹昇騰AI處理器的時候，有介紹過AI Core是有多個的，那我們怎么把多個AI Core充分利用起來呢？常用的并行計算方法中，有一種SPMD（Single-Program Multiple-Data）數(shù)據(jù)并行的方法，簡單說就是將數(shù)據(jù)分片，每片數(shù)據(jù)經(jīng)過完整的一個數(shù)據(jù)處理流程。這個就能和昇騰AI處理器的多核匹配上了，我們將數(shù)據(jù)分成多份，每份數(shù)據(jù)的處理運行在一個核上，這樣每份數(shù)據(jù)并行處理完成，整個數(shù)據(jù)也就處理完了。Ascend C是SPMD（Single-Program Multiple-Data）編程，多個AI Core共享相同的指令代碼，每個核上的運行實例唯一的區(qū)別是就是block_idx（內置變量）不同，這樣我們就可以通過block_idx來區(qū)分不同的核，只要對Global Memory上的數(shù)據(jù)地址進行切分偏移，就可以讓每個核處理自己對應的那部分數(shù)據(jù)了。
算子被調用時，所有的計算核心都執(zhí)行相同的實現(xiàn)代碼，入口函數(shù)的入?yún)⒁彩窍嗤?。每個核上處理的數(shù)據(jù)地址需要在起始地址上增加block_idx*BLOCK_LENGTH（每個block處理的數(shù)據(jù)長度）的偏移來獲取。這樣也就實現(xiàn)了多核并行計算的數(shù)據(jù)切分。
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);
……
}
……
}
2.4 Ascend C API介紹
在整個kernel實現(xiàn)中，最最核心的代碼就是Add(zLocal, xLocal, yLocal, TILE_LENGTH);通過一個Ascend C提供的API接口完成了所有數(shù)據(jù)的加法計算，對，沒看錯，就是這個接口完成了計算。
接下來就介紹下Ascend C提供的API。Ascend C算子采用標準C++語法和一組類庫API進行編程，類庫API主要包含以下幾種，大家可以在核函數(shù)的實現(xiàn)中根據(jù)自己的需求選擇合適的API：
· 計算類API，包括標量計算API、向量計算API、矩陣計算API，分別實現(xiàn)調用Scalar計算單元、Vector計算單元、Cube計算單元執(zhí)行計算的功能。
· 數(shù)據(jù)搬運API，上述計算API基于Local Memory數(shù)據(jù)進行計算，所以數(shù)據(jù)需要先從Global Memory搬運至Local Memory，再使用計算接口完成計算，最后從Local Memory搬出至Global Memory。執(zhí)行搬運過程的接口稱之為數(shù)據(jù)搬移接口，比如DataCopy接口。
· 內存管理API，用于分配管理內存，比如AllocTensor、FreeTensor接口。
· 任務同步API，完成任務間的通信和同步，比如EnQue、DeQue接口。
Ascend C API的計算操作數(shù)都是Tensor類型：GlobalTensor和LocalTensor。
介紹完Ascend C API種類后，下面來解釋下為什么一個Add接口就可以計算所有的數(shù)。原來Ascend C編程模型是基于SIMD（單指令多數(shù)據(jù)）架構的，單條指令可以完成多個數(shù)據(jù)操作，同時在API內部封裝了一些指令的高級功能。
2.5 算子執(zhí)行基本流程
前面有提到，在異構計算架構中，NPU與CPU是協(xié)同工作的，在Ascend C編程模型中，我們需要實現(xiàn)NPU側的代碼和CPU側的代碼。在NPU側的代碼我們通常叫做Kernel實現(xiàn)代碼，CPU側的代碼我們一般叫做Host實現(xiàn)代碼，一份完整的Ascend C代碼，通常包括Host側實現(xiàn)代碼和Kernel側實現(xiàn)代碼。Ascend C算子執(zhí)行的基本流程如下：
1、 初始化Device設備；
2、 創(chuàng)建Context綁定設備；
3、 分配Host內存，并進行數(shù)據(jù)初始化；
4、 分配Device內存，并將數(shù)據(jù)從Host上拷貝到Device上；
5、 用內核調用符<<<>>>調用核函數(shù)完成指定的運算；
6、 將Device上的運算結果拷貝回Host；
7、 釋放申請的資源。
2.6 核函數(shù)介紹
上面的流程中，最重要的一步就是調用核函數(shù)來進行并行計算任務。核函數(shù)（Kernel Function）是Ascend C算子Device側實現(xiàn)的入口。在核函數(shù)中，需要為在AI核上執(zhí)行的代碼規(guī)定要進行的數(shù)據(jù)訪問和計算操作。
extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z);
上面這個是一個核函數(shù)聲明的示例，extern "C"表示核函數(shù)按照類C的編譯和連接規(guī)約來編譯和連接，__global__函數(shù)類型限定符表示它是一個核函數(shù)， __aicore__函數(shù)類型限定符表示該核函數(shù)在device側的AI Core上執(zhí)行。參數(shù)列表中的變量類型限定符__gm__，表明該指針變量指向Global Memory上某處內存地址，注意這里的入?yún)⒅荒苤С种羔樆駽/C++內置數(shù)據(jù)類型，樣例里指針使用的類型為uint8_t，在后續(xù)的使用中需要將其轉化為實際的指針類型。
Ascend C編程模型中的核函數(shù)采用內核調用符<<<...>>>來調用，樣例如下：
kernel_name<<
kernel_name即為上面講的核函數(shù)名稱，argument list是核函數(shù)的函數(shù)入?yún)?，?lt;<<>>>中間，有3個參數(shù)：
· blockDim，規(guī)定了核函數(shù)將會在幾個核上執(zhí)行，我們可以先設置為1；
· l2ctrl，保留參數(shù)，暫時設置為固定值nullptr，我們不用關注；
· stream，使用aclrtCreateStream創(chuàng)建，用于多線程調度。
3 樣例開發(fā)講解
3.1 樣例代碼結構
|-- CMakeLists.txt //編譯工程文件
|-- cmake //編譯工程文件
|-- data_utils.h //數(shù)據(jù)讀入寫出函數(shù)
|-- input //存放腳本生成的輸入數(shù)據(jù)目錄
|-- leakyrelu_custom.cpp //算子kernel實現(xiàn)
|-- leakyrelu_custom.py //輸入數(shù)據(jù)和真值數(shù)據(jù)生成腳本文件
|-- leakyrelu_custom_tiling.h //host側tiling函數(shù)
|-- main.cpp //主函數(shù)，host側調用代碼，含cpu域及npu域調用
|-- output //存放算子運行輸出數(shù)據(jù)和標桿數(shù)據(jù)的目錄
|-- readme.md //執(zhí)行命令說明
|-- run.sh //運行腳本
3.2 主要文件
3.2.1 輸入數(shù)據(jù)和真值數(shù)據(jù)生成腳本文件：KERNEL_NAME.py。
根據(jù)算子的輸入輸出編寫生成輸入數(shù)據(jù)和真值數(shù)據(jù)的腳本。
本例子生成8 * 200 * 1024大小的fp16數(shù)據(jù)：
……
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)
//生成輸入數(shù)據(jù)
input_x = np.random.uniform(-100, 100, [8, 200, 1024]).astype(np.float16)
//生成golden數(shù)據(jù)，功能和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")
3.2.2 編譯工程文件：CMakeLists.txt
用于編譯cpu側或npu側運行的Ascend C算子。主要關注CMakeLists.txt中源文件是否全部列全。
3.2.3 調用算子的應用程序：main.cpp
主要是內存申請，數(shù)據(jù)拷貝和文件讀寫等操作，并最終調用算子，相關API的介紹如下：
1、 AscendCL初始化接口aclInit，用于運行時接口AscendCL的初始化，是程序最先調用的接口；aclrtCreateContext和aclrtCreateStream用于創(chuàng)建Context和Stream，主要用于線程相關的資源管理。
2、 aclrtMallocHost接口，用于在Host上申請內存：
aclError aclrtMallocHost(void **hostPtr, size_t size)
這個函數(shù)和C語言中的malloc類似，用于在Host上申請一定字節(jié)大小的內存，其中hostPtr是指向所分配內存的指針，size是申請的內存大小，如果需要釋放這塊內存的話，使用aclrtFreeHost接口釋放，這和C語言中的free函數(shù)對應。
3、 aclrtMalloc接口，用于在Device上申請內存：
aclError aclrtMalloc(void **devPtr, size_t size, aclrtMemMallocPolicy policy)
和Host上的內存申請接口相比，多了一個policy參數(shù)，用于設置內存分配規(guī)則，一般設置成ACL_MEM_MALLOC_HUGE_FIRST就可以了。使用完畢后可以用對應的aclrtFree接口釋放內存。
4、 aclrtMemcpy接口，用于Host和Device之間數(shù)據(jù)拷貝：
前面申請的內存區(qū)分了Host內存和Device內存，那就會涉及到數(shù)據(jù)同步的問題，aclrtMemcpy就是用于Host和Device之間數(shù)據(jù)通信的接口：
aclError aclrtMemcpy(void *dst, size_t destMax, const void *src, size_t count, aclrtMemcpyKind kind)
其中src指向數(shù)據(jù)源，而dst是目標內存地址，destMax 是目的內存地址的最大內存長度，count是拷貝的字節(jié)數(shù)，其中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上數(shù)據(jù)拷貝到Device上。
5、 核心函數(shù)為CPU側的調用kernel函數(shù)
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數(shù)據(jù)
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
//申請內存用于存放輸入和輸出數(shù)據(jù)
uint8_t* x = (uint8_t*)AscendC::GmAlloc(inputByteSize);
uint8_t* y = (uint8_t*)AscendC::GmAlloc(inputByteSize);
//獲取輸入數(shù)據(jù)
ReadFile("./input/input_x.bin", inputByteSize, x, inputByteSize);
// PrintData(x, 16, printDataType::HALF);
//在AIV上執(zhí)行
AscendC::SetKernelMode(KernelMode::AIV_MODE);
//調用kernel函數(shù)
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數(shù)據(jù)
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上拷貝輸入數(shù)據(jù)和tilling數(shù)據(jù)到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));
//調用核函數(shù)
leakyrelu_custom_do(blockDim, nullptr, stream, xDevice, yDevice, workspaceDevice, tilingDevice);
//等待核函數(shù)運行完成
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;
}
3.2.4 一鍵式編譯運行腳本run.sh
編譯和運行應用程序。
cpu側運行命令：
bash run.sh leakyrelu_custom ascend910B1 VectorCore cpu
npu側運行命令：
bash run.sh leakyrelu_custom ascend910B1 VectorCore npu
參數(shù)含義如下：
bash run.sh
3.3 kernel實現(xiàn)
3.3.1 函數(shù)原型定義
本樣例中，函數(shù)名為leakyrelu_custom，根據(jù)對算子輸入輸出的分析，確定有2個參數(shù)x，y，其中x為輸入內存，y為輸出內存。核函數(shù)原型定義如下所示：
extern "C" __global__ __aicore__ void leakyrelu_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling){ }
使用__global__函數(shù)類型限定符來標識它是一個核函數(shù)，可以被<<<...>>>調用；使用__aicore__函數(shù)類型限定符來標識該核函數(shù)在設備端AI Core上執(zhí)行；為方便起見，統(tǒng)一使用GM_ADDR宏修飾入?yún)?，GM_ADDR宏定義：
#define GM_ADDR __gm__ uint8_t* __restrict__
3.3.2 獲取tilling數(shù)據(jù)，并調用算子類的Init和Process函數(shù)。
算子類的Init函數(shù)，完成內存初始化相關工作，Process函數(shù)完成算子實現(xiàn)的核心邏輯。
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();
}
3.3.3 對核函數(shù)的調用進行封裝
封裝后得到leakyrelu_custom_do函數(shù)，便于主程序調用。#ifndef __CCE_KT_TEST__表示該封裝函數(shù)僅在編譯運行NPU側的算子時會用到，編譯運行CPU側的算子時，可以直接調用add_custom函數(shù)。調用核函數(shù)時，除了需要傳入輸入輸出參數(shù)x，y，切分相關參數(shù)tiling，還需要傳入blockDim（核函數(shù)執(zhí)行的核數(shù)）, l2ctrl（保留參數(shù)，設置為nullptr）, stream（應用程序中維護異步操作執(zhí)行順序的stream）來規(guī)定核函數(shù)的執(zhí)行配置。
#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<<
}
#endif
3.3.4 獲取tiling參數(shù)
主要從tilingPointer中獲取tiling的參數(shù)totalLength(總長度)、tileNum（切分個數(shù)，單核循環(huán)處理數(shù)據(jù)次數(shù)）和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
3.3.5 Init函數(shù)
主要獲取tiling數(shù)據(jù)后，設置單核上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
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));
}
3.3.6 Process函數(shù)
主要實現(xiàn)三個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);
}
}
3.3.7 CopyIn函數(shù)
負責從Global Memory拷貝數(shù)據(jù)到Local Memory，并將數(shù)據(jù)加入Queue
__aicore__ inline void CopyIn(int32_t progress)
{
// alloc tensor from queue memory
LocalTensor
// 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);
}
3.3.8 Compute函數(shù)
負責從Queue中取出數(shù)據(jù)，進行計算，并將結果放入Queue
__aicore__ inline void Compute(int32_t progress)
{
// deque input tensors from VECIN queue
LocalTensor
LocalTensor
// call LeakyRelu instr for computation
LeakyRelu(yLocal, xLocal, scalar, tileLength);
// enque the output tensor to VECOUT queue
outQueueY.EnQue
// free input tensors for reuse
inQueueX.FreeTensor(xLocal);
}
3.3.9 CopyOut函數(shù)
負責從Queue中將數(shù)據(jù)取出，并將數(shù)據(jù)從Local Memory拷貝到Global Memory。
__aicore__ inline void CopyOut(int32_t progress)
{
// deque output tensor from VECOUT queue
LocalTensor
// copy progress_th tile from local tensor to global tensor
DataCopy(yGm[progress * tileLength], yLocal, tileLength);
// free output tensor for reuse
outQueueY.FreeTensor(yLocal);
}
3.4 編譯和執(zhí)行
3.4.1 在CPU側執(zhí)行
執(zhí)行結果如下：
可以看到最后的輸出結果output_y.bin和標桿數(shù)據(jù)golden.bin的MD5值相同，說明計算結果相同。
執(zhí)行完成后，在input下存放輸入數(shù)據(jù)和tiling數(shù)據(jù)，在output下面存放了輸出數(shù)據(jù)和標桿數(shù)據(jù)，npuchk目錄下是每個核的npu_check執(zhí)行結果
在當前目錄還有一個可執(zhí)行二進制文件leakyrelu_custom_cpu，如果執(zhí)行報錯，可以通過gdb調試這個可執(zhí)行文件，具體調試可參考文末官方教程。
3.4.2 在NPU側執(zhí)行
在NPU側執(zhí)行有兩種方式：仿真執(zhí)行和上板運行，命令都相同，只是編譯選項不同，我們可以通過修改編譯選項-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}
……
}
4 參考資料
總之，學習Ascend C，僅需了解C++編程、理解對列通信與內存申請釋放機制、通過調用相應的計算接口與搬運接口，就可以寫出運行在昇騰AI處理器上的高性能算子。
了解更多Ascend C學習資源，請訪問官方教程：Ascend C編程指南（官方教程）
（免責聲明：本網(wǎng)站內容主要來自原創(chuàng)、合作伙伴供稿和第三方自媒體作者投稿，凡在本網(wǎng)站出現(xiàn)的信息，均僅供參考。本網(wǎng)站將盡力確保所提供信息的準確性及可靠性，但不保證有關資料的準確性及可靠性，讀者在使用前請進一步核實，并對任何自主決定的行為負責。本網(wǎng)站對有關資料所引致的錯誤、不確或遺漏，概不負任何法律責任。
任何單位或個人認為本網(wǎng)站中的網(wǎng)頁或鏈接內容可能涉嫌侵犯其知識產(chǎn)權或存在不實內容時，應及時向本網(wǎng)站提出書面權利通知或不實情況說明，并提供身份證明、權屬證明及詳細侵權或不實情況證明。本網(wǎng)站在收到上述法律文件后，將會依法盡快聯(lián)系相關文章源頭核實，溝通刪除相關內容或斷開相關鏈接。 ）
