李 冬, 焦義文, 高澤夫, 楊文革, 毛飛龍, 滕 飛

(航天工程大學(xué)電子與光學(xué)工程系, 北京 101400)

0 引 言

相位干涉儀是無(wú)源定位方法中發(fā)展較為成熟的一種比相法測(cè)向體制,相位干涉儀的關(guān)鍵技術(shù)之一是鑒相算法。而在鑒相算法中, 頻域互相關(guān)(簡(jiǎn)稱(chēng)為FX)鑒相法具有能夠最大限度的抑制噪聲能量、減小噪聲相位對(duì)信號(hào)的干擾、鑒相精度高等顯著優(yōu)點(diǎn)。但在FX鑒相法中,涉及大量并行數(shù)據(jù)的處理與計(jì)算,尤其是算法中的快速傅里葉變換(fast Fourier transform, FFT)部分,其運(yùn)算量尤為巨大,對(duì)如何實(shí)現(xiàn)大量高速并行數(shù)據(jù)的快速處理提出了較高要求。為滿(mǎn)足測(cè)向系統(tǒng)高速實(shí)時(shí)性的要求,需要考慮采用協(xié)加速器對(duì)鑒相過(guò)程進(jìn)行加速處理。

常見(jiàn)的相位干涉儀鑒相器使用現(xiàn)場(chǎng)可編程門(mén)陣列(field programmable gate array,FPGA)和數(shù)字信號(hào)處理器(digital signal processing, DSP),雖然其計(jì)算速度能夠滿(mǎn)足實(shí)時(shí)性要求,但存在以下局限性:一是FPGA和DSP開(kāi)發(fā)周期長(zhǎng)、升級(jí)難度大、程序調(diào)試?yán)щy;二是僅使用中央處理器(central processing unit, CPU)作為計(jì)算核心的信號(hào)處理機(jī),受限于CPU時(shí)鐘頻率和內(nèi)核數(shù)目,在處理實(shí)時(shí)數(shù)據(jù)時(shí)面臨著系統(tǒng)資源消耗過(guò)多、性能下降等突出問(wèn)題。

圖形處理單元(graphic processing uint, GPU)由于具有眾多的運(yùn)算核心,特別適合進(jìn)行大量數(shù)據(jù)的并行處理,特別是NVIDIA公司提出的計(jì)算統(tǒng)一設(shè)備架構(gòu)(compute unified device architecture, CUDA)大大簡(jiǎn)化了基于GPU的系統(tǒng)架構(gòu)設(shè)計(jì)與軟件開(kāi)發(fā)流程,使得GPU在通用計(jì)算領(lǐng)域得到了更為廣泛的應(yīng)用。針對(duì)前文分析的相位干涉儀在鑒相過(guò)程中并行運(yùn)算多、實(shí)時(shí)性要求高等特點(diǎn),采用 GPU 來(lái)進(jìn)行核心鑒相算法的運(yùn)行和處理是非常適合的。下文將對(duì)基于GPU 的相關(guān)技術(shù)進(jìn)行具體研究。

前人基于GPU和CUDA架構(gòu),在信號(hào)處理領(lǐng)域取得了一定的研究成果。文獻(xiàn)[22]中,王可東等針對(duì)FFT捕獲算法的并行運(yùn)算、處理數(shù)據(jù)量較大等特點(diǎn)和GPU平臺(tái)特別適合于進(jìn)行并行處理的優(yōu)勢(shì),設(shè)計(jì)了FFT捕獲算法的GPU實(shí)現(xiàn)方案,試驗(yàn)結(jié)果表明捕獲結(jié)果完全正確,且捕獲時(shí)間得到了大幅度縮短。文獻(xiàn)[23]中,劉燕都等基于CUDA平臺(tái)設(shè)計(jì)了混頻器、鑒相器、濾波器等模塊的并行程序,實(shí)現(xiàn)了二進(jìn)制相移鍵控(binary phase shift keying, BPSK)的信號(hào)解調(diào),測(cè)試結(jié)果表明,基于GPU的通用計(jì)算平臺(tái)與專(zhuān)有硬件解調(diào)信號(hào)的指標(biāo)相當(dāng),但通用計(jì)算機(jī)平臺(tái)的實(shí)現(xiàn)方法更為靈活、易于功能擴(kuò)展。文獻(xiàn)[24]中,楊千禾為解決傳統(tǒng)信號(hào)處理機(jī)在研發(fā)階段調(diào)試?yán)щy、計(jì)算能力受硬件限制及程序復(fù)用性差等問(wèn)題,提出使用GPU作為計(jì)算核心方案,將程序編寫(xiě)為并行代碼,調(diào)用合適的GPU快速計(jì)算庫(kù),并通過(guò)相應(yīng)的優(yōu)化方法使優(yōu)化后的GPU計(jì)算性能相比CPU提升了20～1 000倍。文獻(xiàn)[25]中,謝維華等采用高性能并行CUDA編程方法,設(shè)計(jì)并實(shí)現(xiàn)了一種新的相關(guān)器算法,該算法處理1 ms多通道數(shù)據(jù)的相關(guān)運(yùn)算平均時(shí)間小于0.5 ms,實(shí)現(xiàn)了多通道并行實(shí)時(shí)相關(guān)運(yùn)算。文獻(xiàn)[26]中,陳大強(qiáng)等提出了基于GPU的軟件化警戒雷達(dá),完成對(duì)警戒雷達(dá)信號(hào)處理算法從CPU平臺(tái)到GPU平臺(tái)的CUDA軟件并行移植,很好解決了信號(hào)處理實(shí)時(shí)性、雷達(dá)軟硬性解耦等問(wèn)題。

基于上述研究成果和本文的實(shí)際需求，本文設(shè)計(jì)了一種基于GPU的相位干涉儀FX鑒相算法，完成了相應(yīng)的并行程序設(shè)計(jì)，利用CPU進(jìn)行數(shù)據(jù)調(diào)度、分發(fā)和邏輯控制，利用GPU進(jìn)行大量信號(hào)數(shù)據(jù)的并行計(jì)算，并進(jìn)行了實(shí)測(cè)驗(yàn)證。實(shí)驗(yàn)結(jié)果表明，在較好地保證鑒相精度的條件下，本文設(shè)計(jì)的基于GPU的鑒相系統(tǒng)數(shù)據(jù)處理速度可達(dá)CPU平臺(tái)的140倍左右，鑒相速度得到了明顯提升。

1 相位干涉儀測(cè)向原理研究

傳統(tǒng)一維單基線(xiàn)相位干涉儀通過(guò)測(cè)量目標(biāo)信號(hào)到達(dá)基線(xiàn)兩端接收天線(xiàn)的幾何時(shí)延差,計(jì)算來(lái)波方向角。圖1中所示為傳統(tǒng)的一維單基線(xiàn)相位干涉儀測(cè)向原理模型,陣元間的連線(xiàn)稱(chēng)為基線(xiàn)。

圖1 傳統(tǒng)單基線(xiàn)相位干涉儀測(cè)向模型Fig.1 Direction finding model of traditional single baseline phase interferometer

假設(shè)信號(hào)到達(dá)接收天線(xiàn)前的陣前波為平行波,則天線(xiàn)1和天線(xiàn)2接收到的信號(hào)分別為

(1)

(2)

式中:為信號(hào)波長(zhǎng);為入射信號(hào)與視軸之間的夾角,即入射信號(hào)的俯仰角;c為光速,即信號(hào)傳播的速度;為兩個(gè)接收天線(xiàn)的間距,即基線(xiàn)長(zhǎng)度。則對(duì)應(yīng)的來(lái)波方向角為

(3)

實(shí)際工程應(yīng)用中,由于物理因素限制,相位干涉儀鑒相器輸出的測(cè)量相位差的范圍是[-π,π],所以測(cè)量相位差是一個(gè)以2π為模糊的觀測(cè)值,因此無(wú)模糊理論相位差和鑒相器測(cè)量相位差的關(guān)系滿(mǎn)足:

(4)

式中:為相位模糊數(shù)。由式(2)和式(4)可得

(5)

對(duì)于多基線(xiàn)相位干涉儀接收同一入射信號(hào)時(shí),假設(shè)有條基線(xiàn),基線(xiàn)長(zhǎng)度分別為,,…,,基線(xiàn)長(zhǎng)度滿(mǎn)足:

::…:=::…:

(6)

式中：(=1,2,…,)分別為互質(zhì)正整數(shù),在不存在相位差噪聲干擾的情況下,任意兩個(gè)基線(xiàn)的相位差滿(mǎn)足:

(7)

其中,,=1,2,…,。由式(7)可得,相位干涉儀不同長(zhǎng)度基線(xiàn)的相位差不同,且基線(xiàn)越長(zhǎng)相位模糊數(shù)越大,即鑒相器測(cè)量相位差在[-π,π]范圍內(nèi)翻折次數(shù)越多。

2 FX鑒相算法基本原理

以只有兩個(gè)陣元的單基線(xiàn)干涉儀為例,傳統(tǒng)的FX鑒相算法先將干涉儀接收到的兩路信號(hào)(),()做FFT得到()和(),再求得其相關(guān)譜為

(8)

相關(guān)譜譜峰位置處的相位值即為兩路信號(hào)的相位差:

=angle[()|=2π]

(9)

本系統(tǒng)是四陣元測(cè)向系統(tǒng),算法的流程示意圖如圖2所示。由4個(gè)天線(xiàn)接收四路模擬信號(hào),經(jīng)過(guò)射頻前端下變頻處理后,再對(duì)信號(hào)分別進(jìn)行離散化處理,得到四路數(shù)字信號(hào),接下來(lái)數(shù)字信號(hào)進(jìn)入鑒相系統(tǒng)中開(kāi)始進(jìn)行提相處理。

圖2 四陣元相位干涉儀FX鑒相算法流程Fig.2 Flow of FX phase identification for four-array phase interferometer

首先對(duì)四路數(shù)字信號(hào)分別進(jìn)行FFT,得到每路信號(hào)的頻譜,再由每?jī)陕沸盘?hào)進(jìn)行共軛相乘得到相關(guān)譜。受躁聲的影響,由少數(shù)幾個(gè)數(shù)據(jù)得到的結(jié)果可能不準(zhǔn)確,因此需要多個(gè)樣本平均以提高測(cè)向準(zhǔn)確度,即能量累積。然后，在搜索到載波頻率后利用反正切求出頻譜最大值對(duì)應(yīng)位置處的相位差信息,由此實(shí)現(xiàn)了四陣元相位干涉儀FX鑒相算法。

3 基于GPU的FX鑒相算法

本文設(shè)計(jì)的FX鑒相算法基于GPU+CPU異構(gòu)處理平臺(tái)開(kāi)發(fā),以CPU為控制器,GPU為協(xié)處理器,由于GPU采用并行處理的架構(gòu)設(shè)計(jì),使得其在進(jìn)行大量并行數(shù)據(jù)處理時(shí)計(jì)算性能明顯優(yōu)于CPU。

3.1 本文算法相比于傳統(tǒng)算法的主要改進(jìn)

本文設(shè)計(jì)的基于GPU的FX鑒相算法,相比于傳統(tǒng)基于CPU的算法,主要改進(jìn)如下:

如圖3所示, 利用GPU具有的眾核優(yōu)勢(shì),將本文面對(duì)的FX鑒相這一核心算法加載到基于CUDA架構(gòu)的GPU中實(shí)現(xiàn),運(yùn)用CUDA流結(jié)構(gòu)設(shè)計(jì)了高速并行的實(shí)時(shí)鑒相算法處理流程,通過(guò)調(diào)用GPU中足夠多的線(xiàn)程和線(xiàn)程塊,對(duì)于鑒相算法的各分模塊(特別是FFT模塊)進(jìn)行了加速并行化處理,基于CUDA C語(yǔ)言編寫(xiě)了基于GPU的FX鑒相算法代碼,在實(shí)測(cè)過(guò)程中通過(guò)調(diào)整GPU中的線(xiàn)程和線(xiàn)程塊數(shù)目,實(shí)現(xiàn)最優(yōu)的加速并行效果,最終完成了整個(gè)鑒相算法。

圖3 基于GPU的FX鑒相算法在GPU中的線(xiàn)程調(diào)用情況Fig.3 GPU thread invocation of GPU-based FX phase identification algorithm

本算法通過(guò)將原本需要在CPU中實(shí)現(xiàn)的眾多運(yùn)算單元加載到GPU中,充分發(fā)揮了GPU的多線(xiàn)程并發(fā)處理的優(yōu)勢(shì);通過(guò)設(shè)計(jì)GPU中的鑒相算法流程,有效調(diào)用了GPU中的運(yùn)算資源,利用眾多線(xiàn)程和線(xiàn)程塊實(shí)現(xiàn)了各運(yùn)算單元的高速處理,基本滿(mǎn)足了實(shí)時(shí)化處理需求。

以本算法的核心且運(yùn)算量最大的FFT模塊為例,在進(jìn)行FFT計(jì)算時(shí)可以將任意維數(shù)的矩陣數(shù)據(jù)分成若干層,其中層與層是串行的計(jì)算,每層計(jì)算的次數(shù)和點(diǎn)數(shù)都一樣,所以具有良好的并行特性,因此使用具有高度并行性的GPU進(jìn)行鑒相處理可以極大地提高數(shù)據(jù)處理速度。本文調(diào)用了GPU中的512個(gè)線(xiàn)程塊,充分滿(mǎn)足了FFT模塊的運(yùn)算量和實(shí)時(shí)性要求。并且不同于傳統(tǒng)算法在CPU下先將實(shí)數(shù)轉(zhuǎn)成復(fù)數(shù)再進(jìn)行計(jì)算,GPU下的FFT計(jì)算通過(guò)cufftExecR2C(實(shí)數(shù)到復(fù)數(shù))直接輸出復(fù)數(shù)結(jié)果,因此極大地提高了算法運(yùn)行效率。

3.2 基于GPU的FX鑒相算法的主要流程

基于GPU的FX鑒相算法實(shí)現(xiàn)流程如圖4所示。

圖4 基于GPU的FX鑒相算法實(shí)現(xiàn)流程Fig.4 Implementation process of FX phase identification based on GPU

GPU中的算法流程主要分為以下幾步:原始數(shù)據(jù)轉(zhuǎn)換、FFT、共軛相乘、能量累積和求頻譜最大值處的相位值。其中,FFT調(diào)用CUDA已有的函數(shù)庫(kù)CUFFT中的cufftPlanD函數(shù),其他各步的處理則分別調(diào)用在GPU內(nèi)編寫(xiě)的CUDA架構(gòu)的核函數(shù)(kernel)。當(dāng)調(diào)用的一個(gè)函數(shù)或者加載的核函數(shù)接收到數(shù)據(jù)后,將啟用一定數(shù)量的線(xiàn)程(thread)(線(xiàn)程數(shù)量在前期已經(jīng)過(guò)分析和計(jì)算),工作完成后關(guān)閉已啟用的線(xiàn)程,處理完畢的數(shù)據(jù)經(jīng)過(guò)特定的分發(fā)機(jī)制依照流程傳遞至下一個(gè)函數(shù)或者核函數(shù),不同核函數(shù)之間的線(xiàn)程完全獨(dú)立。

3.3 原始數(shù)據(jù)類(lèi)型轉(zhuǎn)換

相位干涉儀測(cè)向時(shí)需要兩個(gè)或多個(gè)天線(xiàn)接收信號(hào),每次采樣時(shí)在同一時(shí)刻的各個(gè)天線(xiàn)的信號(hào)按照通道號(hào)排列在一起,即多路信號(hào)AD采集后轉(zhuǎn)換為1路數(shù)據(jù)流發(fā)送至GPU顯存。本文的測(cè)向系統(tǒng)共設(shè)有4個(gè)天線(xiàn),CPU將4個(gè)天線(xiàn)接收的信號(hào)數(shù)據(jù)分割成每段16 000個(gè)采樣點(diǎn)拷貝給GPU,因此GPU第一個(gè)核函數(shù)(kernel 1)每次讀取16 000個(gè)點(diǎn)數(shù)據(jù),kernel 1的功能是將1路數(shù)據(jù)流的每段數(shù)據(jù)分成4路、每路4 000個(gè)點(diǎn)存放。

本文設(shè)計(jì)的GPU每個(gè)線(xiàn)程塊(block)中分配線(xiàn)程數(shù)為512,共啟用(≥1)個(gè)線(xiàn)程塊。kernel 1工作時(shí)所有線(xiàn)程塊內(nèi)的線(xiàn)程同時(shí)啟動(dòng),保證每個(gè)線(xiàn)程對(duì)每個(gè)采樣點(diǎn)數(shù)據(jù)同時(shí)進(jìn)行轉(zhuǎn)換,啟用的所有線(xiàn)程完成一次點(diǎn)數(shù)據(jù)轉(zhuǎn)換后,循環(huán)加載進(jìn)行下一段點(diǎn)數(shù)據(jù)類(lèi)型轉(zhuǎn)換。設(shè)本次處理的數(shù)據(jù)在第iW個(gè)線(xiàn)程塊,每個(gè)線(xiàn)程塊大小為uL,數(shù)據(jù)點(diǎn)在線(xiàn)程塊內(nèi)索引為iL,其在整個(gè)數(shù)據(jù)序列的索引值為=iW·uL+iL。由此實(shí)現(xiàn)了GPU中的原始數(shù)據(jù)類(lèi)型的轉(zhuǎn)換。

本文設(shè)計(jì)的原始數(shù)據(jù)類(lèi)型轉(zhuǎn)換核函數(shù)的CUDA代碼如下所示。

/*轉(zhuǎn)置函數(shù)的定義*/__global__ void transpose(float *odata, char *idata, int width, int height){int tid=threadIdx.x+blockIdx.x*blockDim.x;if(tid>=width)return;int lid=tid+blockIdx.y*width;int did=tid*height+blockIdx.y;odata[did]=idata[lid];}/*轉(zhuǎn)置函數(shù)的調(diào)用*/transpose<<>>(m_float2,m_raw,m_channelNum,m_samples);

3.4 FFT

在已有的FFT研究成果中,已知在 CPU 平臺(tái)下,點(diǎn)FFT的運(yùn)算量為2·log次的復(fù)數(shù)乘法和·log次的復(fù)數(shù)加法,表示做傅里葉變換的點(diǎn)數(shù)。本系統(tǒng)每次讀取16 000個(gè)點(diǎn)數(shù)據(jù),在CPU平臺(tái)下每次的計(jì)算量大小為16 000/2·log16 000次的復(fù)數(shù)乘法以及16 000·log16 000次的復(fù)數(shù)加法。對(duì)于如此大量的高速并行數(shù)據(jù),在CPU平臺(tái)中無(wú)法做到實(shí)時(shí)處理。針對(duì)FFT變化中大量高速數(shù)據(jù)的并行性這一顯著特點(diǎn),本文考慮將FFT部分放在GPU中進(jìn)行加速并行處理,利用CUDA中自帶的CUFFT庫(kù),可實(shí)現(xiàn)大量高速數(shù)據(jù)的FFT實(shí)時(shí)同步處理。

在上一步kernel 1完成原始數(shù)據(jù)類(lèi)型轉(zhuǎn)換后關(guān)閉所有線(xiàn)程,將數(shù)據(jù)發(fā)送給下一個(gè)函數(shù),調(diào)用CUFFT函數(shù)庫(kù)內(nèi)的cufftPlanD函數(shù)對(duì)數(shù)據(jù)進(jìn)行FFT。在GPU上進(jìn)行FFT需要完成以下工作。

(1) 用malloc(·)函數(shù)為輸入信號(hào)分配主機(jī)內(nèi)存,并存儲(chǔ)輸入信號(hào);

(2) 用cudaMalloc(·)為輸入信號(hào)分配GPU端顯存;

(3) 用cudaMemcpy(·)將主機(jī)內(nèi)存數(shù)據(jù)傳輸?shù)紾PU端顯存;

(4) 創(chuàng)建一個(gè)plan，調(diào)用cufftPlan創(chuàng)建FFT;

(5) 執(zhí)行plan，使用cufftExecR2C(·)函數(shù)完成 plan 的計(jì)算任務(wù);

(6) 用cudaMemcpy(·)將GPU端顯存數(shù)據(jù)傳輸?shù)街鳈C(jī)內(nèi)存;

(7) 執(zhí)行完成以后不再需要該plan,則調(diào)用cufftDestroy(·)函數(shù)銷(xiāo)毀該plan及為其分配的計(jì)算資源。

以上各步的具體流程和所用的CUDA函數(shù)如圖5所示。

圖5 GPU上進(jìn)行FFT的具體流程圖Fig.5 Specific flow chart of FFT on GPU

在CUFFT庫(kù)中,有3種不同的庫(kù)函數(shù):cufftExecC2C(復(fù)數(shù)到復(fù)數(shù)，簡(jiǎn)稱(chēng)為C2C)、cufftExecC2R(復(fù)數(shù)到實(shí)數(shù),簡(jiǎn)稱(chēng)為C2R)、cufftExecR2C(實(shí)數(shù)到復(fù)數(shù),簡(jiǎn)稱(chēng)為R2C),3種庫(kù)函數(shù)的計(jì)算耗時(shí)由于本身運(yùn)算量的差異而有所不同。因此,為最大限度地縮短運(yùn)算時(shí)間,必須針對(duì)本提相過(guò)程的實(shí)際情況,選擇最合適的庫(kù)函數(shù)類(lèi)型。在本系統(tǒng)輸入信號(hào)的FFT計(jì)算中,由于輸入是實(shí)信號(hào),而輸出是復(fù)數(shù)信號(hào),因此最宜采用R2C類(lèi)型的FFT。

實(shí)驗(yàn)表明,相比于使用C2C或R2C,本庫(kù)函數(shù)可以減少數(shù)據(jù)虛部初始化操作,有效減少數(shù)據(jù)傳輸量和計(jì)算量,節(jié)省運(yùn)算時(shí)間,便于更好地實(shí)現(xiàn)數(shù)據(jù)的實(shí)時(shí)同步處理。

本系統(tǒng)中天線(xiàn)數(shù)為4,cufftPlanD函數(shù)每次處理點(diǎn)數(shù)為,即對(duì)每個(gè)天線(xiàn)接收的數(shù)據(jù)每次處理4個(gè)采樣點(diǎn),為了減小噪音對(duì)信號(hào)的干擾,做FFT前將4個(gè)點(diǎn)等分成段,分段進(jìn)行FFT變換,每段4個(gè)點(diǎn),因此每個(gè)天線(xiàn)每段FFT后會(huì)得到2·4+1點(diǎn)數(shù)據(jù),段數(shù)據(jù)FFT后得到2·4+點(diǎn)復(fù)數(shù)數(shù)據(jù),即每個(gè)天線(xiàn)每次接收的個(gè)實(shí)數(shù)得到2·4+個(gè)復(fù)數(shù),因此4個(gè)天線(xiàn)FFT時(shí)每輸入個(gè)實(shí)數(shù)輸出2+4個(gè)復(fù)數(shù)。由此我們實(shí)現(xiàn)了基于GPU的FFT。

本文調(diào)用的CUFFT庫(kù)代碼如下所示。

/*直接調(diào)用CUDA自帶的cufft庫(kù),選擇cufftExecR2C 函數(shù)*/cufftExecR2C(m_fftHandle_batch,(cufftReal*)m_float2,(cufftComplex*)m_float2FFT);

3.5 共軛相乘

本文設(shè)計(jì)的共軛相乘的核函數(shù)部分的CUDA代碼如下。

/*共軛相乘函數(shù)的定義/__global__ void Conj(float2*src,float2*send,float2*dst,int length){ int tid=threadIdx.x+blockIdx.x*blockDim.x; if(tid>=length) return; int lid=tid+blockIdx.y*length; float signalx=src[lid].x; float signaly=src[lid].y; dst[lid].x=signalx*send[lid].x+signaly*send[lid].y; dst[lid].y=signaly*send[lid].x-signalx*send[lid].y;}/*共軛相乘函數(shù)的調(diào)用(以21該路信號(hào)的處理為例)*/Conj<<>>(m_float2FFT+3*m1_sam-ples, m_float2FFT+m1_samples, m_conj21, NFFT);

3.6 能量累積并求相位

共軛相乘后的數(shù)據(jù)傳輸至kernel 3進(jìn)行能量累積及提相等處理,能量累積即將每條基線(xiàn)的2·4+個(gè)數(shù)據(jù)等分成段進(jìn)行相加平均,累加平均后使用atan2函數(shù)求出每條基線(xiàn)的(2·4+)個(gè)復(fù)數(shù)的相角,所得相角即為每個(gè)采樣點(diǎn)處兩路信號(hào)相位差。接著,在所有相角中,找到最大值處的相位信息,由此完成了GPU中的整個(gè)提相算法過(guò)程。

本文設(shè)計(jì)的能量累積的核函數(shù)部分的CUDA代碼如下所示。

/*能量累積函數(shù)定義*/___global__ void FAdd(float2* src,float* dst,float* dst2,int jft,int length){ int tid=threadIdx.x+blockIdx.x*blockDim.x; if(tid>=length) return; int curpos; float2 sum={0}; for(int i=0;i

/*能量累積函數(shù)調(diào)用,并求相位*/FAdd<<>>(m_conj21,m_phase21,m_amp21,m_jft,NFFT);/*取最大值處相位(以其中一路信號(hào)的處理為例)*/cudaMemcpy(m_ampHost21,m_amp21,501*sizeof(float),cu-daMemcpyDeviceToHost);for(int kk=0;kk<(m_nfft/2);kk++) { if(m_ampHost21[kk]>m_maxVal) { m_maxVal=m_ampHost21[kk]; m_maxIdx=kk; } }

4 實(shí)驗(yàn)驗(yàn)證與分析

4.1 系統(tǒng)及核心參數(shù)設(shè)計(jì)

圖6 本文設(shè)計(jì)的相位干涉儀布陣示意圖Fig.6 Element arrangement of phase interferometer

每條基線(xiàn)的長(zhǎng)度如表1所示。

表1 相位干涉儀各條基線(xiàn)長(zhǎng)度Table 1 Length of each of phase interferometer baseline

系統(tǒng)硬件平臺(tái)選用NVIDIA Tesla V100顯卡,接收信號(hào)頻段采用測(cè)控常用S頻段,系統(tǒng)采樣率為50 MHz。本系統(tǒng)的核心設(shè)計(jì)指標(biāo)如表2所示。

表2 系統(tǒng)總體設(shè)計(jì)指標(biāo)Table 2 Overall system design index MHz

AD采集后的數(shù)據(jù)進(jìn)入鑒相器進(jìn)行FFT時(shí),每個(gè)通道信號(hào)在每個(gè)積分時(shí)間(0.4 ms)內(nèi)處理2×10個(gè)采樣點(diǎn)的數(shù)據(jù),積分后只取相關(guān)譜最大值處對(duì)應(yīng)的相位值,因此每個(gè)通道每2×10個(gè)采樣點(diǎn)求得一個(gè)相位結(jié)果,由此每個(gè)通道的信號(hào)每秒將得到2 500個(gè)相位結(jié)果。

為充分驗(yàn)證本文提出的基于GPU的鑒相算法的優(yōu)越性,本文將該鑒相系統(tǒng)在CPU和GPU平臺(tái)下分別進(jìn)行了鑒相精度和鑒相速度的測(cè)試,并根據(jù)測(cè)試結(jié)果分析了鑒相標(biāo)準(zhǔn)差和鑒相加速比。

4.2 鑒相精度驗(yàn)證與分析

本文分別基于CPU和GPU平臺(tái),對(duì)鑒相算法進(jìn)行了多次長(zhǎng)時(shí)間的實(shí)測(cè)數(shù)據(jù)驗(yàn)證,得到了如下所示的鑒相結(jié)果。圖7所示為6條基線(xiàn)的1 s數(shù)據(jù)在GPU平臺(tái)下的鑒相結(jié)果。圖8所示為6條基線(xiàn)的1 s數(shù)據(jù)在CPU平臺(tái)下鑒相結(jié)果。

圖7 GPU鑒相算法的鑒相結(jié)果Fig.7 Phase identification result based on GPU phase identification algorithm

圖8 傳統(tǒng)鑒相算法的鑒相結(jié)果Fig.8 Phase identification result of traditional phase identification algorithms

由圖7和圖8可知,兩個(gè)平臺(tái)下6條基線(xiàn)提取的相位差的翻折次數(shù)與基線(xiàn)長(zhǎng)度的關(guān)系均滿(mǎn)足式(6),即基線(xiàn)越長(zhǎng),相位翻折次數(shù)越多。

為驗(yàn)證基于GPU的鑒相算法的準(zhǔn)確性,本文在相同的測(cè)試條件下,針對(duì)同樣的實(shí)測(cè)數(shù)據(jù),對(duì)比了兩種算法的鑒相結(jié)果,得到了如圖9所示的同一次測(cè)試中兩平臺(tái)下6條基線(xiàn)鑒相結(jié)果的誤差圖。由圖9可知,各條基線(xiàn)的鑒相誤差基本均維持在[(-2×10)°,(2×10)°]內(nèi)。

圖9 兩種鑒相算法在各基線(xiàn)的鑒相結(jié)果標(biāo)準(zhǔn)差Fig.9 Standard deviations of phase dientification results of two phase identification algorithms shown at each baseline

圖10所示為兩個(gè)平臺(tái)下分別進(jìn)行10次系統(tǒng)測(cè)試之后鑒相誤差的標(biāo)準(zhǔn)差,6條基線(xiàn)的鑒相標(biāo)準(zhǔn)差始終小于2×10°,可知兩種平臺(tái)下的鑒相效果相當(dāng),從而充分驗(yàn)證了本文設(shè)計(jì)的GPU平臺(tái)下的FX鑒相算法的準(zhǔn)確性。

圖10 兩個(gè)平臺(tái)連續(xù)10次測(cè)試的鑒相結(jié)果標(biāo)準(zhǔn)差Fig.10 Standard deviation of phase results of ten consecutive tests on two platforms

4.3 加速效果驗(yàn)證與分析

為了進(jìn)一步研究?jī)煞N算法提取相位差時(shí)各個(gè)關(guān)鍵運(yùn)算單元及整個(gè)系統(tǒng)的耗時(shí)情況,本文利用NVIDIA提供的分析工具NVVP軟件,針對(duì)1 s數(shù)據(jù)對(duì)GPU鑒相算法中關(guān)鍵運(yùn)算單元的執(zhí)行性能和整個(gè)鑒相算法的時(shí)間分配進(jìn)行了分析和測(cè)試。全過(guò)程的數(shù)據(jù)傳輸和核函數(shù)運(yùn)算時(shí)間的NVVP分析結(jié)果如圖11所示。

圖11 基于NVVP的GPU鑒相算法全流程模塊調(diào)用情況分析Fig.11 Analysis of the whole process module call of GPU phase identification algorithm based on NVVP

從圖11可以看出,cudaMalloc(申請(qǐng)?jiān)O(shè)備內(nèi)存)、cudaMemcpy(數(shù)據(jù)傳輸)、Compute(執(zhí)行關(guān)鍵運(yùn)算單元)、cudaFree(釋放設(shè)備內(nèi)存)、cudaDeviceReset(釋放和清空GPU中的資源)等模塊依次執(zhí)行,計(jì)算1 s數(shù)據(jù)的相位差總用時(shí)895 ms,滿(mǎn)足系統(tǒng)實(shí)時(shí)性的要求。其中,并行化程度最高,運(yùn)算量最大的Compute部分,僅耗時(shí)285 ms,占總鑒相過(guò)程的31.84%。Compute中原始數(shù)據(jù)轉(zhuǎn)換(transpose)、FFT、共軛相乘(Conj)、能量累積并求相位(FAdd)這4個(gè)模塊并行執(zhí)行,充分利用了GPU各線(xiàn)程塊之間的并行化特性,極大地減少了時(shí)間消耗。

接著選取并行化程度最高的Compute部分作重點(diǎn)分析,將圖11中4個(gè)關(guān)鍵運(yùn)算單元在處理數(shù)據(jù)時(shí)的總時(shí)長(zhǎng)和各個(gè)單元的用時(shí)占比,以及CPU平臺(tái)下鑒相系統(tǒng)的用時(shí)情況如表3、圖12和圖13所示。在CPU、GPU執(zhí)行同一任務(wù)的條件下,定義在CPU平臺(tái)下的運(yùn)行時(shí)間為,在GPU平臺(tái)下的運(yùn)行時(shí)間為,那么加速比=/。

表3 兩種算法關(guān)鍵運(yùn)算單元的性能比較Table 3 Performance comparison of the key operation units of two algorithms

圖12 GPU鑒相算法中關(guān)鍵運(yùn)算單元在GPU上的耗時(shí)情況Fig.12 Time-consuming proportion of the key computing units in the GPU phase identification algorithm

圖13 兩種鑒相算法關(guān)鍵運(yùn)算單元的耗時(shí)情況對(duì)比Fig.13 Comparison of time-consuming conditions of the key computing units two phase identification algorithms

通過(guò)對(duì)比各個(gè)單元的加速比可知,GPU平臺(tái)下對(duì)于FFT處理加速較快。從圖12中可見(jiàn)，運(yùn)算量最大的FFT在Compute的總耗時(shí)中僅占到了9%。圖13則充分說(shuō)明了調(diào)用CUFFT函數(shù)庫(kù)以及采用GPU進(jìn)行并行算法加速的巨大優(yōu)勢(shì)。這充分證明了GPU的加速效果及其對(duì)于并行化大規(guī)模數(shù)據(jù)量計(jì)算的優(yōu)越性。

表4及圖14所示為兩個(gè)平臺(tái)分別進(jìn)行10次系統(tǒng)測(cè)試的鑒相時(shí)長(zhǎng)以及兩個(gè)平臺(tái)的加速比,由此可知，相比于傳統(tǒng)算法,本文涉及的基于GPU的鑒相算法可以加速140倍左右,充分證明了GPU平臺(tái)下的鑒相算法具有較為強(qiáng)大的實(shí)時(shí)化、高速處理數(shù)據(jù)能力。

表4 兩種算法連續(xù)10次的總運(yùn)行時(shí)間測(cè)試結(jié)果Table 4 Total running time test results of two algorithms for ten consecutive times

圖14 兩種算法在10次測(cè)試下的總運(yùn)行時(shí)間對(duì)比Fig.14 Comparison of total running time of two algorithms for ten tests

5 結(jié) 論

為了實(shí)現(xiàn)鑒相過(guò)程的高速、實(shí)時(shí)、可靠,本文研究了提相算法原理,分析了算法并行性并設(shè)計(jì)了基于GPU的實(shí)時(shí)信號(hào)鑒相算法。針對(duì)GPU的加速效果進(jìn)行分析,研究了GPU的鑒相性能并對(duì)結(jié)果進(jìn)行分析,本文主要所做工作如下。

(1) 設(shè)計(jì)了基于CUDA架構(gòu)的鑒相器提相算法。通過(guò)對(duì)鑒相算法的分析,將模塊運(yùn)算流程分為原始數(shù)據(jù)轉(zhuǎn)換、FFT、共軛相乘、能量累積并求相位等環(huán)節(jié),分析了各環(huán)節(jié)算法的并行性并基于CUDA架構(gòu)設(shè)計(jì)了并行實(shí)現(xiàn)的具體算法及流程;最后用實(shí)驗(yàn)對(duì)算法進(jìn)行了驗(yàn)證,證明了算法的正確性。

(2) 研究了GPU平臺(tái)的加速效果。通過(guò)對(duì)并行運(yùn)算單元中關(guān)鍵運(yùn)算單元以及整個(gè)鑒相過(guò)程的研究,分析了GPU平臺(tái)下鑒相算法的顯著優(yōu)勢(shì)與較為強(qiáng)大的實(shí)時(shí)化、高速處理能力,基于NVVP分析了GPU鑒相算法流程內(nèi)部的具體線(xiàn)程分配與模塊調(diào)用情況,證明了算法的優(yōu)越性。