一种基于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。