• <tr id="yyy80"></tr>
  • <sup id="yyy80"></sup>
  • <tfoot id="yyy80"><noscript id="yyy80"></noscript></tfoot>
  • 99热精品在线国产_美女午夜性视频免费_国产精品国产高清国产av_av欧美777_自拍偷自拍亚洲精品老妇_亚洲熟女精品中文字幕_www日本黄色视频网_国产精品野战在线观看 ?

    基于GPU的寬帶信號時延差與相位差估計方法*

    2023-11-25 13:08:32毛飛龍焦義文高澤夫王育欣
    電訊技術(shù) 2023年11期
    關(guān)鍵詞:子帶線程寬帶

    毛飛龍,焦義文,馬 宏,李 超,高澤夫,王育欣

    (航天工程大學(xué) 電子與光學(xué)工程系,北京 101416)

    0 引 言

    隨著航天技術(shù)與深空探測技術(shù)的迅猛發(fā)展,各國對空間的探測向著更深更遠的范圍拓展,探測距離將從目前月球探測的40萬km拓展到火星探測的4億km和木星探測的10億km[1]。這將導(dǎo)致嚴重的傳輸損耗[2],同時也對深空探測航天器和地面設(shè)備提出了兩點要求:一是遠距離通信能力;二是高速率傳輸能力[3]。由于深空探測航天器的等效全向發(fā)射功率(Effective Isotropic Radiated Power,EIRP)有限,為了實現(xiàn)對低信噪比信號的高質(zhì)量接收,就必須增大地面站的天線口徑。根據(jù)NASA的任務(wù)參數(shù)計算可知,火星探測任務(wù)中,傳輸1 Mb/s的信號需要直徑80 m的天線。而目前NASA深空網(wǎng)(Deep Space Network,DSN)的直徑70 m天線已接近工程極限,伺服機動十分困難,靈活性較差。同時,天線的口徑越大,其信號的波束越窄,越難以捕獲探測器的下行信號。因此,單天線的性能提升已經(jīng)到了緩慢發(fā)展接近停滯的狀態(tài),通過多天線合成的方法獲得更高的深空通信增益是值得大力發(fā)展的方法[4]。

    隨著高速率通信系統(tǒng)的發(fā)展,信號的帶寬、傳輸速率在近年來也有了大幅度提升。信號帶寬可達幾十MHz甚至幾百MHz,這將為地面接收設(shè)備帶來巨大挑戰(zhàn)。寬帶信號多天線合成系統(tǒng)是一個數(shù)據(jù)密集、計算密集、邏輯復(fù)雜的計算系統(tǒng),使用單純的中央處理器(Central Processing Unit,CPU)計算已不能滿足運算要求,需要選用新的計算架構(gòu)來提高系統(tǒng)實時處理的性能。近些年來,圖形處理單元(Graphics Processing Unit,GPU)在計算領(lǐng)域得到了充分應(yīng)用,已發(fā)展成為一種高度并行化、多線程、多核的通用計算設(shè)備,具有杰出的計算能力和極高的存儲器帶寬,并被越來越多地應(yīng)用于圖形處理之外的計算領(lǐng)域?;贕PU的信號處理技術(shù)成為眾多領(lǐng)域的熱點,如射電天文、雷達、無線通信、人工智能等。因此,基于GPU的寬帶信號多天線合成具有一定的工程研究意義。多天線信號合成技術(shù)的本質(zhì)就是在消除不同天線信號間的時延差和相位差后進行相干相加,因此天線間時延差和相位差的估計、時延差與相位差的補償為多天線合成中的兩項關(guān)鍵技術(shù)。本文旨在現(xiàn)有寬帶信號時延差與相位差估計方法的基礎(chǔ)上,設(shè)計基于GPU的寬帶信號時延差與相位差估計方法,通過算法的并行化映射尋求一種精度高且實時性好的估計方法。

    1 寬帶信號估計方法分析

    1.1 多天線信號合成方法

    多天線信號合成的概念最早是由美國噴氣推進實驗室于1965年提出的[25]。該技術(shù)通過將不同天線接收到同一信源的信號進行合成,從而提高接收信號的信噪比(Signal-to-Noise Ratio,SNR)。圖1為多天線信號合成原理示意圖。以雙天線為例,兩天線接收到來自同一信源的信號,通過AD采集為數(shù)字信號,然而由于各天線信號傳播路徑的不同,導(dǎo)致信號間存在各種各樣的差異,這些差異將損失信號間的相干性。此時兩天線信號存在時延差與相位差。將其中值較高的天線當作參考天線,另一天線為待補償天線,對其進行時延差與相位差估計并補償,將兩天線的時延與相位對齊。之后對兩天線信號相加,就得到了信號幅度翻倍、信噪比提升的合成信號。多天線信號合成技術(shù)的本質(zhì)就是在消除不同天線信號間的時延差和相位差后進行相干相加。由于各天線間的噪聲是隨機非相關(guān)的,理論上使用N個天線合成信號的信噪比是單天線接收信號的N倍。

    圖1 多天線信號合成原理示意

    多天線信號合成技術(shù)經(jīng)過幾十年的發(fā)展,逐漸形成了5種合成方案[6]:全頻譜合成、基帶合成(Baseband Combining,BC)、符號流合成(Symbol-Stream Combining,SSC)、載波組陣合成(Carrier Arraying,CA)、復(fù)符號合成(Complex-Symbol Combining,CSC)[7]。除全頻譜合成方法之外,其余四種方法都屬于基于載波跟蹤技術(shù)的合成方法,都需要陣元載波鎖定后合成,不適用于當前大規(guī)模的多天線信號合成。本文在考慮分析了各種方法的利弊后,選用全頻譜合成方法對寬帶信號進行合成。

    1.2 現(xiàn)有寬帶信號時延差與相位差估計方法

    寬帶信號合成的關(guān)鍵是對寬帶信號間的殘余時延和相位進行高精度估計和補償,目前可以分為時域合成方法和頻域合成方法。

    1.2.1 時域合成方法

    NASA深空網(wǎng)在其全頻譜處理陣(Full Spectrum Processing Array,FSPA)設(shè)備中,使用了一種基于上下邊帶的殘余時延和相位差聯(lián)合估計方法,原理如圖2所示。該方法首先通過提取信號濾波器(中心頻率可調(diào)帶通濾波器)提取各天線信號的上邊帶和下邊帶,然后利用相位估計算法分別估計待校準天線與參考天線上下邊帶的相位差,最終兩相位點構(gòu)成的直線的斜率即為殘余時延,截距即為載波相位差。

    圖2 基于上下邊帶的殘余時延和相位估計方法的原理

    在圖2中,L1和L2表示兩天線信號經(jīng)過提取信號濾波器(中心頻率可調(diào)帶通濾波器)得到的下邊帶,其中心頻率為fL;U1和U2表示兩天線信號的上邊帶,其中心頻率為fU;相位差φL和φU為相位差估計算法估計得到的下邊帶間和上邊帶間的相位差。則兩相位點構(gòu)成直線的斜率即為殘余時延,截距即為載波處的殘余相位。設(shè)殘余時延為Δτ,載波處的殘余相位為Δφc,則有

    (1)

    (2)

    該方法的優(yōu)勢在于運算量小,而且將時延估計問題轉(zhuǎn)化為相位估計問題,在低信噪比條件下,各邊帶可以采用高性能相位估計算法來提高估計精度。然而,該方法存在兩個問題,下面進行詳細分析。

    1)相位卷繞問題

    由于相位估計算法只能提取初相差,當上下邊帶頻率差較大或者殘余時延較大時出現(xiàn)相位卷繞問題。而該方法只取了上下邊帶兩個相位差值進行估計,在沒有時延差與相位差先驗知識的情況下無法正確地進行解卷繞處理,進而導(dǎo)致殘余時延和殘余相位與其真實值差距較大。

    本文研究寬帶信號的合成方法,以典型的Ka頻段信號為例,其信號帶寬可達500 MHz,若采用基于上下邊帶的殘余時延和相位差聯(lián)合估計方法,在接收信號的上下邊帶分別進行互相關(guān),當上下邊帶頻率差較大或者殘余時延較大時,得到上下邊帶的相位結(jié)果為其卷繞的值。因此,為了得到正確的殘余時延與殘余相位,就要保證不發(fā)生相位卷繞(φU-φL<2π),此時

    (3)

    當Δf為250 MHz時,Δτ為4×10-9s,即為4 ns,而Ka頻段的寬帶信號通常殘余時延為幾十ns甚至幾百ns,因此引入可消除相位卷繞現(xiàn)象的殘余時延與殘余相位估計方法是寬帶信號合成中的重點之一。

    2)頻帶內(nèi)相位滑動問題

    整數(shù)時延補償都采用整數(shù)時鐘移位的補償方法,其最高精度為時鐘的周期Ts,整數(shù)時延補償之后必然存在Δτ

    Δφ=2π·fU·Δτ-2π·fL·Δτ=2π·Δf·Δτ。

    (4)

    式中:Δφ為相位滑動;Δf為頻率高處與頻率低處的差值。相位滑動現(xiàn)象將影響合成信號的質(zhì)量,當Δφ=0時,信號的帶寬較窄,可近似看作同一頻率,各處的相位補償值相同,信號合成的性能較好,如圖3所示;當Δφ=π時,相位補償為真實值的負值,將導(dǎo)致合成信號出現(xiàn)反向相消的現(xiàn)象,如圖4所示。

    圖3 無相位滑動現(xiàn)象示意圖

    圖4 有相位滑動現(xiàn)象示意圖

    為了抑制反向相消現(xiàn)象,需保證Δφ<<π,即2π·Δf·Δτ<<π,若Δf為250 MHz,需滿足Δτ<<2 ns,同樣地,需要滿足采樣時鐘Ts<<2 ns,即采樣率需遠大于500 MHz??紤]到信號帶寬為500 MHz時,根據(jù)奈奎斯特采樣定理,只需采樣率大于1 000 MHz即可。而遠大于500 MHz的采樣率將帶來采樣效率極大地浪費,同時也增加了系統(tǒng)實時處理的難度。因此,引入可消除相位滑動現(xiàn)象的殘余時延與殘余相位估計方法也是寬帶信號合成中的重點之一。

    由上述分析可知,時域合成方法都存在補償精度受采樣率限制的問題。即使假設(shè)時延估計誤差為零,補償后的小數(shù)時延(也稱作分數(shù)時鐘時延,最大為半個采樣時鐘周期)也將導(dǎo)致頻帶內(nèi)的相位滑動,從而惡化合成性能。尤其當帶寬較大時,上述方法不再適用。因此,引入更合理的小數(shù)時延補償方法是寬帶信號合成的重點和難點問題之一。

    1.2.2 頻域合成方法

    頻域合成方法的處理流程為首先利用分析濾波器組對寬帶信號進行信道化,得到低采樣速率的子帶信號,再利用相位差估計技術(shù)得到不同天線子帶間的相位差,進而估計寬帶信號間的殘余時延和相位并對其進行補償,然后對各天線子帶進行合成,最終利用綜合濾波器組對各天線合成后的子帶進行重構(gòu),得到原始采樣率的寬帶合成信號。在對殘余時延進行補償時,首先采用數(shù)字延遲線進行整數(shù)時延補償,殘余的小數(shù)時延以及相位的補償則在頻域進行,通過對各個子帶進行線性相位偏移來實現(xiàn)。相比較時域合成方法,頻域合成方法具有如下優(yōu)勢:一是通過對子帶構(gòu)成的相頻曲線進行解卷繞,可以有效解決時延差引起的2π模糊問題;二是通過將各子帶中心頻率獨立進行相位對齊,可將由小數(shù)時延引起的整帶內(nèi)相位滑動限制在子帶范圍內(nèi),大大減小了小數(shù)時延引起的合成損失。

    由上述分析可知,頻域合成方法可以有效解決時域合成方法存在的問題,已得到廣泛關(guān)注和研究,是目前寬帶信號全頻譜合成的最佳實現(xiàn)方法??梢灶A(yù)見,該方法將是未來大規(guī)模寬帶天線組陣信號合成的主流方法。

    1.3 本文采用的寬帶信號時延差與相位差估計方法

    由1.2節(jié)分析可知,現(xiàn)有的基于上下邊帶的殘余時延和相位差聯(lián)合估計方法在處理本文中帶寬500 MHz的信號時存在兩方面問題,分別是相位卷繞問題和頻帶內(nèi)相位滑動問題。對相位卷繞問題分析可知,其主要出現(xiàn)在上下邊帶頻率差較大或者殘余時延較大時,而信號的殘余時延差Δτ無法改變,只能通過減小頻率差Δf來改進。對頻帶內(nèi)相位滑動問題分析可知,Δφ<<π時,該現(xiàn)象可消除。而Δφ=2π·Δf·Δτ,同樣地,由于Δτ無法改變,只能通過減小頻率差Δf來減小相位滑動問題帶來的影響。綜合上述因素,減小頻率差Δf是一種改進思路,結(jié)合軟件無線電中信道化的思想,將各天線的信號拆分為多個獨立的子帶,之后對不同天線對應(yīng)的子帶信號進行處理。下面進行詳細分析:

    將寬帶信號拆分為K個子帶,對相應(yīng)的子帶進行相位差估計,得到K個子帶對應(yīng)的相位差信息。此時,只要保證相位差出現(xiàn)卷繞現(xiàn)象是由于某一子帶的相位差值在[-π,π]之外引起的,而并非相鄰兩個相位差的差值大于2π引起的,換句話說,只要保證相鄰兩個相位差不存在卷繞現(xiàn)象就可正確進行解卷繞。在各子帶中,頻率間隔在數(shù)值上為fs/K,要使相鄰兩個相位差不存在卷繞現(xiàn)象,即

    (5)

    式中:τd為子帶內(nèi)的最大時延。式(5)可化簡為

    (6)

    由式(3)可知,K需滿足

    K>2fs·τd。

    (7)

    因此,當時延差τd的精度確定時,子帶的最小數(shù)量即確定。

    對于頻帶內(nèi)相位滑動問題,將寬帶信號拆為子帶信號進行處理,可將相位滑動限制在子帶范圍內(nèi),如圖5所示??梢钥闯?雖然沒有徹底消除相位滑動問題,但相位滑動被限制在了很低的范圍內(nèi),可大幅度降低相位滑動問題對合成性能的影響,使合成信號的功率譜保持在可接收的范圍內(nèi),如圖6所示。

    圖5 相位滑動現(xiàn)象被限制在子帶內(nèi)示意

    圖6 減弱相位滑動前后合成信號功率譜示意

    寬帶信號合成需考慮相位卷繞問題和頻帶內(nèi)相位滑動問題,根據(jù)上述分析可知,將寬帶信號拆為子帶進行處理可有效地降低這兩種問題帶來的影響。同時,寬帶信號合成方案必須考慮相頻特性差異造成的相位非相干性以及采樣率高帶來的數(shù)據(jù)吞吐率大的問題。因此,本文選用寬帶信號頻域合成方法,如圖7所示。

    圖7 寬帶信號頻域全頻譜合成方法

    基本思路為首先根據(jù)頻率范圍對寬帶信號進行子帶拆分,得到低采樣速率的子帶信號。利用相位差估計技術(shù)得到不同天線子帶信號間的相位差,進而估計并補償寬帶信號間的殘余時延和相位,然后對各天線子帶進行合成。最終通過對各天線合成后的子帶進行重構(gòu),得到原始采樣率的寬帶合成信號。

    寬帶信號時延差與相位差估計示意圖如圖8所示。以雙天線為例,天線信號經(jīng)AD采集與子帶拆分之后,每個天線信號得到了頻率范圍不同的K個子帶。對兩個天線對應(yīng)頻率的子帶信號進行相位差估計,得到K個子帶的相位差結(jié)果。由于寬帶信號的相位信息φ(ω)是頻率ω的線性函數(shù),且滿足

    圖8 寬帶信號時延差與相位差估計示意

    φ(ω)=ω·Δτk+θk。

    (8)

    則寬帶信號的殘余延遲差Δτk多項式為一次項參數(shù),即斜率,殘余相位差θk為零次項參數(shù),即截距。因此,在估計得到兩個天線子帶間的相位差信息后,通過線性最小二乘擬合就可以得到兩天線寬帶信號的殘余時延與殘余相位。由于子帶相位差估計方法最后需要進行反正切運算,因此,當殘余時延對應(yīng)整周期內(nèi)的角度大于2π時,就會出現(xiàn)相位卷繞現(xiàn)象。因此,需要在擬合之前進行解卷繞得到正確的相位值。

    圖9為基于GPU的寬帶信號時延差與相位差估計流程,具體步驟如下:

    Step1 參考天線與待補償天線首先經(jīng)過FFT均勻信道化器將寬帶信號拆為多個子帶信號。

    Step2 待補償天線的子帶信號均與參考天線對應(yīng)的子帶信號進行共軛相乘運算,完成互相關(guān)。

    Step3 將多段互相關(guān)的結(jié)果進行累加平均并求反正切,得到對應(yīng)子帶相位差的估計值。

    Step4 待補償天線各子帶的相位值進行解卷繞運算得到正確的相位值。

    Step5 對待補償天線所有子帶的相位差結(jié)果進行最小二乘線性擬合,得到寬帶信號的時延差與相位差。

    其中Step 1~3即為子帶相位差估計算法中典型的Simple算法。

    2 GPU并行優(yōu)化

    2.1 線程并行優(yōu)化

    由1.3節(jié)可知,子帶相位差估計算法中Simple算法主要包括FFT、共軛相乘、累加平均三個模塊,下面對各模塊的GPU實現(xiàn)方法進行研究。

    在CPU平臺下,N點FFT的運算量為N/2×lbN次的復(fù)數(shù)乘法和N×lbN次的復(fù)數(shù)加法,N表示做傅里葉變換的點數(shù)。對于如此大量的高速并行數(shù)據(jù),在CPU平臺中無法做到實時處理。但是由于FFT計算具有高度并行性這一顯著特點,本文考慮將快速傅里葉變換部分放在GPU中進行并行加速處理,利用CUDA中自帶的CUFFT庫,可實現(xiàn)大量高速數(shù)據(jù)的實時FFT計算。調(diào)用CUFFT函數(shù)庫內(nèi)的cufftPlanD函數(shù)對數(shù)據(jù)進行快速傅里葉變換。在GPU上進行快速傅里葉變換需要完成以下幾步工作:

    Step1 使用 malloc()函數(shù)為輸入的數(shù)據(jù)分配主機端(CPU)內(nèi)存,并存儲輸入數(shù)據(jù)。

    Step2 使用 cudaMalloc()函數(shù)為輸入數(shù)據(jù)分配設(shè)備端(GPU)顯存。

    Step3 使用 cudaMemcpy()函數(shù)將主機(CPU)內(nèi)存數(shù)據(jù)傳輸?shù)皆O(shè)備端(GPU)顯存。

    Step4 創(chuàng)建一個plan,調(diào)用 cufftPlanD函數(shù)創(chuàng)建快速傅里葉變換。

    Step5 執(zhí)行plan,使用 cufftExecC2C()函數(shù)完成plan的計算任務(wù)。

    Step6 用cudaMemcpy()將GPU端顯存數(shù)據(jù)傳輸?shù)街鳈C內(nèi)存。

    上述工作執(zhí)行完成以后如果這個plan不再需要,則調(diào)用cufftDestroy()函數(shù)銷毀該plan及分配給它的計算資源。以上各步的具體流程和所調(diào)用的CUDA函數(shù)如圖10所示。在CUFFT庫中,有3種不同的庫函數(shù):cufftExecC2C(復(fù)數(shù)到復(fù)數(shù))、cufftExecC2R(復(fù)數(shù)到實數(shù))、cufftExecR2C(實數(shù)到復(fù)數(shù)),它們的計算耗時由于本身運算量的差異而有所不同。因此,為最大限度地縮短運算時間,必須針對本提相過程的實際情況,選擇最合適的庫函數(shù)類型。在本系統(tǒng)輸入信號的FFT計算中,由于輸入是復(fù)信號,而輸出是復(fù)數(shù)信號,因此最宜采用cufftExecC2C類型的FFT。

    圖11是GPU中FFT與共軛相乘模塊的實現(xiàn)框圖。調(diào)用cufft庫進行FFT、多個線程同時進行共軛相乘運算,代碼如下:

    圖11 FFT與共軛相乘GPU實現(xiàn)框圖

    cufftPlan1d(&plan,NFFT,CUFFT_C2C,nBatch);

    cufftExecC2C(plan,d_signal,d_signal_FFT,CUFFT_FORWARD);

    cufftDestroy(plan);

    Conj<<>>(d_signal_1_FFT,d_signal_4_FFT,d_Conj,nsize);

    __global__ void Conj(float2 *src,float2 *send,float2 *dst,int length)

    { int tid = threadIdx.x + blockIdx.x * blockDim.x;

    if(tid >= length)

    return;

    dst[tid].x =(src[tid].x * send[tid].x + src[tid].y * send[tid].y);

    dst[tid].y =(src[tid].y * send[tid].x - src[tid].x * send[tid].y);}

    圖12為累加平均并求相位核函數(shù)實現(xiàn)框圖。將共軛相乘后數(shù)據(jù)傳輸至累加平均并求相位核函數(shù)中進行能量累積和提相等處理。將每個天線的N個數(shù)據(jù)等分成x段進行相加平均,每段M/x點,累加平均后使用atan 2函數(shù)求出M/x個復(fù)數(shù)的相角,所得相角即為子帶間的相位差。

    圖12 累加平均并求相位GPU實現(xiàn)框圖

    本文設(shè)計的累加平均并求相位核函數(shù)部分的CUDA代碼如下所示:

    FAdd<<>>(d_Conj,d_phase,d_amp,fft_time,NFFT);

    __global__ void FAdd(float2 *src,float *dst,float *dst2,int jft,int length)

    { int tid = threadIdx.x + blockIdx.x * blockDim.x;

    if(tid >= length)

    return;

    int curpos;

    double sum_x = 0.0;

    double sum_y = 0.0;

    for(int i = 0;i

    { curpos = tid + i * length;

    sum_x = sum_x +((double)src[curpos].x);

    sum_y = sum_y +((double)src[curpos].y);}

    sum_x = sum_x / jft;

    sum_y = sum_y / jft;

    dst[tid]= atan2(sum_y,sum_x);

    dst2[tid]= 10 * log10(sum_x * sum_x + sum_y * sum_y);}

    2.2 異步流并發(fā)優(yōu)化

    如圖9所示,基于GPU的寬帶信號時延差與相位差估計流程主要包括在GPU平臺上進行FFT、子帶相位差估計、解卷繞、擬合等。子帶相位差估計具體包括共軛相乘、累加平均等運算,FFT過程可用CUDA的庫函數(shù)CUFFT進行計算,而CUFFT庫提供的二維FFT計算模型為本文的多路FFT高效并行實現(xiàn)提供了可能。

    流(stream)是CUDA環(huán)境中的一個并行概念,在GPU中可使用流來處理并發(fā)操作[8]。流本質(zhì)上是一個操作隊列,該隊列可由主機上不同的線程發(fā)出并以異步方式在GPU上按順序執(zhí)行。不同的流可以并發(fā)執(zhí)行,多個流并發(fā)執(zhí)行不同路數(shù)據(jù)的FFT、子帶相位差估計、解卷繞、擬合等運算,可大大提高運算的并行性?;贕PU的寬帶信號時延差與相位差估計,各路數(shù)據(jù)完全獨立,每一路數(shù)據(jù)的處理包含F(xiàn)FT、共軛相乘、累加平均、解卷繞、擬合等多個操作,因此,可進一步利用異步流并發(fā)技術(shù)對估計方法進行優(yōu)化。

    設(shè)天線數(shù)為N,若將每一路天線數(shù)據(jù)的處理過程用一個CUDA流來處理,則基于CUDA流式架構(gòu)的相位差估計算法實現(xiàn)框圖如圖13所示。由圖中虛線框所示部分可知,在該算法中,同一時刻最多有N個Kernel在同時執(zhí)行。

    圖13 基于CUDA流式架構(gòu)的相位差估計算法實現(xiàn)框圖

    圖14為N=3時流式架構(gòu)相位差估計算法流程。從圖中可知,基于線程的并行方法中,數(shù)據(jù)傳輸和核函數(shù)順序執(zhí)行,即一次操作完全完成后才進行下一項操作,在某次操作執(zhí)行期間其他所有操作都處于等待狀態(tài)。而在流式結(jié)構(gòu)并發(fā)算法中,數(shù)據(jù)傳輸和核函數(shù)的執(zhí)行被拆分為多個并行的操作序列,隨著CUDA流的并發(fā)啟動,數(shù)據(jù)傳輸過程被核函數(shù)執(zhí)行過程較好地隱藏。另外,從圖14的算法結(jié)構(gòu)對比結(jié)果可看出,3路并行流式結(jié)構(gòu)并行算法在執(zhí)行時,同一時刻駐留在GPU上的核函數(shù)最多可達3個,其執(zhí)行效率明顯高于線程并行算法。

    圖14 基于CUDA流式架構(gòu)的4天線相位差估計算法實現(xiàn)框圖

    2.3 小結(jié)

    首先利用線程并行的方式對時延差與相位差估計方法中的共軛相乘、累加平均、解卷繞等模塊的實現(xiàn)算法進行了相應(yīng)的改進和優(yōu)化,實現(xiàn)了多個線程同時處理多個數(shù)據(jù)的運算,提高了每個核函數(shù)本身的運算效率。之后利用異步流并發(fā)的方式對估計方法進一步優(yōu)化,每路信號采用獨立的流進行運算,從結(jié)構(gòu)層面有效提高了數(shù)據(jù)的并行處理效率。

    3 仿真驗證

    本文采用Simple算法對寬帶信號進行合成,利用相位干涉儀中的4個天線進行實測數(shù)據(jù)信號合成。系統(tǒng)由4個陣元構(gòu)成,陣元擺放形式如圖15所示。系統(tǒng)硬件平臺選用NVIDIA Tesla V100顯卡,接收信號頻段采用S頻段,利用相位校正(Phase CALibration,PCAL)信號進行性能分析,PCAL信號的頻率特性具有梳狀頻譜特性和線性相位特性。本系統(tǒng)的核心設(shè)計指標如表1所示。

    表1 仿真實驗參數(shù)

    圖15 相位干涉儀布陣示意

    3.1 正確性驗證

    首先驗證GPU線程并行與GPU流并行估計方法的正確性,設(shè)置天線4為參考天線,天線1與天線4的時延差為2個采樣點、相位差為-π/2。利用CPU串行Simple算法、GPU線程并行Simple算法與GPU流并行Simple算法估計天線1與參考天線的時延差與相位差。每次估計出時延差與相位差后對天線1信號進行修正,按式(9)更新時延和相位補償值:

    (9)

    式中:delaynow和phasenow為本次迭代計算出的時延差與相位差;delayi-1和phasei-1為上一次迭代的時延和相位補償值;delayi和phasei為本次更新后的時延補償值。步長因子為0.5,記錄迭代30次的時延和相位補償值。

    圖16為CPU串行Simple算法、GPU線程并行Simple算法與GPU流并行Simple算法的時延和相位補償值隨迭代次數(shù)變化的曲線。20次迭代后,三種算法的時延和相位補償值都收斂于真值,都可實現(xiàn)高精度時延和相位估計,驗證了GPU線程并行與GPU流并行估計方法的正確性。

    利用本文中的并行Simple算法與傳統(tǒng)串行Simple算法在不同單天線信噪比的條件下分別進行合成,進行蒙特卡洛仿真并得出合成損失與合成效率的平均值,得到不同信噪比條件下兩種方法的合成損失與合成效率。圖17給出了兩種方法在不同單天線信噪比條件下合成損失與合成效率的仿真結(jié)果。

    (a)合成損失

    從圖17可以看出,本文中的并行Simple算法與傳統(tǒng)串行Simple算法的合成損失與合成效率曲線重合,兩者的合成性能一致,與分析得出的結(jié)論一致。同時,兩種算法隨著單天線信噪比的升高,合成損失降低,合成效率升高。仿真結(jié)果表明,本文中的并行Simple算法與傳統(tǒng)串行Simple算法在信噪比高于-15 dB的條件下有較好的合成性能。

    3.2 加速效果驗證

    為了詳細分析核函數(shù)運行過程,利用NVIDIA Nsight Systems對流式架構(gòu)Simple算法進行分析,結(jié)果如圖18和圖19所示。

    圖18 CUDA流式架構(gòu)Simple算法時序圖

    圖19 CUDA流式架構(gòu)Sumple算法時序圖

    圖18中上方為未采用異步流并發(fā)方式優(yōu)化的Simple算法時序圖,可以看到同一時刻最多只有一個核函數(shù)駐留,不同天線的相同運算模塊順序運行。圖18中下方為采用異步流并發(fā)方式優(yōu)化的Simple算法時序圖。由于Simple算法將其中1個天線設(shè)置為參考天線,因此,在天線數(shù)為4的情況下只需進行3路信號流并發(fā)即可。從圖18中可以看到同一時刻最多有3個核函數(shù)駐留,不同天線的相同運算模塊分別在3個不同流上執(zhí)行。

    圖19中上方為未采用異步流并發(fā)方式優(yōu)化的Sumple算法時序圖,可以看到同一時刻最多只有一個核函數(shù)駐留,不同天線的相同運算模塊順序運行。圖19中下方為采用異步流并發(fā)方式優(yōu)化的Sumple算法時序圖。由于Sumple算法將所有天線的加權(quán)值設(shè)置為參考天線,每路天線信號都要與虛擬的參考天線進行時延差與相位差估計,因此,在天線數(shù)為4的情況下需要進行4路信號的流并發(fā)。從圖19中可以看到同一時刻最多有4個核函數(shù)駐留,不同天線的相同運算模塊分別在4個不同流上執(zhí)行。從時序圖中可明顯地看出異步流并發(fā)適合于時延差與相位差估計方法的優(yōu)化加速。

    下面對GPU線程并行與GPU流并行估計方法相比CPU串行方法的加速比進行測試,分別對30次迭代中CPU串行、GPU線程并行與GPU流并行估計方法的耗時進行記錄。共測試4組,每組的區(qū)別只有迭代的數(shù)據(jù)量,分別為512,5 120,51 200,512 000。表2加速比測試結(jié)果。

    表2 加速比測試結(jié)果

    圖20為GPU線程并行與GPU流并行估計方法相比CPU串行方法的加速比測試結(jié)果圖,可以看出,當數(shù)據(jù)量為512時,GPU線程并行方法相比CPU串行方法的加速比為0.98,表明此時的GPU線程并行方法的平均耗時大于CPU串行方法的平均耗時。因為當數(shù)據(jù)量較小時,每個核函數(shù)調(diào)用的線程數(shù)較少,不能充分地發(fā)揮GPU多線程的優(yōu)勢,因此沒有加速效果。GPU流并行方法相比CPU串行方法的加速比為2.01,雖然數(shù)據(jù)量較小時不能發(fā)揮GPU多線程的優(yōu)勢,但是流并行方法在結(jié)構(gòu)上對運算順序做了優(yōu)化可以利用多流并發(fā)的方式同時對多路信號進行運算,因此加速效果明顯。隨著數(shù)據(jù)量的升高,GPU線程并行與GPU流并行估計方法相比CPU串行方法的加速比也升高,表明GPU線程并行與GPU流并行估計方法適合大規(guī)模數(shù)據(jù)的實時處理。從表2可以看出,CPU串行、GPU線程并行、GPU流并行估計方法在數(shù)據(jù)量為512 000的估計耗時約為17.43 ms,0.33 ms與0.14 ms。本文中系統(tǒng)的采樣率為56 000 000,處理512 000數(shù)據(jù)量耗時若小于10 ms即可實現(xiàn)實時估計。因此,CPU串行估計方法無法實現(xiàn)實時估計,本文中的GPU線程并行、GPU流并行估計方法可實現(xiàn)實時估計。此外,本文中的數(shù)據(jù)為仿真數(shù)據(jù),提前讀入GPU,而實際系統(tǒng)中,受限于系統(tǒng)吞吐率與GPU并行性能的限制,實際系統(tǒng)的加速比不能無限增大。

    圖20 加速比測試結(jié)果

    3 結(jié) 論

    針對大規(guī)模多天線信號合成系統(tǒng)的天線數(shù)量越來越多,接收信號的帶寬越來越大,單獨使用CPU已不能滿足運算需求這一問題,本文研究了寬帶信號時延差與相位差估計方法,分析了算法并行的可行性并設(shè)計了基于GPU的寬帶信號時延差與相位差估計方法,通過仿真對比實驗驗證了本文提出的并行算法的優(yōu)越性。

    本文的并行算法可以運用到未來的大規(guī)模寬帶信號多天線信號合成系統(tǒng)中,實現(xiàn)高速、準確、實時的信號合成,有力支持未來的一系列深空探測任務(wù)。

    猜你喜歡
    子帶線程寬帶
    一種基于奇偶判斷WPT的多音干擾抑制方法*
    裝寬帶的人
    文苑(2020年7期)2020-08-12 09:36:04
    子帶編碼在圖像壓縮編碼中的應(yīng)用
    電子制作(2019年22期)2020-01-14 03:16:24
    淺談linux多線程協(xié)作
    基于虛擬孔徑擴展的子帶信息融合寬帶DOA估計
    一種新穎的寬帶大功率分配器
    可否把寬帶作為社會福利
    什么是寬帶?
    Linux線程實現(xiàn)技術(shù)研究
    基于子帶模式的AMC技術(shù)算法分析
    久久天堂一区二区三区四区| 久久精品aⅴ一区二区三区四区| 亚洲第一青青草原| 国产成人影院久久av| 亚洲精品乱久久久久久| 我要看黄色一级片免费的| 999久久久国产精品视频| 国产真人三级小视频在线观看| 国产三级黄色录像| 人人妻人人澡人人看| 久久国产精品人妻蜜桃| 国产视频一区二区在线看| 日韩电影二区| 一级片免费观看大全| 中文字幕色久视频| 老司机影院成人| 女人高潮潮喷娇喘18禁视频| 丝袜美腿诱惑在线| 两人在一起打扑克的视频| 日日摸夜夜添夜夜添小说| 精品少妇一区二区三区视频日本电影| 9191精品国产免费久久| 黄色视频,在线免费观看| 国产精品av久久久久免费| 欧美精品高潮呻吟av久久| 另类亚洲欧美激情| 国产成人精品在线电影| 久久久水蜜桃国产精品网| 一本一本久久a久久精品综合妖精| 中文字幕人妻丝袜一区二区| 国产不卡av网站在线观看| 他把我摸到了高潮在线观看 | 久久久久久久大尺度免费视频| 欧美老熟妇乱子伦牲交| 啪啪无遮挡十八禁网站| 亚洲五月婷婷丁香| 日韩一卡2卡3卡4卡2021年| 亚洲国产中文字幕在线视频| 午夜福利影视在线免费观看| 动漫黄色视频在线观看| 妹子高潮喷水视频| 精品人妻在线不人妻| 俄罗斯特黄特色一大片| 中文字幕另类日韩欧美亚洲嫩草| 国产成人精品久久二区二区91| 99久久国产精品久久久| 精品国产乱子伦一区二区三区 | 亚洲精品美女久久av网站| 亚洲精品粉嫩美女一区| 中亚洲国语对白在线视频| 久9热在线精品视频| 久久精品亚洲av国产电影网| av福利片在线| 老司机午夜十八禁免费视频| 少妇的丰满在线观看| 中文字幕最新亚洲高清| 9191精品国产免费久久| 美女中出高潮动态图| 亚洲国产精品一区二区三区在线| 亚洲精品国产色婷婷电影| 日韩 亚洲 欧美在线| 国产精品影院久久| 男女床上黄色一级片免费看| 丝袜人妻中文字幕| 亚洲人成电影免费在线| 中文字幕高清在线视频| 国产精品麻豆人妻色哟哟久久| 香蕉丝袜av| 菩萨蛮人人尽说江南好唐韦庄| 亚洲中文日韩欧美视频| 国产成人系列免费观看| 老鸭窝网址在线观看| 一级毛片电影观看| 每晚都被弄得嗷嗷叫到高潮| 一个人免费在线观看的高清视频 | 成在线人永久免费视频| 少妇的丰满在线观看| 99国产精品一区二区三区| 国产精品久久久av美女十八| 男女边摸边吃奶| 正在播放国产对白刺激| 999精品在线视频| 在线天堂中文资源库| av福利片在线| 久久精品亚洲熟妇少妇任你| 久久精品熟女亚洲av麻豆精品| 美女大奶头黄色视频| 免费av中文字幕在线| 亚洲精品第二区| 亚洲国产精品999| 亚洲熟女精品中文字幕| 无限看片的www在线观看| 天天影视国产精品| 又大又爽又粗| 天堂8中文在线网| 丝袜在线中文字幕| 国产三级黄色录像| 国产亚洲欧美精品永久| 亚洲专区字幕在线| 色老头精品视频在线观看| 欧美日韩视频精品一区| 各种免费的搞黄视频| 夫妻午夜视频| 欧美另类一区| 一本久久精品| 大片免费播放器 马上看| 高清欧美精品videossex| 黑人操中国人逼视频| 两性午夜刺激爽爽歪歪视频在线观看 | 亚洲第一欧美日韩一区二区三区 | 不卡一级毛片| 国产免费现黄频在线看| 操美女的视频在线观看| av不卡在线播放| 久久精品熟女亚洲av麻豆精品| 18在线观看网站| 精品国产乱子伦一区二区三区 | 日韩大片免费观看网站| 九色亚洲精品在线播放| 亚洲精品粉嫩美女一区| 欧美另类亚洲清纯唯美| 亚洲精品中文字幕在线视频| 免费观看av网站的网址| 亚洲av电影在线观看一区二区三区| 超碰97精品在线观看| 五月开心婷婷网| 亚洲伊人久久精品综合| 国产精品一区二区精品视频观看| 欧美激情 高清一区二区三区| 亚洲一卡2卡3卡4卡5卡精品中文| 午夜日韩欧美国产| 国产主播在线观看一区二区| a在线观看视频网站| 亚洲色图综合在线观看| 精品国产一区二区三区四区第35| 男人操女人黄网站| 中文字幕高清在线视频| 999精品在线视频| 一个人免费看片子| 久久久精品国产亚洲av高清涩受| 真人做人爱边吃奶动态| 极品少妇高潮喷水抽搐| 亚洲精品美女久久久久99蜜臀| 丝袜人妻中文字幕| 久久综合国产亚洲精品| 久久国产精品影院| 午夜免费成人在线视频| 日韩视频一区二区在线观看| 亚洲伊人久久精品综合| 夜夜骑夜夜射夜夜干| 国产精品影院久久| 日本黄色日本黄色录像| 成年人黄色毛片网站| 黑人欧美特级aaaaaa片| 亚洲欧美激情在线| 午夜福利一区二区在线看| 欧美激情高清一区二区三区| 天天操日日干夜夜撸| 日韩一卡2卡3卡4卡2021年| 精品国产乱子伦一区二区三区 | 日韩熟女老妇一区二区性免费视频| 久久精品亚洲av国产电影网| 91精品伊人久久大香线蕉| 亚洲欧美色中文字幕在线| 色精品久久人妻99蜜桃| 国产1区2区3区精品| 免费黄频网站在线观看国产| 亚洲精品自拍成人| 久久精品国产综合久久久| 国产淫语在线视频| 国产麻豆69| 水蜜桃什么品种好| 黄色视频不卡| 午夜久久久在线观看| 美女视频免费永久观看网站| 欧美xxⅹ黑人| 一级毛片精品| av视频免费观看在线观看| 久久香蕉激情| 一级黄色大片毛片| 丝袜美腿诱惑在线| 国产av国产精品国产| 免费黄频网站在线观看国产| 永久免费av网站大全| 欧美另类亚洲清纯唯美| 欧美国产精品va在线观看不卡| 久久精品国产a三级三级三级| 蜜桃在线观看..| 久久免费观看电影| 婷婷丁香在线五月| 丰满少妇做爰视频| 新久久久久国产一级毛片| 美女大奶头黄色视频| 国产99久久九九免费精品| 少妇精品久久久久久久| 69精品国产乱码久久久| 免费日韩欧美在线观看| videosex国产| 欧美久久黑人一区二区| 一区福利在线观看| 老司机午夜十八禁免费视频| 亚洲精品中文字幕在线视频| 伊人亚洲综合成人网| 久久国产亚洲av麻豆专区| 欧美97在线视频| 女人高潮潮喷娇喘18禁视频| 中国美女看黄片| 亚洲av成人一区二区三| 十八禁网站网址无遮挡| 伦理电影免费视频| videos熟女内射| 这个男人来自地球电影免费观看| 欧美 亚洲 国产 日韩一| 亚洲第一青青草原| 日本猛色少妇xxxxx猛交久久| 每晚都被弄得嗷嗷叫到高潮| 亚洲精品久久久久久婷婷小说| 精品卡一卡二卡四卡免费| 他把我摸到了高潮在线观看 | 国产免费一区二区三区四区乱码| 高清视频免费观看一区二区| 美女扒开内裤让男人捅视频| av在线app专区| 一区二区三区精品91| av电影中文网址| 男女床上黄色一级片免费看| 久久久精品94久久精品| 天天影视国产精品| 亚洲精品久久久久久婷婷小说| 看免费av毛片| 欧美亚洲日本最大视频资源| 久久av网站| 老鸭窝网址在线观看| av一本久久久久| 女性被躁到高潮视频| 波多野结衣av一区二区av| 性高湖久久久久久久久免费观看| 午夜激情久久久久久久| 十八禁高潮呻吟视频| 免费看十八禁软件| 精品少妇久久久久久888优播| 999久久久精品免费观看国产| 亚洲第一青青草原| 日本av免费视频播放| 久久 成人 亚洲| 午夜免费成人在线视频| 51午夜福利影视在线观看| 91大片在线观看| www.自偷自拍.com| 一二三四社区在线视频社区8| 亚洲一码二码三码区别大吗| 亚洲avbb在线观看| 国产精品影院久久| 免费女性裸体啪啪无遮挡网站| 777米奇影视久久| 我要看黄色一级片免费的| 精品视频人人做人人爽| 最近最新免费中文字幕在线| 午夜成年电影在线免费观看| 搡老熟女国产l中国老女人| 97在线人人人人妻| 免费观看av网站的网址| 亚洲综合色网址| 老司机深夜福利视频在线观看 | 成人国语在线视频| 欧美+亚洲+日韩+国产| 日韩大码丰满熟妇| 欧美精品亚洲一区二区| 91av网站免费观看| 欧美乱码精品一区二区三区| 国产精品一区二区在线观看99| 国产日韩一区二区三区精品不卡| 午夜影院在线不卡| 性少妇av在线| 亚洲av成人不卡在线观看播放网 | 国产在线一区二区三区精| 久久性视频一级片| 久久国产精品影院| 久久免费观看电影| 97人妻天天添夜夜摸| 美女脱内裤让男人舔精品视频| 久久青草综合色| 热99久久久久精品小说推荐| 91麻豆av在线| 午夜福利视频在线观看免费| 国产成人免费观看mmmm| 亚洲全国av大片| a在线观看视频网站| 色老头精品视频在线观看| 黄色视频,在线免费观看| 脱女人内裤的视频| 在线精品无人区一区二区三| 人人澡人人妻人| 一区二区日韩欧美中文字幕| 可以免费在线观看a视频的电影网站| 999久久久国产精品视频| 国产精品偷伦视频观看了| 成年人午夜在线观看视频| 大陆偷拍与自拍| 日日摸夜夜添夜夜添小说| 人妻一区二区av| 多毛熟女@视频| 午夜影院在线不卡| 极品人妻少妇av视频| 国产高清视频在线播放一区 | 99久久人妻综合| 后天国语完整版免费观看| 女性被躁到高潮视频| 最近中文字幕2019免费版| 窝窝影院91人妻| 人人妻人人爽人人添夜夜欢视频| 久久综合国产亚洲精品| 深夜精品福利| 国产精品一区二区在线不卡| 宅男免费午夜| 亚洲精品国产色婷婷电影| av欧美777| 可以免费在线观看a视频的电影网站| 乱人伦中国视频| 最黄视频免费看| 热re99久久精品国产66热6| 亚洲一区二区三区欧美精品| 国产精品一区二区精品视频观看| 一级片免费观看大全| 岛国在线观看网站| 亚洲黑人精品在线| 午夜激情av网站| 国产免费福利视频在线观看| www日本在线高清视频| 国产有黄有色有爽视频| 美国免费a级毛片| 91成年电影在线观看| 丁香六月欧美| 久久精品人人爽人人爽视色| 亚洲中文日韩欧美视频| 老司机影院毛片| 一级毛片精品| videosex国产| 精品卡一卡二卡四卡免费| 国产三级黄色录像| 真人做人爱边吃奶动态| 亚洲av电影在线进入| 免费一级毛片在线播放高清视频 | 99精品久久久久人妻精品| 激情视频va一区二区三区| a在线观看视频网站| 亚洲专区国产一区二区| 永久免费av网站大全| 热99国产精品久久久久久7| 91av网站免费观看| 老司机亚洲免费影院| 免费观看人在逋| 搡老熟女国产l中国老女人| 少妇被粗大的猛进出69影院| 两人在一起打扑克的视频| 99久久人妻综合| 一区二区三区精品91| 美国免费a级毛片| 国产亚洲欧美在线一区二区| 亚洲av成人一区二区三| 欧美激情高清一区二区三区| 欧美精品啪啪一区二区三区 | 91精品伊人久久大香线蕉| 亚洲中文av在线| 丁香六月欧美| 国产男女超爽视频在线观看| 无限看片的www在线观看| 中文字幕人妻丝袜制服| 老司机亚洲免费影院| 亚洲精品中文字幕一二三四区 | 后天国语完整版免费观看| 人人澡人人妻人| 80岁老熟妇乱子伦牲交| 91麻豆av在线| 久久久国产欧美日韩av| 欧美日韩视频精品一区| 女性被躁到高潮视频| 色94色欧美一区二区| 多毛熟女@视频| 免费少妇av软件| 国产麻豆69| 欧美日韩亚洲高清精品| 日韩大码丰满熟妇| 欧美国产精品va在线观看不卡| 国产精品自产拍在线观看55亚洲 | 亚洲五月婷婷丁香| 亚洲国产精品成人久久小说| 久久女婷五月综合色啪小说| 色视频在线一区二区三区| 蜜桃在线观看..| 法律面前人人平等表现在哪些方面 | 日本精品一区二区三区蜜桃| √禁漫天堂资源中文www| 欧美精品高潮呻吟av久久| av福利片在线| 国产亚洲一区二区精品| 亚洲一码二码三码区别大吗| 色播在线永久视频| 亚洲国产日韩一区二区| 亚洲欧美成人综合另类久久久| 欧美日韩黄片免| 1024视频免费在线观看| 亚洲欧美一区二区三区久久| 欧美黄色片欧美黄色片| 首页视频小说图片口味搜索| 国产一级毛片在线| 人成视频在线观看免费观看| 久久 成人 亚洲| 久久久欧美国产精品| 国产成人免费观看mmmm| 欧美亚洲日本最大视频资源| 久久99一区二区三区| 免费一级毛片在线播放高清视频 | 久久天堂一区二区三区四区| 淫妇啪啪啪对白视频 | 91大片在线观看| 亚洲成av片中文字幕在线观看| 50天的宝宝边吃奶边哭怎么回事| 国产三级黄色录像| 精品久久久精品久久久| 人成视频在线观看免费观看| 国产亚洲一区二区精品| 午夜老司机福利片| 欧美精品高潮呻吟av久久| 美女扒开内裤让男人捅视频| 亚洲av美国av| 精品高清国产在线一区| 亚洲七黄色美女视频| 黄色视频在线播放观看不卡| 亚洲美女黄色视频免费看| 少妇粗大呻吟视频| 日本撒尿小便嘘嘘汇集6| 午夜福利视频在线观看免费| 国产精品 欧美亚洲| 免费在线观看完整版高清| 亚洲精品av麻豆狂野| 国产精品一区二区在线观看99| 亚洲自偷自拍图片 自拍| 伊人亚洲综合成人网| 韩国高清视频一区二区三区| 人人妻人人澡人人看| 国产精品麻豆人妻色哟哟久久| √禁漫天堂资源中文www| 人妻一区二区av| 日日摸夜夜添夜夜添小说| 亚洲性夜色夜夜综合| 国产成+人综合+亚洲专区| 久久99一区二区三区| 秋霞在线观看毛片| 操出白浆在线播放| 亚洲第一av免费看| 制服人妻中文乱码| 大香蕉久久网| 两性午夜刺激爽爽歪歪视频在线观看 | 一本久久精品| 亚洲激情五月婷婷啪啪| 欧美 日韩 精品 国产| 免费不卡黄色视频| 久热这里只有精品99| 我的亚洲天堂| 最近中文字幕2019免费版| 亚洲午夜精品一区,二区,三区| 免费看十八禁软件| 91字幕亚洲| 欧美一级毛片孕妇| 欧美老熟妇乱子伦牲交| 国产成人精品久久二区二区免费| 满18在线观看网站| 日韩一区二区三区影片| 国产极品粉嫩免费观看在线| 国产日韩欧美在线精品| 久久综合国产亚洲精品| 成在线人永久免费视频| 久久久久国产精品人妻一区二区| 成人黄色视频免费在线看| 新久久久久国产一级毛片| 国产亚洲午夜精品一区二区久久| 亚洲精品美女久久久久99蜜臀| 丰满少妇做爰视频| 视频在线观看一区二区三区| 飞空精品影院首页| 91麻豆av在线| 在线精品无人区一区二区三| av电影中文网址| 视频在线观看一区二区三区| 91字幕亚洲| 人妻久久中文字幕网| 1024香蕉在线观看| 黑人操中国人逼视频| 男女下面插进去视频免费观看| 亚洲中文字幕日韩| 菩萨蛮人人尽说江南好唐韦庄| av欧美777| 脱女人内裤的视频| 51午夜福利影视在线观看| 99热全是精品| 性少妇av在线| 国产不卡av网站在线观看| 91精品伊人久久大香线蕉| cao死你这个sao货| 国产高清视频在线播放一区 | 亚洲av美国av| 午夜福利影视在线免费观看| 18在线观看网站| 99久久人妻综合| 午夜成年电影在线免费观看| 国产三级黄色录像| 欧美日韩精品网址| 一区在线观看完整版| 午夜福利视频在线观看免费| 久久热在线av| 中文字幕另类日韩欧美亚洲嫩草| 午夜两性在线视频| 国产区一区二久久| 久久av网站| 咕卡用的链子| 黄片大片在线免费观看| 男女午夜视频在线观看| 国产高清videossex| 91成年电影在线观看| 国产精品一区二区精品视频观看| av福利片在线| 十八禁高潮呻吟视频| 成人国语在线视频| 日韩 欧美 亚洲 中文字幕| 免费久久久久久久精品成人欧美视频| 80岁老熟妇乱子伦牲交| 青春草亚洲视频在线观看| 99精品欧美一区二区三区四区| 国产精品一区二区免费欧美 | 一级毛片女人18水好多| 大香蕉久久网| 久久精品国产综合久久久| 日韩视频一区二区在线观看| 一本大道久久a久久精品| 在线看a的网站| 老司机在亚洲福利影院| 久久天躁狠狠躁夜夜2o2o| 亚洲成人免费电影在线观看| 麻豆国产av国片精品| 天天躁夜夜躁狠狠躁躁| e午夜精品久久久久久久| 国产一区有黄有色的免费视频| 大片电影免费在线观看免费| 精品人妻在线不人妻| 制服人妻中文乱码| 日韩,欧美,国产一区二区三区| 啦啦啦在线免费观看视频4| 国产精品一区二区在线观看99| 国产亚洲精品一区二区www | 国产在线一区二区三区精| 最新在线观看一区二区三区| 亚洲第一av免费看| 国产片内射在线| 欧美在线黄色| 亚洲avbb在线观看| 亚洲精品一二三| 免费在线观看日本一区| 91av网站免费观看| 老司机靠b影院| 大码成人一级视频| 国产精品国产av在线观看| 人人妻人人添人人爽欧美一区卜| 国产在线视频一区二区| 亚洲av片天天在线观看| 久久国产精品大桥未久av| 精品少妇内射三级| 国产伦人伦偷精品视频| 欧美精品亚洲一区二区| 国产黄频视频在线观看| 欧美日韩亚洲高清精品| 一进一出抽搐动态| 黑人巨大精品欧美一区二区蜜桃| 成人国产av品久久久| 九色亚洲精品在线播放| 最新的欧美精品一区二区| 超碰97精品在线观看| 亚洲人成电影免费在线| 午夜免费成人在线视频| 中文字幕色久视频| 国产免费av片在线观看野外av| 久久久久久久久免费视频了| 女性生殖器流出的白浆| 美国免费a级毛片| cao死你这个sao货| 一区二区三区激情视频| 日本一区二区免费在线视频| 在线观看人妻少妇| 欧美另类亚洲清纯唯美| a级片在线免费高清观看视频| 亚洲精品久久午夜乱码| 久久国产精品大桥未久av| 女性被躁到高潮视频| 国产成人影院久久av| 欧美 日韩 精品 国产| 久久久久国产一级毛片高清牌| 日韩,欧美,国产一区二区三区| 国产精品99久久99久久久不卡| 日韩一卡2卡3卡4卡2021年| 欧美黑人欧美精品刺激| 女性被躁到高潮视频| 久久这里只有精品19| 欧美老熟妇乱子伦牲交| 老司机影院成人| 亚洲国产av新网站| 99re6热这里在线精品视频| 老司机在亚洲福利影院| 女人精品久久久久毛片| 黄色视频不卡| 日韩一区二区三区影片| 老司机福利观看| 成人免费观看视频高清| 国产亚洲一区二区精品| 国产激情久久老熟女| 叶爱在线成人免费视频播放|