基于GPU的衛(wèi)星通信基帶處理高吞吐率并行算法*

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

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

1 引言

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

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

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

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

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

3 相關工作

近年來,基于GPU的無線通信并行算法引起學術界和產業(yè)界越來越多的重視。大多數工作聚焦于單個無線通信算法的GPU加速,如基于GPU的同步算法[2]、Viterbi算法[3]、解調算法[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)的全球導航衛(wèi)星系統(tǒng)(GNSS)GPU接收機[14]、全球導航衛(wèi)星系統(tǒng)反射測量GNSS-R(Global Navigation Satellite System-Reflectometry)GPU接收機[15]、基于多GPU負載均衡的GNSS接收機[16,17]、基于GPU的GNSS多徑并行仿真器[18]、基于MPI+CUDA的遙感衛(wèi)星數據處理通用系統(tǒng)[19]和基于GPU的信號相關器的GNSS接收機。以上大多數研究都是針對接收機鏈路中的部分算法進行并行優(yōu)化,本文基于GPU實現GNSS衛(wèi)星通信基帶接收機完整下行鏈路算法群。

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

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

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

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

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

Figure 2 Division of CPU and GPU computing tasks

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

4.2 基于GPU的并行解調算法

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

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

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

4.2.1 并行位同步算法

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

圖3所示的算法是Gardner位同步算法的主要過程。其中,x(t)為輸入信號,經過周期為Ts的固定時鐘采樣后,得到x(mTs),m=1,2,…表示離散序列編號;插值濾波器、環(huán)路濾波器、定時誤差檢測器和數控振蕩器組成了一個NCO環(huán)路,NCO環(huán)路的作用是生成采樣間隔。插值濾波器通過插值運算得到k時刻的結果y(kTi),其中Ti為輸出信號的采樣周期;τ(n)為計算得到的定時誤差;e(n)經過環(huán)路濾波后的校準誤差;mk為NCO得到的采樣步長;μk為根據k時刻插值結果計算得到的小數間隔。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寄存器過零點時,就是對輸入信號進行重采樣的位置,該位置相對于輸入信號的時鐘周期的整數倍多出的部分即為重采樣的小數間隔μ。

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

l=R(k)/wc

(2)

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

l=sTs+μ

(3)

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

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

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

(4)

(5)

(6)

(7)

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

y(k)=Cx

(8)

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

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

并行計算內插值的過程可以分為2步。第1步是通過插值多項式的系數矩陣與內插點小數偏移μ的冪次向量的乘積。第2步是插值系數向量和信號向量的乘積。即4×4的矩陣與長度為4的μ的冪次向量的乘積,即[1,μ,μ2,μ3],以及2個長度為4的向量的乘積?？梢允褂?個線程分別計算4個插值系數,并且分別計算系數與信號值的乘積,經過2次歸約運算后即可得到內插值。

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

每當計算得到2個內插結果后,都要對時鐘相位進行校正,也就是計算更新NCO控制步長wc的值。對于不同的解調方式,校正方法有所區(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路分別計算得到的內插值。除了計算時鐘相位誤差之外,在得到插值結果后,還要計算相應的載波相位誤差。在更新相位誤差后,NCO控制步長wc和載波相位θ都發(fā)生了變化,新的插值結果需要用新的控制步長和載波相位進行后續(xù)的計算。因此,不能一次計算過多的內插值,假設校準間隔為M,當M=1時,每計算得到2個內插結果校正一次相位誤差;當M≥2時,就需要考慮未及時檢測更新相位誤差帶來的影響。

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

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

(10)

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

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

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

4.2.2 并行濾波

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

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

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

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

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

L=N-D+1

(11)

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

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

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

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

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

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

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

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

Figure 4 Filtering algorithm

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

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

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

Figure 5 Design of CPU-GPU heterogeneous buffer

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

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

第3級緩沖區(qū)Buf3位于GPU內存,用于保存超過幀長整數倍的解調結果;

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

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

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

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

當譯碼結束后,處理結果會由GPU內存存入CPU內存,這時首先將結果存入Buf4,當處理過若干輪次后,再將結果由Buf4存入結果數據文件。

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

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

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

(12)

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

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

(13)

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

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

(14)

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

(15)

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

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

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

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

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

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

4.4 基于GPU的并行譯碼算法

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

4.4.1 LDPC譯碼

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

Figure 6 AR4JA check matrix

Figure 7 C2 check matrix

4.4.2 RS譯碼

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

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

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

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

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

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

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

4.4.3 卷積碼

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

本文在實現卷積碼譯碼時采用了如圖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算法更新對應的分支度量值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)沒有依賴關系。同樣,對于狀態(tài)State(n),當33≤n≤64時,BMn的更新與BMp(33≤p≤64,p≠n)沒有依賴關系。

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

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

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

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

5 實驗及結果分析

5.1 實驗設置

本文在基于圖靈架構的NVIDIA?GeForce?RTXTM2080Ti上實現了衛(wèi)星基帶處理算法,該GPU擁有4 352個CUDA核心、68個多處理器和11 GB的GDDR6內存。CPU為Intel?CoreTMi7-8700K。CUDA版本為10.2。實驗調制數據為遞增碼00 01..FF,數據大小采用4字節(jié)浮點數,采樣速率為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展示了不同的解調參數和譯碼參數下基帶的處理速度。

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

5.2 實驗結果

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

Table 3 Performance comparison of different GPU-based decoders

6 結束語

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