面向實時應用的深度學習研究綜述?

2020-11-03 12:25:36張政馗龐為光謝文靜呂鳴松

軟件學報 2020年9期

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

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

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

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

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

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

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

1 背景知識

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

1.1 深度學習簡史

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

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

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

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

1.2 深層神經網絡

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

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

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

1.3 嵌入式GPU

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

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 卡數量的影響.圖2顯示了一塊NVIDIA GPU 芯片的架構模塊示意圖,包括3 種關鍵模塊[35]:流處理器簇(streaming multiprocessor,簡稱SM)、流處理器(streaming processor,簡稱SP)、(全局、常量、共享)內存.“流”是一個GPU 操作隊列,該隊列中的操作(可能由多個主機線程發出)將以添加到流中的先后順序而依次執行.可以將一個流看作是GPU 上的一個任務,不同流里的任務可以并行執行.所以一個流對應于并發的概念,多個流對應并行的概念.GPU 是由多個SM 組成的SM 陣列,每個SM 包含8N 個SP(SP 也稱CUDA(compute unified device architecture)核,G80/GT200 中有8 個SP,RTX2080 中有64 個SP),每個SP 都是一個設計簡單的處理器.全局內存(global memory)就是GPU 卡的顯存.紋理內存(texture memory)和常量內存(constant memory)都是針對全局內存的一個特殊視圖.其中,紋理內存是在顯示2D 或3D 圖像時存儲插值計算所需要的數據,常量內存用于存儲只讀數據.每個SM 通過總線可以獨立訪問3 種內存.每個SM 內部都有共享內存(shared memory).與CPU 不同,它沒有自動完成數據替換的硬件邏輯,而是完全由程序員控制,所以它是一種程序可控的高速緩存.

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

以上是PC 平臺GPU 的基本結構、CUDA 程序的基本概念和基本運行規則.下面以NVIDIA Jetson TX2[15,37]為例,介紹嵌入式CPU+GPU SoC 的結構特點.NVIDIA Jetson TX2 SoC 計算卡(圖3)集成了一枚64 位四核ARMv8@2.0 GHz 的A57 微處理器、一枚雙核超標量ARMv8@2.0GHz 的Denver 微處理器和一枚嵌入式的Pascal 架構GPU.兩枚CPU 各自擁有2MB L2 cache.GPU 內包含執行引擎(execution engine,簡稱EE)和拷貝引擎(copy engine,簡稱CE),其中,EE 由兩個流處理器簇(SM)構成,每個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 的數量比Jetson TX2 高出一個數量級,所以其全精度計算性能也高出后者一個數量級,達到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 神經網絡加速器

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

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

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

2 DNN 任務在實時嵌入式系統中面臨的挑戰

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

2.1 DNN模型帶來的關鍵問題

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

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

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

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

(3) DNN 模型結構更新后不能保證系統的實時性

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

2.2 硬件計算平臺所面臨的關鍵問題

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

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

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

(2) GPU 調度機制的細節信息不公開

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

(3) GPU 上執行混合關鍵任務會發生任務間干涉

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

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

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

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

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

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

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

3 DNN 任務性能分析與優化

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

3.1 DNN任務的性能分析

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

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

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

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

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

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

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

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

3.2 輕量化神經網絡

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

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

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

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

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

(2) 神經網絡輕量化設計

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

Table 3 Lightweight network development timeline表3 輕量化網絡發展時間軸

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

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

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

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

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

(3) 自動機器學習

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

3.3 實時神經網絡

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

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

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

4 GPU 時間分析與調度管理技術

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

4.1 GPU WCET分析

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

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

此后,Berezovskyi 等學者轉向用基于測量的動態方法[88]分析GPU 程序的WCET.動態方法本質上要實際執行被分析的kernel 程序,因此該方法能夠考慮程序對各級存儲結構的訪問延時,并適用于有多SM 的GPU.實際上,在只有單個kernel 任務執行的情況下,動態分析方法獲得的kernel 代碼執行時間與靜態分析出的WCET 偏差并不大.原因是:其一,kernel 程序遵循SIMT 模型,其各個線程的訪存模式和代碼執行模式是非常規律的;其二,GPU 調度器為每個SM 預分配足夠的資源以后,SM 在執行時是完全隔離的.但是,動態方法的問題仍然不夠安全,無法用于硬實時系統.此外,動態方法難以深入剖析程序內部的行為特征,因此難以得到有效信息,指導程序性能優化.

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

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

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

4.2 GPU實時調度策略

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

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

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

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

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

4.3 GPU資源管理策略

調度策略的實現離不開資源管理策略的配合.通過管理kernel 任務訪問和使用GPU 的資源,能夠很好地提高多任務并行系統的時間確定性.多核CPU 系統研究經驗也表明:如果能夠更多地管控程序和系統的行為,可以大大降低分析層面的挑戰[54].GPU 資源的管控的主要目的是協調并行程序對各類資源訪問階段,有效降低對資源的訪問沖突,提高程序對硬件訪問的時間可預測性,可以分為“空間隔離”和“時間隔離”:空間隔離機制針對的是系統內存、GPU 內存、GPU 片上的各級共享緩存,時間隔離機制針對的是程序的“訪存階段”和“運算階段”.

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

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

5 DNN 與加速器的協同設計技術

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

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

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

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

6 總結與展望

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

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

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

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

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

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

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

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