本文分享自華為云社區(qū)《Ascend C保姆級(jí)教程：我的第一份Ascend C代碼》，作者：昇騰CANN 。

Ascend C是昇騰AI異構(gòu)計(jì)算架構(gòu)CANN針對(duì)算子開發(fā)場(chǎng)景推出的編程語言，原生支持C和C++標(biāo)準(zhǔn)規(guī)范，最大化匹配用戶開發(fā)習(xí)慣；通過多層接口抽象、自動(dòng)并行計(jì)算、孿生調(diào)試等關(guān)鍵技術(shù)，極大提高算子開發(fā)效率，助力AI開發(fā)者低成本完成算子開發(fā)和模型調(diào)優(yōu)部署。

本文提供Ascend C保姆級(jí)教程，從一個(gè)簡(jiǎn)單的實(shí)例出發(fā)，帶你體驗(yàn)Ascend C算子開發(fā)的基本流程。

完成實(shí)例開發(fā)之前，需要先了解一些必備的背景知識(shí)。

1. 背景知識(shí)

• 多核并行

使用Ascend C開發(fā)的算子運(yùn)行在AI Core上，AI Core是昇騰NPU硬件平臺(tái)的計(jì)算核心，NPU內(nèi)部有多個(gè)AI Core。Ascend C編程過程中會(huì)將需要處理的數(shù)據(jù)拆分同時(shí)在多個(gè)AI Core上運(yùn)行，從而獲取更高的性能。多個(gè)AI Core共享相同的指令代碼，每個(gè)核上的運(yùn)行實(shí)例唯一的區(qū)別是block_idx不同，開發(fā)者只需要關(guān)注單核上的處理程序,也就是核函數(shù)。

• 流水并行

上文提到，開發(fā)者只需要關(guān)注單核處理程序（核函數(shù)），那么如何實(shí)現(xiàn)核函數(shù)的具體邏輯呢？Ascend C提供流水線式的編程范式，基于編程范式可以快速搭建算子實(shí)現(xiàn)的代碼框架，實(shí)現(xiàn)流水并行。

流水線并行的概念和工業(yè)生產(chǎn)中的流水線是類似的，任務(wù)1完成對(duì)某片數(shù)據(jù)的處理后，將其加入到通信隊(duì)列，任務(wù)2空閑時(shí)就會(huì)從隊(duì)列中取出數(shù)據(jù)繼續(xù)處理；可以類比為生產(chǎn)流水線中的工人只完成某一項(xiàng)固定工序，完成后就交由下一項(xiàng)工序負(fù)責(zé)人繼續(xù)處理。

Ascend C編程范式是一種流水線式的編程范式，把算子核內(nèi)的處理程序，分成多個(gè)流水任務(wù)：“搬入、計(jì)算、搬出”，通過隊(duì)列（Queue）完成任務(wù)間通信和同步，并通過統(tǒng)一的內(nèi)存管理模塊（Pipe）管理任務(wù)間通信內(nèi)存。開發(fā)者只需聚焦實(shí)現(xiàn)“搬入、計(jì)算、搬出”內(nèi)容。

• 孿生調(diào)試

基于NPU域算子的調(diào)用接口編寫程序，通過畢昇編譯器編譯后運(yùn)行，可以完成算子NPU域的運(yùn)行驗(yàn)證；基于CPU域算子的調(diào)用接口編寫程序，通過標(biāo)準(zhǔn)的GCC編譯器進(jìn)行編譯后運(yùn)行，并通過GDB通用調(diào)試工具進(jìn)行單步調(diào)試，精準(zhǔn)驗(yàn)證程序執(zhí)行流程是否符合預(yù)期。孿生調(diào)試的能力，大大提升了算子的調(diào)試效率。下文的示例開發(fā)，僅介紹核函數(shù)CPU側(cè)和NPU側(cè)的運(yùn)行驗(yàn)證，具體的調(diào)試步驟將會(huì)在后續(xù)的文章中詳細(xì)介紹。

2. 開發(fā)流程

本文將引導(dǎo)你完成以下任務(wù)，體驗(yàn)Ascend C算子開發(fā)的基本流程。

1. 使用Ascend C完成Add算子核函數(shù)開發(fā)；
2. 使用ICPU_RUN_KF CPU調(diào)測(cè)宏完成算子核函數(shù)CPU側(cè)運(yùn)行驗(yàn)證；
3. 使用<<<>>>內(nèi)核調(diào)用符完成算子核函數(shù)NPU側(cè)運(yùn)行驗(yàn)證。

在正式的開發(fā)之前，還需要先完成環(huán)境準(zhǔn)備和算子分析工作，開發(fā)Ascend C算子的基本流程如下圖所示：

參考本文進(jìn)行開發(fā)之前請(qǐng)先獲取樣例代碼目錄quick-start，該樣例代碼只保留了部分代碼框架，核心代碼在下文的指導(dǎo)步驟中體現(xiàn)。您可以在閱讀本文時(shí)，將指導(dǎo)步驟中的代碼拷貝至對(duì)應(yīng)位置，即可快速完成Ascend C算子的開發(fā)。

3. 環(huán)境準(zhǔn)備

• CANN軟件安裝

開發(fā)算子前，需要先準(zhǔn)備好開發(fā)環(huán)境和運(yùn)行環(huán)境，開發(fā)環(huán)境和運(yùn)行環(huán)境的介紹和具體的安裝步驟可參見昇騰社區(qū)文檔的CANN軟件安裝指南。

• 環(huán)境變量配置

安裝CANN軟件后，使用CANN運(yùn)行用戶編譯、運(yùn)行時(shí)，需要以CANN運(yùn)行用戶登錄環(huán)境，執(zhí)行source ${install_path}/set_env.sh命令設(shè)置環(huán)境變量，其中${install_path}為CANN軟件的安裝目錄。

4. 算子分析

主要分析算子的數(shù)學(xué)表達(dá)式、輸入、輸出以及計(jì)算邏輯的實(shí)現(xiàn)，明確需要調(diào)用的Ascend C接口。

1. 明確算子的數(shù)學(xué)表達(dá)式及計(jì)算邏輯。

Add算子的數(shù)學(xué)表達(dá)式為：

z = x + y

計(jì)算邏輯是：要完成AI Core上的數(shù)據(jù)計(jì)算，輸入數(shù)據(jù)需要先從外部存儲(chǔ)Global Memory搬運(yùn)進(jìn)AI Core的內(nèi)部存儲(chǔ)Local Memory，然后使用計(jì)算接口完成兩個(gè)輸入?yún)?shù)相加，得到最終結(jié)果，再搬出到外部存儲(chǔ)Global Memory上。

2. 明確輸入和輸出。

• Add算子有兩個(gè)輸入：x與y，輸出為z。
• 本樣例中算子的輸入支持的數(shù)據(jù)類型為half（float16），算子輸出的數(shù)據(jù)類型與輸入數(shù)據(jù)類型相同。
• 算子輸入支持shape（8，2048），輸出shape與輸入shape相同。
• 算子輸入支持的format為：ND。

3. 確定核函數(shù)名稱和參數(shù)。

• 您可以自定義核函數(shù)名稱，本樣例中核函數(shù)命名為add_custom。
• 根據(jù)對(duì)算子輸入輸出的分析，確定核函數(shù)有3個(gè)參數(shù)x，y，z；x，y為輸入在Global Memory上的內(nèi)存地址，z為輸出在Global Memory上的內(nèi)存地址。
• 確定算子實(shí)現(xiàn)所需接口。
  • 實(shí)現(xiàn)涉及外部存儲(chǔ)和內(nèi)部存儲(chǔ)間的數(shù)據(jù)搬運(yùn)，查看Ascend C API參考中的數(shù)據(jù)搬移接口，需要使用DataCopy來實(shí)現(xiàn)數(shù)據(jù)搬移。
  • 本樣例只涉及矢量計(jì)算的加法操作，查看Ascend C API參考中的矢量計(jì)算接口，初步分析可使用雙目指令A(yù)dd接口實(shí)現(xiàn)x+y。
  • 計(jì)算中使用到的Tensor數(shù)據(jù)結(jié)構(gòu)（數(shù)據(jù)操作的基礎(chǔ)數(shù)據(jù)結(jié)構(gòu)），使用AllocTensor、FreeTensor進(jìn)行申請(qǐng)和釋放。
  • 并行流水任務(wù)之間使用Queue隊(duì)列完成通信和同步，會(huì)使用到EnQue、DeQue等接口。

通過以上分析，得到Ascend C Add算子的設(shè)計(jì)規(guī)格如下：

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

完成環(huán)境準(zhǔn)備和初步的算子分析后，即可開始Ascend C核函數(shù)的開發(fā)。開發(fā)之前請(qǐng)先獲取樣例代碼目錄quick-start，以下核函數(shù)開發(fā)的樣例代碼在add_custom.cpp中實(shí)現(xiàn)。

本樣例中使用多核并行計(jì)算，即把數(shù)據(jù)進(jìn)行分片，分配到多個(gè)核上進(jìn)行處理。Ascend C核函數(shù)是在一個(gè)核上的處理函數(shù)，所以只處理部分?jǐn)?shù)據(jù)。分配方案是：數(shù)據(jù)整體長(zhǎng)度TOTAL_LENGTH為8* 2048，平均分配到8個(gè)核上運(yùn)行，每個(gè)核上處理的數(shù)據(jù)大小BLOCK_LENGTH為2048。下文的核函數(shù)，只關(guān)注長(zhǎng)度為BLOCK_LENGTH的數(shù)據(jù)應(yīng)該如何處理。

5.1 核函數(shù)的定義

進(jìn)行核函數(shù)的定義，并在核函數(shù)中調(diào)用算子類的Init和Process函數(shù)。請(qǐng)將下文代碼添加至add_custom.cpp的“核函數(shù)實(shí)現(xiàn)”注釋處。

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();

}

1. 使用__global__函數(shù)類型限定符來標(biāo)識(shí)它是一個(gè)核函數(shù)，可以被<<<...>>>調(diào)用；使用__aicore__函數(shù)類型限定符來標(biāo)識(shí)該核函數(shù)在設(shè)備端AI Core上執(zhí)行。指針入?yún)⒆兞啃枰黾幼兞款愋拖薅ǚ鸰_gm__，表明該指針變量指向Global Memory上某處內(nèi)存地址為了統(tǒng)一表達(dá)，使用GM_ADDR宏來修飾入?yún)?，GM_ADDR宏定義如下：

#define GM_ADDR __gm__ uint8_t* __restrict__

2. 算子類的Init函數(shù)，完成內(nèi)存初始化相關(guān)工作，Process函數(shù)完成算子實(shí)現(xiàn)的核心邏輯。

5.2 算子類定義

本樣例中定義KernelAdd算子類，其具體成員如下。請(qǐng)將下文代碼添加至add_custom.cpp的“算子類實(shí)現(xiàn)”注釋處。

class KernelAdd {

public:

__aicore__ inline KernelAdd(){}

// 初始化函數(shù)，完成內(nèi)存初始化相關(guān)操作

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

// 核心處理函數(shù)，實(shí)現(xiàn)算子邏輯，調(diào)用私有成員函數(shù)CopyIn、Compute、CopyOut完成矢量算子的三級(jí)流水操作

__aicore__ inline void Process(){}

private:

// 搬入函數(shù)，完成CopyIn階段的處理，被核心Process函數(shù)調(diào)用

__aicore__ inline void CopyIn(int32_t progress){}

// 計(jì)算函數(shù)，完成Compute階段的處理，被核心Process函數(shù)調(diào)用

__aicore__ inline void Compute(int32_t progress){}

// 搬出函數(shù)，完成CopyOut階段的處理，被核心Process函數(shù)調(diào)用

__aicore__ inline void CopyOut(int32_t progress){}

private:

TPipe pipe; //Pipe內(nèi)存管理對(duì)象

TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY; //輸入數(shù)據(jù)Queue隊(duì)列管理對(duì)象，QuePosition為VECIN

TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueZ; //輸出數(shù)據(jù)Queue隊(duì)列管理對(duì)象，QuePosition為VECOUT

GlobalTensor<half> xGm, yGm, zGm; //管理輸入輸出Global Memory內(nèi)存地址的對(duì)象，其中xGm, yGm為輸入，zGm為輸出

};

內(nèi)部函數(shù)的調(diào)用關(guān)系示意圖如下：

由此可見除了Init函數(shù)完成初始化外，Process中完成了對(duì)流水任務(wù)：“搬入、計(jì)算、搬出”的調(diào)用，開發(fā)者可以重點(diǎn)關(guān)注三個(gè)流水任務(wù)的實(shí)現(xiàn)。

5.3 Init 函數(shù)實(shí)現(xiàn)

初始化函數(shù)Init主要完成以下內(nèi)容：設(shè)置輸入輸出Global Tensor的Global Memory內(nèi)存地址，通過Pipe內(nèi)存管理對(duì)象為輸入輸出Queue分配內(nèi)存。

上文我們介紹到，本樣例將數(shù)據(jù)切分成8塊，平均分配到8個(gè)核上運(yùn)行，每個(gè)核上處理的數(shù)據(jù)大小BLOCK_LENGTH為2048。那么我們是如何實(shí)現(xiàn)這種切分的呢？

每個(gè)核上處理的數(shù)據(jù)地址需要在起始地址上增加GetBlockIdx()*BLOCK_LENGTH（每個(gè)block處理的數(shù)據(jù)長(zhǎng)度）的偏移來獲取。這樣也就實(shí)現(xiàn)了多核并行計(jì)算的數(shù)據(jù)切分。

以輸入x為例，x + BLOCK_LENGTH * GetBlockIdx()即為單核處理程序中x在Global Memory上的內(nèi)存偏移地址，獲取偏移地址后，使用GlobalTensor類的SetGlobalBuffer接口設(shè)定該核上Global Memory的起始地址以及長(zhǎng)度。具體示意圖如下。

?

上面已經(jīng)實(shí)現(xiàn)了多核數(shù)據(jù)的切分，那么單核上的處理數(shù)據(jù)如何進(jìn)行切分？

對(duì)于單核上的處理數(shù)據(jù)，可以進(jìn)行數(shù)據(jù)切塊（Tiling），在本示例中，僅作為參考，將數(shù)據(jù)切分成8塊（并不意味著8塊就是性能最優(yōu)）。切分后的每個(gè)數(shù)據(jù)塊再次切分成2塊，即可開啟double buffer，實(shí)現(xiàn)流水線之間的并行。

這樣單核上的數(shù)據(jù)（2048個(gè)數(shù)）被切分成16塊，每塊TILE_LENGTH（128）個(gè)數(shù)據(jù)。Pipe為inQueueX分配了兩塊大小為TILE_LENGTH * sizeof(half)個(gè)字節(jié)的內(nèi)存塊，每個(gè)內(nèi)存塊能容納TILE_LENGTH（128）個(gè)half類型數(shù)據(jù)。數(shù)據(jù)切分示意圖如下。

?

具體的初始化函數(shù)代碼如下：

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

{

// 多核并行，設(shè)定當(dāng)前核上Global Memory的起始地址以及長(zhǎng)度 xGm.SetGlobalBuffer((__gm__ half*)x + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);

yGm.SetGlobalBuffer((__gm__ half*)y + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);

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

// 通過pipe為queue分配內(nèi)存，單位為Bytes

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));

}

5.4 核心處理函數(shù)實(shí)現(xiàn)

基于矢量編程范式，將核函數(shù)的實(shí)現(xiàn)分為3個(gè)基本任務(wù)：CopyIn，Compute，CopyOut。任務(wù)之間通過隊(duì)列進(jìn)行通信，交互示意圖如下：

?

Process函數(shù)中通過如下方式調(diào)用這三個(gè)函數(shù)。

__aicore__ inline void Process()

{

// 開啟double buffer后循環(huán)次數(shù)需要乘以2

constexpr int32_t loopCount = TILE_NUM * BUFFER_NUM;

// 多個(gè)任務(wù)實(shí)現(xiàn)流水并行

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

CopyIn(i);

Compute(i);

CopyOut(i);

}

}

• CopyIn函數(shù)實(shí)現(xiàn)。

__aicore__ inline void CopyIn(int32_t progress)

{

// 1、從隊(duì)列中分配Tensor

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

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

// 2、使用DataCopy接口將GlobalTensor數(shù)據(jù)拷貝到LocalTensor

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

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

// 3、將LocalTensor放入搬入數(shù)據(jù)的存放位置VecIn的Queue中

inQueueX.EnQue(xLocal);

inQueueY.EnQue(yLocal);

}

• Compute函數(shù)實(shí)現(xiàn)。

__aicore__ inline void Compute(int32_t progress)

{

// 1、使用DeQue從VecIn中取出LocalTensor

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

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

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

// 2、調(diào)用Add指令完成雙目矢量計(jì)算

Add(zLocal, xLocal, yLocal, TILE_LENGTH);

// 3、使用EnQue將計(jì)算結(jié)果LocalTensor放入到搬出數(shù)據(jù)的存放位置VECOUT的Queue中

outQueueZ.EnQue<half>(zLocal);

// 4、使用FreeTensor將釋放不再使用的LocalTensor

inQueueX.FreeTensor(xLocal);

inQueueY.FreeTensor(yLocal);

}

• CopyOut函數(shù)實(shí)現(xiàn)。

__aicore__ inline void CopyOut(int32_t progress)

{

// 1、使用DeQue接口從VecOut的Queue中取出LocalTensor

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

// 2、使用DataCopy接口將LocalTensor拷貝到GlobalTensor上

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

// 3、使用FreeTensor將不再使用的LocalTensor進(jìn)行回收

outQueueZ.FreeTensor(zLocal);

}

6 核函數(shù)運(yùn)行驗(yàn)證

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

除了上文核函數(shù)實(shí)現(xiàn)文件add_custom.cpp外，核函數(shù)的調(diào)用與驗(yàn)證還需要需要準(zhǔn)備以下文件：

• 調(diào)用算子的應(yīng)用程序：main.cpp。
• 輸入數(shù)據(jù)和真值數(shù)據(jù)生成腳本文件：add_custom.py。
• 編譯cpu側(cè)或npu側(cè)運(yùn)行的算子的編譯工程文件：CMakeLists.txt。
• 編譯運(yùn)行算子的腳本：run.sh。

本文僅介紹調(diào)用算子的應(yīng)用程序的編寫，該應(yīng)用程序在main.cpp中體現(xiàn)，其他內(nèi)容您可以在quick-start中直接獲取。

6.1 host側(cè)應(yīng)用程序框架編寫

內(nèi)置宏__CCE_KT_TEST__ 是區(qū)分運(yùn)行CPU模式或NPU模式邏輯的標(biāo)志，在同一個(gè)main函數(shù)中通過對(duì)__CCE_KT_TEST__宏定義的判斷來區(qū)分CPU和NPU側(cè)的運(yùn)行程序。

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__

// 用于CPU調(diào)試的調(diào)用程序



#else

// NPU側(cè)運(yùn)行算子的調(diào)用程序

#endif

return 0;

}

6.2 CPU運(yùn)行驗(yàn)證

完成算子核函數(shù)CPU側(cè)運(yùn)行驗(yàn)證的步驟如下：

1. 分配共享內(nèi)存，并進(jìn)行數(shù)據(jù)初始化；
2. 調(diào)用ICPU_RUN_KF調(diào)測(cè)宏，完成核函數(shù)CPU側(cè)的調(diào)用；
3. 釋放申請(qǐng)的資源。

請(qǐng)將下文代碼添加至上面代碼框架的“用于CPU調(diào)試的調(diào)用程序”注釋處。

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);

6.3 NPU側(cè)運(yùn)行驗(yàn)證

完成算子核函數(shù)NPU側(cè)運(yùn)行驗(yàn)證的步驟如下：

1. 初始化Device設(shè)備；
2. 創(chuàng)建Context綁定設(shè)備；
3. 分配Host內(nèi)存，并進(jìn)行數(shù)據(jù)初始化；
4. 分配Device內(nèi)存，并將數(shù)據(jù)從Host上拷貝到Device上；
5. 用內(nèi)核調(diào)用符<<<>>>調(diào)用核函數(shù)完成指定的運(yùn)算；
6. 將Device上的運(yùn)算結(jié)果拷貝回Host；
7. 釋放申請(qǐng)的資源。

請(qǐng)將下文代碼添加至上面代碼框架的“NPU側(cè)運(yùn)行算子的調(diào)用程序”注釋處。

// AscendCL初始化

CHECK_ACL(aclInit(nullptr));

// 創(chuàng)建Context綁定設(shè)備

aclrtContext context;

int32_t deviceId = 0;

CHECK_ACL(aclrtSetDevice(deviceId));

CHECK_ACL(aclrtCreateContext(&context, deviceId));

aclrtStream stream = nullptr;

CHECK_ACL(aclrtCreateStream(&stream));

// 分配Host內(nèi)存，并進(jìn)行數(shù)據(jù)初始化

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內(nèi)存，并將數(shù)據(jù)從Host上拷貝到Device上

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

CHECK_ACL(aclrtMalloc((void**)&yDevice, outputByteSize, 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));

// 用內(nèi)核調(diào)用符<<<>>>調(diào)用核函數(shù)完成指定的運(yùn)算

add_custom_do(blockDim, nullptr, stream, xDevice, yDevice, zDevice);

CHECK_ACL(aclrtSynchronizeStream(stream));

// 將Device上的運(yùn)算結(jié)果拷貝回Host

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

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

// 釋放申請(qǐng)的資源

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(aclrtDestroyContext(context));

CHECK_ACL(aclrtResetDevice(deviceId));

CHECK_ACL(aclFinalize());

6.4 執(zhí)行一鍵式編譯運(yùn)行腳本，編譯和運(yùn)行應(yīng)用程序

腳本執(zhí)行方式如下：

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

1. <kernel_name>表示需要運(yùn)行的算子。
2. <soc_version>表示算子運(yùn)行的AI處理器型號(hào)。
3. <core_type>表示在AiCore上或者VectorCore上運(yùn)行。
4. <run_mode>表示算子以cpu模式或npu模式運(yùn)行。

1. CPU模式下執(zhí)行如下命令（算子運(yùn)行的AI處理器型號(hào)以Ascend 910為例）：

bash run.sh add_custom ascend910 AiCore cpu

運(yùn)行結(jié)果如下，當(dāng)前使用md5sum對(duì)比了所有輸出bin文件，md5值一致表示實(shí)際的輸出數(shù)據(jù)和真值數(shù)據(jù)相符合。

2. NPU模式下執(zhí)行如下命令：bash run.sh add_custom ascend910 AiCore npu

運(yùn)行結(jié)果如下，當(dāng)前使用md5sum對(duì)比了所有輸出bin文件，md5值一致表示實(shí)際的輸出數(shù)據(jù)和真值數(shù)據(jù)相符合。

至此，你已經(jīng)完成了Ascend C算子開發(fā)的快速入門，更多內(nèi)容請(qǐng)參考：《Ascend C 官方教程》

?文章來源地址http://www.zghlxwxcb.cn/news/detail-684765.html

點(diǎn)擊關(guān)注，第一時(shí)間了解華為云新鮮技術(shù)~

?

到了這里，關(guān)于Ascend C保姆級(jí)教程：我的第一份Ascend C代碼的文章就介紹完了。如果您還想了解更多內(nèi)容，請(qǐng)?jiān)谟疑辖撬阉鱐OY模板網(wǎng)以前的文章或繼續(xù)瀏覽下面的相關(guān)文章，希望大家以后多多支持TOY模板網(wǎng)！