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

智能計算系統(tǒng)

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

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

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

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

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

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

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

……

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

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

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

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

⑥運行結果獲取

⑦資源釋放60智能編程模型實例：BANG異構編程BANG語言是針對MLU(MachineLearningUnit)硬件提出的編程語言，兼顧云邊端等不同目標平臺，提供高性能機器學習計算支持提供通用的異構編程模型，方便用戶擴展自己的應用程序提供高效的編程接口，充分發(fā)揮底層硬件特性基于C/C++語言的擴展，簡單易用61CNNLMLUBANGMLU上的異構編程語言，可用于定制開發(fā)編程框架/機器學習庫所需的算子MLU上的機器學習庫，用于智能芯片上加速各種機器學習算法BANG異構編程：流程BANG異構程序同樣采用分離式編程，即主機端與設備端程序分開編程并分別編譯，最后鏈接成一個可執(zhí)行程序主機端為C/C++程序，通常需調(diào)用CNRT運行時接口完成以下步驟準備輸入數(shù)據(jù)拷貝輸入數(shù)據(jù)到MLU準備Kernel參數(shù)創(chuàng)建Queue指定Kernel任務規(guī)模以及調(diào)度類型啟動KernelMLU到主機的輸出數(shù)據(jù)拷貝資源的釋放設備端使用BANG語言特定的語法規(guī)則和接口進行編程62常用CNRT接口名稱功能描述cnrtGetDevice獲取設備cnrtQueueCreate創(chuàng)建一個新的Queue，默認異步運行cnrtQueueDestroy銷毀QueuecnrtQueueSync直到之前Queue中所有的Function都完成，阻塞其他FunctioncnrtInvokeKernel通過在MLU上給定的參數(shù)塊，啟動Kernelfoo<<<>>>()編譯器和運行時配合完成的語法糖，可替代cnrtInvokeKernelcnrtMalloc分配給定空間的設備內(nèi)存cnrtFree釋放指針指向的空間cnrtMemcpy從源地址拷貝指定字節(jié)數(shù)據(jù)到目的地址BANG異構編程示例：主機端31頭文件bang.h包含主機和設備端的接口聲明創(chuàng)建異步Queue配置Kernel的并行規(guī)模申請設備端內(nèi)存配置Kernel的并行任務類型同步拷貝輸入數(shù)據(jù)至設備內(nèi)存異步啟動Kernel函數(shù)同步Queue等待Kernel計算完成同步拷貝輸出到主機內(nèi)存釋放設備內(nèi)存和QueueBANG異構編程示例：設備端64設備端源碼和主機端源碼可以放在同一個源碼文件中，文件名后綴為mluBANG程序kernel的入口函數(shù)需要用__mlu_entry__或__mlu_global__標記__nram__標記nram空間的變量將數(shù)據(jù)從GDRAM拷貝到NRAM上將NRAM上數(shù)據(jù)調(diào)用BANGC的transpose、mul和reduce_sum等接口計算將計算的結果從NRAM拷貝到GDRAM，以供主機端程序讀取調(diào)用設備端的計時函數(shù)統(tǒng)計硬件時間并打印耗時BANG程序編譯與鏈接65BANG的源碼可以寫在同一份源碼文件中，編譯器自動完成主機端和設備端的編譯和鏈接foo.mlu混合編程源碼被cncc前端driver拆分成主機和設備源碼主機端源碼使用clang進行編譯，得到主機端的二進制對象文件設備端源碼根據(jù)指定的單一arch或多arch，分別編譯成目標arch的對象文件設備端多arch的bin對象被設備端鏈接器鏈接成fatbin，并包裝成主機端可識別的對象文件主機端鏈接器將主機和設備端對象文件鏈接為可執(zhí)行文件張量計算語句實例：BANG數(shù)學庫BANG語言將MLU的數(shù)學類和神經(jīng)網(wǎng)絡類指令封裝成庫函數(shù)BANG數(shù)學庫函數(shù)是在MLU架構上進行高性能編程的關鍵BANG數(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ù)進行操作源和目的都是NRAM上的數(shù)據(jù)66BANG數(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ù)學庫實例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兩個張量進行element-wise的eq操作，并將結果存入dst中接口：__bang_select(half*dst,half*src0,halfsrc1,int32_tNumOfEle)語義：如果src1中某個元素非0，則從src0中選擇對應位置元素存入dst中。dst中的結果有兩部分，第一部分是選出的元素個數(shù)，第二部分是選出的元素集合示例：BANG數(shù)學庫實例接口：__bang_max(half*dst,half*src,int32_tNumOfEle)語義：從張量src中選擇最大值的元素，存入dst中，結果有兩部分，第一部分是選出的元素，第二部分是選擇元素在src中index示例:接口：__bang_count(uint16_t*dst,half*src,int32_tNumOfEle)語義：計算src中非0元素的個數(shù)，結果存入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ù)學庫實例接口：__bang_cycle_add(half*dst,half*src0,half*src1,int32_tNumOfEle0,int32_tNumOfEle1)語義：將src0元素分成整數(shù)段，每段元素與src1元素個數(shù)相等，然后將每一段與src1進行向量add操作，結果存入dst中接口：__bang_transpose(half*dst,half*src,constint32_th,constint32_tw)語義：將src[h][w]轉置成dst[w][h]70BANG數(shù)學庫實例接口：__bang_maxpool(half*dst,half*src,constint32_tic,constint32_tih,constint32_tiw,constint32_tkh,constint32_tkw[,constint32_tsx,constint32_tsy])語義：對src進行maxpool操作，結果存入dst中，其中：src[ih,iw,ic]slide-window[kh,kw]stride[sx,sy]，如果沒有指定，則使用[kh,kw]作為stridedst[h,w,c]71BANG數(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編程接口實例：BANG程序BANG單核向量加法BANG多核向量加法BANG矩陣乘法涉及MLU上的CNRT的編程接口，包括Kernel控制和運行時接口等73設備端代碼：BANG實現(xiàn)74BANG單核向量加法

LENBANG單核向量加法

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

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

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

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

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

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

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

0x2333:666888功能調(diào)試實踐以智能編程語言BCL和相應調(diào)試器BCL-GDB為例介紹串行及并行程序的調(diào)試示例BCL實現(xiàn)的快速排序程序88入口函數(shù)核心代碼串行調(diào)試程序使用設備端編譯器編譯出帶調(diào)試信息的設備端二進制kernel.o，并用gcc編譯鏈接出主機端的可執(zhí)行程序89指定在函數(shù)名SplitMiddle處插入斷點顯示在kernel.dlp中的48行運行到函數(shù)入口處暫停，并打印出相應代碼串行調(diào)試程序采用backtrace（bt）命令查看當前函數(shù)的調(diào)用棧：使用layoutsrc命令查看源碼和當前斷點：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í)行結果串行調(diào)試程序使用up命令返回上一級調(diào)用棧，然后打印NRAM地址空間變量nBuff中數(shù)據(jù)

92從SplitMiddle的遞歸調(diào)用中返回，直到kernel.dlp中的第一次調(diào)用并行程序調(diào)試以kernel.dlp為例，可以通過4個核并行執(zhí)行該程序，主機端調(diào)用時設置UNION1任務執(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ào)試狀態(tài)Continue繼續(xù)執(zhí)行顯示代碼片段第八章智能編程語言為什么需要智能編程語言智能計算系統(tǒng)抽象架構智能編程模型智能編程語言基礎智能應用編程接口智能編程語言實例智能應用功能調(diào)試智能應用性能調(diào)優(yōu)智能編程語言的應用95智能應用性能調(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)控程序的運行狀態(tài)，分析其執(zhí)行瓶頸，找到優(yōu)化空間?？梢苑譃閮深悾簯眉壭阅芷饰龉ぞ呦到y(tǒng)級性能監(jiān)控工具應用級性能剖析工具具體使用流程分為兩個階段：1）采用record命令來運行可執(zhí)行程序并生成相應的性能分析報告；2）采用report或者kernel命令查看性能分析報告，獲取包括執(zhí)行時間、調(diào)用關系以及性能計數(shù)器等信息。98應用級性能剖析工具運行bcl-perfrecord./vecMult記錄用戶程序運行數(shù)據(jù)，生成日志文件運行bcl-perfkernel解析日志文件，輸出核函數(shù)執(zhí)行的硬件性能計數(shù)器數(shù)據(jù)

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

100主機端函數(shù)耗時設備端函數(shù)耗時內(nèi)存拷貝耗時內(nèi)存拷貝量系統(tǒng)級性能監(jiān)控工具系統(tǒng)級性能監(jiān)控工具主要利用驅動通過讀取寄存器的方式來收集硬件的靜態(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向量化張量化的基本原理是將大量標量計算合并為張量計算，使用智能編程語言的張量計算語句改寫代碼，充分利用硬件的張量計算單元，提升程序運行速度。

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

104流水啟動軟件流水

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

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