【2023 · CANN訓(xùn)練營第一季】——Ascend C算子代碼分析—Add算子（內(nèi)核調(diào)用符方式）

前言：Ascend C算子（TIK C++）使用C/C++作為前端開發(fā)語言椭符，通過四層接口抽象恨胚、并行編程范式、孿生調(diào)試等技術(shù)徘键，極大提高算子開發(fā)效率练对，助力AI開發(fā)者低成本完成算子開發(fā)和模型調(diào)優(yōu)部署。學(xué)習(xí)完理論后吹害，上代碼螟凭，通過實(shí)踐理解Ascend C算子的概念，掌握開發(fā)流程它呀，以及內(nèi)核調(diào)用符方式的調(diào)試方法螺男。

一、算子分析

????????Add算子的數(shù)學(xué)公式:z= x+y纵穿，為簡單起見下隧，設(shè)定輸入張量x, y,z為固定shape(8,2048)，數(shù)據(jù)類型dtype為half類型谓媒，數(shù)據(jù)排布類型format為ND淆院。

????????確定如下內(nèi)容：

????????1、計(jì)算邏輯:輸入數(shù)據(jù)需要先搬入到片上存儲句惯，然后使用計(jì)算接口（TIK C++ API/矢量計(jì)算/雙目/ADD土辩，采用2級接口）完成兩個加法運(yùn)算，得到最終結(jié)果抢野，再搬出到外部存儲拷淘。

????????2、輸入與輸出

????????輸入：x,y：固定shape(8指孤，2048)辕棚，數(shù)據(jù)排布類型為ND。

????????輸出：z：與輸入相同邓厕，固定shape(8逝嚎，2048)，數(shù)據(jù)排布類型為ND详恼。

????????3补君、核函數(shù)名稱和入?yún)?/p>

????????核函數(shù)名稱：定義為add_tik2

????????入?yún)?個，x昧互，y挽铁，z：x，y為輸入向量在Global Memory上的內(nèi)存地址敞掘，z為計(jì)算結(jié)果輸出到Global Memory上的內(nèi)存地址叽掘。?

二、代碼分析

????代碼結(jié)構(gòu)：

一）算子實(shí)現(xiàn)——Add_tik2.cpp

1玖雁、核函數(shù)定義

extern "C" __global__ __aicore__ void add_tik2(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z)

2更扁、核函數(shù)實(shí)現(xiàn)——算子類的init()和process()

1）在核函數(shù)里實(shí)例化算子類KernelAdd，并調(diào)用init()實(shí)現(xiàn)初始化；調(diào)用process()實(shí)現(xiàn)流水操作

extern "C" __global__ __aicore__ void add_tik2(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z)

{

? ? KernelAdd op;

? ? op.Init(x, y, z);

? ? op.Process();

}

2）KernelAdd算子類定義

class KernelAdd {

public:

? ? __aicore__ inline KernelAdd() {}

? ? __aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z)

? ? {

? ? ? ? // get start index for current core, core parallel

? ? ? ? xGm.SetGlobalBuffer((__gm__ half*)x + block_idx * BLOCK_LENGTH, BLOCK_LENGTH);

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

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

? ? ? ? // pipe alloc memory to queue, the unit is 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));

? ? }

? ? __aicore__ inline void Process()

? ? {

? ? ? ? // loop count need to be doubled, due to double buffer

? ? ? ? constexpr int32_t loopCount = TILE_NUM * BUFFER_NUM;

? ? ? ? // tiling strategy, pipeline parallel

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

? ? ? ? ? ? CopyIn(i);

? ? ? ? ? ? Compute(i);

? ? ? ? ? ? CopyOut(i);

? ? ? ? }

? ? }

private:

? ? __aicore__ inline void CopyIn(int32_t progress)

? ? {

? ? ? ? // alloc tensor from queue memory

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

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

? ? ? ? // copy progress_th tile from global tensor to local tensor

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

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

? ? ? ? // enque input tensors to VECIN queue

? ? ? ? inQueueX.EnQue(xLocal);

? ? ? ? inQueueY.EnQue(yLocal);

? ? }

? ? __aicore__ inline void Compute(int32_t progress)

? ? {

? ? ? ? // deque input tensors from VECIN queue

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

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

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

? ? ? ? // call Add instr for computation

? ? ? ? Add(zLocal, xLocal, yLocal, TILE_LENGTH);

? ? ? ? // enque the output tensor to VECOUT queue

? ? ? ? outQueueZ.EnQue<half>(zLocal);

? ? ? ? // free input tensors for reuse

? ? ? ? inQueueX.FreeTensor(xLocal);

? ? ? ? inQueueY.FreeTensor(yLocal);

? ? }

? ? __aicore__ inline void CopyOut(int32_t progress)

? ? {

? ? ? ? // deque output tensor from VECOUT queue

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

? ? ? ? // copy progress_th tile from local tensor to global tensor

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

? ? ? ? // free output tensor for reuse

? ? ? ? outQueueZ.FreeTensor(zLocal);

? ? }

private:

? ? TPipe pipe;

? ? // create queues for input, in this case depth is equal to buffer num

? ? TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;

? ? // create queue for output, in this case depth is equal to buffer num

? ? TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueZ;

? ? GlobalTensor<half> xGm, yGm, zGm;

};

3）算子類——init()

__aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z)

{

? ? // get start index for current core, core parallel

? ? xGm.SetGlobalBuffer((__gm__ half*)x + block_idx * BLOCK_LENGTH, BLOCK_LENGTH);

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

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

? ? // pipe alloc memory to queue, the unit is 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));

}

4）算子類——process()

__aicore__ inline void Process()

{

? ? // loop count need to be doubled, due to double buffer

? ? constexpr int32_t loopCount = TILE_NUM * BUFFER_NUM;

? ? // tiling strategy, pipeline parallel

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

? ? ? ? CopyIn(i);

? ? ? ? Compute(i);

? ? ? ? CopyOut(i);

? ? }

}

__aicore__ inline void CopyIn(int32_t progress)

{

? ? // alloc tensor from queue memory

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

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

? ? // copy progress_th tile from global tensor to local tensor

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

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

? ? // enque input tensors to VECIN queue

? ? inQueueX.EnQue(xLocal);

? ? inQueueY.EnQue(yLocal);

}

__aicore__ inline void Compute(int32_t progress)

{

? ? // deque input tensors from VECIN queue

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

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

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

? ? // call Add instr for computation

? ? Add(zLocal, xLocal, yLocal, TILE_LENGTH);

? ? // enque the output tensor to VECOUT queue

? ? outQueueZ.EnQue<half>(zLocal);

? ? // free input tensors for reuse

? ? inQueueX.FreeTensor(xLocal);

? ? inQueueY.FreeTensor(yLocal);

}

__aicore__ inline void CopyOut(int32_t progress)

{

? ? // deque output tensor from VECOUT queue

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

? ? // copy progress_th tile from local tensor to global tensor

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

? ? // free output tensor for reuse

? ? outQueueZ.FreeTensor(zLocal);

}

二）算子驗(yàn)證

1浓镜、算子調(diào)用——main.c

1）CPU方式——通過ICPU_RUN_KF宏調(diào)用

#ifdef __CCE_KT_TEST__

? ? uint8_t* x = (uint8_t*)tik2::GmAlloc(inputByteSize);

? ? uint8_t* y = (uint8_t*)tik2::GmAlloc(inputByteSize);

? ? uint8_t* z = (uint8_t*)tik2::GmAlloc(outputByteSize);

? ? ReadFile("./input/input_x.bin", inputByteSize, x, inputByteSize);

? ? // PrintData(x, 16, printDataType::HALF);

? ? ReadFile("./input/input_y.bin", inputByteSize, y, inputByteSize);

? ? // PrintData(y, 16, printDataType::HALF);

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

? ? // PrintData(z, 16, printDataType::HALF);

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

? ? tik2::GmFree((void *)x);

? ? tik2::GmFree((void *)y);

? ? tik2::GmFree((void *)z);

2）NPU方式——內(nèi)核調(diào)用符方式

使用NPU方式溃列，需要按照AscendCL的編程流程調(diào)用。

#ifdef __CCE_KT_TEST__

//cpu 方式

#else

? ? aclInit(nullptr);

? ? aclrtContext context;

? ? aclError error;

? ? int32_t deviceId = 0;

? ? aclrtCreateContext(&context, deviceId);

? ? aclrtStream stream = nullptr;

? ? aclrtCreateStream(&stream);

? ? uint8_t *xHost, *yHost, *zHost;

? ? uint8_t *xDevice, *yDevice, *zDevice;

? ? aclrtMallocHost((void**)(&xHost), inputByteSize);

? ? aclrtMallocHost((void**)(&yHost), inputByteSize);

? ? aclrtMallocHost((void**)(&zHost), outputByteSize);

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

? ? aclrtMalloc((void**)&yDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST);

? ? aclrtMalloc((void**)&zDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST);

? ? ReadFile("./input/input_x.bin", inputByteSize, xHost, inputByteSize);

? ? // PrintData(xHost, 16, printDataType::HALF);

? ? ReadFile("./input/input_y.bin", inputByteSize, yHost, inputByteSize);

? ? // PrintData(yHost, 16, printDataType::HALF);

? ? aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE);

? ? aclrtMemcpy(yDevice, inputByteSize, yHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE);

? ? add_tik2_do(blockDim, nullptr, stream, xDevice, yDevice, zDevice); // call kernel in this function

? ? aclrtSynchronizeStream(stream);

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

? ? // PrintData(zHost, 16, printDataType::HALF);

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

? ? aclrtFree(xDevice);

? ? aclrtFree(yDevice);

? ? aclrtFree(zDevice);

? ? aclrtFreeHost(xHost);

? ? aclrtFreeHost(yHost);

? ? aclrtFreeHost(zHost);

? ? aclrtDestroyStream(stream);

? ? aclrtResetDevice(deviceId);

? ? aclFinalize();

#endif

實(shí)質(zhì)上膛薛，使用的是內(nèi)核調(diào)用符方式：<<<>>>

#ifndef __CCE_KT_TEST__

// call of kernel function

void add_tik2_do(uint32_t blockDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y, uint8_t* z)

{

? ? add_tik2<<<blockDim, l2ctrl, stream>>>(x, y, z);

}

#endif

2听隐、算子驗(yàn)證

????????通過numpy生成輸入x，y的值哄啄，并計(jì)算出x+y的值作為精度比對基準(zhǔn)雅任，上述三個數(shù)據(jù)落盤存儲，然后調(diào)用寫好的add算子在CPU模式和npu模式下分別以落盤的x咨跌，y作為輸入椿访，計(jì)算出結(jié)果z，并于numpy的計(jì)算結(jié)果進(jìn)行對比虑润，驗(yàn)證成玫。采用計(jì)算md5方式比較add算子和numpy對相同輸入的計(jì)算結(jié)果，兩者md5相同拳喻，則兩個文件完全相同哭当。

1）生成基準(zhǔn)數(shù)據(jù)——add_tik2.py

????????用numpy的隨機(jī)生成輸入：input_x和input_y，并計(jì)算出input_x+input_y的值golden作為比對基準(zhǔn)數(shù)據(jù)冗澈，并落盤存儲钦勘。

import numpy as np

def gen_golden_data_simple():

? ? input_x = np.random.uniform(-100, 100, [8, 2048]).astype(np.float16)

? ? input_y = np.random.uniform(-100, 100, [8, 2048]).astype(np.float16)

? ? golden = (input_x + input_y).astype(np.float16)

? ? input_x.tofile("./input/input_x.bin")

? ? input_y.tofile("./input/input_y.bin")

? ? golden.tofile("./output/golden.bin")

if __name__ == "__main__":

? ? gen_golden_data_simple()

2）數(shù)據(jù)比對

????????直接比較算子計(jì)算結(jié)果和基準(zhǔn)數(shù)據(jù)的md5，兩者相同亚亲，則數(shù)據(jù)完全相同彻采。在run.sh的末尾處。

# 驗(yàn)證計(jì)算結(jié)果

echo "md5sum: ";md5sum output/*.bin

三捌归、運(yùn)行調(diào)試

????????本次訓(xùn)練營沒有提供開發(fā)環(huán)境肛响，提供了一個沙箱，沙箱已經(jīng)安裝好了開發(fā)環(huán)境惜索。首先把代碼搞沙箱里面特笋。老師為了簡化操作，提前將cpu和npu模式下的編譯和運(yùn)行巾兆，封裝到腳本run.sh中猎物。使用腳本命令分別執(zhí)行CPU或NPU模式下的調(diào)試。

????????一）CPU模式下運(yùn)行角塑、調(diào)試

????????1蔫磨、編譯、運(yùn)行：

bash run.sh add_tik2 ascend910 aicore cpu

????????編譯及運(yùn)行結(jié)果：

?????2圃伶、gdb調(diào)試：

????????使用gdb單步調(diào)試算子計(jì)算精度堤如，也可以在代碼中直接編寫printf(...)來觀察數(shù)值的輸出蒲列。由于cpu調(diào)測已轉(zhuǎn)為多進(jìn)程調(diào)試，每個核都是一個獨(dú)立的子進(jìn)程煤惩，故gdb需要轉(zhuǎn)換成子進(jìn)程調(diào)試的方式。

????????在gdb啟動后炼邀，首先設(shè)置跟蹤子進(jìn)程魄揉，之后再打斷點(diǎn)，就會停留在子進(jìn)程中拭宁，設(shè)置的命令為：

set follow-fork-mode child

? ? ? ? 這樣洛退，停留在遇到斷點(diǎn)的第一個子進(jìn)程中。其余不再贅述杰标。

????????二）NPU模式下運(yùn)行兵怯、調(diào)試

????????1、運(yùn)行：

bash run.sh add_tik2 ascend910 aicore npu

????????編譯及運(yùn)行結(jié)果：

????????2腔剂、調(diào)試：

????????在真實(shí)芯片上獲取profiling數(shù)據(jù)媒区，進(jìn)行性能精細(xì)調(diào)優(yōu)。

msprof --application="./add_tik2_npu" --output="./out" --ai-core=on --aic-metrics="PipeUtilization"

? ? ? ? 執(zhí)行過程如下：

?????????執(zhí)行后掸犬，對Profiling數(shù)據(jù)進(jìn)行解析與導(dǎo)出袜漩，存放在工程的下述目錄下。

?著作權(quán)歸作者所有,轉(zhuǎn)載或內(nèi)容合作請聯(lián)系作者

人面猴
序言：七十年代末湾碎，一起剝皮案震驚了整個濱河市宙攻，隨后出現(xiàn)的幾起案子，更是在濱河造成了極大的恐慌介褥，老刑警劉巖座掘，帶你破解...
沈念sama閱讀 222,807評論 6贊 518
死咒
序言：濱河連續(xù)發(fā)生了三起死亡事件，死亡現(xiàn)場離奇詭異柔滔，居然都是意外死亡溢陪，警方通過查閱死者的電腦和手機(jī)，發(fā)現(xiàn)死者居然都...
沈念sama閱讀 95,284評論 3贊 399
救了他兩次的神仙讓他今天三更去死
文/潘曉璐我一進(jìn)店門睛廊，熙熙樓的掌柜王于貴愁眉苦臉地迎上來嬉愧，“玉大人，你說我怎么就攤上這事喉前∶缓ǎ” “怎么了？”我有些...
開封第一講書人閱讀 169,589評論 0贊 363
道士緝兇錄：失蹤的賣姜人
文/不壞的土叔我叫張陵卵迂，是天一觀的道長裕便。經(jīng)常有香客問我，道長见咒，這世上最難降的妖魔是什么偿衰？我笑而不...
開封第一講書人閱讀 60,188評論 1贊 300
?港島之戀（遺憾婚禮）
正文為了忘掉前任，我火速辦了婚禮，結(jié)果婚禮上下翎，老公的妹妹穿的比我還像新娘缤言。我一直安慰自己，他們只是感情好视事，可當(dāng)我...
茶點(diǎn)故事閱讀 69,185評論 6贊 398
惡毒庶女頂嫁案：這布局不是一般人想出來的
文/花漫我一把揭開白布胆萧。她就那樣靜靜地躺著，像睡著了一般俐东。火紅的嫁衣襯著肌膚如雪跌穗。梳的紋絲不亂的頭發(fā)上，一...
開封第一講書人閱讀 52,785評論 1贊 314
城市分裂傳說
那天虏辫，我揣著相機(jī)與錄音蚌吸，去河邊找鬼。笑死砌庄，一個胖子當(dāng)著我的面吹牛羹唠，可吹牛的內(nèi)容都是我干的。我是一名探鬼主播娄昆，決...
沈念sama閱讀 41,220評論 3贊 423
雙鴛鴦連環(huán)套：你想象不到人心有多黑
文/蒼蘭香墨我猛地睜開眼肉迫，長吁一口氣：“原來是場噩夢啊……” “哼！你這毒婦竟也來了稿黄？” 一聲冷哼從身側(cè)響起喊衫，我...
開封第一講書人閱讀 40,167評論 0贊 277
萬榮殺人案實(shí)錄
序言：老撾萬榮一對情侶失蹤，失蹤者是張志新（化名）和其女友劉穎杆怕，沒想到半個月后族购，有當(dāng)?shù)厝嗽跇淞掷锇l(fā)現(xiàn)了一具尸體，經(jīng)...
沈念sama閱讀 46,698評論 1贊 320
?護(hù)林員之死
正文獨(dú)居荒郊野嶺守林人離奇死亡陵珍，尸身上長有42處帶血的膿包…… 初始之章·張勛以下內(nèi)容為張勛視角年9月15日...
茶點(diǎn)故事閱讀 38,767評論 3贊 343
?白月光啟示錄
正文我和宋清朗相戀三年寝杖，在試婚紗的時候發(fā)現(xiàn)自己被綠了。大學(xué)時的朋友給我發(fā)了我未婚夫和他白月光在一起吃飯的照片互纯。...
茶點(diǎn)故事閱讀 40,912評論 1贊 353
活死人
序言：一個原本活蹦亂跳的男人離奇死亡瑟幕，死狀恐怖，靈堂內(nèi)的尸體忽然破棺而出留潦，到底是詐尸還是另有隱情只盹，我是刑警寧澤，帶...
沈念sama閱讀 36,572評論 5贊 351
?日本核電站爆炸內(nèi)幕
正文年R本政府宣布兔院，位于F島的核電站殖卑，受9級特大地震影響，放射性物質(zhì)發(fā)生泄漏坊萝。R本人自食惡果不足惜孵稽，卻給世界環(huán)境...
茶點(diǎn)故事閱讀 42,254評論 3贊 336
男人毒藥：我在死后第九天來索命
文/蒙蒙一许起、第九天我趴在偏房一處隱蔽的房頂上張望。院中可真熱鬧菩鲜，春花似錦园细、人聲如沸。這莊子的主人今日做“春日...
開封第一講書人閱讀 32,746評論 0贊 25
一樁弒父案猛频，背后竟有這般陰謀
文/蒼蘭香墨我抬頭看了看天上的太陽。三九已至馅笙，卻和暖如春伦乔，著一層夾襖步出監(jiān)牢的瞬間厉亏，已是汗流浹背董习。一陣腳步聲響...
開封第一講書人閱讀 33,859評論 1贊 274
情欲美人皮
我被黑心中介騙來泰國打工，沒想到剛下飛機(jī)就差點(diǎn)兒被人妖公主榨干…… 1. 我叫王不留爱只，地道東北人皿淋。一個月前我還...
沈念sama閱讀 49,359評論 3贊 379
代替公主和親
正文我出身青樓，卻偏偏與公主長得像恬试，于是被迫代替她去往敵國和親窝趣。傳聞我的和親對象是個殘疾皇子，可洞房花燭夜當(dāng)晚...
茶點(diǎn)故事閱讀 45,922評論 2贊 361

【2023 · CANN訓(xùn)練營第一季】——Ascend C算子代碼分析—Add算子（內(nèi)核調(diào)用符方式）

推薦閱讀更多精彩內(nèi)容