孫兆鵬,周寬久
(大連理工大學(xué)軟件學(xué)院,遼寧 大連 116620)
5G的到來(lái)再次推動(dòng)了物聯(lián)網(wǎng)的發(fā)展,隨之而來(lái)的是海量數(shù)據(jù)。海量數(shù)據(jù)由于保存占用大量空間和資源而需要及時(shí)處理。異構(gòu)計(jì)算[1]是一種特殊新穎的并行計(jì)算方式,它能夠根據(jù)不同計(jì)算單元的結(jié)構(gòu)特點(diǎn)為其分配不同的計(jì)算任務(wù),在提高服務(wù)器計(jì)算性能、能效比和計(jì)算實(shí)時(shí)性方面顯示出傳統(tǒng)架構(gòu)所不具備的優(yōu)勢(shì),因此異構(gòu)計(jì)算技術(shù)是解決未來(lái)數(shù)據(jù)中心能效問(wèn)題的重要手段。
陳左寧院士等專家認(rèn)為多維異構(gòu)硬件資源是由通用計(jì)算資源(如CPU)、多用計(jì)算資源(如FPGA)、專用計(jì)算資源(如GPU)以及存儲(chǔ)資源(主存、外存)、互連資源(內(nèi)部互連、外部接口)所構(gòu)成的復(fù)雜高性能異構(gòu)計(jì)算資源集合,如圖1所示。
Figure 1 Heterogeneous computing resource set圖1 異構(gòu)計(jì)算資源集合
本文的主要貢獻(xiàn)包括3個(gè)方面:
(1) 結(jié)合現(xiàn)有的CUDA和Vivado編譯器與函數(shù)庫(kù),提出了基于狀態(tài)遷移矩陣STM(State Transition Matrix)的異構(gòu)統(tǒng)一自動(dòng)化編程方法。
(2) 利用PCIe實(shí)現(xiàn)GPU到FPGA的直連通信,通過(guò)GPUDirect RDMA實(shí)現(xiàn)FPGA作為主控器的PCIe通信,突破了GPU作為主控器的PCIe通信當(dāng)中讀取操作的短板。
(3) 以行人檢測(cè)為應(yīng)用場(chǎng)景,實(shí)現(xiàn)fastHOG+SVM的異構(gòu)設(shè)計(jì),并在異構(gòu)平臺(tái)上進(jìn)行一系列實(shí)驗(yàn),與基于共享內(nèi)存的間接通信方式進(jìn)行性能比較,證明了直接通信的優(yōu)越性。
本文的其余結(jié)構(gòu)如下:第2節(jié)介紹異構(gòu)計(jì)算方面相關(guān)的工作,第3節(jié)介紹異構(gòu)編程環(huán)境的整個(gè)流程和實(shí)現(xiàn)方案,第4節(jié)描述實(shí)現(xiàn)直接GPU-FPGA傳輸?shù)臋C(jī)制,第5,6節(jié)介紹實(shí)驗(yàn)方案和實(shí)驗(yàn)結(jié)果,第7節(jié)描述未來(lái)的工作,并討論異構(gòu)編程架構(gòu)的實(shí)際應(yīng)用價(jià)值。
現(xiàn)如今在大量不同領(lǐng)域,如深度學(xué)習(xí)[2 - 4]、圖像處理[5 - 7]等領(lǐng)域,有相關(guān)研究比較了FPGA和GPU的性能。這些研究都體現(xiàn)出FPGA和GPU在計(jì)算能力方面有著獨(dú)特的表現(xiàn),F(xiàn)PGA擅長(zhǎng)浮點(diǎn)算術(shù)運(yùn)算,而GPU在矩陣運(yùn)算方面有更好的性能。由于這些特性,研究人員開(kāi)始研究異構(gòu)體系結(jié)構(gòu)來(lái)提高系統(tǒng)的計(jì)算能力。
最初的工作當(dāng)中,大部分研究人員是用不同計(jì)算單元處理不同任務(wù)模塊,例如文獻(xiàn)[5]提出了一種運(yùn)動(dòng)實(shí)時(shí)定位的FPGA-GPU-CPU異構(gòu)平臺(tái),文獻(xiàn)[8]使用組合FPGA-GPU架構(gòu)進(jìn)行心臟生理光學(xué)映射。但是,在這些研究中,F(xiàn)PGA只是起到數(shù)據(jù)采集的作用,GPU始終作為主要的計(jì)算單元,并沒(méi)有充分發(fā)揮各個(gè)計(jì)算單元的計(jì)算性能,算不上真正的異構(gòu)計(jì)算。
后續(xù)的研究中,通過(guò)比較FPGA和GPU的性能給不同計(jì)算單元分配計(jì)算任務(wù)。文獻(xiàn)[9]提出了一種FPGA和GPU性能對(duì)比的系統(tǒng)方法,并且選取了5種不同的算法進(jìn)行實(shí)驗(yàn)比較,最終給出了目標(biāo)設(shè)備的吞吐性能和硬件特性。文獻(xiàn)[10]使用Roofline模型[11]來(lái)檢測(cè)適合不同算法加速的加速器,然后基于fastHOG行人檢測(cè)程序進(jìn)行了驗(yàn)證。文獻(xiàn)[12]以計(jì)算機(jī)視覺(jué)算法為例對(duì)FPGA和GPU進(jìn)行了非常全面的對(duì)比,包括性能、成本和可移植性。后續(xù)一些工作大多基于這些研究提出了任務(wù)級(jí)或代碼級(jí)的并行計(jì)算[13 - 15]以及其他的優(yōu)化方法。但是,這些工作大部分都是基于共享內(nèi)存實(shí)現(xiàn)通信,實(shí)驗(yàn)中并不會(huì)遇到傳輸瓶頸,而實(shí)際應(yīng)用中面對(duì)實(shí)時(shí)的數(shù)據(jù)處理,主機(jī)的內(nèi)存限制了異構(gòu)平臺(tái)的計(jì)算效率。
異構(gòu)系統(tǒng)中涉及多種計(jì)算單元,要想發(fā)揮出各計(jì)算單元的性能,需要以高帶寬和低延遲實(shí)現(xiàn)它們之間的實(shí)時(shí)數(shù)據(jù)傳輸。文獻(xiàn)[16]實(shí)現(xiàn)了一種基于共享內(nèi)存的GPU-FPGA架構(gòu)的桌面編程環(huán)境,初步實(shí)現(xiàn)了工業(yè)級(jí)應(yīng)用的工具鏈,并且進(jìn)行了性能分析。文獻(xiàn)[17]使用FPGA通過(guò)自定義互連實(shí)現(xiàn)了對(duì)等GPU通信。文獻(xiàn)[18]通過(guò)PCIe總線實(shí)現(xiàn)了直接的雙向GPU-FPGA通信方案。但是,這些工作僅僅實(shí)現(xiàn)了GPU作為主控器的傳輸方式,F(xiàn)PGA向GPU傳輸數(shù)據(jù)的效率低下,只有在特殊情景下才能得到傳輸效率的優(yōu)化,無(wú)法體現(xiàn)出異構(gòu)計(jì)算系統(tǒng)的優(yōu)勢(shì)。
經(jīng)過(guò)幾年大范圍多行業(yè)的應(yīng)用和對(duì)行業(yè)內(nèi)其他異構(gòu)應(yīng)用的調(diào)研,我們總結(jié)了用戶與科研人員反饋的一些實(shí)際問(wèn)題:(1)OpenCL通用性更強(qiáng),但是針對(duì)實(shí)際計(jì)算任務(wù),編程復(fù)雜困難,計(jì)算結(jié)果正確性難以保證;(2)高負(fù)載運(yùn)算時(shí)共享內(nèi)存的通信方式對(duì)PCIe設(shè)備通信壓力過(guò)大,限制了各計(jì)算單元性能的發(fā)揮;(3)基本沒(méi)有考慮計(jì)算單元調(diào)度策略,無(wú)法適應(yīng)未來(lái)數(shù)據(jù)中心應(yīng)用。從CUDA 5.0開(kāi)始,NVIDIA發(fā)布了遠(yuǎn)程直接內(nèi)存訪問(wèn)(GPUDirect RDMA)[19],使其他PCIeP[20]設(shè)備可以完全繞開(kāi)CPU內(nèi)存直接訪問(wèn)GPU內(nèi)存。
針對(duì)以上這些問(wèn)題,本文提出了基于STM的異構(gòu)統(tǒng)一自動(dòng)化編程方法,并且利用PCIe實(shí)現(xiàn)了GPU到FPGA的直連通信,通過(guò)實(shí)現(xiàn)FPGA作為主控器的PCIe通信,突破了GPU作為主控器的PCIe通信當(dāng)中讀取操作的短板。進(jìn)行了文件傳輸和行人檢測(cè)的一系列對(duì)比實(shí)驗(yàn),表明了異構(gòu)架構(gòu)的必要性。并且將本文研究應(yīng)用于CPU-GPU-FPGA多維異構(gòu)大數(shù)據(jù)處理平臺(tái)MATRIX,實(shí)現(xiàn)了虛擬化管理。目前該研究成果已在信息安全、人工智能應(yīng)用、集成電路設(shè)計(jì)等諸多領(lǐng)域得到了應(yīng)用。
圖2所示是本文提出的基于STM的異構(gòu)統(tǒng)一自動(dòng)化編程流程。
Figure 2 Process of heterogeneous programming圖2 異構(gòu)編程流程
通過(guò)STM對(duì)異構(gòu)計(jì)算所需的標(biāo)準(zhǔn)C代碼進(jìn)行建模與生成,提高了編程效率,保證了代碼的可靠性。具體流程如下所示:
(1) 建立狀態(tài)遷移矩陣實(shí)現(xiàn)需求的功能,并且確定由GPU或FPGA加速的應(yīng)用程序部分。
(2) 通過(guò)狀態(tài)遷移表進(jìn)行邏輯驗(yàn)證后,生成相應(yīng)的標(biāo)準(zhǔn)C程序。
①GPU代碼→GPU編譯器(CUDA 10.0);
②FPGA代碼→高層次綜合(Vivado HLS)。
(3) 編譯程序綜合FPGA設(shè)計(jì),生成連接CPU、GPU和FPGA二進(jìn)制文件的可執(zhí)行文件。
(4) 加載GPU、CPU代碼二進(jìn)制文件和FPGA配置二進(jìn)制文件。
(5) 執(zhí)行程序。
狀態(tài)遷移矩陣STM[21 - 23],也叫做狀態(tài)遷移表,它是一種基于表格形式的建模方法,可以對(duì)系統(tǒng)的動(dòng)態(tài)行為進(jìn)行有效建模,如圖3所示。
Figure 3 Structure of state transition matrix圖3 狀態(tài)遷移矩陣結(jié)構(gòu)
定義1STM由 (S,E,C) 3元組表示,滿足以下條件:
(1)S為有限狀態(tài)集。STM表的列表示系統(tǒng)模型中狀態(tài)的集合scolumn,狀態(tài)scolumn∈S,column∈{1,2,3,…,M},在STM運(yùn)行過(guò)程中,只有唯一的初始激活狀態(tài)。
(2)E為有限事件集。STM表的行表示系統(tǒng)模型中即將觸發(fā)事件的集合erow,事件erow∈E,row∈{1,2,3,…,N}。
(3)STM表中狀態(tài)scolumn與事件erow相交處的單元格集合為C(si,ej,g(i,j),a(c)),i,j∈{1,2,3,…,N},即系統(tǒng)在si狀態(tài)下,同時(shí)觸發(fā)目標(biāo)ej事件時(shí),處理動(dòng)作表達(dá)式C中的判定條件g(i,j)并進(jìn)入目標(biāo)執(zhí)行狀態(tài)a(c)。
根據(jù)上述定義,對(duì)圖3所示模型進(jìn)行如下說(shuō)明:
(1)狀態(tài)集合S={電源OFF,電源ON};
(2)事件集合E={ON,OFF};
(3)單元格集合C={C00,C11},其中:
C00= {cTransition00,cactions00},其中cTransition00為“電源ON處理”,cactions00為“電源ON”;
C11= {cTransition11,cactions11},其中cTransition11為“電源OFF處理”,cactions11為“電源OFF”。
STM中表現(xiàn)的if分支、switch選擇和for/while循環(huán)結(jié)構(gòu),將控制流轉(zhuǎn)換成STM模型,若從生成的STM模型中發(fā)現(xiàn)一些空白STM單元格,則說(shuō)明代碼可能存在邏輯缺陷,針對(duì)STM模型驗(yàn)證可發(fā)現(xiàn)C代碼邏輯錯(cuò)誤。此時(shí),通過(guò)完善STM模型可生成邏輯完善的標(biāo)準(zhǔn)C代碼,部分生成規(guī)則如圖4所示。
Figure 4 STM to C code conversion example圖4 STM到C代碼轉(zhuǎn)換示例
在本文提出的異構(gòu)編程架構(gòu)中,頂層通過(guò)STM集成CUDA和Vivado的接口來(lái)控制GPU和FPGA的計(jì)算資源,底層各計(jì)算單元之間的通信還需要實(shí)現(xiàn)CPU-GPU-FPGA設(shè)備之間的兩兩互相通信來(lái)提高數(shù)據(jù)傳輸效率。本文的工作是基于PCIe實(shí)現(xiàn)整個(gè)底層通信過(guò)程。NVIDIA和XILINX在各自的產(chǎn)品中已經(jīng)提供了高效的CPU-GPU和CPU-FPGA通信方法,本文的主要貢獻(xiàn)就是實(shí)現(xiàn)高效的GPU-FPGA通信。
現(xiàn)有GPU-FPGA通信的主流方法主要有2種,第1種是通過(guò)CPU內(nèi)存?zhèn)鬏敂?shù)據(jù)來(lái)實(shí)現(xiàn)GPU到FPGA的通信,這種通信方式被稱為間接GPU-FPGA傳輸或基于共享內(nèi)存的通信方式,如圖5中粗實(shí)線所示。使用基于共享內(nèi)存的間接方法,數(shù)據(jù)要經(jīng)過(guò)PCIe交換機(jī)2次,并且系統(tǒng)和內(nèi)存操作也會(huì)需要必要的延遲,因此間接通信方式需要額外的通信等待時(shí)間。
Figure 5 Heterogeneous communication implementation圖5 異構(gòu)架構(gòu)通信實(shí)現(xiàn)方式
第2種則是通過(guò)PCIe總線創(chuàng)建直接的GPU-FPGA通信,如圖5中虛線所示,每次通信數(shù)據(jù)僅通過(guò)PCIe交換機(jī)1次,不會(huì)復(fù)制到CPU內(nèi)存中,GPU-FPGA通信具有更高效的通信效率,稱之為直接GPU-FPGA傳輸。這種方法按照主從設(shè)備又可以劃分為GPU作為主控器的通信和FPGA作為主控器的通信。表1匯總了每種傳輸類型的主/從關(guān)系。本文實(shí)現(xiàn)了這2種通信方式并且分別選取了GPU作為主控器的寫(xiě)入操作和FPGA作為主控器的寫(xiě)入操作實(shí)現(xiàn)GPU與FPGA的雙向通信過(guò)程優(yōu)化整個(gè)異構(gòu)計(jì)算系統(tǒng)的通信效率。下面詳細(xì)介紹如何實(shí)現(xiàn)本文的異構(gòu)架構(gòu)。
Table 1 Subordination in communication表1 通信中的主從關(guān)系
Figure 6 XILINX PCIe structure圖6 XILINX PCIe結(jié)構(gòu)圖
Figure 7 XILINX PCIe Communication speed measurement圖7 XILINX PCIe通信測(cè)速
圖6所示的PCIe內(nèi)核是XILINX設(shè)計(jì)的可免費(fèi)下載的FPGA內(nèi)核。XILINX PCIe IP核[24]為FPGA設(shè)計(jì)人員提供了將PCIe總線封裝成類似于存儲(chǔ)器接口的功能,并將該存儲(chǔ)器接口轉(zhuǎn)換為高速DMA引擎;同時(shí),XILINX也提供了Windows環(huán)境下的驅(qū)動(dòng)程序,可以實(shí)現(xiàn)CPU存儲(chǔ)器和FPGA本地的DDR3存儲(chǔ)器之間的通信,通信速度高達(dá)4.0 GB/s(PCIe 2.0×8)。圖7為本文實(shí)現(xiàn)的CPU和FPGA之間通信的測(cè)速結(jié)果。
由于圖像領(lǐng)域的蓬勃發(fā)展,有關(guān)GPU各方面的研究在現(xiàn)階段都已經(jīng)比較成熟,各種工具封裝得非常完善,所以對(duì)GPU硬件可以進(jìn)行的操作很少,當(dāng)然這本身也非常困難,因?yàn)榇蟛糠值挠布δ芏挤庋b在驅(qū)動(dòng)程序中。通常,CUDA僅支持GPU和CPU內(nèi)存間的傳輸,不支持GPU內(nèi)存與其他任意設(shè)備之間的傳輸。但是,在CUDA 4.0之后,NVIDIA為Quadro和Tesla的專業(yè)級(jí)GPU提供了對(duì)等內(nèi)存?zhèn)鬏敼δ?GPU-GPU),CUDA 5.0之后,NVIDIA又為專業(yè)級(jí)的GPU提供了GPUDirect RDMA[19]的接口,通過(guò)此接口可以獲取GPU memory的總線地址,這樣就可以作為FPGA DMA讀/寫(xiě)的目的地址。
針對(duì)消費(fèi)者級(jí)別GPU不支持GPUDirect RDMA API的問(wèn)題,本文分別實(shí)現(xiàn)了2種GPU與FPGA直連通信解決方案,第1種是GPU始終是總線主控器;第2種是GPU和FPGA分別作為總線控制器,二者僅作為主控器進(jìn)行寫(xiě)入操作,不進(jìn)行讀取操作。因?yàn)槲墨I(xiàn)[18]的工作以及本文后續(xù)的實(shí)驗(yàn)結(jié)果都表明讀取操作的復(fù)雜程度和耗時(shí)遠(yuǎn)遠(yuǎn)超過(guò)寫(xiě)入操作。將FPGA作為主控器啟動(dòng)PCIe通信進(jìn)行寫(xiě)入操作需要訪問(wèn)GPU的內(nèi)部結(jié)構(gòu),但是消費(fèi)者級(jí)別的GPU并不支持這種方式。
第1種實(shí)現(xiàn)方案中,GPU是主控器,F(xiàn)PGA是從屬器。這種方式要求將FPGA的內(nèi)存映射到PCIe總線上,GPU直接對(duì)FPGA內(nèi)存進(jìn)行讀寫(xiě)操作。CUDA API支持CPU內(nèi)存頁(yè)面鎖定,頁(yè)面鎖定后的內(nèi)存頁(yè)具有恒定的物理地址,可以由GPU的總線主控DMA控制器有效地訪問(wèn)。CUDA提供了cudaMalloc(),用于分配和釋放此類內(nèi)存塊。經(jīng)過(guò)實(shí)驗(yàn)可以發(fā)現(xiàn),CUDA提供的API并不區(qū)分操作的是哪個(gè)計(jì)算單元的虛擬地址。因此,CUDA中的這些功能同樣可以直接操作FPGA的虛擬地址。通過(guò)這種方式,本文實(shí)現(xiàn)了直接訪問(wèn)FPGA存儲(chǔ)器的驅(qū)動(dòng)程序,從而使GPU直接向FPGA的內(nèi)存寫(xiě)入數(shù)據(jù)或讀取數(shù)據(jù)。文獻(xiàn)[18]使用了開(kāi)源的Speedy PCIe內(nèi)核,將DDR3物理內(nèi)存地址映射到PCIe總線上,并將這些CUDA API應(yīng)用于GPU-FPGA傳輸。
GPU-FPGA通信中DMA代碼示例如下所示:
Example Code 1:
gpu_ptr=cudaMalloc(MEM_SIZE);
//Allocate GPU memory
run_kermel(gpu_ptr,… );
所有子系統(tǒng)的控制器均通過(guò)其自身的MVB-EMD(EMD為用于中距離傳輸?shù)碾娊橘|(zhì))通信接口接入到MVB網(wǎng)絡(luò)。其中,關(guān)鍵子系統(tǒng)如牽引控制系統(tǒng)、輔助電源控制系統(tǒng)、制動(dòng)控制系統(tǒng)、信號(hào)系統(tǒng)、車門(mén)系統(tǒng)等均具有硬線接口,以便在網(wǎng)絡(luò)控制系統(tǒng)故障時(shí),進(jìn)行緊急牽引操作。
/*Perform GPU computation that modify GPU memory*/
fpga_ptr=DeviceIoControl(IOCTL_SPEEDYPCIE_GET_DIRECT_MAPPED_POINTER);
//Map FPGA memory to GPU virtual address space
cudaBostRegister(fpga_ptr,MEM_SIZE );
//Present FPGA memory to CUDA as locked pages
cudaMemcpy(fpga_ptr,gpu_ptr,MEM_SIZE,cudaMemcpyDeviceToHost);
第2種實(shí)現(xiàn)方案中,首先使用第1種方案實(shí)現(xiàn)GPU作為主控器在FPGA的DDR寫(xiě)入操作;然后還需要利用GPUDirect RDMA[25]的API實(shí)現(xiàn)FPGA作為主控器在GPU的DDR寫(xiě)入操作。
FPGA-GPU通信中GPUDirect RDMA的API接口如下所示:
Example Code 2:
Lock the physical page pointed to bypage_ptr
intnvidia_p2p_get_pages(
uint64_tp2p_token;
uint32_tvtl_space;
//GPU memory virtual address
uint64_tvirtual_address;
uint64_tlength;
/*Returns the number of mapped physical pages and the physical address of each page*/
structnvidia_p2p_page_table**page_ptr,
void (*free_callback)(void*data),
void*data)
//Release the physical page pointed to bypage_ptr
intnvidia_p2p_put_pages(
uint64_tp2p_token;
uint32_tva_space;
uint64_tvirtual_address;
structnvidia_p2p_page_table*page_ptr)
本文搭建的實(shí)驗(yàn)環(huán)境分別選擇了一款高端游戲顯卡GPU NVIDIA GeForce GTX 2070和一款專業(yè)顯卡GPU Leadtek Quadro RTX 4000,二者都支持CUDA 10.0 API,唯一的區(qū)別是Quadro RTX 4000支持GPUDirect RDMA的API,而NVIDIA GeForce GTX 2070不支持。二者都具有16個(gè)PCIe 3.0通道,吞吐量高達(dá)15.8 GB/s。
實(shí)驗(yàn)中使用的FPGA平臺(tái)是XILINX Kintex7 AX7325開(kāi)發(fā)板,芯片為XC7K325TFFG900。支持PCIe 2.0×8的接口,吞吐量可達(dá)4 GB/s。(GPU的吞吐量是它的4倍左右)。GPU和FPGA加速卡插在支持PCIe 3.0和2.0的定制主板上,CPU采用的是Intel Xeon E5-2600V3處理器。異構(gòu)計(jì)算平臺(tái)的具體配置如圖8所示。
Figure 8 Heterogeneous computing platform configuration圖8 異構(gòu)計(jì)算平臺(tái)配置
CPU和FPGA之間的通信以及CPU和GPU之間的通信都是使用供應(yīng)商的驅(qū)動(dòng)程序?qū)崿F(xiàn)的。XILINX 7Series PCIe的驅(qū)動(dòng)程序?yàn)橛脩籼峁┑腇PGA文件句柄實(shí)現(xiàn)了CPU和FPGA之間的通信,cudaMemcpy()則在CPU內(nèi)存的用戶緩沖區(qū)中實(shí)現(xiàn)了CPU和GPU之間的通信。GPU/FPGA分別在通信過(guò)程中充當(dāng)PCIe總線主控器,根據(jù)任務(wù)需求對(duì)CPU的內(nèi)存進(jìn)行寫(xiě)入或讀取操作。而在GPU和FPGA直連通信的2種方式中,分別利用GPUDirect RDMA實(shí)現(xiàn)了FPGA作為主控器的通信和利用修改的Speedy PCIe實(shí)現(xiàn)了GPU作為主控器的通信。具體實(shí)現(xiàn)方案在第4節(jié)已做詳細(xì)描述。
本文進(jìn)行了不同大小文件傳輸實(shí)驗(yàn),并統(tǒng)計(jì)了每個(gè)傳送方向上的不同大小文件傳輸帶寬的曲線圖。每種實(shí)驗(yàn)條件下進(jìn)行10次傳輸,計(jì)算平均速度作為每個(gè)傳輸方向的帶寬。如圖9所示,隨著單個(gè)傳輸文件大小的增加,傳輸效率逐漸提高,直到達(dá)到理論帶寬的漸近值為止。
Figure 9 Experimental results of transmission bandwidth圖9 傳輸帶寬實(shí)驗(yàn)結(jié)果
圖9a和圖9c的5條曲線顯示了本文實(shí)現(xiàn)的GPU到FPGA通信的帶寬,其中,Direct GtoF(write)表示GPU到FPGA直接路徑的帶寬(GPU作為主控器),Direct GtoF(read)表示GPU到FPGA直接路徑的帶寬(FPGA作為主控器),GtoC和CtoF分別表示GPU到CPU和CPU到FPGA的帶寬,Indirect GtoF表示GPU到FPGA間接路徑的帶寬。
Figure 10 Speedup of GPU-FPGA relative to GPU-CPU-FPGA transmission圖10 GPU-FPGA相對(duì)于GPU-CPU-FPGA傳輸?shù)募铀俦?/p>
圖9b和圖9d的5條曲線顯示了FPGA到GPU通信的帶寬,圖中各符號(hào)的含義與圖9中的相同,只有數(shù)據(jù)傳輸方向是反向的。
從圖9中可以看出,大型文件傳輸中黑色實(shí)線代表的是CPU與GPU間的通信效率遠(yuǎn)遠(yuǎn)高于其他計(jì)算單元之間的通信效率,因?yàn)镚PU支持生成3.0×16通道的PCIe接口,F(xiàn)PGA僅支持PCIe 2.0×8。所以,GPU到FPGA間接路徑的帶寬會(huì)受到CPU到FPGA帶寬的限制,導(dǎo)致間接路徑通信效率比較低。
由于在FPGA和GPU與CPU的通信中,始終是CPU作為從屬,因此CPU到FPGA/GPU的傳輸效率遠(yuǎn)低于FPGA/GPU到CPU傳輸數(shù)據(jù)的效率。但是,GPU和CPU之間的帶寬較高,因此間接路徑當(dāng)中,GPU-CPU-FPGA方向的傳輸效率要低于FPGA-CPU-GPU的傳輸效率。如果僅僅是采用實(shí)現(xiàn)GPU作為主控器的傳輸方式,那么在FPGA到GPU的傳輸效率反而會(huì)降低。如果僅僅是采用實(shí)現(xiàn)FPGA作為主控器的傳輸方式,GPU到FPGA的傳輸效率就會(huì)降低。這就是本文使用更多的硬件資源來(lái)實(shí)現(xiàn)2種傳輸方式的目的,使得雙向的傳輸都得到優(yōu)化。
為了更加清晰地體現(xiàn)出各種通信方式的效率比較,本文根據(jù)比值繪制了圖10,4條曲線分別代表以下含義:Direct GtoF(write)表示GPU到FPGA的直接通信(GPU執(zhí)行寫(xiě)入操作)與GPU到FPGA間接通信效率的比值。Direct FtoG(read)表示FPGA到GPU的直接通信(GPU執(zhí)行讀取操作)與FPGA到GPU間接通信效率的比值。Direct GtoF(read)表示GPU到FPGA的直接通信(FPGA執(zhí)行讀取操作)與GPU到FPGA間接通信效率的比值。Direct FtoG(write)表示FPGA到GPU的直接通信(FPGA執(zhí)行寫(xiě)入操作)與FPGA到GPU間接通信效率的比值。
實(shí)驗(yàn)中的主要限制因素是FPGA支持的帶寬,因?yàn)镕PGA是PCIe 2.0×8的接口,因此無(wú)論其他計(jì)算單元傳輸性能如何,它的最高性能都是4 GB/s。由于實(shí)驗(yàn)條件所限,沒(méi)有支持PCIe 3.0的FPGA加速卡,但是無(wú)論如何這都不影響對(duì)直接通信和間接通信的效率進(jìn)行對(duì)比實(shí)驗(yàn),而且如果獲得PCIe 3.0的板卡,本文方案的通信效率還會(huì)大大提升。從圖9中上下圖像的相同類型曲線對(duì)比可以看出,數(shù)據(jù)傳輸?shù)膶?xiě)入操作和讀取操作具有不對(duì)稱的帶寬特性。圖10比較了直接GPU-FPGA傳輸和間接GPU-CPU-FPGA傳輸?shù)南鄬?duì)速度??v軸表示加速比,加速比由直接傳輸與間接傳輸速度的比值計(jì)算得出,小于1表示性能降低,大于1表示性能提高。
在GPU到FPGA的傳輸路徑中,對(duì)于大型文件的傳輸,GPU到FPGA的直接通信(GPU執(zhí)行寫(xiě)入操作)比GPU到FPGA的間接通信效率提高了36.58%,而GPU到FPGA的直接通信(FPGA執(zhí)行讀取操作)與GPU到FPGA的間接通信相比通信效率降低了8.71%。FPGA到GPU的傳輸路徑中,F(xiàn)PGA到GPU的直接通信(FPGA執(zhí)行寫(xiě)入操作)比FPGA到GPU的間接通信效率提高了14.53%,F(xiàn)PGA到GPU的直接通信(GPU執(zhí)行讀取操作)與FPGA到GPU的間接通信效率相比降低了20.52%。
在編寫(xiě)驅(qū)動(dòng)程序時(shí),主控寫(xiě)入的PCIe協(xié)議開(kāi)銷較低。所以,在GPU到FPGA的傳輸路徑中,GPU作為總線主控器,當(dāng)數(shù)據(jù)從GPU傳輸?shù)紽PGA時(shí),GPU啟動(dòng)寫(xiě)入操作請(qǐng)求,F(xiàn)PGA可以全速(3.41 GB/s)接收寫(xiě)入操作的請(qǐng)求。反之,數(shù)據(jù)從FPGA傳輸?shù)紾PU時(shí),GPU發(fā)起讀取請(qǐng)求,但是在FPGA中讀取數(shù)據(jù)導(dǎo)致了額外的PCIe協(xié)議開(kāi)銷,所以實(shí)現(xiàn)的傳輸效率較低(2.40 GB/s)。
GPU-CPU傳輸中同樣體現(xiàn)出這樣的不對(duì)稱特性。在GPU到CPU的傳輸中,GPU作為主控器啟動(dòng)總線寫(xiě)入操作請(qǐng)求,最大速度為14.2 GB/s, 反之,GPU啟動(dòng)總線申請(qǐng)讀取,最大速度降至12.1 GB/s。在PCIe通信中,由于協(xié)議開(kāi)銷和實(shí)現(xiàn)機(jī)制復(fù)雜,總線主控寫(xiě)入操作通常比讀取操作效率更高。共享內(nèi)存中解決這種問(wèn)題的另一種方法是使用CPU作為主控器主控寫(xiě)入操作實(shí)現(xiàn)CPU到GPU的數(shù)據(jù)傳輸,但這樣做會(huì)比本文方法產(chǎn)生更多的硬件開(kāi)銷。同時(shí),本文實(shí)現(xiàn)GPU到FPGA的直接通信并不僅僅是因?yàn)樽x寫(xiě)操作效率問(wèn)題,更為嚴(yán)峻的問(wèn)題是在整個(gè)異構(gòu)系統(tǒng)中PCIe是通信中樞,當(dāng)同時(shí)處理大量數(shù)據(jù)并且系統(tǒng)并行程度非常高時(shí),間接通信需要經(jīng)過(guò)2次PCIe,占用大量帶寬,導(dǎo)致通信阻塞,而直接通信會(huì)大大降低PCIe的占用率。
本文選擇了fastHOG+SVM的行人檢測(cè)應(yīng)用作為GPU-FPGA直接通信的測(cè)試程序。圖11所示為fastHOG算法的流程。
Figure 11 Flow chart of fastHOG algorithm圖11 fastHOG算法流程
在異構(gòu)系統(tǒng)中進(jìn)行算法設(shè)計(jì)首先要對(duì)不同任務(wù)進(jìn)行分析,確定適合進(jìn)行加速的計(jì)算單元,這有助于確定最有效的加速方案。文獻(xiàn)[26,27]的研究表明了在FPGA中進(jìn)行HOG計(jì)算是非常好的選擇,同時(shí),在本文的計(jì)算流程當(dāng)中應(yīng)該既包含GPU到FPGA方向的數(shù)據(jù)通信,又包含F(xiàn)PGA到GPU方向的數(shù)據(jù)通信,根據(jù)這2個(gè)原則本文確定了如圖12所示的異構(gòu)計(jì)算流程。
Figure 12 Data flow of person detection in heterogeneous圖12 異構(gòu)計(jì)算中行人檢測(cè)數(shù)據(jù)流向
確定加速模塊后將不同任務(wù)分配到不同計(jì)算單元就得到了異構(gòu)環(huán)境下的算法流程和數(shù)據(jù)傳輸過(guò)程,首先攝像頭從包含用戶行人的場(chǎng)景中捕獲到原始的視頻數(shù)據(jù)。攝像頭通過(guò)USB連接到異構(gòu)系統(tǒng),視頻數(shù)據(jù)首先傳輸?shù)紺PU的DDR當(dāng)中,經(jīng)過(guò)CPU解碼處理后,根據(jù)需求創(chuàng)建HOG圖像,再傳輸?shù)紾PU的內(nèi)存。GPU進(jìn)行圖像縮小、梯度計(jì)算、SVM等計(jì)算,將得到的處理結(jié)果經(jīng)過(guò)PCIe直接傳輸?shù)紽PGA,然后FPGA進(jìn)行直方圖計(jì)算與標(biāo)準(zhǔn)化,最終將結(jié)果傳輸?shù)紺PU再進(jìn)行標(biāo)準(zhǔn)化輸出。
為了使直接和間接路徑之間的性能比較更有說(shuō)服力,本文使用相同的算法對(duì)相同的數(shù)據(jù)進(jìn)行處理,采用相同傳輸路徑,數(shù)據(jù)傳輸?shù)钠瘘c(diǎn)和終點(diǎn)都是CPU的DDR,這樣能夠排除其他因素的影響,獲得更精確的測(cè)量結(jié)果。如圖12所示,左側(cè)為使用間接GPU-FPGA通信機(jī)制的流程,右側(cè)為使用直接GPU-FPGA通信機(jī)制的流程。本文將對(duì)2種情況進(jìn)行實(shí)驗(yàn)對(duì)比。
Figure 13 fastHOG communication delay圖13 fastHOG通信延遲
本文統(tǒng)計(jì)了程序運(yùn)行中的數(shù)據(jù)傳輸時(shí)間,平均了處理300幀圖像后的結(jié)果。各計(jì)算單元之間的傳輸延遲如圖13所示。圖13中位于上方的路徑是間接傳輸?shù)慕Y(jié)果,間接傳輸時(shí)數(shù)據(jù)經(jīng)過(guò)GPU-CPU-FPGA的路徑通過(guò)CPU的內(nèi)存作為媒介進(jìn)行數(shù)據(jù)傳輸。圖13中位于下方的路徑是直接傳輸?shù)慕Y(jié)果。2種機(jī)制下其傳輸延遲總耗時(shí)平均值分別為263 μs和151 μs。因此,在閑置情況下,最佳的應(yīng)用程序加速比可能是1.74倍。隨著GPU和FPGA處理時(shí)間的增加,加速比會(huì)降低,并且在本實(shí)驗(yàn)中,GPU和FPGA的處理時(shí)間在5~20 ms,整個(gè)應(yīng)用程序可以得到15%~20%的性能改善。
本文提出了一種基于狀態(tài)變遷矩陣(STM)的編程框架,利用CUDA和Vivado提供的端口對(duì)GPU和FPGA資源進(jìn)行了集成。而且通過(guò)STM自動(dòng)生成異構(gòu)計(jì)算所需要的標(biāo)準(zhǔn)C代碼同時(shí)可以進(jìn)行形式化驗(yàn)證,既保證了算法的性能,又提高了異構(gòu)環(huán)境的易用性。本文還通過(guò)PCIe實(shí)現(xiàn)了GPU-FPGA直接通信的機(jī)制,同時(shí)分析了其性能特征,并且通過(guò)實(shí)現(xiàn)FPGA和GPU均可以作為主控器,在通信過(guò)程中僅使用寫(xiě)入操作來(lái)提高PCIe通信的效率。本文使用行人識(shí)別算法fastHOG作為案例進(jìn)行研究,使用GPU始終作為主控器的GPU-FPGA直接通信的機(jī)制,數(shù)據(jù)傳輸效率與基于共享內(nèi)存的間接通信方式相比提高了35%;實(shí)現(xiàn)使用FPGA作為主控器的GPU-FPGA直接通信的機(jī)制優(yōu)化FPGA到GPU的數(shù)據(jù)傳輸效率后,異構(gòu)平臺(tái)的整體通信效率進(jìn)一步提高了20%。希望本文工作能為異構(gòu)計(jì)算和體系結(jié)構(gòu)的發(fā)展做出一定貢獻(xiàn)。
雖然本文使用專業(yè)級(jí)顯卡解決了FPGA到GPU傳輸?shù)钠款i,但是顯然這樣的成本不是所有人都可以承擔(dān)的。因此,隨著后續(xù)各種工具的開(kāi)發(fā),希望可以使用消費(fèi)者級(jí)的GPU來(lái)實(shí)現(xiàn)FPGA到GPU的傳輸帶寬增長(zhǎng)。更進(jìn)一步將本文工作融入OpenCL,將這種方法移植到非NVIDIA的GPU與FPGA的通信當(dāng)中。同時(shí),希望在今后異構(gòu)計(jì)算的研究中不斷有更多的研究人員和廠商加入,提供更多的技術(shù)支持。另一方面本文已經(jīng)使用計(jì)算機(jī)視覺(jué)方面的應(yīng)用程序作為案例進(jìn)行了異構(gòu)計(jì)算加速性能的研究,下一步希望能應(yīng)用在神經(jīng)網(wǎng)絡(luò)訓(xùn)練和匹配中。