一種基于GPU的PCM/FM信號多符號檢測算法

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

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

0 引言

飛行器遙測是導彈、火箭和衛星等航天器試驗和運行過程中必不可少的支持系統，能夠實時監測航天器內部工作狀態、電氣性能和環境參數等重要信息，為航天器性能檢測、效能評估及故障分析提供依據[1]。隨著航天事業的深入發展，全球各個航天大國對宇宙的探索和對宇宙資源的爭奪愈加強烈[2]，遙測技術突顯出更為重要的作用。

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

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

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

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

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

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

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

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

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

(1)

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

(2)

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

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

(3)

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

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

(4)

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

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

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

IB+jQB。

(5)

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

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

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

(6)

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

S=Nc2×(IB2+QB2)。

(7)

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

為了方便，圖1僅給出第k個本地復信號與接收到的正交基帶復信號進行處理的框圖，其他組合的本地信號與其類似。基帶多符號檢測算法流程大致如下，各個本地信號分別與接收到的正交基帶復信號進行相關平方處理，然后對各路輸出進行比較判決。

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

在一次運算過程中，求相關運算有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數據需要進行2×106次多符號檢測運算，還需要將每次乘加計算量乘以2×106才能得到1 s數據的計算量。

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

2 基于GPU的位同步算法設計

2.1 Gardner位同步算法研究

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

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

(a)同步

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

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

圖3 鑒頻結果與位同步信號示意

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

2.2 基于GPU的Gardner位同步算法設計

根據上一節介紹的利用Gardner算法實現位同步的方法，設計在GPU上進行Gardner位同步算法的核函數實現方式如圖4所示。

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

圖4 Gardner算法的核函數實現框圖

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

2.3 實驗驗證與結果分析

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

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

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

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

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

表1 位同步運算用時統計對比

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

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

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

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

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

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

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

輸入:接收數據,S;本地參考序列,R;接收數據長度,N;本地參考序列長度,L;本地參考序列數目,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循環，每個循環內的計算以及計算所用的數據是獨立的，所以可以將計算過程進行并行化實現。

3.2 基于GPU的滑動相關運算設計

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

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

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

圖7 并行滑動相關方案1

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

圖8 并行滑動相關方案2

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

采用方案2設計的并行滑動相關的偽代碼如下：

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

輸入:接收信號,S;參考序列,R;接收信號長度,N;參考序列長度,L;參考序列數目,M。輸出:Amp,接收信號S和參考序列R滑動相關并求模平方的結果。__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的多符號檢測算法設計

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

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

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

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

② 在主機端分別用函數CudaMalloc和Malloc分配顯存和內存，使用函數cudaMemcpyHostToDevice將N個數據從主機內存拷貝到設備顯存；

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

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

⑤ 對復數乘法運算得到的結果進行累加求積分，得到2l個結果，即短序列與2l組本地參考序列的相關結果；

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

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

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

⑨ 用函數CudaFree和Free釋放顯存和內存，結束運算。

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

輸入:接收信號,S;參考序列,R;接收信號長度,N;參考序列長度,L;參考序列數目,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;

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

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

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

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

4 實驗驗證與結果分析

4.1 算法實時性測試

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

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

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

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

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

4.2 算法誤碼率測試

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

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

圖12 誤碼率測試結果

5 結束語

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

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