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

    計(jì)算機(jī)圖形處理器加速的光學(xué)航空影像正射校正

    2016-12-19 02:35:42全吉成王宏偉
    光學(xué)精密工程 2016年11期
    關(guān)鍵詞:并行算法插值法線程

    全吉成,王 平,王宏偉

    (空軍航空大學(xué) 航空航天情報(bào)系,吉林 長(zhǎng)春 130022)

    ?

    計(jì)算機(jī)圖形處理器加速的光學(xué)航空影像正射校正

    全吉成,王 平*,王宏偉

    (空軍航空大學(xué) 航空航天情報(bào)系,吉林 長(zhǎng)春 130022)

    提出了計(jì)算機(jī)圖形處理器(GPU)加速的光學(xué)航空影像正射校正并行算法,以滿足獲取光學(xué)航空影像對(duì)實(shí)時(shí)性的要求并提高對(duì)海量影像數(shù)據(jù)在CPU上串行正射校正的效率。介紹了光學(xué)影像正射校正算法原理以及正射校正算法的并行化處理。為減少GPU執(zhí)行的計(jì)算負(fù)載,引入“有效像素區(qū)域”概念,設(shè)計(jì)了改進(jìn)的GPU并行校正算法。通過配置選擇以及存儲(chǔ)器訪問優(yōu)化進(jìn)一步提高了算法的執(zhí)行效率。最后,分析了GPU并行算法的精度,并驗(yàn)證了噪聲干擾對(duì)算法的影響。實(shí)驗(yàn)結(jié)果表明,優(yōu)化的改進(jìn)GPU并行算法顯著提高了正射校正的速度,影像大小為5 000×5 000時(shí),加速比最高可達(dá)CPU串行算法的223倍以上。雖然GPU單精度計(jì)算和噪聲干擾會(huì)使影像校正精度有所下降,但尚在誤差允許范圍之內(nèi)。該算法能夠快速實(shí)現(xiàn)光學(xué)航空影像的正射校正,校正后的影像滿足實(shí)際應(yīng)用需要。

    航空影像;正射校正;計(jì)算機(jī)圖形處理器(GPU);并行算法;有效像素區(qū)域

    1 引 言

    在應(yīng)急救災(zāi),軍事偵察以及環(huán)境監(jiān)測(cè)等應(yīng)用領(lǐng)域,航空攝影為及時(shí)、準(zhǔn)確地獲取數(shù)據(jù)信息發(fā)揮著重要的作用。在影像成像過程中,由于受地形起伏,地球曲率等因素的影響,原始影像各像素將產(chǎn)生不同程度的幾何畸變,因此需要對(duì)航空影像進(jìn)行正射校正。隨著航空光學(xué)影像分辨率和數(shù)據(jù)量的增加,針對(duì)單個(gè)像元的正射校正計(jì)算量也急劇增多,導(dǎo)致傳統(tǒng)的串行處理難以滿足影像實(shí)時(shí)處理的要求[1]。

    近年來,計(jì)算機(jī)圖形處理器(Graphics Processing Unit,GPU)以超過摩爾定律的速度在不斷更新,在通用計(jì)算領(lǐng)域已取得了長(zhǎng)足發(fā)展[2]。2007年,NVIDIA公司提出了統(tǒng)一設(shè)備計(jì)算架構(gòu)(Compute Unified Device Architecture, CUDA)[3]。作為一種新的編程模型,CUDA采用了類C語言的軟件開發(fā)環(huán)境。CUDA核心具有單指令多數(shù)據(jù)流的特性,用來加速的算法涉及領(lǐng)域[4-7]非常廣泛,其中包括一些圖形圖像處理算法的加速。例如國(guó)外學(xué)者將圖像處理器應(yīng)用到加速圖像去噪[8]和可視化[9]算法中;國(guó)內(nèi),李仕等利用圖像處理器加速實(shí)現(xiàn)模糊視頻圖像的實(shí)時(shí)恢復(fù)[10]、航空成像的位移補(bǔ)償[11]以及航空斜視成像異速像移的實(shí)時(shí)恢復(fù)[12]等算法。

    目前,不少學(xué)者[13-14]將GPU引入到影像的正射校正中,大大加快了影像處理的速度。文獻(xiàn)[15]利用GPU加速了航拍影像數(shù)據(jù)的正射校正,達(dá)到了實(shí)時(shí)處理。但該方法采用了OpenGL傳統(tǒng)編程方式,實(shí)現(xiàn)比較復(fù)雜。文獻(xiàn)[16]利用CUDA并行處理幾何校正,實(shí)現(xiàn)了UAV拍攝影像數(shù)據(jù)的快速處理,與CPU性能相比,加速比最大可達(dá)12倍以上。文獻(xiàn)[3]使用了“層次性分塊”策略設(shè)計(jì)了CPU和GPU協(xié)同處理的正射校正模式,通過配置優(yōu)化和存儲(chǔ)層次性訪問提高了校正效率,處理時(shí)間壓縮至了5 s以內(nèi)。

    為了快速實(shí)現(xiàn)影像的正射校正,本文采用劃分“有效數(shù)據(jù)區(qū)域”的方法來減少計(jì)算量,并且針對(duì)GPU架構(gòu)特點(diǎn)進(jìn)行性能優(yōu)化,充分利用GPU強(qiáng)大的并行處理和浮點(diǎn)運(yùn)算能力,提高影像幾何校正的效率。

    2 正射校正算法原理及其并行化

    2.1 正射校正算法原理

    根據(jù)成像時(shí)設(shè)備記錄的攝影參數(shù),基于共線方程建立影像坐標(biāo)與地物坐標(biāo)之間的對(duì)應(yīng)關(guān)系,進(jìn)而對(duì)影像進(jìn)行正射校正。其中共線方程分為正解和反解形式[17],分別為:

    (1)

    (2)

    式中:(x,y)為像點(diǎn)在像平面坐標(biāo)系中的坐標(biāo),(X,Y,Z)為對(duì)應(yīng)地面點(diǎn)在物空間坐標(biāo)系中的坐標(biāo),(XS,YS,ZS)為攝站在物空間坐標(biāo)系中的坐標(biāo),(ai,bi,ci)為圖像的外方位元素確定的旋轉(zhuǎn)矩陣的9個(gè)元素(3個(gè)為獨(dú)立元素)。

    以投影中心為原點(diǎn),x,y軸與像平面坐標(biāo)系的x,y軸平行,z軸與主光軸重合,建立像空間坐標(biāo)系,以大地空間直角坐標(biāo)系作為物空間坐標(biāo)系。利用攝影參數(shù)以及大地測(cè)量相關(guān)公式計(jì)算物像坐標(biāo)系之間的旋轉(zhuǎn)矩陣,代入共線方程中。將原始影像4個(gè)角點(diǎn)的坐標(biāo),代入式(2),求得原始影像在物空間坐標(biāo)系中的覆蓋范圍,即糾正后影像的范圍。然后利用正射校正的間接法方案,對(duì)糾正后影像的每個(gè)像素代入式(1)求得對(duì)應(yīng)像點(diǎn)在像空間坐標(biāo)系中坐標(biāo),同時(shí),在原始影像上進(jìn)行灰度值內(nèi)插并賦予糾正影像空間的像素[18]。

    2.2 正射校正算法的并行化

    正射校正算法是對(duì)圖像中單個(gè)像素的操作,物像空間坐標(biāo)的轉(zhuǎn)換以及灰度值內(nèi)插、賦值都是逐個(gè)像素完成的,且各個(gè)像素之間沒有相關(guān)性,計(jì)算量大,因此非常適合映射到GPU中進(jìn)行處理。

    GPU中的線程是按照線程網(wǎng)格和線程塊的形式進(jìn)行組織的。常規(guī)的GPU正射校正并行化一般將整個(gè)糾正后的影像映射到線程網(wǎng)格上,再將糾正后影像邏輯分塊,使每一塊與線程塊的大小一致,保證線程塊的每個(gè)線程對(duì)影像分塊中的每一個(gè)像素進(jìn)行處理[3],如圖1所示。

    圖1 常規(guī)的GPU正射校正并行化映射

    圖2給出了原始影像與糾正后影像的對(duì)應(yīng)關(guān)系,將對(duì)應(yīng)位置在原始影像范圍內(nèi)的糾正影像像素稱為有效像素。在坐標(biāo)轉(zhuǎn)換和重采樣過程中只需計(jì)算有效像素的灰度值,其余像素賦值為零。常規(guī)GPU正射校正并行處理不考慮有效像素,而是將糾正影像的每個(gè)像素都進(jìn)行計(jì)算,故增加了GPU的工作量。

    (a)原始影像

    (b)糾正后影像

    Fig.2 Mapping relationship between original image and corrected image

    本文將考慮糾正后影像的有效像素,將改進(jìn)后并行算法映射到GPU中,減少GPU的負(fù)載,從而降低影像糾正的處理時(shí)間。

    3 改進(jìn)的正射校正算法及其并行化

    3.1 算法描述

    改進(jìn)的正射校正算法通過正射校正直接法計(jì)算出有效像素區(qū)域,即通過計(jì)算圖2中原始影像的4個(gè)角點(diǎn)a,b,c,d,得出糾正影像的點(diǎn)a*,b*,c*,d*,利用4條線段a*b*,b*c*,c*d*,d*a*即可確定有效像素區(qū)域的邊界。通過每條掃描線與4個(gè)線段的交點(diǎn),可以得出有效像素區(qū)域每一行像素的起始點(diǎn)和結(jié)束點(diǎn)。

    圖3 有效像素區(qū)域邊界的求解過程

    3.2 改進(jìn)算法的并行化

    3.2.1 粗粒度并行化分析

    糾正影像的每一行都有各自的有效像素區(qū)域的起始點(diǎn)和結(jié)束點(diǎn),且各行之間不相干。每行之間可以同時(shí)進(jìn)行處理,因此非常適合用線程塊實(shí)現(xiàn)粗粒度并行。如圖4所示,將每一行像素與一維線程塊Block建立映射關(guān)系,由每個(gè)線程塊的索引號(hào)blockIdx.x對(duì)應(yīng)糾正影像的行號(hào)。

    圖4 改進(jìn)的GPU正射校正粗粒度并行化映射

    3.2.2 細(xì)粒度并行化分析

    線程塊中的線程以線程束warp的形式在多處理器上運(yùn)行,各個(gè)線程映射到標(biāo)量處理器核心。如圖5所示,線程塊Block的每一個(gè)線程映射到糾正影像對(duì)應(yīng)行的一個(gè)像素點(diǎn),實(shí)現(xiàn)對(duì)像素的并行處理,實(shí)現(xiàn)線程間的細(xì)粒度并行化。

    圖5 改進(jìn)的GPU正射校正細(xì)粒度并行化映射

    4 改進(jìn)的GPU算法優(yōu)化加速

    4.1 執(zhí)行配置選擇

    GPU進(jìn)行核函數(shù)計(jì)算時(shí),都要對(duì)線程網(wǎng)格和線程塊的大小進(jìn)行配置。線程網(wǎng)格內(nèi)的多個(gè)線程塊以任意順序獨(dú)立執(zhí)行,線程塊中線程以warp為單位在流多處理器調(diào)度運(yùn)行。線程塊網(wǎng)格和線程塊的大小設(shè)置,需考慮以下幾個(gè)限制:①線程塊中線程數(shù)小于1 024。②流多處理器中線程塊數(shù)量小于8。③流式多處理器中線程數(shù)小于1 536。為了有效隱藏延長(zhǎng),應(yīng)至少滿足active warps≥6,active blocks≥2[19]。

    正射校正算法中選擇一維線程塊Block處理一行像素,大小為256=8 warps>6 warps,同時(shí)1 536/256=6>2,這樣一個(gè)流多處理器可以同時(shí)投入6個(gè)線程塊,1 536個(gè)線程,調(diào)度性能最優(yōu),能很好地隱藏計(jì)算延遲。由于線程塊小于糾正影像一行的像素個(gè)數(shù),因此每個(gè)線程通過線程索引threadIdx.x加threadDim.x進(jìn)行循環(huán)迭代計(jì)算,直至一行所有像素計(jì)算完。

    線程網(wǎng)格大小一般由問題的規(guī)模和處理方式?jīng)Q定,不能小于流多處理數(shù)量的2倍。在本文并行算法中,線程網(wǎng)格配置為一維的線程網(wǎng)格,大小為糾正影像的行數(shù),遠(yuǎn)遠(yuǎn)大于流多處理器數(shù)量的2倍。

    4.2 存儲(chǔ)器優(yōu)化

    GPU線程在運(yùn)行過程中可以訪問設(shè)備中多種存儲(chǔ)器,不同存儲(chǔ)空間的訪問延遲周期不同,因此如何改進(jìn)計(jì)算訪問比和各存儲(chǔ)層次的延遲隱藏是存儲(chǔ)優(yōu)化的核心[20]。

    4.2.1 共享存儲(chǔ)器訪問優(yōu)化

    每個(gè)線程塊都對(duì)應(yīng)一個(gè)共享存儲(chǔ)器且對(duì)塊內(nèi)所有線程可見,塊內(nèi)線程可通過共享存儲(chǔ)器進(jìn)行通信。共享內(nèi)存提供了快速的讀寫速度,在無存儲(chǔ)體沖突時(shí)訪存速度在4個(gè)時(shí)鐘周期以內(nèi),而全局存儲(chǔ)器訪問延遲為400~600個(gè)時(shí)鐘周期。所以將全局內(nèi)存中數(shù)據(jù)讀入到共享內(nèi)存中可大大提高核函數(shù)的執(zhí)行速度。

    并行算法將有效區(qū)域邊界數(shù)組ValidPixel[height][2]讀取到共享存儲(chǔ)器,當(dāng)索引號(hào)為blockIdx.x的線程塊中的線程索引與ValidPixel[i][0]、ValidPixel[i][1]比較時(shí),無需重復(fù)訪問全局內(nèi)存,造成效率低下。當(dāng)warp內(nèi)線程同時(shí)對(duì)共享內(nèi)存同一地址ValidPixel[i][0]或ValidPixel[i][1]發(fā)出讀請(qǐng)求時(shí),可被廣播到所有線程中。

    4.2.2 常量存儲(chǔ)器訪問優(yōu)化

    常量存儲(chǔ)器是GPU設(shè)備中的只讀存儲(chǔ)器,存儲(chǔ)空間為64 kB,同時(shí)含有8 kB的cache,訪問速度遠(yuǎn)遠(yuǎn)快于全局內(nèi)存。當(dāng)half-warp都需要相同的讀取請(qǐng)求時(shí),讀取操作會(huì)廣播到half-warp中的所有線程中。此外,線程通過從常量?jī)?nèi)存緩存中讀取數(shù)據(jù),這也加快了訪問速度。

    在正射校正并行算法中,物空間坐標(biāo)系和像空間坐標(biāo)系的變換參數(shù)在每個(gè)像素點(diǎn)進(jìn)行坐標(biāo)變換時(shí)都會(huì)使用到,而且在計(jì)算過程中不會(huì)變化。因此將物、像空間坐標(biāo)系的變換參數(shù)放入只讀常量存儲(chǔ)器中可以提高核函數(shù)的執(zhí)行效率。

    5 實(shí)驗(yàn)結(jié)果與分析

    本文實(shí)驗(yàn)用CPU型號(hào)為AMD Athlon(tm) Ⅱ 640,4核處理器,主頻3.0 GHz,內(nèi)存為10.0 G;GPU為Fermi架構(gòu)的 NVIDIA GTX580,16個(gè)流多處理器,512顆計(jì)算核心,1.5 G的全局存儲(chǔ)器。

    本文從一次拍攝任務(wù)完成后的所有航空光學(xué)影像中隨機(jī)選取10張攝影參數(shù)相同的航空幾何畸變影像(3波段),每張影像經(jīng)過重采樣仿真出大小分別為1 000×1 000、2 000×2 000、3 000×3 000、4 096×4 096、5 000×5 000的航空影像,從而得到不同大小的畸變影像各10張,以確保相同大小的畸變影像糾正后的影像大小完全相同。

    5.1 效率分析

    對(duì)相同大小的10張畸變影像,分別采用基于CPU、常規(guī)GPU并行算法、改進(jìn)GPU并行算法進(jìn)行正射校正處理(重采樣分別采用最近鄰插值法、雙線性插值法、雙三次插值法)。取10次運(yùn)行結(jié)果的平均值作為相同大小影像正射校正的處理時(shí)間,實(shí)驗(yàn)結(jié)果如表1、表2所示。

    表1 在實(shí)驗(yàn)環(huán)境下CPU和常規(guī)GPU算法正射校正所用時(shí)間對(duì)比

    Tab.1 Comparison of orthorectification times between CPU and conventional GPU algorithm in experimental environment

    圖像大小/pixel最近鄰插值法CPU/ms常規(guī)GPU/ms雙線性插值法CPU/ms常規(guī)GPU/ms雙三次插值法CPU/ms常規(guī)GPU/ms1000×1000208513.64265217.01330520.492000×2000792946.77830948.601112363.583000×300019338112.0723331131.9929279160.314096×409641691236.1245465251.7449517267.285000×500056356316.2669831380.1974272393.38

    表2 在實(shí)驗(yàn)環(huán)境下CPU和改進(jìn)GPU算法正射校正所用時(shí)間對(duì)比

    Tab.2 Comparison of orthorectification times between CPU and improved GPU algorithm in experimental environment

    圖像大小/pixel最近鄰插值法CPU/ms改進(jìn)GPU/ms雙線性插值法CPU/ms改進(jìn)GPU/ms雙三次插值法CPU/ms改進(jìn)GPU/ms1000×1000208512.32265215.02330517.732000×2000792941.88830942.611112354.613000×30001933899.9023331114.9129279137.214096×409641691210.1645465218.3049517228.385000×500056356280.1969831328.0274272334.95

    從表1、表2可以明顯看出,使用CPU串行算法進(jìn)行光學(xué)影像正射校正執(zhí)行時(shí)間很長(zhǎng),而使用GPU并行算法(常規(guī)算法和改進(jìn)算法)進(jìn)行光學(xué)影像正射校正執(zhí)行時(shí)間較前者大幅度減小,并且改進(jìn)的GPU算法的減小幅度更大一些。

    為了更好地對(duì)算法進(jìn)行評(píng)價(jià),本文使用時(shí)間加速比來評(píng)價(jià)算法的效果。定義如下[21]:

    (3)

    式中:Ts為表1、2中CPU串行算法進(jìn)行光學(xué)航空影像正射校正的執(zhí)行時(shí)間;Tp為表1、2中GPU并行算法進(jìn)行光學(xué)航空影像正射校正的執(zhí)行時(shí)間(包括CPU和GPU之間數(shù)據(jù)傳輸時(shí)間);Sp為算法時(shí)間加速比。圖6、圖7分別列出了常規(guī)GPU并行算法和改進(jìn)GPU并行算法分別相對(duì)于CPU串行算法正射校正的時(shí)間加速比??梢钥闯觯?/p>

    (1)使用GPU多線程并行算法對(duì)影像進(jìn)行正射校正相對(duì)于CPU串行校正加速效果非常明顯。常規(guī)的GPU并行算法的時(shí)間加速比為152.86~188.80倍,改進(jìn)的GPU并行算法的時(shí)間加速比達(dá)169.24~221.74倍。

    (2)當(dāng)重采樣方法一定時(shí),常規(guī)的GPU并行校正算法和改進(jìn)的GPU并行校正算法的加速比都隨著影像的增大而增大。當(dāng)影像大小為5 000×5 000時(shí),加速比達(dá)到了最大。這是由于GPU在進(jìn)行正射校正時(shí),算法加速比與校正時(shí)間成正比[3],因此光學(xué)航空影像越大,GPU處理時(shí)間越長(zhǎng),時(shí)間加速比也就越大。

    (3)在影像大小一定的情況下,當(dāng)使用雙三次插值法進(jìn)行灰度重采樣時(shí),算法復(fù)雜度最大,GPU并行算法的校正處理時(shí)間最長(zhǎng),故算法加速比也最大。相應(yīng)的,在使用最近鄰插值法時(shí)GPU并行算法的校正時(shí)間最短,算法復(fù)雜度最小,算法加速比也就最小。

    圖6 常規(guī)GPU并行算法相對(duì)CPU串行算法的加速比

    Fig.6 Speed-up ratio of conventional GPU parallel algorithm relative to CPU serial algorithm

    圖7 改進(jìn)的GPU并行算法相對(duì)CPU串行算法的加速比

    Fig.7 Speed-up ratio of improved GPU parallel algorithm relative to CPU serial algorithm

    將改進(jìn)的GPU并行算法與常規(guī)的GPU并行算法進(jìn)行比較,定義改進(jìn)的GPU并行算法的加速比與常規(guī)的GPU并行算法的加速比之間的比值為性能提升比,如表3所示??梢钥闯觯?/p>

    (1)相對(duì)于常規(guī)的GPU并行校正算法,無論采用何種重采樣方式,改進(jìn)的GPU并行校正算法的性能提升比都大于1,執(zhí)行效率也得到了提升,從而驗(yàn)證了改進(jìn)的GPU并行算法的有效性。

    (2)由于Fermi架構(gòu)的GPU內(nèi)含有L1、L2緩存[22],從而提高了訪存速度,另外非有效區(qū)域像素占糾正影像全部像素的比例不大,因此改進(jìn)后的GPU并行算法相對(duì)于常規(guī)的GPU并行算法性能提升比不是很大,最大沒有超過1.2倍。

    (3)在影像大小一定的情況下,在3種重采樣方式中,雙三次插值法的性能提升比最大,最高可達(dá)1.17倍以上。

    表3 改進(jìn)GPU并行算法與常規(guī)GPU并行算法試驗(yàn)結(jié)果比較

    Tab.3 Comparison of experimental results of improved GPU parallel algorithm and conventional GPU algorithm

    圖像大小/pixel性能提升比/times最近鄰插值法雙線性插值法雙三次插值法1000×10001.1071.1331.1562000×20001.1171.1411.1643000×30001.1221.1491.1684096×40961.1241.1531.1705000×50001.1291.1591.174

    將改進(jìn)的GPU并行算法參照第4節(jié)進(jìn)行進(jìn)一步優(yōu)化,實(shí)驗(yàn)結(jié)果如表4所示。

    表4 優(yōu)化的改進(jìn)GPU并行算法正射校正執(zhí)行時(shí)間

    Tab.4 Execution time of optimized GPU-improved parallel algorithm

    圖像大小/pixel執(zhí)行時(shí)間/ms最近鄰插值法雙線性插值法雙三次插值法1000×100011.7614.3617.032000×200040.9341.7753.683000×300099.12113.49136.134096×4096208.56216.17227.215000×5000277.27326.13332.38

    從表4可以看出:無論對(duì)于何種重采樣方式,優(yōu)化的改進(jìn)GPU并行算法的執(zhí)行時(shí)間均進(jìn)一步減小,相對(duì)常規(guī)GPU并行算法的性能提升比也得到了進(jìn)一步提高。

    與CPU串行處理方法相比,優(yōu)化的改進(jìn)GPU并行算法加速比如圖8所示,加速比最大達(dá)到了220倍以上,可以滿足光學(xué)航空影像正射校正的快速處理。

    圖8 優(yōu)化的改進(jìn)GPU算法相對(duì)CPU串行算法的加速比

    Fig.8 Speed-up ratio of optimized GPU-improved parallel algorithm relative to CPU serial algorithm

    5.2 精度分析

    為了驗(yàn)證改進(jìn)GPU并行校正算法的校正精度,將改進(jìn)GPU并行算法糾正的影像與CPU串行算法糾正的影像相減,得到兩者的差值影像,如圖9所示。從圖中可以看出,改進(jìn)GPU并行算法和CPU串行算法糾正后的影像中一些區(qū)域像素之間存在著灰度值差異。這主要由于改進(jìn)GPU并行算法進(jìn)行影像校正時(shí)采用的是單精度浮點(diǎn)計(jì)算,因此造成GPU并行算法校正的影像校正精度稍微下降,但是這些差異在允許的范圍之內(nèi),糾正后的影像仍然能夠滿足實(shí)際應(yīng)用的需要,進(jìn)而驗(yàn)證了改進(jìn)的GPU并行校正算法的有效性。

    圖9 GPU校正影像與CPU校正影像的差值影像

    另外本文采用的Fermi架構(gòu)GPU,具有和CPU一樣的雙精度浮點(diǎn)數(shù)運(yùn)算能力。為了提高正射校正的精度,可以在GPU并行校正算法中使用雙精度浮點(diǎn)運(yùn)算。通過統(tǒng)計(jì)改進(jìn)GPU并行算法糾正的影像與CPU串行算法糾正的影像之間差值影像的灰度直方圖,可以看出兩者的影像灰度值完全一致,但是算法的執(zhí)行效率較單精度下降很多,因?yàn)镚PU的雙精度浮點(diǎn)運(yùn)算速度只是單精度浮點(diǎn)運(yùn)算速度的1/10左右,大大降低了影像校正的速度。

    5.3 噪聲干擾

    為了進(jìn)一步驗(yàn)證噪聲對(duì)改進(jìn)GPU并行算法校正幾何畸變影像的影響,本文將實(shí)驗(yàn)幾何畸變影像加入高斯噪聲,利用改進(jìn)GPU并行校正算法進(jìn)行處理并優(yōu)化,實(shí)驗(yàn)結(jié)果如表5所示。

    通過對(duì)比表4和表5可以看出,改進(jìn)GPU并行算法的執(zhí)行時(shí)間基本沒有變化,噪聲對(duì)改進(jìn)GPU并行校正算法的執(zhí)行效率沒有影響,這是因?yàn)閳D像噪聲并沒有影響幾何畸變影像校正的計(jì)算量,因此利用改進(jìn)GPU并行算法進(jìn)行影像校正的執(zhí)行時(shí)間幾乎不變。

    表5 優(yōu)化的改進(jìn)GPU并行算法對(duì)噪聲畸變影像正射校正的執(zhí)行時(shí)間

    Tab.5 Execution time of noise distortion orthorectification for optimizated GPU-improved parallel algorithm

    圖像大小/pixel執(zhí)行時(shí)間/ms最近鄰插值法雙線性插值法雙三次插值法1000×100011.2514.7316.872000×200041.3842.2753.963000×300098.21114.62137.744096×4096207.92215.64227.735000×5000278.16325.86333.63

    將改進(jìn)GPU并行算法糾正的噪聲影像與CPU串行算法糾正的無噪聲影像做差,得到了兩者的差值影像,可以看出兩者存在著很明顯的像素灰度值差異,相較圖9差值影像的差異更大,說明改進(jìn)GPU并行算法糾正的噪聲影像精度下降很大。這主要由于噪聲嚴(yán)重影響了糾正影像重采樣的計(jì)算結(jié)果,因此在對(duì)幾何畸變影像做正射校正之前需要盡量消除噪聲以及其他畸變的干擾。

    6 結(jié) 論

    本文針對(duì)光學(xué)航空影像獲取信息實(shí)時(shí)性的要求以及海量影像數(shù)據(jù)在CPU上串行正射校正效率慢的問題,探討了利用GPU并行算法進(jìn)行正射校正。首先介紹了正射校正原理,并利用常規(guī)GPU并行算法對(duì)光學(xué)影像進(jìn)行校正。在此基礎(chǔ)上,引入了“有效像素區(qū)域”以減少GPU的計(jì)算量,設(shè)計(jì)了改進(jìn)的GPU并行算法,并且通過配置選擇以及存儲(chǔ)器優(yōu)化進(jìn)一步提高了算法的效率,最后分析了GPU并行算法的精度,并利用噪聲圖像對(duì)GPU并行算法進(jìn)行驗(yàn)證。實(shí)驗(yàn)結(jié)果表明:優(yōu)化后的改進(jìn)GPU并行算法顯著提高了正射校正的速度,相對(duì)于CPU串行算法加速比最大達(dá)223倍以上,基本能夠滿足海量影像數(shù)據(jù)快速校正處理的要求,但是由于算法的單精度計(jì)算導(dǎo)致影像校正精度下降,誤差在允許范圍之內(nèi),糾正后的影像仍然能夠滿足實(shí)際應(yīng)用的需要。

    [1] 蔣艷凰,楊學(xué)軍,易會(huì)戰(zhàn).衛(wèi)星遙感圖像并行幾何校正算法研究[J].計(jì)算機(jī)學(xué)報(bào),2004,27(7):944-951. JIANG Y H, YANG X J, YI H ZH. Parallel algorithm of geometrical correction for satellite images [J].ChineseJournalofComputer, 2004, 27 (7): 994-951. (in Chinese)

    [2] 王海峰,陳慶奎.圖形處理器通用計(jì)算關(guān)鍵技術(shù)研究綜述[J].計(jì)算機(jī)學(xué)報(bào),2013,36(4):757-772. WANG H F, CHEN Q K. General purpose computing of graphics processing unit: a survey [J].ChineseJournalofComputer, 2013,36(4):757-772. (in Chinese)

    [3] 方留楊,王密,李德仁.CPU和GPU協(xié)同處理的光學(xué)衛(wèi)星遙感影像正射校正方法[J].測(cè)繪學(xué)報(bào),2013,42(5):668-675. FANG L Y, WANG M, LI D R. A CPU-GPU co-processing orthographic rectification approach for optical satellite imagery [J].ActaGeodaeticaetCartographicasinica,2013,42(5):668-675. (in Chinese)

    [4] 吳鑫,張建奇,楊琛.Jetson TK1 平臺(tái)實(shí)現(xiàn)快速紅外圖像背景預(yù)測(cè)算法[J].紅外與激光工程,2015,44(9):2615-2621. WU X, ZHANG J Q,YANG CH. Efficient infrared image background prediction with Jetson TK1 [J].InfraredandLaserEngineering, 2015,44(9):2615-2621. (in Chinese)

    [5] 石坤,郝穎明,王明明,等.海面背景紅外實(shí)時(shí)仿真[J]. 紅外與激光工程,2012,41(1):25-29. SHI K, HAO Y M,WANG M M,etal.. Real-time simulation method of infrared sea background [J].InfraredandLaserEngineering, 2012,41(1):25-29. (in Chinese)

    [6] 朱兵,付飛蚺.基于GPU的虛擬內(nèi)窺鏡場(chǎng)景實(shí)時(shí)繪制算法[J].液晶與顯示,2013,28(1):127-131. ZHU B,F(xiàn)U F R. Real-time rendering algorithm of scene in virtual endoscopy based on GPU [J].ChineseJournalofLiquidCrystalsandDisplays,2013,28(1):127-131. (in Chinese)

    [7] 李大禹.基于多GPU的液晶自適應(yīng)光學(xué)波前處理器[J].液晶與顯示,2016,31(5):491-496. LI D Y. Liquid crystal adaptive optics wavefront processor based on multi-GPU [J].ChineseJournalofLiquidCrystalsandDisplays, 2016,31(5):491-496. (in Chinese)

    [8] GRANATA D, AMATO U, ALFANO B. MRI denoising by nonlocal means on multi-GPU [J].JournalofReal-TimeImageProcessing, 2016:1-11.

    [9] HERMANN M, SCHULTZ A C,SCHULTZ T,etal.. Accurate interactive visualization of large deformations and variablility in biomedical image ensembles [J].IEEETransactionsonVisualization&ComputerGraphics, 2016, 22(1): 708-717.

    [10] 王晶,李仕.運(yùn)動(dòng)模糊視頻圖像在圖形處理器平臺(tái)上的實(shí)時(shí)恢復(fù)[J].光學(xué) 精密程,2010,18(10):2262-2268. WANG J,LI SH.Real-time restoration of motion-blurred video images on GPU [J].Opt.PrecisionEng.,2010, 18(10): 2262-2268. (in Chinese)

    [11] 李仕,孫輝,張葆.航空成像像移補(bǔ)償?shù)牟⑿杏?jì)算[J].光學(xué) 精密工程, 2009, 17(1) : 226-230. LI SH, SUN H, ZHANG B. Parallel restoration for motion-blurred aerial image [J].Opt.PrecisionEng., 2009,17(1): 226-230. (in Chinese)

    [12] 李仕,張葆,孫輝.航空斜視成像異速像移的實(shí)時(shí)恢復(fù)[J].光學(xué) 精密工程,2009,17(4):895-900. LI SH, ZHANG B, SUN H. Real-time restoration for aerial side-oblique images with different motion rates [J].Opt.PrecisionEng., 2009, 17(4): 895-900. (in Chinese)

    [13] 楊靖宇,張永生,李正國(guó),等.遙感影像正射糾正的GPU-CPU協(xié)同處理研究[J].武漢大學(xué):信息科學(xué)版,2011,36(9):1043-1046. YANG J Y, ZHANG Y SH, LI ZH G,etal.. GPU-CPU cooperate processing of RS image ortho-rectification [J].GeomaticsandInformationScienceofWuhanUniversity, 2011, 36(9): 1043-1046.(in Chinese)

    [14] JAVIER R S, MARIA C C, JULIO M H. GPU geocorrection for airborne pushbroom imagers [J].IEEETransactionsonGeoscienceandRemoteSensing, 2012,50(11): 4409-4419.

    [15] ULRIKE THOMAS,F(xiàn)RANZ KURZ,RUPERT MULLER,etal.. GPU-Based orthorectification of digital airbororne camera images in real time [J].TheInternationalArchivesofthePhotogrammetry,RemoteSensingandSpatialInformationSciences,2008:589-594.

    [16] LIN J, LI Y CH LI D L. The orthorectified technology for UAV aerial remote sensing image based on the Programmable GPU [J].35thInternationalSymposiumonRemoteSensingofEnvironment,2014:1-6.

    [17] 李德仁,王樹跟,周月琴.攝影測(cè)量與遙感概論[M].第2版.北京:測(cè)繪出版社,2008. LI D R, WANG SH G, ZH Y Q.AnIntroductiontoPhotogrammetryandRemoteSensing[M]. 2nd ed. Beijing: Surveying and Mapping Press,2008. (in Chinese)

    [18] 張祖勛,張劍清.數(shù)字?jǐn)z影測(cè)量學(xué)[M].武漢:武漢測(cè)繪科技大學(xué)出版社,1997:218-219. ZHANG Z X,ZHANG J Q.AnIntroductiontoDigitalPhotogrammetry[M]. Wuhan:Wuhan Surveying and Mapping Science Technology University Press, 1997: 218-219. (in Chinese)

    [19] NVIDIA Corporation. CUDA Programming Guide 1.0,http://www.nvidia.com,2009(8):11.

    [20] 馬安國(guó),成玉,唐遇星,等.GPU異構(gòu)系統(tǒng)中的存儲(chǔ)層次和負(fù)載均衡策略研究[J].國(guó)防科技大學(xué)報(bào),2009,31(5):38-43. MA A G,CHENG Y, TANG Y X,etal.. Research on memory hierarchy and load balance strategy in heterogeneous system based on GPU[J].JournalofNationalUniversityDefenseTechnology, 2009,31(5):38-43. (in Chinese)

    [21] 鄒賢才,李建成,汪海洪,等.OpenMP并行計(jì)算在衛(wèi)星重力數(shù)據(jù)處理中的應(yīng)用[J].測(cè)繪學(xué)報(bào),2010,39(6):636-641. ZOU X C, LI J CH, WANG H H,etal.. Application of parallel computing with OpenMP in data processing for satellite gravity [J].ActaGeodaeticaetCartographicaSinica,2010,39(6):636-641. (in Chinese)

    [22] DAGA MAYANK, ASHWIN M AJI FENG WU-CHUN. On the efficacy of a fused CPU+GPU processor for parallel computer [C].Proceedingofthe2011SymposiumonApplicationAcceleratorsinHigh-PerformanceComputing.Knoxvile,TN,USA, 2011:141-149.

    全吉成(1960-),男,吉林和龍人,博士,教授,博士生導(dǎo)師,1983年、2011年于空軍工程學(xué)院分別獲得學(xué)士、博士學(xué)位,主要從事三維可視化技術(shù)方面的研究。E-mail: jicheng_quan@126.com

    王 平(1991-),男,山東濟(jì)寧人,碩士研究生,2014年于空軍航空大學(xué)獲得學(xué)士學(xué)位,主要從事數(shù)字圖像處理和GPU并行計(jì)算的研究。E-mail: wpsdjnws@163.com

    (版權(quán)所有 未經(jīng)許可 不得轉(zhuǎn)載)

    Orthorectification of optical aerial images by GPU acceleration

    QUAN Ji-cheng, WANG Ping*, WANG Hong-wei

    (Department of Aeronautic and Astronautic Intelligence,AviationUniversityofAirForce,Changchun130022,China)

    An optical aerial image orthorectification parallel algorithm by Graphic Processing Unit(GPU) acceleration was presented to improve the image real-time processing ability and the serial orthorectification efficiency for massive image data on a CPU. The principle of optical image orthorectification algorithm was introduced, and the parallel processing of orthorectification algorithm was described. To reduce the computational load of GPU execution, the concept of “effective pixel region” was introduced and an improved GPU parallel correction algorithm was designed. Then, the efficiency of the algorithm was improved through configuration options and a memory access optimization. Finally, the algorithm precision was analyzed, and impact of the noise on the algorithm was verified. The experimental results show that the optimized improved-GPU parallel algorithm significantly improves the speed of the correction. When the image size is 5 000×5 000, the speed up is 223 times as compared with the CPU serial algorithm. Although the GPU single precision calculation method and noise interference will cause the serious decline of correction precision, it is still in an allowable error range. As a result, the GPU algorithm implements the orthorectification of optical aerial images rapidly and the corrected images satisfy the need of practical applications.

    orthorectification; Graphic Processing Unit(GPU); parallel algorithm; effective pixel region

    2016-06-13;

    2016-08-05.

    吉林省自然科學(xué)基金資助項(xiàng)目(No.20130101069JC);軍內(nèi)武器裝備重點(diǎn)科研計(jì)劃資助項(xiàng)目(No.KJ2012240)

    1004-924X(2016)11-2863-09

    P237

    A

    10.3788/OPE.20162411.2863

    *Correspondingauthor,E-mail:wpsdjnws@163.com

    猜你喜歡
    并行算法插值法線程
    地圖線要素綜合化的簡(jiǎn)遞歸并行算法
    《計(jì)算方法》關(guān)于插值法的教學(xué)方法研討
    淺談linux多線程協(xié)作
    基于GPU的GaBP并行算法研究
    基于二次插值法的布谷鳥搜索算法研究
    Newton插值法在光伏發(fā)電最大功率跟蹤中的應(yīng)用
    基于GPU的分類并行算法的研究與實(shí)現(xiàn)
    Linux線程實(shí)現(xiàn)技術(shù)研究
    無網(wǎng)格局部徑向點(diǎn)插值法求解Helmholtz方程
    么移動(dòng)中間件線程池并發(fā)機(jī)制優(yōu)化改進(jìn)
    a级毛片在线看网站| 久久午夜综合久久蜜桃| 成人高潮视频无遮挡免费网站| 两个人的视频大全免费| 日本精品一区二区三区蜜桃| 色播亚洲综合网| 精品无人区乱码1区二区| 国产aⅴ精品一区二区三区波| 老熟妇乱子伦视频在线观看| av欧美777| 丝袜美腿诱惑在线| 老司机午夜福利在线观看视频| 中文字幕高清在线视频| 国产男靠女视频免费网站| 床上黄色一级片| 两性夫妻黄色片| 国产av一区在线观看免费| 老司机深夜福利视频在线观看| 人人妻,人人澡人人爽秒播| 亚洲18禁久久av| 久久久久国内视频| 国产在线精品亚洲第一网站| 黄色 视频免费看| 免费电影在线观看免费观看| 国产精品久久久久久人妻精品电影| 国产69精品久久久久777片 | 欧美又色又爽又黄视频| 精品国产乱码久久久久久男人| 精品一区二区三区四区五区乱码| 女人爽到高潮嗷嗷叫在线视频| 麻豆久久精品国产亚洲av| 一a级毛片在线观看| 美女午夜性视频免费| 麻豆一二三区av精品| 精品高清国产在线一区| 日韩大尺度精品在线看网址| 老熟妇乱子伦视频在线观看| 亚洲中文字幕一区二区三区有码在线看 | 99热6这里只有精品| 欧美成狂野欧美在线观看| 成年免费大片在线观看| 国产成人aa在线观看| 女人被狂操c到高潮| 欧洲精品卡2卡3卡4卡5卡区| 精品一区二区三区av网在线观看| 国产激情欧美一区二区| 757午夜福利合集在线观看| 亚洲国产高清在线一区二区三| 久99久视频精品免费| 午夜激情福利司机影院| 欧洲精品卡2卡3卡4卡5卡区| 中文字幕高清在线视频| 手机成人av网站| 麻豆久久精品国产亚洲av| 亚洲欧美日韩无卡精品| 天堂动漫精品| 最近视频中文字幕2019在线8| 母亲3免费完整高清在线观看| 一个人观看的视频www高清免费观看 | 97碰自拍视频| av有码第一页| 夜夜夜夜夜久久久久| 欧美高清成人免费视频www| 国内少妇人妻偷人精品xxx网站 | 99久久国产精品久久久| 国产午夜精品论理片| 亚洲自拍偷在线| 免费看美女性在线毛片视频| 99热这里只有精品一区 | 久久久久久人人人人人| 我要搜黄色片| 亚洲国产精品久久男人天堂| 我的老师免费观看完整版| www.999成人在线观看| 亚洲欧美日韩高清专用| 欧美乱码精品一区二区三区| 国产精品香港三级国产av潘金莲| 亚洲黑人精品在线| 高清在线国产一区| 亚洲一区二区三区色噜噜| 天堂√8在线中文| 午夜精品久久久久久毛片777| 黄色毛片三级朝国网站| 99国产极品粉嫩在线观看| 亚洲性夜色夜夜综合| 欧美丝袜亚洲另类 | 亚洲专区中文字幕在线| 免费观看人在逋| 成人国语在线视频| 欧美激情久久久久久爽电影| 亚洲成人国产一区在线观看| 成人av在线播放网站| 麻豆一二三区av精品| av中文乱码字幕在线| 91av网站免费观看| www.精华液| 一二三四在线观看免费中文在| 他把我摸到了高潮在线观看| 在线观看免费午夜福利视频| 老司机深夜福利视频在线观看| 一a级毛片在线观看| 中国美女看黄片| 91大片在线观看| 人妻夜夜爽99麻豆av| 亚洲精品国产精品久久久不卡| bbb黄色大片| 女人爽到高潮嗷嗷叫在线视频| 欧美在线黄色| 成人手机av| 日日夜夜操网爽| 国产三级中文精品| 人成视频在线观看免费观看| 国产在线精品亚洲第一网站| 亚洲乱码一区二区免费版| 亚洲七黄色美女视频| 精品久久久久久成人av| 亚洲成人久久性| 九九热线精品视视频播放| 欧美中文日本在线观看视频| 麻豆成人av在线观看| 日韩欧美三级三区| 中文字幕精品亚洲无线码一区| 国产精品亚洲美女久久久| 在线观看一区二区三区| 亚洲欧美日韩高清专用| 欧美 亚洲 国产 日韩一| av欧美777| 在线观看免费日韩欧美大片| 亚洲国产日韩欧美精品在线观看 | 国产男靠女视频免费网站| а√天堂www在线а√下载| 亚洲国产中文字幕在线视频| 12—13女人毛片做爰片一| 午夜福利免费观看在线| 国内久久婷婷六月综合欲色啪| 中文亚洲av片在线观看爽| 九色成人免费人妻av| 国产日本99.免费观看| 欧美精品啪啪一区二区三区| 一二三四社区在线视频社区8| 三级毛片av免费| 无限看片的www在线观看| 男人舔女人下体高潮全视频| 成人av在线播放网站| 国产伦一二天堂av在线观看| 精品乱码久久久久久99久播| 成人国语在线视频| 不卡一级毛片| 国产成人精品久久二区二区91| 国产高清videossex| 亚洲成人久久性| 久久人妻福利社区极品人妻图片| 九色国产91popny在线| 老熟妇仑乱视频hdxx| 久久中文字幕人妻熟女| 最新美女视频免费是黄的| 日本免费一区二区三区高清不卡| 母亲3免费完整高清在线观看| 怎么达到女性高潮| av欧美777| 1024手机看黄色片| 又黄又爽又免费观看的视频| 亚洲熟妇熟女久久| www.精华液| 悠悠久久av| 欧美日韩亚洲综合一区二区三区_| 欧美中文综合在线视频| 亚洲 国产 在线| 亚洲精品中文字幕一二三四区| 最新美女视频免费是黄的| 亚洲av成人不卡在线观看播放网| 麻豆av在线久日| 啦啦啦韩国在线观看视频| 亚洲成人精品中文字幕电影| 1024视频免费在线观看| 丁香六月欧美| 99久久精品热视频| 老司机午夜十八禁免费视频| 亚洲欧美精品综合一区二区三区| 搡老岳熟女国产| a级毛片在线看网站| 午夜激情福利司机影院| 亚洲av成人av| 国产精品,欧美在线| 中文字幕熟女人妻在线| 久久性视频一级片| √禁漫天堂资源中文www| 国产99白浆流出| 亚洲狠狠婷婷综合久久图片| 精品第一国产精品| 天天躁狠狠躁夜夜躁狠狠躁| 亚洲欧美一区二区三区黑人| 麻豆av在线久日| 色综合亚洲欧美另类图片| 在线观看美女被高潮喷水网站 | 最近视频中文字幕2019在线8| 在线观看日韩欧美| 可以在线观看毛片的网站| 色哟哟哟哟哟哟| 女生性感内裤真人,穿戴方法视频| 国产av又大| 人妻丰满熟妇av一区二区三区| 国产一区二区三区视频了| 黄色视频不卡| 国产成人av教育| 男男h啪啪无遮挡| 女人高潮潮喷娇喘18禁视频| 听说在线观看完整版免费高清| xxxwww97欧美| 欧美乱色亚洲激情| 村上凉子中文字幕在线| 亚洲精品中文字幕一二三四区| 老司机福利观看| 国产精品一区二区三区四区免费观看 | 99re在线观看精品视频| 精品免费久久久久久久清纯| 一卡2卡三卡四卡精品乱码亚洲| 熟妇人妻久久中文字幕3abv| 美女黄网站色视频| 欧美另类亚洲清纯唯美| 村上凉子中文字幕在线| 国产麻豆成人av免费视频| 曰老女人黄片| 国产精品日韩av在线免费观看| 国产精品久久电影中文字幕| 国产精品影院久久| 好男人在线观看高清免费视频| 岛国在线观看网站| 国产精品精品国产色婷婷| 中文字幕最新亚洲高清| 国产真人三级小视频在线观看| 亚洲色图 男人天堂 中文字幕| 又黄又粗又硬又大视频| av国产免费在线观看| 18禁裸乳无遮挡免费网站照片| 国内精品久久久久精免费| 在线a可以看的网站| 国产激情久久老熟女| 国产伦在线观看视频一区| 国产一区二区在线观看日韩 | 亚洲成av人片免费观看| 美女扒开内裤让男人捅视频| 日本熟妇午夜| 国产免费av片在线观看野外av| 色精品久久人妻99蜜桃| 丰满人妻熟妇乱又伦精品不卡| 女同久久另类99精品国产91| 他把我摸到了高潮在线观看| 午夜福利视频1000在线观看| 夜夜躁狠狠躁天天躁| 久久久久亚洲av毛片大全| 国产精品自产拍在线观看55亚洲| 久久精品国产99精品国产亚洲性色| 婷婷亚洲欧美| 国产精品av视频在线免费观看| 高清毛片免费观看视频网站| 久久久精品大字幕| 少妇粗大呻吟视频| 亚洲狠狠婷婷综合久久图片| 久久草成人影院| 欧美日韩瑟瑟在线播放| 亚洲av美国av| 欧美色欧美亚洲另类二区| 777久久人妻少妇嫩草av网站| 香蕉av资源在线| 精品无人区乱码1区二区| 在线观看www视频免费| 久久精品国产亚洲av高清一级| 日日爽夜夜爽网站| 欧美黄色片欧美黄色片| www.熟女人妻精品国产| 男人舔女人下体高潮全视频| 欧美丝袜亚洲另类 | 两个人免费观看高清视频| 三级毛片av免费| 亚洲国产高清在线一区二区三| av视频在线观看入口| 亚洲国产欧美网| 亚洲av成人av| 精品国产乱子伦一区二区三区| 草草在线视频免费看| 欧美3d第一页| 亚洲乱码一区二区免费版| 一级作爱视频免费观看| 两个人的视频大全免费| 高清毛片免费观看视频网站| 精品不卡国产一区二区三区| 啪啪无遮挡十八禁网站| 亚洲av五月六月丁香网| 欧美zozozo另类| 国产亚洲精品综合一区在线观看 | 香蕉久久夜色| 国产精品av久久久久免费| 窝窝影院91人妻| 岛国视频午夜一区免费看| 蜜桃久久精品国产亚洲av| 女人高潮潮喷娇喘18禁视频| 免费在线观看亚洲国产| 一区二区三区高清视频在线| 欧美人与性动交α欧美精品济南到| 亚洲 欧美一区二区三区| 国产成人系列免费观看| 久久精品国产99精品国产亚洲性色| 成人国产一区最新在线观看| 香蕉国产在线看| 小说图片视频综合网站| 亚洲精品久久国产高清桃花| 一二三四在线观看免费中文在| 亚洲av第一区精品v没综合| 成人欧美大片| 久久久久久国产a免费观看| 欧美一级a爱片免费观看看 | 搡老熟女国产l中国老女人| 久久久水蜜桃国产精品网| 国产又黄又爽又无遮挡在线| 午夜a级毛片| www国产在线视频色| 日本五十路高清| 一进一出抽搐动态| 又粗又爽又猛毛片免费看| 后天国语完整版免费观看| 小说图片视频综合网站| 久久久久九九精品影院| 黄色成人免费大全| a级毛片在线看网站| 欧美激情久久久久久爽电影| 亚洲av第一区精品v没综合| 久久久久免费精品人妻一区二区| 亚洲九九香蕉| 欧美日本视频| 丝袜美腿诱惑在线| 国产精品亚洲美女久久久| 中文字幕熟女人妻在线| 女人爽到高潮嗷嗷叫在线视频| 亚洲无线在线观看| 久久久久久久久中文| 国产黄a三级三级三级人| 亚洲欧美日韩高清在线视频| 亚洲熟妇熟女久久| 亚洲av中文字字幕乱码综合| 久久香蕉国产精品| 亚洲在线自拍视频| 18禁国产床啪视频网站| 级片在线观看| 日韩欧美精品v在线| 欧美黑人巨大hd| 淫秽高清视频在线观看| 在线观看午夜福利视频| 欧美最黄视频在线播放免费| 丁香欧美五月| 午夜福利视频1000在线观看| 国产成+人综合+亚洲专区| 99热只有精品国产| 亚洲欧美激情综合另类| av福利片在线观看| 久久久久久国产a免费观看| 99热这里只有精品一区 | 国产1区2区3区精品| av欧美777| 一进一出抽搐动态| 国产精品1区2区在线观看.| 国产午夜精品久久久久久| 亚洲男人的天堂狠狠| 亚洲精品在线观看二区| 一本一本综合久久| 色老头精品视频在线观看| 亚洲无线在线观看| 在线观看舔阴道视频| 久久久精品欧美日韩精品| 波多野结衣巨乳人妻| videosex国产| 9191精品国产免费久久| 俄罗斯特黄特色一大片| 一本综合久久免费| 久久久久久久精品吃奶| 麻豆成人av在线观看| 一边摸一边抽搐一进一小说| 精品国产亚洲在线| 国产熟女午夜一区二区三区| 18禁观看日本| 99久久精品国产亚洲精品| av中文乱码字幕在线| 亚洲免费av在线视频| 波多野结衣高清无吗| 免费在线观看成人毛片| 日韩成人在线观看一区二区三区| 99久久综合精品五月天人人| 少妇被粗大的猛进出69影院| 淫妇啪啪啪对白视频| 中文字幕人妻丝袜一区二区| 99精品在免费线老司机午夜| 嫩草影院精品99| 亚洲国产日韩欧美精品在线观看 | 国产三级黄色录像| 在线国产一区二区在线| 精品免费久久久久久久清纯| 国语自产精品视频在线第100页| 中国美女看黄片| 欧美成人午夜精品| 香蕉国产在线看| 久久久久久人人人人人| 禁无遮挡网站| 手机成人av网站| 久久精品国产综合久久久| 国产高清视频在线播放一区| 欧美一区二区国产精品久久精品 | 男人的好看免费观看在线视频 | 999精品在线视频| 99国产精品一区二区蜜桃av| 一区二区三区国产精品乱码| 99久久精品热视频| 热99re8久久精品国产| 久久精品亚洲精品国产色婷小说| 久久久久久久久久黄片| 欧美成人免费av一区二区三区| 午夜福利高清视频| 这个男人来自地球电影免费观看| 欧美日韩国产亚洲二区| 亚洲精品粉嫩美女一区| 午夜成年电影在线免费观看| 女人高潮潮喷娇喘18禁视频| 午夜精品久久久久久毛片777| 在线a可以看的网站| 18禁美女被吸乳视频| 看黄色毛片网站| 日韩欧美三级三区| 国产久久久一区二区三区| 欧美日本亚洲视频在线播放| 免费在线观看黄色视频的| 动漫黄色视频在线观看| 免费人成视频x8x8入口观看| 一本大道久久a久久精品| a级毛片a级免费在线| 日日夜夜操网爽| 亚洲av成人av| 亚洲 欧美一区二区三区| 麻豆国产97在线/欧美 | 男人舔奶头视频| 在线观看日韩欧美| 一级黄色大片毛片| 观看免费一级毛片| 天堂影院成人在线观看| 欧美在线黄色| 国产视频内射| 久久婷婷成人综合色麻豆| 久久精品91无色码中文字幕| 亚洲午夜理论影院| 黄片小视频在线播放| 欧美中文综合在线视频| 国产野战对白在线观看| 国产成人aa在线观看| 亚洲中文字幕一区二区三区有码在线看 | 9191精品国产免费久久| 免费在线观看完整版高清| svipshipincom国产片| 97人妻精品一区二区三区麻豆| 免费在线观看成人毛片| 大型黄色视频在线免费观看| a在线观看视频网站| 在线观看66精品国产| 免费搜索国产男女视频| 欧美 亚洲 国产 日韩一| 91av网站免费观看| 欧美日韩国产亚洲二区| 亚洲性夜色夜夜综合| 男男h啪啪无遮挡| 午夜福利免费观看在线| 欧美精品啪啪一区二区三区| 久久性视频一级片| 国产av在哪里看| 欧美另类亚洲清纯唯美| 一二三四社区在线视频社区8| 欧美日本视频| 麻豆av在线久日| 国产高清视频在线播放一区| 国产v大片淫在线免费观看| 亚洲激情在线av| 亚洲 欧美一区二区三区| svipshipincom国产片| 久久久久久久久久黄片| 亚洲国产日韩欧美精品在线观看 | 我要搜黄色片| 国产精品国产高清国产av| 嫁个100分男人电影在线观看| 好男人电影高清在线观看| 特大巨黑吊av在线直播| 亚洲熟妇中文字幕五十中出| 搞女人的毛片| a在线观看视频网站| 欧美+亚洲+日韩+国产| 亚洲av电影不卡..在线观看| 国产成年人精品一区二区| 亚洲国产日韩欧美精品在线观看 | 嫩草影院精品99| 久久久久免费精品人妻一区二区| 精品乱码久久久久久99久播| 亚洲av电影在线进入| 可以在线观看毛片的网站| 日本精品一区二区三区蜜桃| 制服人妻中文乱码| 少妇被粗大的猛进出69影院| 欧洲精品卡2卡3卡4卡5卡区| 日本 av在线| 国产激情久久老熟女| 日韩精品免费视频一区二区三区| 人妻久久中文字幕网| 国产一区二区在线观看日韩 | 国产熟女xx| 久久中文字幕人妻熟女| 亚洲午夜理论影院| 嫩草影院精品99| 香蕉久久夜色| 午夜亚洲福利在线播放| 老司机深夜福利视频在线观看| 久久人妻福利社区极品人妻图片| 女人爽到高潮嗷嗷叫在线视频| 叶爱在线成人免费视频播放| 国产伦在线观看视频一区| 久久人妻福利社区极品人妻图片| 亚洲自偷自拍图片 自拍| 91av网站免费观看| 亚洲欧美日韩东京热| 亚洲一码二码三码区别大吗| 曰老女人黄片| 国产精品久久视频播放| 国产亚洲av高清不卡| 日本一二三区视频观看| 老汉色∧v一级毛片| av福利片在线| 男人舔奶头视频| 国产区一区二久久| 国产精品久久久av美女十八| www.精华液| 十八禁网站免费在线| 人人妻人人看人人澡| 日韩欧美免费精品| 欧美日韩瑟瑟在线播放| 国产精品久久久久久人妻精品电影| 日韩国内少妇激情av| 亚洲国产精品成人综合色| 欧美精品亚洲一区二区| 午夜福利在线观看吧| 熟女电影av网| 久久九九热精品免费| 免费在线观看影片大全网站| 一进一出抽搐动态| 女人高潮潮喷娇喘18禁视频| 国产单亲对白刺激| 亚洲片人在线观看| 丰满人妻熟妇乱又伦精品不卡| 国产精品免费一区二区三区在线| 亚洲在线自拍视频| 黄色视频不卡| 亚洲成av人片在线播放无| 国产黄色小视频在线观看| 亚洲av五月六月丁香网| 日韩大码丰满熟妇| 色哟哟哟哟哟哟| 精品无人区乱码1区二区| 无遮挡黄片免费观看| 成人三级做爰电影| 黄色女人牲交| 亚洲人成网站高清观看| 欧美日韩中文字幕国产精品一区二区三区| 欧美日韩国产亚洲二区| 人成视频在线观看免费观看| 欧美日韩黄片免| 最新在线观看一区二区三区| 久久这里只有精品中国| 黑人欧美特级aaaaaa片| 九色成人免费人妻av| 欧美日韩瑟瑟在线播放| 狂野欧美白嫩少妇大欣赏| 久久天堂一区二区三区四区| 国产精品亚洲美女久久久| 欧美色视频一区免费| 亚洲片人在线观看| 久久精品国产综合久久久| 青草久久国产| 午夜a级毛片| 亚洲午夜理论影院| 亚洲成av人片免费观看| avwww免费| 国产亚洲精品综合一区在线观看 | 99国产极品粉嫩在线观看| 最近最新中文字幕大全电影3| 午夜亚洲福利在线播放| 久久久久久九九精品二区国产 | 欧美性猛交黑人性爽| 全区人妻精品视频| 欧美成人免费av一区二区三区| 色综合站精品国产| 日韩中文字幕欧美一区二区| 人妻久久中文字幕网| 一个人免费在线观看的高清视频| 在线观看www视频免费| 99热这里只有精品一区 | 欧美成人一区二区免费高清观看 | 亚洲精品中文字幕一二三四区| 日韩免费av在线播放| 天堂影院成人在线观看| 日本a在线网址| 天天躁夜夜躁狠狠躁躁| 九色国产91popny在线| 69av精品久久久久久| 亚洲精品一区av在线观看| 国产探花在线观看一区二区| 中国美女看黄片| 美女大奶头视频| 看黄色毛片网站| 制服诱惑二区| 欧美黑人精品巨大| 欧美乱色亚洲激情| bbb黄色大片| 国产成年人精品一区二区|