孫寬飛，楊文革，滕 飛，焦義文，高澤夫

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

0 引言

飛行器遙測是導(dǎo)彈、火箭和衛(wèi)星等航天器試驗和運行過程中必不可少的支持系統(tǒng)，能夠?qū)崟r監(jiān)測航天器內(nèi)部工作狀態(tài)、電氣性能和環(huán)境參數(shù)等重要信息，為航天器性能檢測、效能評估及故障分析提供依據(jù)[1]。隨著航天事業(yè)的深入發(fā)展，全球各個航天大國對宇宙的探索和對宇宙資源的爭奪愈加強烈[2]，遙測技術(shù)突顯出更為重要的作用。

脈沖編碼調(diào)制/調(diào)頻(Pulse Code Modulation/Frequency Modulation，PCM/FM)技術(shù)具有較強的抗尾焰效應(yīng)能力、抗噪聲性能強、抗多徑衰落、抗相位和干擾發(fā)射機功率高等特點[3]，成為國內(nèi)外航空航天遙測領(lǐng)域長期采用的一種主流體制[4-5]。近年來，隨著航天事業(yè)和國防工業(yè)的快速發(fā)展，對遙測技術(shù)的要求也越來越高。具體來說是遙測目標越來越多，需要測量的參數(shù)和遙測信號傳輸?shù)木嚯x也在逐漸增加，這就需要遙測數(shù)據(jù)的傳輸速率更快、數(shù)據(jù)傳輸?shù)目煽啃愿遊6]。傳輸速率的提高和遙測距離的增加會加大接收端數(shù)據(jù)解調(diào)的壓力，造成解調(diào)性能的可靠性降低。

為解決以上問題，有學(xué)者在研究頻譜效率更高的遙測體制，如Multi-h CPM體制[7]和FQPSK體制[8]。但是這些新型的體制復(fù)雜度高、技術(shù)尚不成熟，還停留在實驗室階段，所以現(xiàn)在和將來的一段時間內(nèi)PCM/FM體制還是遙測應(yīng)用的主要體制[9]。因此，提高PCM/FM遙測信號的解調(diào)性能的研究一直是熱點[10-12]，并且提出了很多提高性能的方法。其中，Mark Geoghegan[13-14]將多符號檢測技術(shù)引入PCM/FM遙測信號的解調(diào)中，可以提高約3 dB的解調(diào)增益，在實際工程中獲得廣泛應(yīng)用。

傳統(tǒng)MSD均采用現(xiàn)場可編程邏輯門陣列(Field Programmable Gate Array，F(xiàn)PGA)實現(xiàn)系統(tǒng)功能，存在硬件設(shè)計周期較長、硬件平臺通用性低、升級難度大等問題，距離軟件無線電設(shè)計理念差距較大。隨著高性能計算技術(shù)的發(fā)展，圖形處理器件(Graphic Processing Unit,GPU)從專用于圖像領(lǐng)域的處理器逐漸向著通用并行計算平臺轉(zhuǎn)變[15]。GPU具有的運算核心數(shù)遠多于CPU，比較適合用于數(shù)據(jù)密集型計算的并行加速處理[16]。2007年NVIDIA推出的計算統(tǒng)一設(shè)備架構(gòu)(Compute Unified Device Architecture,CUDA)允許開發(fā)者使用C語言進行編程，簡化了GPU系統(tǒng)的開發(fā)流程，降低了使用GPU并行編程的難度，使得GPU通用計算技術(shù)在信號處理領(lǐng)域得到更為廣泛的應(yīng)用[17]?；贑PU+GPU的信號處理系統(tǒng)成為眾多領(lǐng)域的研究熱點，如雷達[15,18]、射電天文[19-20]和無線電通信[21-22]等。

本文在通用計算機上實現(xiàn)多符號檢測算法，用GPU進行加速處理以助力實現(xiàn)實時解調(diào)?，F(xiàn)有的在硬件系統(tǒng)中實現(xiàn)的MSD算法不能直接移植到GPU上，而在GPU上實現(xiàn)MSD算法的研究較少且不夠深入，如成亞勇[23]和李梓博[24]在GPU上實現(xiàn)了MSD算法但卻不能實現(xiàn)實時解調(diào)。本文首先對MSD算法串行運算過程進行分析，得出算法核心為滑動相關(guān)運算，設(shè)計了基于GPU的滑動相關(guān)運算，依據(jù)并行滑動相關(guān)的實現(xiàn)方法設(shè)計了基于GPU的多符號檢測算法；實現(xiàn)精確的位同步是多符號檢測運算提高解調(diào)增益的前提，因此，設(shè)計了基于GPU的Gardner位同步算法，并對算法進行了驗證。

1 多符號檢測算法原理及計算量分析

多符號檢測算法又可分為中頻多符號檢測和基帶多符號檢測2種，在中頻直接進行多符號檢測處理速率較高，具有較大的難度且會造成計算資源的浪費。如果將PCM/FM遙測中頻信號搬移到基帶進行多符號檢測，處理速率降低更容易實現(xiàn)，所以本文采用基帶多符號檢測算法。

基帶多符號檢測算法實現(xiàn)時，需要將PCM/FM遙測信號從中頻搬移到基帶，用到數(shù)字下變頻、FIR濾波、抽取，本文不再對以上步驟進行闡釋，而只對多符號檢測算法的實現(xiàn)進行研究?；鶐Ф喾枡z測的實現(xiàn)框圖如圖1所示。

圖1 基帶多符號檢測原理框圖

接收到的中頻信號經(jīng)過下變頻等一系列的處理之后，得到同相分量I(t)與正相分量Q(t)：

(1)

式中，kf為調(diào)制頻偏；f(τ)為調(diào)制信號；θ1為初始相位。同相分量I和正相分量Q的表達式可寫為數(shù)字域的表達形式：

(2)

式(2)可以寫成復(fù)信號的形式：

R(nTs)=I(nTs)+jQ(nTs)=

(3)

式中，R(nTs)表示接收到的基帶復(fù)信號。

進行多符號檢測時，本地參考信號為觀測長度內(nèi)碼元序列的所有組合，設(shè)本地參考信號組合中和接收到的基帶復(fù)信號相同的一組序列為L(nTs)，則L(nTs)可以表示為：

(4)

式中，θ2表示本地參考復(fù)信號的初始相位。

將本地參考復(fù)信號與接收到的基帶復(fù)信號R(nTs)進行相關(guān)運算，本地參考復(fù)信號為L(nTs)的相關(guān)結(jié)果為：

R(nTs)×L(nTs)=cos(θ1-θ2)+jsin(θ1-θ2)=

IB+jQB。

(5)

由式(5)可以看出，用和接收到的基帶復(fù)信號相同的一組序列L(nTs)與基帶復(fù)信號R(nTs)進行相關(guān)運算的結(jié)果與nTs無關(guān)，是一個常數(shù)。而其他的情況下，由于與基帶復(fù)信號不同，相關(guān)運算后的結(jié)果均與nTs有關(guān)。

假設(shè)觀測時段內(nèi)的采樣點數(shù)為Nc，對式(5)在觀測時間內(nèi)進行積分，得到：

M=Nc×IB+j×Nc×QB。

(6)

對式(6)得到的復(fù)信號進行模平方運算，得到：

S=Nc2×(IB2+QB2)。

(7)

式(7)表示的是在觀測時段內(nèi)，和接收到的基帶復(fù)信號相同的一組序列L(nTs)與接收到的基帶正交復(fù)信號進行相關(guān)和平方運算后的取值。其他情況下本地信號與接收到的基帶信號相關(guān)之后得到的復(fù)信號都是時變的，因此，平方之后得到的實數(shù)值也都小于式(7)的計算結(jié)果。因此，可以通過比較平方后的結(jié)果判決接收到的基帶信號的取值，即模平方最大的一組本地參考信號的中間一位的值。

為了方便，圖1僅給出第k個本地復(fù)信號與接收到的正交基帶復(fù)信號進行處理的框圖，其他組合的本地信號與其類似?；鶐Ф喾枡z測算法流程大致如下，各個本地信號分別與接收到的正交基帶復(fù)信號進行相關(guān)平方處理，然后對各路輸出進行比較判決。

基于以上多符號檢測算法的推導(dǎo)過程可以看出，多符號檢測算法是求相關(guān)運算并求平方比較大小的計算過程，設(shè)每段信號有N個碼元，每個碼元有Ns個采樣點，觀測長度為L，則每段信號要進行N-L+1次多符號檢測運算，每次運算有L×Ns個采樣點參與。

在一次運算過程中，求相關(guān)運算有M1次乘法和A1次加法，則M1和A1分別為：

M1=4×L×Ns×2L，

(8)

A1=(4×L×Ns-2)×2L。

(9)

在一次運算過程中，求模平方運算有M2次乘法和A2次加法，則M2和A2分別為：

M2=2×2L，

(10)

A2=2L。

(11)

每次多符號檢測運算共有M次乘法和A次加法，則M和A分別為：

M=M1+M2=4×L×Ns×2L+2×2L=

(4×L×Ns+2)×2L,

(12)

A=A1+A2=(4×L×Ns-2)×2L+2L=

(4×L×Ns-1)×2L。

(13)

多符號檢測每次運算過程需要(4×L×Ns-1)×2L次加法計算，(4×L×Ns+2)×2L次乘法計算。以56 MHz采樣率、2 Mb/s碼速率、觀測長度為5、每個碼元8個采樣點為例，每次運算需要207 360次乘法計算和203 520次加法計算，1 s數(shù)據(jù)需要進行2×106次多符號檢測運算，還需要將每次乘加計算量乘以2×106才能得到1 s數(shù)據(jù)的計算量。

可以看出，多符號檢測運算計算量巨大，串行運算需要消耗大量時間。GPU比較擅長處理計算密集型任務(wù)，性能達到了TFLOPS級別，并且具有大量的算術(shù)單元，可以滿足多符號檢測過程中的各種算術(shù)運算，提高系統(tǒng)的計算性能。每個數(shù)據(jù)點多符號檢測運算之間是獨立不相關(guān)的，因此，可以把解調(diào)運算映射到GPU上實現(xiàn)，以降低延遲，增加算法的運行效率。

2 基于GPU的位同步算法設(shè)計

2.1 Gardner位同步算法研究

多符號檢測運算能夠提高解調(diào)增益的前提是能夠?qū)崿F(xiàn)精確的位同步，PCM/FM遙測信號的位同步信息能夠從調(diào)頻信號中提取，常用的位同步算法有基于“早遲門”時延的多符號檢測位同步方法[25]、數(shù)字鎖相環(huán)法[26]、O&M位同步算法[27]和Gardner位同步算法[28]?；凇霸邕t門”時延的多符號檢測位同步方法和數(shù)字鎖相環(huán)法類似，都是采用鎖相環(huán)技術(shù)，但是這種同步方法抗干擾能力差，并且在低信噪比和高碼率的情況下性能較差。O&M位同步算法和Gardner位同步算法的應(yīng)用，需要先將PCM/FM調(diào)頻遙測信號進行非相干鑒頻解調(diào)，利用鑒頻之后的信息進行位同步。O&M位同步算法需要對信號進行平方和傅里葉變換等操作，計算比較復(fù)雜。Gardner位同步算法具有運算簡單、容易實現(xiàn)的特點，因此本文采用Gardner位同步算法。

Gardner位同步算法在求位定時誤差時，每個碼元內(nèi)需要2個采樣點，算法的原理圖如圖2所示。當相鄰的碼元不同時，如果不存在位定時誤差，中間采樣點的值應(yīng)該為0。如果中間采樣點的值不為0，則表示存在位定時誤差，且位定時誤差的大小可以由中間采樣點的值求出。利用中間采樣點求出位定時誤差值之后，還需要判斷此誤差是超前還是滯后，可以利用2個采樣點差值和中間采樣點的乘積的正負來判斷。乘積為負數(shù)表示位同步脈沖比信號超前，與之相反的是，乘積為正數(shù)表示位同步脈沖比信號滯后。圖2(a)表示位同步脈沖和信號是同步的，此時中間采樣點的數(shù)值y(n-1/2)為0；圖2(b)表示位同步脈沖比信號超前，此時中間采樣點的數(shù)值y(n-1/2)不為0，而且2個采樣點的差值和中間采樣點的乘積y(n-1/2)[y(n)-y(n-1)]小于0；圖2(c)表示位同步脈沖比信號滯后，此時中間采樣點的數(shù)值y(n-1/2)不為0，而且乘積y(n-1/2)[y(n)-y(n-1)]大于0。

(a)同步

但實際工程中，鑒頻結(jié)果中不可避免地存在噪聲，會對采樣點的數(shù)值造成影響。因此，將一個調(diào)制數(shù)據(jù)的所有采樣點相加來代替采樣點判斷位定時誤差。對下變頻后的信號進行了重采樣處理，目的是使得采樣頻率為碼速率的固定整數(shù)倍，本文設(shè)定的是8倍，即對一個數(shù)據(jù)而言有8個采樣點，這也是后面進行多符號檢測的需要。因此，可利用重采樣頻率來產(chǎn)生一個位同步信號，頻率與碼速率相同。圖3為鑒頻結(jié)果與位同步信號示意圖，位同步信號的上升沿對應(yīng)一個調(diào)制數(shù)據(jù)的開始，由于一個數(shù)據(jù)有8個采樣點，因此從位同步信號的上升沿開始對采樣點進行累加，則一個周期內(nèi)累加8個點，累加結(jié)果代表該調(diào)制數(shù)據(jù)；同時，從位同步信號的下降沿開始對采樣點進行累積，一個周期內(nèi)也累加了8個點，累加結(jié)果代表了相鄰2個調(diào)制數(shù)據(jù)的平均值。

圖3中，當相鄰碼元不同時，例如t1～t2時刻累加結(jié)果為a1，t2～t3時刻累加結(jié)果為a2，t1’～t2’時刻累加結(jié)果為b1。若位同步信號與碼元同步，則a1<0，a2>0，且b1=0；若b1>0，則說明位同步信號比碼元滯后，反之，b1<0，則說明位同步信號比碼元超前，就可以據(jù)此來調(diào)整位同步信號，使其與碼元同步。當相鄰碼元相同時，如t3～t4時刻與t4～t5時刻，由于a3，a4和b3正負極性相同，且b3的極性不受位同步信號超前滯后的影響，無法反映位同步信號與碼元之間的同步關(guān)系，因此這種情況時無法得到誤差信號。

圖3 鑒頻結(jié)果與位同步信號示意

對一段時間的鑒頻結(jié)果進行如上分析，并將相鄰碼元取不同值時得到的誤差信號累加，通過環(huán)路濾波器后得到一誤差控制信號，反饋給位同步信號產(chǎn)生單元，從而實現(xiàn)位同步。

2.2 基于GPU的Gardner位同步算法設(shè)計

根據(jù)上一節(jié)介紹的利用Gardner算法實現(xiàn)位同步的方法，設(shè)計在GPU上進行Gardner位同步算法的核函數(shù)實現(xiàn)方式如圖4所示。

GPU上的線程以及線程塊均為一維，假設(shè)每個線程塊中線程個數(shù)為TPB，則線程塊個數(shù)根據(jù)數(shù)據(jù)長度進行計算。圖4展示了位定時誤差檢測的計算過程，分為BitSycnkernel和TimeErrkernel兩個核函數(shù)實現(xiàn)。核函數(shù)BitSycnkernel實現(xiàn)對采樣點的累加，此核函數(shù)的輸入是鑒頻結(jié)果d_subphase_demodual，采樣數(shù)據(jù)累加后得到2路數(shù)據(jù)，分別為d_data_tongxiang和d_data_orthogonal。d_data_tongxiang表示每個碼元的累加值也即上一節(jié)中的a；d_data_orthogonal表示2個碼元的平均值也即上一節(jié)中的b。核函數(shù)TimeErrkernel用來判斷信號的超前或滯后，判斷結(jié)果得到誤差信號記為d_data_timeerr。

圖4 Gardner算法的核函數(shù)實現(xiàn)框圖

將誤差信號d_data_timeerr從GPU傳輸?shù)紺PU中，通過IPPS庫中的ippsSum_32f()函數(shù)來加速計算誤差信號的累加。將累加結(jié)果通過環(huán)路濾波器后得到誤差控制信號，反饋給位同步信號產(chǎn)生單元，即重采樣模塊，從而實現(xiàn)位同步。

2.3 實驗驗證與結(jié)果分析

為了驗證位同步模塊設(shè)計的正確性，按照圖4中算法的實現(xiàn)流程進行位同步仿真，重采樣頻率為16 MHz，碼速率為2 Mb/s，每個碼元有8個采樣點。分析1 s數(shù)據(jù)，分為2 000段處理，則每段會產(chǎn)生1個誤差頻率控制字。將GPU產(chǎn)生的誤差頻率控制字與Matlab結(jié)果對比，結(jié)果如圖5所示。

從圖5可以看出，利用GPU運算得到的誤差頻率控制字和Matlab的運算結(jié)果相比有誤差。在初始50段數(shù)據(jù)處理時，位同步尚未完成跟蹤，誤差起伏較大,在10-2左右。在50段數(shù)據(jù)之后，位同步完成跟蹤，誤差也趨于穩(wěn)定,在10-5左右。因此，基于GPU的位同步運算滿足計算精度的要求，也驗證了本文所提方法的有效性和正確性。

(a)GPU計算誤差頻率控制字結(jié)果

按照圖4中算法的實現(xiàn)流程進行位同步運算，每段數(shù)據(jù)為0.5 ms，鑒相之后數(shù)據(jù)長度為8×103。表1給出了位同步算法在Matlab和GPU上的運算時間對比，Matlab time表示用Matlab進行位同步運算所用的時間；kernel time表示用GPU進行位同步運算所用的時間；kernel+IPPS time表示用GPU加速誤差檢測的基礎(chǔ)上，用IPPS庫加速計算誤差信號累加所用的時間；kernel+IPPS +memcpy time表示加上數(shù)據(jù)傳輸?shù)臅r間。

由表1可以看出，若只計算位同步運算所用的時間，GPU相比于Matlab取得了143.57倍的加速比；再通過IPPS庫中的ippsSum_32f()函數(shù)來加速計算誤差信號的累加，能夠使加速比提高到194.5倍，把數(shù)據(jù)傳輸?shù)臅r間考慮在內(nèi)，可以取得60.79倍的加速比。以上分析可以說明，對于位同步算法，GPU并行處理過程可以取得顯著的加速效果。

表1 位同步運算用時統(tǒng)計對比

3 基于GPU的多符號檢測算法研究

3.1 多符號檢測算法串行過程分析

為了在GPU上實現(xiàn)并行多符號檢測算法，首先研究多符號檢測算法的串行實現(xiàn)過程。假設(shè)接收到的PCM/FM遙測中頻信號經(jīng)過數(shù)字下變頻、FIR濾波和重采樣處理，得到同相分量I和正相分量Q，I和Q又分別叫做實數(shù)部分和虛數(shù)部分。觀測長度通常取5,7,9，現(xiàn)取觀測長度5，而觀測長度又決定著參考序列的組合數(shù)，所以參考序列的組合數(shù)為32(25)。c表示碼速率，s表示采樣率，則s/c為碼長。本例中碼速率為2 Mb/s，經(jīng)過重采樣后采樣率16 MHz，則碼長為8，每個組合的參考序列長度為40(5×8)。

多符號檢測算法的串行執(zhí)行過程如圖6所示。

圖6 多符號檢測算法串行計算過程

首先計算接收到信號實數(shù)部分SIn與每一個參考序列的點乘的結(jié)果，然后在虛數(shù)部分執(zhí)行類似的點乘計算過程。將復(fù)數(shù)點乘計算結(jié)果進行累加積分，然后進行模平方運算得到模值。比較模值大小，并選取得出最大模值的那組參考序列組合，將參考序列的中間位作為第n位數(shù)據(jù)的判定結(jié)果。然后將接收數(shù)據(jù)向前滑動一位，重復(fù)以上復(fù)數(shù)點乘、累加、模平方和判決的過程，直到完成所有數(shù)據(jù)的判決。

多符號檢測算法的串行運算偽代碼如下：

輸入:接收數(shù)據(jù),S;本地參考序列,R;接收數(shù)據(jù)長度,N;本地參考序列長度,L;本地參考序列數(shù)目,M。輸出:索引,最大模值的本地參考序列的索引號。for (k=0;k amp_max){ index = i; amp_max = amp; } } Index_channel[k] = index;}ComplexMul(a,b){c.x = a.x * b.x - a.y * b.y;c.y = a.x * b.y + a.y * b.x;}

由以上偽代碼可以看出，多符號檢測計算過程有很多for循環(huán)，每個循環(huán)內(nèi)的計算以及計算所用的數(shù)據(jù)是獨立的，所以可以將計算過程進行并行化實現(xiàn)。

3.2 基于GPU的滑動相關(guān)運算設(shè)計

由串行運算過程可知，多符號檢測算法的核心為滑動相關(guān)運算，這部分也是計算量最大的過程。因此，將滑動相關(guān)運算在GPU平臺上高效地實現(xiàn)是接下來研究的重點?；瑒酉嚓P(guān)運算實際上是一個短數(shù)據(jù)序列和一個長數(shù)據(jù)序列的相關(guān)運算過程。滑動相關(guān)運算實質(zhì)上是點乘運算，2段數(shù)據(jù)的點乘運算已經(jīng)在GPU上得到了較為高效的實現(xiàn)，甚至CUDA已經(jīng)內(nèi)置了點乘的API函數(shù)。但是，點乘的API函數(shù)實現(xiàn)的是2個長度相同序列的相關(guān)運算，因此，不能直接把有關(guān)點乘的GPU實現(xiàn)方法移植到本文的多符號檢測算法中。

若觀測長度為l，則會有2l個本地參考信號的組合。長度為N的一段長序列數(shù)據(jù)需要滑動N-l次，得到N-l+1段短序列數(shù)據(jù)，這些由長序列數(shù)據(jù)滑動得到的多個短序列數(shù)據(jù)是相互獨立的。每段長序列數(shù)據(jù)的滑動相關(guān)運算需要計算2l·(N-l+1)次相關(guān)，每次相關(guān)運算也是相互獨立的。由以上分析可知，這些滑動相關(guān)運算都是可以并行執(zhí)行的?？紤]到相關(guān)運算的數(shù)據(jù)序列長度不一樣，本文提出2種基于GPU實現(xiàn)滑動相關(guān)的方案。

① 每個GPU線程計算一次滑動產(chǎn)生的短序列和一組本地參考數(shù)據(jù)的相關(guān)值，如圖7所示。若長序列的長度為N、觀測長度為l，則滑動產(chǎn)生N-l+1段短序列數(shù)據(jù)，需要N-l+1個線程。每個線程塊負責(zé)一組本地參考短序列組合，與滑動產(chǎn)生N-l+1段短序列數(shù)據(jù)的相關(guān)運算，需要2l個線程塊。在每個線程塊內(nèi)可以將數(shù)據(jù)存儲在共享內(nèi)存中以提高內(nèi)存訪問效率，本方案中可以將長序列數(shù)據(jù)和用于計算的那組短序列數(shù)據(jù)組合存儲在共享內(nèi)存中。

圖7 并行滑動相關(guān)方案1

② 每個GPU線程計算一組本地參考序列和滑動產(chǎn)生的短序列的相關(guān)值，如圖8所示。若觀測長度為l，需要2l個線程，本例中l(wèi)=5，則需要32個線程。每個線程塊負責(zé)一次滑動產(chǎn)生的短序列，與2l個本地參考短序列組合的相關(guān)運算。若長序列的長度為N，則滑動產(chǎn)生N-l+1段短序列數(shù)據(jù)，需要N-l+1個線程塊。本方案可以將用于計算的那段滑動產(chǎn)生的短序列數(shù)據(jù)和所有本地參考序列數(shù)據(jù)組合存儲在共享內(nèi)存中，以提高內(nèi)存訪問效率。

圖8 并行滑動相關(guān)方案2

方案1中每個線程塊內(nèi)的線程數(shù)比方案2中更多，所以方案1有更高的占用率，相應(yīng)的資源利用效率也比方案2高。但是，當每段長序列數(shù)據(jù)的長度過長時，方案1中每個線程塊中的線程數(shù)N-l+1會超過線程數(shù)的上限1 024；而且長度較長的一段長序列全部存儲在共享內(nèi)存中，會超過共享內(nèi)存的容量限制而溢出到本地內(nèi)存，進而影響程序的運行效率。相比之下，方案2雖然在觀測長度較短時效率低一些，但是隨著觀測長度變長時，每個線程塊中的線程數(shù)2l會呈指數(shù)增加。如l=7時線程數(shù)為128，由第3節(jié)可知，這時的資源利用效率已趨于穩(wěn)定，不再隨著線程數(shù)的增加而增加。在使用共享內(nèi)存時，方案2數(shù)據(jù)溢出的風(fēng)險也更小,結(jié)構(gòu)設(shè)計更利于程序的擴展，當觀測長度改變時，只需要修改參數(shù)l的大小即可。綜上所述，本文采用方案2設(shè)計的并行滑動相關(guān)方法進行運算。

采用方案2設(shè)計的并行滑動相關(guān)的偽代碼如下：

偽代碼中，ComplexMul()和ComplexAdd()函數(shù)分別實現(xiàn)2個復(fù)數(shù)的乘法和加法。在每個線程塊迭代時，長序列會向前滑動一位，并將滑動得到的短序列存儲到共享內(nèi)存中，然后調(diào)用ComplexMul()和ComplexAdd()函數(shù)來實現(xiàn)2個序列的相關(guān)運算。

輸入:接收信號,S;參考序列,R;接收信號長度,N;參考序列長度,L;參考序列數(shù)目,M。輸出:Amp,接收信號S和參考序列R滑動相關(guān)并求模平方的結(jié)果。__shared__ s_S[],s_R[];tid = threadIdx.x;idx = blockIdx.x * gridDim.x + tid;s_S[tid]=S[ blockIdx.x + tid ];s_R[tid]=R[tid ];__Sync();for(i = 0;i

3.3 基于GPU的多符號檢測算法設(shè)計

多符號檢測算法的目的是根據(jù)滑動相關(guān)的結(jié)果判決長序列各碼元的符號，3.2中介紹的并行相關(guān)算法將所有的相關(guān)結(jié)果輸出保存，會造成資源的浪費，進而影響算法的效率。實際上，對于每段滑動得到的短序列，需要計算2l×L×(s/c)次相關(guān)，經(jīng)過積分和模平方計算得到2l個模值；最后比較模值大小，并根據(jù)最大的模值判決序列中間位的碼元符號。由以上過程可以看出，在進行多符號檢測運算時，可以避免相關(guān)運算結(jié)果的存儲，而直接輸出碼元判決結(jié)果。因此，本文按照上述分析過程設(shè)計出基于GPU的多符號檢測算法，如圖9所示。

圖9 基于GPU的多符號檢測算法流程

由圖9可知，基于GPU的多符號檢測算法的實現(xiàn)流程如下：

① 設(shè)備初始化，并且讀入一段長度為N的重采樣后的數(shù)據(jù)；

② 在主機端分別用函數(shù)CudaMalloc和Malloc分配顯存和內(nèi)存，使用函數(shù)cudaMemcpyHostToDevice將N個數(shù)據(jù)從主機內(nèi)存拷貝到設(shè)備顯存；

③ 將輸入數(shù)據(jù)進行滑動，得到一段與本地參考序列一樣長的短序列，長度為l；

④ 將滑動得到的短序列與2l組本地參考序列的組合分別進行復(fù)數(shù)乘法運算；

⑤ 對復(fù)數(shù)乘法運算得到的結(jié)果進行累加求積分，得到2l個結(jié)果，即短序列與2l組本地參考序列的相關(guān)結(jié)果；

⑥ 對累加值進行模平方求幅度，得到2l個模值；

⑦ 比較2l個模值，得出最大模值對應(yīng)的本地參考序列，并以此判決短序列中間位的碼元序號；

⑧ 重復(fù)步驟③～⑦，直至將輸入的長度為N的數(shù)據(jù)全部計算完成;將所有判決結(jié)果通過函數(shù)cudaMemcpyHostToDevice從設(shè)備顯存拷貝到主機內(nèi)存；

⑨ 用函數(shù)CudaFree和Free釋放顯存和內(nèi)存，結(jié)束運算。

基于GPU的多符號檢測算法流程的偽代碼如下：

輸入:接收信號,S;參考序列,R;接收信號長度,N;參考序列長度,L;參考序列數(shù)目,M。輸出:Out,0或1。__shared__ s_S[],s_R[];tid = threadIdx.x;idx = blockIdx.x * gridDim.x + tid;s_S[tid] = S[ blockIdx.x + tid ];s_R[tid] = R[tid ];__Sync();for(i = 0;iAmp_A) Amp_A= Amp[ blockIdx.x * gridDim.x+j];for(j=M/2;jAmp_B) Amp_B= Amp[ blockIdx.x * gridDim.x+j];if(Amp_A>Amp_B) Out[ blockIdx.x]=-1;else if( amp_A < amp_B) Out[ blockIdx.x]=1;

由偽代碼可知，將每段當前觀測數(shù)據(jù)的模平方結(jié)果劃分為長度相同的兩部分，獲取前一半數(shù)據(jù)中幅度的最大值，以及后一半數(shù)據(jù)中幅度的最大值；比較這2個幅度，進行符號位判決，若前一半數(shù)據(jù)對應(yīng)的幅度最大值較大，則當前觀測數(shù)據(jù)的中間符號位輸出-1；若后一半數(shù)據(jù)對應(yīng)的幅度的最大值較大，則此符號位輸出1。

要實現(xiàn)上述直接輸出中間符號的計算，需要將參考序列按照圖10中的方式排列。以觀測長度等于5、每個符號8個采樣點為例，則參考序列的個數(shù)為32(25)，分為2部分，每部分16個參考序列，中間一位分別為-1和1。這樣得出每部分相關(guān)之后的最大模值amp_A,amp_B之后，再比較二者大小，就能直接確定觀測數(shù)據(jù)的中間符號位的取值。

圖10中也展示了本地參考序列的存儲結(jié)構(gòu)，每個縱列表示一組本地參考序列，以觀測長度等于5為例，則共有32個縱列代表32組本地參考序列。每個符號又有8個采樣點，則每縱列存儲40個點。在讀取和存儲數(shù)據(jù)時，是按照橫行進行的，這樣可以保證一個線程束讀取32位連續(xù)的地址，進而提高內(nèi)存的加載效率。

圖10 參考序列排列存儲方式

4 實驗驗證與結(jié)果分析

4.1 算法實時性測試

為了對基于GPU的多符號檢測算法加速性能進行實驗驗證，利用PCM/FM遙測基帶設(shè)備產(chǎn)生多種速率的調(diào)頻遙測信號。為了降低實驗過程中時間統(tǒng)計數(shù)據(jù)極端值對實驗結(jié)論的影響，本文將使用截尾平均數(shù)法對多次實驗數(shù)據(jù)求平均值作為最終結(jié)果，即直接將明顯不正確的極值結(jié)果剔除，本實驗取20次實驗結(jié)果剔除3次。將數(shù)據(jù)分為0.5 ms每段，每段數(shù)據(jù)長度2.8×104，采用真實的調(diào)頻遙測基帶產(chǎn)生實時PCM/FM遙測信號，信號中頻為70 MHz，采樣頻率為56 MHz。

在利用并行多符號檢測算法進行解調(diào)時，要首先對接收到的PCM/FM遙測信號進行浮點數(shù)轉(zhuǎn)換、并行下變頻、并行時域濾波和并行重采樣運算，以獲得符合多符號算法輸入的基帶固定采樣點的數(shù)據(jù)。因此，利用并行多符號檢測算法進行解調(diào)運算的時間包括以上下變頻等運算的總時間。

圖11顯示了在不同碼速率和不同觀測長度時的運算用時，其中觀測長度用N表示，N=3,5,7。由圖11可以看出，在觀測長度為5時，隨著碼速率的增加，運算用時也在增加，在碼速率10 Mb/s時運算時間達到最大為718.793 ms。因此，在觀測長度為5時，碼速率10 Mb/s以下的數(shù)據(jù)均可以實現(xiàn)實時解調(diào)。隨著觀測長度的增加，運算量急劇增加，運算時間也快速增加。特別是在觀測長度為9時，僅在碼速率為1 Mb/s的情況下，可以在1 s內(nèi)完成解調(diào)。

圖11 并行多符號檢測時間對比

因此，本文設(shè)計的并行多符號檢測算法在觀測長度為5的情況下，10 Mb/s以下碼速率均可實現(xiàn)實時解調(diào)；在觀測長度為7的情況下，碼速率低于5 Mb/s時可實現(xiàn)實時解調(diào)。

4.2 算法誤碼率測試

對PCM/FM遙測信號并行解調(diào)算法的性能進行測試，調(diào)頻遙測基帶輸出的信號經(jīng)過可變衰減器，由信號源輸出產(chǎn)生高斯白噪聲，用功分器將以上2路信號合成一路信號。再將這一路信號平均分為2路，一路輸出給數(shù)字處理端進行并行解調(diào)運算；另一路輸出給調(diào)頻遙測基帶進行解調(diào)。通過設(shè)置可變衰減器的值和噪聲功率，使信噪比達到固定大小，統(tǒng)計并記錄此信噪比下的誤碼率。

誤碼率測試結(jié)果如圖12所示，可以看出，多符號檢測算法的解調(diào)性能比非相干鑒頻算法有較大的提升，且隨著觀測長度的增加，解調(diào)增益也在增加。在誤碼率為10-4的情況下，觀測長度為5時，多符號檢測算法比非相干鑒頻算法提高約2.6 dB的解調(diào)增益；觀測長度為7時，提高約3.2 dB的解調(diào)增益。本文基于GPU的并行解調(diào)算法的性能與FPGA的誤碼率曲線近似，并且由于GPU浮點運算相比FPGA精度更高，因此GPU解調(diào)增益稍優(yōu)于FPGA,大約有0.1 dB的增益。

圖12 誤碼率測試結(jié)果

5 結(jié)束語

傳統(tǒng)多符號檢測均采用FPGA實現(xiàn)系統(tǒng)功能，存在硬件設(shè)計周期較長、硬件平臺通用性低和升級難度大等問題。針對傳統(tǒng)硬件實現(xiàn)的問題，本文提出了基于GPU的并行多符號檢測算法。通過對多符號檢測算法串行運算過程的分析，得出算法核心為滑動相關(guān)運算，即一個短數(shù)據(jù)序列和一個長數(shù)據(jù)序列的相關(guān)運算過程。提出了滑動相關(guān)的并行化實現(xiàn)方案，并據(jù)此設(shè)計了基于GPU的多符號檢測算法。針對多符號檢測運算能夠提高解調(diào)增益的前提是能夠?qū)崿F(xiàn)精確的位同步，本文提出了基于GPU的Gardner位同步算法。

最后進行了實驗驗證，在觀測長度為5的情況下，10 Mb/s以下碼速率均可實現(xiàn)實時解調(diào)；在觀測長度為7的情況下，碼速率低于5 Mb/s時可實現(xiàn)實時解調(diào)?；贕PU的并行解調(diào)算法與基于FPGA實現(xiàn)解調(diào)的誤碼率曲線近似，并且GPU解調(diào)增益稍優(yōu)于FPGA。