基于PCIe的高性能FPGA-GPU-CPU異構(gòu)編程架構(gòu)*

孫兆鵬,周寬久

(大連理工大學(xué)軟件學(xué)院，遼寧 大連 116620)

1 引言

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à)值。

2 相關(guān)工作

現(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)用。

3 異構(gòu)編程環(huán)境

3.1 異構(gòu)編程流程

圖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í)行程序。

3.2 統(tǒng)一編程模型定義

狀態(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)換示例

4 異構(gòu)架構(gòu)直連通信實(shí)現(xià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)系

4.1 PCIe內(nèi)核

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é)果。

4.2 GPU與FPGA直連通信的實(shí)現(xiàn)

由于圖像領(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)

5 實(shí)驗(yàn)環(huán)境與實(shí)驗(yàn)結(jié)果

5.1 實(shí)驗(yàn)環(huán)境

本文搭建的實(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ì)描述。

5.2 實(shí)驗(yàn)結(jié)果

本文進(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)化。

5.3 實(shí)驗(yàn)分析

為了更加清晰地體現(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的占用率。

6 行人檢測(cè)應(yīng)用

本文選擇了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%的性能改善。

7 結(jié)束語(yǔ)

本文提出了一種基于狀態(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)練和匹配中。