龔 昊,劉 瑩,馮建周,趙仁良,冷佳旭

(1.中國科學(xué)院大學(xué)計(jì)算機(jī)科學(xué)與技術(shù)學(xué)院，北京 100089；2.中國科學(xué)院大學(xué)數(shù)據(jù)挖掘與高性能計(jì)算實(shí)驗(yàn)室，北京 100089；3.燕山大學(xué)信息科學(xué)與工程學(xué)院，河北 秦皇島 066004)

1 引言

雷達(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ā)周期長、調(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ī)。國內(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í)性。

2 脈沖多普勒雷達(dá)信號(hào)處理算法分析

2.1 脈沖多普勒雷達(dá)信號(hào)處理流程

如圖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)處理流程圖

2.2 熱點(diǎn)測(cè)試

將采樣率設(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 算法并行性分析

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)鍵。

3 基于GPU的脈沖多普勒雷達(dá)信號(hào)處理算法并行化技術(shù)

3.1 基于網(wǎng)格跨步并行的細(xì)粒度并行化

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ì)算的真正并行。

3.2 基于多CUDA流的粗粒度并行化

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)存以鎖頁內(nèi)存的形式進(jìn)行分配時(shí)，可以避免操作系統(tǒng)請(qǐng)求分頁管理將內(nèi)存以頁面的形式換出到磁盤，從而為內(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ā)

3.3 基于并行掃描的數(shù)據(jù)預(yù)處理

在恒虛警率模塊中，對(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)鬏敗lelloch算法[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)。

4 實(shí)驗(yàn)與評(píng)估

本文所使用的測(cè)試數(shù)據(jù)為仿真線性調(diào)頻信號(hào)，具體參數(shù)如表2所示。

Table 2 Radar simulation signal parameters

4.1 性能測(cè)試

CPU實(shí)現(xiàn)中的傅里葉變換操作使用了FFTW3.5加速庫，GPU實(shí)現(xiàn)中的傅里葉變換操作使用了cuFFT10.2加速庫。測(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é)果誤差分析

4.2 誤差分析

圖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ì)比

5 結(jié)束語

本文提出了一種基于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倍的加速比。

致謝：

本文研究得到國家自然科學(xué)基金(71671178)、中國科學(xué)院大學(xué)優(yōu)秀青年教師科研能力提升重點(diǎn)項(xiàng)目以及總裝備部裝備預(yù)先研究基金項(xiàng)目的支持。同時(shí)，感謝北京雷鷹科技有限公司對(duì)本文研究提供的幫助和支持。