李榮春,周 鑫,王慶林,梅松竹

(國防科技大學(xué)計算機學(xué)院并行與分布處理國防重點實驗室,湖南 長沙 410073)

1 引言

近年來,隨著近地衛(wèi)星數(shù)量越來越多,衛(wèi)星通信也越來越被廣泛應(yīng)用于航海、軍事、應(yīng)急通信,甚至是手機通信中。衛(wèi)星通信因為其全覆蓋、高穩(wěn)定性和高魯棒性,更加被人們所接受。衛(wèi)星數(shù)字傳輸信號的處理通常由數(shù)字信號處理器DSP(Digital Signal Processor)或者現(xiàn)場可編程邏輯門陣列FPGA(Field Programmable Gate Array)完成。近年來,GPU(Graphics Processing Unit)憑借其強大的高速并行計算能力和豐富便捷的編程接口等優(yōu)勢逐漸開始在數(shù)字信號處理領(lǐng)域扮演十分重要的角色。在衛(wèi)星高速數(shù)字傳輸?shù)膽?yīng)用場景下,要對接收到的信號進行解調(diào)、同步和譯碼等一系列計算,為了實現(xiàn)高速傳輸?shù)哪康?應(yīng)最大程度降低處理設(shè)備的計算時延,提高信號處理的吞吐率。利用GPU的并行多線程的計算能力,能夠?qū)崿F(xiàn)衛(wèi)星信號多路、多幀的并行處理,還能利用多線程運算對信號處理算法進行并行優(yōu)化,降低處理時延。與專用元器件和FPGA相比,基于GPU的數(shù)字傳輸信號處理器開發(fā)周期短,能夠適配的信號制式類型完全由軟件來定義,還可以根據(jù)需求靈活調(diào)整配置,從而達到最佳性能,具有較好的應(yīng)用前景。

利用GPU進行基帶處理并行算法的設(shè)計有以下難點。首先,GPU是批處理模式,利用多個并行計算單元對大批量數(shù)據(jù)進行并行計算,而通信協(xié)議中基帶處理算法的延遲要求在毫秒級,速率要求極高,這就對GPU并行算法的延遲和吞吐率提出了一定的要求;其次,很多基帶處理算法是按照位運算進行設(shè)計的,數(shù)據(jù)序列之間存在一定的數(shù)據(jù)依賴性,如果采用GPU進行并行計算,需要重新設(shè)計并行算法,打破數(shù)據(jù)序列之間的依賴關(guān)系,從而可以將這些數(shù)據(jù)分布到多個計算單元上實現(xiàn)并行計算,這對并行算法的改造提出了一定的挑戰(zhàn);最后,多個基帶算法之間需要低延遲的快速切換,這對算法群之間的緩沖區(qū)設(shè)計提出了更高的要求。

本文針對上述3個難點,提出了基于GPU的衛(wèi)星通信基帶處理高吞吐率并行算法,通過巧妙的并行算法設(shè)計,打破數(shù)據(jù)序列的順序依賴關(guān)系,并通過性能優(yōu)化、緩存區(qū)優(yōu)化,解決吞吐率與延遲問題,有效支持衛(wèi)星通信協(xié)議在GPU上的高速執(zhí)行。

2 衛(wèi)星通信基帶鏈路

衛(wèi)星通信下行鏈路主要對衛(wèi)星信號進行解析,其基帶算法負責(zé)對數(shù)字傳輸信號進行處理,并將其還原成發(fā)送端原始數(shù)據(jù)。本文優(yōu)化的基帶鏈路是全球?qū)Ш叫l(wèi)星系統(tǒng)GNSS(Global Navigation Satellite System)的衛(wèi)星通信下行鏈路[1]。衛(wèi)星通信下行鏈路主要包括重采樣、匹配濾波、解相位模糊、幀同步和解擾等過程,解調(diào)后的數(shù)據(jù)還需要進行判決和譯碼,最后得到原始數(shù)據(jù)。解調(diào)包括BPSK/QPSK/8PSK/16QAM這4種星座映射模式。譯碼過程包括了基于國際空間數(shù)據(jù)系統(tǒng)咨詢委員會CCSDS(Consultative Committee for Space Data Systems)標準的低密度奇偶校驗LDPC(Low Density Parity Check)碼、里德-所羅門RS(Reed-Solomon)碼、卷積碼以及RS碼+卷積碼級聯(lián)碼。

3 相關(guān)工作

近年來,基于GPU的無線通信并行算法引起學(xué)術(shù)界和產(chǎn)業(yè)界越來越多的重視。大多數(shù)工作聚焦于單個無線通信算法的GPU加速,如基于GPU的同步算法[2]、Viterbi算法[3]、解調(diào)算法[4]、LDPC算法[5-8]、極化碼[9]、信道估計算法[10]、MIMO-OFDM(Multiple-Input Multiple-Output,Orthogonal Frequency Division Multiplexing)檢測器[11]等。

目前很多研究人員開始研究基于GPU的完整的通信鏈路,如WiMAX[12]和WiFi[13]。近些年,隨著衛(wèi)星通信的興起,越來越多的研究人員開始聚焦基于GPU的衛(wèi)星通信基帶并行算法研究[14-20],如基于直接位置估計DPE(Direct Position Estimation)的全球?qū)Ш叫l(wèi)星系統(tǒng)(GNSS)GPU接收機[14]、全球?qū)Ш叫l(wèi)星系統(tǒng)反射測量GNSS-R(Global Navigation Satellite System-Reflectometry)GPU接收機[15]、基于多GPU負載均衡的GNSS接收機[16,17]、基于GPU的GNSS多徑并行仿真器[18]、基于MPI+CUDA的遙感衛(wèi)星數(shù)據(jù)處理通用系統(tǒng)[19]和基于GPU的信號相關(guān)器的GNSS接收機。以上大多數(shù)研究都是針對接收機鏈路中的部分算法進行并行優(yōu)化,本文基于GPU實現(xiàn)GNSS衛(wèi)星通信基帶接收機完整下行鏈路算法群。

4 基于GPU的衛(wèi)星通信基帶并行算法

4.1 基于GPU衛(wèi)星通信基帶信號處理框架

基于GPU基帶信號處理框架主要面向NVIDIA系列GPU,使用CUDA C/C++開發(fā),采用了B/S架構(gòu),便于遠程監(jiān)控處理程序運行狀態(tài)。如圖1所示,該框架的功能模塊可分為解調(diào)模塊、譯碼模塊和輔助模塊,其中每個模塊又根據(jù)各自特點分成多個子模塊。解調(diào)模塊負責(zé)完成接收到的下變頻后的數(shù)字傳輸信號重采樣、匹配濾波、幀同步、解擾、判決等譯碼之前的所有步驟。經(jīng)過解調(diào)判決的數(shù)據(jù)繼續(xù)留存在GPU內(nèi)存中參與后續(xù)的譯碼運算。譯碼模塊根據(jù)事先設(shè)置好的譯碼參數(shù),對解調(diào)后的程序進行譯碼。圖1中,r表示譯碼碼率,Lk表示信息位長度。輔助模塊包括了參數(shù)配置、文件讀寫、狀態(tài)顯示和日志輸出等子模塊,負責(zé)配合處理程序完成處理任務(wù)。解調(diào)模塊和譯碼模塊既可聯(lián)合使用,又可獨立運行,程序的運行次序和功能均可以根據(jù)參數(shù)來確定。

Figure 1 Framework of baseband signal processing for satellite communication based on GPU

衛(wèi)星基帶信號處理包含了復(fù)雜的控制邏輯、密集的浮點計算和頻繁的不規(guī)則訪存,為了提高處理速度,框架通過結(jié)合算法的計算規(guī)律和訪存特性,設(shè)計并實現(xiàn)了基于GPU的高速并行解調(diào)算法和并行譯碼算法。設(shè)計的并行位同步算法和并行濾波算法,對原有串行算法進行了粗粒度和細粒度的并行優(yōu)化,在低誤碼率損失的情況下實現(xiàn)了BPSK/QPSK/8PSK/16QAM的高速解調(diào)。該加速器還實現(xiàn)了基于CCSDS標準的LDPC碼、RS碼、卷積碼以及RS碼+卷積碼級聯(lián)碼的并行譯碼。其中,LDPC碼采用了MSA(Min-Sum Algorithm)譯碼算法,RS碼采用了BMA(Berlekamp-Massey Algorithm)譯碼算法,卷積碼采用了Viterbi譯碼算法。在實現(xiàn)MSA并行譯碼算法的過程中,使用了兩段式求解最小值的并行算法,提高了譯碼速率。在實現(xiàn)BMA并行譯碼算法的過程中,對Chien搜索算法進行了并行優(yōu)化,大幅提高了尋找錯誤位置的速度。Viterbi譯碼算法實現(xiàn)了并行butterfly算法來計算路徑度量值,并且使用截斷-重疊譯碼的方式實現(xiàn)了卷積碼多碼塊并行譯碼。如圖2所示,該解調(diào)器的解調(diào)、幀同步和譯碼等計算過程均在GPU上實現(xiàn),CPU僅負責(zé)控制計算流程和數(shù)據(jù)文件讀寫操作,該方案有效避免了CPU和GPU通信帶來的開銷。

Figure 2 Division of CPU and GPU computing tasks

CPU和GPU的任務(wù)分工如圖2所示?？蚣茌斎胧菂?shù)文件和對應(yīng)的數(shù)據(jù)文件,數(shù)據(jù)文件先由CPU從磁盤讀入主機內(nèi)存,同時CPU會根據(jù)參數(shù)文件中的配置參數(shù)進行一系列的初始化操作,然后將主機內(nèi)存中的數(shù)據(jù)分批次傳入GPU內(nèi)存。當數(shù)據(jù)傳入GPU內(nèi)存時,即開啟GPU的工作周期。首先由解調(diào)模塊的GPU核函數(shù)依次對傳入信號的正交分量I路數(shù)據(jù)和Q路數(shù)據(jù)進行處理運算。當解調(diào)模塊各個步驟結(jié)束之后,由幀同步子模塊將處理結(jié)果排列成固定幀長的連續(xù)碼字,并送入譯碼模塊進行譯碼。譯碼模塊按照參數(shù)文件中的配置開始并行譯碼。譯碼結(jié)束后,譯碼結(jié)果被傳入CPU內(nèi)存,最終由CPU將結(jié)果寫入磁盤。

4.2 基于GPU的并行解調(diào)算法

解調(diào)模塊支持對BPSK、QPSK、8PSK、16QAM這4類調(diào)制方式的并行解調(diào)。信號在解調(diào)的重采樣過程中使用了Gardner位同步算法來確定最佳采樣時刻。FPGA和DSP在處理信號時可以將信號按照連續(xù)序列依次處理,通過高速計算單元和并行計算邏輯來加快處理速度。GPU處理信號數(shù)據(jù)的時候通常需要一次性將大量的數(shù)據(jù)傳入GPU內(nèi)存,經(jīng)過計算后再將結(jié)果傳回CPU內(nèi)存。這樣避免了CPU-GPU的頻繁通信造成的處理時延。GPU還能夠利用豐富的并行計算資源一次性計算多組結(jié)果,實現(xiàn)更高的處理吞吐率。由于GPU內(nèi)存有限,這時候的信號數(shù)據(jù)往往是被截斷的、不連續(xù)的。在位同步的計算過程中,接收載波的時鐘相位不斷變化,數(shù)字控制振蕩器NCO(Numerically Controlled Oscillator)環(huán)路生成的采樣間隔也在不斷校準,這樣累計校準的算法使得位同步的計算成為一個基于時間序列的算法。對于t時刻的時鐘誤差只能從起始時刻t0一步一步校準得到,而無法從任意中間時刻的采樣點開始計算,這給位同步計算的并行化帶來了極大的挑戰(zhàn),使得GPU豐富的內(nèi)存資源和計算資源沒有了用武之地。然而對于單路信號而言,位同步計算的并行化還是可以針對重采樣過程中拉格朗日插值和采樣間隔的計算這2個環(huán)節(jié)進一步優(yōu)化,以實現(xiàn)更低的計算時延。

重采樣后的信號要經(jīng)過匹配濾波來獲取更好的解調(diào)效果。匹配濾波本質(zhì)上就是對采樣結(jié)果進行線性卷積運算。雖然線性卷積依然是基于時間序列的算法,但是已經(jīng)有方法能夠通過循環(huán)卷積來代替線性卷積,將計算的時間范圍縮小到固定的多個時間段中。利用重疊相加法或重疊保留法可以將連續(xù)的線性卷積運算轉(zhuǎn)換成分段的循環(huán)卷積運算。這樣的轉(zhuǎn)換已被證明對結(jié)果的精度和誤差不會帶來任何損失,轉(zhuǎn)換最大的好處是使濾波運算擺脫了時間序列的束縛,能夠?qū)Χ喽螖?shù)據(jù)同時計算卷積,充分利用了GPU的計算資源。本文基于GPU平臺設(shè)計并實現(xiàn)了重疊保留法的并行卷積算法,實現(xiàn)了高速濾波運算。

在使用GPU處理信號數(shù)據(jù)的時候通常是將信號數(shù)據(jù)文件讀取到GPU內(nèi)存中,在一般情況下,數(shù)據(jù)文件的大小能達到數(shù)GB。由于GPU的內(nèi)存是有限的,這些數(shù)據(jù)難以一次讀入GPU內(nèi)存,只能分批次讀入GPU內(nèi)存。分批讀入會造成截斷,為了避免截斷造成的數(shù)據(jù)丟失或者誤碼,需要針對GPU處理流程設(shè)計一定的重疊區(qū)域和對齊功能。重疊區(qū)域保證了數(shù)據(jù)的連續(xù)性,例如在進行Gardner位同步計算時,計算一個插值結(jié)果常常需要用到連續(xù)4個時刻的采樣值,如果這4個時刻的采樣值恰好被分到了不同的批次讀入,會造成數(shù)據(jù)計算錯誤。因此,需要設(shè)置一定的重疊區(qū)域,以避免程序設(shè)計導(dǎo)致的誤碼現(xiàn)象。此外,在進行幀同步的時候,如果解調(diào)結(jié)果不是幀長的整數(shù)倍,對這些多出幀長整數(shù)倍的部分需要進行相應(yīng)處理,以避免丟幀。綜上所述,針對GPU的解調(diào)需要專門設(shè)計一種緩沖對齊機制,以保證數(shù)據(jù)的連續(xù)性和解調(diào)的正確性。本節(jié)將主要介紹上述專門為GPU設(shè)計的并行解調(diào)算法和相應(yīng)的處理技術(shù)。

4.2.1 并行位同步算法

Gardner位同步單元的計算過程可以由NCO環(huán)路、插值濾波器和定時誤差檢測器組成。為了確保能夠在最大值處采樣,Gardner位同步算法通過內(nèi)插的方法計算出采樣值,同時用NCO環(huán)路生成采樣間隔。采用內(nèi)插方法計算得到采樣值后,對應(yīng)的采樣時刻也需要重新計算。這個采樣時刻與輸入信號采樣周期Ts的整數(shù)倍相差小數(shù)間隔μ。

圖3所示的算法是Gardner位同步算法的主要過程。其中,x(t)為輸入信號,經(jīng)過周期為Ts的固定時鐘采樣后,得到x(mTs),m=1,2,…表示離散序列編號;插值濾波器、環(huán)路濾波器、定時誤差檢測器和數(shù)控振蕩器組成了一個NCO環(huán)路,NCO環(huán)路的作用是生成采樣間隔。插值濾波器通過插值運算得到k時刻的結(jié)果y(kTi),其中Ti為輸出信號的采樣周期;τ(n)為計算得到的定時誤差;e(n)經(jīng)過環(huán)路濾波后的校準誤差;mk為NCO得到的采樣步長;μk為根據(jù)k時刻插值結(jié)果計算得到的小數(shù)間隔。NCO寄存器的初始值R(0)=1,NCO寄存器深度為1,相鄰2個周期的NCO寄存器的值如式(1)所示:

Figure 3 Synchronization algorithm

R(k+1)←[R(k)-wc]mod 1

(1)

其中wc是控制步長。當NCO寄存器的值過零點時,寄存器的值模1,即得到下一個符號周期NCO寄存器的值。當NCO寄存器過零點時,就是對輸入信號進行重采樣的位置,該位置相對于輸入信號的時鐘周期的整數(shù)倍多出的部分即為重采樣的小數(shù)間隔μ。

根據(jù)幾何關(guān)系,2次重采樣的間隔可以由式(2)得出:

l=R(k)/wc

(2)

該采樣間隔也可由式(3)表示:

l=sTs+μ

(3)

其中,s是采樣間隔的整數(shù)部分,下一次采樣則從k+s處開始計算。

采樣過程中(k-1)Ts,kTs,(k+1)Ts,(k+2)Ts時刻的信號值將被作為參數(shù)進行內(nèi)插運算,采樣時刻相對于采樣周期整數(shù)倍的小數(shù)間隔μ也將作為內(nèi)插運算的參數(shù)。

本文采用了立方插值函數(shù)作為拉格朗日插值函數(shù)。立方插值濾波器的插值多項式如式(4)～式(7)所示:

(4)

(5)

(6)

(7)

其中,μ是內(nèi)插點的時間偏移,C-2、C-1、C0、C1是立方插值濾波器的系數(shù)。當計算得到C=[C-2(k),C-1(k),C0(k),C1(k)]后,即由式(8)可得內(nèi)插值:

y(k)=Cx

(8)

其中,x為(k-1)Ts、kTs、(k+1)Ts、(k+2)Ts時刻的信號的采樣值。

在GPU計算過程中,更新NCO寄存器值、計算小數(shù)間隔都可以由一個線程完成,當進行拉格朗日插值時,則可以由多個線程同時參與計算。

并行計算內(nèi)插值的過程可以分為2步。第1步是通過插值多項式的系數(shù)矩陣與內(nèi)插點小數(shù)偏移μ的冪次向量的乘積。第2步是插值系數(shù)向量和信號向量的乘積。即4×4的矩陣與長度為4的μ的冪次向量的乘積,即[1,μ,μ2,μ3],以及2個長度為4的向量的乘積。可以使用4個線程分別計算4個插值系數(shù),并且分別計算系數(shù)與信號值的乘積,經(jīng)過2次歸約運算后即可得到內(nèi)插值。

根據(jù)Gardner算法,在得到連續(xù)3次的內(nèi)插采樣值后,即可進行一次時鐘誤差的檢測和校正。在計算內(nèi)插值的時候,可以申請更多的線程參與運算,從而一次得到多個內(nèi)插采樣值。假設(shè)同時計算的內(nèi)插值數(shù)量為I,則申請4I個線程同時計算上述乘積運算。

每當計算得到2個內(nèi)插結(jié)果后,都要對時鐘相位進行校正,也就是計算更新NCO控制步長wc的值。對于不同的解調(diào)方式,校正方法有所區(qū)別。對于QPSK來說,時鐘相位誤差由式(9)確定:

e=sign(yI(k-2))-sign(yI(k))*yI(k-1)+

sign(yQ(k-2))-sign(yQ(k))*yQ(k-1)

(9)

其中,yI和yQ為I、Q路分別計算得到的內(nèi)插值。除了計算時鐘相位誤差之外,在得到插值結(jié)果后,還要計算相應(yīng)的載波相位誤差。在更新相位誤差后,NCO控制步長wc和載波相位θ都發(fā)生了變化,新的插值結(jié)果需要用新的控制步長和載波相位進行后續(xù)的計算。因此,不能一次計算過多的內(nèi)插值,假設(shè)校準間隔為M,當M=1時,每計算得到2個內(nèi)插結(jié)果校正一次相位誤差;當M≥2時,就需要考慮未及時檢測更新相位誤差帶來的影響。

在并行計算插值結(jié)果時,需要先由主線程計算得到M次插值對應(yīng)的NCO寄存器值和小數(shù)間隔。對于k時刻的一次插值來說,下一次插值對應(yīng)的時刻由式(10)計算得到:

R(k+s)=R(k)-s×wc

(10)

為了實現(xiàn)并行計算,需要一次計算多次采樣的NCO寄存器的值,通過事先迭代計算多次得到寄存器結(jié)果R1,R2,…,RM、時鐘偏移μ1,μ2,…,μM及對應(yīng)的采樣時刻k1,k2,…,kM。通過4M個線程將對應(yīng)M個時刻的4M個對應(yīng)的信號采樣值提取出來,利用提前計算好的參數(shù)實現(xiàn)并行內(nèi)插運算。

并行插值運算減少了條件判斷,有利于降低GPU運算延遲。當M≥2時,需要考慮邊界判斷帶來的結(jié)果誤差。如果k1到kM中的某一個時刻ki所需要的4個采樣值沒有在這一輪讀入GPU內(nèi)存,則會造成結(jié)果錯誤。解決辦法是:在插值計算結(jié)束后設(shè)置同步柵欄來判斷所需采樣值是否超過信號讀取范圍,如果超過范圍則舍棄相應(yīng)的插值結(jié)果,僅保留所需采樣值處于本輪信號值讀取范圍的插值結(jié)果;對于超過范圍的結(jié)果,記錄最后一個超過讀取范圍的采樣時刻偏移和NCO寄存器等各個寄存器狀態(tài),在下一輪讀取信號值時,通過重疊讀取和更新寄存器狀態(tài)的方式來繼續(xù)計算重采樣內(nèi)插值。

校準間隔可以根據(jù)信號情況進行動態(tài)調(diào)整,當M較大時,可以提高解調(diào)速率,但是誤碼性能也會相應(yīng)地受到影響。實際測試過程中發(fā)現(xiàn),當M≥64后會出現(xiàn)誤碼增多的情況。

4.2.2 并行濾波

濾波運算實質(zhì)上就是將濾波器的時域信號值與輸入信號進行線性卷積。線性卷積的操作十分適合如FPGA這樣串行計算效率較高的設(shè)備。對于GPU而言,如果按照線性卷積中的時間序列逐步計算卷積,則無法發(fā)揮GPU的資源優(yōu)勢。通過對輸入信號進行一系列的分段和重疊等預(yù)處理操作,可以用多個循環(huán)卷積計算結(jié)果的拼接來得到和一段線性卷積計算相同的結(jié)果。這樣十分有利于發(fā)揮GPU的線程資源和計算資源優(yōu)勢,達到更快的處理速率。

替代的方法主要有重疊相加法和重疊保留法。這2種方法的步驟十分類似,本文采用了重疊保留法。對于重疊保留法而言,需要先對輸入信號進行分段,在分段的時候?qū)γ恳欢蔚念^部重疊填充前一個分段的數(shù)值,在計算卷積之后,再將每一段結(jié)果的頭部刪去,最后把各個分段的結(jié)果拼接起來,就得到最終結(jié)果。

基于GPU的并行卷積運算步驟如下所示:

(1)將濾波器時域信號值存入GPU內(nèi)存,并進行FFT(Fast Fourier Transform)卷積。

(2)對接收信號值進行分段,接收信號的分段長度為Ls,分段長度由事先設(shè)定的FFT長度確定。設(shè)FFT長度為N,N?Ls,濾波器長度為D,則將信號分為長度為L的段,其中L的計算如式(11)所示:

L=N-D+1

(11)

總共分段數(shù)BS=Ls/L。至少需要申請BS×N長度的GPU內(nèi)存,為每一段信號申請長度為N的GPU內(nèi)存,將長度為L的數(shù)值填入L段的后部。

(3)將每一分段的最后一段長度為D-1的數(shù)值存入后一個分段。對于第1個分段填充值需要區(qū)分是否是第1輪讀取的數(shù)據(jù),如果是第1輪讀取的數(shù)據(jù),則填充零;如果是中間輪讀取的數(shù)據(jù),則需要在上一輪計算時將最后D-1長度數(shù)據(jù)保存起來,以給下一輪填充使用。

(4)對填充好的BS個分段數(shù)據(jù)進行批量并行FFT運算。

(5)將FFT運算的結(jié)果與之前計算濾波器的FFT結(jié)果相乘。

(6)對相乘結(jié)果進行并行IFFT(Inverse Fast Fourier Transform)計算。

(7)對每個分段進行內(nèi)存移動操作,只取每個分段的后L個數(shù)據(jù),舍棄前D-1個數(shù)據(jù)。

(8)對計算結(jié)果進行拼接調(diào)整并且作最后抽樣。

圖4展示了通過分段卷積進行濾波運算的流程。在GPU執(zhí)行上述步驟時,可以最大程度調(diào)用計算資源,實現(xiàn)并行FFT、并行IFFT和并行乘法操作,在整個過程中有大量內(nèi)存移動操作,所以需要注意線程同步問題,在實現(xiàn)的過程中通過合理調(diào)整計算和訪存順序,將對同一個內(nèi)存區(qū)域的讀寫操作步驟分隔開,避免了寫后讀沖突。

Figure 4 Filtering algorithm

4.2.3 多級緩沖區(qū)對齊技術(shù)

針對基于GPU數(shù)字信號處理的特點,在充分確保處理結(jié)果的連續(xù)性和正確性的基礎(chǔ)上,最大程度實現(xiàn)運算的并行度,需要對數(shù)據(jù)處理的各個節(jié)點,特別是處理邊界部分進行專門的優(yōu)化。設(shè)置緩沖區(qū)能夠減少不必要的訪存操作,降低處理時延。與此同時,對于大規(guī)模并行處理操作,計算單元支持的最大長度往往是固定的,而輸入數(shù)據(jù)的長度無法事先確定,需要對這些超過計算單元支持的最大長度的部分進行相應(yīng)的處理。本文通過軟件編程設(shè)置專門的緩沖區(qū)和對齊單元,解決訪存時延和并行計算的邊界問題。

在整個解調(diào)周期中共設(shè)置了4級緩沖區(qū),如圖5所示,每個緩沖區(qū)的功能如下:

Figure 5 Design of CPU-GPU heterogeneous buffer

第1級緩沖區(qū)Buf1位于CPU內(nèi)存,用于保存讀入的信號數(shù)據(jù);

第2級緩沖區(qū)Buf2位于GPU內(nèi)存,用于保存超過FFT長度整數(shù)倍的重采樣結(jié)果;

第3級緩沖區(qū)Buf3位于GPU內(nèi)存,用于保存超過幀長整數(shù)倍的解調(diào)結(jié)果;

第4級緩沖區(qū)Buf4位于CPU內(nèi)存,用于保存譯碼結(jié)果。

當處理程序?qū)?shù)據(jù)文件讀入CPU內(nèi)存后,將信號數(shù)據(jù)存入Buf1中,分多輪讀入GPU內(nèi)存,減少磁盤到CPU內(nèi)存的訪存時延,同時節(jié)省GPU內(nèi)存。

當重采樣結(jié)束后,為了保證輸入濾波器的數(shù)據(jù)長度是并行卷積分段長度的整數(shù)倍,設(shè)置對齊單元,將多出的部分存入Buf2。在下一輪重采樣結(jié)果進入對齊單元時,將Buf2中保存的數(shù)據(jù)放在新數(shù)據(jù)的頭部,同時將尾部多余的部分放入Buf2。對齊方式通過截取幀長整數(shù)倍來完成。

當解調(diào)結(jié)束之后幀同步之前,將多出幀長整數(shù)倍的部分存入Buf3。在幀同步開始時,首先將Buf3中上一輪解調(diào)結(jié)果放入新解調(diào)結(jié)果的頭部。在幀同步結(jié)束后,將多出幀長整數(shù)倍的部分存入Buf3。

當譯碼結(jié)束后,處理結(jié)果會由GPU內(nèi)存存入CPU內(nèi)存,這時首先將結(jié)果存入Buf4,當處理過若干輪次后,再將結(jié)果由Buf4存入結(jié)果數(shù)據(jù)文件。

4.3 基于GPU的并行幀同步算法

并行幀同步的任務(wù)是找到幀同步頭的位置,并將同步后的數(shù)據(jù)放入連續(xù)的內(nèi)存空間中以備譯碼或者直接存入文件。一些譯碼算法屬于軟判決譯碼算法,即需要對解調(diào)后的結(jié)果計算LLR(Log Likelihood Ratio)值,然后進行譯碼運算,本文中的LDPC譯碼算法采用的就是軟判決譯碼算法。另一些譯碼算法屬于硬判決譯碼算法,即在譯碼之前需要對解調(diào)后的信號值進行硬判決,將浮點數(shù)值轉(zhuǎn)換成0或者1的二進制數(shù)值。本文根據(jù)譯碼算法的需求,分別實現(xiàn)了用于軟判決譯碼的幀同步算法和用于硬判決譯碼的并行幀同步算法?；谲浥袥Q譯碼的幀同步算法可以根據(jù)信號的浮點值計算幀同步結(jié)果,保留信號值的原始精度,給譯碼器提供更精確的輸入值?；谟才袥Q譯碼的幀同步算法根據(jù)信號的二進制數(shù)值計算幀同步結(jié)果,占用內(nèi)存更小,速度更快。

基于GPU的幀同步算法采用了argmax算法,設(shè)幀同步頭長度為H,碼字長度為E,組成的幀長度F=H+E,計算每個偏移位置q的分數(shù),如式(12)所示:

(12)

其中,si是幀同步序列,yi是幀同步信號序列。在計算出每個偏移位置的分數(shù)后,選出得分最高或者最低的偏移位置,如式(13)所示:

m=arg max{Sq|0≤q≤F-1}

(13)

對于硬判決結(jié)果的幀同步,匹配函數(shù)如式(14)所示:

f(s,y)=s·sign(y)

(14)

對于軟判決結(jié)果的幀同步,匹配函數(shù)如式(15)所示:

(15)

4.3.1 并行滑塊分區(qū)搜索幀同步算法

在利用argmax算法尋找?guī)筋^的過程中,需要有一個移動的長度為S的滑塊,這個滑塊滑到哪里,就要計算出這個范圍內(nèi)的信號與幀同步頭之間的分值,最終根據(jù)分值的高低來確定幀同步頭的位置。

并行幀同步采用并行滑塊的方式搜索幀同步頭。假設(shè)輸入信號長度為Linput,信號的每一個位置都設(shè)置一個相應(yīng)的分值SLinput。通過申請P個線程來計算每個位置的分值,相當于建立了P個滑塊,每個線程維護一個滑塊游標,當滑塊滑動時,從當前位置開始,取長度為S的信號計算與幀同步頭的分數(shù)。計算完畢當前位置的分值后,滑塊的游標長度增加P,即跳轉(zhuǎn)到下一個位置計算相關(guān)分值,由此往復(fù)直到信號末尾。

當計算得到每個位置的分數(shù)后,需要通過比較得到的相關(guān)分值確定幀同步頭的位置。并行幀同步算法采用了分區(qū)計算最終偏移的方法。將所有位置的分數(shù)分成Z個區(qū)段,每個區(qū)段的長度為Lblock,即幀同步頭加碼字的長度。由于一個幀長范圍內(nèi)最多可能有一個幀同步頭,通過尋找最小值的方法找到幀同步頭的偏移。在比較相關(guān)分數(shù)的時候,由Z個線程分別尋找本區(qū)段長度為Lblock的相關(guān)分數(shù)最小值對應(yīng)的偏移量。計算偏移的過程由Z個線程并行完成,其中Z≥Linput/Lblock。

算法1是并行幀同步算法的偽代碼。其中,r表示信號數(shù)據(jù)向量,asm表示幀同步頭向量,Lasm表示幀同步頭長度,S表示每次參與計算的數(shù)據(jù)長度,SLinput表示分數(shù)向量,C為經(jīng)過幀同步的數(shù)據(jù)序列。

算法1并行幀同步算法

輸入：r,asm,Lasm,S,Linput。

輸出：C。

1.tx←threadIdx.x;

2.fork=tx→Linput-1paralleldo

3.fori=0→Lasm-1do

4.s=asm[i]⊕r[k+i];

5.s′=(1-asm[i])⊕r[k+i];

6.endfor

7.SLinput[k]=max(s,s′);

8.PLinput[k]=s

9.endfor

10.count←Linput/S;

11.MAX←-1;

12.iftx

13.fori=0→S-1paralleldo

14.ifSLinput[tx×S+i]

15.MAX←SLinput[tx×S+i];

16.H[tx]←tx×S+i;

17.endif

18.endfor

19.endif

20.fork=tx→Linput-1paralleldo

21.fid←k/S;

22.bid←kmodS;

23.ifPLinput[k]==0then

24.C[k]=r[H[fid]+bid];

25.else

26.C[k]=1-r[H[fid]+bid];

27.endfor

4.3.2 相位模糊識別同步算法

在MPSK進行解調(diào)的時候,容易發(fā)生相位模糊。如果輸入信號存在相位模糊,在幀同步的時候就無法得到正確的相關(guān)分值,也就無法實現(xiàn)幀同步。為了解決相位模糊的問題,本文在幀同步的過程中為輸入信號增加了一個相位向量,并且在計算相關(guān)分值的時候同時計算同相情況下的分值和反相情況下的分值,如果反向分值小于同相分值,則將反向分值保存到分值向量中,同時在相位向量中記錄分數(shù)向量對應(yīng)的相位。這樣在計算分值的時候就能夠確定是否發(fā)生了相位模糊的情況。在幀同步的最后一步通過偏移量排列幀數(shù)據(jù)的時候,根據(jù)相位向量所記錄的相位情況對幀數(shù)據(jù)的相位進行修正,最終得到原始相位的幀同步結(jié)果。

4.4 基于GPU的并行譯碼算法

本文基于CCSDS推薦的編碼方案設(shè)計了相應(yīng)的GPU并行譯碼算法。在處理程序進行初始化的時候,就已經(jīng)事先根據(jù)參數(shù)的配置為相應(yīng)的譯碼器申請好了相應(yīng)的內(nèi)存資源和計算資源,當完成幀同步運算后,進入譯碼器的數(shù)據(jù)就是一幀幀已編碼的碼字,這些碼字數(shù)據(jù)以LLR浮點數(shù)值或者二進制字符存入GPU的連續(xù)內(nèi)存中,可以批量送入譯碼器進行譯碼運算。在CCSDS推薦的幾種碼字類型中,LDPC碼和RS碼屬于線性分組碼,即所有的碼字可以按照碼長分塊,在譯碼的時候也是以碼塊作為基本的譯碼單元。而卷積碼不屬于線性分組碼,對卷積碼譯碼時沒有固定的碼長。LDPC碼和RS碼在譯碼時需要對譯碼算法進行并行優(yōu)化,保證單個碼字的計算時延達到最小,再利用GPU多線程的優(yōu)勢增加譯碼的吞吐率從而實現(xiàn)高速并行譯碼。對于卷積碼來說,在不改變譯碼算法的情況下,多線程并行帶來的速度增益并不明顯。即使如此,本文采用重疊-截斷譯碼算法,通過基于序列串行譯碼的算法并行化,實現(xiàn)了基于GPU的高速并行譯碼。

4.4.1 LDPC譯碼

LDPC譯碼算法的復(fù)雜度主要取決于在進行LDPC編碼時采用的校驗矩陣。CCSDS推薦的LDPC編碼算法可以分為2類:AR4JA和C2。校驗矩陣都是由基本矩陣組成的,圖6和圖7中的小方格就是基本矩陣通過不同步數(shù)的循環(huán)移位得到的結(jié)果。AR4JA類型的基本矩陣每一列只有一個非零值,這類LDPC碼在利用MSA算法譯碼的時候可以以基本矩陣的維度為基本并行維度實現(xiàn)并行譯碼。對于C2類型的校驗矩陣,其基本矩陣每一列有2個非零值,而且這2個非零值的距離不固定,因此無法以基本矩陣的維度作為并行維度。通過觀察可以發(fā)現(xiàn),C2類型的校驗矩陣2個非零值之間的距離有最小值,只要譯碼的時候并行維度小于這個最小值,就可以保證在譯碼時更新每個VN(Variable Node)節(jié)點的度量值時不發(fā)生沖突。例如,對于7/8碼率的C2類型的LDPC碼,基本矩陣的維度是511×511,2個非零值之間距離的最小值大于64且小于128,因此可以選擇64作為譯碼的并行維度。基于GPU的LDPC算法在以前的工作中已有闡述[5],在此不做贅述。

Figure 6 AR4JA check matrix

Figure 7 C2 check matrix

4.4.2 RS譯碼

CCSDS推薦的RS碼標準根據(jù)其糾錯性能分為(255,223)和(255,239)2種。RS碼的運算主要是伽羅華域中的運算,所有的數(shù)據(jù)都是用8位數(shù)表示,所以剛好對應(yīng)GPU中的一個char類型的字母。RS碼的主要譯碼步驟可以分為以下幾步:

(1)由接收到的R(x)求得伴隨式。

(2)利用BMA算法得錯誤位置多項式。

(3)利用Chien搜索算法求出錯誤位置多項式的根,并通過根的倒數(shù)確定錯誤位置。

(4)由錯誤位置多項式求得錯誤值,從而得到錯誤圖樣。

(5)R(X)-E(X)=C(X),完成糾錯過程。

在GPU實現(xiàn)并行RS譯碼的過程中,對求伴隨式、BMA算法和Chien搜索算法進行了并行優(yōu)化。假設(shè)RS碼的糾錯能力為t,不難發(fā)現(xiàn),對于求解伴隨式、BMA算法和Chien搜索算法來說,都依次計算2t個相關(guān)的結(jié)果,因此譯碼過程對于每個RS碼字申請了2t個線程,這樣計算2t次的值只需要在一個周期內(nèi)就能給出結(jié)果。在多碼字并行方面,通過調(diào)用多組2t線程實現(xiàn)更大程度的并行,每個線程塊至少能夠支持16個碼字的并行譯碼。

4.4.3 卷積碼

卷積碼譯碼就是根據(jù)接收序列來預(yù)測可能的發(fā)送序列。對于卷積解碼來說,就是要通過維護每條譯碼路徑的路徑度量值,最終選擇最有可能正確的路徑,最后通過回溯操作獲取譯碼結(jié)果。

本文在實現(xiàn)卷積碼譯碼時采用了如圖8所示的Viterbi譯碼算法。Viterbi譯碼算法在計算分支路徑度量值時采用了Butterfly算法,Butterfly算法的優(yōu)化是本文進行并行加速的重點。

Figure 8 Viterbi decoding algorithm

并行Butterfly算法利用了以下幾個特點:

(1)譯碼過程中一共64種狀態(tài),狀態(tài)State(n),其中n=1,2,… 64。當1≤n≤32時,在使用Butterfly算法更新對應(yīng)的分支度量值BMn的過程中,需要依賴狀態(tài)State(n+32)的分支度量值BMn+32。同理,State(n+32)的分支度量值BMn+32的更新也依賴State(n)的分支度量值BMn。

(2)對于狀態(tài)State(n),當1≤n≤32時,BMn的更新與BMp(1≤p≤32,p≠n)沒有依賴關(guān)系。同樣,對于狀態(tài)State(n),當33≤n≤64時,BMn的更新與BMp(33≤p≤64,p≠n)沒有依賴關(guān)系。

(3)采用32個線程并行計算Butterfly中的分支度量值,實現(xiàn)算法并行。

(4)通過4倍循環(huán)展開的方法,減少循環(huán)次數(shù)。

由于卷積碼基于輸入信號的序列進行譯碼,無法直接進行分段譯碼。已經(jīng)有方法證明可以通過重疊-截斷的方式進行并行分段譯碼,只需要在每個分段的頭部與前一個分段的最后部分有12 bit的重疊,最后將分段譯碼的結(jié)果進行拼接,即可得到正確的譯碼序列。

對于每個分段使用32個線程實現(xiàn)并行譯碼,以32個線程為1個線程組,通過申請多個32線程組,可以實現(xiàn)更多分段的Viterbi譯碼。分段的數(shù)量可以由用戶根據(jù)GPU支持的最大線程數(shù)和內(nèi)存資源自行設(shè)定。

5 實驗及結(jié)果分析

5.1 實驗設(shè)置

本文在基于圖靈架構(gòu)的NVIDIA?GeForce?RTXTM2080Ti上實現(xiàn)了衛(wèi)星基帶處理算法,該GPU擁有4 352個CUDA核心、68個多處理器和11 GB的GDDR6內(nèi)存。CPU為Intel?CoreTMi7-8700K。CUDA版本為10.2。實驗調(diào)制數(shù)據(jù)為遞增碼00 01..FF,數(shù)據(jù)大小采用4字節(jié)浮點數(shù),采樣速率為40 MHz,符號速率為5 Mbps,幀同步頭為4字節(jié)1ACFFC1D。

表1展示了不同部分的處理速率和處理延遲。從圖1可以看到,LDPC譯碼速度達到1 998 Mbps,RS譯碼器達到3 853 Mbps,Viterbi譯碼器達到822 Mbps。

Table 1 Delay of different parts of baseband

表2展示了不同的解調(diào)參數(shù)和譯碼參數(shù)下基帶的處理速度。

Table 2 Comparison of processing rates under different demodulation and decoding parameters

5.2 實驗結(jié)果

目前還沒有基于GPU的衛(wèi)星通信基帶的相關(guān)工作。本文比較了基于GPU的LDPC譯碼器的實現(xiàn)性能。表3是本文實現(xiàn)的LDPC譯碼器和其他LDPC譯碼器的性能比較。由表3可知,本文提出的譯碼器實現(xiàn)了1.998 Gbps的吞吐率,超過了其他相關(guān)工作[5-8]的。為了體現(xiàn)不同GPU平臺比較的公平性,本文利用“利用率參數(shù)=吞吐率/峰值性能”來衡量加速器的性能。從表3可以看出,本文針對GPU峰值性能的利用率也是最高的。

Table 3 Performance comparison of different GPU-based decoders

6 結(jié)束語

本文提出了一種基于GPU的衛(wèi)星通信基帶算法群,實現(xiàn)了重采樣、匹配濾波、解相位模糊、幀同步、解擾、LDPC碼、RS碼和卷積碼等并行算法,最終實現(xiàn)的下行鏈路基帶處理速率在170 Mbps～978 Mbps,有效提升了衛(wèi)星通信信號處理速度。