龔 昊,劉 瑩,馮建周,趙仁良,冷佳旭
(1.中國(guó)科學(xué)院大學(xué)計(jì)算機(jī)科學(xué)與技術(shù)學(xué)院,北京 100089;2.中國(guó)科學(xué)院大學(xué)數(shù)據(jù)挖掘與高性能計(jì)算實(shí)驗(yàn)室,北京 100089;3.燕山大學(xué)信息科學(xué)與工程學(xué)院,河北 秦皇島 066004)
雷達(dá)信號(hào)處理算法的高性能加速是現(xiàn)代雷達(dá)技術(shù)中至關(guān)重要的一環(huán)。傳統(tǒng)的雷達(dá)信號(hào)處理技術(shù)包括數(shù)字信號(hào)處理器DSP、現(xiàn)場(chǎng)可編程邏輯陣列FPGA和專用集成電路ASIC等[1 - 3],雖然這些技術(shù)已經(jīng)得到了大量的應(yīng)用,但是它們都屬于專用設(shè)備,具有開發(fā)周期長(zhǎng)、調(diào)試難度大等缺點(diǎn),在開發(fā)和測(cè)試過程中需要耗費(fèi)大量的資源。而GPU屬于通用設(shè)備,隨著NVIDIA、AMD等GPU廠商芯片技術(shù)的快速發(fā)展,CPU-GPU異構(gòu)計(jì)算體系結(jié)構(gòu)已經(jīng)走進(jìn)了千家萬戶,CUDA的出現(xiàn)更是大大簡(jiǎn)化了程序員和科學(xué)家們對(duì)GPU進(jìn)行編程的難度,基于CUDA的GPU開發(fā)環(huán)境逐步發(fā)展為成熟的貨架產(chǎn)品。
與CPU相比,GPU的眾核結(jié)構(gòu)特別適合處理雷達(dá)信號(hào)這種大規(guī)模數(shù)據(jù)[4 - 6],它的出現(xiàn)讓雷達(dá)信號(hào)處理領(lǐng)域煥發(fā)了生機(jī)。國(guó)內(nèi)外關(guān)于使用GPU加速脈沖多普勒雷達(dá)這種目標(biāo)探測(cè)雷達(dá)的研究比較少[7 - 19],大部分研究成果都集中在SAR成像這種對(duì)實(shí)時(shí)性要求高的情景之中。文獻(xiàn)[8]針對(duì)脈沖多普勒雷達(dá)信號(hào)處理中的脈沖壓縮環(huán)節(jié),采用FFTW(Fastest Fourier Transform in the West)和CUFFT(CUDA Fast Fourier Transform library)2種實(shí)現(xiàn)方式進(jìn)行了性能基準(zhǔn)測(cè)試,并分析了時(shí)域脈沖壓縮和頻域脈沖壓縮2種實(shí)現(xiàn)方式的性能差異。雖然實(shí)驗(yàn)結(jié)果表明頻域脈沖壓縮可以取得更好的性能收益,但是并沒有針對(duì)脈沖壓縮模塊進(jìn)行有針對(duì)性的性能提升與優(yōu)化。文獻(xiàn)[11]研究了脈沖多普勒雷達(dá)中恒虛警率算法的GPU實(shí)現(xiàn),雖然文中的方法降低了算法的復(fù)雜度,提高了并行性,但是總的計(jì)算量并沒有減少,提升空間有限。文獻(xiàn)[14]提出了一種基于GPU共享內(nèi)存的恒虛警率算法,雖然使用共享內(nèi)存可以優(yōu)化程序的效率,但還是不可避免地引入了過多的重復(fù)運(yùn)算。文獻(xiàn)[15]研究了脈沖壓縮模塊在不同數(shù)據(jù)處理規(guī)模和結(jié)構(gòu)下,GPU硬件上的計(jì)算資源分布與基于CPU平臺(tái)的加速性能之間存在相關(guān)性,并基于實(shí)證數(shù)據(jù)提出了一種基于GPU的資源優(yōu)化配置方法。但是,在大部分情況下,脈沖壓縮只是整個(gè)雷達(dá)信號(hào)處理系統(tǒng)中的一個(gè)流程,對(duì)其資源的過度傾向勢(shì)必會(huì)影響到其他模塊。
本文提出了一套完整的脈沖多普勒雷達(dá)信號(hào)處理GPU并行化方法。通過對(duì)脈沖多普勒雷達(dá)信號(hào)處理程序進(jìn)行熱點(diǎn)測(cè)試來分析程序的性能瓶頸并進(jìn)行針對(duì)性優(yōu)化。本文的貢獻(xiàn)主要體現(xiàn)在以下幾點(diǎn):(1)針對(duì)大規(guī)模雷達(dá)數(shù)據(jù)并行所造成的線程過度利用問題,提出了一種利用網(wǎng)格跨步并行技術(shù)[20]優(yōu)化雷達(dá)信號(hào)處理的方法。該方法能有效地解決GPU創(chuàng)建和撤銷線程所造成的時(shí)間開銷,提高硬件資源利用率。(2)針對(duì)大規(guī)模雷達(dá)數(shù)據(jù)所造成的數(shù)據(jù)傳輸開銷,提出了一種利用多流異步處理技術(shù)加速脈沖多普勒雷達(dá)信號(hào)處理的方法。該方法能顯著降低大規(guī)模雷達(dá)數(shù)據(jù)的數(shù)據(jù)傳輸延遲,減小性能瓶頸。(3)針對(duì)恒虛警率CFAR(Constant False-Alarm Rate)算法中存在的冗余計(jì)算問題,提出了一種使用并行掃描來進(jìn)行算法模塊數(shù)據(jù)預(yù)處理的方法。該方法能有效地避免重復(fù)計(jì)算,對(duì)于大規(guī)模的數(shù)據(jù)有明顯優(yōu)勢(shì)。(4)設(shè)計(jì)了對(duì)比實(shí)驗(yàn),從性能測(cè)試和誤差分析等多角度來評(píng)估并行算法的準(zhǔn)確性和實(shí)時(shí)性。
如圖1所示,脈沖多普勒雷達(dá)信號(hào)處理包含脈沖壓縮模塊、動(dòng)目標(biāo)顯示模塊、動(dòng)目標(biāo)檢測(cè)模塊和恒虛警率模塊4個(gè)基本功能模塊。
Figure 1 Flow chart of pulse Doppler radar signal processing圖1 脈沖多普勒雷達(dá)信號(hào)處理流程圖
將采樣率設(shè)置為1 GHz,產(chǎn)生40個(gè)雷達(dá)脈沖,總的采樣點(diǎn)數(shù)為9 600 000,對(duì)CPU程序利用Intel VTune進(jìn)行熱點(diǎn)測(cè)試,結(jié)果如表1所示。
回波信號(hào)仿真模塊包括產(chǎn)生線性調(diào)頻信號(hào)和產(chǎn)生模擬回波信號(hào)2個(gè)部分,是程序中最耗時(shí)的部分。但是,它并不屬于真正的雷達(dá)信號(hào)處理系統(tǒng)的一部分,它是為了方便進(jìn)行算法測(cè)試而實(shí)現(xiàn)的,在實(shí)際應(yīng)用中這部分會(huì)被來自Rapid I/O等高速接口的真實(shí)數(shù)據(jù)所替代。
Table 1 CPU program hotspot test
初始化和內(nèi)存釋放這2個(gè)步驟在整個(gè)程序的生命周期中只需要進(jìn)行一次,而脈沖壓縮模塊、動(dòng)目標(biāo)顯示模塊、動(dòng)目標(biāo)檢測(cè)模塊和恒虛警率模塊這4部分是雷達(dá)信號(hào)處理系統(tǒng)的核心,需要進(jìn)行針對(duì)性的加速。其中脈沖壓縮模塊和恒虛警率模塊是雷達(dá)信號(hào)處理系統(tǒng)中耗時(shí)最多的2個(gè)部分,本文將對(duì)這2個(gè)部分進(jìn)行算法并行性分析。
2.3.1 脈沖壓縮模塊
脈沖壓縮模塊的作用是將信號(hào)的寬脈沖壓縮成窄脈沖,從而提高信號(hào)的探測(cè)距離和距離分辨率,其最終目的是提高信號(hào)的信噪比。脈沖壓縮的本質(zhì)是計(jì)算回波信號(hào)對(duì)匹配濾波器的沖激響應(yīng)。
脈沖壓縮模塊需要計(jì)算匹配濾波系數(shù),它的計(jì)算方法是將原始線性調(diào)頻信號(hào)進(jìn)行左右翻轉(zhuǎn)變換之后取共軛復(fù)數(shù),其表達(dá)式如式(1)所示:
h(t)=s*(-t)
(1)
其中,h(t)為系統(tǒng)沖激響應(yīng)函數(shù),s*(t)為線性調(diào)頻信號(hào)復(fù)共軛。
計(jì)算沖激響應(yīng)有2種方法,分別是時(shí)域處理和頻域處理。對(duì)于時(shí)域處理,直接計(jì)算回波信號(hào)對(duì)匹配濾波系數(shù)的離散卷積即可,如式(2)所示:
s0(t)=h(t)?sr(t)
(2)
其中,s0(t)為脈沖壓縮結(jié)果,sr(t)為回波信號(hào)。
對(duì)于頻域處理,先對(duì)回波信號(hào)和匹配濾波器的系數(shù)進(jìn)行傅里葉變換,然后將這2種信號(hào)在頻域進(jìn)行點(diǎn)乘,如式(3)所示:
S0(f)=FFT{h(t)?sr(t)}=H(f)·Sr(f)
(3)
其中,S0(f)為脈沖壓縮結(jié)果頻譜,H(f)為系統(tǒng)函數(shù)頻譜,Sr(f)為回波信號(hào)頻譜。
在這一過程中可以對(duì)H(f)進(jìn)行頻域加窗,加窗之后的頻率響應(yīng)輸出峰值會(huì)變小,信噪比也有一定的提升,將點(diǎn)乘的結(jié)果再進(jìn)行傅里葉逆變換即可得到脈沖壓縮的結(jié)果,如式(4)所示:
s0(t)=IFFT(S0(f))
(4)
時(shí)域處理和頻域處理這2種方法都可以進(jìn)行脈沖壓縮,但是頻域處理可以在計(jì)算過程中對(duì)頻域響應(yīng)加窗。時(shí)域脈沖壓縮的計(jì)算效率不高,頻域脈沖壓縮相比較于時(shí)域脈沖壓縮具有更好的并行性。
在脈沖壓縮模塊中,多個(gè)脈沖可以同時(shí)進(jìn)行脈沖壓縮,彼此之間不存在數(shù)據(jù)相關(guān)。對(duì)于單個(gè)脈沖來說,所涉及的操作也僅有FFT、IFFT和復(fù)數(shù)乘法這幾個(gè)基本算子。雖然脈沖壓縮算法是程序的瓶頸,但是在GPU并行化的基礎(chǔ)上可提升優(yōu)化的程度有限。
2.3.2 恒虛警率模塊
恒虛警率模塊也叫做CFAR模塊,用于在信號(hào)處理的結(jié)果中判定目標(biāo)。CFAR模塊會(huì)根據(jù)每一個(gè)采樣點(diǎn)周圍的幅度信息確定一個(gè)門限,超過這個(gè)門限的采樣點(diǎn)會(huì)被判決為目標(biāo)。圖2展示了恒虛警率模塊的工作原理。
Figure 2 Schematic diagram of of CFAR module圖2 恒虛警率模塊原理圖
恒虛警率模塊主要由平方率檢波和雜波功率估計(jì)并判決這2個(gè)部分組成。其中平方律檢波部分是針對(duì)單個(gè)采樣點(diǎn)的變換過程,不存在與其他采樣點(diǎn)數(shù)據(jù)的相互依賴。而雜波功率估計(jì)并判決部分在計(jì)算每一個(gè)采樣點(diǎn)的門限值時(shí)需要用到相鄰采樣點(diǎn)的幅度信息,這就在并行之后的算法中引入了歸約與掃描計(jì)算。能否妥善處理算法中的歸約計(jì)算是制約算法性能提升的關(guān)鍵。
GPU的并行度是有限的,其并行度由硬件執(zhí)行模型來決定。以Maxwell架構(gòu)為例,每個(gè)SM(Streaming Multiprocessor)單元上最多可調(diào)度2個(gè)尺寸為1 024的線程塊,當(dāng)線程塊數(shù)量遠(yuǎn)大于所有SM單元上可調(diào)度的線程塊數(shù)量時(shí),就需要頻繁地進(jìn)行線程的激活和撤銷,造成了不必要的時(shí)間開銷?;诰W(wǎng)格跨步并行的細(xì)粒度并行算法如算法1所示。
算法1基于網(wǎng)格跨步并行的細(xì)粒度并行算法
輸入:雷達(dá)回波信號(hào)。
輸出:雷達(dá)信號(hào)處理結(jié)果。
Step1初始化核函數(shù)網(wǎng)格尺寸和線程塊尺寸。
Step2主機(jī)端調(diào)用核函數(shù)。
Step3對(duì)每一個(gè)GPU線程:
Step3.1計(jì)算當(dāng)前線程在網(wǎng)格中的索引;
Step3.2計(jì)算當(dāng)前線程塊在網(wǎng)格中的跨度;
Step3.3對(duì)每一個(gè)分配給當(dāng)前線程的任務(wù):
Step3.3.1計(jì)算雷達(dá)回波信號(hào)處理結(jié)果。
當(dāng)使用傳統(tǒng)方法時(shí)需要啟用更多的線程塊才能完成大規(guī)模的計(jì)算任務(wù),但是當(dāng)硬件資源不充裕時(shí),這種方法會(huì)造成額外的時(shí)間開銷,線程塊會(huì)被頻繁地調(diào)用。因此,在CUDA編程中應(yīng)盡可能增加每一個(gè)線程的計(jì)算量,減少核函數(shù)調(diào)用的次數(shù)。圖3展示了數(shù)據(jù)元素?cái)?shù)量大于網(wǎng)格中線程數(shù)的情況。圖3中下半部分的深色方塊代表當(dāng)前活躍的線程,在這種情況下要想完成計(jì)算任務(wù)需要喚醒足夠多的線程塊。
Figure 3 Schematic diagram of the number of data elements is greater than the number of threads in the grid圖3 數(shù)據(jù)元素?cái)?shù)量大于網(wǎng)格中線程數(shù)示意圖
Figure 4 Schematic diagram of single-thread grid striding parallism圖4 單線程網(wǎng)格跨步并行示意圖
圖4所示的網(wǎng)格跨步并行可以很好地解決該問題。對(duì)于單個(gè)線程而言,在網(wǎng)格跨度循環(huán)中,線程計(jì)算的第1個(gè)元素使用threadIdx.x+blockIdx.x*blockDim.x得出。其中,threadIdx.x表示線程塊內(nèi)線程索引,blockIdx.x表示網(wǎng)格內(nèi)線程塊索引,blockDim.x表示線程塊內(nèi)線程數(shù)。然后,線程會(huì)按網(wǎng)格中的線程數(shù)blockDim.x*gridDim.x向前推進(jìn),其中g(shù)ridDim.x表示網(wǎng)格內(nèi)線程塊數(shù),直到其數(shù)據(jù)索引超出數(shù)據(jù)元素的數(shù)量為止。
圖5展示了當(dāng)所有線程均按照網(wǎng)格跨步并行的方式運(yùn)作時(shí),所有計(jì)算元素均被覆蓋的情況。
Figure 5 Schematic diagram of all threads grid striding parallism圖5 所有線程網(wǎng)格跨步并行示意圖
網(wǎng)格跨步并行所使用的核函數(shù)尺寸是固定的,一般根據(jù)具體的硬件執(zhí)行模型來設(shè)定而不會(huì)根據(jù)數(shù)據(jù)規(guī)模自動(dòng)調(diào)整。對(duì)于每一個(gè)線程而言,它會(huì)持續(xù)進(jìn)行運(yùn)算,當(dāng)所有的運(yùn)算結(jié)束之后線程才會(huì)被撤銷。
網(wǎng)格跨步并行還可靈活分配核函數(shù)對(duì)GPU的利用率,當(dāng)以多線程形式控制GPU完成工作時(shí),GPU會(huì)被每一個(gè)線程以CUDA流的方式管制。GPU的計(jì)算資源是有限的,它并不能保證所有CUDA流的線程塊同時(shí)運(yùn)行,必須減小同一時(shí)刻每一個(gè)CUDA流上核函數(shù)對(duì)GPU的占有率才能實(shí)現(xiàn)多流計(jì)算的真正并行。
GPU的計(jì)算資源是有限的,而核函數(shù)一般使用盡可能多的計(jì)算資源以達(dá)到最佳效率,這就意味著2個(gè)控制流的計(jì)算不會(huì)重疊,雖然網(wǎng)格跨步并行可以靈活分配每一個(gè)控制流的計(jì)算量,但是這樣做并不能使GPU持續(xù)保持較高的利用率。
CUDA并行計(jì)算由CPU和GPU配合完成,在GPU計(jì)算時(shí)CPU可以以同步的形式等待GPU的計(jì)算結(jié)果,也可以進(jìn)行和GPU沒有數(shù)據(jù)相關(guān)的計(jì)算,這樣就可以重疊CPU和GPU的計(jì)算,提高程序運(yùn)行效率。
CPU和GPU的數(shù)據(jù)流和控制流通信必須經(jīng)過PCIE總線,這是CUDA程序最大的性能瓶頸。當(dāng)內(nèi)存以鎖頁(yè)內(nèi)存的形式進(jìn)行分配時(shí),可以避免操作系統(tǒng)請(qǐng)求分頁(yè)管理將內(nèi)存以頁(yè)面的形式換出到磁盤,從而為內(nèi)存和顯存之間異步的數(shù)據(jù)傳輸創(chuàng)造條件,主機(jī)端與設(shè)備端之間異步的數(shù)據(jù)傳輸可以與計(jì)算重疊起來,以隱藏?cái)?shù)據(jù)傳輸延遲。
如圖6所示的多CUDA流并發(fā)是粗粒度并行方案。圖6中H2D代表主機(jī)端到設(shè)備端的數(shù)據(jù)傳輸,D2H代表設(shè)備端到主機(jī)端的數(shù)據(jù)傳輸。多CUDA流并發(fā)的重要用途是重疊數(shù)據(jù)傳輸和計(jì)算的時(shí)間,無論是主機(jī)端向設(shè)備端的數(shù)據(jù)傳輸,還是設(shè)備端向主機(jī)端的數(shù)據(jù)傳輸,都可以進(jìn)行重疊;無論是CPU計(jì)算的時(shí)間,還是GPU計(jì)算的時(shí)間,都可以進(jìn)行重疊。重疊可以增加程序的吞吐量,更進(jìn)一步地降低延遲。
Figure 6 Multi-CUDA stream concurrency圖6 多CUDA流并發(fā)
在恒虛警率模塊中,對(duì)于每一個(gè)采樣點(diǎn)都要完成一次求和運(yùn)算。采用并行掃描的方法進(jìn)行預(yù)處理可以避免重復(fù)運(yùn)算。在每一次求和運(yùn)算中,涉及到的數(shù)據(jù)元素?cái)?shù)量非常少,只有幾十個(gè)數(shù)據(jù)元素參與歸約。而采樣點(diǎn)的數(shù)據(jù)量非常龐大,甚至能夠達(dá)到吉比特?cái)?shù)量級(jí)。掃描是一種典型的以空間換時(shí)間的方法,當(dāng)有內(nèi)存空間可復(fù)用時(shí),掃描能夠帶來很好的收益,如式(5)所示:
sum(i)=sum(i-1)+a(i)
(5)
其中,sum(i)為掃描處理結(jié)果,a(i)為采樣點(diǎn)數(shù)據(jù)。
掃描預(yù)處理的時(shí)間復(fù)雜度為O(n),其中n表示參與并行掃描運(yùn)算的數(shù)據(jù)元素個(gè)數(shù)。當(dāng)需要獲取子區(qū)間之和時(shí),可以使用式(6)通過O(1)的時(shí)間復(fù)雜度計(jì)算得出。
ans(l,r)=sum(r)-sum(l-1)
(6)
以上展示了串行掃描的方法。當(dāng)使用CPU完成掃描預(yù)處理時(shí),必然會(huì)造成主機(jī)端和設(shè)備端之間的內(nèi)存?zhèn)鬏?。Blelloch算法[21]是一種在設(shè)備端計(jì)算掃描的方法,該算法分為Reduce和Down-Sweep 2個(gè)過程。
圖7展示了Reduce過程,該過程是使用相鄰配對(duì)實(shí)現(xiàn)的并行歸約,為Down-Sweep過程準(zhǔn)備了數(shù)據(jù)。Reduce過程利用所有線程計(jì)算出了所有2的整數(shù)次冪索引值的和。
Figure 7 Reduce process圖7 Reduce過程
圖8展示了Down-Sweep過程,該過程是Reduce過程的逆過程。圖8中實(shí)線箭頭代表求和運(yùn)算,虛線箭頭代表賦值運(yùn)算。
Figure 8 Down-Sweep process圖8 Down-Sweep過程
該并行算法的階躍時(shí)間復(fù)雜度為O(2 logn),工作時(shí)間復(fù)雜度為O(n)。
本文所使用的測(cè)試數(shù)據(jù)為仿真線性調(diào)頻信號(hào),具體參數(shù)如表2所示。
Table 2 Radar simulation signal parameters
CPU實(shí)現(xiàn)中的傅里葉變換操作使用了FFTW3.5加速庫(kù),GPU實(shí)現(xiàn)中的傅里葉變換操作使用了cuFFT10.2加速庫(kù)。測(cè)試所使用的CPU和GPU配置如下所示:
(1)CPU:Intel Core i7-6700HQ。Intel Core i7-6700HQ CPU基于Skylake架構(gòu),具有4個(gè)物理核心,超線程8線程。它含有2個(gè)FMA(Fused Multiply-Add)單元,可同時(shí)發(fā)射2條256 bit融合乘加指令。Skylake架構(gòu)的單精度浮點(diǎn)執(zhí)行單元數(shù)的計(jì)算公式為:
2(融合乘加指令)× 2(FMA單元數(shù))× 8(1條指令可以處理256/32=8個(gè)單精度浮點(diǎn)數(shù))= 32
Intel Core i7-6700HQ CPU的理論單精度峰值浮點(diǎn)運(yùn)算能力計(jì)算公式為:
2.6 GHz (默認(rèn)主頻,超頻3.5 GHz) × 32(單精度浮點(diǎn)執(zhí)行單元數(shù))× 4(物理核心數(shù))= 332.8 GFlop/s = 0.3328 TFlop/s = 0.3328萬億次浮點(diǎn)計(jì)算
(2)GPU:NVIDIA GeForce GTX 950M。NVIDIA GeForce GTX 950M GPU基于Maxwell架構(gòu),它含有5個(gè)SM單元,每個(gè)SM單元包含了128個(gè)CUDA核心,其理論單精度峰值浮點(diǎn)運(yùn)算能力計(jì)算公式為:
1124 MHz(GPU Boost主頻)× 640(CUDA核心數(shù)量)× 2(單個(gè)時(shí)鐘周期內(nèi)能處理的浮點(diǎn)計(jì)算次數(shù))= 1438.72 GFlop/s = 1.43872 TFlop/s = 1.43872萬億次浮點(diǎn)計(jì)算
不同型號(hào)的CPU的物理核心數(shù)差異很大。在實(shí)踐中往往通過比較單核CPU和GPU的性能來衡量加速效果的優(yōu)劣。
4.1.1 網(wǎng)格跨步并行實(shí)驗(yàn)結(jié)果
幾乎所有的核函數(shù)都可以使用網(wǎng)格跨步并行進(jìn)行優(yōu)化。為了體現(xiàn)網(wǎng)格跨步并行的性能優(yōu)勢(shì),本文以核函數(shù)運(yùn)算量最大的CFAR模塊為例來進(jìn)行網(wǎng)格跨步并行的對(duì)比測(cè)試。使用不同的策略在32個(gè)雷達(dá)脈沖,每個(gè)雷達(dá)脈沖8 192個(gè)采樣點(diǎn)的條件下分別進(jìn)行3次測(cè)試,結(jié)果如表3所示。
Table 3 Experiment results of grid striding parallism
GPU運(yùn)算時(shí)間分為數(shù)據(jù)傳輸時(shí)間和內(nèi)核時(shí)間2部分,在本算法中數(shù)據(jù)傳輸時(shí)間對(duì)加速比有顯著的影響。本文所統(tǒng)計(jì)的時(shí)間為這2部分時(shí)間之和,網(wǎng)格跨步并行的優(yōu)化效果只體現(xiàn)在內(nèi)核時(shí)間上。本實(shí)驗(yàn)采用了較小的數(shù)據(jù)量,以更好地體現(xiàn)網(wǎng)格跨步并行的性能提升結(jié)果。
4.1.2 多CUDA流并發(fā)實(shí)驗(yàn)結(jié)果
以整個(gè)雷達(dá)信號(hào)處理流程來進(jìn)行多CUDA流并行的對(duì)比測(cè)試,在測(cè)試中使用1 GHz的采樣率,產(chǎn)生40個(gè)雷達(dá)脈沖,總采樣點(diǎn)數(shù)為9 600 000,調(diào)用脈沖壓縮、MTI(Moving Target Indication)、MTD(Moving Target Detection)和CFAR模塊100次來模擬算法模塊實(shí)時(shí)工作的情況。由于顯存容量的限制,本實(shí)驗(yàn)中CUDA流的最大數(shù)量為8。在不同的CUDA流數(shù)量下進(jìn)行實(shí)驗(yàn)得到的實(shí)驗(yàn)結(jié)果如表4所示。
在理想情況下多CUDA流并行最多可取得近3倍的速度提升,那是從主機(jī)端到設(shè)備端的數(shù)據(jù)傳輸、從設(shè)備端到主機(jī)端的數(shù)據(jù)傳輸和GPU運(yùn)算這三者的時(shí)間完全重疊的情況。進(jìn)行不同的任務(wù)數(shù)據(jù)傳輸和計(jì)算所占的比重不同,在本實(shí)驗(yàn)中數(shù)據(jù)傳輸所占的比重不到1/10,并不能達(dá)到理想的時(shí)間重疊效果,多CUDA流并行的性能提升只有12%。
Table 4 Experiment results of multi-CUDA stream parallism
4.1.3 并行掃描預(yù)處理實(shí)驗(yàn)結(jié)果
通過調(diào)整采樣率來增大數(shù)據(jù)規(guī)模,針對(duì)CFAR模塊采用非并行掃描和并行掃描2種策略,分別設(shè)計(jì)程序并進(jìn)行實(shí)驗(yàn),結(jié)果如表5所示。
Table 5 Experiment results of parallel scan preprocessing
并行掃描通過對(duì)數(shù)據(jù)的預(yù)處理減少了CFAR模塊中門限計(jì)算時(shí)的重復(fù)計(jì)算,從根本上提升了算法的工作效率。
4.1.4 不同硬件平臺(tái)性能測(cè)試加速比
表6整理了不同硬件平臺(tái)上和不同采樣率下算法CPU實(shí)現(xiàn)和GPU實(shí)現(xiàn)所消耗的時(shí)間之比。實(shí)驗(yàn)結(jié)果顯示,程序在PC機(jī)上達(dá)到了10倍左右的加速比,在嵌入式設(shè)備中達(dá)到了42倍左右的加速比。對(duì)于PC機(jī)來說,GPU運(yùn)算能力是CPU運(yùn)算能力的4.3倍;對(duì)于嵌入式設(shè)備來說,GPU運(yùn)算能力是CPU運(yùn)算能力的2.5倍,無論是哪種情況,加速比都超過了計(jì)算能力差距,這說明針對(duì)GPU的性能優(yōu)化取得了很好的效果。
Table 6 Acceleration ratio of performance test
本次測(cè)試使用的PC機(jī)是HP Pavilion Gaming Notebook,其CPU和GPU的性能差距較小,其中NVIDIA GeForce GTX 950M是移動(dòng)定制版,顯存帶寬為28.8 GB/s,較低的峰值浮點(diǎn)運(yùn)算能力和顯存帶寬使得其在性能方面無法與桌面級(jí)顯卡相媲美。
在生產(chǎn)環(huán)境部署時(shí)使用的CPU為雙路Intel Xeon E5-2683 v3,每一路CPU都具有14個(gè)物理核心,其理論峰值浮點(diǎn)運(yùn)算能力為896 GFlop/s。生產(chǎn)環(huán)境使用的GPU為NVIDIA GeForce GTX 1080 Ti,其理論峰值浮點(diǎn)運(yùn)算能力為11 339.78 GFlop/s,顯存帶寬為484.44 GB/s,較大的性能差距使得算法在生產(chǎn)環(huán)境上能夠獲得更加出色的加速比。
Figure 9 Stage result error analysis圖9 階段結(jié)果誤差分析
圖9展示了算法CPU實(shí)現(xiàn)和GPU實(shí)現(xiàn)的相對(duì)誤差,其中脈沖壓縮模塊的絕對(duì)誤差小于3.5×10-5,動(dòng)目標(biāo)顯示模塊絕對(duì)誤差小于4.5×10-5,動(dòng)目標(biāo)檢測(cè)模塊絕對(duì)誤差小于6×10-4。從雷達(dá)信號(hào)處理中間結(jié)果的誤差分析圖中可以看出,相比于CPU實(shí)現(xiàn),GPU實(shí)現(xiàn)具有極高的精度。
圖10展示了CFAR模塊的結(jié)果對(duì)比情況,圖10中橫縱坐標(biāo)分別表示目標(biāo)的距離信息和速度信息。從圖10中可以看到,使用CPU和GPU處理的結(jié)果其距離信息和速度信息完全相同,表明GPU實(shí)現(xiàn)產(chǎn)生了正確的結(jié)果。
Figure 10 Comparison of CFAR results圖10 CFAR結(jié)果對(duì)比
本文提出了一種基于GPU加速的脈沖多普勒雷達(dá)信號(hào)處理方法,通過網(wǎng)格跨步并行、多CUDA流并發(fā)和并行掃描等多種優(yōu)化策略來實(shí)現(xiàn)加速。本文提出的方法既滿足了雷達(dá)信號(hào)處理大吞吐量和高實(shí)時(shí)性的要求,又充分發(fā)揮了GPU設(shè)備眾核并行的優(yōu)勢(shì)。本文所闡述的方法最終在生產(chǎn)環(huán)境上達(dá)到了300倍的加速比。
致謝:
本文研究得到國(guó)家自然科學(xué)基金(71671178)、中國(guó)科學(xué)院大學(xué)優(yōu)秀青年教師科研能力提升重點(diǎn)項(xiàng)目以及總裝備部裝備預(yù)先研究基金項(xiàng)目的支持。同時(shí),感謝北京雷鷹科技有限公司對(duì)本文研究提供的幫助和支持。