劉創(chuàng)，李智

（四川大學 電子信息學院，四川成都，610000）

0 引言

人工智能在視頻圖像領域已經(jīng)從理論研究進入了工業(yè)和生活，例如自動駕駛、智能視頻監(jiān)控分析等。這項技術的實現(xiàn)與應用就是讓計算機系統(tǒng)能快速地讀取大量的視頻圖像信息，并及時做出一系列判斷和反饋。因此，如何在有限的時間內(nèi)分析處理大批量圖像數(shù)據(jù)是實際產(chǎn)業(yè)應用系統(tǒng)設計中的難點。為了滿足駕駛員監(jiān)控系統(tǒng)（DMS）的應用需求，各類計算機視覺算法及框架層出不窮，各種異構計算架構及高性能計算平臺百花齊放。與此同時，面對特定場景需求的復雜性，學術界和產(chǎn)業(yè)界陸續(xù)展開了關于如何將算法模型輕松高效部署在特定計算平臺的研究與應用工作[1-4]。

OpenCL 是專門為異構計算制定的標準，它可以協(xié)調(diào)具有不同架構的處理器同時進行工作，這就解決了傳統(tǒng)的同構架構下資源利用率低、處理速率慢的問題。同時，OpenCL可以充分發(fā)揮不同架構處理器的性能，比如將C 與GPU 和FPGA 強大的并行處理能力相結合，可以完成高性能和低功耗的產(chǎn)品設計以及物聯(lián)網(wǎng)應用等[5]。

以駕駛員監(jiān)控系統(tǒng)應用為項目背景，針對深度學習模型計算量大，視頻圖像識別實時性要求高，應用場景設備資源受限等問題，利用 OpenCL 在基于 CPU+GPU 的異構計算平臺上研究實現(xiàn) YOLOv3 算法的并行化，并且結合最新自動化的端到端的深度學習優(yōu)化編譯器 TVM 解決算法的多平臺移植和部署問題[6-7]。

1 背景與相關研究

1.1 車載監(jiān)控系統(tǒng)

駕駛員監(jiān)控系統(tǒng)，縮寫DMS，主要是實現(xiàn)對駕駛員的身份識別、駕駛員疲勞駕駛以及危險行為的檢測功能[8]。在現(xiàn)階段開始量產(chǎn)的L2-L3 級自動駕駛中，其實都只有在特定條件下才可以實行，很多實際情況需要駕駛員能及時接管車輛進行處置。因此，在駕駛員過于信任自動駕駛而放棄或減弱對駕駛過程的掌控時可能會導致某些事故的發(fā)生，而DMS 的引入可以有效減輕這一問題的出現(xiàn)。因此，近年來各國的政策法規(guī)等多方面開始推進DMS 的上車：歐盟和中國均出臺法律法規(guī)。國內(nèi)已率先對“兩客一危”等商用車車型安裝DMS 系統(tǒng)作出強制要求，乘用車搭載要求也在推進制定中。而歐盟則將DMS 納入EuroNCAP 五星安全評級的關鍵要素，而且是必要條件。幾乎是從2020 年開始，DMS系統(tǒng)的裝車率快速提升，行業(yè)進入發(fā)展快車道。

DMS 的核心功能主要是疲勞監(jiān)測、分心監(jiān)測、危險行為監(jiān)測。早期的DMS 方案主要通過非生物特征的技術來實現(xiàn)，比如通過方向盤及轉向傳感器，監(jiān)測任何不穩(wěn)定的方向盤運動、車道偏離或無故改變速度等。但由于系統(tǒng)復雜，非直觀感知，整體的搭載率一直很低?，F(xiàn)階段的DMS 則大多是基于攝像頭的面部識別和眼球跟蹤技術，通過紅外光等采集駕駛員面部信息再經(jīng)過算法分析出人員當下的身體狀態(tài)，在檢測到駕駛員處于不安全狀態(tài)時，再通過閃爍紅光或是方向盤震動等方案對駕駛者進行提醒。整個系統(tǒng)的硬件部分是由攝像頭+集成座艙車機/域控制器解決方案組成；軟件部分則主要涉及視覺加速算法。

1.2 異構編程模型設計

OpenCL 在高性能計算領域具有許多優(yōu)勢，而最重要的優(yōu)勢之一就是可移植性，允許使用多種加速器，包括多核CPU，GPU，DSP，F(xiàn)PGA 和專用硬件[9]。OpenCL 由三個模塊構成：實現(xiàn)執(zhí)行在OpenCL 設備上的內(nèi)核程序的編程語言，定義和控制平臺的應用編程接口和運行時系統(tǒng)。OpenCL 支持任務與數(shù)據(jù)兩種并行化計算模式，很大程度上增強了GPU 的計算性能，其整體核心架構包括：平臺模型，執(zhí)行模型，存儲模型，編程模型四種模型?；贠penCL的異構編程模型設計如圖1 所示。

圖1 基于OpenCL 的異構編程模型設計

在CPU+GPU 異構硬件平臺上，通過OpenCL API 調(diào)用查詢平臺和設備屬性，選擇合適的平臺和設備進行初始化。主機端創(chuàng)建上下文、命令隊列，分配內(nèi)存，并在主機與設備之間進行數(shù)據(jù)傳輸和計算。異構平臺上，主機封裝設備內(nèi)存為內(nèi)存對象以實現(xiàn)數(shù)據(jù)管理，通過命令隊列向設備發(fā)送命令，使用上下文與設備進行信息交互。設備端實現(xiàn)并行算法核函數(shù)，創(chuàng)建程序對象執(zhí)行內(nèi)核，最后將設備執(zhí)行的數(shù)據(jù)結果映射到主機內(nèi)存，以生成最終的結果文件。

基于CPU+GPU 異構架構，將計算任務劃分為塊。CPU 負責任務調(diào)配、復雜邏輯處理和事務管理，GPU 處理簡單邏輯、計算密集、大規(guī)模并行的任務。通過并發(fā)執(zhí)行映射到CPU+GPU 多個計算單元的子任務，進一步細粒度劃分模塊以提高CPU 和GPU 的協(xié)同計算效率。

1.3 TVM 深度學習自動化端到端優(yōu)化編譯器

TVM 是一種編譯器，支持計算圖級和運算符級優(yōu)化，能將深度學習任務映射到各種硬件原語，實現(xiàn)性能在不同硬件設備上的可移植性。通過機器學習方法解決高級算子融合和內(nèi)存延遲等優(yōu)化問題，并提出了基于成本建模的高效搜索方法，自動優(yōu)化生成滿足底層硬件特性的程序。

TVM 系統(tǒng)框架與執(zhí)行流程如下：導入現(xiàn)有框架中的網(wǎng)絡模型，轉換為計算圖，利用高級數(shù)據(jù)流對計算圖進行優(yōu)化。運算符級優(yōu)化生成高效可執(zhí)行代碼，其中運算符的定義簡化為使用張量描述語言宏觀指定。TVM 結合用戶設定的硬件目標將運算符映射到可能的代碼優(yōu)化集合，通過基于機器學習的成本模型在優(yōu)化空間中搜索運算符的優(yōu)化。最終，系統(tǒng)將訓練得到的優(yōu)化代碼整合為可部署的編譯運行時模塊，包括優(yōu)化處理的計算圖、映射及搜索生成的運算符庫和目標設備的運行參數(shù)。

2 實現(xiàn)與優(yōu)化

2.1 性能模型的設計與實現(xiàn)

設計一種基于GPU 架構的自調(diào)優(yōu)性能模型[10~11]。參數(shù)化影響GPU 程序性能的因素并確定取值范圍，構建參數(shù)集合空間。通過在GPU 平臺上配置所有可能的參數(shù)并測量實際的kernel 執(zhí)行時間，從測試結果中選取最小值，對應于最優(yōu)配置。首先，向Host 端主機內(nèi)存輸入并初始化數(shù)據(jù)。利用搜索空間優(yōu)化算法選擇一組最優(yōu)參數(shù)配置。然后，將輸入數(shù)據(jù)和kernel 函數(shù)的參數(shù)配置傳輸?shù)紻evice 端顯存內(nèi)，初始化設備平臺，并在GPU 上執(zhí)行kernel 函數(shù)進行自調(diào)優(yōu)。最后，將最優(yōu)配置和實際執(zhí)行時間返回到Host 端并輸出結果。

基于GPU 架構的OpenCL 性能模型實現(xiàn)如圖2 所示，分為兩方面：一是從并行粒度出發(fā)，包括設置OpenCL 核函數(shù)中work-group 大小和每個線程處理任務量，workgroup 的維度根據(jù)具體應用確定。二是布爾型變量，評估GPU 平臺的優(yōu)化方法，如局部內(nèi)存使用、循環(huán)展開、避免bank conflict 等，以確定這些優(yōu)化方法對特定算法應用的有效性。

圖2 基于 GPU 架構的自調(diào)優(yōu)性能模型實現(xiàn)

步驟1：讀取輸入數(shù)據(jù)，對平臺配置參數(shù)化；

步驟2：將輸入數(shù)據(jù)從host 內(nèi)存拷貝到device 的全局內(nèi)存中；

步驟3：通過搜索空間優(yōu)化算法選取參數(shù)配置，所選取的參數(shù)配置用來初始化OpenCL 核函數(shù)；

步驟4：在device 端執(zhí)行kernel 函數(shù)，得出測試時間并將結果返回host 中；

步驟5：重復步驟3、4，直至遍歷完全部的搜索空間；

步驟6：對所有測試時間進行排序，即可得出最小時間與最優(yōu)配置。

2.2 YOLOv3 算法在異構計算的實現(xiàn)

在上文模型搭建完成的基礎上，設計實現(xiàn)基于 OpenCL的 YOLOv3 視頻圖像識別算法的并行加速[12~13]，根據(jù)用 C語言和 CUDA 編寫的開源神經(jīng)網(wǎng)絡框架 DarkNet 項目，然后結合視頻圖像識別任務需求、硬件平臺特征和 OpenCL編程模型，利用OpenCL 的可移植性，在異構系統(tǒng)上實現(xiàn)基于OpenCL 的YOLOv3 卷積神經(jīng)網(wǎng)絡算法模型。

使用并行編程模型OpenCL 設計規(guī)范實現(xiàn)算法，主機應完成OpenCL 平臺設備選擇，內(nèi)存聲明，建立上下文并創(chuàng)建命令執(zhí)行隊列等工作，為設備創(chuàng)建緩沖內(nèi)存及內(nèi)存對象，將數(shù)據(jù)從主機端緩存區(qū)發(fā)送到目標設備的緩存區(qū)。設備端應編寫算法中需要并行化設計的內(nèi)核代碼，然后創(chuàng)建對應的程序對象并在設備上執(zhí)行內(nèi)核，在主機代碼中，需要使用clSetKernelArg()設置內(nèi)核參數(shù)，然后調(diào)用clEnqueueNDRangeKernel()劃分安排NDRange 工作項和工作組的大小，調(diào)用OpenCL 內(nèi)核函數(shù)啟動內(nèi)核。內(nèi)核運行后得到的結果數(shù)據(jù)仍然存儲在設備內(nèi)存空間，主機需要將數(shù)據(jù)映射主機內(nèi)存空間中，最后在設備任務執(zhí)行完成后，由主機清理工作期間創(chuàng)建的內(nèi)存緩沖區(qū)并關閉OpenCL 對象等。主機函數(shù)設計流程圖如圖3 所示。

圖3 主機函數(shù)設計流程圖

本文基于clBLAS 庫（OpenCLBLAS，基于OpenCL 內(nèi)核的基礎線性代數(shù)操作數(shù)值庫），采用im2col 法將整個卷積過程轉換為GEMM（通用矩陣乘法）過程實現(xiàn)卷積層內(nèi)核函數(shù)。采用NDRange模式實現(xiàn)數(shù)據(jù)并行，利用系統(tǒng)的多級存儲結構和程序執(zhí)行的局部性來充分加速運算。

傳統(tǒng)卷積計算的復雜度很高，需要7 層循環(huán)遍歷圖像batch 數(shù)、batch 大小、圖像通道、圖像尺寸、卷積尺寸。采用矩陣乘法的方式將卷積計算轉化，將圖像通道和卷積核按矩陣拼接，減少循環(huán)層數(shù)。通過OpenCL 執(zhí)行模型中的單指令多線程（SIMT）特性，利用多個工作項同時計算，實現(xiàn)并行加速。

在二維卷積運算中，每個輸出點的計算是獨立且不依賴的，因此，通過將OpenCL 工作項一對一映射到輸出點，可以輕松實現(xiàn)多輸出點的并行計算。在卷積神經(jīng)網(wǎng)絡中，卷積層、批量歸一化和激活函數(shù)層通常形成一個固定的結構。為了進一步優(yōu)化數(shù)據(jù)流，采用算子融合，將固定結構的計算集中處理，減少數(shù)據(jù)移動和內(nèi)核啟動關閉帶來的額外性能開銷。下面為卷積計算convolutional_kernels_cl.cpp 函數(shù)設計：

2.3 基于TVM 的加速優(yōu)化設計實現(xiàn)

為了描述算法模型中計算操作指定的張量輸出大小以及每個元素的計算表達式，TVM 采用一種張量描述語言來描述張量在索引空間中的每個操作，張量描述語言不僅支持常見的數(shù)學運算表示，而且實現(xiàn)了常見的深度學習運算符的表示。借用Halide 將計算算法和調(diào)度邏輯進行抽象并分離的思想優(yōu)化神經(jīng)網(wǎng)絡算子，然后采用一些搜索算法來找到較優(yōu)的調(diào)度方案，從而自動生成最終的執(zhí)行代碼。其中Halide 是C++實現(xiàn)的圖像處理領域的領域專用語言（DomainSpecifiedLanguage，DSL）。它的特點是實現(xiàn)了圖像算法的運算（包含函數(shù)及表達式），這些運算在計算硬件單元上以函數(shù)為單位進行可分離性的調(diào)度。

2.3.1 TVM 運行環(huán)境搭建

在異構平臺上構建TVM 運行環(huán)境，運行系統(tǒng)為Windows10 x64，由于TVM 需要將張量表達式映射到特定的低級代碼以便部署在異構平臺上，因此需要采用低級編譯器中間表示（IR），準備visualstudio2017，CMake。由于TVM 在CPU 平臺的編譯會依賴LLVM，下載LLVM source code 和Clang source code 并使用CMake 編譯，再添加到系統(tǒng)路徑下配置系統(tǒng)環(huán)境變量，配置CUDA 及OpenCL。

接下來安裝TVM，從GitHub 上下載整個安裝包，修改tvm 源碼下面的CMakeLists.txt，把USE_LLVM、USE_OPENCL、USE_CUDA 等設置修改為ON。使用CMake 編譯生成tvm.sln，打開tvm.sln，確認編譯的平臺和版本release x64，編譯成功后，獲取Windows 動態(tài)庫libtvm.dll，libtvm_topi.dll， 進 入tvm/python，tvm/topi/python，運行pythonsetup.pyinstall，安裝成功便可以導入tvm 包文件。

2.3.2 YOLOv3 模型優(yōu)化部署

使用TVM 進行模型部署的完整流程：

(1)導入DarkNet 深度學習框架的YOLOv3 模型，以實現(xiàn)計算圖 iR（中間表示）的轉換。

(2)對原始計算圖中間表示（IR）進行計算圖優(yōu)化，得到優(yōu)化的計算圖。

(3)對計算圖中的每個計算操作用張量表示描述語言描述的張量計算表達式，并針對異構硬件平臺，選擇最小計算原語生成具體的調(diào)度。

(4)使用基于成本模型的機器學習自動優(yōu)化器生成經(jīng)過優(yōu)化的特定的低級代碼。

(5)生成特定于硬件設備的二進制程序。

(6)構建異構平臺可部署的模型。

訓練好的模型編譯為TVM 模型，TVM 模型由deploy.dll、deploy.json、deploy.params 三個文件組成。將生成的TVM 部署庫文件deploy.dll 添加到動態(tài)鏈接庫，deploy.json、deploy.params 作為資源文件加入工程項目中。部分核心代碼代碼如下：

3 實驗結果與分析

本文分為三組實驗組進行對比實驗，實驗組一從MagicData 發(fā)布的開源DMS 駕駛員行為數(shù)據(jù)集中隨機抽5000 張圖，劃分為5 個圖片集，將準備好的每組測試集圖片的路徑全部存放在一個txt 文件里建立測試文件，修改detector.c 文件，在異構平臺上用GPU+CUDA cuDNN 庫執(zhí)行批量測試，重新make Darknet 框架，進行五次重復測試，并記錄圖片識別推理時間，得到原模型在 GPU 并行運行的性能數(shù)據(jù)，該實驗為基準實驗組。

實驗組二將YOLOv3 算法移植到基于 OpenCL 的異構平臺上實現(xiàn)優(yōu)化卷積計算過程，融合卷積層+批量歸一化+激活函數(shù)層固定的組成結構，在主機函數(shù)和內(nèi)核函數(shù)設計實現(xiàn)過程中運用循環(huán)展開、向量化、數(shù)據(jù)重排、多線程并行和內(nèi)存訪問優(yōu)化等運算并行化策略，同時將計算負載合理分配給 CPU 和 GPU，通過 OpenCL 實現(xiàn)異構并行計算并解決移植性問題?；?OpenCL 的異構并行編程模型，在CPU+GPU 異構計算平臺上運行 YOLOv3 算法模型，得到手動優(yōu)化算法模型運行性能數(shù)據(jù)。

手動憑借經(jīng)驗在異構平臺對算法模型優(yōu)化設計，存在局限性，無法實現(xiàn)全局最優(yōu)以及負載均衡，對于算法的優(yōu)化，涉及高性能張量分解，數(shù)據(jù)布局，低開銷負載平衡調(diào)度，內(nèi)存分配、通信、同步等多重優(yōu)化方案。引入TVM 自動優(yōu)化算法模型，并面向目標平臺編譯生成部署代碼。以實驗組一和實驗組二為參考，在實驗組二算法實現(xiàn)的基礎上植入TVM 的 CPU/GPU 自動編譯優(yōu)化部署運行測試，得到測試結果。對比結果如表1，圖4 所示。

表1 優(yōu)化對比實驗

圖4 優(yōu)化對比實驗結果

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

4 結束語

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