基于嵌入式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衛星信號處理。