基于嵌入式GPU的三天線GNSS基帶信號處理加速技術

楊智博，金 天

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

0 引言

近年來，由于全球定位系統(tǒng)(Global Positioning System，GPS)、俄羅斯全球衛(wèi)星導航系統(tǒng)(Global Navigation Satellite System，GLONASS)的現(xiàn)代化升級和伽利略衛(wèi)星導航系統(tǒng)(Galileo Satellite Navigation System)、中國北斗衛(wèi)星導航系統(tǒng)(BeiDou Navigation Satellite System，BDS)的不斷建設和完善，全球衛(wèi)星導航系統(tǒng)(Global Navigation Satellite System，GNSS)正在發(fā)生著巨大的變化。未來四大衛(wèi)星導航(GPS、BDS、GLONASS、GALILEO)以及準天頂衛(wèi)星系統(tǒng)(Quasi-Zenith Satellite System，QZSS)、印度區(qū)域?qū)Ш叫l(wèi)星系統(tǒng)(Indian Regional Navigation Satellite System， IRNSS)和廣域增強系統(tǒng)(Wide Area Augmentation System，WAAS)等眾多增強型導航系統(tǒng)將播發(fā)多種頻率的民用信號[1]，在提高衛(wèi)星導航性能的同時，也對接收機的設計提出了新的挑戰(zhàn)。

傳統(tǒng)的以現(xiàn)場可編程門陣列(Field Programmable Gate Array，F(xiàn)PGA)或者專用集成電路(Application Specific Integrated Circuit，ASIC)為代表的硬件接收機，由于其開發(fā)周期長和不夠靈活的缺點，難以滿足用戶和運營商不斷變化的需求[2]。在這種背景下，軟件接收機受到越來越多來自業(yè)內(nèi)的關注。軟件接收機具有很高的靈活性，能夠以很少的額外組件實現(xiàn)低成本，并且可以根據(jù)用戶的需求便捷的更改和升級[3]。但傳統(tǒng)的軟件接收機存在計算量大、耗時長的捕獲和跟蹤基帶信號處理過程，使得軟件接收機往往跟蹤通道較少，且難以在嵌入式系統(tǒng)上運行，因此基帶信號處理過程的加速是軟件接收機研究的重點。

隨著圖形處理器(Graphics Processing Unit，GPU)技術的不斷發(fā)展，其應用領域由圖形領域擴展到高性能計算領域，使得利用GPU加速軟件接收機成為可能，相關研究也逐漸增多。Hobiger T等使用GPU處理GPS C/A碼信號的捕獲和跟蹤，基于Intel Core 2 Q9450 CPU 和NVIDIA GeForce GTX 280 GPU實現(xiàn)了8MHz采樣率、4bit量化且支持12個衛(wèi)星通道的實時GPS軟件接收機[4]。Huang B等使用Intel i7-3770K CPU和NVIDIA GeForce GTX 580 GPU實現(xiàn)了多系統(tǒng)實時軟件接收機，可跟蹤11種衛(wèi)星信號[3]。Xu L等基于NVIDIAGeForce GTX 750 Ti GPU實現(xiàn)了一款模塊化GPS軟件接收機NAVSDR，支持300余個跟蹤通道以9.75MHz采樣率和4bit量化實時運行[5]。Park K W等基于GPU和CPU異構的嵌入式平臺Jetson TK1實現(xiàn)了一款軟件接收機，支持10個跟蹤通道以10MHz采樣率實時運行[6]。

雖然在利用GPU加速軟件接收機方面的研究越來越多，但多數(shù)研究均針對PC，針對嵌入式平臺的相關研究較少。因此，在嵌入式平臺上實現(xiàn)支持高采樣率、多衛(wèi)星通道的實時GNSS軟件接收機仍亟待研究。

本文對基于嵌入式平臺的GNSS軟件接收機的實現(xiàn)進行了研究。為滿足軟件接收機的實時性、高采樣率以及多衛(wèi)星通道的要求，本文采用嵌入式GPU對耗時長且并行性明顯的數(shù)據(jù)讀取模塊、捕獲模塊和跟蹤模塊進行了加速，研制了三天線GNSS定位和測姿接收機的樣機，并對其進行了加速性能和樣機性能的測試。

1 三天線GNSS軟件接收機整體架構

1.1 總體結構設計

如圖1所示，三天線GNSS信號定位和測姿接收機包含如下4個模塊：射頻前端、基帶信號處理、導航定位解算和姿態(tài)測量。射頻前端將射頻模擬信號轉(zhuǎn)換成數(shù)字基帶信號，3個天線的信號通過采集設備轉(zhuǎn)換成數(shù)字基帶信號；基帶信號處理模塊包括捕獲、跟蹤和位同步模塊，從基帶信號中獲得導航電文數(shù)據(jù)比特；導航定位解算模塊從導航電文中提取出相關信息，進而解算出接收機的位置、速度和時間信息；姿態(tài)測量模塊根據(jù)載波相位觀測值求差，得到載體的姿態(tài)信息。

圖1 接收機框架圖Fig.1 Receiver frame

1.2 射頻前端設計

采集設備的主要任務是將3個天線接收到的射頻模擬信號離散成包含衛(wèi)星信號成分的、頻率較低的數(shù)字中頻信號，并且在此過程中進行必要的濾波和增益控制。

采集設備選用 MAX2771 芯片，能夠以較低的成本提供最高性能和集成度，可以覆蓋 E5/L5、 L2、E6、E1/L1 波段，支持 GPS、GLONASS、Galileo、QZSS、IRNSS 和 BDS。通過采集設備上的MicroUSB接口虛擬串口，可以配置衛(wèi)星信號類型、頻點、帶寬，滿足多個GNSS的衛(wèi)星信號在信號形式、頻點、帶寬等方面的不同要求。

3個天線的模擬信號經(jīng)采集設備轉(zhuǎn)換為數(shù)字中頻信號后，數(shù)據(jù)讀取模塊讀取采樣數(shù)據(jù)到緩存中，然后根據(jù)采樣數(shù)據(jù)的存儲格式拆分為接收機各個系統(tǒng)需要的中頻數(shù)據(jù)。

1.3 基帶信號處理模塊設計

基帶信號處理是軟件接收機中重要的一環(huán)，其中包括捕獲、跟蹤和位同步模塊。通過基帶信號處理模塊，接收機從數(shù)字中頻信號中獲得導航電文數(shù)據(jù)比特，為后續(xù)導航定位解算做準備。

1.3.1 捕獲模塊設計

捕獲是針對碼相位和多普勒頻移的二維搜索過程，即捕獲的目的是確定能觀測到的衛(wèi)星，并估計出粗略的偽碼相位和載波多普勒頻移。目前常用的捕獲算法有串行捕獲、并行頻率捕獲和并行碼相位捕獲。其中，并行碼相位捕獲算法將數(shù)字相關器的相關運算替換為快速傅立葉變換[7]，減少了運算時間，在實時軟件接收機中被廣泛使用。

并行碼相位捕獲算法的流程如圖2所示。數(shù)字中頻輸入信號與本地載波信號相乘，分別得到I、Q支路。I、Q支路經(jīng)傅里葉變換，得到頻域信號，與本地偽碼的頻域復共軛相乘，再進行傅里葉逆變換，取模，查找出最大相關峰值。該最大相關峰值對應的偽碼相位和載波多普勒頻移即為捕獲模塊的目標值。

圖2 并行碼相位算法圖Fig.2 Phase algorithm diagram of parallel code

軟件接收機的捕獲模塊使用并行碼相位算法。各系統(tǒng)第一次調(diào)用捕獲模塊時，對其對應的導航系統(tǒng)的全部衛(wèi)星遍歷捕獲，之后每1s捕獲1次，每次捕獲1顆未被成功跟蹤的衛(wèi)星。

1.3.2 跟蹤模塊設計

經(jīng)由捕獲模塊，可以得到可見衛(wèi)星信號的載波頻率(包括載波多普勒偏移)和偽碼相位的粗略估計值。但它們的精度無法達到解調(diào)導航信息的要求，同時由于在接收衛(wèi)星信號的過程中，如衛(wèi)星持續(xù)運動、接收機運動、時鐘晶振鐘漂和隨機抖動等多種因素都會引起信號偽碼相位和載波頻率的變化，因此接收機在捕獲到某顆衛(wèi)星的信號后，將會在跟蹤模塊對載波頻率和偽碼相位值進行進一步處理，得到它們的精確值。

如圖3所示，軟件接收機的跟蹤模塊分為4個子模塊：相關器子模塊、PLL子模塊、DLL子模塊和NCO子模塊。跟蹤模塊的更新周期設置為40ms，子模塊的更新周期設置為1ms。即每次調(diào)用跟蹤模塊時，子模塊循環(huán)運行40次。

圖3 跟蹤模塊架構圖Fig.3 Tracking module architecture

相關器子模塊根據(jù)載波頻率和載波相位得到本地載波同相、正交三角函數(shù)表；根據(jù)偽碼相位和碼間隔得到本地三路超前、即時、滯后偽碼碼表；采用比特并行快速算法進行載波剝離和相關解擴， 最終得到E路、P路、L路(超前、即時、滯后)的同相正交共6個相關值。PLL跟蹤子模塊采用對導航電文翻轉(zhuǎn)不敏感的科斯塔(Costas)鎖相環(huán)，根據(jù)接收信號載波相位和本地載波相位的差值對本地載波相位進行調(diào)整，最終實現(xiàn)對載波相位的跟蹤，達到接收信號載波和本地載波相位的動態(tài)同步。DLL跟蹤子模塊選用延遲鎖定環(huán)，根據(jù)接收信號偽碼和本地偽碼的相位差異對本地偽碼相位進行調(diào)整，最終實現(xiàn)對偽碼相位的跟蹤。在經(jīng)由以上模塊對基帶信號中的載波和偽碼進行剝離之后，NCO跟蹤子模塊存儲即時的同相相關值，組成1ms寬數(shù)據(jù)比特流，以供后續(xù)位同步模塊進行處理。

1.3.3 位同步模塊設計

接收機的位同步模塊采用直方圖算法，通過對跟蹤環(huán)路輸出的1ms寬數(shù)據(jù)比特流進行處理，得到數(shù)據(jù)比特邊緣。之后將對應于同一個導航電文數(shù)據(jù)比特的20個1ms寬的數(shù)據(jù)比特合并起來，組成20ms寬的數(shù)據(jù)比特，完成導航電文數(shù)據(jù)比特的解調(diào)。

1.4 導航定位解算模塊設計

導航定位解算模塊包含幀同步、導航數(shù)據(jù)解碼、衛(wèi)星位置及偽距計算和導航參數(shù)解算4個子模塊。經(jīng)此4個子模塊，導航定位解算模塊最終得到天線的位置速度和時間(Position Velocity and Time, PVT)信息。

1.5 姿態(tài)測量模塊設計

姿態(tài)測量模塊基于實時動態(tài)(Real-Time Kinematic，RTK)載波相位差分技術，通過對安裝在載體上的3個天線測得的載波相位進行差分處理，實時確定載體的姿態(tài)。

姿態(tài)測量模塊采用基線長度約束的LAMBDA算法，即CLAMBDA算法。CLAMBDA算法通過基線約束，可以在觀測歷元較少的情況下得到精確的雙差模糊度整數(shù)解[8]，進而得到精確的基線矢量和姿態(tài)角。

2 基于嵌入式GPU的GNSS信號處理加速技術

2.1 數(shù)據(jù)讀取模塊加速技術

前文中提到，數(shù)據(jù)讀取模塊的功能是將采集設備輸出的采樣數(shù)據(jù)讀取到緩存中，然后根據(jù)存儲格式將其拆分和重組為接收機各個系統(tǒng)需要的中頻數(shù)據(jù)。現(xiàn)以量化位數(shù)為1bit，采集設備同時采集3個天線的衛(wèi)星信號為例，對采樣數(shù)據(jù)的存儲格式以及數(shù)據(jù)拆分和重組的過程進行說明。

如圖4所示，來自3個天線的采樣數(shù)據(jù)以比特相鄰的方式存儲。由于接收機針對3個天線的3個系統(tǒng)是相互獨立的，數(shù)據(jù)讀取模塊需對采樣數(shù)據(jù)進行拆分，并重組成各個系統(tǒng)對應的中頻數(shù)據(jù)?？梢钥闯?，其組成中頻數(shù)據(jù)的1個字節(jié)，需要對采樣數(shù)據(jù)的4個字節(jié)進行拆分。

圖4 采樣數(shù)據(jù)的存儲格式圖Fig.4 Storage format of the sampled data

以62MHz采樣率和1bit量化位數(shù)為例，數(shù)據(jù)讀取模塊每次對40ms的數(shù)據(jù)進行操作，進行34410000次位運算、7440000次乘法運算和7440000次加法運算。如此繁重的運算量直接影響著接收機的實時處理性能，因此需針對數(shù)據(jù)讀取模塊進行加速。

加速后數(shù)據(jù)讀取模塊的實現(xiàn)流程如圖5所示。

圖5 數(shù)據(jù)讀取模塊的實現(xiàn)流程圖Fig.5 Implementation flow of data reading module

一方面，中頻數(shù)據(jù)的每個字節(jié)的重組過程是相互獨立的，具有很好的數(shù)據(jù)并行性，可以由GPU上的內(nèi)核函數(shù)并行的執(zhí)行。內(nèi)核函數(shù)開啟的線程格中包含大量線程，通過這些線程能夠?qū)崿F(xiàn)多字節(jié)并行，加速中頻數(shù)據(jù)的重組過程。

另一方面，由于中頻數(shù)據(jù)的重組過程在GPU上的內(nèi)核函數(shù)內(nèi)執(zhí)行，執(zhí)行完畢后還需要進行由GPU到CPU的內(nèi)存復制操作，將GPU上的中頻數(shù)據(jù)復制到CPU上，方便后續(xù)對數(shù)據(jù)的處理。而支持設備重疊過程的GPU能夠在執(zhí)行一個內(nèi)核函數(shù)的同時，在設備和主機之間執(zhí)行復制操作，這種計算與數(shù)據(jù)傳輸?shù)闹丿B是通過CUDA流來實現(xiàn)的。CUDA流表示一個GPU操作隊列，該隊列中的操作將按指定的順序執(zhí)行，可以將每個流視為GPU上的一個任務，并且這些任務可以并行執(zhí)行。如圖5所示，通過CUDA流的使用，減少了2組中頻數(shù)據(jù)復制的時間，進一步加速了數(shù)據(jù)讀取模塊。

2.2 捕獲模塊加速技術

捕獲模塊使用并行碼相位捕獲算法，其算法流程如圖2所示。通過觀察可以發(fā)現(xiàn)，算法中混頻、乘法、傅里葉變換、傅里葉逆變換等操作都以采樣點為基本單位，且對每個采樣點的操作具有明顯的并行性，因此對捕獲模塊的加速采用多采樣點并行的方式。

針對捕獲中的混頻和乘法運算，使用內(nèi)核函數(shù)開啟的并行線程加速，一個線程處理一個或多個采樣點。對捕獲中的傅里葉變換和傅里葉逆變換運算，使用CUFFT庫中提供的專門用于CUDA傅里葉變換的函數(shù)，CUFFT 庫在計算性能上對連續(xù)的快速傅里葉變換(Fast Fourier Transformation, FFT)處理進行了批處理優(yōu)化，通過批處理可以掩蓋GPU 訪問顯存的時間，從而降低整體運算時間[9]。

針對捕獲中的求和運算和峰值查找運算，使用基于共享內(nèi)存的并行規(guī)約算法進行加速。以8個元素為例，并行規(guī)約的流程如圖6所示，規(guī)約的每個步驟都使得數(shù)據(jù)量變?yōu)橐话耄词沟脮r間復雜度由O(N)降為O(logN)。并行規(guī)約算法需要線程間的通信和協(xié)作，相較于全局內(nèi)存，共享內(nèi)存的性能更加優(yōu)越。一方面，SM對全局內(nèi)存的訪問需要通過L2 cache，而自動管理的cache具有不確定性；另一方面，作為用戶手工管理的緩存，共享內(nèi)存在性能上堪比CPU的L1 cache，相較于全局內(nèi)存具有更高的帶寬和更低的延遲。

圖6 并行規(guī)約算法示意圖Fig.6 Schematic diagram of parallel reduction algorithm

2.3 跟蹤模塊加速技術

三天線軟件接收機需要對多個信號通道的衛(wèi)星進行跟蹤，其數(shù)量可達幾十個甚至上百個。因此跟蹤模塊的加速使用多衛(wèi)星并行的方式，其中計算任務最繁重的相關器跟蹤子模塊以采樣點數(shù)為單位對中頻數(shù)據(jù)進行相關操作，采用多衛(wèi)星、多采樣點并行的方式。

前文中提到，相關器子模塊采用比特并行快速算法進行載波剝離和相關解擴。比特并行快速算法將32個采樣點進行的相關運算轉(zhuǎn)化為整型數(shù)據(jù)的按位異或操作，每32個采樣點的相關結果的累加通過統(tǒng)計整型變量的比特1的數(shù)目實現(xiàn)[10]。線程間相關結果的累加，與捕獲過程的求和運算和峰值查找運算類似，使用基于共享內(nèi)存的并行規(guī)約算法。

另外，由于跟蹤模塊需要在CPU和GPU之間傳遞的數(shù)據(jù)較多且結構復雜，對除中頻數(shù)據(jù)以外的其他數(shù)據(jù)的傳輸，使用統(tǒng)一尋址[11]，無需顯式拷貝，而是使用函數(shù)cudaMallocManaged()開辟一塊存儲空間，主機代碼和設備代碼均能對其訪問。這種方式避免了人為的數(shù)據(jù)拷貝，能夠降低代碼復雜度。

3 樣機研制與實驗驗證

3.1 樣機研制

樣機使用的硬件平臺是NVIDIA推出的Jetson TX2，它是一種超小型模塊化的超級計算機，尺寸僅有信用卡大小，非常適合于 GPU計算的嵌入式系統(tǒng)。Jetson TX2平臺的操作系統(tǒng)為64bit的Ubuntu 16.04，其性能如表1所示。

表1 Jetson TX2性能概況

瑞泰新時代(北京)科技有限公司針對Jetson TX2開發(fā)了配套的超小型開發(fā)系統(tǒng)，型號為RTS-ASG003。樣機如圖7所示。

圖7 樣機圖Fig.7 Prototype diagram

3.2 加速性能測試

在62MHz采樣率、1bit量化條件下，采集BDS B3I的三天線實際衛(wèi)星信號，長度為200s。以40ms的數(shù)據(jù)塊為單位，統(tǒng)計數(shù)據(jù)讀取模塊的處理時間，測試結果對5000個樣本值取平均值；統(tǒng)計三系統(tǒng)在第1次調(diào)用捕獲模塊時，對BDS的37顆衛(wèi)星遍歷搜索的時間，一顆衛(wèi)星的捕獲時間對3×37個樣本值取平均值；以40ms的數(shù)據(jù)塊為單位，統(tǒng)計后100s接收機穩(wěn)定跟蹤24顆衛(wèi)星的情況下跟蹤模塊的處理時間，測試結果對2500個樣本值取平均值。

根據(jù)上述測試條件，在樣機上對加速前后的數(shù)據(jù)讀取模塊、捕獲模塊和跟蹤模塊的性能進行對比測試，測試結果如表2所示。

表2 加速性能對比

表2顯示，在62MHz采樣率、1bit量化條件下，數(shù)據(jù)讀取模塊達到了3.43倍的加速比；捕獲模塊對單個衛(wèi)星通道的捕獲達到了16.83倍的加速比；跟蹤模塊對24個衛(wèi)星通道的跟蹤達到了11.28倍的加速比。

3.3 樣機性能測試

3.3.1 定位、測速和測姿性能測試

分別在軟件接收機的定位和測速結果穩(wěn)定輸出后，每間隔1s采集其輸出的位置測量值和速度測量值，共采集300s。在軟件接收機的測姿結果穩(wěn)定輸出后，每間隔1s采集其輸出的姿態(tài)測量值，共采集100s。測試信號為BDS B3I信號，測試采樣頻率為62MHz。

圖8所示為位置測量值與其均值的差。

圖8 位置誤差Fig.8 Position errors

定位結果的標準差如表3所示。

表3 定位結果的標準差

圖9所示為速度測量值與其均值的差。

圖9 速度誤差Fig.9 Velocity errors

測速結果的標準差如表4所示。

表4 測速結果的標準差

圖10所示為姿態(tài)測量值與其均值的差。

圖10 姿態(tài)誤差Fig.10 Attitude errors

測姿結果的標準差如表5所示。

表5 測姿結果的標準差

3.3.2 跟蹤通道數(shù)測試

考慮到實際可見的衛(wèi)星數(shù)量有限，無法測出接收機的峰值性能。本文利用接收機對100s的BDS B3I衛(wèi)星的62MHz采樣信號進行處理，分別統(tǒng)計其在不同跟蹤通道數(shù)量下的程序總運行時間，由此估算不同跟蹤通道數(shù)量所需的GPU處理能力。圖11所示為3、6、9、12、15、18、21和24個跟蹤通道的處理時間，以及對更多通道性能的擬合。

圖11 運行時間圖Fig.11 Program run time

由圖11可知，接收機能夠在不超過95%處理器負載的情況下，支持超過90個62MHz采樣的GNSS衛(wèi)星信號處理。而參考文獻[6]中Park K W等研制的基于Jetson TK1的軟件接收機，僅支持10個跟蹤通道以10MHz采樣率實時運行。因此，本文中研制的樣機在更高采樣率條件下，實現(xiàn)了90個跟蹤通道的實時處理，接收機性能得到了顯著提升。

4 結論

本文利用GPU強大的浮點計算性能和并行運算能力，對GNSS軟件接收機中耗時長且并行性明顯的數(shù)據(jù)讀取模塊、捕獲模塊和跟蹤模塊進行了加速，并對三系統(tǒng)在62MHz采樣率條件下的加速性能進行了測試。測試結果如下：

1)對數(shù)據(jù)讀取模塊進行GPU加速，實現(xiàn)了多個采樣數(shù)據(jù)的并行拆分與重組，并采用CUDA流實現(xiàn)了內(nèi)核函數(shù)與復制操作的并行執(zhí)行，達到了3.43倍的加速比。

2)對捕獲模塊進行GPU加速，采用多采樣點并行、CUFFT庫以及基于共享內(nèi)存的并行規(guī)約算法，達到了16.83倍的加速比。

3)對跟蹤模塊中計算任務最繁重的相關器跟蹤子模塊進行GPU加速，采用多采樣點、多衛(wèi)星并行以及基于共享內(nèi)存的并行規(guī)約算法，達到了11.28倍的加速比。

利用上述GPU平臺研制的三天線GNSS定位和測姿接收機樣機，定位內(nèi)符合精度達到0.60m，測速內(nèi)符合精度達到0.15m/s，測姿內(nèi)符合精度達到0.16°，可以支持超過90個62MHz采樣的GNSS衛(wèi)星信號處理。