本文分享自華為雲社群《3天上手Ascend C程式設計 | Day2 通過Ascend C程式設計正規化實現一個運算元範例》，作者：昇騰CANN 。

一、Ascend C程式設計正規化

Ascend C程式設計正規化把運算元內部的處理程式，分成多個流水任務( stage )，以張量( Tensor)為資料載體，以佇列 ( Queue ) 進行任務之間的通訊與同步，以記憶體管理模組( Pipe ) 管理任務間的通訊記憶體。

1、流水任務

流水任務指的是單核處理程式中主程式排程的並行任務。在核函數內部，可以通過流水任務實現資料的並行處理，進一步提升效能。下面舉例來說明，流水任務如何進行並行排程。以下面的流水任務示意圖為例，單核處理程式的功能被拆分成3個流水任務：Stage1、Stage2、Stage3，每個任務專注於完成單一功能；需要處理的資料被切分成n片，使用Progress1~n表示，每個任務需要依次完成n個資料切片的處理。Stage間的箭頭表達資料間的依賴關係，比如Stage1處理完Progress1之後，Stage2才能對Progress1進行處理。

若n=3，即待處理的資料被切分成3片，則上圖中的流水任務執行起來的示意圖如下，從執行圖中可以看出，對於同一片資料，Stage1、Stage2、Stage3之間的處理具有依賴關係，需要序列處理；不同的資料切片，同一時間點，可以有多個任務在並行處理，由此達到任務並行、提升效能的目的。

向量（Vector）程式設計正規化把運算元的實現流程分為3個基本任務：CopyIn，Compute，CopyOut。CopyIn負責搬入操作，Compute負責向量計算操作，CopyOut負責搬出操作。

2、任務間通訊與同步

不同的流水任務之間存在資料依賴，需要進行資料傳遞。Ascend C中使用Queue佇列完成任務之間的資料通訊和同步，提供EnQue、DeQue等基礎API。Queue佇列管理NPU上不同層級的實體記憶體時，用一種抽象的邏輯位置（QuePosition）來表達各級別的儲存，代替了片上物理儲存的概念，開發者無需感知硬體架構。

向量程式設計中使用到的邏輯位置（QuePosition）定義如下：

搬入資料的存放位置：VECIN搬出資料的存放位置：VECOUT向量程式設計主要分為CopyIn、Compute、CopyOut三個任務：

CopyIn任務中將輸入資料從Global記憶體搬運至Local記憶體後，需要使用EnQue將LocalTensor放入VECIN的Queue中；

Compute任務等待VECIN的Queue中LocalTensor出隊之後才可以完成向量計算，計算完成後使用EnQue將計算結果LocalTensor放入到VECOUT的Queue中；

CopyOut任務等待VECOUT的Queue中LocalTensor出隊，再將其拷貝到Global記憶體。

Ascend C使用GlobalTensor和LocalTensor作為資料的基本操作單元，它是各種指令API直接呼叫的物件，也是資料的載體。

3、記憶體管理機制

任務間資料傳遞使用到的記憶體統一由記憶體管理模組Pipe進行管理。Pipe作為片上記憶體管理者，通過InitBuffer介面對外提供Queue記憶體初始化功能，開發者可以通過該介面為指定的Queue分配記憶體。

Queue佇列記憶體初始化完成後，需要使用記憶體時，通過呼叫AllocTensor來為LocalTensor分配記憶體，當建立的LocalTensor完成相關計算無需再使用時，再呼叫FreeTensor來回收LocalTensor的記憶體。

程式設計過程中使用到的臨時變數記憶體同樣通過Pipe進行管理。臨時變數可以使用TBuf資料結構來申請指定QuePosition上的儲存空間，並使用Get()來將分配到的儲存空間分配給新的LocaLTensor從TBuf上獲取全部長度，或者獲取指定長度的LocalTensor。

使用TBuf申請的記憶體空間只能參與計算，無法執行Queue佇列的入隊出隊操作。

二、使用Ascend C程式設計正規化實現一個運算元範例

向量運算元開發一般開發流程如下：

下面以add作為例子介紹Ascend C向量運算元的開發流程。完整樣例大家可以參考程式碼樣例。

1、運算元分析

分析運算元的數學表示式、輸入、輸出以及計算邏輯的實現，明確需要呼叫的Ascend C介面。

例子以Add運算元為例，數學公式:z= x+y，為簡單起見，設定輸入張量x, y,z為固定shape(8,2048)，資料型別dtype為half型別，資料排布型別format為ND，核函數名稱為add_custom。

• 運算元的數學表示式及計算邏輯。Add運算元的數學表示式為:z = x + y；計算邏輯：輸入資料需要先搬入到片上儲存，然後使用計算介面完成兩個加法運算，得到最終結果，再搬出到外部儲存。

• 輸入和輸出。Add運算元有兩個輸入:x與y，輸出為z。輸入資料型別為half，輸出資料型別與輸入資料型別相同。輸入支援固定shape(8，2048)輸出shape與輸入shape相同，輸入資料排布型別為ND。

• 確定核函數名稱和引數。自定義核函數名，如add_custom。根據輸入輸出，確定核函數有3個入參x，y，z。x，y為輸入在Global Memory上的記憶體地址，z為輸出在Global Memory上的記憶體地址。

• 確定運算元實現所需介面。涉及內外部儲存間的資料搬運，使用資料搬移介面:Datacopy實現；涉及向量計算的加法操作，使用向量雙目指令:Add實現；使用到LocalTensor，使用Queue佇列管理，會使用到EnQue、DeQue等介面。

2、核函數定義

在add_custom核函數的實現中範例化kernelAdd運算元類，呼叫Init()數完成記憶體初始化，呼叫Process()函數完成核心邏輯。

// 實現核函數
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();
}

3、根據向量程式設計正規化實現運算元類

根據前面的知識，運算元實現三個流水任務CopyIn、Compute、CopyOut。任務間通過佇列VECIN、VECOUT進行通訊和同步，由pipe記憶體管理物件對任務間互動使用到的記憶體、臨時變數使用到的記憶體統一進行管理。如下圖所示：

• CopyIn任務:將Global Memory上的輸入Tensor xGm和yGm搬運至Local Memory，分別儲存在xLocal,yLocal；

• Compute任務:對xLocal,yLocal執行加法操作，計算結果儲存在zLocal中；

• CopyOut任務:將輸出資料從zLocal搬運至Global Memory上的輸出Tensor zGm中

CopyIn，Compute任務間通過VECIN佇列inQueuex，inQueuer進行通訊和同步；compute，copyout任務間通過VECOUT佇列outQueuez進行通訊和同步。

第一步，我們進行運算元類定義：​

第三步，實現Process()函數——CopyIn，Compute、CopyOut三個流水任務：

​最後，基於核心呼叫符方式進行運算元驗證

先使用python指令碼生成x,y，並計算出z（golden）並落盤。然後再用相同的x，y，在cpu和npu模式下呼叫add運算元，計算出結果z，並與python指令碼採用計算md5sum的方式進行對比，完全一樣，則表示結果正確。

為了執行方便，我們使用一個run.sh，寫有cpu和npu模式的編譯命令，通過輸入引數進行選擇cpu或npu模式進行編譯，執行。

1）CPU模式下：

使用ICPU_RUN_KF宏進行CPU偵錯。

ICPU_RUN_KF(add_custom, blockDim, x, y, z); // use this macro for cpu debug

bash run.sh add_custom ascend910 AiCore cpu執行結果：​

2）NPU模式下：

NPU模式使用<<<>>>方式呼叫，由於CPU模式g++沒有<<<>>>的表達，需要使用內建宏 __CCE_KT_TEST。

#ifndef __CCE_KT_TEST__
//call of kernel function
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);
}
#endif

bash run.sh add_custom ascend910 AiCore npu執行結果：​

​更多學習資源

好啦，本次分享結束啦，Ascend C的學習資源還有很多，想深入學習的可以參考官網教學：Ascend C官方教學。

3天上手Ascend C程式設計 | Day1 Ascend C基本概念及常用介面3天上手Ascend C程式設計 | Day2 通過Ascend C程式設計正規化實現一個運算元範例3天上手Ascend C程式設計 | Day3 Ascend C運算元偵錯調優方法

點選關注，第一時間瞭解華為雲新鮮技術~