劉創,李智
(四川大學 電子信息學院,四川成都,610000)
人工智能在視頻圖像領域已經從理論研究進入了工業和生活,例如自動駕駛、智能視頻監控分析等。這項技術的實現與應用就是讓計算機系統能快速地讀取大量的視頻圖像信息,并及時做出一系列判斷和反饋。因此,如何在有限的時間內分析處理大批量圖像數據是實際產業應用系統設計中的難點。為了滿足駕駛員監控系統(DMS)的應用需求,各類計算機視覺算法及框架層出不窮,各種異構計算架構及高性能計算平臺百花齊放。與此同時,面對特定場景需求的復雜性,學術界和產業界陸續展開了關于如何將算法模型輕松高效部署在特定計算平臺的研究與應用工作[1-4]。
OpenCL 是專門為異構計算制定的標準,它可以協調具有不同架構的處理器同時進行工作,這就解決了傳統的同構架構下資源利用率低、處理速率慢的問題。同時,OpenCL可以充分發揮不同架構處理器的性能,比如將C 與GPU 和FPGA 強大的并行處理能力相結合,可以完成高性能和低功耗的產品設計以及物聯網應用等[5]。
以駕駛員監控系統應用為項目背景,針對深度學習模型計算量大,視頻圖像識別實時性要求高,應用場景設備資源受限等問題,利用 OpenCL 在基于 CPU+GPU 的異構計算平臺上研究實現 YOLOv3 算法的并行化,并且結合最新自動化的端到端的深度學習優化編譯器 TVM 解決算法的多平臺移植和部署問題[6-7]。
駕駛員監控系統,縮寫DMS,主要是實現對駕駛員的身份識別、駕駛員疲勞駕駛以及危險行為的檢測功能[8]。在現階段開始量產的L2-L3 級自動駕駛中,其實都只有在特定條件下才可以實行,很多實際情況需要駕駛員能及時接管車輛進行處置。因此,在駕駛員過于信任自動駕駛而放棄或減弱對駕駛過程的掌控時可能會導致某些事故的發生,而DMS 的引入可以有效減輕這一問題的出現。因此,近年來各國的政策法規等多方面開始推進DMS 的上車:歐盟和中國均出臺法律法規。國內已率先對“兩客一危”等商用車車型安裝DMS 系統作出強制要求,乘用車搭載要求也在推進制定中。而歐盟則將DMS 納入EuroNCAP 五星安全評級的關鍵要素,而且是必要條件。幾乎是從2020 年開始,DMS系統的裝車率快速提升,行業進入發展快車道。
DMS 的核心功能主要是疲勞監測、分心監測、危險行為監測。早期的DMS 方案主要通過非生物特征的技術來實現,比如通過方向盤及轉向傳感器,監測任何不穩定的方向盤運動、車道偏離或無故改變速度等。但由于系統復雜,非直觀感知,整體的搭載率一直很低。現階段的DMS 則大多是基于攝像頭的面部識別和眼球跟蹤技術,通過紅外光等采集駕駛員面部信息再經過算法分析出人員當下的身體狀態,在檢測到駕駛員處于不安全狀態時,再通過閃爍紅光或是方向盤震動等方案對駕駛者進行提醒。整個系統的硬件部分是由攝像頭+集成座艙車機/域控制器解決方案組成;軟件部分則主要涉及視覺加速算法。
OpenCL 在高性能計算領域具有許多優勢,而最重要的優勢之一就是可移植性,允許使用多種加速器,包括多核CPU,GPU,DSP,FPGA 和專用硬件[9]。OpenCL 由三個模塊構成:實現執行在OpenCL 設備上的內核程序的編程語言,定義和控制平臺的應用編程接口和運行時系統。OpenCL 支持任務與數據兩種并行化計算模式,很大程度上增強了GPU 的計算性能,其整體核心架構包括:平臺模型,執行模型,存儲模型,編程模型四種模型。基于OpenCL的異構編程模型設計如圖1 所示。

圖1 基于OpenCL 的異構編程模型設計
在CPU+GPU 異構硬件平臺上,通過OpenCL API 調用查詢平臺和設備屬性,選擇合適的平臺和設備進行初始化。主機端創建上下文、命令隊列,分配內存,并在主機與設備之間進行數據傳輸和計算。異構平臺上,主機封裝設備內存為內存對象以實現數據管理,通過命令隊列向設備發送命令,使用上下文與設備進行信息交互。設備端實現并行算法核函數,創建程序對象執行內核,最后將設備執行的數據結果映射到主機內存,以生成最終的結果文件。
基于CPU+GPU 異構架構,將計算任務劃分為塊。CPU 負責任務調配、復雜邏輯處理和事務管理,GPU 處理簡單邏輯、計算密集、大規模并行的任務。通過并發執行映射到CPU+GPU 多個計算單元的子任務,進一步細粒度劃分模塊以提高CPU 和GPU 的協同計算效率。
TVM 是一種編譯器,支持計算圖級和運算符級優化,能將深度學習任務映射到各種硬件原語,實現性能在不同硬件設備上的可移植性。通過機器學習方法解決高級算子融合和內存延遲等優化問題,并提出了基于成本建模的高效搜索方法,自動優化生成滿足底層硬件特性的程序。
TVM 系統框架與執行流程如下:導入現有框架中的網絡模型,轉換為計算圖,利用高級數據流對計算圖進行優化。運算符級優化生成高效可執行代碼,其中運算符的定義簡化為使用張量描述語言宏觀指定。TVM 結合用戶設定的硬件目標將運算符映射到可能的代碼優化集合,通過基于機器學習的成本模型在優化空間中搜索運算符的優化。最終,系統將訓練得到的優化代碼整合為可部署的編譯運行時模塊,包括優化處理的計算圖、映射及搜索生成的運算符庫和目標設備的運行參數。
設計一種基于GPU 架構的自調優性能模型[10~11]。參數化影響GPU 程序性能的因素并確定取值范圍,構建參數集合空間。通過在GPU 平臺上配置所有可能的參數并測量實際的kernel 執行時間,從測試結果中選取最小值,對應于最優配置。首先,向Host 端主機內存輸入并初始化數據。利用搜索空間優化算法選擇一組最優參數配置。然后,將輸入數據和kernel 函數的參數配置傳輸到Device 端顯存內,初始化設備平臺,并在GPU 上執行kernel 函數進行自調優。最后,將最優配置和實際執行時間返回到Host 端并輸出結果。
基于GPU 架構的OpenCL 性能模型實現如圖2 所示,分為兩方面:一是從并行粒度出發,包括設置OpenCL 核函數中work-group 大小和每個線程處理任務量,workgroup 的維度根據具體應用確定。二是布爾型變量,評估GPU 平臺的優化方法,如局部內存使用、循環展開、避免bank conflict 等,以確定這些優化方法對特定算法應用的有效性。

圖2 基于 GPU 架構的自調優性能模型實現
步驟1:讀取輸入數據,對平臺配置參數化;
步驟2:將輸入數據從host 內存拷貝到device 的全局內存中;
步驟3:通過搜索空間優化算法選取參數配置,所選取的參數配置用來初始化OpenCL 核函數;
步驟4:在device 端執行kernel 函數,得出測試時間并將結果返回host 中;
步驟5:重復步驟3、4,直至遍歷完全部的搜索空間;
步驟6:對所有測試時間進行排序,即可得出最小時間與最優配置。
在上文模型搭建完成的基礎上,設計實現基于 OpenCL的 YOLOv3 視頻圖像識別算法的并行加速[12~13],根據用 C語言和 CUDA 編寫的開源神經網絡框架 DarkNet 項目,然后結合視頻圖像識別任務需求、硬件平臺特征和 OpenCL編程模型,利用OpenCL 的可移植性,在異構系統上實現基于OpenCL 的YOLOv3 卷積神經網絡算法模型。
使用并行編程模型OpenCL 設計規范實現算法,主機應完成OpenCL 平臺設備選擇,內存聲明,建立上下文并創建命令執行隊列等工作,為設備創建緩沖內存及內存對象,將數據從主機端緩存區發送到目標設備的緩存區。設備端應編寫算法中需要并行化設計的內核代碼,然后創建對應的程序對象并在設備上執行內核,在主機代碼中,需要使用clSetKernelArg()設置內核參數,然后調用clEnqueueNDRangeKernel()劃分安排NDRange 工作項和工作組的大小,調用OpenCL 內核函數啟動內核。內核運行后得到的結果數據仍然存儲在設備內存空間,主機需要將數據映射主機內存空間中,最后在設備任務執行完成后,由主機清理工作期間創建的內存緩沖區并關閉OpenCL 對象等。主機函數設計流程圖如圖3 所示。

圖3 主機函數設計流程圖
本文基于clBLAS 庫(OpenCLBLAS,基于OpenCL 內核的基礎線性代數操作數值庫),采用im2col 法將整個卷積過程轉換為GEMM(通用矩陣乘法)過程實現卷積層內核函數。采用NDRange模式實現數據并行,利用系統的多級存儲結構和程序執行的局部性來充分加速運算。
傳統卷積計算的復雜度很高,需要7 層循環遍歷圖像batch 數、batch 大小、圖像通道、圖像尺寸、卷積尺寸。采用矩陣乘法的方式將卷積計算轉化,將圖像通道和卷積核按矩陣拼接,減少循環層數。通過OpenCL 執行模型中的單指令多線程(SIMT)特性,利用多個工作項同時計算,實現并行加速。
在二維卷積運算中,每個輸出點的計算是獨立且不依賴的,因此,通過將OpenCL 工作項一對一映射到輸出點,可以輕松實現多輸出點的并行計算。在卷積神經網絡中,卷積層、批量歸一化和激活函數層通常形成一個固定的結構。為了進一步優化數據流,采用算子融合,將固定結構的計算集中處理,減少數據移動和內核啟動關閉帶來的額外性能開銷。下面為卷積計算convolutional_kernels_cl.cpp 函數設計:

為了描述算法模型中計算操作指定的張量輸出大小以及每個元素的計算表達式,TVM 采用一種張量描述語言來描述張量在索引空間中的每個操作,張量描述語言不僅支持常見的數學運算表示,而且實現了常見的深度學習運算符的表示。借用Halide 將計算算法和調度邏輯進行抽象并分離的思想優化神經網絡算子,然后采用一些搜索算法來找到較優的調度方案,從而自動生成最終的執行代碼。其中Halide 是C++實現的圖像處理領域的領域專用語言(DomainSpecifiedLanguage,DSL)。它的特點是實現了圖像算法的運算(包含函數及表達式),這些運算在計算硬件單元上以函數為單位進行可分離性的調度。
2.3.1 TVM 運行環境搭建
在異構平臺上構建TVM 運行環境,運行系統為Windows10 x64,由于TVM 需要將張量表達式映射到特定的低級代碼以便部署在異構平臺上,因此需要采用低級編譯器中間表示(IR),準備visualstudio2017,CMake。由于TVM 在CPU 平臺的編譯會依賴LLVM,下載LLVM source code 和Clang source code 并使用CMake 編譯,再添加到系統路徑下配置系統環境變量,配置CUDA 及OpenCL。
接下來安裝TVM,從GitHub 上下載整個安裝包,修改tvm 源碼下面的CMakeLists.txt,把USE_LLVM、USE_OPENCL、USE_CUDA 等設置修改為ON。使用CMake 編譯生成tvm.sln,打開tvm.sln,確認編譯的平臺和版本release x64,編譯成功后,獲取Windows 動態庫libtvm.dll,libtvm_topi.dll, 進 入tvm/python,tvm/topi/python,運行pythonsetup.pyinstall,安裝成功便可以導入tvm 包文件。
2.3.2 YOLOv3 模型優化部署
使用TVM 進行模型部署的完整流程:
(1)導入DarkNet 深度學習框架的YOLOv3 模型,以實現計算圖 iR(中間表示)的轉換。
(2)對原始計算圖中間表示(IR)進行計算圖優化,得到優化的計算圖。
(3)對計算圖中的每個計算操作用張量表示描述語言描述的張量計算表達式,并針對異構硬件平臺,選擇最小計算原語生成具體的調度。
(4)使用基于成本模型的機器學習自動優化器生成經過優化的特定的低級代碼。
(5)生成特定于硬件設備的二進制程序。
(6)構建異構平臺可部署的模型。
訓練好的模型編譯為TVM 模型,TVM 模型由deploy.dll、deploy.json、deploy.params 三個文件組成。將生成的TVM 部署庫文件deploy.dll 添加到動態鏈接庫,deploy.json、deploy.params 作為資源文件加入工程項目中。部分核心代碼代碼如下:

本文分為三組實驗組進行對比實驗,實驗組一從MagicData 發布的開源DMS 駕駛員行為數據集中隨機抽5000 張圖,劃分為5 個圖片集,將準備好的每組測試集圖片的路徑全部存放在一個txt 文件里建立測試文件,修改detector.c 文件,在異構平臺上用GPU+CUDA cuDNN 庫執行批量測試,重新make Darknet 框架,進行五次重復測試,并記錄圖片識別推理時間,得到原模型在 GPU 并行運行的性能數據,該實驗為基準實驗組。
實驗組二將YOLOv3 算法移植到基于 OpenCL 的異構平臺上實現優化卷積計算過程,融合卷積層+批量歸一化+激活函數層固定的組成結構,在主機函數和內核函數設計實現過程中運用循環展開、向量化、數據重排、多線程并行和內存訪問優化等運算并行化策略,同時將計算負載合理分配給 CPU 和 GPU,通過 OpenCL 實現異構并行計算并解決移植性問題。基于 OpenCL 的異構并行編程模型,在CPU+GPU 異構計算平臺上運行 YOLOv3 算法模型,得到手動優化算法模型運行性能數據。
手動憑借經驗在異構平臺對算法模型優化設計,存在局限性,無法實現全局最優以及負載均衡,對于算法的優化,涉及高性能張量分解,數據布局,低開銷負載平衡調度,內存分配、通信、同步等多重優化方案。引入TVM 自動優化算法模型,并面向目標平臺編譯生成部署代碼。以實驗組一和實驗組二為參考,在實驗組二算法實現的基礎上植入TVM 的 CPU/GPU 自動編譯優化部署運行測試,得到測試結果。對比結果如表1,圖4 所示。

表1 優化對比實驗

圖4 優化對比實驗結果
通過實驗數據結果,計算三個實驗組的平均時間分別53.635(ms)、37.542(ms)、34.200(ms),實驗表明 YOLOv3 算法在基于OpenCL 的 CPU+GPU 異構計算平臺上相對于GPU 加速比達到 1.42。TVM 優化部署在GPU加速比達到 1.56,并且相對于手動優化也能達到1.10 的加速比。結果表明,本文提出的基于OpenCL 異構平臺的視頻監控圖像處理加速方案和自調優化編譯方法具有可行性和有效性,突破了原算法的應用平臺局限性,有利于與其他設備擴展結合及移植且自動編譯優化方案可快速部署在各種設備端。

本文主要從汽車智能化的駕駛員監控系統實際應用出發,構建了OpenCL 框架下的自調編程模型和TVM 優化編譯器以及算法研究,利用 OpenCL 在基于GPU 的異構計算平臺上實現視頻圖像識別 YOLO 算法的并行化加速,并進一步結合TVM 進行自動編譯優化部署,完成了三個對照組實驗,結果表明相較于基準對照實驗,植入YOLOv3 算法的OpenCL 異構并行編程模型的加速比能達到1.42,基于TVM 優化部署后加速比能達到1.56。表明本文提出的基于異構平臺的圖像識別加速方案和基于 TVM 的端到端自動優化編譯方法具有可行性和有效性,且OpenCL 框架突破了原始基于 CUDA 的應用平臺局限性,利于與其他設備擴展結合及移植且自動編譯優化方案可快速部署在各種設備端。