【2023 · CANN訓(xùn)練營第一季】第一節(jié)課程筆記-算子開發(fā)基礎(chǔ)

一、算子開發(fā)優(yōu)勢簡介

1.語言：C/C++

2.編程模式屏蔽硬件差異，編程范式提高編程效率

3.多層級API封裝，從簡單到靈活，兼顧易用與高效

4.孿生調(diào)試，CPU側(cè)模仿NPU側(cè)行為，可現(xiàn)在CPU側(cè)調(diào)試

二、核函數(shù)簡介

基本概念：

核函數(shù) (Kernel Function)是TIK C++算子設(shè)備側(cè)的入口。TIK C++允許用戶使用核函數(shù)這種C/C++函數(shù)的語法擴展來管理設(shè)備側(cè)的運行代碼，用戶在核函數(shù)中實現(xiàn)算子邏輯的編寫，例如自定義算子類及其成員函數(shù)以實現(xiàn)該算子的所有功能。核函數(shù)是主機側(cè)和設(shè)備側(cè)連接的橋梁。

核函數(shù)是直接在設(shè)備側(cè)執(zhí)行的代碼。在核函數(shù)中，需要為在一個核上執(zhí)行的代碼規(guī)定要進行的數(shù)據(jù)訪問和計算操作，當(dāng)核函數(shù)被調(diào)用時，多個核將并行執(zhí)行同一個計算任務(wù)

基本構(gòu)成圖如下：

使用函數(shù)類型限定符

上圖的__global__和__aicore__是其兩種類型限定符，用來標(biāo)識它是一個核函數(shù)，使用<<<...>>>來進行調(diào)用,aicore是用來表示該核函數(shù)是用來在設(shè)備側(cè)AI Core上執(zhí)行?

還有一種入?yún)⒆兞康南薅ǚ?

為了方便，統(tǒng)一使用__gm__uint8_t*。用戶可統(tǒng)一使用uint8_t,并在使用時轉(zhuǎn)化為實際的指針類型;亦可直接傳入實際的指針類型.

其他規(guī)則：

1.必須具有void返回類型

2.使用extern"C"

3.僅支持入?yún)橹羔橆愋突駽/C++內(nèi)置數(shù)據(jù)類型(Primitive Data Tvpes)，如: half* s0、float* s1、int32_t

函數(shù)調(diào)用方式：

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

• blockDim，規(guī)定了核函數(shù)將會在幾個核上執(zhí)行，每個執(zhí)行該核函數(shù)的核會被分配一個邏輯ID，表現(xiàn)為內(nèi)置變量block idx，編號從0開始，可為不同的邏輯核定義不同的行為，可以在算子實現(xiàn)中使用

• l2ctrl，保留參數(shù)，暫時設(shè)置為固定值nullptr

• stream，類型為acltStream，stream是一個任務(wù)隊列，應(yīng)用程序通過stream來管理任務(wù)的并行

eg:

HelloWorld<<<8，nullptr，stream>>>(fooDevice);

blockDim設(shè)置為8，表示在8個核上調(diào)用了HelLoworld核函數(shù)，每個核都會獨立且并行地執(zhí)行該核函數(shù)

aclrtstreamStream可以通過aclrtCreatestream來創(chuàng)建,它的作用是在當(dāng)前進程或線程中顯式創(chuàng)建一個arqument List設(shè)置為fooDevice這1個入?yún)?/p>

三、算子執(zhí)行的模式

算子的執(zhí)行分為CPU模式和NPU模式

代碼可以使用宏定義來進行分割CPU運行還是NPU運行

四、接口簡介

常用數(shù)據(jù)定義：

GlobalTensor：

用來存放Global Memory(外部存儲)的全局數(shù)據(jù)。

LocalTensor：

用于存放核上Local Memory (內(nèi)部存儲)的數(shù)據(jù)

矢量計算指令接口簡介：

矢量計算指令接口，能夠啟動AI Core中的Vector單元執(zhí)行計算為了降低開發(fā)者的使用門檻，指令按照由易到難，分成了3級到0級接口。其中3級接口最為簡單，0級接口最為復(fù)雜，(1級接口還末發(fā)布)

多層級API封裝的作用:

• 降低復(fù)雜指令的使用難度

• 跨代兼容性保障

• 保留最大靈活度的可能

3級接口

3級接口，運算符重載，支持+，-,*,/，|, &, ^,等C/C++運算符實現(xiàn)2級接口的簡化表達

允許用戶使用形如: dst = src*sr1，針對整個Tensor進行計算

2級接口

2級連續(xù)計算接口，針對源操作數(shù)srcLocal的連續(xù)COUNT個數(shù)據(jù)進行計算，并連續(xù)寫入目的操作數(shù)dstLocal，提供了一維Tensor的連續(xù)COUNT個數(shù)據(jù)的計算支持。

0級接口

0級功能靈活計算接口，是最底層的開發(fā)接口，可以完整發(fā)揮硬件優(yōu)勢的計算API，可以進行非連續(xù)的計算該功能可以充分發(fā)揮CANN系列芯片的強大功能指令，支持對每個操作數(shù)的Block stride，Repeat stride，MASK的操作，允許用戶使用諸多的通用參數(shù)來定制化所需要的操作

1.Repeat times：迭代的次數(shù)

矢量計算單元，一次最多可以計算256Bytes的數(shù)據(jù)，每次讀取連續(xù)的8個block (每個block 32Bytes，共256Bytes)數(shù)據(jù)進行計算，為完成對輸入數(shù)據(jù)的處理，必須通過多次迭代(repeat)才能完成所有數(shù)據(jù)的讀取與計算。

2.Repeat stride：相鄰迭代間相同block的地址步長

當(dāng)Repeat times大于1，需要多次迭代完成矢量計算時，可以根據(jù)不同的使用場景合理設(shè)置相鄰迭代間相同block的地址步長Repeat stride的值

• 連續(xù)計算場景:

假設(shè)定義一個Tensor供目的操作數(shù)和源操作數(shù)同時使用(即地址重疊)，Repeat stride取值為8.此時，矢量計算單元第一次迭代讀取連續(xù)8個block，第二輪迭代讀取下一個連續(xù)的8個block，通過多次選代即可完成所有輸入數(shù)據(jù)的計算

• 非連續(xù)計算場景:

Repeat stride取值大8(如取10)時，則相鄰選代間矢量計算單元讀取的數(shù)據(jù)在地址上不連續(xù)，出現(xiàn)2個block的間隔

• 反復(fù)計算場景:

Repeat stride取值為0時，矢量計算單元會對首個連續(xù)的8個block進行反復(fù)讀取和計算

• 部分重復(fù)計算:

Repeat stride取值大于且小于8時，相鄰選代間部分數(shù)據(jù)會被矢量計算單元重復(fù)讀取和計算，此種情形一般場景不涉及

3.Block stride：表示同一迭代內(nèi)不同block的地址步長

如果需要控制單次迭代內(nèi)，數(shù)據(jù)處理的步長，可以通過設(shè)置同一迭代內(nèi)不同block的地址步長Block stride來實現(xiàn)連續(xù)計算，Block stride 設(shè)置為1，對同一迭代內(nèi)的8個block數(shù)據(jù)連續(xù)進行處理非連續(xù)計算，Block stride值大于1(如取2)，同一選代內(nèi)不同block之間在讀取數(shù)據(jù)時出現(xiàn)一個block的間隔

4.Mask參數(shù)

Mask用于控制每次迭代內(nèi)參與計算的元素?？赏ㄟ^連續(xù)模式和逐比特模式兩種方式進行設(shè)置

• 連續(xù)模式:

表示前面連續(xù)的多少個元素參與計算。數(shù)據(jù)類型為uint64_t。取值范圍和操作數(shù)的數(shù)據(jù)類型有關(guān)，數(shù)據(jù)類型不同，每次迭代內(nèi)能夠處理的元素個數(shù)最大值不同(當(dāng)前數(shù)據(jù)類型單次迭代時能處理的元素個數(shù)最大值為: 256/sizeof(數(shù)據(jù)類型))。當(dāng)操作數(shù)的數(shù)據(jù)類型占比特位16位時(如half，uint16_t)，mask[1,128];當(dāng)操作數(shù)為32位時(如float, int32_t) ，maske取值范圍[1, 64]。

• 逐比特模式:

可以按位控制哪些元素參與計算，比特位的值為1表示參與計算，0表示不參與。參數(shù)類型為長度為2的uint64_t類型數(shù)組 參數(shù)取值范圍和操作數(shù)的數(shù)據(jù)類型有關(guān)，數(shù)據(jù)類型不同，每次迭代內(nèi)能夠處理的元素個數(shù)最大值不同。當(dāng)操作數(shù)為16位時，mask[0]、mask[1]取值范圍[0, 264-11; 當(dāng)dst/src為32位時，mask[1]為0，maski[0]取值范圍為[0,(0.2的64次方)-1]

三等級通用接口

ps:該文僅是為了記錄CANN訓(xùn)練營的學(xué)習(xí)過程所用，不參與任何商業(yè)用途