基于嵌入式GPU的三天线GNSS基带信号处理加速技术

杨智博，金 天

(北京航空航天大学电子信息工程学院，北京 100191)

0 引言

近年来，由于全球定位系统(Global Positioning System，GPS)、俄罗斯全球卫星导航系统(Global Navigation Satellite System，GLONASS)的现代化升级和伽利略卫星导航系统(Galileo Satellite Navigation System)、中国北斗卫星导航系统(BeiDou Navigation Satellite System，BDS)的不断建设和完善，全球卫星导航系统(Global Navigation Satellite System，GNSS)正在发生着巨大的变化。未来四大卫星导航(GPS、BDS、GLONASS、GALILEO)以及准天顶卫星系统(Quasi-Zenith Satellite System，QZSS)、印度区域导航卫星系统(Indian Regional Navigation Satellite System， IRNSS)和广域增强系统(Wide Area Augmentation System，WAAS)等众多增强型导航系统将播发多种频率的民用信号[1]，在提高卫星导航性能的同时，也对接收机的设计提出了新的挑战。

传统的以现场可编程门阵列(Field Programmable Gate Array，FPGA)或者专用集成电路(Application Specific Integrated Circuit，ASIC)为代表的硬件接收机，由于其开发周期长和不够灵活的缺点，难以满足用户和运营商不断变化的需求[2]。在这种背景下，软件接收机受到越来越多来自业内的关注。软件接收机具有很高的灵活性，能够以很少的额外组件实现低成本，并且可以根据用户的需求便捷的更改和升级[3]。但传统的软件接收机存在计算量大、耗时长的捕获和跟踪基带信号处理过程，使得软件接收机往往跟踪通道较少，且难以在嵌入式系统上运行，因此基带信号处理过程的加速是软件接收机研究的重点。

随着图形处理器(Graphics Processing Unit，GPU)技术的不断发展，其应用领域由图形领域扩展到高性能计算领域，使得利用GPU加速软件接收机成为可能，相关研究也逐渐增多。Hobiger T等使用GPU处理GPS C/A码信号的捕获和跟踪，基于Intel Core 2 Q9450 CPU 和NVIDIA GeForce GTX 280 GPU实现了8MHz采样率、4bit量化且支持12个卫星通道的实时GPS软件接收机[4]。Huang B等使用Intel i7-3770K CPU和NVIDIA GeForce GTX 580 GPU实现了多系统实时软件接收机，可跟踪11种卫星信号[3]。Xu L等基于NVIDIAGeForce GTX 750 Ti GPU实现了一款模块化GPS软件接收机NAVSDR，支持300余个跟踪通道以9.75MHz采样率和4bit量化实时运行[5]。Park K W等基于GPU和CPU异构的嵌入式平台Jetson TK1实现了一款软件接收机，支持10个跟踪通道以10MHz采样率实时运行[6]。

虽然在利用GPU加速软件接收机方面的研究越来越多，但多数研究均针对PC，针对嵌入式平台的相关研究较少。因此，在嵌入式平台上实现支持高采样率、多卫星通道的实时GNSS软件接收机仍亟待研究。

本文对基于嵌入式平台的GNSS软件接收机的实现进行了研究。为满足软件接收机的实时性、高采样率以及多卫星通道的要求，本文采用嵌入式GPU对耗时长且并行性明显的数据读取模块、捕获模块和跟踪模块进行了加速，研制了三天线GNSS定位和测姿接收机的样机，并对其进行了加速性能和样机性能的测试。

1 三天线GNSS软件接收机整体架构

1.1 总体结构设计

如图1所示，三天线GNSS信号定位和测姿接收机包含如下4个模块：射频前端、基带信号处理、导航定位解算和姿态测量。射频前端将射频模拟信号转换成数字基带信号，3个天线的信号通过采集设备转换成数字基带信号；基带信号处理模块包括捕获、跟踪和位同步模块，从基带信号中获得导航电文数据比特；导航定位解算模块从导航电文中提取出相关信息，进而解算出接收机的位置、速度和时间信息；姿态测量模块根据载波相位观测值求差，得到载体的姿态信息。

图1 接收机框架图Fig.1 Receiver frame

1.2 射频前端设计

采集设备的主要任务是将3个天线接收到的射频模拟信号离散成包含卫星信号成分的、频率较低的数字中频信号，并且在此过程中进行必要的滤波和增益控制。

采集设备选用 MAX2771 芯片，能够以较低的成本提供最高性能和集成度，可以覆盖 E5/L5、 L2、E6、E1/L1 波段，支持 GPS、GLONASS、Galileo、QZSS、IRNSS 和 BDS。通过采集设备上的MicroUSB接口虚拟串口，可以配置卫星信号类型、频点、带宽，满足多个GNSS的卫星信号在信号形式、频点、带宽等方面的不同要求。

3个天线的模拟信号经采集设备转换为数字中频信号后，数据读取模块读取采样数据到缓存中，然后根据采样数据的存储格式拆分为接收机各个系统需要的中频数据。

1.3 基带信号处理模块设计

基带信号处理是软件接收机中重要的一环，其中包括捕获、跟踪和位同步模块。通过基带信号处理模块，接收机从数字中频信号中获得导航电文数据比特，为后续导航定位解算做准备。

1.3.1 捕获模块设计

捕获是针对码相位和多普勒频移的二维搜索过程，即捕获的目的是确定能观测到的卫星，并估计出粗略的伪码相位和载波多普勒频移。目前常用的捕获算法有串行捕获、并行频率捕获和并行码相位捕获。其中，并行码相位捕获算法将数字相关器的相关运算替换为快速傅立叶变换[7]，减少了运算时间，在实时软件接收机中被广泛使用。

并行码相位捕获算法的流程如图2所示。数字中频输入信号与本地载波信号相乘，分别得到I、Q支路。I、Q支路经傅里叶变换，得到频域信号，与本地伪码的频域复共轭相乘，再进行傅里叶逆变换，取模，查找出最大相关峰值。该最大相关峰值对应的伪码相位和载波多普勒频移即为捕获模块的目标值。

图2 并行码相位算法图Fig.2 Phase algorithm diagram of parallel code

软件接收机的捕获模块使用并行码相位算法。各系统第一次调用捕获模块时，对其对应的导航系统的全部卫星遍历捕获，之后每1s捕获1次，每次捕获1颗未被成功跟踪的卫星。

1.3.2 跟踪模块设计

经由捕获模块，可以得到可见卫星信号的载波频率(包括载波多普勒偏移)和伪码相位的粗略估计值。但它们的精度无法达到解调导航信息的要求，同时由于在接收卫星信号的过程中，如卫星持续运动、接收机运动、时钟晶振钟漂和随机抖动等多种因素都会引起信号伪码相位和载波频率的变化，因此接收机在捕获到某颗卫星的信号后，将会在跟踪模块对载波频率和伪码相位值进行进一步处理，得到它们的精确值。

如图3所示，软件接收机的跟踪模块分为4个子模块：相关器子模块、PLL子模块、DLL子模块和NCO子模块。跟踪模块的更新周期设置为40ms，子模块的更新周期设置为1ms。即每次调用跟踪模块时，子模块循环运行40次。

图3 跟踪模块架构图Fig.3 Tracking module architecture

相关器子模块根据载波频率和载波相位得到本地载波同相、正交三角函数表；根据伪码相位和码间隔得到本地三路超前、即时、滞后伪码码表；采用比特并行快速算法进行载波剥离和相关解扩， 最终得到E路、P路、L路(超前、即时、滞后)的同相正交共6个相关值。PLL跟踪子模块采用对导航电文翻转不敏感的科斯塔(Costas)锁相环，根据接收信号载波相位和本地载波相位的差值对本地载波相位进行调整，最终实现对载波相位的跟踪，达到接收信号载波和本地载波相位的动态同步。DLL跟踪子模块选用延迟锁定环，根据接收信号伪码和本地伪码的相位差异对本地伪码相位进行调整，最终实现对伪码相位的跟踪。在经由以上模块对基带信号中的载波和伪码进行剥离之后，NCO跟踪子模块存储即时的同相相关值，组成1ms宽数据比特流，以供后续位同步模块进行处理。

1.3.3 位同步模块设计

接收机的位同步模块采用直方图算法，通过对跟踪环路输出的1ms宽数据比特流进行处理，得到数据比特边缘。之后将对应于同一个导航电文数据比特的20个1ms宽的数据比特合并起来，组成20ms宽的数据比特，完成导航电文数据比特的解调。

1.4 导航定位解算模块设计

导航定位解算模块包含帧同步、导航数据解码、卫星位置及伪距计算和导航参数解算4个子模块。经此4个子模块，导航定位解算模块最终得到天线的位置速度和时间(Position Velocity and Time, PVT)信息。

1.5 姿态测量模块设计

姿态测量模块基于实时动态(Real-Time Kinematic，RTK)载波相位差分技术，通过对安装在载体上的3个天线测得的载波相位进行差分处理，实时确定载体的姿态。

姿态测量模块采用基线长度约束的LAMBDA算法，即CLAMBDA算法。CLAMBDA算法通过基线约束，可以在观测历元较少的情况下得到精确的双差模糊度整数解[8]，进而得到精确的基线矢量和姿态角。

2 基于嵌入式GPU的GNSS信号处理加速技术

2.1 数据读取模块加速技术

前文中提到，数据读取模块的功能是将采集设备输出的采样数据读取到缓存中，然后根据存储格式将其拆分和重组为接收机各个系统需要的中频数据。现以量化位数为1bit，采集设备同时采集3个天线的卫星信号为例，对采样数据的存储格式以及数据拆分和重组的过程进行说明。

如图4所示，来自3个天线的采样数据以比特相邻的方式存储。由于接收机针对3个天线的3个系统是相互独立的，数据读取模块需对采样数据进行拆分，并重组成各个系统对应的中频数据。可以看出，其组成中频数据的1个字节，需要对采样数据的4个字节进行拆分。

图4 采样数据的存储格式图Fig.4 Storage format of the sampled data

以62MHz采样率和1bit量化位数为例，数据读取模块每次对40ms的数据进行操作，进行34410000次位运算、7440000次乘法运算和7440000次加法运算。如此繁重的运算量直接影响着接收机的实时处理性能，因此需针对数据读取模块进行加速。

加速后数据读取模块的实现流程如图5所示。

图5 数据读取模块的实现流程图Fig.5 Implementation flow of data reading module

一方面，中频数据的每个字节的重组过程是相互独立的，具有很好的数据并行性，可以由GPU上的内核函数并行的执行。内核函数开启的线程格中包含大量线程，通过这些线程能够实现多字节并行，加速中频数据的重组过程。

另一方面，由于中频数据的重组过程在GPU上的内核函数内执行，执行完毕后还需要进行由GPU到CPU的内存复制操作，将GPU上的中频数据复制到CPU上，方便后续对数据的处理。而支持设备重叠过程的GPU能够在执行一个内核函数的同时，在设备和主机之间执行复制操作，这种计算与数据传输的重叠是通过CUDA流来实现的。CUDA流表示一个GPU操作队列，该队列中的操作将按指定的顺序执行，可以将每个流视为GPU上的一个任务，并且这些任务可以并行执行。如图5所示，通过CUDA流的使用，减少了2组中频数据复制的时间，进一步加速了数据读取模块。

2.2 捕获模块加速技术

捕获模块使用并行码相位捕获算法，其算法流程如图2所示。通过观察可以发现，算法中混频、乘法、傅里叶变换、傅里叶逆变换等操作都以采样点为基本单位，且对每个采样点的操作具有明显的并行性，因此对捕获模块的加速采用多采样点并行的方式。

针对捕获中的混频和乘法运算，使用内核函数开启的并行线程加速，一个线程处理一个或多个采样点。对捕获中的傅里叶变换和傅里叶逆变换运算，使用CUFFT库中提供的专门用于CUDA傅里叶变换的函数，CUFFT 库在计算性能上对连续的快速傅里叶变换(Fast Fourier Transformation, FFT)处理进行了批处理优化，通过批处理可以掩盖GPU 访问显存的时间，从而降低整体运算时间[9]。

针对捕获中的求和运算和峰值查找运算，使用基于共享内存的并行规约算法进行加速。以8个元素为例，并行规约的流程如图6所示，规约的每个步骤都使得数据量变为一半，即使得时间复杂度由O(N)降为O(logN)。并行规约算法需要线程间的通信和协作，相较于全局内存，共享内存的性能更加优越。一方面，SM对全局内存的访问需要通过L2 cache，而自动管理的cache具有不确定性；另一方面，作为用户手工管理的缓存，共享内存在性能上堪比CPU的L1 cache，相较于全局内存具有更高的带宽和更低的延迟。

图6 并行规约算法示意图Fig.6 Schematic diagram of parallel reduction algorithm

2.3 跟踪模块加速技术

三天线软件接收机需要对多个信号通道的卫星进行跟踪，其数量可达几十个甚至上百个。因此跟踪模块的加速使用多卫星并行的方式，其中计算任务最繁重的相关器跟踪子模块以采样点数为单位对中频数据进行相关操作，采用多卫星、多采样点并行的方式。

前文中提到，相关器子模块采用比特并行快速算法进行载波剥离和相关解扩。比特并行快速算法将32个采样点进行的相关运算转化为整型数据的按位异或操作，每32个采样点的相关结果的累加通过统计整型变量的比特1的数目实现[10]。线程间相关结果的累加，与捕获过程的求和运算和峰值查找运算类似，使用基于共享内存的并行规约算法。

另外，由于跟踪模块需要在CPU和GPU之间传递的数据较多且结构复杂，对除中频数据以外的其他数据的传输，使用统一寻址[11]，无需显式拷贝，而是使用函数cudaMallocManaged()开辟一块存储空间，主机代码和设备代码均能对其访问。这种方式避免了人为的数据拷贝，能够降低代码复杂度。

3 样机研制与实验验证

3.1 样机研制

样机使用的硬件平台是NVIDIA推出的Jetson TX2，它是一种超小型模块化的超级计算机，尺寸仅有信用卡大小，非常适合于 GPU计算的嵌入式系统。Jetson TX2平台的操作系统为64bit的Ubuntu 16.04，其性能如表1所示。

表1 Jetson TX2性能概况

瑞泰新时代(北京)科技有限公司针对Jetson TX2开发了配套的超小型开发系统，型号为RTS-ASG003。样机如图7所示。

图7 样机图Fig.7 Prototype diagram

3.2 加速性能测试

在62MHz采样率、1bit量化条件下，采集BDS B3I的三天线实际卫星信号，长度为200s。以40ms的数据块为单位，统计数据读取模块的处理时间，测试结果对5000个样本值取平均值；统计三系统在第1次调用捕获模块时，对BDS的37颗卫星遍历搜索的时间，一颗卫星的捕获时间对3×37个样本值取平均值；以40ms的数据块为单位，统计后100s接收机稳定跟踪24颗卫星的情况下跟踪模块的处理时间，测试结果对2500个样本值取平均值。

根据上述测试条件，在样机上对加速前后的数据读取模块、捕获模块和跟踪模块的性能进行对比测试，测试结果如表2所示。

表2 加速性能对比

表2显示，在62MHz采样率、1bit量化条件下，数据读取模块达到了3.43倍的加速比；捕获模块对单个卫星通道的捕获达到了16.83倍的加速比；跟踪模块对24个卫星通道的跟踪达到了11.28倍的加速比。

3.3 样机性能测试

3.3.1 定位、测速和测姿性能测试

分别在软件接收机的定位和测速结果稳定输出后，每间隔1s采集其输出的位置测量值和速度测量值，共采集300s。在软件接收机的测姿结果稳定输出后，每间隔1s采集其输出的姿态测量值，共采集100s。测试信号为BDS B3I信号，测试采样频率为62MHz。

图8所示为位置测量值与其均值的差。

图8 位置误差Fig.8 Position errors

定位结果的标准差如表3所示。

表3 定位结果的标准差

图9所示为速度测量值与其均值的差。

图9 速度误差Fig.9 Velocity errors

测速结果的标准差如表4所示。

表4 测速结果的标准差

图10所示为姿态测量值与其均值的差。

图10 姿态误差Fig.10 Attitude errors

测姿结果的标准差如表5所示。

表5 测姿结果的标准差

3.3.2 跟踪通道数测试

考虑到实际可见的卫星数量有限，无法测出接收机的峰值性能。本文利用接收机对100s的BDS B3I卫星的62MHz采样信号进行处理，分别统计其在不同跟踪通道数量下的程序总运行时间，由此估算不同跟踪通道数量所需的GPU处理能力。图11所示为3、6、9、12、15、18、21和24个跟踪通道的处理时间，以及对更多通道性能的拟合。

图11 运行时间图Fig.11 Program run time

由图11可知，接收机能够在不超过95%处理器负载的情况下，支持超过90个62MHz采样的GNSS卫星信号处理。而参考文献[6]中Park K W等研制的基于Jetson TK1的软件接收机，仅支持10个跟踪通道以10MHz采样率实时运行。因此，本文中研制的样机在更高采样率条件下，实现了90个跟踪通道的实时处理，接收机性能得到了显著提升。

4 结论

本文利用GPU强大的浮点计算性能和并行运算能力，对GNSS软件接收机中耗时长且并行性明显的数据读取模块、捕获模块和跟踪模块进行了加速，并对三系统在62MHz采样率条件下的加速性能进行了测试。测试结果如下：

1)对数据读取模块进行GPU加速，实现了多个采样数据的并行拆分与重组，并采用CUDA流实现了内核函数与复制操作的并行执行，达到了3.43倍的加速比。

2)对捕获模块进行GPU加速，采用多采样点并行、CUFFT库以及基于共享内存的并行规约算法，达到了16.83倍的加速比。

3)对跟踪模块中计算任务最繁重的相关器跟踪子模块进行GPU加速，采用多采样点、多卫星并行以及基于共享内存的并行规约算法，达到了11.28倍的加速比。

利用上述GPU平台研制的三天线GNSS定位和测姿接收机样机，定位内符合精度达到0.60m，测速内符合精度达到0.15m/s，测姿内符合精度达到0.16°，可以支持超过90个62MHz采样的GNSS卫星信号处理。