歡迎關注我的公眾號 ［極智視界］，獲取我的更多筆記分享
O_o
>_<
o_O
O_o
~_~
o_O
本文介紹一下寒武紀 BANG C 的程式設計方式。
相信接觸過 CUDA C 的同學比較多， 但接觸過 BANG C 的同學會少很多，這主要源於 GPU 的使用人群量是 MLU 不可比擬的，不過多學習一些異構程式設計方面的東西還是好的。這裡分享一下寒武紀 BANG C 的程式設計方式，看起來和 CUDA C 會比較類似。我們知道一般的異構程式設計會分為 Host 端和 Device 端，在 Host 端做邏輯判斷，在 Device 端做平行計算。在 BANG C 程式設計中也是一樣，分為 Host 端和 MLU 端，下面從這兩端程式設計方式分別進行介紹。
1、Host 端
Host 端呼叫 CNRT 介面，CNRT 是 MLU 的推理執行時庫，需要包 cnrt。h 頭：
#include “cnrt。h”
進行裝置初始化：
cnrtInit（0）；cnrtDev_t dev；cnrtGetDeviceHandle（&dev， 0）；cnrtSetCurrentDevice（dev）；
建立 Queue：
cnrtQueue_t queue；cnrtCreateQueue（&queue）；
指定任務規模和排程型別：
cnrtDim3_t dim； dim。x = 1；dim。y = 1；dim。z = 1；cnrtFunctionType_t func_type = CNRT_FUNC_TYPE_BLOCK；
準備 Kernel 的引數：
cnrtKernelParamsBuffer_t params；cnrtKernelInitParam_t init_param；cnrtCreateKernelInitParam（&init_param）；cnrtInitKernelMemory（kernel， init_param）；。。。uint32_t param1 = SIZE_INPUT；cnrtGetKernelParamsBuffer（¶ms）；cnrtKernelParamsBufferAddParam（params， ¶m1， sizeof（uint32_t））；。。。
在 GDRAM 上為輸入資料分配記憶體：
half *ptr_mlu_input；cnrtMalloc（（void **）&ptr_mlu_input， N * sizeof（half））；
輸入資料拷入 GDRAM：
cnrtMemcpy（&ptr_mlu_input prt_host_input， CNRT_MEM_TRANS_DIR_HOST2DEV）；
啟動 Kernel，並繫結到 Queue：
cnrtInvokeKernel_V3（（void *）&kernel， init_param， dim， params， func_type， queue， NULL）；
同步 Queue：
cnrtSyncQueue（queue）；
資料拷出：
cnrtMemcpy（ptr_host_output， ptr_mlu_output， SIZE_OUTPUT， CNRT_MEM_TRANS_DIR_DEV2HOST）；
資源釋放：
cnrtDestroyQueue（queue）；cnrtDestoryKernelParamsBuffer（params）；cnrtDestroy（）；
2、MLU 端
MLU 端的程式碼就是基於 BANG C 來寫的了，類似 CUDA C，是基於 C 語言的擴充套件，其檔名字尾為 。mlu。
MLU 端程式需要包含 mlu。h 頭：
#inlcude “mlu。h”
再包含一些其他需要的頭：
#include “macro。h”#define s1 N1#define s2 N2
BANG C Kernel 的入口函式需要用
__mlu_entry__
標記，類似於 CUDA C 中的
__global__
：
__mlu__entry__ void kernel（uint32_t size1， uint32_t size2， half* c， half* a， half* b）{ /// __nram__ 是標記 NRAM 空間的變數 __nram__ half a_tmp［s1］； __nram__ half b_tmp［s2］； __nram__ half c_tmp［s3］； /// 將資料從 GRAM 複製到 NRAM 上 __memcpy（a_tmp， a， s1 * sizeof（half）， GDRAM2NRAM）； __memcpy（b_tmp， b， s2 * sizeof（half）， GDRAM2NRAM）； /// 將 NRAM 上資料進行迴圈 add 計算 __bang_cycle_add（c_tmp， b_tmp， a_tmp， size2， size1）； /// 將計算結果從 NRAM 複製到 GDRAM，HOST端可以從 RDRAM 拿資料 __memcpy（c， c_tmp， s2 * sizeof（half）， NRAM2GDRAM）；}
以上分享了一下 BANG C 異構程式設計的方式，關於什麼是 NRAM、GDRAM 這些，下次我再寫一篇關於寒武紀架構的文章說明一下。
好了，收工~
【公眾號傳送】
《【程式設計藝術】寒武紀 BANG C 異構程式設計方式》
