張政馗, 龐為光, 謝文靜, 呂鳴松, 王 義

(東北大學 計算機科學與工程學院 智慧系統(tǒng)實驗室,遼寧 沈陽 110819)

深度學習(deep learning,簡稱DL)[1]是機器學習(machine learning,簡稱ML)[2]的一個分支,都是能夠讓計算機根據(jù)數(shù)據(jù)進行預測,并且改進其預測或行為的一組方法.其核心思想是:在訓練階段,以最小化損失函數(shù)(loss function)為引導,通過梯度下降算法(gradient descent)來調(diào)整計算模型的權(quán)重(weight)和偏置(bias)兩種參數(shù);在推理階段,則使用輸入數(shù)據(jù)和訓練好的模型參數(shù)來計算預測值.深度學習的主要特征是采用了分層的深層神經(jīng)網(wǎng)絡(deep neural network,簡稱DNN)模型:每一層抽象出不同的簡單特征,不同層的簡單特征可以疊加組合成更復雜的特征,并使用這些組合的復雜特征解決問題[3].相比而言,傳統(tǒng)機器學習算法很難抽象出足夠有效的特征.現(xiàn)在的DNN 模型已經(jīng)有百萬個人工神經(jīng)元,深度達到了幾十層.例如2015 年,ImageNet 圖像分類比賽(ImageNet large scale visual recognition challenge,簡稱ILSVRC)的冠軍是ResNet,其最大深度為152 層,屬于超深神經(jīng)網(wǎng)絡.

隨著大數(shù)據(jù)、深度學習算法與高性能計算技術(shù)的不斷發(fā)展,深度學習在人工智能(artificial intelligence,簡稱AI)領(lǐng)域一些最困難的問題上取得了重大突破,并且被應用于圖像識別、語音識別、自然語言處理、自動駕駛、生物信息處理、電腦游戲、搜索引擎、廣告推薦、醫(yī)學自動診斷和金融分析等各大領(lǐng)域.現(xiàn)如今,物聯(lián)網(wǎng)、可穿戴設備正在從概念變成現(xiàn)實,5G 網(wǎng)絡建設已全面起步,這些全新的信息技術(shù)必將促進人工智能技術(shù)從工作站和服務器端遷移到各種嵌入式終端.這其中也包括了安全攸關(guān)系統(tǒng),典型案例是全球互聯(lián)網(wǎng)巨頭和汽車巨頭正如火如荼地研究以深度學習為基礎的自動駕駛系統(tǒng).特斯拉公司在2015 年已經(jīng)將自動駕駛技術(shù)(autopilot)投入了商用.然而,近兩年接連發(fā)生的幾起關(guān)于自動駕駛汽車的嚴重交通事故表明,深度學習技術(shù)的成熟度還遠未達到安全攸關(guān)應用的要求.事實上,無論從功能層面還是非功能層面,深度學習賦能的安全攸關(guān)系統(tǒng)亟待有效的驗證技術(shù)來保障整個系統(tǒng)計算邏輯的正確性和響應外部事件的實時性[4].

在功能層面保證DNN 的正確性會面臨3 大困難[5,6]:(1) DNN 具有魯棒性缺陷,對抗樣本的輕微擾動就可能導致DNN 給出錯誤的判斷,這對安全攸關(guān)系統(tǒng)是一個致命的問題;(2) DNN 不具備可解釋性,這導致無法對DNN 模型進行(精確的)形式化建模和屬性描述;(3) DNN 無窮的輸入空間、大量的神經(jīng)元、非線性激活函數(shù)會導致驗證時發(fā)生狀態(tài)空間爆炸.針對這些挑戰(zhàn),學術(shù)界在DNN 的可解釋性(interpretability)[2]、建模(modeling)[7]、測試(testing,把神經(jīng)網(wǎng)絡看作一個組件進行白盒測試,力求讓測試案例最大化地覆蓋整個網(wǎng)絡結(jié)構(gòu),包括神經(jīng)元、條件和分支等)[8]、證偽(falsification,又稱半形式化驗證(semi-formal verification),其目的是產(chǎn)生讓整個系統(tǒng)違反設計規(guī)范的極端測試案例(corner test case))[9,10]、驗證(verification)[11-13]等方面展開了大量的研究.此外,歐美政府也開始重視可信人工智能方向的研究.2017 年,美國國防部高級研究計劃局(Defense Advanced Research Projects Agency,簡稱DARPA)發(fā)起了“XAI 計劃”,旨在研究如何更好地理解AI 系統(tǒng)的預測邏輯.2018 年,歐盟頒布的GDPR(general data protection regulation)要求AI 算法具有透明性和可解釋性.賓夕法尼亞大學成立了安全AI 研究中心(PRECISE center for safe AI)[14],并致力于研究安全AI 系統(tǒng)的設計和形式化驗證技術(shù).

在非功能層面保證DNN 的實時性是指:在實時嵌入式系統(tǒng)中,DNN 任務(特別說明:本文中的深度學習和DNN 任務都是指深度學習的推理階段(inference phase),即用訓練好的DNN 模型和輸入數(shù)據(jù)進行預測)的運行時間不能超過該任務所指定的時限.而實現(xiàn)這一目標同樣面臨3 個方面的挑戰(zhàn):第一,嵌入式平臺的計算資源和內(nèi)存資源是非常有限的,電源供給也受電池容量的限制,散熱條件也遠比不上配備了大型冷卻設備的數(shù)據(jù)中心;第二,DNN 任務已經(jīng)被普遍部署在多核CPU+GPU SoC 異構(gòu)嵌入式平臺(例如NVIDIA Jetson TX2[15],NVIDIA DRIVE PX2[16]),但是學術(shù)界對調(diào)度異構(gòu)多核嵌入式系統(tǒng)的理論研究還不夠成熟;第三,目前為了保證實時性,一般讓GPU 只負責運行某種特定的DNN 任務,例如圖像識別,但在未來,需要在GPU 上并行運行多個混合關(guān)鍵任務,例如未來的自動駕駛汽車上,GPU 可能同時并行處理多個DNN 推理任務(圖像識別、語音識別、自然語言處理)、DNN 在線訓練任務[17]、容錯備份任務等關(guān)鍵任務[18],同時還需要執(zhí)行儀表盤的信息顯示任務、多媒體或游戲系統(tǒng)的2D/3D 渲染任務等非關(guān)鍵任務,這種多任務需求更增加了實時嵌入式操作系統(tǒng)實施任務調(diào)度和資源管理的復雜度.

近兩年,實時系統(tǒng)學術(shù)界開始關(guān)注深度學習所帶來的新問題.其中,德克薩斯大學Liu 等學者在DNN 網(wǎng)絡的實時優(yōu)化設計和DNN 網(wǎng)絡的能耗分析方面做了開創(chuàng)性的研究[19,20],北卡羅萊納大學教堂山分校Smith 教授對NIVIDA GPU 的調(diào)度機制進行了深入細致的探究[15,21].但是,目前還缺乏對實時深度學習推理系統(tǒng)的主要設計難點以及分析的方法進行系統(tǒng)且深入闡述的文獻資料.針對上述問題,本綜述介紹了深度學習推理過程應用于實時嵌入式系統(tǒng)所面臨的關(guān)鍵設計問題,并比較了現(xiàn)有主要解決方案的優(yōu)缺點.針對深層神經(jīng)網(wǎng)絡的輕量化設計、GPU 時間分析與任務調(diào)度、CPU+GPU SoC 異構(gòu)平臺的資源管理、深層神經(jīng)網(wǎng)絡與網(wǎng)絡加速器的協(xié)同設計等多個方面,對現(xiàn)有的研究工作進行了分析和總結(jié),旨在為相關(guān)領(lǐng)域的研究者提供參考.

1 背景知識

本節(jié)簡要介紹深度學習的發(fā)展史、基本概念、深層神經(jīng)網(wǎng)絡DNN、DNN 運行框架以及運行DNN 實時任務的嵌入式硬件平臺,包括嵌入式GPU 和網(wǎng)絡加速器.

1.1 深度學習簡史

眾所周知,深度學習已經(jīng)成為了當下最流行的技術(shù)詞匯,似乎是一項才剛剛發(fā)明的新技術(shù).但是事實上,有關(guān)深度學習的研究歷史(見表1)可以追溯到1943 年由McCulloch 教授和Pitts 教授提出的McCulloch-Pitts Neuron 單神經(jīng)元計算結(jié)構(gòu)[22].該結(jié)構(gòu)會通過N個權(quán)重與N個輸入來計算加權(quán)和作為輸出,是現(xiàn)代神經(jīng)元結(jié)構(gòu)的雛形.1958 年,Rosenblatt 教授提出的感知器(perceptron)模型[23]是首個可以根據(jù)樣本數(shù)據(jù)來學習特征權(quán)重的模型,對現(xiàn)代機器學習的發(fā)展產(chǎn)生了巨大影響.1969 年,Minsky 教授在文獻[24]中證明:感知器模型只能解決線性可分問題,無法解決異或問題,并且給出了“基于感知器的研究注定將失敗”的結(jié)論.這導致了神經(jīng)網(wǎng)絡的第一次重大低潮期,在之后的10 多年內(nèi),神經(jīng)網(wǎng)絡的研究幾乎處于停滯狀態(tài).

Table 1 Scientific milestones and big news on deep learning history[3,25]表1 深度學習發(fā)展史上的重大科學及新聞事件[3,25]

這期間,人們在認知生物學領(lǐng)域取得了重要進展,發(fā)現(xiàn)了認知的兩個重要機制:一個是抽象,另一個是迭代.從原始信號,做低層抽象,逐漸向高層抽象迭代,在迭代中抽象出更高層的模式[25].到了20 世紀80 年代末,分布式知識表達(distributed representation)[26]和反向傳播算法(back propagation)[27]的提出,開啟了神經(jīng)網(wǎng)絡研究的第二階段的序幕.分布式知識表達是深度學習的重要性質(zhì),基本思想是:先讓每個神經(jīng)元表達簡單的特征,再把神經(jīng)元組合起來用于描述復雜的特征.相比本地式知識表達(即一個特征對應一個神經(jīng)元),分布式表達描述相同大小的特征空間需要的神經(jīng)數(shù)量少很多,極大節(jié)約了存儲空間.1986 年,Rumelhart,Hinton 和Williams 這3 位教授提出的反向傳播算法則大幅降低了訓練神經(jīng)網(wǎng)絡所需要的時間.時至今日,反向傳播算法仍是訓練深層神經(jīng)網(wǎng)絡的主要方法.同期,計算機的飛速發(fā)展也促進了卷積神經(jīng)網(wǎng)絡(convolutional neural network,簡稱CNN)、遞歸神經(jīng)網(wǎng)絡(recurrent neural network,簡稱RNN)、長短期記憶模型(long short-term memory,簡稱LSTM)等模型的發(fā)展.

到了2010 年左右,計算機性能的極大提升和大數(shù)據(jù)互聯(lián)網(wǎng)+的發(fā)展,使得曾經(jīng)阻礙神經(jīng)網(wǎng)絡發(fā)展的計算力和訓練樣本量問題得到了解決,從此深度學習的發(fā)展一日千里.2012 年,ImageNet 舉辦的ILSVRC 圖像分類競賽中,由Krizhevsky 教授實現(xiàn)的深度學習系統(tǒng)AlexNet[28]贏得了冠軍.此后,深度學習算法的性能在圖像識別領(lǐng)域已經(jīng)完全碾壓了傳統(tǒng)機器學習算法[29](如支持向量機SVM 等),并且在2013 年之后的ILSVRC 中基本就只有深度學習算法參賽了.2013 年,深度學習被麻省理工(MIT)評為“年度十大科技突破之一”.2015 年10 月,特斯拉在Model 系列車型上開啟了自動駕駛功能(autopilot),標志著基于深度學習的自動駕駛技術(shù)已開始進入了商用階段.2016 年和2017 年,Google 的子公司DeepMind 基于深度學習研發(fā)的AlphaGo 擊敗了圍棋大師李世石和柯潔,一度引發(fā)了全世界對人工智能的恐慌.2019 年1 月,DeepMind 研發(fā)的AlphaStar 擊敗了經(jīng)典戰(zhàn)略游戲《星際爭霸》的職業(yè)電競選手,標志著人工智能在部分信息博弈中已經(jīng)可以戰(zhàn)勝人類了.

1.2 深層神經(jīng)網(wǎng)絡

深度學習是“一類通過多層非線性變換對高復雜性數(shù)據(jù)建模算法的合集”,深層神經(jīng)網(wǎng)絡(DNN)是實現(xiàn)“多層非線性變換”的最常用的一種方式,兩者互為代名詞[3].DNN 的兩個非常重要的特征是多層和非線性[30]:多層是為了符合分布式知識表達(第1.1 節(jié))的要求,非線性是為了解決更加復雜的問題.因為在現(xiàn)實世界中,絕大部分的問題都是無法線性分割的,而任何線性模型的組合仍然還是線性的.為DNN 提供非線性表達能力的是激活函數(shù)(activation function).圖1 展示了一個神經(jīng)元的輸出是輸入數(shù)據(jù)加權(quán)和與偏置加和之后經(jīng)過激活函數(shù)非線性變換得到的結(jié)果.激活函數(shù)的特點是可微分并且單調(diào).常用的激活函數(shù)有 Sign,Sigmoid,Tanh,ReLU,P-ReLU,Leaky-ReLU,ELU,Maxout 等.損失函數(shù)(loss function)是用于度量DNN 輸出結(jié)果向量與樣本期望向量之間差距的函數(shù).常用的損失函數(shù)有交叉熵(cross entropy)、均方差(mean square error,簡稱MSE)、Log、L1 Loss、L2 Loss、Elastic Net 等.構(gòu)造一個深層神經(jīng)網(wǎng)絡就是確定網(wǎng)絡的3 個組成部分:DNN 的架構(gòu)(或稱為拓撲結(jié)構(gòu))、激活函數(shù)與損失函數(shù)、訓練DNN 的算法.DNN 的使用一般分為訓練和推理兩個階段.訓練DNN 即為網(wǎng)絡中的神經(jīng)元找到最優(yōu)權(quán)值配置,主流的訓練算法是梯度下降算法[1]和反向傳播算法(第1.1 節(jié)).訓練得到的網(wǎng)絡也稱為推理網(wǎng)絡(inference network),可以用于對測試數(shù)據(jù)集或?qū)嶋H數(shù)據(jù)的推理.

神經(jīng)網(wǎng)絡主要分為兩種類型[1]:(1) 前饋神經(jīng)網(wǎng)絡(feedforward neural networks),即有向無環(huán)圖,信號只能沿著最終輸出的那個方向傳播;(2) 反饋神經(jīng)網(wǎng)絡(feedback neural networks),即網(wǎng)絡中有回路,輸出層的信號經(jīng)過一步時移(shift)再接入到輸入層.常用的DNN 模型如卷積神經(jīng)網(wǎng)絡(CNN)屬于前饋神經(jīng)網(wǎng)絡,遞歸神經(jīng)網(wǎng)絡(RNN)和長短期記憶模型(LSTM)都屬于反饋神經(jīng)網(wǎng)絡.CNN 可以有效地降低傳統(tǒng)前饋神經(jīng)網(wǎng)絡的復雜性,并廣泛應用于圖像分類和物體識別等場景.CNN 網(wǎng)絡的架構(gòu)可以用公式[3]來表示:輸入層→(卷積層+→池化層?)+→全連接層+,其中,“卷積層+”表示一層或多層卷積層(CONV),“池化層?”表示沒有或一層池化層(POOL).卷積層的核心是卷積核,尺寸一般為3×3,5×5,7×7.相比全連接方式,卷積核的參數(shù)非常少,各層通過卷積核共享機制可以極大減少訓練階段需要優(yōu)化的總參數(shù)量.池化層可以非常有效地縮減矩陣的尺寸(主要用于減小矩陣的長和寬),從而減少最后全連接層中的參數(shù),并有防止過擬合的作用.常用的池化策略有最大池化和平均池化等.RNN引入了“記憶”的概念,主要用途是處理輸出內(nèi)容和歷史內(nèi)容有關(guān)聯(lián)的場景,比如時間序列分析、語音識別、機器翻譯、自然語言處理等.RNN 實現(xiàn)遞歸的結(jié)構(gòu)非常簡單,一般是若干個激活函數(shù)層疊加在一起組成隱藏層(DNN可以劃分成輸入層、隱藏層、輸出層.隱藏層就是輸入層和輸出層之間的多層神經(jīng)元結(jié)構(gòu)),循環(huán)元素在隱藏層上執(zhí)行相同的任務,但是輸出依賴于當前的輸入和歷史記憶.RNN 的推理過程可以用兩個公式[1]表示:St=f(UXt+WSt-1),Ot=Softmax(VSt).其中:St代表時刻t的隱藏層狀態(tài);Ot代表時刻t的輸出;Xt是時刻t的輸入;f是激活函數(shù);U表示輸入層到隱藏層的權(quán)重;V表示隱藏層到輸出層的權(quán)重;W表示隱藏層到隱藏層的權(quán)重,負責RNN 的記憶調(diào)度.LSTM 是一種特殊的RNN,用來解決RNN 存在的長期依賴問題,即相關(guān)信息與預測位置的間隔很大而導致RNN 無法學習連接信息.LSTM 的隱藏層比傳統(tǒng)RNN 復雜,是一種擁有3 種“門”的特殊網(wǎng)絡結(jié)構(gòu),從而更有效地保存長期記憶.其中:“遺忘門”的作用是讓循環(huán)神經(jīng)網(wǎng)絡“忘記”之前沒有用的信息,“輸入門”決定哪些部分記憶進入當前時刻的狀態(tài),“輸出門”決定當前時刻的輸出.

當前主流的DNN 開發(fā)及運行框架包括TensorFlow(Google)[31],PyTorch(Facebook)[32],Caffe(Berkeley 大學)[33],其他DNN 框架如Theano(Montreal 大學),Keras(Keras-Team),MXNet(Amazon),CNTK(Microsoft)的用戶基礎遠比不上前3 種.DNN 框架的運行原理(以TensorFlow 為例[34]):首先,將用戶基于應用層API 編寫的、以神經(jīng)網(wǎng)絡算法為代表訓練或推理過程表達為數(shù)據(jù)流圖計算;在運行時,把數(shù)據(jù)流圖轉(zhuǎn)換成C++核心層的細粒度、抽象化運行時狀態(tài),進而在計算設備(CPU 或GPU)上以一致而有效的方式調(diào)度執(zhí)行.主流的DNN 框架能夠支持在個人電腦、大型數(shù)據(jù)中心的服務器、嵌入式設備等多種平臺上運行.

1.3 嵌入式GPU

圖像處理器(graphics processing unit,簡稱GPU)最初是計算機系統(tǒng)中加速圖形圖像運算的專用芯片,如今的GPU 因集成了大量的處理器核心,具有了非常強大的并行處理能力,并且已經(jīng)被廣泛應用于通用計算領(lǐng)域,例如游戲開發(fā)、圖像處理、視頻處理、科學計算、大數(shù)據(jù)、深度學習等領(lǐng)域.當前,國際主流的GPU 廠商主要有通用計算領(lǐng)域的 NVIDIA,AMD(ATI),Intel(CPU 內(nèi)置顯示核心)以及移動計算領(lǐng)域的 ARM,Qualcomm,PowerVR 等.本節(jié)選擇最為廣泛的NVIDIA GPU 系列,先介紹普通PC 平臺GPU,了解GPU 的一般架構(gòu)和工作原理.在此基礎上,再介紹嵌入式GPU 的架構(gòu)和工作特點.

PC 平臺的GPU 計算卡一般通過高速的PCI-E(peripheral communications interconnect express)總線與主板上的北橋芯片相連.PCI-E 是連接GPU 卡與CPU 的全雙工高速總線.PCI-E 2.0 的傳輸速率為5GB/s,PCI-E 3.0提升到了8GB/s.PCI-E 總線能夠為每塊GPU 卡提供確定的讀寫帶寬,而不受所掛載的GPU 卡數(shù)量的影響.圖2顯示了一塊NVIDIA GPU 芯片的架構(gòu)模塊示意圖,包括3 種關(guān)鍵模塊[35]:流處理器簇(streaming multiprocessor,簡稱SM)、流處理器(streaming processor,簡稱SP)、(全局、常量、共享)內(nèi)存.“流”是一個GPU 操作隊列,該隊列中的操作(可能由多個主機線程發(fā)出)將以添加到流中的先后順序而依次執(zhí)行.可以將一個流看作是GPU 上的一個任務,不同流里的任務可以并行執(zhí)行.所以一個流對應于并發(fā)的概念,多個流對應并行的概念.GPU 是由多個SM 組成的SM 陣列,每個SM 包含8N 個SP(SP 也稱CUDA(compute unified device architecture)核,G80/GT200 中有8 個SP,RTX2080 中有64 個SP),每個SP 都是一個設計簡單的處理器.全局內(nèi)存(global memory)就是GPU 卡的顯存.紋理內(nèi)存(texture memory)和常量內(nèi)存(constant memory)都是針對全局內(nèi)存的一個特殊視圖.其中,紋理內(nèi)存是在顯示2D 或3D 圖像時存儲插值計算所需要的數(shù)據(jù),常量內(nèi)存用于存儲只讀數(shù)據(jù).每個SM 通過總線可以獨立訪問3 種內(nèi)存.每個SM 內(nèi)部都有共享內(nèi)存(shared memory).與CPU 不同,它沒有自動完成數(shù)據(jù)替換的硬件邏輯,而是完全由程序員控制,所以它是一種程序可控的高速緩存.

主流的GPU 通用編程接口包括NVIDIA 公司開發(fā)的CUDA 運算平臺、蘋果公司持有但是保持開放的OpenCL(open computing language)標準、微軟公司開發(fā)的Direct(direct compute)標準.CUDA[36]是C 語言的一種擴展,支持基于PTX(parallel thread execution)虛擬指令集的運行時編譯,即CUDA 代碼先被編譯成并行線程執(zhí)行(parallel thread execution,簡稱PTX)這種中間形式,然后再次編譯為原生的GPU 微碼;并且向前兼容(forwards compatibility),即無論GPU 的硬件結(jié)構(gòu)如何改變,為早期CUDA 設備編寫的程序依然能夠運行在最新的CUDA設備上.CUDA 編程模型是一種異構(gòu)模型.CUDA 程序可以在CPU 和GPU 上并行運行,在CPU 上運行的代碼段叫Host code,在GPU 上運行的代碼段叫Device code,其中包含若干kernel 函數(shù).每一個kernel 在GPU 上執(zhí)行時會啟動很多線程運行同樣的代碼指令,即單指令多線程(single-instruction multiple-thread,簡稱SIMT)的并行計算模型.Kernel 線程的數(shù)量取決于程序員為kernel 函數(shù)指定的參數(shù)num_blocks(線程塊數(shù))和num_threads(每個線程塊內(nèi)執(zhí)行kernel 函數(shù)的線程數(shù)).這兩個參數(shù)的配置會影響kernel 函數(shù)的執(zhí)行性能.每個SM 最多能處理的線程數(shù)是有上界,也就間接影響每個SM 最多能容納的線程塊(thread block)的數(shù)量.另外,每32 個線程組成一個線程束(wrap),線程束是GPU 調(diào)度和執(zhí)行的基本單元.一個線程塊所管理的線程束等于num_threads 除以32 并向上取整,所以應該把 num_threads 設置成 32 的整數(shù)倍,否則會造成最后一個線程束中有部分線程被閑置.CUDA 主程序啟動后,會將操作指令和數(shù)據(jù)提供給線程塊,線程塊以鎖步(lock-step)的方式廣播到每個線程束所占用的SP 中;線程塊一旦被調(diào)度到GPU 的某個SM,就必須從開始執(zhí)行到結(jié)束.執(zhí)行期間,線程束之間的數(shù)據(jù)交換也由CUDA 主程序負責管理.

以上是PC 平臺GPU 的基本結(jié)構(gòu)、CUDA 程序的基本概念和基本運行規(guī)則.下面以NVIDIA Jetson TX2[15,37]為例,介紹嵌入式CPU+GPU SoC 的結(jié)構(gòu)特點.NVIDIA Jetson TX2 SoC 計算卡(圖3)集成了一枚64 位四核ARMv8@2.0 GHz 的A57 微處理器、一枚雙核超標量ARMv8@2.0GHz 的Denver 微處理器和一枚嵌入式的Pascal 架構(gòu)GPU.兩枚CPU 各自擁有2MB L2 cache.GPU 內(nèi)包含執(zhí)行引擎(execution engine,簡稱EE)和拷貝引擎(copy engine,簡稱CE),其中,EE 由兩個流處理器簇(SM)構(gòu)成,每個SM 有128 個SP 核@1.3GHz,并共享512 KB的L2 cache.所有的CPU 與GPU 共享8GB DRAM@1.866GHz.

表2 比較了NVIDIA 的RTX 2080 GPU(PC 平臺)與Jetson TX2 中嵌入式GPU(SoC 平臺)在計算資源、計算性能、顯存、功耗等方面的差別.RTX 2080 擁有46 個SM 和8GB 獨立顯存,每個SM 包含64 個SP(即CUDA核)、8 個Tensor 核用于深度學習加速、1 個RT 核用于光線處理加速.相比而言,TX2 SoC 平臺中GPU 的計算資源要少得多,只有256 個SP 核,并且沒有獨立顯存.由于RTX 2080 的SP 的數(shù)量比Jetson TX2 高出一個數(shù)量級,所以其全精度計算性能也高出后者一個數(shù)量級,達到10.1 TFLOPs.不過,TX2 中GPU 的功率為7.5～15 瓦,遠小于RTX 2080 GPU 的功率(215 瓦).

Table 2 Comparing RTX 2080 GPU and Jetson TX2 embedded GPU表2 比較RTX 2080 GPU 與Jetson TX2 嵌入式GPU

1.4 神經(jīng)網(wǎng)絡加速器

近10 年以來,深度學習的飛速發(fā)展得益于CPU 和GPU 算力的巨大進步.反過來,DNN 計算時的高負載特性,也促使CPU 和GPU 改進其體系結(jié)構(gòu)以進一步提高計算效率.例如,Intel 最新的Xeon Phi Knights Mill 處理器增加了對浮點運算可變精度的支持,NIVIDA 最新的Volte 架構(gòu)增加了專門的Tensor Core 用于快速處理DNN任務中的被密集調(diào)用的矩陣乘與累加(multiply and accumulate,簡稱MAC)操作.不過,這兩種通用計算芯片所包含的部分功能模塊(如多級緩存、分支處理、內(nèi)存管理、線程調(diào)度等)在處理DNN 任務時并不會被用到,反而占用芯片的面積,這就限制了在處理DNN 任務時每單位芯片面積計算性能與能效比的提升[38].于是,人們研發(fā)了專用于處理深度學習任務的神經(jīng)網(wǎng)絡加速器,包括基于DSP 或FPGA 兩種通用型芯片改進而來的加速器,還有TPU(tensor processing unit),NPU(neural network processing unit)等采用為深度學習任務而定制體系架構(gòu)的專用神經(jīng)網(wǎng)絡加速器.這里僅概述各加速器的主要特性和優(yōu)缺點,詳細信息可以參考文獻[38].

DSP(digial signal processor)即數(shù)字信號處理芯片,具有強大的浮點運算能力.DSP 芯片通過在架構(gòu)中加入神經(jīng)網(wǎng)絡處理核心成為了一種實現(xiàn)神經(jīng)網(wǎng)絡加速的可編程可擴展的有效平臺.主流的DSP 加速產(chǎn)品包括EV6x(Synopsys),VIP8000(VeriSilicon)等,它們都支持當前主流的網(wǎng)絡模型和框架.FPGA(field programmable gate array)即現(xiàn)場可編程門陣列芯片,能夠?qū)崿F(xiàn)快速的硬件功能驗證和評估,從而加快設計的迭代速度.FPGA 憑借可編程陣列的特性,利用VHDL 或Verilog 語言可以設計新的硬件結(jié)構(gòu),更好地匹配神經(jīng)網(wǎng)絡的計算特點,比如針對CNN 卷積層硬件實現(xiàn)各類優(yōu)化的卷積方法[39,40].FPGA 在能耗方面更優(yōu)于GPU,但浮點運算性能遜于GPU,因為目前的DNN 計算還是重度依賴密集的浮點矩陣乘法(general matrix multiplication,簡稱GEMM),而GEMM 更利于映射到GPU 上(常規(guī)并行性).

相比通用芯片,專用神經(jīng)網(wǎng)絡加速芯片(TPU,NPU 等)在處理DNN 計算時能耗更低,性能更高(芯片的面積和功耗僅為GPU 的百分之一量級).專用加速器片內(nèi)主要包括針對DNN 計算特點而特別設計的運算單元和存儲單元,并且一般會設計專用的復雜指令集(complex instruction set,簡稱CISC)來處理復雜的神經(jīng)元操作,提升了速度,同時減少了功耗.運算單元有兩種主流的結(jié)構(gòu):(1) 樹狀結(jié)構(gòu),包括 DianNao/DaDianNao/PuDianNao NPU(寒武紀)等;(2) 脈動陣列結(jié)構(gòu),包括TPU(Google)、Scaledeep(普渡大學)、Eyeriss(MIT)、ShiDianNao(寒武紀)、昇騰Atlas NPU(華為海思)等.存儲單元用來存儲神經(jīng)網(wǎng)絡每一層的輸入值、權(quán)重和輸出值等參數(shù).隨著神經(jīng)網(wǎng)絡規(guī)模變大,傳統(tǒng)2D 存儲結(jié)構(gòu)的DDR 技術(shù)已經(jīng)不能適應高帶寬的要求,研究者已開始把最新的3D 存儲技術(shù)引入到加速器設計中[41,42].同時,為了最小化數(shù)據(jù)搬移,緩解帶寬瓶頸,運算器和存儲器正朝著一體化的方向發(fā)展(憶阻器[43]).除此以外,專用神經(jīng)網(wǎng)絡加速器未來還將支持更多的神經(jīng)網(wǎng)絡框架,并采用非馮·諾依曼體系架構(gòu)來大幅度提高整體運算性能、降低功耗.

2 DNN 任務在實時嵌入式系統(tǒng)中面臨的挑戰(zhàn)

嵌入式系統(tǒng)的共性特征是“實時性”,從應用的角度,實時性是指“不僅要保證運算的邏輯正確性,還必須保證運算在規(guī)定的時間內(nèi)完成,否則將產(chǎn)生災難性后果”.DNN 任務(主要是推理過程)在實時嵌入式系統(tǒng)上的成功部署與運行,既要在功能層面保證DNN 推理結(jié)果的正確性和精確度,又要在非功能層面確保滿足系統(tǒng)的實時性、資源和功耗的要求.本節(jié)將從DNN 模型、深度學習框架以及硬件計算平臺等方面總結(jié)和剖析DNN 任務應用于實時嵌入式系統(tǒng)帶來的問題和挑戰(zhàn).首先從DNN 模型性能層面討論在資源受限的實時嵌入式系統(tǒng)上部署復雜DNN 任務所面臨的問題,第3 節(jié)針對這些問題從網(wǎng)絡的性能分析、網(wǎng)絡輕量化設計、實時神經(jīng)網(wǎng)絡方面做了詳細的調(diào)研;然后討論了計算平臺(主要是GPU、SoC、操作系統(tǒng))保障DNN 任務的實時性方面所面臨的挑戰(zhàn),第4 節(jié)針對這些問題從GPU 時間分析、任務調(diào)度、資源管理策略方面做了詳細的調(diào)研.

2.1 DNN模型帶來的關(guān)鍵問題

(1) 傳統(tǒng)DNN 在實時嵌入式系統(tǒng)中具有局限性

當今,嵌入式信息物理融合系統(tǒng)正與深度神經(jīng)網(wǎng)絡領(lǐng)域互相融合,并且未來有朝著增強自主性方向發(fā)展的趨勢.DNN 應用部署在高性能硬件平臺(例如GPU 集群)上,并且具有良好的可擴展性[19].如果把現(xiàn)有的網(wǎng)絡模型部署在資源受限的實時嵌入式系統(tǒng)中,這將很難滿足時序要求[21].其主要原因是受到DNN 模型復雜度、計算和內(nèi)存資源的巨大需求等約束,DNN 在低功耗的終端設備上運行,無法滿足系統(tǒng)的實時性和可預測性.因此,在計算資源與時間方面,DNN 任務本身的需求與實際目標硬件平臺的能力存在著較大的差異.在硬件層面,由于硬件的發(fā)展和更新速度明顯落后于軟件方面,僅僅通過提升硬件性能來解決該問題仍然具有局限性.在軟件層面,探索輕量化網(wǎng)絡的設計和優(yōu)化也是一種有效的解決方案,即:在DNN 網(wǎng)絡的精度與時間、資源等方面進行權(quán)衡,綜合考慮系統(tǒng)的功能層面要求與非功能層面要求.

(2) 通用深度學習框架的設計并未考慮嵌入式平臺實時性要求

當前流行的深度學習框架有TensorFlow,Caffe,CNTK,Theano,Torch 等,這些框架各具特色,基于這些深度學習框架可以較為容易地實現(xiàn)各種前沿的深度學習算法.深度學習框架的設計面向高性能計算,是為了提高任務吞吐量和系統(tǒng)可擴展性.然而,深度學習框架所實現(xiàn)的數(shù)據(jù)或任務并行(data or task parallelism)加速并未充分考慮到底層硬件,并未考慮嵌入式平臺實時性要求[44](如資源分配時間有上界等).另一方面,不同的深度學習框架之間存在著性能差異,并且同一網(wǎng)絡在不同框架上實現(xiàn)也存在著明顯的性能差異,所以從深度學習框架角度優(yōu)化網(wǎng)絡運行時間,對于確保深度學習任務的實時性具有較大的提升空間.Kim 等學者[45]通過實驗的手段分析了AlexNet 在上述5 種流行的框架上性能差異,只是通過實驗的手段簡單地揭示了不同框架之間存在性能差異的現(xiàn)象,但卻沒有從本質(zhì)上揭露造成性能差異的原因.

(3) DNN 模型結(jié)構(gòu)更新后不能保證系統(tǒng)的實時性

深度學習模型是基于設計的網(wǎng)絡結(jié)構(gòu),通過大規(guī)模數(shù)據(jù)進行訓練得到的網(wǎng)絡參數(shù),且滿足樣本分類和預測的精度.為了保證網(wǎng)絡模型的預測精度和系統(tǒng)性能,網(wǎng)絡模型往往需要優(yōu)化,需要對擴充后的樣本數(shù)據(jù)重新訓練,更新后的網(wǎng)絡模型會修改模型參數(shù),甚至會更新網(wǎng)絡結(jié)構(gòu).然而,當網(wǎng)絡模型結(jié)構(gòu)發(fā)生變化后,任務重新部署在資源受限的實時嵌入式系統(tǒng)上,為了滿足系統(tǒng)實時性,其網(wǎng)絡模型更新過程需要滿足一系列的約束[46].系統(tǒng)更新時,既要保證新的網(wǎng)絡模型不會干擾現(xiàn)有的系統(tǒng),又要確保非功能的正確性.換句話說,計算平臺要有足夠的計算資源來運行新應用程序,而不會違反任何時間限制[47].對于實時嵌入式系統(tǒng)的有限資源下,任務模型更新后需要重新驗證與分析系統(tǒng)的可信性.

2.2 硬件計算平臺所面臨的關(guān)鍵問題

實時系統(tǒng)要求在系統(tǒng)實際執(zhí)行之前對時間行為進行分析,以確保在運行時系統(tǒng)的時間約束能夠得到滿足.一般通過任務級時間行為分析[48],即最壞執(zhí)行時間(worst-case execution time,簡稱WCET)分析,與系統(tǒng)級實時調(diào)度與調(diào)度分析[49]加以保證.WCET 分析的主要功能是分析程序的執(zhí)行路徑信息以及程序?qū)τ布脑L問行為,從而求得最壞執(zhí)行時間.實時調(diào)度分析的目標是:在給定每個任務的WCET 以及系統(tǒng)的實時調(diào)度算法的情況下,利用數(shù)學手段分析系統(tǒng)中的所有任務是否能在截止期之前完成.在GPU 目標硬件平臺上部署實時DNN 任務需要解決以下6 個主要問題:

(1) 針對GPU 的WCET 分析和實時調(diào)度分析尚不成熟

CPU 和GPU 在體系結(jié)構(gòu)上存在本質(zhì)不同,導致CPU 程序與GPU 程序的執(zhí)行特性、訪存特性差別很大,這對程序的時間行為有巨大影響[50].現(xiàn)有的實時系統(tǒng)研究大多針對于CPU,相關(guān)技術(shù)不能直接用于GPU 程序的分析.現(xiàn)有的GPU 性能分析與優(yōu)化相關(guān)研究主要考慮系統(tǒng)的平均性能,而非實時性能(即最壞情況下的性能)與可調(diào)度性,因此也無法滿足實時系統(tǒng)的要求.相對于GPU 硬件的快速發(fā)展,面向GPU 的實時系統(tǒng)時間分析與優(yōu)化已嚴重滯后,這成為了GPU 面向?qū)崟r嵌入式系統(tǒng)應用的主要障礙.

(2) GPU 調(diào)度機制的細節(jié)信息不公開

文獻[51]指出:當前主流的GPU 生產(chǎn)廠商(例如NVIDIA)出于商用機密的原因不公開GPU 內(nèi)部調(diào)度器的調(diào)度邏輯,客觀上對實時GPU 任務負載的可調(diào)度性分析造成了障礙.而且廠商有可能在沒有任何通知的情況下,在新版本的GPU 上修改調(diào)度邏輯.此外,由于GPU 廠商的主要市場是高性能計算領(lǐng)域,提高任務的吞吐量和減小執(zhí)行時間延遲是GPU 調(diào)度器要追求的目標,這種調(diào)度策略必然不適合實時系統(tǒng)中任務的執(zhí)行時間需要具有可預測性的要求.毫無疑問,嵌入式GPU 調(diào)度策略的透明性給自動駕駛的安全性埋下了隱患.

(3) GPU 上執(zhí)行混合關(guān)鍵任務會發(fā)生任務間干涉

GPU 上運行單個任務時,該任務獨享GPU 上的所有硬件資源,包括流處理器簇(SM)、cache、寄存器文件、內(nèi)存、PCI-E 總線等.當GPU 上執(zhí)行混合關(guān)鍵任務時,由于需要共享資源,一個任務的執(zhí)行時間將受到其他并行任務的干涉[52].干涉會增加系統(tǒng)執(zhí)行時間的不確定性,如果設計不當,會造成系統(tǒng)整體執(zhí)行性能的嚴重下降.同時,干涉會導致任務的狀態(tài)空間隨并行任務的增加呈現(xiàn)指數(shù)爆炸,這為任務級時間分析帶來了巨大的挑戰(zhàn).還需要注意一點:一個GPU 程序往往包含若干個kernel 函數(shù),在運行時會為每個kernel 函數(shù)分配線程塊(第1.3 節(jié)).線程塊是GPU 資源分配的單位,線程束是GPU 調(diào)度和執(zhí)行的單位.這是比GPU 任務更加細粒度的執(zhí)行元素,增加了分析與調(diào)度的難度.

(4) GPU 不支持對混合關(guān)鍵任務的搶占

早期的NIVDIA GPU(如Tesla,Fermi,Kepler,Maxwell 架構(gòu))不提供搶占機制.Pascal 架構(gòu)開始為圖像任務和計算任務提供像素級(pixel-level)搶占和線程級(thread-level)搶占支持[53],但這是針對單GPU 程序使用場景,目的是提高單任務的實時性.但是GPU 調(diào)度器所提供的搶占粒度太小,不足以在任務級實現(xiàn)支持混合關(guān)鍵任務的動態(tài)優(yōu)先級調(diào)度如EDF(earliest-deadline first,最早時限優(yōu)先)等.

(5) CPU+GPU SoC 平臺上存在CPU 和GPU 之間協(xié)同的不確定性

在NVIDIA Jetson TX2 嵌入式SoC 平臺中,CPU 和GPU 共享了8GB 的DDR 內(nèi)存(圖3).文獻[54]指出,CPU與GPU 在訪問共享內(nèi)存可能會發(fā)生沖突造成CPU 任務和GPU 任務訪存延遲的不確定性.由于GPU 沒有中斷線路,所以響應CPU 通知事件的唯一方式是通過CUDA 的同步機制.文獻[51]指出:CPU 與GPU 之間的同步操作存在時間不確定性,同步操作完成的時間取決于在GPU 上同時運行的其他任務,從而造成GPU 任務的響應時間無確定上界.除了CUDA 的顯式同步操作(如cudaDeviceSynchronize,cudaStreamSynchronize)完成時間的不確定性,一些共享內(nèi)存訪問的操作也會引發(fā)隱式同步行為,更增加了時間不確定性和分析的難度.

(6) GPU 工作溫度過高會自動觸發(fā)GPU 降頻運行

NVIDIA 開發(fā)者社區(qū)的一篇技術(shù)文章[55]指出:在不做任何設置的情況下,NVIDIA GPU 的工作頻率會隨著GPU 的工作溫度變化、能耗管理策略自動調(diào)整GPU 工作頻率,并呈現(xiàn)頻繁的波動變化,導致同樣的GPU 程序在不同時間段運行所測量的運行時間會有較顯著的差異.通過GPU 驅(qū)動所提供的接口可鎖定GPU 的運行頻率為GPU BIOS(basic input output system)中的預設頻率,從而讓GPU 程序的運行時間保持恒定.但是當GPU 溫度超過最高工作溫度閾值時,GPU 的硬件保護機制仍然會自動強制GPU 降低工作頻率.因此,當在GPU 上運行實時任務時,需要在初始化時鎖定GPU 的工作頻率.在散熱能力有限的情況下,要特別注意GPU 任務負載的調(diào)度,避免GPU 長時間處于高負荷工作狀態(tài)而產(chǎn)出高熱,并引發(fā)GPU 自動降頻.GPU 的自動降頻可能會導致實時系統(tǒng)調(diào)度嚴重失效,引發(fā)災難性后果.

3 DNN 任務性能分析與優(yōu)化

DNN 任務性能分析與優(yōu)化的重點在于分析DNN 任務的性能特性,基于系統(tǒng)瓶頸提出相關(guān)優(yōu)化策略.現(xiàn)有研究工作旨在提升網(wǎng)絡執(zhí)行的平均性能.但遺憾的是,當前研究工作并沒有把網(wǎng)絡實時性納入考慮范疇.在實時系統(tǒng)中僅考慮平均性能而不考慮實時性能,是無法滿足安全攸關(guān)實時系統(tǒng)的要求的,這是實時系統(tǒng)中亟待解決的問題,也是未來研究的重點和難點.本節(jié)就DNN 任務平均性能方面,首先總結(jié)前人在DNN 任務性能分析方面的主要工作,綜合闡述了DNN 任務性能分析方法以及現(xiàn)有系統(tǒng)的瓶頸;然后,從輕量化網(wǎng)絡角度闡述了DNN 網(wǎng)絡的性能優(yōu)化工作;最后,在DNN 任務實時性能方面,整理了當前實時神經(jīng)網(wǎng)絡的最新進展與成果.

3.1 DNN任務的性能分析

程序的性能分析(performance analysis)指的是通過靜態(tài)/動態(tài)程序分析方法計算/收集程序運行時的性能數(shù)據(jù),用于分析程序的負載特性和性能瓶頸,用戶(程序開發(fā)者、編譯器或者硬件計算平臺)根據(jù)性能分析提供的反饋信息綜合優(yōu)化系統(tǒng)整體性能.GPU 作為主要的硬件加速設備,已經(jīng)被廣泛的應用于深度學習領(lǐng)域.DNN 任務在GPU 上運行時會遇到各種性能障礙,例如低效的DNN 算法、低精度的結(jié)果、并行的開銷、負載的不均衡、內(nèi)存緩存和帶寬的低效使用、計算延遲超出響應時間要求等.為了識別DNN 任務在GPU 上執(zhí)行時的負載特性和性能瓶頸,需要進行性能分析.目前,對基于GPU 的DNN 任務進行性能分析可以概括為3 個發(fā)展階段:

性能分析的第1 階段是借助性能分析工具,基于實驗測試的方式獲取DNN 任務的性能數(shù)據(jù).基于特定的GPU 目標硬件平臺,選取特定的深度學習框架或DNN 算法,采用實驗的手段,結(jié)合NVIDIA 官方提供的性能分析工具(例如nvprof,NVVP,Nsight)讀取GPU 上有限的硬件計數(shù)器獲取程序運行時間、資源利用等性能數(shù)據(jù).Kim 等學者[45]通過實驗的手段分析了同一種DNN 網(wǎng)絡模型在5 種流行的深度學習框架(Tensorflow,Caffe,Torch,CNTK,Theano)上的性能差異,他們基于DNN 模型的每一層(layer)的前向/后向傳播進行“端到端”的測量,獲取每一層操作的執(zhí)行時間和資源占用情況,識別每種框架下的DNN 任務的“Hot Layer”(指計算量大、耗時長、消耗資源多的層).此外,Shams 與Platania 等學者在文獻[56]中,通過對不同的高性能計算硬件環(huán)境(NVLink,Knsight Landing)進行測量評估,比較了Caffe,TensorFlow 與Apache SINGA 在運行速度與可擴展性等方面的性能差異.雖然該研究工作揭示了不同深度學習框架或DNN 算法之間存在著性能差異的客觀事實,但是卻沒有揭示導致差異的本質(zhì)原因.為探究DNN 任務在特定GPU 硬件平臺上的工作性能,Mojumder 等學者基于NVIDIA DGX-1 多GPU 系統(tǒng)開展了一些研究工作[57].他們選取了GoogleNet,AlexNet 等圖像分類應用中5 種典型的網(wǎng)絡模型,在DGX-1 系統(tǒng)中進行訓練,對網(wǎng)絡訓練過程中的各個階段的訓練時間進行獲取和分析,探究DNN 模型在訓練階段的時間行為特性.該階段的研究工作旨在通過基礎的實驗測量手段獲取性能數(shù)據(jù),用于識別DNN 的負載特性和性能瓶頸,輔助程序開發(fā)人員或者DNN 模型的設計人員設計出高效的DNN 的訓練和推理方法.但由于實驗手段與性能工具的局限性,獲取得到的性能數(shù)據(jù)和性能分析結(jié)果也存在著一定的局限性.

性能分析的第2 階段是采用程序插樁的方法獲取細粒度的性能數(shù)據(jù).基于GPU 硬件平臺,結(jié)合編譯器和程序插樁方法,獲取細粒度的性能數(shù)據(jù),并基于分析結(jié)果進行DNN 任務的優(yōu)化工作.Stephenson 等學者在文獻[58]中提出了一種基于后臺編譯的插樁工具SASSI,該工具允許用戶指定用于插樁的指令代碼,自定義插樁位置和插樁內(nèi)容,相對于其他插樁工具更加靈活.SASSI 內(nèi)置于NVIDIA GPU 的編譯器中,對PTX 代碼進行插樁,實現(xiàn)對待檢測程序的控制流行為、訪存行為以及程序執(zhí)行過程中的寄存器狀態(tài)進行性能數(shù)據(jù)獲取和分析.Du Shen等學者提出了CUDAAdvisor 性能分析框架[59],旨在對現(xiàn)代GPU 上的代碼優(yōu)化提供指導作用.該工具利用LLVM 框架對CUDA 程序代碼進行插樁,在程序的執(zhí)行過程中收集性能數(shù)據(jù),包括CPU 與GPU 兩者之間的交互作用,基于性能數(shù)據(jù)進行訪存行為分析、控制流分析以及代碼調(diào)試.與其他工具不同的是,CUDAAdvisor 具有較好的擴展性,能夠支持不同版本的CUDA 與不同架構(gòu)的GPU.Farooqui 等學者[60]針對于GPU 并行平臺的性能檢測和分析問題,提出并開發(fā)了動態(tài)插樁系統(tǒng)Lynx,基于GPU Ocelot 平臺,使用JIT 編譯器對待檢測程序在PTX 代碼級別進行轉(zhuǎn)化、插樁與優(yōu)化操作,具有較高的效率,能夠獲取kernel 函數(shù)的執(zhí)行時間、branch divergence情況以及不同粒度(例如kernel,thread block 等)的性能數(shù)據(jù).雖然針對GPU 的性能分析問題,插樁是一種有效的技術(shù)手段,可以在程序源碼級別、中間代碼級別以及目標代碼級別實施插樁,通過插樁代碼來檢測程序運行時的性能數(shù)據(jù),但是該方法會引入額外的負載,需要對引入的負載進行評估.

性能分析的第3 階段是結(jié)合DNN 任務特性和GPU 目標硬件平臺特性構(gòu)建性能模型,然后采用數(shù)學的方法分析和評估DNN 任務的性能,識別性能瓶頸.Qi 等學者在文獻[61]中提出了一種可分析的性能模型PALEO,可以對簡單的、具有特定聲明規(guī)范的神經(jīng)網(wǎng)絡架構(gòu)進行性能評估.PALEO 從指定的架構(gòu)中提取需求并將其映射到軟件、硬件和通信策略等設計空間中,將神經(jīng)網(wǎng)絡建模成一個有向圖,利用數(shù)學方法計算網(wǎng)絡的執(zhí)行時間.由于DNN 任務與GPU 硬件兩者的復雜性,構(gòu)建性能模型的方法存在較大的難度.目前,相關(guān)的研究工作還未取得突破性進展,是未來進一步的研究方向.此外,還有一些學者進行了其他方面的探索.Dong Shi 在文獻[62]中探究了卷積神經(jīng)網(wǎng)絡的微結(jié)構(gòu)對DNN 任務在GPU 上執(zhí)行所造成的影響,他們使用微架構(gòu)設計逐層分析CNN 模型的性能影響,并在典型的GPU 存儲架構(gòu)中描述訪存行為特性以及硬件資源的利用情況,識別單個GPU 的潛在性能瓶頸.Madougou 等學者在文獻[63]中提出了一個GPU 應用程序的性能分析工具,該工具綜合采用了隨機森林的統(tǒng)計方法、GPU 硬件計數(shù)器以及機器學習的方法,構(gòu)建了一個用于性能預測的隨機森林模型.

利用上述的分析方法或工具,可以對GPU 上DNN 模型的訓練和推理過程進行性能分析,并得出以下結(jié)論.

(1) 不同的深度學習框架會導致較大的性能差異.在相同的硬件配置條件下,同一個Alexnet 模型在主流深度學習框架下的訓練時間存在著較大的差異:在CNTK 中訓練最快,在Theano 中訓練最慢.不同的深度學習框架實現(xiàn)機制與CUDA 驅(qū)動的融合程度不同,所以其性能具有較大差異.因此在實際的應用中,針對指定的DNN 模型,可以選擇更加匹配的深度學習框架,提升模型訓練和推理的性能;

(2) 卷積層運算涉及到大量的密集型計算,需要大量的計算資源,占據(jù)網(wǎng)絡模型全部訓練/推理時間的大部分.一個完整的DNN 網(wǎng)絡模型包含多個層,例如數(shù)據(jù)層、卷積層、池化層、激勵層、全連接層,每一層都執(zhí)行特定的操作.其中,卷積層對輸入數(shù)據(jù)進行卷積運算,并映射到特征空間中,是DNN 模型中“Hot Layer”,提高卷積層的性能,可以大幅度的提高系統(tǒng)性能;

(3) 在實際應用中,DNN 模型包含大量的神經(jīng)元,具有數(shù)千萬的訓練參數(shù),計算量大,存在兩個潛在的性能瓶頸:計算瓶頸和通信瓶頸.一方面,由于GPU 等硬件加速平臺計算資源是有限的,DNN 模型訓練以及推理過程中,大量的數(shù)據(jù)計算就容易在硬件平臺上造成性能瓶頸;另一方面,隨著DNN 網(wǎng)絡規(guī)模的不斷增大,為了加快網(wǎng)絡的訓練速度,經(jīng)常需要在分布式CPU/GPU 集群中完成整個訓練.在主流的基于數(shù)據(jù)并行的分布式深度學習中,各個計算單元并發(fā)訓練一個DNN 模型,各個單元的訓練數(shù)據(jù)不同,每一次迭代結(jié)束后,各個計算單元需要同步DNN 參數(shù)或梯度,更新參數(shù)服務器上的模型,之后再將最新模型推送到各個數(shù)據(jù)并行單元,用于下一輪計算.因此,數(shù)據(jù)并行中參數(shù)的交換容易成為整體性能瓶頸,如何解決這種通信瓶頸以提高性能,成為并行方法設計的研究重點.

3.2 輕量化神經(jīng)網(wǎng)絡

本節(jié)陳述的DNN 網(wǎng)絡層性能優(yōu)化相關(guān)研究工作主要是提高網(wǎng)絡的平均性能,而非實時性能.基于上述DNN 性能分析結(jié)論,DNN 任務性能優(yōu)化方向是減少模型中冗余參數(shù),優(yōu)化卷積核的計算結(jié)構(gòu).設計輕量化的神經(jīng)網(wǎng)絡方法包括人工手動優(yōu)化DNN 模型和自動化設計方式優(yōu)化DNN 模型:手動優(yōu)化DNN 模型包括深度學習模型壓縮、設計輕量化網(wǎng)絡、優(yōu)化卷積計算方式和網(wǎng)絡中卷積層池化層搭配的網(wǎng)絡結(jié)構(gòu)設計;自動化設計方式是指采用神經(jīng)網(wǎng)絡架構(gòu)搜索方式,基于多目標優(yōu)化約束自動化設計網(wǎng)絡.

(1) 深度學習模型壓縮技術(shù)

將現(xiàn)有神經(jīng)網(wǎng)絡模型部署在資源受限的實時嵌入式系統(tǒng)的主要挑戰(zhàn)之一是同時滿足嚴格的運行時限要求和受限的SWaP(size,weight,and power)要求[64],該問題尚未得到有效解決.近年來,許多研究機構(gòu)從軟件層面研究復雜深度學習網(wǎng)絡的精簡問題,提出了一些輕量化的神經(jīng)網(wǎng)絡模型.DNN 輕量化技術(shù)是指采用模型的近似化技術(shù)保證網(wǎng)絡輸出精度不損失或損失在可接受范圍內(nèi)的情況下,使得網(wǎng)絡結(jié)構(gòu)和網(wǎng)絡參數(shù)盡可能的精簡,實現(xiàn)系統(tǒng)模型的輕量化[65].網(wǎng)絡模型近似操作常采用的3 種方式是:模型壓縮、剪枝、矩陣分解.經(jīng)過模型近似技術(shù)處理過的模型稱為近似模型,該模型在許多實際環(huán)境中已被證明表現(xiàn)出足夠的精度[66,67].

模型近似操作的目標是為了探索更高效的基礎架構(gòu).已有的工作聚焦于剪枝、壓縮以及低比特方法來表示基本的網(wǎng)絡架構(gòu)[68-71].Denil 等學者[72]證明了深度學習模型的參數(shù)存在明顯的冗余,并且參數(shù)值可以通過預測得到.在MNIST 數(shù)據(jù)集上測試,該方法在最好情況下可以預測出多層神經(jīng)網(wǎng)絡的95%的參數(shù),同時驗證了網(wǎng)絡模型參數(shù)的冗余性.Han 等學者[17]提出使用10%的精度關(guān)鍵數(shù)據(jù)快速訓練DNN,可以實現(xiàn)網(wǎng)絡98%的精度.這個結(jié)論再次證明了DNN 網(wǎng)絡存在一定的冗余性,說明了DNN 模型壓縮技術(shù)具有一定的可行性.Han 等學者[73,74]嘗試將訓練好的模型通過剪枝和權(quán)重共享增加權(quán)重稀疏性、降低存儲量的方法,使得神經(jīng)網(wǎng)絡的存儲需求減少35 倍～49 倍卻不損失網(wǎng)絡的準確性.Jaderberg 等學者[69]嘗試使用降秩技術(shù)來加速卷積神經(jīng)網(wǎng)絡,并且在場景字符分類數(shù)據(jù)集上訓練的4 層CNN 精度下降不到1%,卻獲得了4.5 倍的加速.Zhang 等學者[70]提出了一種多層逼近時減小累積誤差的算法,將非線性響應的重構(gòu)誤差降至最小,并采用低秩約束來降低了濾波器的復雜度.ImageNet 數(shù)據(jù)集上的標準測試結(jié)果表明:通過該算法壓縮SPPNet 得到的加速模型,比原SPPNet 的推理速度快了4 倍.雖然加速模型的top-5 誤差率比原模型增加了0.9%,但是其推理精度相比AlexNet 還是提高了4.7%.

在DNN 網(wǎng)絡參數(shù)的線性量化方面,二值化神經(jīng)網(wǎng)絡(binary neural network,簡稱BNN)[75]具有高模型壓縮率和極快計算速度的優(yōu)點,近幾年格外受到重視,成為深度學習的熱門研究方向.二值化網(wǎng)絡方法通過將單精度浮點型權(quán)重矩陣二值化,其中一個權(quán)重值只用一個比特來表示.二值化方法根據(jù)權(quán)重值和激活函數(shù)值的符號決定在{+1,-1}取值:數(shù)值大于等于0 則取+1,否則取為-1.對于原來32 位浮點型數(shù)的乘加運算,二值化之后的權(quán)重值和激活函數(shù)值可以通過一次異或運算(xnor)和一次POPCNT 位運算指令實現(xiàn).二值化方法在網(wǎng)絡模型的內(nèi)存消耗理論上能減少32 倍,可以看出,二值化神經(jīng)網(wǎng)絡在模型壓縮上具有很大的優(yōu)勢.研究二值化神經(jīng)網(wǎng)絡對解決當前浮點型神經(jīng)網(wǎng)絡應用到實時嵌入式系統(tǒng)中存在的模型過大、計算密度過高等問題,具有很重大的意義.

(2) 神經(jīng)網(wǎng)絡輕量化設計

然而,上述在已訓練模型上進行因式分解和稀疏卷積技術(shù)沒有從根本上實現(xiàn)網(wǎng)絡的高效輕量化.研究者們從網(wǎng)絡結(jié)構(gòu)設計出發(fā),設計出了適用于移動設備的輕量級神經(jīng)網(wǎng)絡.表3 中,最早于2016 年2 月,加州大學伯克利分校和斯坦福大學提出輕量化網(wǎng)絡SqueezeNet[65]模型.該模型具有新的網(wǎng)絡結(jié)構(gòu),通過降低大型卷積核的輸入通道數(shù)量,新的網(wǎng)絡結(jié)構(gòu)實現(xiàn)了不同大小卷積核的通道連接來特征提取.同年10 月,Google 提出了Xception模型.2017 年4 月,Google 提出的MobileNet[76]模型具有一種新穎的卷積結(jié)構(gòu),其特點是在保證特征的非線性表示情況下,深度分離卷積式的卷積結(jié)構(gòu),充分解耦了傳統(tǒng)卷積模型,將卷積分為深度卷積和逐點卷積.這種卷積改進方式可以極大降低計算量和參數(shù)量,并且適配于移動設備.2018 年,Google 和 Face++分別相繼發(fā)布MobileNetV2 和ShuffleNetV2,兩個輕量化網(wǎng)絡均是在之前模型的改進模型.輕量化網(wǎng)絡極大地降低了深度學習模型的參數(shù)量和計算量,同時,在很大程度上保證了特征的高度抽象提取,在一定程度上保證了模型的精確度.

Table 3 Lightweight network development timeline表3 輕量化網(wǎng)絡發(fā)展時間軸

實現(xiàn)輕量化網(wǎng)絡的技術(shù)方法主要包括:利用卷積核分解方法使用1×N網(wǎng)絡和N×1 網(wǎng)絡代替N×N卷積核;使用深度壓縮(deep compression)方法,包括網(wǎng)絡剪枝、量化、哈弗曼編碼、奇異值分解、硬件加速器等方法.因此,以SqueezeNet 為例,其設計使用以下3 個策略來減少參數(shù):

(a) 使用1×1 卷積代替3×3 卷積,該設計使得計算參數(shù)減少為原來的1/9;

(b) 減少輸入通道數(shù)量,該部分使用squeeze 層來實現(xiàn);

(c) 在網(wǎng)絡中,將欠采樣(downsample)操作延后,可以給卷積層提供更大的激活圖(activation maps).因為更大的激活圖保留了更多的信息,可以提供更高的分類準確率.

其中,策略(a)和策略(b)可以顯著減少參數(shù)數(shù)量,策略(c)可以在參數(shù)數(shù)量受限的情況下提高準確率.輕量化網(wǎng)絡與傳統(tǒng)DNN 網(wǎng)絡相比具有諸多優(yōu)點[65],網(wǎng)絡模型的精簡也有助于整個系統(tǒng)的優(yōu)化.比如:在分布式訓練中,輕量化網(wǎng)絡與服務器通訊需求較小,由于網(wǎng)絡模型的參數(shù)少,從云端下載模型也更快.

當然,在使用輕量化技術(shù)縮短網(wǎng)絡推理時間,提升系統(tǒng)性能的同時,人們還必須保證網(wǎng)絡的預測精度足夠高.優(yōu)秀的代表性網(wǎng)絡包括SqueezeNet 和SqueezeNext 等.SqueezeNet 模型和Alexnet 模型在ImageNet 數(shù)據(jù)集上的推理精度相當,但是前者比后者的參數(shù)數(shù)量減少了50 個,并且SqueezeNet 模型的尺寸小于0.5MB.隨后,Gholami 等學者總結(jié)了已有輕量化網(wǎng)絡的結(jié)構(gòu)優(yōu)點,并根據(jù)神經(jīng)網(wǎng)絡加速器上的仿真結(jié)果作為指導,提出新的神經(jīng)網(wǎng)絡體系結(jié)構(gòu)SqueezeNext[81].它采用濾波器降秩(low rank filters)的方法進一步減少權(quán)值參數(shù),并且采用瓶頸模塊(bottleneck module)減少全連接層參數(shù).有別于上述的單從軟件層對網(wǎng)絡模型進行優(yōu)化,SqueezeNext是從網(wǎng)絡設計和硬件實現(xiàn)兩個角度綜合優(yōu)化網(wǎng)絡,其思想很值得借鑒.與MobileNet 相比,SqueezeNext 在獲得相似的top-5 的分類精度下,其參數(shù)減少了1.3 倍.

(3) 自動機器學習

為嵌入式系統(tǒng)設計深度學習網(wǎng)絡,是一項具有挑戰(zhàn)性的工作,因為嵌入式平臺要求網(wǎng)絡模型不僅體積小、運行速度快,還要確保網(wǎng)絡預測的準確率.盡管研究者已經(jīng)做了許多輕量化模型的設計和改進工作,例如上述SqueezeNet,MobileNet 系列,但是手動設計高效模型仍然是一項挑戰(zhàn),因為網(wǎng)絡設計要考慮的因素太多.AutoML(automated machine learning)[82]和神經(jīng)架構(gòu)搜索(neural architecture search,簡稱NAS)[83]的發(fā)展,促進了深度學習模型的自動化設計.AutoML 是模型選擇、特征抽取和超參數(shù)調(diào)優(yōu)等一系列自動化方法,可以實現(xiàn)自動訓練有價值的模型.機器學習最耗費人力的部分主要是數(shù)據(jù)清洗和模型調(diào)參,而這部分過程如果采用自動化方式實現(xiàn),將會加快網(wǎng)絡模型的開發(fā)過程.NAS 實現(xiàn)了用神經(jīng)網(wǎng)絡設計神經(jīng)網(wǎng)絡,代表了機器學習發(fā)展的未來方向.NAS 是AutoML 的子領(lǐng)域,在超參數(shù)優(yōu)化和元學習(meta learning)等領(lǐng)域高度重疊.與以往自動化搜索方法不同,MnasNet[84]的設計綜合權(quán)衡了網(wǎng)絡輸出精度和實時性的多目標優(yōu)化問題.它是Google 提出探索使用架構(gòu)搜索和強化學習設計模型的一種方法,并且在模型準確率和運算速率上均有突破.MnasNet 模型的運行速度比MobileNet V2[79]快1.5 倍、比NASNet[83]快2.4 倍,同時達到同樣的ImageNet top-1 的準確率.

3.3 實時神經(jīng)網(wǎng)絡

對于包含了人工智能應用的實時嵌入式系統(tǒng),實時系統(tǒng)的設計除了考慮網(wǎng)絡推理任務的時間可預測性以及硬實時性,還要考慮網(wǎng)絡模型的輸出精度.如果能夠?qū)ι疃葘W習網(wǎng)絡的性能進行建模,就能分析出推理網(wǎng)絡的執(zhí)行時間,得出任務執(zhí)行時間的上界,最終設計出調(diào)度策略來保證實時任務在截止期之前正確執(zhí)行完.深度學習網(wǎng)絡的輸出結(jié)果本質(zhì)上是基于一定概率的輸出,因此,如何權(quán)衡網(wǎng)絡的預測精度與網(wǎng)絡的執(zhí)行時間,即在犧牲可接受的網(wǎng)絡精度來換取網(wǎng)絡推理過程的實時性,引起了實時領(lǐng)域研究者的關(guān)注.Bateni 等學者[19]提出近似意識的實時神經(jīng)網(wǎng)絡ApNet,對AlexNet 模型以layer 為單位應用降秩(low rank)近似方法,實驗數(shù)據(jù)顯示,以損失5%的精度,換取網(wǎng)絡層任務執(zhí)行時間減少50%～80%.時間可預測運行時系統(tǒng)ApNet 通過逐層設計有效的近似壓縮,確保了DNN 任務的硬截止期.ApNet 是建立在多層DNN 端到端理論分析的調(diào)度框架,基于每一層基礎上有效近似技術(shù)來滿足層任務的子截止期.基于權(quán)衡網(wǎng)絡計算精度和運行時間的理論分析,在NVIDIA Jetson TX2 平臺上,設計并實現(xiàn)了一個運行時系統(tǒng),提高了運行時多任務性能.

Zhou 等學者針對GPU 加速的實時DNN 任務提出監(jiān)督流調(diào)度方案S3DNN.S3DNN 綜合考慮DNN 任務的實時性能、系統(tǒng)的吞吐量和資源利用率等問題[20].通過在數(shù)據(jù)采集處理的輸入端進行并行設計,S3DNN 提出將多傳感器的數(shù)據(jù)進行融合,合成數(shù)據(jù)后,減少了DNN 任務實例數(shù)量.同時還提出GPU 的kernel 調(diào)度與并行,通過將kernel 切分為3 部分,基于最小slack 策略,設計CUDA 流調(diào)度機制實現(xiàn)kernel 并行執(zhí)行.Yang 等學者從CNN實際智能駕駛應用提出了共享CNN 網(wǎng)絡的策略,將非關(guān)鍵級的攝像機4 幀圖像壓縮合并為一張圖片共享一個Tiny YOLO 網(wǎng)絡,相當于4 路CNN 網(wǎng)絡并行;同時還提出對網(wǎng)絡進行分段處理,在多段過程中增加任務流的并行性[16].該方案在既不增加DNN 任務處理延遲又不降低網(wǎng)絡識別精度情況下,還增加了網(wǎng)絡吞吐量.雖然圖像壓縮后網(wǎng)絡精度有一定損失,但對于非關(guān)鍵級別任務,其精度在可接受范圍,并且經(jīng)過再次訓練后精度有所提升.

除此之外,在DNN 的運行框架層的實時性優(yōu)化方面,Hang 等學者[85]對Caffe 框架進行了修改,增加了內(nèi)存管理和垃圾回收機制,提升了處理速度.優(yōu)化后的Caffe 框架使得機器學習程序具有時間可預測性.面向?qū)崟r應用的DNN 網(wǎng)絡方面的工作目前比較少,在人工智能全面應用的背景下,如何確保具有深度學習網(wǎng)絡功能的實時系統(tǒng)的可靠性,是當下的研究熱點和難點.深度學習網(wǎng)絡應用的實時系統(tǒng)分析挑戰(zhàn)一方面來自網(wǎng)絡建模的復雜性,另一方面來自底層硬件加速平臺的透明性.不同于只執(zhí)行傳統(tǒng)任務的實時系統(tǒng),執(zhí)行DNN 任務的實時系統(tǒng)不僅要滿足任務的實時性約束,還必須滿足DNN 網(wǎng)絡模型預測精度的約束.總之,由于DNN 模型的存在,該系統(tǒng)的實時性和可靠性將面臨來自系統(tǒng)內(nèi)部和外部環(huán)境的雙重不確定性挑戰(zhàn).

4 GPU 時間分析與調(diào)度管理技術(shù)

隨著GPU 在嵌入式領(lǐng)域的應用愈加廣泛,特別是深度學習(推理階段)在CPU+GPU SoC 上的應用,對GPU的時間分析與調(diào)度管理策略的研究已經(jīng)越來越受到學術(shù)界的關(guān)注.但當前,GPU 軟硬件設計的目標仍然是盡可能提高軟件的并發(fā)度,最大可能榨取GPU 的運算能力,但這并不能確保系統(tǒng)滿足實時約束,例如,一個關(guān)鍵的實時任務可能頻繁受到其他非關(guān)鍵任務執(zhí)行的干擾,導致其時間行為不可預測.這些客觀事實給GPU 實時分析與調(diào)度的研究帶來了挑戰(zhàn).學術(shù)界對GPU 實時應用的研究主要包括:(1) 對GPU 的WCET 分析;(2) GPU 的調(diào)度策略;(3) 對GPU 的計算和內(nèi)存資源管理.

4.1 GPU WCET分析

在傳統(tǒng)實時系統(tǒng)研究中,時間分析通常分為兩步:首先分析每個程序的WCET;然后把所有程序的WCET 作為輸入,結(jié)合實時調(diào)度算法,分析系統(tǒng)的可調(diào)度性(即每個程序是否能夠在截止期前完成).所以,WCET 分析是可調(diào)度性分析的基礎.起初,學術(shù)界對GPU 上執(zhí)行的kernel 代碼段的WCET 分析也沿用了類似傳統(tǒng)實時系統(tǒng)的研究思路:不考慮調(diào)度和數(shù)據(jù)傳輸?shù)挠绊?并假設數(shù)據(jù)已經(jīng)存在于GPU 上,然后讓WCET 模型的分析結(jié)果盡量逼近kernel 代碼的實際執(zhí)行時間.

Berezovskyi 等學者早在2012 年之前就開始了GPU 上的WCET 分析,并提出了一個靜態(tài)分析模型[86,87].該模型根據(jù)一個SM 上的可分配Wrap 數(shù)量、kernel 程序的線程數(shù)、kernel 程序指令數(shù)以及指令集的時鐘周期表,在一些簡化假設下,把WCET 求解轉(zhuǎn)化成一個整數(shù)線性規(guī)劃(integer linear programming,簡稱ILP)問題.該方法可以在幾個小時內(nèi)分析出一個kernel 程序的WCET,具有可接受的分析效率.但是該研究存在一些問題:首先,該方法假設所有的數(shù)據(jù)已經(jīng)在GPU 上,忽略了數(shù)據(jù)在系統(tǒng)中的搬運過程;其次,該方法只能分析一個SM 上的一個GPU kernel 程序,而顯然,實際系統(tǒng)中一個SM 上會運行多個kernel,多個SM 上的大量kernel 也是并行運行的.

此后,Berezovskyi 等學者轉(zhuǎn)向用基于測量的動態(tài)方法[88]分析GPU 程序的WCET.動態(tài)方法本質(zhì)上要實際執(zhí)行被分析的kernel 程序,因此該方法能夠考慮程序?qū)Ω骷壌鎯Y(jié)構(gòu)的訪問延時,并適用于有多SM 的GPU.實際上,在只有單個kernel 任務執(zhí)行的情況下,動態(tài)分析方法獲得的kernel 代碼執(zhí)行時間與靜態(tài)分析出的WCET 偏差并不大.原因是:其一,kernel 程序遵循SIMT 模型,其各個線程的訪存模式和代碼執(zhí)行模式是非常規(guī)律的;其二,GPU 調(diào)度器為每個SM 預分配足夠的資源以后,SM 在執(zhí)行時是完全隔離的.但是,動態(tài)方法的問題仍然不夠安全,無法用于硬實時系統(tǒng).此外,動態(tài)方法難以深入剖析程序內(nèi)部的行為特征,因此難以得到有效信息,指導程序性能優(yōu)化.

Betts 與Donaldson 提出一種混合分析技術(shù)[89]來分析GPU kernel 程序的WCET.該方法利用控制流程圖(control flow graph,簡稱CFG)來建模kernel 程序的行為,對GPU 的硬件調(diào)度器分配線程塊的行為進行建模,利用kernel 的執(zhí)行路徑(execution traces)來幫助分析多個wrap 之間的干涉,并估算干涉對執(zhí)行時間產(chǎn)生的影響.干涉的開銷會累積并沿著執(zhí)行路徑傳導,最終獲得SM 上最后一個wrap 的執(zhí)行時間.遺憾的是,實驗表明:這種干涉分析方法非常不精確,分析得到的WCET 過于悲觀.Punniyamurthy 等學者提出了基于混合模型模擬的GPU kernel 運行時間分析技術(shù)[90],并基于開源GPU 模擬器GPGPU-Sim[91,92]開發(fā)了抽象時間仿真工具GATSim.GATSim 使用功能模型(function model)與時間模型(time model)來快速模擬GPU 程序在SM 上的動態(tài)運行情況,并報告較準確的預測運行時間.其中,功能模型描述了wrap 對應的指令序列塊和每條指令的執(zhí)行周期.時間模型描述GPU 對負載的映射與調(diào)度、SM 上wrap 的并行與干涉,從而實現(xiàn)對指令序列塊動態(tài)運行狀況的模擬.實驗表明:GATSim的平均預測準確率達96%,平均運算速度高于80MIPS(million instructions per second).這個速度遠高于被學術(shù)界廣泛使用的GPGPU-Sim(2MIPS).

Zhang 等學者對GPU WCET 分析的研究主要于GPU 的片上存儲器(L1/L2 cache 等),研究目標是提高GPU程序?qū)ζ洗鎯ζ髟L問的時間和可預測性.文獻[93]研究了把緩存鎖定(cache locking)與緩存分區(qū)(cache partitioning)技術(shù)應用于CPU/GPU 共享緩存管理.他們研究了64KB～1024KB 不同大小末級緩存上鎖定和分區(qū)技術(shù)下的程序性能,結(jié)果表明,這兩種技術(shù)對于提高GPU 程序性能的作用甚微.在文獻[94]中,Zhang 等學者研究了對GPU 片上L2 緩存使用鎖定技術(shù)的效果.實驗結(jié)果表明:對于部分程序,L2 緩存鎖定能夠顯著提高程序性能.實際上,這一研究結(jié)果隱含的另一層結(jié)論是:如果多個wrap 的線程在L2 緩存上大量沖突,可能嚴重降低執(zhí)行性能.此外,Zhang 等學者研究了通過調(diào)整wrap 線程對內(nèi)存的讀寫操作順序來提高程序的時間可預測性與總體性能[95,96].

當前,GPU WCET 分析框架所遇到的瓶頸與多核CPU 時間分析遇到的問題[97]相似.當GPU 上運行多個kernel 程序時,一個程序的執(zhí)行時間勢必受到并發(fā)執(zhí)行的其他程序的嚴重干擾.為了保證WCET 分析的安全性,通常假設最壞的干擾情況發(fā)生,其結(jié)果是WCET 分析過于悲觀,造成系統(tǒng)運行資源的過量預留和嚴重浪費.而在GPU 上,由于程序執(zhí)行的并行度要遠遠高于多核CPU 系統(tǒng),上述問題會更加突出.如果能夠獲取并行kernel 程序執(zhí)行交疊信息,則可以大大提高干涉分析的準確性.但是問題的關(guān)鍵在于:并行kernel 程序的執(zhí)行交疊取決于任務調(diào)度,而WCET 的結(jié)果又是任務調(diào)度的輸入(即程序級和系統(tǒng)級時間分析存在耦合關(guān)系).因此,學術(shù)界亟待在WCET 分析上突破原有的兩級分析框架,把程序級WCET 與可調(diào)度分析結(jié)合在一個分析框架中,從而提高時間分析的精確性.

4.2 GPU實時調(diào)度策略

GPU 實時調(diào)度需要解決的問題是調(diào)度不同優(yōu)先級且相互競爭的作業(yè)完成數(shù)據(jù)移動任務和GPU 計算任務,保證每個作業(yè)能在截止期之前完成.該方向的研究主要聚焦在為GPU 設計實時調(diào)度算法和框架,并建立用于可調(diào)度性分析的數(shù)學模型.由于GPU 的硬件調(diào)度器不提供搶占機制,所以早期研究的GPU 調(diào)度算法都是固定優(yōu)先級的.有限的優(yōu)化措施是把大的kernel 任務切分成小的任務,增加調(diào)度的靈活性.近幾年,有學者提出了基于軟件機制實現(xiàn)kernel 任務搶占,從而有可能實現(xiàn)動態(tài)優(yōu)先級調(diào)度.

GPU 實時調(diào)度領(lǐng)域較早期的工作是Kato 等學者于2011 年提出的TimeGraph 調(diào)度器[98].TimeGraph 可被內(nèi)嵌到開源的GPU 驅(qū)動中,其所提供的PRT(predictable response time)調(diào)度算法支持固定優(yōu)先級調(diào)度,通過預分配給每個任務固定的執(zhí)行預算以及監(jiān)視每個任務已使用的執(zhí)行預算來調(diào)整資源分配.TimeGraph 下,任務有可能運行超時,所以它本質(zhì)上支持的是軟實時調(diào)度.此后,Kato 等學者又提出了RGEM[99],一個GPGPU 實時調(diào)度器,支持固定優(yōu)先級調(diào)度,主要貢獻是在固定優(yōu)先級調(diào)度分析中考慮了DMA 操作的時間延遲.此外,RGEM 把GPU程序的執(zhí)行分為了邏輯執(zhí)行和數(shù)據(jù)拷貝兩個階段,這符合GPU 程序的執(zhí)行特點.Basaran 和Kang 等學者研究了如何把一個運行時間較長且不可搶占的GPU 程序進行切分,從而提高系統(tǒng)的可調(diào)度性[100].但是該工作中,GPU kernel 的切分需要用戶參與而非自動化進行.此后,Zhong 和He 等學者提出了Kernelet 框架[101],該框架能夠自動化分析kernel 程序代碼,但是代碼優(yōu)化是在運行階段進行的,因此無法支持靜態(tài)分析,也難以用于硬實時系統(tǒng).

Verner 等學者研究了如何把sporadic 任務模型下的多個作業(yè)(運行在GPU 上)綜合成一個大的作業(yè),并讓這個大的作業(yè)按照一個四級流水線的方式工作,以約束訪存行為[102～104].所提出的調(diào)度算法能夠用于硬實時系統(tǒng).但是該工作的局限是只能夠處理GPU 程序,而無法分析CPU+GPU SoC 系統(tǒng)計算的應用.CPU 與GPU 在協(xié)同過程中會發(fā)生相互等待,Kim 等學者指出:如果存在任務自我等待(self-suspension)的行為,則傳統(tǒng)的速率單調(diào)調(diào)度分析就不再成立[105],還研究了如何把大的任務切分成小的任務,從而降低由于任務相互等待造成的GPU 時間浪費.但是,如何為這些切分的任務分配固定優(yōu)先級仍是一個主要挑戰(zhàn),這個問題的復雜度是NP-hard.

鑒于NVIDIA 公司不公開GPU 調(diào)度邏輯的細節(jié),在文獻[15]中,Smith 等學者通過在NVIDIA Jetson TX2 上開展廣泛深入的“黑盒實驗”,揭示了該款嵌入式GPU 調(diào)度器的25 條隱含知識,并被歸為8 個大類:通用調(diào)度規(guī)則、非搶占執(zhí)行條件、線程資源約束、GPU 片內(nèi)共享內(nèi)存約束、寄存器資源約束、拷貝操作條件、流(stream)規(guī)則、其他注意事項.其中,CUDA 流表示一個GPU 操作隊列,該隊列中的操作(可能由多個主機線程發(fā)出)將以添加到流中的先后順序而依次執(zhí)行.可以將一個流看作是GPU 上的一個任務,不同流里的任務可以并行執(zhí)行.所以一個流對應于并發(fā)的概念,多個流對應并行的概念.這些在NVIDIA 官方開發(fā)手冊上找不到的調(diào)度細節(jié),對CUDA 程序的優(yōu)化以及對GPU 調(diào)度器的建模都非常有幫助.另一方面,GPU 廠商提供的硬件調(diào)度算法盡管能將計算任務盡可能快速地分配給SM 處理器完成計算任務,但是并沒有提供實時系統(tǒng)動態(tài)優(yōu)先級調(diào)度算法所必須的搶占機制.由于硬件條件的限制,基于軟件機制實現(xiàn)GPU 的kernel 搶占的有效調(diào)度策略較為困難,傳統(tǒng)的方法是對kernel 進一步進行劃分,增加調(diào)度的靈活性.Chen 等學者在文獻[106]中提出EffiSha 搶占調(diào)度框架,使用了一種新穎的以SM 為中心的轉(zhuǎn)換方法,對kernel 的形式進行轉(zhuǎn)換,將kernel 內(nèi)部用于計算ID 的blockIdx.x 用taskId 進行替換,解除了block task 與thread block 的綁定關(guān)系,實現(xiàn)了block 級別的搶占調(diào)度.但是該框架目前只支持單個stream 的情況,在擴展性和效率等方面有所欠缺.Wang 等學者在文獻[107]中研究了一種動態(tài)的thread block 發(fā)射策略,實現(xiàn)了對thread block 的動態(tài)分配.他們對基礎的GPU 架構(gòu)進行了擴展,增加了必要的部件和數(shù)據(jù)結(jié)構(gòu)用于跟蹤GPU 上的thread block 的分配和執(zhí)行情況,從而根據(jù)當前的執(zhí)行狀態(tài)進行動態(tài)的調(diào)整和分配.該方法雖然實現(xiàn)了有效的動態(tài)資源分配和調(diào)度,但是需要改進GPU 的基礎架構(gòu),并且引入了較大的額外負載.

2017 年,隨著NVIDIA Pascal 架構(gòu)GPU 引入了像素級與線程級搶占,一些研究者就如何在GPU 上實現(xiàn)可預測的實時動態(tài)優(yōu)先級調(diào)度展開了研究.由于得到了NVIDIA 公司的支持,Capodieci 等學者在NVIDIA Tegra SoC 的驅(qū)動程序中實現(xiàn)了首個EDF 實時調(diào)度算法[53].該調(diào)度策略作為NVIDIA 虛擬機管理程序上層的軟件分區(qū)運行,利用新一代GPU 支持的線程級搶占特性,實現(xiàn)了GPU 任務的EDF 調(diào)度.不過,這項研究成果還未走出NVIDIA 的實驗室,所以市場上銷售的GPU 卡并不支持GPU 任務的EDF 調(diào)度.針對第2.2 節(jié)問題(6)提出的GPU工作溫度過高會影響任務的時間確定性問題,文獻[108]提出了基于熱量感知的GPU 調(diào)度策略.其思路是:在GPU 任務時間模型里加入熱量模型作為約束,然后通過調(diào)度分析確定任務周期性占用/釋放GPU 的時間長度,理論上保證了在GPU 散熱系統(tǒng)工作正常的情況下,GPU 任務的實時性,并且GPU 溫度不會超過閥值.最后,通過真實SoC 系統(tǒng)測試檢驗了該模型的有效性.

4.3 GPU資源管理策略

調(diào)度策略的實現(xiàn)離不開資源管理策略的配合.通過管理kernel 任務訪問和使用GPU 的資源,能夠很好地提高多任務并行系統(tǒng)的時間確定性.多核CPU 系統(tǒng)研究經(jīng)驗也表明:如果能夠更多地管控程序和系統(tǒng)的行為,可以大大降低分析層面的挑戰(zhàn)[54].GPU 資源的管控的主要目的是協(xié)調(diào)并行程序?qū)Ω黝愘Y源訪問階段,有效降低對資源的訪問沖突,提高程序?qū)τ布L問的時間可預測性,可以分為“空間隔離”和“時間隔離”:空間隔離機制針對的是系統(tǒng)內(nèi)存、GPU 內(nèi)存、GPU 片上的各級共享緩存,時間隔離機制針對的是程序的“訪存階段”和“運算階段”.

在傳統(tǒng)實時系統(tǒng)研究中,緩存層面的空間隔離技術(shù)包括緩存鎖定、緩存旁路、緩存劃分.目前,大多數(shù)GPU不支持緩存劃分和緩存鎖定,而NVIDIA 公司的CUDA 平臺提供了在編譯階段旁路(bypass)L1 cache 和共享內(nèi)存的能力,但是還不支持旁路L2 cache.文獻[109]提出一種分析GPU 程序緩存復用距離(cache reuse distance,簡稱CRD)的技術(shù),如果一個程序的CRD 大于某個閾值,則緩存對于該程序的加速效果將基本不存在,因此可以旁路掉.文獻[110]提出了優(yōu)化并行kernel 任務的線程級并行化(thread-level parallelism,簡稱TLP)來提高GPU 資源利用率.kernel 任務對GPU 資源的需求是多維度的(寄存器、共享內(nèi)存、線程和線程塊),如果放任所有的kernel任務運行盡可能多的線程,則必將導致嚴重的資源競爭沖突而降低GPU 性能.更合理的做法是:根據(jù)kernel 任務申請資源的互補性,調(diào)整各個任務分配的線程塊數(shù),從而最小化對資源競爭的沖突.該文還提出kernel 級動態(tài)緩存旁路技術(shù)來調(diào)和并行kernel 任務對L1 cache 的競爭,即,只讓部分kernel 任務使用L1 cache 并旁路掉其他kernel 任務使用L1 cache.不過,該技術(shù)需要GPU 硬件中增加一個位向量,用于標記應該被旁路掉的kernel 任務.最后,利用GPGPU-Sim 仿真實驗展示了兩種調(diào)度優(yōu)化技術(shù)能夠帶來平均1.42 倍的GPU 性能提升和1.33 倍的能效比提升.Park 等學者在文獻[111]中提出一種多任務GPU 上的動態(tài)資源管理方法,GPU 是由多個SM 組成,任務負載在運行的過程中,在每個SM 上會有不同的分配情況.該方法先對任務的每次運行情況進行性能的監(jiān)聽和測量,根據(jù)不同設置產(chǎn)生的測試結(jié)果,選擇最佳的資源劃分方案.

美國北卡羅萊納大學的Elliott 與Anderson 等學者圍繞GPU 實時調(diào)度開展了近5 年的研究工作,提出了GPUSync 這一實時調(diào)度框架[112].GPUSync 框架的主要創(chuàng)新是把GPU 當作一個共享資源對待,則GPU 上的實時調(diào)度問題轉(zhuǎn)化為帶有資源共享的實時調(diào)度問題.該研究團隊在GPU 實時調(diào)度領(lǐng)域還開展了大量實時層面的工作,相關(guān)的實時調(diào)度技術(shù)已經(jīng)集成到了該團隊開發(fā)的LITMUSRT實時操作系統(tǒng)中.Anderson 團隊的研究成果目前是GPU 實時調(diào)度領(lǐng)域最領(lǐng)先的工作.2011 年,Pellizzoni 等學者提出了用于多核CPU 訪存隔離的可預測執(zhí)行模型(predictable execution model,簡稱 PREM)[113,114].PREM 的思想是:把程序劃分為沖突敏感的訪存階段(contension-sensitive memory phase)和無沖突的計算階段(contension-free computation phase),并在任務調(diào)度時保證兩個階段不會發(fā)生重疊,從而避免了訪存沖突.受PREM 啟發(fā),瑞士蘇黎世聯(lián)邦理工學院的Forsberg 等學者提出了GPUguard 軟件框架[54]來控制SoC 平臺中CPU 與GPU 的訪存沖突.GPUguard 提供了CPU/GPU 對共享內(nèi)存的分時訪問策略、內(nèi)存帶寬控制機制、訪存與計算階段的同步機制.如圖4 所示,GPUguard 為程序的訪存階段和運算階段設置同步點(sync),并保障CPU 訪存階段和GPU 訪存階段在相鄰同步分區(qū)中交替執(zhí)行,從而消除了SoC 上因CPU/GPU 對共享內(nèi)存的無序訪問造成的沖突,提高了程序訪存的時間確定性.Forsberg 的研究團隊把GPUguard 框架實現(xiàn)為一個Linux 內(nèi)核模塊和一個CUDA 應用模塊,并在NVIDIA Tegra TX1 SoC 平臺上得到了良好的實驗結(jié)果.Ali 等學者在文獻[64]中提出了一種軟件框架BWLOCK++實現(xiàn)對SoC 內(nèi)存帶寬的控制,確保實時任務能夠不受非實時且訪存密集型任務的干擾.該方法是一種軟件機制,利用任務的訪存情況與GPU 的內(nèi)存帶寬,在kernel 級別對任務進行訪存帶寬的限制,從而實現(xiàn)優(yōu)先確保實時任務的執(zhí)行.

5 DNN 與加速器的協(xié)同設計技術(shù)

深度學習與網(wǎng)絡加速器以兩個各自獨立的陣營在快速發(fā)展,于是有學者指出:這種不匹配的獨立發(fā)展的現(xiàn)狀會導致模型的設計目標難免會有未充分考慮最新硬件的情況;反過來,網(wǎng)絡加速器的設計也存在未充分考慮最新深度網(wǎng)絡特征的情況[115,116].網(wǎng)絡模型與網(wǎng)絡加速器的協(xié)同設計(co-design),則能夠充分發(fā)掘硬件潛力,提高網(wǎng)絡的計算性能(速度、精度、能效比等),從而避免過度設計而導致的成本增加.協(xié)同設計可以從兩個方面進行:(1) 以網(wǎng)絡為固定標的,優(yōu)化加速器的架構(gòu);(2) 以加速器為固定標的,優(yōu)化網(wǎng)絡的結(jié)構(gòu).

Kwon 等學者展示了如何從這兩個方面進行SqueezeNet 網(wǎng)絡和加速器(基于硬件模擬器)之間協(xié)同設計[116].加速器的架構(gòu)設計常用的執(zhí)行流包括:權(quán)重固定流(weight stationary,簡稱WS)和輸出固定流(output stationary,簡稱OS).針對方向(1),Kwon 通過實驗6 種DNN 模型發(fā)現(xiàn):WS 流和OS 流對于不同大小的卷積操作,速度有明顯的不同(在1×1 卷積中,WS 流比OS 流快1.4 倍～7 倍;在Depthwise 卷積中,OS 流比WS 流快19 倍～96 倍).于是,為SqueezeNet 定制了可變分層執(zhí)行流架構(gòu)的加速器Squeezelerator,并獲得了很好的加速效果(比單一OS 流加速了26%,比單一WS 流加速了106%).針對方向(2),Kwon 又以Squeezelerator 加速器為基礎,對SqueezeNet進行優(yōu)化,得到SqueezeNext[81,116]網(wǎng)絡.優(yōu)化方向包括:把第一個卷積層的卷積核從7×7 縮小到5×5,減少了推理時間;減少SqueezeNet 前面階段的層數(shù)并增加后面階段的層數(shù),因為前面階段的硬件利用率低,而后面階段硬件利用率高.SqueezeNext 獲得了2.59 倍的加速和2.25 倍的能效比提升,并且比原SqueezeNet 模型在圖像分類基準測試上中的準確率高出2%.

Yang 等學者也從兩個方向上進行了協(xié)同設計[117].先從原 ShuffleNetV2 模型出發(fā)提出了優(yōu)化模型DiracDeltaNet,其采用了4 項“激進”的優(yōu)化策略:(1) 把所有的3×3 卷積操作替換成移位(shift)運算[80]加1×1 卷積操作;(2) 最大池化核從3×3 縮小為2×2;(3) 修改了通道交換(channel shuffle)的順序;(4) 采用量化方法(quantization)把浮點型參數(shù)壓縮成整型.這些模型層面的優(yōu)化策略在硬件層面得到了來自協(xié)同設計的FPGA 加速器Synetgy 的支持:由FPGA 負責執(zhí)行“1×1 Conv-Pooling-Shift-Shuffle”操作.該協(xié)同設計方案在ImageNet 圖像分類測試中達到88.2%的精度(top-5),推理速度達到了96.5 FPS(frames per second),這刷新了所有FPGA 分類器的最好成績.Abdelouahab 等學者綜述了通過FPGA 加速CNN 推理網(wǎng)的相關(guān)研究[115],分析了CNN 計算負載、并行性和內(nèi)存訪問情況,闡述了卷積層和全連接層的優(yōu)化問題、近似計算和數(shù)據(jù)路徑優(yōu)化方法.

Gao 等學者提出了通過優(yōu)化數(shù)據(jù)流模型來實現(xiàn)可擴展的NPU 加速器陣列設計[118].如果把Eyeriss NPU 以瓦片架構(gòu)(tiled architecture)加以連接,可以獲得更大的硬件計算能力,從而可以計算更復雜的DNN 任務.然而,隨著NPU 陣列的增大,數(shù)據(jù)冗余增大,數(shù)據(jù)移動和內(nèi)存訪問的開銷也會隨之增大.該研究團隊通過使用兩種技術(shù)——Buffer sharing dataflow 和alternate layer loop ordering 優(yōu)化了數(shù)據(jù)流模型的層內(nèi)并發(fā)性(intra-layer parallelism)和層間流水線(inter-layer pipelining).基準測試實驗顯示該協(xié)同設計方案實現(xiàn)了2 倍的加速比,并減少了45%的能耗.

6 總結(jié)與展望

隨著深度學習算法、嵌入式計算硬件、5G、物聯(lián)網(wǎng)的不斷發(fā)展,以深度學習為主要手段的人工智能技術(shù)必將在嵌入式應用領(lǐng)域得到更加廣泛的應用,這其中也包括了安全攸關(guān)系統(tǒng).因此,如何構(gòu)建可信人工智能系統(tǒng),已經(jīng)成為了學術(shù)界和工業(yè)界的一個研究熱點.本文對現(xiàn)有的面向?qū)崟r應用的深度學習研究工作進行了綜述,介紹了深度學習技術(shù)(主要是推理過程)應用于實時嵌入式系統(tǒng)所面臨的挑戰(zhàn),并從深層神經(jīng)網(wǎng)絡的輕量化設計、GPU 時間分析與任務調(diào)度、CPU+GPU SoC 異構(gòu)平臺的資源管理、深層神經(jīng)網(wǎng)絡與網(wǎng)絡加速器的協(xié)同設計等4 個方面綜述了近5 年來的研究進展.雖然學術(shù)界在以上4 個方取得了一定的成果,但仍然存在一些問題需要進一步研究和完善.下面將總結(jié)這些具體問題并展望進一步的研究方向.

DNN 輕量化設計問題:當前,對DNN 的性能分析和優(yōu)化以實驗為主要手段,缺少基于形式化的DNN 任務建模與分析框架.在DNN 輕量化的研究方面,缺乏描述精度和實時性之間權(quán)衡關(guān)系的量化表達方法,也沒有建立DNN 輕量化技術(shù)的方法論.在DNN 運行框架方面,現(xiàn)有的主流框架,如TensorFlow,Caffe 等,尚無法滿足實時系統(tǒng)對DNN 任務運行時間確定性的要求.針對這些問題,需要進一步研究DNN 任務性能分析模型、DNN 輕量化設計方法論、DNN 實時運行框架.其中,實時框架的開發(fā)涉及硬件、驅(qū)動、庫、框架等多個層次,是個復雜的系統(tǒng)工程;

GPU 的實時分析問題:學者們大多延續(xù)了傳統(tǒng)單核CPU 實時系統(tǒng)的研究路線,無法應對GPU 這種復雜的大規(guī)模并行器件,表現(xiàn)為兩方面.

(1) 研究方法重分析輕設計,強調(diào)分析技術(shù)(WCET 分析、可調(diào)度性分析)在保障系統(tǒng)實時性中的作用.但近幾年多核CPU 分析遇到的困境表明:如果能夠管控GPU 程序的行為,提高時間的可預測性,則可以大大降低分析層面的復雜度;

(2) 現(xiàn)有的時間分析框架中,程序級分析和可調(diào)度性分析是獨立的,沒有考慮到GPU 并行任務之間的干涉,導致分析不精確.新的分析框架應該融合程序級分析與系統(tǒng)級分析,提高時間分析的精確性;

GPU 調(diào)度和資源管理問題:科技巨頭NVIDIA 公司不公開其GPU 調(diào)度邏輯的詳細資料,阻礙了對NVIDIA GPU 開展實時調(diào)度研究和實驗.雖然通過黑盒實驗的方法可以獲得很多不公開的調(diào)度規(guī)則,但是并不能確定這份調(diào)度規(guī)則清單是否足夠完備,在新架構(gòu)的GPU 上是否依然有效.對SoC 平臺上CPU 和GPU 訪問共享內(nèi)存的分時隔離技術(shù)的研究已經(jīng)取得了很大的進展,但是CPU 與GPU 之間顯式或隱式的同步仍然會導致時間不確定性問題.由于AMD 公司對其GPU 技術(shù)細節(jié)的曝露要開放許多,并提供開源驅(qū)動GPU Open[119],因此,一個可行的研究方向是以AMD GPU+OpenCL[120]為平臺來研究GPU 實時調(diào)度[121]和資源管理技術(shù),并研發(fā)用于實時DNN計算的基礎軟件.此外,前面綜述過的調(diào)度或資源管理優(yōu)化的研究工作存在技術(shù)路線不夠系統(tǒng)化的問題,可以從GPU 程序建模分析出發(fā),結(jié)合系統(tǒng)的調(diào)度和資源分配,綜合研究實時性能優(yōu)化技術(shù);

面向?qū)崟r系統(tǒng)的網(wǎng)絡加速器協(xié)同設計問題:無論是通用還是專用網(wǎng)絡加速器僅能在一定程度上改善網(wǎng)絡性能,并難以設計普適性的網(wǎng)絡加速器結(jié)構(gòu).DNN 和網(wǎng)絡加速器的協(xié)同設計可以提高兩者的契合度,從而設計出性能特征高度匹配的網(wǎng)絡與網(wǎng)絡加速器整體解決方案,并降低硬件成本.不過,這個方向的研究主要集中在提高系統(tǒng)的平均性能,還未建立起滿足實時系統(tǒng)要求的協(xié)同設計理論、性能建模與分析方法.考慮到神經(jīng)網(wǎng)絡加速器在未來必將廣泛應用于安全攸關(guān)領(lǐng)域,面向?qū)崟r應用的協(xié)同設計理論是一個非常有意義的研究方向;

智能實時嵌入式系統(tǒng)可更新問題:傳統(tǒng)的實時嵌入式系統(tǒng)基于量體裁衣的方式設計程序,而很少考慮應用或系統(tǒng)更新之后可能會帶來違背時間約束的問題(第2.1 節(jié)中關(guān)鍵問題(3)).Wang 教授在文獻[46]中指出了CPS安全攸關(guān)系統(tǒng)安全可更新問題和解決該問題的必要性、理論方向以及技術(shù)路線,同理,在AI 賦能的實時嵌入式系統(tǒng)中,DNN 模型也會不斷地更新迭代,那么如何保證模型更新之后的人工智能應用仍然能夠滿足原初設計的實時約束,將會是一個更具挑戰(zhàn)性的理論問題,解決該問題無疑將大大促進人工智能實時嵌入式系統(tǒng)的發(fā)展.