劉燕都，鄭海昕，王小平

(1.裝備學(xué)院,北京 101416；2.76160部隊(duì),廣東 廣州 510000)

?

基于CUDA的數(shù)字調(diào)相信號并行解調(diào)程序設(shè)計(jì)*

劉燕都1，鄭海昕1，王小平2

(1.裝備學(xué)院,北京 101416；2.76160部隊(duì),廣東 廣州 510000)

摘要：在通用計(jì)算機(jī)上實(shí)現(xiàn)數(shù)字通信的接收解調(diào)是近年來信號處理領(lǐng)域的重要研究方向。就實(shí)時性要求較高的數(shù)字調(diào)相信號解調(diào)問題，分析了算法的可并行性，研究了在通用計(jì)算機(jī)上實(shí)現(xiàn)的算法。針對“CPU+GPU”異構(gòu)計(jì)算平臺特點(diǎn)，提出了數(shù)字調(diào)相信號并行計(jì)算模型，基于CUDA平臺設(shè)計(jì)了混頻器、鑒相器、濾波器等模塊的并行程序，實(shí)現(xiàn)了BPSK信號解調(diào)。測試結(jié)果表明，計(jì)算時間比為1:1.7，在SNR=9.6 dB時誤碼率可以達(dá)到10-5，與專有硬件解調(diào)數(shù)字調(diào)相信號的指標(biāo)相當(dāng)，但通用計(jì)算機(jī)平臺實(shí)現(xiàn)方法更為靈活、易于功能擴(kuò)展。

關(guān)鍵詞：數(shù)字調(diào)相信號；解調(diào)；CUDA；并行

0引言

近年來，隨著通用計(jì)算機(jī)性能的不斷提高，數(shù)字通信系統(tǒng)中信號處理平臺在經(jīng)歷了由模擬器件構(gòu)建的硬件平臺向軟件無線電技術(shù)的數(shù)字化平臺的轉(zhuǎn)變后，正在開始向純軟件化的方向發(fā)展，將采樣后的信號直接送入通用計(jì)算機(jī)，以軟件處理的方式完成信號的解調(diào)、解碼等。

數(shù)字相位調(diào)制，也就是相移鍵控(Phase Shift Keying，PSK)，是一種十分重要的基本數(shù)字調(diào)制技術(shù)，也是一種用載波相位表示輸入信號信息的調(diào)制技術(shù)。在時不變信道中，調(diào)相信號比調(diào)幅和調(diào)頻信號具有更高的抗噪聲性能和頻帶利用率，即使在有衰落信道、多徑和強(qiáng)干擾的情況下也有較好的效果。因此，數(shù)字相位調(diào)制是一種性能優(yōu)良的調(diào)制方式，在衛(wèi)星通信、高速數(shù)據(jù)傳輸中得到了廣泛應(yīng)用。根據(jù)理想軟件無線電的發(fā)展趨勢[1]，通信過程中，各項(xiàng)功能都應(yīng)由軟件來實(shí)現(xiàn)。本文基于高性能計(jì)算平臺，研究了調(diào)相信號的軟件化并行解調(diào)算法，以及基于CUDA的并行程序設(shè)計(jì)方法，并對調(diào)相信號并行解調(diào)程序進(jìn)行測試，驗(yàn)證了調(diào)相信號并行解調(diào)的可行性。此處填入前言內(nèi)容 。

1調(diào)相信號解調(diào)

調(diào)相信號解調(diào)方法一般可分為相干解調(diào)和非相干解調(diào)[2]。相干解調(diào)需要恢復(fù)與輸入信號同頻同相的信號，而非相干解調(diào)只需提取同頻信號，雖然非相干解調(diào)實(shí)現(xiàn)方式簡單，但相干解調(diào)因其恢復(fù)的信號與輸入信號相位嚴(yán)格對齊，因此解調(diào)效果更好，應(yīng)用更為廣泛。但調(diào)相信號在調(diào)制過程中抑制了載波分量，若要提取同頻同相的載波進(jìn)行解調(diào)一般采用特殊設(shè)計(jì)的鎖相環(huán)[3]。如平方環(huán)法、Costas環(huán)[4]等。還可以根據(jù)相鄰碼元之間相位跳變情況采用差分解調(diào)方法[5]，算法相對簡單，但其抗噪聲性能要明顯劣于相干解調(diào)。數(shù)字Costas環(huán)結(jié)構(gòu)見圖1。

圖1　數(shù)字Costas環(huán)結(jié)構(gòu)

以二進(jìn)制數(shù)字調(diào)相信號為例，其信號經(jīng)采樣后[6]：

(1)

(2)

(3)

(4)

式中，K11,K12為低通濾波器系數(shù)，經(jīng)濾波后的I、Q兩路信號經(jīng)相乘鑒相后，由環(huán)路濾波，可得：

(5)

式中，Kp為鑒相增益，Kd為環(huán)路增益，環(huán)路濾波器輸出為NCO的頻率控制字。

當(dāng)環(huán)路鎖定后，提取出同頻同相的相干載波后，將其與輸入的已調(diào)信號直接相乘，并濾波輸出，即可得到基帶信號波形，如圖2所示，解調(diào)和鎖相環(huán)跟蹤如圖3所示。

圖2　數(shù)字調(diào)相信號時域波形及頻譜

圖3　解調(diào)結(jié)果及鎖相環(huán)路頻率跟蹤曲線

2基于CUDA的并行計(jì)算模型

分析Costas環(huán)結(jié)構(gòu)，可知不僅I、Q兩路可以實(shí)現(xiàn)任務(wù)級的并行計(jì)算，其中的鑒相器、混頻器和濾波器還可以實(shí)現(xiàn)數(shù)據(jù)級的并行計(jì)算，因此，研究在特定并行計(jì)算平臺上的計(jì)算方法。

2.1CUDA

CUDA[7]是由NVIDIA公司推出的一種通用并行計(jì)算架構(gòu)。起初是為加速圖像實(shí)時處理而設(shè)計(jì)的一種運(yùn)行在GPU上的開發(fā)平臺，其充分運(yùn)用了GPU的高存儲帶寬和超大規(guī)模的浮點(diǎn)計(jì)算單元，現(xiàn)在被廣泛應(yīng)用于大型并行化問題，如氣象模擬、地震預(yù)報、分子計(jì)算等。CUDA硬件架構(gòu)如圖4所示。

圖4　CPU和GPU結(jié)構(gòu)對比

從圖4不難看出，GPU是特別為計(jì)算密集，高并行度計(jì)算設(shè)計(jì)的，因此將更多的片上資源用于計(jì)算而不是數(shù)據(jù)緩存和邏輯。特別地，從GPU結(jié)構(gòu)分析，其非常適合處理SIMD(單指令多處理)并行問題，即同一程序在多個數(shù)據(jù)上并行執(zhí)行的問題，而數(shù)字信號處理具備這種特征，所以在CUDA平臺非常適合進(jìn)行數(shù)字信號處理。

2.2并行計(jì)算模型

根據(jù)GPU的硬件設(shè)計(jì)特點(diǎn)，CUDA在并行算法設(shè)計(jì)層做出了較為細(xì)致的約束[8]。模型假設(shè)CUDA線程在物理上獨(dú)立的GPU上執(zhí)行，GPU作為主機(jī)的協(xié)處理器，采取異構(gòu)并行的模式，并行計(jì)算的內(nèi)核程序在GPU上執(zhí)行，程序的其它部分在CPU上執(zhí)行。因?yàn)镚PU設(shè)備不具備顯示功能，因此數(shù)據(jù)需要在顯存和內(nèi)存之間由PCIe總線進(jìn)行交互，受限于通用計(jì)算機(jī)的速度限制，在大規(guī)模數(shù)值計(jì)算中，數(shù)據(jù)傳輸時間占了程序大部分執(zhí)行時間，圖5所示為不同數(shù)據(jù)規(guī)模下并行加的計(jì)算時間，在數(shù)據(jù)規(guī)模較小時，數(shù)據(jù)幾乎占了程序總執(zhí)行時間99%以上，因此，只有在計(jì)算規(guī)模較大時，才能體現(xiàn)GPU計(jì)算的優(yōu)勢。

圖5　并行求和運(yùn)算時間對比

根據(jù)CUDA平臺這一特點(diǎn)，數(shù)字通信系統(tǒng)并行計(jì)算模型中，應(yīng)盡量減少數(shù)據(jù)傳輸，充分發(fā)揮GPU高性能計(jì)算能力。如圖6所示，程序執(zhí)行開始，將需要處理的數(shù)據(jù)全部轉(zhuǎn)至顯存，全部大規(guī)模計(jì)算由GPU完成，CPU和GPU在程序執(zhí)行過程中，只進(jìn)行少量數(shù)據(jù)傳遞，CPU只進(jìn)行小規(guī)模的計(jì)算和數(shù)據(jù)監(jiān)視和顯示功能。

圖6　基于GPU的數(shù)字通信并行計(jì)算模型

3數(shù)字調(diào)相信號并行解調(diào)程序設(shè)計(jì)

3.1數(shù)字調(diào)相信號解調(diào)算法結(jié)構(gòu)

根據(jù)2.2節(jié)的并行計(jì)算模型和Costas環(huán)解調(diào)結(jié)構(gòu)，數(shù)字調(diào)相信號并行解調(diào)算法[9-11]如圖7示。

圖7　基于CUDA的數(shù)字調(diào)相信號并行解調(diào)算法模型

中頻采樣數(shù)據(jù)從CPU讀取并傳輸?shù)紾PU，在GPU中完成數(shù)字正交下變頻、低通濾波、位同步、鑒相、環(huán)路濾波和譯碼；環(huán)路濾波后的相位誤差信號傳輸回CPU計(jì)算多普勒頻移，多普勒頻移再傳回GPU修正NCO輸出的正、余弦波形；GPU譯碼后的數(shù)據(jù)再傳回CPU進(jìn)行顯示和存儲。見圖8。

圖8　算法流程

3.2混頻器設(shè)計(jì)

混頻器就是把信號從中頻搬移到基頻，是軟件無線電的核心，通常硬件上采用數(shù)控振蕩器(NCO)產(chǎn)生本地數(shù)字載波進(jìn)行混頻。并行混頻時，對應(yīng)數(shù)據(jù)點(diǎn)與對應(yīng)相位的正弦和余弦波形采樣點(diǎn)做乘法，并行程序算法如圖9所示。

圖9　變頻器算法流程

混頻器偽代碼如下：

__global__ voidDownfreqKernel()

Begin

int tid = blockIdx.x * blockDim.x + threadIdx.x;

if (tid

downfreq[tid].x = chandata[tid]*cos(2*PI*(fb+phasefd));

downfreq[tid].y = chandata[tid]*sin(2*PI*(fb+phasefd));

End

3.3濾波器設(shè)計(jì)

FIR濾波器以其良好的群延遲性被廣泛應(yīng)用在數(shù)字通信系統(tǒng)中，其可以保證任意幅頻特性的同時具有嚴(yán)格的線性相頻特性，同時其具有有限長沖擊響應(yīng)，算法如圖10所示。

圖10　濾波器算法流程

濾波器偽代碼如下：

__global__ void GPUFilterKernel()

Begin

__shared__ float2 cache[];

inttid = threadIdx.x + blockIdx.x * blockDim.x;

cache[threadIdx.x].x=signal[tid].x;

cache[threadIdx.x].y=signal[tid].y;

__syncthreads();

float sumx=(cache[threadIdx.x].x+

……+cache[threadIdx.x+16].x)*f.a；

float sumx=(cache[threadIdx.y].y+

……+cache[threadIdx.x+16].x)*f.a；

result[tid].x=sumx;

result[tid].y=sumy;

End

3.4鑒相器設(shè)計(jì)

鑒相器主要完成鑒別輸入信號相差的功能，是鎖相環(huán)路的關(guān)鍵，在并行程序設(shè)計(jì)中，主要依靠求解前后采樣點(diǎn)相差的辦法，算法如圖11所示。

圖11　鑒相器算法流程

鑒相器偽代碼如下：

__global__ void ComputeSubphaseKernel()

Begin

inttid = threadIdx.x + blockIdx.x * blockDim.x;

data[tid] = atan2(-src[tid].y,src[tid].x);

data[tid] = src[tid+1]-src[sid];

if (data[tid]>PI)

data[tid] -= 2*PI;

else if (data[tid]<-PI)

data[tid] += 2*PI;

End

3.5測試結(jié)果

測試硬件平臺選用NIVIDIA Tesla K20顯卡，輸入數(shù)據(jù)1 ms模擬數(shù)據(jù)，計(jì)算時間在1.7 ms以內(nèi)，如圖12所示，程序能夠正確解調(diào)原數(shù)據(jù)，在Eb/N0=9.6 dB時誤碼率可以達(dá)到10-5。

圖12　基于CUDA的BPSK信號并行解調(diào)結(jié)果

4結(jié)語

在通用計(jì)算機(jī)平臺上實(shí)現(xiàn)數(shù)字調(diào)相信號的接收解調(diào)，降低了系統(tǒng)設(shè)計(jì)、開發(fā)的難度和成本，軟件化的處理方式增加了系統(tǒng)的靈活性，通過加載不同的軟件，可以實(shí)現(xiàn)更多的功能。通過硬件的升級和重組，還可以進(jìn)一步提高系統(tǒng)性能?？梢娀谕ㄓ糜?jì)算機(jī)平臺，尤其是基于CUDA的數(shù)字信號處理是信號處理的一個重要發(fā)展方向，也是計(jì)算機(jī)應(yīng)用的新趨勢和研究的新領(lǐng)域。

參考文獻(xiàn)：

[1]王曉琴,黑勇.軟件無線電硬件體系結(jié)構(gòu)研究[J].科學(xué)技術(shù)與工程,2006,16(13):1820-1824.

WANG Xiao-qin,HEI Yong.A Study of Software Radio Hardware Architecture[J].Science Technology and Engineering,2006,16(13):1820-1824.

[2]Sklar Bernard.Digital Communications[M].北京: 電子工業(yè)出版社,2006:538-542.Sklar Bernard.Digital Communications[M].Beijing: Publishing House of Electronics Industry,2006:538-542.

[3]劉艷華.Costas環(huán)法BPSK信號解調(diào)的研究與實(shí)現(xiàn)[J].通信技術(shù),2012,45(01):16-21.

LIU Yan-hua.Research and Implementation of BPSK signal Demodulation based on Costas[J].Communications Technology.2012,45(01):16-21.

[4]季仲梅,楊洪生,王大鳴等.通信中的同步技術(shù)及應(yīng)用[M].北京:清華大學(xué)出版社,2008:111-117.

JI Zhong-mei,YANG Hong-sheng,WANG Da-ming,et al.Synchronization Technology and Application in Communication[M].Beijing: Tsinghua University Press,2008:126-128.

[5]方浩華,王躍林,徐會勤等.基于DSP的DPSK差分解調(diào)的實(shí)現(xiàn)與研究[J].移動通信,2003增刊:79-84.

FANG Hao-hua,WANG Yue-lin,XU Hui-qin,et al.Research on DPSK Difference Demodulation based on the DSP Implementation[J].Mobile Communications,2003(Supplement.):79-83.

[6]Riter S.An Optimum Phase Reference Detector for Fully Modulated Phase Shift Keyed Signal[C].IEEE AES-5,1969,4(7):11-17.

[7]NVIDIA Corp.NVIDIA CUDA Programming Guide 5.0[S].http://www.nvidia.com/object/cuda_Develop.html.

[8]陳國良,孫廣中,徐云等.并行算法研究方法學(xué)[J].計(jì)算機(jī)學(xué)報,2008,31(09):1493-1502.

CHEN Guo-liang,SUN Guang-zhong,Xu Yun,et al.Methodology of Research on Parallel Algorithms[J].Chinese Journal of Computers,2008,31(09):1493-1502.

[9]Core Mark T,Tan Harry H.BER for Optical Heterodyne DPSK Receivers Using Delay Demodulation and Integration Detection[J].IEEE Transactions on Communications,2002,50(1):1451-1459.

[10]LI Gui-xin.,AN Zhi-qi.,YUAN Si-jie.Study on Software Demodulation of DQPSK Signal based on Digital Phase Measurement[J].Journal of Spacecraft TT&C Technology.2008,27(2): 1105-1110.

[11]Mitra S K.Digital Signal Processing,A Computer-Based Approach[M].Second Edition,MeGraw-Hill Companies,Inc.,2001.

Parallel Programming of PSK Signal Demodulation based on CUDA

LIU Yan-du1，ZHENG Hai-xin1，WANG Xiao-ping2

(1.Equipment Academy,Beijing 101416,China;2.Unit 76160 of PLA,Guangzhou Guangdong 510000,China)

Abstract：To implement digital communication receiving/demodulation on the general computer is an important research direction in the field of signal processing in recent years.Parallelism of the algorithm for BPSK signal demodulation with high real-time requirements is analyzed and implemented on the general computer.Based on the characteristics of "CPU + GPU" heterogeneous computing platform,the parallel computation model for digital phase modulation signal is proposed,and based on CUDA platform,the parallel-program modules of mixer,phase discriminator and filter are designed,thus to realize the BPSK signal demodulation.Experiment results show that the computing time ratio of 1:1.7,the BER of 10-5when SNR=9.6 dB could be reached,well-matched with those by special hardware demodulation,and in particular,the implementation method by general computer is more flexible,and easier for function extension.

Key words：digital phase modulation signal; demodulation ; CUDA; parallel

doi:10.3969/j.issn.1002-0802.2016.02.020

* 收稿日期：2015-09-10；修回日期：2015-12-20Received date:2015-09-10；Revised date：2015-12-20

中圖分類號：TN911.7

文獻(xiàn)標(biāo)志碼：A

文章編號：1002-0802(2013)08-0227-06

作者簡介：

劉燕都(1986—)，男，碩士研究生，主要研究方向?yàn)楦咚贁?shù)字信號處理，航天測控技術(shù)；

鄭海昕(1974—)，女，碩士，副教授，主要研究方向?yàn)楹教鞙y控技術(shù)，空間數(shù)據(jù)傳輸；

王小平(1987—)，男，學(xué)士，主要研究方向?yàn)楦咚傩盘柼幚?，空間信息傳輸。