1.基本概念

1.1 Ascend C

什么是Ascend C？
Ascend c是CANN針對算子開發(fā)場景推出的編程語言，原生支持C和C 標(biāo)準(zhǔn)規(guī)范，最大化匹配用戶開發(fā)習(xí)慣;通過多層接口抽象、自動并行計算、李生調(diào)試等關(guān)鍵技術(shù)，極大提高算子開發(fā)效率，助力AI開發(fā)者低成本完成算子開發(fā)和模型調(diào)優(yōu)部署。
使用Ascend C開發(fā)自定義算子的優(yōu)勢:

• C/C 原語編程，最大化匹配用戶的開發(fā)習(xí)慣
• 編程模型屏蔽硬件差異，編程范式提高開發(fā)效率
• 多層級API封裝，從簡單到靈活，兼顧易用與高效
• 李生調(diào)試，CPU側(cè)模擬NPU側(cè)的行為，可優(yōu)先在CPU側(cè)調(diào)試

1.2 CANN

CANN是華為針對AI場景推出的異構(gòu)計算架構(gòu)，本次活動主要聚焦的是其中的算子開發(fā)的部分。

1.3應(yīng)用場景

將host == 服務(wù)器，Device就是華為的NPU。而一張NPU中有多個Aicore核心。

Ascend C能夠為華為AI加速卡在大規(guī)模神經(jīng)網(wǎng)絡(luò)計算加速。

1.4 AIcore

首先，既然提到了AIcore,那具體AIcore能做什么呢？

Aicore支持核心計算，分別是：

• 標(biāo)量（scalar)
• 向量（vector）
• 矩陣（cube）

以上圖為例，AI Core中包含計算單元、存儲單元、搬運單元等核心組件。

• 計算單元包括了三種基礎(chǔ)計算資源：*Cube計算單元、Vector計算單元和Scalar計算單元。
• 存儲單元即為AI Core的內(nèi)部存儲，統(tǒng)稱為Local Memory，與此相對應(yīng)，AI Core的外部存儲稱之為Global Memory。
• DMA搬運單元負(fù)責(zé)在Global Memory和Local Memory之間搬運數(shù)據(jù)。

且針對存在在不同區(qū)域中的數(shù)據(jù)類型，不論其原本的數(shù)據(jù)類型是什么（int,float…)

我們將用于存放AI Core中Local Memory（內(nèi)部存儲）的數(shù)據(jù)成為Local Tensor,

將用于存放AI Core中Gocal Memory（內(nèi)部存儲）的數(shù)據(jù)成為Gocal Tensor,

1.5 并行計算常見模型

并行計算常見模型有兩種，SPMD（Single-Program Multiple-Data）數(shù)據(jù)并行 和 流水線并行。

前者SPMD將數(shù)據(jù)切分成不同部分，經(jīng)多個進程處理，最好一同輸出。

后者流水線同樣是將數(shù)據(jù)切分，同時將進程的任務(wù)拆分成多個任務(wù)，全部數(shù)據(jù)如流水線操作一般，與SPMD不同，每個進程只會專注于一個任務(wù)的處理，會處理所有的數(shù)據(jù)分片。

2.編程模型與范式

上面介紹的主要是Ascend C中的一些基礎(chǔ)概念。接下來主要介紹編程模型與范式。

編程模型主要由三個部分組成：

• 并行編程SPMD
• 核函數(shù)
• API

2.1 并行編程SPMD

Ascend C算子編程是SPMD（Single-Program Multiple-Data）編程，具體到Ascend C編程模型中的應(yīng)用，是將需要處理的數(shù)據(jù)被拆分并同時在多個計算核心（類比于上文介紹中的多個進程）上運行，從而獲取更高的性能。多個AI Core共享相同的指令代碼，每個核上的運行實例唯一的區(qū)別是block_idx不同，每個核通過不同的block_idx來識別自己的身份。block的概念類似于進程的概念，block_idx就是標(biāo)識進程唯一性的進程ID。編程中使用函數(shù)GetBlockldx()獲取ID。并行計算過程的示意圖如下圖所示。

2.2 核函數(shù)

從SPMD模型可以得知，使用Ascend C進行編程時，我們編寫一份算子實現(xiàn)代碼，算子被調(diào)用時，將啟動N個運行示例，在N個核上運行。本節(jié)將介紹算子實現(xiàn)的入口函數(shù)。

核函數(shù)（Kernel Function）是Ascend C算子設(shè)備側(cè)實現(xiàn)的入口。在核函數(shù)中，需要為在一個核上執(zhí)行的代碼規(guī)定要進行的數(shù)據(jù)訪問和計算操作，當(dāng)核函數(shù)被調(diào)用時，多個核都執(zhí)行相同的核函數(shù)代碼，具有相同的參數(shù)，并行執(zhí)行。

Ascend C允許用戶使用核函數(shù)這種C/C 函數(shù)的語法擴展來管理設(shè)備端的運行代碼，用戶在核函數(shù)中進行算子類對象的創(chuàng)建和其成員函數(shù)的調(diào)用，由此實現(xiàn)該算子的所有功能。核函數(shù)是主機端和設(shè)備端連接的橋梁。

//核函數(shù)的聲明
extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z);

這里其實可以看出，核函數(shù)的聲明和普通C 函數(shù)聲明大有不同。其中

global和aicore是函數(shù)類型限定符，使用global函數(shù)類型限定符來標(biāo)識它是一個核函數(shù)，可以被<<<…>>>調(diào)用；使用aicore函數(shù)類型限定符來標(biāo)識該核函數(shù)在設(shè)備端AI Core上執(zhí)行。參數(shù)中的gm則表示存儲在Global memory中。

編程中使用到的函數(shù)可以分為三類：核函數(shù)（device側(cè)執(zhí)行）、host側(cè)執(zhí)行函數(shù)、device側(cè)執(zhí)行函數(shù)（除核函數(shù)之外的）。三者的調(diào)用關(guān)系如下圖所示：

• host側(cè)執(zhí)行函數(shù)可以調(diào)用同類的host執(zhí)行函數(shù)，也就是通用C/C 編程中的函數(shù)調(diào)用；也可以通過<<<>>>調(diào)用核函數(shù)。
• device側(cè)執(zhí)行函數(shù)（除核函數(shù)之外的）可以調(diào)用調(diào)用同類的device執(zhí)行函數(shù)。
• 核函數(shù)可以調(diào)用device側(cè)執(zhí)行函數(shù)（除核函數(shù)之外的）。

這里也可以看出核函數(shù)是作為host側(cè)核Device側(cè)之間的橋梁，讓兩邊的執(zhí)行函數(shù)連接起來。

除此之外，還有兩條核函數(shù)應(yīng)該遵守的規(guī)則：

1. 核函數(shù)必須具有void返回類型。
2. 僅支持入?yún)橹羔樆駽/C 內(nèi)置數(shù)據(jù)類型（Primitive data types），如：half* s0，float* s1、int32_t c。

//調(diào)用核函數(shù)
kernel_name<<<blockDim, l2ctrl, stream>>>(argument list);
// blockDim設(shè)置為8表示在8個核上調(diào)用了add_custom核函數(shù)，每個核都會獨立且并行地執(zhí)行該核函數(shù)，該核函數(shù)的參數(shù)列表為x，y，z。
add_custom<<<8, nullptr, stream>>>(x, y, z);

執(zhí)行配置由3個參數(shù)決定：

• blockDim，規(guī)定了核函數(shù)將會在幾個核上執(zhí)行。每個執(zhí)行該核函數(shù)的核會被分配一個邏輯ID，即blockidx，可以在核函數(shù)的實現(xiàn)中調(diào)用[GetBlockIdx](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/70RC1alpha003/operatordevelopment/ascendcopdevg/atlasascendcapi070129.html)來獲取block_idx；
• l2ctrl，保留參數(shù)，暫時設(shè)置為固定值nullptr，開發(fā)者無需關(guān)注；
• stream，類型為aclrtStream，stream是一個任務(wù)隊列，應(yīng)用程序通過stream來管理任務(wù)的并行。

下方是Add算子的例子：

// 實現(xiàn)核函數(shù)
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
// 初始化算子類，算子類提供算子初始化和核心處理等方法
KernelAdd op;
// 初始化函數(shù)，獲取該核函數(shù)需要處理的輸入輸出地址，同時完成必要的內(nèi)存初始化工作
op.Init(x, y, z);
// 核心處理函數(shù)，完成算子的數(shù)據(jù)搬運與計算等核心邏輯
op.Process();
}
// 調(diào)用核函數(shù)
void add_custom_do(uint32_t blockDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y, uint8_t* z)
{
add_custom<<<blockDim, l2ctrl, stream>>>(x, y, z);
}

2.3 API

Ascend C算子采用標(biāo)準(zhǔn)C 語法和一組類庫API進行編程，類庫API主要包含以下幾種，您可以在核函數(shù)的實現(xiàn)中根據(jù)自己的需求選擇合適的API：

• 計算類API，包括標(biāo)量計算API、向量計算API、矩陣計算API，分別實現(xiàn)調(diào)用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接口。
• 內(nèi)存管理API，用于分配管理內(nèi)存，比如AllocTensor、FreeTensor接口。
• 任務(wù)同步API，完成任務(wù)間的通信和同步，比如EnQue、DeQue接口。不同的API指令間有可能存在依賴關(guān)系，從AI Core內(nèi)部并行計算架構(gòu)抽象可知，不同的指令異步并行執(zhí)行，為了保證不同指令隊列間的指令按照正確的邏輯關(guān)系執(zhí)行，需要向不同的組件發(fā)送同步指令。任務(wù)同步類API內(nèi)部即完成這個發(fā)送同步指令的過程，開發(fā)者無需關(guān)注內(nèi)部實現(xiàn)邏輯，使用簡單的API接口即可完成。

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

這里簡單理解，就是在調(diào)用同樣功能的API時，0級的計算性能會比其它等級的API更好。（0>1>2>3)

2.4 編程范式

在有了上述的核函數(shù)，API，并行編程等作為工具之后，編程范式描述了算子實現(xiàn)的固定流程。

把算子核內(nèi)的處理程序，分成多個流水任務(wù)，通過隊列（Queue）完成任務(wù)間**通信和同步，并通過統(tǒng)一的內(nèi)存管理**模塊（Pipe）管理任務(wù)間通信內(nèi)存。流水編程范式應(yīng)用了流水線并行計算方法。它提供了：

• 快速開發(fā)編程的固定步驟
• 統(tǒng)一代碼框架的開發(fā)捷徑
• 使用者總結(jié)出的開發(fā)經(jīng)驗
• 面向特定場景的編程思想
• 定制化的方法論開發(fā)體驗

以下圖為例解釋流水任務(wù)，其可以看作兩種并行計算常見方法的組合，即將數(shù)據(jù)切分后，將線程任務(wù)也切分，通過多線程快速處理切分的數(shù)據(jù)。

Ascend C分別針對Vector、Cube編程設(shè)計了不同的流水任務(wù)，

• Vector編程范式把算子的實現(xiàn)流程分為3個基本任務(wù)：CopyIn，Compute，CopyOut。CopyIn負(fù)責(zé)搬入操作，Compute負(fù)責(zé)矢量計算操作，CopyOut負(fù)責(zé)搬出操作。
• Cube編程范式把算子的實現(xiàn)流程分為5個基本任務(wù)：CopyIn，Split，Compute，Aggregate，CopyOut。CopyIn負(fù)責(zé)搬入操作，Split負(fù)責(zé)數(shù)據(jù)切分操作，Compute負(fù)責(zé)矩陣指令計算操作，Aggregate負(fù)責(zé)數(shù)據(jù)匯聚操作，CopyOut負(fù)責(zé)搬出操作。

上文中提到，進行編程范式需要將數(shù)據(jù)切分，不同的流水任務(wù)之間存在數(shù)據(jù)依賴，那如何保持任務(wù)間通信和同步？

Ascend C中使用Queue隊列完成任務(wù)之間的數(shù)據(jù)通信和同步，提供EnQue、DeQue等基礎(chǔ)API。我們以矢量（vector)編程中的流程為例。矢量編程中使用到的邏輯位置（QuePosition）定義如下：

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

1. Stage1：CopyIn任務(wù)。

• 使用DataCopy接口將GlobalTensor數(shù)據(jù)拷貝到LocalTensor。使用EnQue將LocalTensor放入VECIN的Queue中。

1. Stage2：Compute任務(wù)。

• 使用DeQue從VECIN中取出LocalTensor。使用Ascend C接口完成矢量計算。使用EnQue將計算結(jié)果LocalTensor放入到VECOUT的Queue中。

1. Stage3：CopyOut任務(wù)。

• 使用DeQue接口從VECOUT的Queue中取出LocalTensor。使用DataCopy接口將LocalTensor拷貝到GlobalTensor上。

cube的編程范式與Vector類似，只是多了spilt和aggreagte的環(huán)節(jié)。

且對于VECIN 和 VECOUT 等queue的創(chuàng)建和刪除，任務(wù)間數(shù)據(jù)傳遞使用到的內(nèi)存統(tǒng)一由內(nèi)存管理模塊Pipe進行管理。如下圖所示，Pipe作為片上內(nèi)存管理者，通過InitBuffer接口對外提供Queue內(nèi)存初始化功能，開發(fā)者可以通過該接口為指定的Queue分配內(nèi)存。

Queue隊列內(nèi)存初始化完成后，需要使用內(nèi)存時，通過調(diào)用AllocTensor來為LocalTensor分配內(nèi)存，當(dāng)創(chuàng)建的LocalTensor完成相關(guān)計算無需再使用時，再調(diào)用FreeTensor來回收LocalTensor的內(nèi)存。

這里和C/C 的內(nèi)存管理有相似的地方，即new 和 delete[] 需要成對出現(xiàn)。這樣對資源的管理在Ascend C 的編程中會經(jīng)?？匆?。

3.Helloworld實例

//在代碼中，由于需要分別在CPU和NPU中調(diào)式，所以我們會使用__CCE_KT_TEST__來表示不同的調(diào)用程序。
#ifdef __CCE_KT_TEST__
// 用于CPU調(diào)試的調(diào)用程序
#else
// NPU側(cè)運行算子的調(diào)用程序
#endif

宏是個很好用的工具，在之后我們還會接觸到宏函數(shù)和其它的一些宏定義。