張?chǎng)斡睿瑮罴咨?,?聰,陳志敏,智 佳,陳 托
(1.中國(guó)科學(xué)院國(guó)家空間科學(xué)中心,北京 100190;2.中國(guó)科學(xué)院大學(xué) 計(jì)算機(jī)科學(xué)與技術(shù)學(xué)院,北京 100049)
數(shù)傳預(yù)處理是地面應(yīng)用系統(tǒng)生產(chǎn)數(shù)據(jù)產(chǎn)品、獲取探測(cè)成果的前提和基礎(chǔ)[1],其處理內(nèi)容包括傳輸幀同步、幀校驗(yàn)、分虛擬信道、源包同步、分應(yīng)用過程、工程遙測(cè)參數(shù)提取與物理量轉(zhuǎn)換、科學(xué)數(shù)據(jù)解壓縮等。高分辨率、高靈敏度、多頻段覆蓋的載荷探測(cè)趨勢(shì)下,Gbps 級(jí)數(shù)傳預(yù)處理需求日益迫切。
我國(guó)航天器數(shù)傳大多采用的空間數(shù)據(jù)系統(tǒng)咨詢委員會(huì)(Consultative Committee for Space Data Systems,CCSDS)相關(guān)協(xié)議具有典型的分層特征[2],層間數(shù)據(jù)相互依賴,層內(nèi)不同虛擬信道、應(yīng)用過程數(shù)據(jù)相對(duì)獨(dú)立。因此,目前國(guó)內(nèi)外學(xué)者大多采用中央處理器(Central Processing Unit,CPU)單機(jī)、集群等同構(gòu)平臺(tái),以虛擬信道、應(yīng)用過程為任務(wù)劃分依據(jù),并行提高處理性能。包括基于消息傳遞接口(Message-Passing Interface,MPI)[3]、基于MPI+OpenMP[4]在分布式集群上實(shí)現(xiàn)衛(wèi)星遙感數(shù)據(jù)處理;基于Storm 流式計(jì)算框架實(shí)現(xiàn)遙感衛(wèi)星快視處理[5]、空間科學(xué)衛(wèi)星數(shù)據(jù)處理[6-7]。CPU 適用于處理復(fù)雜控制邏輯、較小數(shù)據(jù)規(guī)模問題[8],限制了上述方法部分處理步驟性能提升。圖形處理器(Graphics Processing Unit,GPU)以 其TFLOPS 量級(jí)的計(jì)算能力,在信號(hào)處理[9-10]、譯碼[11-12]、遙感圖像處理[13-18]、衛(wèi)星軌道計(jì)算[19-20]等并行程度較大的數(shù)據(jù)密集型領(lǐng)域應(yīng)用中取得了良好的加速效果。
本文面向高性能數(shù)傳預(yù)處理需求,在分析處理性能瓶頸的基礎(chǔ)上,針對(duì)預(yù)處理步驟層間依賴、各虛擬信道保序并行的特點(diǎn),提出一種層間流程CPU控制、層內(nèi)瓶頸步驟GPU 加速的協(xié)同處理新方法。對(duì)數(shù)傳數(shù)據(jù)預(yù)處理的CPU-GPU 協(xié)同任務(wù)劃分、GPU 線程分配進(jìn)行設(shè)計(jì),實(shí)現(xiàn)基于GPU 的幀校驗(yàn)、工程參數(shù)提取與物理量轉(zhuǎn)換的加速,為數(shù)傳預(yù)處理性能的提升提供一種新的解決途徑。
如圖1 所示,在地面接收分系統(tǒng)完成譯碼后,數(shù)傳預(yù)處理完成高級(jí)在軌系統(tǒng)(Advanced Orbiting Systems,AOS)幀處理、源包處理,將結(jié)果分發(fā)至后續(xù)環(huán)節(jié)。涉及到CCSDS 協(xié)議包括空間包協(xié)議[21]、AOS 空間數(shù)據(jù)鏈路協(xié)議[22]、遙測(cè)同步和信道編碼協(xié)議[23]等,其中AOS 傳輸幀、空間包結(jié)構(gòu)見表1和表2。
表1 AOS 傳輸幀格式Tab.1 AOS transfer frame format
表2 空間包結(jié)構(gòu)Tab.2 Structural components of the space packet
圖1 地面端數(shù)據(jù)處理的一般過程Fig.1 General process of ground data processing
利用文件回放模擬數(shù)據(jù)輸入,測(cè)試CPU 單線程模式下預(yù)處理各操作速率見表3。
表3 CPU 單線程熱點(diǎn)測(cè)試Tab.3 Hotpot test of the CPU single-threaded program
其中幀校驗(yàn)和參數(shù)提取與物理量轉(zhuǎn)換計(jì)算量大,是處理瓶頸且可并行性高,作為本文主要研究對(duì)象。
基于CCSDS 協(xié)議的分層模型和多路復(fù)用機(jī)制,預(yù)處理過程中數(shù)據(jù)結(jié)構(gòu)變化如圖2 所示。處理存在如下特點(diǎn):
1)分層模型使各層數(shù)據(jù)的處理可任務(wù)并行;
2)上層協(xié)議數(shù)據(jù)處理依賴下層輸出,不同層級(jí)數(shù)據(jù)間存在關(guān)聯(lián);
3)多路復(fù)用機(jī)制使同一物理信道中分屬不同虛擬信道數(shù)據(jù)幀間、同一虛擬信道分屬不同應(yīng)用過程的源包間無依賴關(guān)系,可完全并行;
4)存在保序約束,須保持各虛擬信道、各應(yīng)用過程內(nèi)數(shù)據(jù)的順序,發(fā)揮計(jì)數(shù)字段檢測(cè)數(shù)據(jù)丟失、重復(fù)、失序的作用,可保證源包及應(yīng)用過程數(shù)據(jù)正確拼接。
依據(jù)預(yù)處理步驟層間依賴、各虛擬信道保序并行的特點(diǎn),設(shè)計(jì)AOS 幀、源包處理步驟的CPU、GPU任務(wù)分劃,并基于GPU 統(tǒng)一計(jì)算設(shè)備架構(gòu)(Compute Unified Device Architecture,CUDA),實(shí)現(xiàn)幀校驗(yàn)、工程參數(shù)提取與物理量轉(zhuǎn)換步驟的加速。
GPU 含有上千個(gè)CUDA 核心,一次可支持幾萬個(gè)線程并發(fā)。GPU 的寄存器帶寬為TB·s-1量級(jí),遠(yuǎn)高于GB·s-1量級(jí)的GPU 顯存帶寬。CPU 與GPU間傳輸數(shù)據(jù)的高速串行計(jì)算機(jī)擴(kuò)展總線(Peripheral Component Interconnect Express,PCIe)帶寬遠(yuǎn)又小于GPU 顯存帶寬。因此使用GPU 獲得良好加速效果,CPU、GPU 任務(wù)分劃要遵循以下原則[24]:
1)增大核函數(shù)并行規(guī)模,避免核心空閑;
2)提高GPU 中核函數(shù)算數(shù)強(qiáng)度,降低GPU 內(nèi)存讀寫占比;
3)降低CPU-GPU 間數(shù)據(jù)傳輸量。
2.1.1 AOS 幀處理流程設(shè)計(jì)
分路分包前碼速率較高,而幀校驗(yàn)又為計(jì)算量較大的瓶頸步驟;各AOS 幀獨(dú)立校驗(yàn),幀間無依賴關(guān)系;高碼速率同時(shí)降低了等待數(shù)據(jù)積累到較大規(guī)模供GPU 處理的時(shí)延,基于上述考慮,如圖3 所示,設(shè)計(jì)GPU 加速幀校驗(yàn)與解析。使用CPU 對(duì)解擾后的碼流數(shù)據(jù)同步,得到AOS 幀寫入幀緩存,經(jīng)過PCIe 總線傳輸至GPU 全局內(nèi)存。GPU 處理完畢后不再向CPU 端傳回完整AOS 幀,而是傳回8 bit/幀的后續(xù)處理必須信息,數(shù)據(jù)傳輸量降低99.1%。傳回信息包括:幀格式符合性、VCID、虛擬信道幀計(jì)數(shù)、數(shù)據(jù)域有效數(shù)據(jù)結(jié)束地址。CPU 端異常幀報(bào)警器在檢測(cè)到幀格式不符合時(shí),重新解析該幀,生成異常幀報(bào)告;幀數(shù)據(jù)域提取處理器提取數(shù)據(jù),按VCID 寫入各源包碼流數(shù)據(jù)環(huán)形隊(duì)列存儲(chǔ)器,同步處理后得到源包(空間包)。
圖3 CPU-GPU 協(xié)同的AOS 幀處理流程Fig.3 Processing flow of AOS frame parsing based on the CPU-GPU collaboration
2.1.2 源包處理流程設(shè)計(jì)
源包按裝載的數(shù)據(jù)類型,可分為業(yè)務(wù)數(shù)據(jù)源包、遙測(cè)參數(shù)源包、工程參數(shù)源包。空間科學(xué)衛(wèi)星的科學(xué)數(shù)據(jù)、遙感衛(wèi)星的圖像數(shù)據(jù)等業(yè)務(wù)數(shù)據(jù),數(shù)據(jù)量大、處理方法復(fù)雜、不同業(yè)務(wù)處理差異大。遙測(cè)數(shù)據(jù)是與航天器工作狀態(tài)有關(guān)的測(cè)量數(shù)據(jù)(如電壓、電流、溫度、壓力等),反映航天器的工作過程、工作結(jié)果和健康狀態(tài)[25],實(shí)時(shí)性要求高,數(shù)據(jù)量較小。工程參數(shù)主要是有效載荷的工作狀態(tài)數(shù)據(jù),相比遙測(cè)參數(shù),種類更多、數(shù)據(jù)量更大。
根據(jù)不同類型的源包特點(diǎn),設(shè)計(jì)源包處理任務(wù)流程如圖4 所示。CPU 解析源包后生成異常源包報(bào)告,根據(jù)APID 將格式符合的源包寫入相應(yīng)的環(huán)形隊(duì)列存儲(chǔ)器。業(yè)務(wù)數(shù)據(jù)在預(yù)處理階段只做拼接處理,由CPU 完成。參數(shù)配置信息處理器讀取含有參數(shù)位置、源數(shù)據(jù)長(zhǎng)度、物理量轉(zhuǎn)換方法的配置文件,計(jì)算掩碼,在程序啟動(dòng)時(shí)向GPU 全局內(nèi)存?zhèn)鬏斠淮巍_b測(cè)參數(shù)數(shù)據(jù)量較小,為降低時(shí)延,由CPU完成提取與物理量轉(zhuǎn)換。工程參數(shù)由GPU 完成提取與物理量轉(zhuǎn)換。
圖4 CPU-GPU 協(xié)同的源包處理流程Fig.4 Processing flow of space packet parsing based on the CPU-GPU collaboration
2.2.1 幀校驗(yàn)的CUDA 加速
AOS 協(xié)議的差錯(cuò)控制使用生成多項(xiàng)式為C(x)=x16+x12+x5+1 的CRC 校 驗(yàn),CRC 校 驗(yàn)的基本思想是接收端的數(shù)據(jù)無誤碼則能被給定的生成多項(xiàng)式C(x)整除。CRC 校驗(yàn)以幀為基本單位,各幀之間不影響,設(shè)計(jì)每個(gè)GPU 線程校驗(yàn)一幀。從主機(jī)端緩存每次讀取N個(gè)長(zhǎng)度均為L(zhǎng)字節(jié)的AOS 幀至GPU 全局內(nèi)存中線性存儲(chǔ)。將網(wǎng)格內(nèi)的線程塊和線程塊內(nèi)的線程均組織成一維,線程塊數(shù)量為Nb;每個(gè)網(wǎng)格中線程數(shù)量Ng為
每個(gè)線程在核函數(shù)中的唯一身份標(biāo)識(shí)idx為
式 中:Tx∈[0,Db-1]為各線程在其線程塊中的索引。
第idx個(gè)線程處理第i個(gè)傳輸幀,該幀的數(shù)據(jù)在全局內(nèi)存地址為:[idx?L,(idx+1)?L-1],每個(gè)線程完成校驗(yàn)后,解析幀主導(dǎo)頭各字段,提取所需信息寫入幀解析結(jié)果結(jié)構(gòu)體。所有線程計(jì)算完畢,使用cudaMemcpy()函數(shù)將幀解析結(jié)果結(jié)構(gòu)體數(shù)組傳輸至CPU 端內(nèi)存。
2.2.2 工程參數(shù)提取與物理量轉(zhuǎn)換的CUDA 加速
參數(shù)在源包中以二進(jìn)制信號(hào)形式存在,需要將參數(shù)提取、編碼、轉(zhuǎn)換后,才能直觀地用于載荷狀態(tài)監(jiān)視、參數(shù)判讀等任務(wù),可分為以下3 個(gè)步驟:
步驟一參數(shù)提取。提取源包數(shù)據(jù)域碼流中各參數(shù)的二進(jìn)制信號(hào)源碼。
參數(shù)位置由參數(shù)開始的字節(jié)地址Ab、開始字節(jié)的位地址Ab、比特長(zhǎng)度Lb唯一確定。某參數(shù)在源包中所占字節(jié)長(zhǎng)度Lb為
將幀中第Ab至Ab+Lb-1 字節(jié)數(shù)據(jù)與位掩碼做按位與運(yùn)算并位移后得到參數(shù)源碼。
步驟二參數(shù)編碼。將各參數(shù)的二進(jìn)制信號(hào)源碼編碼為實(shí)際數(shù)字或字符串。
步驟三物理量轉(zhuǎn)換。參數(shù)經(jīng)某種方法的運(yùn)算后,轉(zhuǎn)換為直觀表達(dá)真實(shí)物理含義的值。常用的物理量轉(zhuǎn)換方法有源碼顯示、進(jìn)制轉(zhuǎn)換、含義映射、多項(xiàng)式變換、熱敏電阻分壓值轉(zhuǎn)換為溫度等。
從CPU 端緩存每次讀取特定種類長(zhǎng)度為L(zhǎng)字節(jié),含M個(gè)參數(shù)的工程參數(shù)源包N個(gè),每個(gè)參數(shù)的處理結(jié)果使用R字節(jié)保存。將網(wǎng)格內(nèi)的線程塊和線程塊內(nèi)的線程均組織成一維,線程塊數(shù)量為Nb。設(shè)計(jì)了每個(gè)線程處理一個(gè)參數(shù)、每個(gè)線程處理一個(gè)源包2 種線程分配方案,通過實(shí)驗(yàn)分析2 種方案的特性。
方案一每個(gè)線程處理一個(gè)參數(shù)。
每個(gè)網(wǎng)格中x方向線程數(shù)量Ng為
第idx個(gè)線程處理第個(gè)源包的第(idx%M)個(gè)參數(shù),將結(jié)果寫入GPU 全局內(nèi)存,內(nèi)存地址為[idx*R,(idx+1)*R-1]。
方案二每個(gè)線程處理一個(gè)源包。
每個(gè)網(wǎng)格中x方向線程數(shù)量Ng計(jì)算方法與公式(1)相同,第idx個(gè)線程處理第idx個(gè)源包的所有參數(shù),將結(jié)果寫入GPU 全局內(nèi)存,內(nèi)存地址為[idx*M*R,(idx+1)*M*R-1]。
2.2.3 基于CUDA 流并發(fā)執(zhí)行核函數(shù)和數(shù)據(jù)傳輸
CUDA 流是從主機(jī)發(fā)出在GPU 中執(zhí)行的一系列異步CUDA 操作,不同CUDA 流中的操作可能并發(fā)或交錯(cuò)執(zhí)行[14]。由于數(shù)據(jù)傳輸、計(jì)算分別通過GPU 的讀寫單元和流處理器進(jìn)行,如圖5 所示,主機(jī)向一個(gè)流發(fā)出數(shù)據(jù)傳輸命令后可立即獲得控制權(quán),到另一個(gè)流調(diào)用核函數(shù),從而重疊數(shù)據(jù)傳輸與核函數(shù)執(zhí)行[15]。在CUDA 操作完全并行執(zhí)行的理想情況下,數(shù)據(jù)傳輸時(shí)間可完全隱藏,隨著流數(shù)量增加,加速比趨近3。
圖5 CUDA 流并發(fā)示意Fig.5 CUDA stream concurrency
CUDA 流實(shí)現(xiàn)異步數(shù)據(jù)傳輸,需要在程序運(yùn)行期間將主機(jī)端數(shù)據(jù)緩存定義為不可分頁內(nèi)存,保持內(nèi)存物理地址不變。設(shè)創(chuàng)建NS個(gè)CUDA 流,將每次從緩存中讀取的N個(gè)AOS 幀或源包劃分NS個(gè)子集,平均分配到各個(gè)流中,使用cudaMemcpyAsync()函數(shù)異步傳輸數(shù)據(jù),使用循環(huán)調(diào)度各個(gè)流的數(shù)據(jù)傳輸和計(jì)算。
CPU 使用AMD Ryzen 5 2600X,基頻3.6 GHz。GPU 使用NVIDIA GeForce RTX 2080ti,顯卡為 圖靈架構(gòu),有68 個(gè)流多處理器,4 352 個(gè)CUDA 核心,11 GB GDDR6 顯存,顯存帶寬616 GB·s-1,單精度浮點(diǎn)數(shù)運(yùn)算峰值13TFLOPS,CUDA 版本10.1。CPU 與GPU 之間使用PCIe3.0 x16 總線連接,理 論帶寬15.754 GB·s-1。內(nèi)存使用48GB DDR4 雙通道內(nèi)存。硬盤使用512 GB NVMe 固態(tài)硬盤。
由于不同型號(hào)CPU 物理核心數(shù)的差距,實(shí)踐中評(píng)價(jià)加速效果常使用CPU 單核與GPU 比較[9]。
傳輸幀幀長(zhǎng)898 字節(jié),前4 字節(jié)為同步碼,幀差錯(cuò)控制域2 字節(jié)。Nb設(shè)為256。如圖6 所示,使用單個(gè)CUDA 流,一次處理215個(gè)幀時(shí)達(dá)到最高處理速 率8.053 6 GB·s-1,加速比38.003;使 用16 個(gè)CUDA 流,一次處理219個(gè)幀時(shí)達(dá)到最高處理速率11.449 6 GB·s-1,加速比53.866。CPU 單線程方式,最高處理速率0.214 7 GB·s-1,串行計(jì)算性能基本與幀數(shù)量無關(guān)。
圖6 速率隨AOS 幀個(gè)數(shù)的變化-CRCFig.6 Variation of the processing speed with the AOS frame number-CRC
如圖7 所示,一次處理的幀數(shù)量少時(shí),GPU 沒有滿負(fù)荷工作,有空閑的計(jì)算資源,增加一次處理的幀數(shù)量,GPU 核函數(shù)耗時(shí)基本不變;當(dāng)一次處理的幀個(gè)數(shù)較大時(shí),GPU 滿負(fù)荷工作,GPU 核函數(shù)的耗時(shí)正比于一次處理的幀個(gè)數(shù)??筛鶕?jù)不同GPU性能,配置單次計(jì)算規(guī)模,達(dá)到最佳性能。
圖7 GPU 核函數(shù)計(jì)算與數(shù)據(jù)傳輸耗時(shí)-CRCFig.7 Time of the kernel function and data transmission by GPU-CRC
工程參數(shù)源包包長(zhǎng)390 字節(jié),含467 個(gè)參數(shù),每個(gè)參數(shù)轉(zhuǎn)換結(jié)果8 字節(jié)保存。Nb設(shè)為256。如圖8所 示,CPU 最高速 率0.004 4 GB·s-1。使用單 個(gè)CUDA 流,最高速率0.506 GB·s-1,加速比115.98;使用16 個(gè)CUDA 流,最高速率可達(dá)0.902 4 GB·s-1,加速比211.021。
圖8 速率隨源包個(gè)數(shù)的變化Fig.8 Variation of the processing speed with the packet number
如圖9 和圖10 所示,一次處理的源包個(gè)數(shù)小于215時(shí),每個(gè)線程處理一個(gè)源包的方案并行規(guī)模較小,速率比每個(gè)線程處理一個(gè)參數(shù)的方案低;源包個(gè)數(shù)大于215時(shí),兩方案均能使GPU 滿負(fù)荷工作,每個(gè)線程處理一個(gè)源包的算數(shù)強(qiáng)度更大,速率更高。在實(shí)際應(yīng)用中,可根據(jù)數(shù)據(jù)的時(shí)延性要求,配置一次處理的源包個(gè)數(shù),進(jìn)而確定線程分配方案。
圖10 GPU 核函數(shù)計(jì)算與數(shù)據(jù)傳輸耗時(shí)-1 線程處理1 參數(shù)Fig.10 Time of the kernel function and data transmission by GPU-1 parameter for 1 thread
數(shù)傳數(shù)據(jù)預(yù)處理是衛(wèi)星數(shù)據(jù)處理的前置必須環(huán)節(jié),本文提出的CPU-GPU 協(xié)同的衛(wèi)星數(shù)傳數(shù)據(jù)預(yù)處理方法,利用GPU 并行計(jì)算加速處理過程的幀校驗(yàn)、參數(shù)提取與物理量轉(zhuǎn)換等耗時(shí)瓶頸步驟,實(shí)驗(yàn)結(jié)果表明,該方法可滿足Gbps 的衛(wèi)星數(shù)傳數(shù)據(jù)預(yù)處理需求。后續(xù)可開展的工作如下。
1)增加處理深度,引入基于GPU 的數(shù)據(jù)解壓縮、科學(xué)數(shù)據(jù)快視、數(shù)據(jù)判讀,進(jìn)一步加速數(shù)據(jù)處理過程。
2)研究適用于數(shù)傳數(shù)據(jù)處理的多CPU、多GPU 場(chǎng)景中有硬件和任務(wù)優(yōu)先級(jí)約束下的多目標(biāo)優(yōu)化調(diào)度算法。