智能計算系統(tǒng)：從深度學(xué)習(xí)到大模型 第2版課件 8、第八章-智能編程語言

智能計算系統(tǒng)

第八章智能編程語言中國科學(xué)院計算技術(shù)研究所李威副研究員liwei2017@本章內(nèi)容定位2編程框架智能編程語言

輸入輸出建模實(shí)現(xiàn)運(yùn)行本章將學(xué)習(xí)到智能編程語言的基礎(chǔ)知識及相應(yīng)的應(yīng)用開發(fā)、調(diào)試和調(diào)優(yōu)方法等學(xué)習(xí)了實(shí)現(xiàn)深度學(xué)習(xí)算法所使用的編程框架的基本原理、簡單用法等提綱為什么需要智能編程語言智能計算系統(tǒng)抽象架構(gòu)智能編程模型智能編程語言基礎(chǔ)智能應(yīng)用編程接口智能編程語言實(shí)例智能應(yīng)用功能調(diào)試智能應(yīng)用性能調(diào)優(yōu)智能編程語言的應(yīng)用3為什么需要智能編程語言4編程語言智能應(yīng)用A…..智能應(yīng)用B智能應(yīng)用N智能硬件A智能硬件B智能硬件N語義鴻溝硬件鴻溝平臺鴻溝…..執(zhí)行效率低開發(fā)效率低可移植性差語義鴻溝以深度學(xué)習(xí)中最為核心的卷積運(yùn)算為例5C++：標(biāo)量計算Python：向量語義（Array）7重循環(huán)4重循環(huán)有Conv語義和Tensor類型的編程語言6一條語句完成計算Tensor類型硬件鴻溝智能計算硬件在控制、存儲和計算等方面有獨(dú)特性傳統(tǒng)編程語言難以有效描述上述硬件特點(diǎn)不同層次編程語言和硬件特性帶來的性能影響71、C語言實(shí)現(xiàn)相較Python/JAVA實(shí)現(xiàn)有47倍的性能提升2、考慮了底層硬件提供的并行度、存儲層次以及向量指令后，相較Python/JAVA有62,806倍的性能提升如何在語言層面填補(bǔ)硬件鴻溝存儲邏輯上一般采用程序員可見的ScratchpadMemory（SPM），而不是通用平臺上程序員透明的Cache計算邏輯上提供了面向智能計算的定制運(yùn)算單元，如16位浮點(diǎn)、Brain浮點(diǎn)等，其精度損失幾乎可以忽略8傳統(tǒng)編程語言中主要提供的是INT和FP32等數(shù)據(jù)類型，導(dǎo)致難以利用智能計算系統(tǒng)中更加豐富和高效的運(yùn)算單元，如FP16和BF16等，甚至是INT4及Binary的數(shù)據(jù)類型。平臺鴻溝9功能可移植性：采用特定平臺專用語言所編寫的程序能夠在別的平臺上正常運(yùn)行矩陣乘法的例子調(diào)用了AVX的intrinsic函數(shù)在ARM上無法運(yùn)行性能可移植性：在特定平臺上優(yōu)化好的程序，在新的硬件平臺上仍然保證有較高的執(zhí)行效率理想的編程語言需要抽取不同硬件平臺的共性特征，并在此基礎(chǔ)上提取性能關(guān)鍵特征作為語言特性提供給用戶小結(jié)10面向語義、硬件和平臺三大鴻溝，傳統(tǒng)編程語言難以滿足需求領(lǐng)域?qū)Ｓ镁幊陶Z言是滿足智能計算高開發(fā)效率、高性能和高可移植性的重要途徑內(nèi)容組織11硬件抽象架構(gòu)智能編程模型語言基礎(chǔ)

編程接口功能調(diào)試性能調(diào)優(yōu)編程框架深度學(xué)習(xí)處理器集成智能應(yīng)用運(yùn)行映射定義介紹智能編程語言原理基礎(chǔ)上進(jìn)一步介紹BCL語言的實(shí)例—MLU上的BANG語言系統(tǒng)開發(fā)提綱為什么需要智能編程語言智能計算系統(tǒng)抽象架構(gòu)智能編程模型智能編程語言基礎(chǔ)智能應(yīng)用編程接口智能編程語言實(shí)例智能應(yīng)用功能調(diào)試智能應(yīng)用性能調(diào)優(yōu)智能編程語言的應(yīng)用12智能計算系統(tǒng)抽象架構(gòu)抽象硬件架構(gòu)典型智能計算系統(tǒng)控制模型計算模型定制運(yùn)算單元并行計算架構(gòu)存儲模型13抽象硬件架構(gòu)層次化的智能計算系統(tǒng)抽象硬件架構(gòu)14智能計算系統(tǒng)中每一層都包含存儲單元、控制單元和若干個計算單元每個計算單元又進(jìn)一步分解為子控制單元、子計算單元和子存儲單元三部分，整個系統(tǒng)以這樣的方式遞歸構(gòu)成在最底層，每個葉節(jié)點(diǎn)都是具體的加速器，用于完成最基本的計算任務(wù)。典型智能計算系統(tǒng)多卡的DLP服務(wù)器抽象為五個層次，即服務(wù)器級（Server）、板卡級（Card）、芯片級（Chip）、處理器簇級（Cluster）和處理器核級（Core）。15可以方便地通過增加各層次的規(guī)模來提升整個系統(tǒng)算力控制模型指令是實(shí)現(xiàn)對計算和存儲進(jìn)行控制的關(guān)鍵。為了設(shè)計高效的指令集、需要充分分析智能領(lǐng)域的典型計算模式，提煉最具代表性的操作，并進(jìn)行針對性設(shè)計。對智能算法進(jìn)行抽象控制數(shù)據(jù)傳輸計算：標(biāo)量、向量和矩陣運(yùn)算等邏輯操作：標(biāo)量和向量運(yùn)算等關(guān)注計算與存儲的交互：盡可能將計算與存儲并行，例如可以將控制計算和訪存的指令分開在不同的隊(duì)列中發(fā)射執(zhí)行，以提高并行度16計算模型程序員可見的主要包括定制運(yùn)算單元和并行計算架構(gòu)定制運(yùn)算單元智能應(yīng)用具有一定誤差容忍度。通過利用智能應(yīng)用誤差容忍的特性，一般在智能計算系統(tǒng)中會采用定制的低位寬運(yùn)算單元（如FP16、INT8、BF16甚至是INT4等）以提升處理能效。由于智能應(yīng)用的多樣性和復(fù)雜性，目前對于哪種低位寬最為合適并未形成統(tǒng)一結(jié)論。例如推斷和訓(xùn)練對于精度的要求可能不一樣，圖像/視頻類應(yīng)用和語音類應(yīng)用對于精度要求可能也不一樣。17智能編程語言中需要有和各種類型運(yùn)算單元對應(yīng)的數(shù)據(jù)類型并行計算架構(gòu)任務(wù)切分與同步18存儲模型智能應(yīng)用中存在大量數(shù)據(jù)密集的內(nèi)存訪問，因此合理地組織存儲層次和計算單元同樣重要，需要兩者協(xié)同設(shè)計以平衡計算與訪存，實(shí)現(xiàn)高效的智能計算分為全局存儲和本地存儲19片外存儲DDR/HBM片內(nèi)存儲SPML1CacheL2Cache本地存儲全局存儲NRAMWRAM提綱為什么需要智能編程語言智能計算系統(tǒng)抽象架構(gòu)智能編程模型智能編程語言基礎(chǔ)智能應(yīng)用編程接口智能編程語言實(shí)例智能應(yīng)用功能調(diào)試智能應(yīng)用性能調(diào)優(yōu)智能編程語言的應(yīng)用20智能編程模型異構(gòu)編程模型分類及流程編譯器支持運(yùn)行時支持通用智能編程模型Kernel定義編譯器支持運(yùn)行時支持21異構(gòu)編程模型異構(gòu)計算系統(tǒng)組成通用處理器：控制設(shè)備（簡稱主機(jī)端），負(fù)責(zé)控制和調(diào)度等工作領(lǐng)域處理器：從設(shè)備（簡稱設(shè)備端），負(fù)責(zé)大規(guī)模的并行計算或領(lǐng)域?qū)Ｓ糜嬎闳蝿?wù)二者協(xié)同完成完整計算任務(wù)22典型異構(gòu)計算系統(tǒng)（a）GPU為核心（b）FPGA為核心（c）TPU為核心（d）DLP為核心PCIeSwitch/NVSwitchPCIeSwitch/NVSwitch異構(gòu)編程模型：分類異構(gòu)并行編程模型從用戶接口角度大致可分為兩類：構(gòu)建全新的異構(gòu)并行編程對現(xiàn)有編程語言進(jìn)行異構(gòu)并行擴(kuò)展23典型異構(gòu)并行編程模型對比類別主要編程考量OpenCL語言擴(kuò)展任務(wù)劃分+數(shù)據(jù)分布、通信、同步CUDA語言擴(kuò)展任務(wù)劃分+數(shù)據(jù)分布、通信、同步Copperhead新語言任務(wù)劃分為主Merge新語言任務(wù)劃分為主異構(gòu)編程模型：流程異構(gòu)編程模型的編譯和鏈接流程整體采用分離式編程方式：主機(jī)端代碼和設(shè)備端代碼24異構(gòu)編程模型：編譯器支持編譯器支持是異構(gòu)并行編程模型的核心25異構(gòu)并行編程編譯器任務(wù)劃分?jǐn)?shù)據(jù)分布數(shù)據(jù)通信并行同步編程模型需要向程序員提供并行編程接口，方便程序員定義和劃分任務(wù)。編譯器負(fù)責(zé)底層的任務(wù)劃分，使得程序可以在并行架構(gòu)上高效執(zhí)行。對于編譯器和底層運(yùn)行時系統(tǒng)而言，需要根據(jù)算法和硬件架構(gòu)的特點(diǎn)，通過合適的數(shù)據(jù)分布指導(dǎo)后續(xù)編譯和運(yùn)行時優(yōu)化。由于設(shè)備端通常有多級存儲空間、編譯器需要支持各種地址空間聲明，以方便程序員顯式控制存儲數(shù)據(jù)的地址空間。設(shè)備端程序一般要求感知多個核的并行處理，因此需要提供對同步機(jī)制的支持。異構(gòu)編程模型：運(yùn)行時支持完成任務(wù)映射及調(diào)度，即指定任務(wù)具體在哪個設(shè)備或計算單元上以何種順序執(zhí)行分為主機(jī)端和設(shè)備端

主機(jī)端：控制部分和串行任務(wù)在主機(jī)端執(zhí)行

設(shè)備端：計算部分和并行任務(wù)在設(shè)備端執(zhí)行26通用智能編程模型通用智能編程模型以前述異構(gòu)編程模型為基礎(chǔ)Kernel函數(shù)定義定義在設(shè)備端DLP上的核心計算任務(wù)編譯器支持指定DLP上的核心計算任務(wù)如何高效地翻譯成目標(biāo)代碼運(yùn)行時支持指定DLP上的核心計算任務(wù)以何種方式映射到計算單元27通用智能編程模型：Kernel定義與異構(gòu)編程模型中的概念一致，DLP上執(zhí)行的任務(wù)叫Kernel，資源允許情況下DLP可以同時執(zhí)行多個并行的Kernel設(shè)備端，每個Kernel有一個入口函數(shù)，BCL中用__dlp_entry__來指定設(shè)備端，程序默認(rèn)的函數(shù)類型：Device函數(shù)，以__dlp_device__來修飾主機(jī)端，Kernel的啟動使用InvokeKernel接口或語法糖foo<<<…>>>()或28ret=InvokeKernel((constvoid*)MatrixMul,dim,ktype,args,0,queue);__dlp_entry__voidMatrixMul(half*dA,half*dB,half*dC,…){…}MatrixMul<<<dim,ktype,queue>>>(dA,dB,dC,…);__dlp_device__voidVectorMul(half*dA,half*dB,half*dC,…){…}通用智能編程模型：編譯器支持任務(wù)劃分并行內(nèi)建變量硬件：clusterDim（維度），clusterId（序號），coreDim，coreId任務(wù)：taskDim[XYZ]，taskId[XYZ]表示Kernel啟動task的規(guī)模，有XYZ三個維度，用戶根據(jù)需求進(jìn)行指定任務(wù)調(diào)度類型表示Kernel運(yùn)行調(diào)度時需要的硬件核BLOCK類型：Kernel為單核任務(wù)，按單核進(jìn)行調(diào)度UNIONx類型：Kernel為多核并行任務(wù)（其中x可以為1/2/4，UNION1對應(yīng)1個cluster4個核）29taskDimZtaskDimYtaskDimXtaskID[XYZ]數(shù)據(jù)通信：為DLP的復(fù)雜存儲層次提供支持隱式數(shù)據(jù)管理：GPR標(biāo)量數(shù)據(jù)，由編譯器隱式插入Load/Store指令顯式數(shù)據(jù)管理：DRAM/NRAM/WRAM/SRAM間向量及張量數(shù)據(jù)主機(jī)-DLP間DRAM數(shù)據(jù)30顯式管理HostDRAM顯式管理隱式管理同步支持：為并行計算架構(gòu)提供支持

抽象硬件架構(gòu)中有Chip-Cluster-Core的層次結(jié)構(gòu)，Core內(nèi)還支持指令隊(duì)列之間的并行，因此智能編程語言需要提供至少三種不同類型的同步操作：__sync:同步一個Core內(nèi)所有的指令隊(duì)列，只有所有指令隊(duì)列的指令都執(zhí)行完畢才能繼續(xù)執(zhí)行后續(xù)的指令。__sync_cluster：同步一個Cluster內(nèi)部的所有核，當(dāng)一個Cluster內(nèi)所有核都達(dá)到同步點(diǎn)時才繼續(xù)往下執(zhí)行。與CUDA相比，GPU同一個線程塊內(nèi)的thread可以同步，而線程塊間的thread無法同步。__sync_all：同步任務(wù)執(zhí)行的所有核，只有所有核到達(dá)同步點(diǎn)時才繼續(xù)往下執(zhí)行。參與同步的核數(shù)由任務(wù)的調(diào)度類型確定，例如當(dāng)任務(wù)類型為UNION1時，只有4個核參與同步（假定一個Cluster內(nèi)包含4個核），當(dāng)任務(wù)類型為UNION2時，參與同步的核數(shù)為8。31內(nèi)建運(yùn)算：為用戶編程提供支持，提高開發(fā)效率通用智能編程語言提供并實(shí)現(xiàn)了__matmul和__mlp等內(nèi)建函數(shù)接口，分別對應(yīng)卷積和全連接等典型神經(jīng)網(wǎng)絡(luò)運(yùn)算上述接口是對C/C++語言的擴(kuò)展，深度學(xué)習(xí)處理器端Kernel程序編寫時可以調(diào)用這些接口，通過編譯器將這些接口翻譯為底層硬件指令通用智能編程模型直接實(shí)現(xiàn)了神經(jīng)網(wǎng)絡(luò)計算的內(nèi)建接口，能更好地支持智能應(yīng)用。

32通用智能編程模型：運(yùn)行時支持任務(wù)調(diào)度單位（BLOCK/UNIONx）以調(diào)度單位將Kernel中的任務(wù)在時間或空間維度展開BLOCK：單核調(diào)度，當(dāng)有一個核空閑時，調(diào)度一個任務(wù)執(zhí)行UNION1：調(diào)度時需要1個cluster，當(dāng)有1個cluster空閑時，調(diào)度任務(wù)執(zhí)行UNION2：調(diào)度時需要2個cluster，當(dāng)有2個cluster空閑時，調(diào)度任務(wù)執(zhí)行調(diào)度單位需要用戶在編程時指定。運(yùn)行時只有當(dāng)空閑的硬件資源數(shù)大于調(diào)度單位時，Kernel才會被調(diào)度隊(duì)列（Queue）管理需要執(zhí)行的任務(wù)，隊(duì)列既可以單獨(dú)工作，也可以協(xié)同工作運(yùn)行時（或硬件）不斷把任務(wù)放到隊(duì)列中，一旦硬件計算資源有空閑，就從隊(duì)列中取出一個任務(wù)執(zhí)行331、主機(jī)端異步發(fā)射3個Kernel到Queue中。用戶可以根據(jù)同步和通信的需要，在三次發(fā)射之間或之后任意位置調(diào)用同步接口SyncQueue。假設(shè)在三次發(fā)射之后調(diào)用，則等待Queue中的任務(wù)全部完成后再繼續(xù)執(zhí)行主機(jī)端SyncQueue后面的程序；2、第一個任務(wù)Kernel1在Time1被發(fā)射后立即進(jìn)入Queue，設(shè)備端發(fā)現(xiàn)當(dāng)前全部核心空閑則立即執(zhí)行Kernel1。Kernel1的任務(wù)類型為UNION2，會從Time1開始占用2個Cluster執(zhí)行計算；34主機(jī)端執(zhí)行3個Kernel任務(wù)設(shè)備端有4個Cluster，16個Core353、因沒有調(diào)用SyncQueue，所以主機(jī)端發(fā)射Kernel1后立即發(fā)射Kernel2，設(shè)備端調(diào)度器在調(diào)度執(zhí)行Kernel1后發(fā)現(xiàn)隊(duì)列中有了新的Kernel2，也幾乎在Time1時刻開始執(zhí)行Kernel2。Kernel2的任務(wù)類型為UNION1，會從Time1開始占用1個Cluster計算；4、Kernel1和Kernel2幾乎同時被調(diào)度器執(zhí)行且假設(shè)同時在Time2時刻結(jié)束；5、當(dāng)Kernel2被發(fā)射后，因?yàn)闆]有執(zhí)行同步，所以Kernel3也會立即被發(fā)射，此時刻也幾乎為Time1；6、Kernel3的任務(wù)類型是UNION4，需要4個Cluster，但在Time1時刻到Time2時刻硬件的4個Cluster被占用了3個（假設(shè)只有4個Cluster），那么設(shè)備端調(diào)度器會一直等待Time2時刻有4個Cluster的空閑時才執(zhí)行Kernel3。7、假設(shè)Kernel1和Kernel2的任務(wù)并行總規(guī)模taskDim超過了任務(wù)類型表示的核數(shù)（例如kernel1的taskDim=16，而調(diào)度類型是UNION2，即一次要占用8個Core），則調(diào)度器會將同一份Kernel程序在時間序上執(zhí)行多次；36提綱為什么需要智能編程語言智能計算系統(tǒng)抽象架構(gòu)智能編程模型智能編程語言基礎(chǔ)智能應(yīng)用編程接口智能編程語言實(shí)例智能應(yīng)用功能調(diào)試智能應(yīng)用性能調(diào)優(yōu)智能編程語言的應(yīng)用37智能編程語言基礎(chǔ)語法概述數(shù)據(jù)類型宏、常量與內(nèi)置變量I/O操作語句標(biāo)量計算語句張量計算語句控制流語句串行程序示例并行程序示例38語法概述智能編程語言考慮基于過程式語言當(dāng)前大多數(shù)語言都是過程式的，可以減少用戶學(xué)習(xí)成本當(dāng)前主流智能算法可以描述為明確的過程，適合采用過程式語言描述參考經(jīng)典的過程式語言C/C++，我們所定義的智能編程語言同樣具有數(shù)據(jù)和函數(shù)兩個基本要素例如：定義add_func函數(shù)的函數(shù)體和C語言一致39基本數(shù)據(jù)類型40基本數(shù)據(jù)類型長度說明int8_t1byte1字節(jié)整數(shù)uint8_t1byte1字節(jié)無符號整數(shù)int16_t2byte2字節(jié)整數(shù)uint16_t2byte2字節(jié)無符號整數(shù)int32_t4byte4字節(jié)整數(shù)uint32_t4byte4字節(jié)無符號整數(shù)half2byte半精度浮點(diǎn)數(shù)據(jù)類型，采用IEEE-754fp16格式float4byteIEEE-754fp32格式浮點(diǎn)類型,目前僅支持類型轉(zhuǎn)換計算char1byte對應(yīng)C語言char類型bool1byte對應(yīng)C語言bool類型指針8byte指針類型宏、常量與內(nèi)置變量宏/常量由用戶定義，內(nèi)置變量是語言既有的宏和常量，宏不僅可以定義常量數(shù)據(jù)，也可定義一段代碼；常量是不可修改的數(shù)據(jù)，只能在初始化時被賦值。內(nèi)置變量，編程語言本身包含的常量和變量，不需用戶定義即可直接使用。41內(nèi)置變量名具體含義coreIdDLP中核的編號clusterIdDLP中簇的編號taskId程序運(yùn)行時分配的任務(wù)編號I/O操作語句42典型的NUMA（Non-UniformMemoryAccessArchitecture）架構(gòu)不同處理器核對不同位置存儲器的訪問速度不同。對于核0而言，其訪問設(shè)備內(nèi)存DDR0的速度比訪問DDR1的速度更快。DDR0看作是本地存儲，DDR1作為全局存儲片上存儲，除了單核內(nèi)NRAM和權(quán)值WRAM，還有一類共享存儲，可用于簇內(nèi)的多核共享不同層次的智能處理節(jié)點(diǎn)有各自的本地存儲，需要提供不同存儲層次間的數(shù)據(jù)搬移針對上述典型架構(gòu)（三種片上存儲、兩種設(shè)備內(nèi)存），可以有多種不同的數(shù)據(jù)搬移操作類型針對上述搬移操作類型，可以在智能編程語言中定義相應(yīng)的內(nèi)建函數(shù)__memcpy，方便用戶進(jìn)行不同類型的數(shù)據(jù)搬移43標(biāo)量計算語句標(biāo)量即單個數(shù)據(jù)的計算，標(biāo)量計算是編程語言的基本功能智能編程語言的標(biāo)量計算語句有兩種形式：運(yùn)算符號（如+，-，*，/等）內(nèi)建函數(shù)（如abs，max，min等）智能編程語言的標(biāo)量計算語句由編譯器映射到標(biāo)量計算單元，雖然吞吐上不及張量運(yùn)算，但具有良好的通用性和靈活性44張量計算語句張量計算是智能編程語言的主要特點(diǎn)，可以通過內(nèi)建函數(shù)直接映射到張量計算單元張量計算直接對精度類型和語義類型等數(shù)據(jù)直接進(jìn)行操作45一維張量計算語句具體功能__vector_add(DType*dst,constDType*lhs,constDType*rhs,intelem_num)向量對位加__vector_sub(DType*dst,constDType*lhs,constDType*rhs,intelem_num)向量對位減__vector_mul(DType*dst,constDType*lhs,constDType*rhs,intelem_num)向量對位乘__vector_relu(DType*dst,constDType*src,intelem_num)向量Relu__vector_argmax(DType*dst,constDType*src,intelem_num)向量ArgMax控制流語句與通用編程語言一樣，智能編程語言同樣需要有分支和循環(huán)等控制流語句分支語句：用于處理程序的選擇邏輯，由判斷條件與分支代碼段組成，與傳統(tǒng)編程語言類似循環(huán)語句：用于處理程序循環(huán)邏輯，由循環(huán)執(zhí)行條件和循環(huán)代碼段組成，與傳統(tǒng)編程語言類似同步語句：解決多核間并行數(shù)據(jù)依賴問題，保證最終計算結(jié)果正確。主要分為三類：核內(nèi)同步（__sync）、Cluster內(nèi)同步（__sync_cluster）與全局同步（__sync_all）。46串行程序示例向量中每個數(shù)求平方，每次處理64個數(shù)47并行程序示例矩陣乘法示例，在4個核上并行執(zhí)行，每個核上代碼一致計算1*32和32*32的矩陣乘法，最終得到4*32的結(jié)果通過運(yùn)行時API來指定任務(wù)規(guī)模（4*1*1）及調(diào)度方式（UNION1）48設(shè)備端主機(jī)端…

……

…任務(wù)規(guī)模調(diào)度類型通過core0把輸出內(nèi)存區(qū)域?qū)?提綱為什么需要智能編程語言智能計算系統(tǒng)抽象架構(gòu)智能編程模型智能編程語言基礎(chǔ)智能應(yīng)用編程接口智能編程語言實(shí)例智能應(yīng)用功能調(diào)試智能應(yīng)用性能調(diào)優(yōu)智能編程語言的應(yīng)用49智能應(yīng)用編程接口Kernel函數(shù)接口運(yùn)行時接口使用示例50Kernel函數(shù)接口為了充分利用并行資源，需要在Kernel內(nèi)部對任務(wù)進(jìn)行有效切分，同時在主機(jī)端配置和調(diào)用相應(yīng)的Kernel函數(shù)接口任務(wù)切分的內(nèi)置變量51變量說明coreDim（核維數(shù)）內(nèi)置變量，等于單個Cluster內(nèi)部的計算核個數(shù)。coreId（核序號）內(nèi)置變量，執(zhí)行核函數(shù)的核心的邏輯編號，取值范圍為[0,coreDim-1]。clusterDim（簇維數(shù)）內(nèi)置變量，核函數(shù)運(yùn)行時所需的邏輯核心簇的數(shù)量clusterId（簇序號）內(nèi)置變量，執(zhí)行核函數(shù)運(yùn)行的核心簇的邏輯編號，取值范圍為[0,clusterDim-1]。taskDim（任務(wù)維數(shù)）內(nèi)置變量，核函數(shù)運(yùn)行時分配的任務(wù)規(guī)模，taskDim=taskDimX×taskDimY×taskDimZ。taskId（任務(wù)序號）內(nèi)置變量，核函數(shù)運(yùn)行時分配的任務(wù)編號，取值范圍為[0,taskDim-1]。taskId的值對應(yīng)邏輯規(guī)模降維后的任務(wù)ID，即taskId=taskIdZ×taskDimY×taskDimX+taskIdY×taskDimX+taskIdX。主機(jī)端Kernel函數(shù)接口用于將智能編程語言編寫的程序加載到深度學(xué)習(xí)處理器上執(zhí)行；與Kernel函數(shù)相關(guān)的接口主要關(guān)注Kernel參數(shù)設(shè)置和Kernel調(diào)用InvokeKernel(constvoid*kernel,dlpDim3_tdim,FunctionType_tktype,void**args,size_treserved,Queue_tqueue);通過在設(shè)備上給定的參數(shù)塊，調(diào)用Kernel，成功返回RET_SUCCESS，否則返回相應(yīng)的錯誤碼。其中FunctionType_t包含BLOCK和UNIONx（如UNION1和UNION2，取決于系統(tǒng)的規(guī)模）等類型，分別表示單核任務(wù)和多核并行任務(wù)。在主機(jī)端異步執(zhí)行。52主機(jī)端Kernel函數(shù)接口DLP還可以通過編譯器和運(yùn)行時的配合，實(shí)現(xiàn)foo<<<>>>()語法糖53運(yùn)行時接口包括設(shè)備管理、隊(duì)列管理和內(nèi)存管理等接口設(shè)備管理：主要涉及設(shè)備獲取和設(shè)置、屬性獲取操作Init和Destroy可由運(yùn)行時庫隱式自動完成，無需用戶感知GetDevice(int*pOrdinal);獲取主機(jī)端當(dāng)前線程上下文所使用的設(shè)備序號SetDevice(intordinal);為當(dāng)前主機(jī)端線程設(shè)置所使用的設(shè)備序號DeviceGetAttribute(int*pValue,DeviceAttr_tattr,intdevice);獲取設(shè)備屬性54隊(duì)列管理隊(duì)列是用于執(zhí)行任務(wù)的環(huán)境。計算任務(wù)可以下發(fā)到隊(duì)列中執(zhí)行。同一個隊(duì)列可以容納多個任務(wù)。具體來說，隊(duì)列具有以下屬性：串行性：下發(fā)到同一個隊(duì)列中的任務(wù)，按下發(fā)順序串行執(zhí)行異步性：任務(wù)下發(fā)到隊(duì)列是異步過程，即下發(fā)完成后程序控制流回到主機(jī)，主機(jī)程序繼續(xù)往下執(zhí)行。運(yùn)行時環(huán)境提供隊(duì)列的同步接口SyncQueue用于等待隊(duì)列中所有任務(wù)完成。并行性：不同隊(duì)列中的任務(wù)并行執(zhí)行。如果希望任務(wù)間并行執(zhí)行，用戶可以創(chuàng)建多個隊(duì)列并將任務(wù)分配到不同的隊(duì)列中。QueueCreate(Queue_t*pQueue);QueueSync(Queue_tqueue);QueueDestroy(Queue_tqueue);55內(nèi)存管理內(nèi)存管理主要分為主機(jī)端內(nèi)存管理、設(shè)備端內(nèi)存管理和主機(jī)與設(shè)備端內(nèi)存拷貝三類：hostMalloc(void**ptr,size_tbytes,…)hostFree(void*ptr)devMalloc(void**ptr,size_tbytes)devFree(void*ptr)Memcpy(void*dst,void*src,size_tbytes,MemTransDir_tdir)56使用實(shí)例首先完成Device端的Kernel函數(shù)編寫57注意：對每個智能編程語言編寫的程序，有且僅有一個標(biāo)記為__dlp_entry__的核函數(shù)，表示整個Kernel函數(shù)的入口，其返回值類型必須是voidHost端代碼①設(shè)備隱式初始化和獲取設(shè)備屬性58設(shè)備的初始化和銷毀由運(yùn)行時庫隱式自動完成使用GetDevice可以獲取當(dāng)前的設(shè)備序號使用DeviceGetAttribute可以獲取當(dāng)前設(shè)備的屬性Host端代碼②主機(jī)/設(shè)備端數(shù)據(jù)準(zhǔn)備

③設(shè)備端內(nèi)存空間分配

④數(shù)據(jù)至設(shè)備端拷貝

59Host端代碼⑤調(diào)用Kernel啟動設(shè)備

⑥運(yùn)行結(jié)果獲取

⑦資源釋放60智能編程模型實(shí)例：BANG異構(gòu)編程BANG語言是針對MLU(MachineLearningUnit)硬件提出的編程語言，兼顧云邊端等不同目標(biāo)平臺，提供高性能機(jī)器學(xué)習(xí)計算支持提供通用的異構(gòu)編程模型，方便用戶擴(kuò)展自己的應(yīng)用程序提供高效的編程接口，充分發(fā)揮底層硬件特性基于C/C++語言的擴(kuò)展，簡單易用61CNNLMLUBANGMLU上的異構(gòu)編程語言，可用于定制開發(fā)編程框架/機(jī)器學(xué)習(xí)庫所需的算子MLU上的機(jī)器學(xué)習(xí)庫，用于智能芯片上加速各種機(jī)器學(xué)習(xí)算法BANG異構(gòu)編程：流程BANG異構(gòu)程序同樣采用分離式編程，即主機(jī)端與設(shè)備端程序分開編程并分別編譯，最后鏈接成一個可執(zhí)行程序主機(jī)端為C/C++程序，通常需調(diào)用CNRT運(yùn)行時接口完成以下步驟準(zhǔn)備輸入數(shù)據(jù)拷貝輸入數(shù)據(jù)到MLU準(zhǔn)備Kernel參數(shù)創(chuàng)建Queue指定Kernel任務(wù)規(guī)模以及調(diào)度類型啟動KernelMLU到主機(jī)的輸出數(shù)據(jù)拷貝資源的釋放設(shè)備端使用BANG語言特定的語法規(guī)則和接口進(jìn)行編程62常用CNRT接口名稱功能描述cnrtGetDevice獲取設(shè)備cnrtQueueCreate創(chuàng)建一個新的Queue，默認(rèn)異步運(yùn)行cnrtQueueDestroy銷毀QueuecnrtQueueSync直到之前Queue中所有的Function都完成，阻塞其他FunctioncnrtInvokeKernel通過在MLU上給定的參數(shù)塊，啟動Kernelfoo<<<>>>()編譯器和運(yùn)行時配合完成的語法糖，可替代cnrtInvokeKernelcnrtMalloc分配給定空間的設(shè)備內(nèi)存cnrtFree釋放指針指向的空間cnrtMemcpy從源地址拷貝指定字節(jié)數(shù)據(jù)到目的地址BANG異構(gòu)編程示例：主機(jī)端31頭文件bang.h包含主機(jī)和設(shè)備端的接口聲明創(chuàng)建異步Queue配置Kernel的并行規(guī)模申請設(shè)備端內(nèi)存配置Kernel的并行任務(wù)類型同步拷貝輸入數(shù)據(jù)至設(shè)備內(nèi)存異步啟動Kernel函數(shù)同步Queue等待Kernel計算完成同步拷貝輸出到主機(jī)內(nèi)存釋放設(shè)備內(nèi)存和QueueBANG異構(gòu)編程示例：設(shè)備端64設(shè)備端源碼和主機(jī)端源碼可以放在同一個源碼文件中，文件名后綴為mluBANG程序kernel的入口函數(shù)需要用__mlu_entry__或__mlu_global__標(biāo)記__nram__標(biāo)記nram空間的變量將數(shù)據(jù)從GDRAM拷貝到NRAM上將NRAM上數(shù)據(jù)調(diào)用BANGC的transpose、mul和reduce_sum等接口計算將計算的結(jié)果從NRAM拷貝到GDRAM，以供主機(jī)端程序讀取調(diào)用設(shè)備端的計時函數(shù)統(tǒng)計硬件時間并打印耗時BANG程序編譯與鏈接65BANG的源碼可以寫在同一份源碼文件中，編譯器自動完成主機(jī)端和設(shè)備端的編譯和鏈接foo.mlu混合編程源碼被cncc前端driver拆分成主機(jī)和設(shè)備源碼主機(jī)端源碼使用clang進(jìn)行編譯，得到主機(jī)端的二進(jìn)制對象文件設(shè)備端源碼根據(jù)指定的單一arch或多arch，分別編譯成目標(biāo)arch的對象文件設(shè)備端多arch的bin對象被設(shè)備端鏈接器鏈接成fatbin，并包裝成主機(jī)端可識別的對象文件主機(jī)端鏈接器將主機(jī)和設(shè)備端對象文件鏈接為可執(zhí)行文件張量計算語句實(shí)例：BANG數(shù)學(xué)庫BANG語言將MLU的數(shù)學(xué)類和神經(jīng)網(wǎng)絡(luò)類指令封裝成庫函數(shù)BANG數(shù)學(xué)庫函數(shù)是在MLU架構(gòu)上進(jìn)行高性能編程的關(guān)鍵BANG數(shù)學(xué)庫和其他設(shè)備端接口見頭文件/usr/local/neuware/lib/clang/x.x.x/include/bang_device_functions_decls.h/usr/local/neuware/lib/clang/x.x.x/include/bang_math_decls.h使用約束：張量計算通常是對批量數(shù)據(jù)進(jìn)行操作源和目的都是NRAM上的數(shù)據(jù)66BANG數(shù)學(xué)庫實(shí)例接口：__bang_add(half*dst,half*src0,half*src1,int32_tNumOfEle)語義：兩個向量相加，得到新的向量

約束：源和目的都是half類型接口：__bang_mul_scalar(half*dst,half*src0,halfscale,int32_tNumOfEle)語義：對向量的每一個元素乘以一個常數(shù)52BANG數(shù)學(xué)庫實(shí)例53__bang_select(dst,src0,src1,NumOfEle);…//getthenumberofselectedelementsfromthefirstlineuint16_tnumberOfSelectedNumbers=*((uint16_t*)dst);//dst[0]//gettheselectedelementsfromthesecondlineselectedNum[0]=dst[16];selectedNum[1]=dst[17];接口：__bang_eq(half*dst,half*src0,half*src1,int32_tNumOfEle)語義：對src0和src1兩個張量進(jìn)行element-wise的eq操作，并將結(jié)果存入dst中接口：__bang_select(half*dst,half*src0,halfsrc1,int32_tNumOfEle)語義：如果src1中某個元素非0，則從src0中選擇對應(yīng)位置元素存入dst中。dst中的結(jié)果有兩部分，第一部分是選出的元素個數(shù)，第二部分是選出的元素集合示例：BANG數(shù)學(xué)庫實(shí)例接口：__bang_max(half*dst,half*src,int32_tNumOfEle)語義：從張量src中選擇最大值的元素，存入dst中，結(jié)果有兩部分，第一部分是選出的元素，第二部分是選擇元素在src中index示例:接口：__bang_count(uint16_t*dst,half*src,int32_tNumOfEle)語義：計算src中非0元素的個數(shù)，結(jié)果存入dst中示例：54__bang_max(dst,src,NumOfEle);…h(huán)alfmaxNumber=dst[0];uint16_tindexOfMaxNumber=((uint16_t*)dst)[1];__bang_count(dst,src,NumOfEle);…uint16_tcounter=((uint16_t*)dst)[0];BANG數(shù)學(xué)庫實(shí)例接口：__bang_cycle_add(half*dst,half*src0,half*src1,int32_tNumOfEle0,int32_tNumOfEle1)語義：將src0元素分成整數(shù)段，每段元素與src1元素個數(shù)相等，然后將每一段與src1進(jìn)行向量add操作，結(jié)果存入dst中接口：__bang_transpose(half*dst,half*src,constint32_th,constint32_tw)語義：將src[h][w]轉(zhuǎn)置成dst[w][h]70BANG數(shù)學(xué)庫實(shí)例接口：__bang_maxpool(half*dst,half*src,constint32_tic,constint32_tih,constint32_tiw,constint32_tkh,constint32_tkw[,constint32_tsx,constint32_tsy])語義：對src進(jìn)行maxpool操作，結(jié)果存入dst中，其中：src[ih,iw,ic]slide-window[kh,kw]stride[sx,sy]，如果沒有指定，則使用[kh,kw]作為stridedst[h,w,c]71BANG數(shù)學(xué)庫實(shí)例接口：__bang_maxpool_index(uint16_t*dst,half*src,constint32_tic,constint32_tih,constint32_tiw,constint32_tkh,constint32_tkw[,constint32_tsx,constint32_tsy])語義：

與__bang_maxpool語義類似，不同是每次滑窗時選擇最大元素的index而非最大元素本身，src[ih,iw,ic]slide-window[kh,kw]stride[sx,sy]，如果沒有指定，則使用[kh,kw]作為stridedst[h,w,c]72編程接口實(shí)例：BANG程序BANG單核向量加法BANG多核向量加法BANG矩陣乘法涉及MLU上的CNRT的編程接口，包括Kernel控制和運(yùn)行時接口等73設(shè)備端代碼：BANG實(shí)現(xiàn)74BANG單核向量加法

LENBANG單核向量加法

主機(jī)端代碼：數(shù)據(jù)準(zhǔn)備創(chuàng)建Queue選擇0號設(shè)備傳入計算數(shù)據(jù)并行規(guī)模為1，只啟用1個核75任務(wù)調(diào)度類型為BLOCK設(shè)備端申請內(nèi)存主機(jī)端構(gòu)造輸入和結(jié)果準(zhǔn)備Notifier掐時間主機(jī)端代碼：Kernel調(diào)用76BANG單核向量加法

調(diào)用Kernel并在前后插入Notifier時間戳等待Kernel執(zhí)行完畢將計算結(jié)果拷回校驗(yàn)計算結(jié)果打印端到端耗時釋放資源77BANG多核向量加法

設(shè)備端BANG代碼主機(jī)端任務(wù)規(guī)模和調(diào)度單位BANG矩陣乘法bang_mulbang_transposebang_add66bang_mulbang_transposebang_add對位乘法BANG矩陣乘法第八章智能編程語言為什么需要智能編程語言智能計算系統(tǒng)抽象架構(gòu)智能編程模型智能編程語言基礎(chǔ)智能應(yīng)用編程接口智能編程語言實(shí)例智能應(yīng)用功能調(diào)試智能應(yīng)用性能調(diào)優(yōu)智能編程語言的應(yīng)用80智能應(yīng)用功能調(diào)試功能調(diào)試接口功能調(diào)試工具功能調(diào)試實(shí)踐81功能調(diào)試接口82__assert(boolflag)：abortthekernelifflagisfalse__abort():exitthekernelwitherrorcode-1exit(intstatus):exitthekernelwithcodestatusprintf(“<格式化字符串>”,<參數(shù)變量>):將字符串打印到屏幕BANG調(diào)試：格式化輸出在設(shè)備側(cè)代碼中可以使?printf實(shí)現(xiàn)格式化輸出，輸出結(jié)果默認(rèn)打印到控制臺。在兼容C語言?部分格式化字符規(guī)范的前提下，新增%hf等格式化字符?持。

%[標(biāo)志字段][寬度字段][.精度字段][?度字段]類型字段83功能調(diào)試工具面向智能編程語言BCL的調(diào)試器整體流程：調(diào)試前準(zhǔn)備、調(diào)試器托管、狀態(tài)查看及錯誤分析調(diào)試前準(zhǔn)備配置調(diào)試目標(biāo)設(shè)備號，配置設(shè)備號為1的命令：增加調(diào)試信息，編譯時（編譯器為bclc）使用-g選項(xiàng)增加調(diào)試信息，使用1/2/3等選項(xiàng)指定調(diào)試信息的等級84調(diào)試器托管通過調(diào)試器執(zhí)行被調(diào)試程序來實(shí)現(xiàn)托管。由于是異構(gòu)程序，必須從主機(jī)端程序啟動，如果要在設(shè)備端Kernel程序的入口處停住，可以使用：針對多核架構(gòu)DLP考慮在不同核間切換，可以用info命令查看當(dāng)前調(diào)試焦點(diǎn)并使用focus命令進(jìn)行切換

85顯示當(dāng)前斷點(diǎn)情況切換到cluster2core1調(diào)試器托管采用break命令可以根據(jù)函數(shù)名、代碼行號、指令地址以及Kernel入口來增加斷點(diǎn)。在break命令中可以使用if語句配置條件斷點(diǎn)。斷點(diǎn)的查看和刪除則可以分別使用info和delete命令來完成

86條件斷點(diǎn)斷點(diǎn)刪除狀態(tài)查看查看變量內(nèi)容：調(diào)試時可以直接根據(jù)變量名采用print命令來打印相關(guān)內(nèi)容查看寄存器內(nèi)容：寄存器內(nèi)容則可以采用inforegisters命令進(jìn)行查看查看地址內(nèi)容：指定地址中的數(shù)據(jù)內(nèi)容可以通過examine（或x）命令進(jìn)行查看

87(bcl-gdb)printa$1=233(bcl-gdb)printarray$2=(1,2,3,4)(bcl-gdb)x/2wd0x2333

0x2333:666888功能調(diào)試實(shí)踐以智能編程語言BCL和相應(yīng)調(diào)試器BCL-GDB為例介紹串行及并行程序的調(diào)試示例BCL實(shí)現(xiàn)的快速排序程序88入口函數(shù)核心代碼串行調(diào)試程序使用設(shè)備端編譯器編譯出帶調(diào)試信息的設(shè)備端二進(jìn)制kernel.o，并用gcc編譯鏈接出主機(jī)端的可執(zhí)行程序89指定在函數(shù)名SplitMiddle處插入斷點(diǎn)顯示在kernel.dlp中的48行運(yùn)行到函數(shù)入口處暫停，并打印出相應(yīng)代碼串行調(diào)試程序采用backtrace（bt）命令查看當(dāng)前函數(shù)的調(diào)用棧：使用layoutsrc命令查看源碼和當(dāng)前斷點(diǎn)：90SplitMiddle在kernel.dlp中的61行被調(diào)用SplitMiddle的源碼串行調(diào)試程序使用display命令和單步執(zhí)行觀察變量狀態(tài)：

91(bcl-gdb)display*m1:*m=869(bcl-gdb)n1:*m=869(bcl-gdb)nBreakpoint1,SplitMiddle(left=0,m=0x200440,right=54)atkernel.dlp:481:*m=128(bcl-gdb)n1:*m=128(bcl-gdb)nBreakpoint1,SplitMiddle(left=0,m=0x200440,right=6)atkernel.dlp:481:*m=112(bcl-gdb)單步執(zhí)行結(jié)果串行調(diào)試程序使用up命令返回上一級調(diào)用棧，然后打印NRAM地址空間變量nBuff中數(shù)據(jù)

92從SplitMiddle的遞歸調(diào)用中返回，直到kernel.dlp中的第一次調(diào)用并行程序調(diào)試以kernel.dlp為例，可以通過4個核并行執(zhí)行該程序，主機(jī)端調(diào)用時設(shè)置UNION1任務(wù)執(zhí)行93通過taskId指定更新：local[0][0]local[1][0]local[2][0]local[3][0]并行程序調(diào)試通過break、continue、info、focus、list等命令觀察多核計算的執(zhí)行及調(diào)試核心的切換94調(diào)試核心切換到：cluster0，core3打印斷點(diǎn)信息打印調(diào)試狀態(tài)Continue繼續(xù)執(zhí)行顯示代碼片段第八章智能編程語言為什么需要智能編程語言智能計算系統(tǒng)抽象架構(gòu)智能編程模型智能編程語言基礎(chǔ)智能應(yīng)用編程接口智能編程語言實(shí)例智能應(yīng)用功能調(diào)試智能應(yīng)用性能調(diào)優(yōu)智能編程語言的應(yīng)用95智能應(yīng)用性能調(diào)優(yōu)性能分析方法性能分析工具性能調(diào)優(yōu)方法96性能分析方法識別瓶頸的過程先找到耗時長的部分再通過硬件計數(shù)器（PerformanceCounter）分析硬件執(zhí)行特征

97通知（Notifier）接口硬件計數(shù)器接口使用Notifier接口獲取vector_mult的執(zhí)行時間性能調(diào)優(yōu)工具在程序外部監(jiān)控程序的運(yùn)行狀態(tài)，分析其執(zhí)行瓶頸，找到優(yōu)化空間?？梢苑譃閮深悾簯?yīng)用級性能剖析工具系統(tǒng)級性能監(jiān)控工具應(yīng)用級性能剖析工具具體使用流程分為兩個階段：1）采用record命令來運(yùn)行可執(zhí)行程序并生成相應(yīng)的性能分析報告；2）采用report或者kernel命令查看性能分析報告，獲取包括執(zhí)行時間、調(diào)用關(guān)系以及性能計數(shù)器等信息。98應(yīng)用級性能剖析工具運(yùn)行bcl-perfrecord./vecMult記錄用戶程序運(yùn)行數(shù)據(jù)，生成日志文件運(yùn)行bcl-perfkernel解析日志文件，輸出核函數(shù)執(zhí)行的硬件性能計數(shù)器數(shù)據(jù)

99向量功能單元利用率向量功能單元執(zhí)行時間片外訪存量和帶寬應(yīng)用級性能剖析工具運(yùn)行bcl-perfreport查看函數(shù)的執(zhí)行時間：

100主機(jī)端函數(shù)耗時設(shè)備端函數(shù)耗時內(nèi)存拷貝耗時內(nèi)存拷貝量系統(tǒng)級性能監(jiān)控工具系統(tǒng)級性能監(jiān)控工具主要利用驅(qū)動通過讀取寄存器的方式來收集硬件的靜態(tài)和動態(tài)信息

101性能調(diào)優(yōu)方法DLP核函數(shù)調(diào)優(yōu)的核心思想：如何充分利用大規(guī)模的并行計算單元使用片上存儲

優(yōu)化前（完成長度為16384的向量對位乘法）：

使用NRAM優(yōu)化后：

102原始的兩個輸入和輸出都在GDRAM上GDRAM的延遲太高，計算單元大部分時間在等數(shù)使用片上的低延遲NRAM向量化張量化的基本原理是將大量標(biāo)量計算合并為張量計算，使用智能編程語言的張量計算語句改寫代碼，充分利用硬件的張量計算單元，提升程序運(yùn)行速度。

103把標(biāo)量的for循環(huán)替換成一條張量計算語句，使用張量計算單元來加速軟件流水智能處理器的計算和訪存單元可以并行工作，編程時可以顯式將無依賴的計算和訪存指令放在一起，從而提高硬件的利用率和程序性能。計算和訪存并行最常用的方法是三級流水。

104流水啟動軟件流水

105流水排空流水循環(huán)多核并行針對（程序員可見的）多核，可以將一個任務(wù)分拆到多個核上并行計算，進(jìn)一步提升程序性能。

106把向量乘法拆分到4個核上，每個核計算4096大小的向量乘法，用taskId進(jìn)行索引第八章智能編程語言為什么需要智能編程語言智能計算系統(tǒng)抽象架構(gòu)智能編程模型智能編程語言基礎(chǔ)智能應(yīng)用編程接口智能編程語言實(shí)例智能應(yīng)用功能調(diào)試智能應(yīng)用性能調(diào)優(yōu)智能編程語言的應(yīng)用107智能編程語言的應(yīng)用高性能庫算子開發(fā)實(shí)踐：基于BANG的高性能庫算子開發(fā)編程框架算子開發(fā)實(shí)踐：基于BANG開發(fā)開發(fā)算子并集成到PyTorch108高性能庫算子開發(fā)高性能庫（如MKL、cuDNN和CNNL等）提供了常見算子在特定平臺上的高性能實(shí)現(xiàn)，方便用戶以API的形式直接調(diào)用如何用智能編程語言擴(kuò)展高性能庫算子109高性能庫算子開發(fā)的關(guān)鍵在于：1）Kernel代碼邏輯的開發(fā)與優(yōu)化2）高性能庫算子接口API的使用高性能庫算子開發(fā)：原理及流程高性能庫一般提供C或C++的API，通過智能編程語言開發(fā)并經(jīng)AOT編譯（Ahead-Of-Time編譯）封裝為動態(tài)庫，其API實(shí)現(xiàn)主體運(yùn)行在主機(jī)端，內(nèi)部的算子核函數(shù)運(yùn)行在DLP硬件。算子一般支持可變的輸入規(guī)模，即編程框架或用戶直接調(diào)用算子時，不會觸發(fā)重新編譯，已經(jīng)編譯的二進(jìn)制可動態(tài)適配可變的輸入形狀，這種AOT編譯且可變的高性能算子庫可以為上層用戶節(jié)省運(yùn)行時編譯的時間。高性能庫中的算子開發(fā)流程一般為：1）設(shè)計算子的API，根據(jù)輸入輸出的Tensor描述符設(shè)計異構(gòu)并行Kernel的拆分策略；2）使用智能編程語言進(jìn)行算子的邏輯開發(fā)，編寫相應(yīng)的Kernel代碼；3）通過異構(gòu)混合編譯工具鏈進(jìn)行編譯，包含Kernel二進(jìn)制的對象文件；4）通過主機(jī)端的編譯器將Kernel二進(jìn)制對象文件、主機(jī)端對象文件鏈接成為動態(tài)庫；5）開發(fā)應(yīng)用程序調(diào)用動態(tài)庫和依賴的運(yùn)行時庫執(zhí)行高性能算子。110高性能庫算子開發(fā)：算子庫結(jié)構(gòu)一個由BCL開發(fā)的典型高性能算子庫工程目錄分為cmake、core、kernels、scripts、test等5個模塊，此外還有算子的API頭文件dlp_ops.hcmake目錄存放異構(gòu)混合編譯相關(guān)的配置，需要處理*.dlp文件的編譯和*.cpp文件的編譯鏈接core目錄存放了算子庫和運(yùn)行時API銜接的代碼，負(fù)責(zé)管理context，queue，memory，devicekernels目錄存放算子Device端的Kernel實(shí)現(xiàn)，還有Host端多核拆分策略、<<<>>>語法糖調(diào)用等代碼scripts目錄存放工程開發(fā)必須的腳本和工具test目錄存放gtest等算子測試框架或經(jīng)典對照的cpu算子實(shí)現(xiàn)111高性能庫算子開發(fā)：算子庫API一個高性能算子庫的頭文件dlp_ops.h需要張量描述符，運(yùn)算描述符等數(shù)據(jù)結(jié)構(gòu)，還需要提供算子運(yùn)算接口，以及與之配套的數(shù)據(jù)結(jié)構(gòu)和運(yùn)行時相關(guān)的接口等。112高性能庫算子開發(fā):實(shí)踐核函數(shù)性能優(yōu)化的方法（