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

(國(guó)防科技大學(xué)計(jì)算機(jī)學(xué)院并行與分布處理國(guó)防重點(diǎn)實(shí)驗(yàn)室,湖南 長(zhǎng)沙 410073)

1 引言

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

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

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

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

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

3 相關(guān)工作

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

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

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

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

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

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

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

Figure 2 Division of CPU and GPU computing tasks

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

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

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

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

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

4.2.1 并行位同步算法

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

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

Figure 3 Synchronization algorithm

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

(1)

其中wc是控制步長(zhǎng)。當(dāng)NCO寄存器的值過(guò)零點(diǎn)時(shí),寄存器的值模1,即得到下一個(gè)符號(hào)周期NCO寄存器的值。當(dāng)NCO寄存器過(guò)零點(diǎn)時(shí),就是對(duì)輸入信號(hào)進(jìn)行重采樣的位置,該位置相對(duì)于輸入信號(hào)的時(shí)鐘周期的整數(shù)倍多出的部分即為重采樣的小數(shù)間隔μ。

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

l=R(k)/wc

(2)

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

l=sTs+μ

(3)

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

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

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

(4)

(5)

(6)

(7)

其中,μ是內(nèi)插點(diǎn)的時(shí)間偏移,C-2、C-1、C0、C1是立方插值濾波器的系數(shù)。當(dāng)計(jì)算得到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時(shí)刻的信號(hào)的采樣值。

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

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

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

每當(dāng)計(jì)算得到2個(gè)內(nèi)插結(jié)果后,都要對(duì)時(shí)鐘相位進(jìn)行校正,也就是計(jì)算更新NCO控制步長(zhǎng)wc的值。對(duì)于不同的解調(diào)方式,校正方法有所區(qū)別。對(duì)于QPSK來(lái)說(shuō),時(shí)鐘相位誤差由式(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路分別計(jì)算得到的內(nèi)插值。除了計(jì)算時(shí)鐘相位誤差之外,在得到插值結(jié)果后,還要計(jì)算相應(yīng)的載波相位誤差。在更新相位誤差后,NCO控制步長(zhǎng)wc和載波相位θ都發(fā)生了變化,新的插值結(jié)果需要用新的控制步長(zhǎng)和載波相位進(jìn)行后續(xù)的計(jì)算。因此,不能一次計(jì)算過(guò)多的內(nèi)插值,假設(shè)校準(zhǔn)間隔為M,當(dāng)M=1時(shí),每計(jì)算得到2個(gè)內(nèi)插結(jié)果校正一次相位誤差;當(dāng)M≥2時(shí),就需要考慮未及時(shí)檢測(cè)更新相位誤差帶來(lái)的影響。

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

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

(10)

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

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

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

4.2.2 并行濾波

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

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

基于GPU的并行卷積運(yùn)算步驟如下所示:

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

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

L=N-D+1

(11)

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

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

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

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

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

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

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

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

Figure 4 Filtering algorithm

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

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

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

Figure 5 Design of CPU-GPU heterogeneous buffer

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

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

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

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

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

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

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

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

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

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

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

(12)

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

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

(13)

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

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

(14)

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

(15)

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

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

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

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

算法1是并行幀同步算法的偽代碼。其中,r表示信號(hào)數(shù)據(jù)向量,asm表示幀同步頭向量,Lasm表示幀同步頭長(zhǎng)度,S表示每次參與計(jì)算的數(shù)據(jù)長(zhǎng)度,SLinput表示分?jǐn)?shù)向量,C為經(jīng)過(guò)幀同步的數(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 相位模糊識(shí)別同步算法

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

4.4 基于GPU的并行譯碼算法

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

4.4.1 LDPC譯碼

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

Figure 6 AR4JA check matrix

Figure 7 C2 check matrix

4.4.2 RS譯碼

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

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

(2)利用BMA算法得錯(cuò)誤位置多項(xiàng)式。

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

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

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

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

4.4.3 卷積碼

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

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

Figure 8 Viterbi decoding algorithm

并行Butterfly算法利用了以下幾個(gè)特點(diǎn):

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

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

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

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

由于卷積碼基于輸入信號(hào)的序列進(jìn)行譯碼,無(wú)法直接進(jìn)行分段譯碼。已經(jīng)有方法證明可以通過(guò)重疊-截?cái)嗟姆绞竭M(jìn)行并行分段譯碼,只需要在每個(gè)分段的頭部與前一個(gè)分段的最后部分有12 bit的重疊,最后將分段譯碼的結(jié)果進(jìn)行拼接,即可得到正確的譯碼序列。

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

5 實(shí)驗(yàn)及結(jié)果分析

5.1 實(shí)驗(yàn)設(shè)置

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

表1展示了不同部分的處理速率和處理延遲。從圖1可以看到,LDPC譯碼速度達(dá)到1 998 Mbps,RS譯碼器達(dá)到3 853 Mbps,Viterbi譯碼器達(dá)到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 實(shí)驗(yàn)結(jié)果

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

Table 3 Performance comparison of different GPU-based decoders

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

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