基于CUDA架構(gòu)的AMTI算法實(shí)現(xiàn)

金 莉，王仁志，宋萬(wàn)杰

(西安電子科技大學(xué) 雷達(dá)信號(hào)處理國(guó)家重點(diǎn)實(shí)驗(yàn)室，陜西 西安710071)

雷達(dá)所處的大環(huán)境中存在各種雜波，與目標(biāo)信號(hào)相比，雜波往往具有更大的功率，所以雷達(dá)信號(hào)處理的關(guān)鍵在于從雜波背景中檢測(cè)出目標(biāo)。雷達(dá)信號(hào)處理系統(tǒng)中，自適應(yīng)動(dòng)目標(biāo)顯式(AMTI)因?yàn)閷?duì)靜雜波和動(dòng)雜波都有抑制能力，成為常用的雜波抑制方式。實(shí)現(xiàn)AMTI技術(shù)的方法主要有兩種：一種是兩級(jí)動(dòng)目標(biāo)顯示級(jí)聯(lián)形式實(shí)現(xiàn)，讓第一級(jí)凹口位于多普勒零頻處，抑制靜止雜波；第二級(jí)設(shè)計(jì)凹口抑制動(dòng)雜波；另一種是設(shè)計(jì)雙凹口雜波抑制濾波器[1]。

本文采用NVIDIA的圖形處理單元(Graphics Processing Unit ,GPU)來(lái)實(shí)現(xiàn)AMTI算法。自GPU概念提出以來(lái)，GPU峰值性能一直以超過(guò)摩爾定律的速度增加，平均每6個(gè)月翻一番[2-3]。相對(duì)CPU而言，GPU具有浮點(diǎn)計(jì)算能力強(qiáng)、性?xún)r(jià)比高、能耗低等優(yōu)點(diǎn)，被廣泛用于通用計(jì)算中。特別是統(tǒng)一設(shè)備架構(gòu)(Compute Unified Device Architecture ,CUDA)的發(fā)布，將圖形硬件和應(yīng)用接口程序進(jìn)行封裝，使開(kāi)發(fā)人員可以直接采用C語(yǔ)言進(jìn)行編程，方便易學(xué)[4]。以GPU為平臺(tái)，采用CUDA技術(shù)，在單指令多數(shù)據(jù)(Single Instruction Multiple Data，SIMD)編程模型下實(shí)現(xiàn)AMTI算法，發(fā)揮GPU強(qiáng)大的計(jì)算能力，提高運(yùn)算效率。

1 AMTI算法

為了濾除雜波，需要在雜波信號(hào)位置形成凹口。因?yàn)閯?dòng)雜波不位于零頻附近，所以可以在動(dòng)雜波譜位置形成凹口，也可以將動(dòng)雜波移動(dòng)至凹口附近來(lái)濾除動(dòng)雜波。即有兩種方法實(shí)現(xiàn)AMTI技術(shù)：一種是將雜波信號(hào)譜中心移到零頻凹口處，然后采用MTI濾波器進(jìn)行濾波處理；另一種是將抑制雜波的凹口移動(dòng)至雜波譜中心直接實(shí)現(xiàn)濾波，這種方法需要事先設(shè)計(jì)凹口位于不同頻率點(diǎn)的濾波器，并且將濾波器權(quán)系數(shù)存儲(chǔ)在一個(gè)庫(kù)里，稱(chēng)為權(quán)系數(shù)庫(kù)法。在工程中常使用這種方法，首先設(shè)計(jì)出凹口位于零頻處的MTI濾波器，然后在頻率軸平移得到濾波器權(quán)系數(shù)庫(kù)，當(dāng)估計(jì)到雜波頻譜中心后，直接調(diào)用對(duì)應(yīng)的濾波器進(jìn)行對(duì)消處理[5-6]。本文采用基于權(quán)系數(shù)庫(kù)實(shí)現(xiàn)的AMTI算法，圖1為該算法的原理框圖。

圖1 基于權(quán)庫(kù)法的AMTI原理框圖

基于權(quán)系數(shù)庫(kù)的AMTI算法的思想是：先估計(jì)動(dòng)雜波的多普勒范圍f∈(fdmin,fdmax)。然后，在這個(gè)估計(jì)的多普勒范圍內(nèi)，等間隔的設(shè)計(jì)一組MTI雜波抑制濾波器，最后將該雜波抑制濾波器的濾波器系數(shù)直接放入存儲(chǔ)設(shè)備中備用[7-8]。具體實(shí)現(xiàn)過(guò)程中，首先需要確定調(diào)用的MTI濾波器，具體方法為：先估計(jì)雜波的多普勒中心頻率f，然后依據(jù)f調(diào)用對(duì)應(yīng)的MTI雜波抑制濾波器濾波。因此，基于權(quán)系數(shù)庫(kù)的AMTI濾波器的設(shè)計(jì)重點(diǎn)在于準(zhǔn)確估計(jì)動(dòng)雜波的多普勒中心頻率f。在實(shí)際的工程實(shí)現(xiàn)中，動(dòng)雜波多普勒中心頻率f通過(guò)計(jì)算相鄰脈沖間相位差得到。

考慮到雷達(dá)回波信號(hào)x(t)的表達(dá)式為

x(t)=A(t)ej(wdt+φ0)+n(t)

(1)

式中，A(t)為回波信號(hào)的幅度；wd為雜波的多普勒頻率，并有wd=2πfd；w0表示初回波信號(hào)的初始相位；加性噪聲n(t)在不同脈沖間互不相關(guān)。假設(shè)Tr為脈沖的重復(fù)周期，那么雷達(dá)接收到的下一個(gè)Tr的發(fā)射脈沖回波為

x(t-Tr)=A(t-Tr)ej(wd(t-Tr)+φ0)+n(t-Tr)

(2)

那么，兩次回波信號(hào)的相關(guān)函數(shù)可表示為

Rxx=E[x(t)x*(t-Tr)]=E[A(t)A(t-Tr)]ejwdTr

(3)

式(3)中A(t)為窄帶信號(hào)，因此，可以得出式E[A(t)A(t-Tr)]=E[|A(t)|2]的結(jié)果為實(shí)數(shù)。因此，雜波的多普勒頻率wd為

(4)

進(jìn)一步

(5)

式(5)中，i表示不同脈沖回波信號(hào)序列。采用相鄰兩個(gè)脈沖間相位差的方法，得到f表達(dá)式為

(6)

為了準(zhǔn)確估計(jì)雜波多普勒中心頻率，實(shí)際工程中常采用質(zhì)心法，質(zhì)心法就是將回波信息中的幅度值也考慮在內(nèi)，多次估計(jì)的雜波多普勒中心頻率，最后進(jìn)行平滑處理[9]。進(jìn)一步提高f估計(jì)的準(zhǔn)確性。

2 CUDA編程模型

CUDA編程模型將CPU作為主機(jī)(Host),GPU作為協(xié)處理器(Co-processor)或者設(shè)備(Device)。在一個(gè)系統(tǒng)中可以存在若干個(gè)主機(jī)和設(shè)備[10]。在CUDA編程模中，CPU和GPU協(xié)同工作，CPU主要負(fù)責(zé)整個(gè)串行運(yùn)算的執(zhí)行，控制算法執(zhí)行流程，初始化數(shù)據(jù)將其傳遞到GPU，GPU接收CPU傳遞過(guò)來(lái)的數(shù)據(jù)，進(jìn)行并行計(jì)算，計(jì)算完成后，將計(jì)算結(jié)果傳遞回主機(jī)[11-12]。另外，GPU和CPU擁有各自獨(dú)立的存儲(chǔ)空間，CPU端為主機(jī)內(nèi)存，GPU端為顯存，GPU對(duì)顯存的操作一般與CPU對(duì)內(nèi)存操作一樣，需要調(diào)用CUDA API中的內(nèi)存管理函數(shù)。

運(yùn)行在GPU上的函數(shù)稱(chēng)為核函數(shù)(Kernel),核函數(shù)并不是一個(gè)完整的程序，而是整個(gè)CUDA程序中一個(gè)可以被并行執(zhí)行的步驟[13-15]。如圖2所示，一個(gè)完整的CUDA程序是由主機(jī)端的串行代碼和設(shè)備端的并行函數(shù)Kernel(·)共同組成，各個(gè)并行函數(shù)之間為串行關(guān)系。

圖2 CUDA并行編程模型

3 AMTI算法實(shí)現(xiàn)流程

為滿(mǎn)足雷達(dá)系統(tǒng)要求，本算法采用5參差A(yù)MTI，濾波器間隔是5 Hz，從0～90 Hz共19組AMTI雙凹口濾波器，圖3為AMTI算法實(shí)現(xiàn)流程圖。

圖3 基于權(quán)系數(shù)庫(kù)的AMTI算法實(shí)現(xiàn)流程圖

基于CUDA架構(gòu)的AMTI算法實(shí)現(xiàn)，主要流程如下：

步驟1在CPU上進(jìn)行數(shù)據(jù)初始化，內(nèi)存以及顯存分配，將準(zhǔn)備就緒的數(shù)據(jù)通過(guò)CUDA API中的拷貝函數(shù)拷貝至GPU；

步驟2首先在GPU上完成雜波譜中心f的估計(jì)，并判斷頻率值的正負(fù)值，將該值存放在單獨(dú)數(shù)組中；

步驟3接下來(lái)在GPU上進(jìn)行AMTI處理，具體處理流程如圖3所示，最終實(shí)現(xiàn)抑制雜波的目的；

步驟4再次調(diào)用CUDA API中的拷貝函數(shù)就計(jì)算結(jié)果從設(shè)備拷貝回主機(jī)；

步驟5在CPU上釋放分配的內(nèi)存以及顯存。

圖4為采用CUDA編程架構(gòu)的AMTI算法實(shí)現(xiàn)流程圖。

GPU上實(shí)現(xiàn)AMT處理，主要是抽取AMTI中的并行部分，分配到各線程塊中的線程進(jìn)行執(zhí)行。就雜波譜中心估計(jì)來(lái)說(shuō)，包含脈沖數(shù)和數(shù)據(jù)量?jī)蓪友h(huán)，但GPU中只要合理分配線程及線程塊，讓每個(gè)線程執(zhí)行一個(gè)脈沖一組數(shù)據(jù)的運(yùn)算，就可以大幅提高運(yùn)算效率。本實(shí)現(xiàn)中，分別用x維線程控制數(shù)據(jù)量，y維線程控制脈沖數(shù)，運(yùn)算完成后批量增加循環(huán)過(guò)程中的累加值，代碼如下：

int tidx=threadIdx.x+blockIdx.x*blockDim.x;∥數(shù)據(jù)量

int tidy=threadIdx.y+blockIdx.y*blockDim.y;∥脈沖數(shù)

while(tidx

while(tidy

{

∥計(jì)算得到濾波器系數(shù)相關(guān)地址及符號(hào)

tidy+=blockDim.y*gridDim.y; 距離單元/個(gè)

}

∥計(jì)算系數(shù)偏移地址去除系數(shù)完成AMTI處理

tidx+=blockDim.x*gridDim.x; ∥批量并行運(yùn)算

}

需要說(shuō)明的是，在AMTI實(shí)現(xiàn)過(guò)程中，需要計(jì)算反正切值。反正切值是通過(guò)查表法得到。事先計(jì)算好得到每個(gè)角度對(duì)應(yīng)的正切值，存放在內(nèi)存中以備使用。為了提高正切值的精度，可以采用先乘數(shù)再取整的方法。具體實(shí)現(xiàn)方法為，將求得的值首先乘以一個(gè)倍數(shù)，然后進(jìn)行取整，放入表格中，根據(jù)具體雷達(dá)系統(tǒng)對(duì)數(shù)據(jù)精度的要求，確定該倍數(shù)大小。

4 仿真結(jié)果及分析

仿真采用的操作系統(tǒng)為Windows 7 SP1，顯卡是計(jì)算能力為2.0的NVIDIA Tesla C2050，該顯卡支持雙精度浮點(diǎn)運(yùn)算，滿(mǎn)足雷達(dá)系統(tǒng)對(duì)精度的要求。為了提高運(yùn)算精度，數(shù)據(jù)處理均采用雙精度浮點(diǎn)數(shù)，圖5為AMTI前脈壓數(shù)據(jù)。

圖5 脈壓原始數(shù)據(jù)

實(shí)驗(yàn)分別在CPU和GPU上實(shí)現(xiàn)AMTI算法，其結(jié)果分別如圖6和圖7所示。

圖6 CPU平臺(tái)實(shí)現(xiàn)AMTI

從圖6和圖7可以看到濾除了地雜波，目標(biāo)幅度達(dá)到103數(shù)量級(jí)，可以認(rèn)為雜波抑制效果良好且CPU和GPU處理結(jié)果一致，但兩者的處理速度有很大差別。

圖7 GPU平臺(tái)實(shí)現(xiàn)AMT

表1 CPU與GPU所運(yùn)行時(shí)間對(duì)比/ms

從表1可以看出， GPU比CPU的數(shù)據(jù)運(yùn)算效率更高。需要說(shuō)明的是，表格中括號(hào)內(nèi)數(shù)字表示用于主機(jī)和設(shè)備之間數(shù)據(jù)傳輸?shù)臅r(shí)間，例如AMTI算法中，總時(shí)間0.645 ms中僅0.02 ms用于數(shù)據(jù)處理，括號(hào)中為用于數(shù)據(jù)傳輸?shù)臅r(shí)間，即GPU有很大一部分時(shí)間都是用在數(shù)據(jù)傳輸上，所以在GPU運(yùn)算過(guò)程中，應(yīng)當(dāng)盡量減數(shù)據(jù)傳輸時(shí)間。

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

實(shí)驗(yàn)表明，通過(guò)在GPU上實(shí)現(xiàn)AMTI算法，不僅達(dá)到了和CPU上一樣的效果，而且加快了運(yùn)算速度，提高了運(yùn)算效率。但是在CUDA程序中，主機(jī)與設(shè)備之間的數(shù)據(jù)傳輸會(huì)耗費(fèi)很多時(shí)間，所以可以采用以下3種方法減少傳輸時(shí)間：(1)由于在設(shè)備和主機(jī)之間進(jìn)行數(shù)據(jù)傳輸會(huì)耗費(fèi)相當(dāng)多的時(shí)間，所以應(yīng)該在兩次數(shù)據(jù)傳輸之間在GPU上進(jìn)行盡可能多的運(yùn)算，或者采用流的處理方式，將數(shù)據(jù)進(jìn)行分塊處理，傳輸?shù)耐瑫r(shí)進(jìn)行運(yùn)算；(2)因?yàn)榧y理存儲(chǔ)器的存儲(chǔ)方式為只讀，處理速度相對(duì)其他存儲(chǔ)器速度較快，可以采用紋理存儲(chǔ)器來(lái)保存一些反三角函數(shù)表，濾波器系數(shù)等，以?xún)?yōu)化數(shù)據(jù)處理效率；(3)合理安排線程塊以及線程數(shù)量，例如當(dāng)使用二維線程塊以及線程索引時(shí)，可以利用dim3這個(gè)變量進(jìn)行定義，最好以一個(gè)線程束的倍數(shù)定義線程和線程塊數(shù)量來(lái)最大限度提高運(yùn)算速率?？偟膩?lái)說(shuō)，利用CUDA處理數(shù)據(jù)相比CPU而言，較好地滿(mǎn)足了目前信息處理對(duì)高精度、實(shí)時(shí)性要求。