《電子技術應用》
您所在的位置:首頁 > 嵌入式技術 > 設計應用 > CUDA技術及其在數字圖像拼接中的應用
CUDA技術及其在數字圖像拼接中的應用
來源:微型機與應用2013年第2期
王亮亮, 趙曙光
(東華大學 信息科學與技術學院, 上海 201620)
摘要: 將CUDA技術應用于數字圖像拼接領域,闡述了圖像拼接的基本理論及其關鍵技術、多分辨率圖像融合的關鍵算法以及CUDA技術的基本原理和開發方法,并編寫了軟件以實現圖像快速拼接。采用對于尺度具有魯棒性的SIFT 算法進行特征點的提取與匹配,使用穩健的RANSAC 算法求出圖像間變換矩陣的值,并將圖像映射到拼接平面,最后使用基于CUDA的SIFT算法實現了圖像的無縫拼接。該方法提高了圖像拼接的效率,克服了傳統圖像拼接方法因計算量大而等待時間長的缺點。實驗結果表明,CUDA在數字圖像處理的實際應用中卓有成效,有廣闊的應用前景。
Abstract:
Key words :

摘  要:CUDA技術應用于數字圖像拼接領域,闡述了圖像拼接的基本理論及其關鍵技術、多分辨率圖像融合的關鍵算法以及CUDA技術的基本原理和開發方法,并編寫了軟件以實現圖像快速拼接。采用對于尺度具有魯棒性的SIFT 算法進行特征點的提取與匹配,使用穩健的RANSAC 算法求出圖像間變換矩陣的值,并將圖像映射到拼接平面,最后使用基于CUDA的SIFT算法實現了圖像的無縫拼接。該方法提高了圖像拼接的效率,克服了傳統圖像拼接方法因計算量大而等待時間長的缺點。實驗結果表明,CUDA在數字圖像處理的實際應用中卓有成效,有廣闊的應用前景。
關鍵詞: CUDA; 圖像拼接; SIFT; 多分辨率融合

    圖像拼接是計算機視覺領域的一個重要分支,它是一種將多幅相關的部分重疊圖像進行無縫拼接從而獲得寬視角圖像的技術。利用計算機進行匹配,將多幅具有重疊關系的圖像拼合成為一幅具有更大視野范圍的圖像,這就是圖像拼接的目的。CUDA是英偉達(NVIDIA)公司傾力開發和推廣的并行計算架構,該架構通過利用圖形處理器(GPU)的處理能力,能夠大幅提升計算性能。隨著微軟Windows 7與蘋果Snow Leopard操作系統的問世,GPU計算必將成為主流。本文基于SIFT算法,使用最新的CUDA并行計算技術重編算法并編制軟件,不僅可以克服傳統圖像拼接技術中的局限性(如光照、尺度變化的影響等),實現光照和尺度變化條件下的多視角無縫圖像拼接,而且將提高圖像拼接的速度和效率。
1 數字圖像拼接的基本理論和方法
    圖像拼接的基本流程如圖1所示,主要分為圖像預處理、圖像配準和圖像融合與邊界平滑3個步驟。圖像預處理主要指對圖像進行幾何畸變校正和噪聲點的抑制等,使參考圖像和待拼接圖像不存在明顯的幾何畸變。圖像預處理主要是為下一步圖像配準做準備,使圖像質量能夠滿足圖像配準的要求。圖像配準主要指對參考圖像和待拼接圖像中的匹配信息進行提取,在提取出的信息中尋找最佳的匹配,完成圖像間的對齊。圖像拼接的成功與否主要是圖像的配準。待拼接的圖像之間可能存在平移、旋轉和縮放等多種變換或者大面積的同色區域等很難匹配的情況,一個好的圖像配準算法應該能夠在各種情況下準確找到圖像間的對應信息,將圖像對齊。圖像融合指在完成圖像匹配以后對圖像進行縫合,并對縫合的邊界進行平滑處理,使縫合自然過渡。由于任何兩幅相鄰圖像在采集條件上都不可能做到完全相同,因此,對于一些本應該相同的圖像特性(如圖像的光照特性等),在兩幅圖像中就不會表現得完全一樣。圖像拼接縫隙就是從一幅圖像的圖像區域過渡到另一幅圖像的圖像區域時,由于圖像中的某些相關特性發生了躍變而產生的。圖像融合就是使圖像間的拼接縫隙不明顯,拼接更自然。

2 CUDA技術的基本原理和開發方法
    圖像處理的本質是大規模矩陣運算,特別適合并行處理。但CPU通用計算很難利用該特性。與此相反,GPU在并行數據運算上具有強大的計算能力,特別適合作運算符相同而運算數據不同的運算,當執行具有高運算密度的多數據元素時,內存訪問的延遲可以被忽略。
    CUDA編程模型將CPU作為主機(Host),GPU作為協處理器(Coprocessor)或設備(Device),一個系統中可以存在多個設備。在這個模型中,CPU與GPU共同工作,CPU負責邏輯性強的事務處理和串行計算,GPU則專注于執行高度線程化的并行處理任務。圖 2顯示了CUDA的流程架構。

   CUDA對內存的操作與一般的C程序基本相同,操作顯存則需要調用CUDA API中的存儲器管理函數。一旦確定好程序中的并行部分,就可以將這部分計算交給GPU。運行在GPU上的CUDA并行計算函數稱為Kernel,即內核函數,它并不是一個完整的程序,而是CUDA程序中可以被并行執行的步驟。內核函數必須通過_global_函數類型限定符定義,且只能在主機端代碼中調用。
    CUDA線程結構如圖3所示。Kernel以線程網格(Grid)為組織形式,每個Grid由若干個線程塊(Block)組成,每個Block又由若干個線程(Thread)組成。在程序運行過程中,Kernel是以Block為單位執行的,Grid只是一系列可以被執行的Block的集合。不同Block是并行執行的,沒有執行的先后順序,而且相互無法通信。

    CUDA軟件體系由CUDA Library、CUDA Runtime API和CUDA Driver API構成,如圖 4所示。CUDA的核心是CUDA C語言,需要通過nvcc編譯器進行編譯。編譯得到的僅是GPU端的代碼,要在GPU上分配顯存并啟動Kernel就需要借助CUDA Runtime API或者CUDA Driver API來實現。CUDA Runtime API和CUDA Driver API 提供了實現設備管理、上下文管理、存儲器管理、代碼塊管理和執行控制等操作的應用程序接口。CUDA Runtime API在CUDA Driver API的基礎上進行了封裝,使得編程方便代碼簡潔, 因此通常采用CUDA Runtime API進行項目開發。

 

 

 3 SIFT特征匹配算法
    SIFT算法首先在尺度空間進行特征檢測,并確定關鍵點(Key Points)的位置和關鍵點所處的尺度,然后使用關鍵點鄰域梯度的主方向作為該點的方向特征,以實現算子對尺度和方向的無關性。
    SIFT特征匹配算法包括兩個階段:(1)SIFT特征的生成,即從多幅待匹配圖像中提取出對尺度縮放、旋轉和亮度變化無關的特征向量;(2)SIFT特征向量的匹配。
    在實際的尺度不變特征點提取中,SIFT算法將圖像金字塔引入了尺度空間。首先采用不同尺度因子的高斯核對圖像進行卷積以得到圖像的不同尺度空間,將這一組圖像作為金字塔圖像的第一階。接著對其中的2倍尺度圖像(相對于該階第一幅圖像的2倍尺度)以2倍像素距離進行下采樣來得到金字塔圖像第二階的第一幅圖像,對該圖像采用不同尺度因子的高斯核進行卷積,以獲得金字塔圖像第二階的一組圖像。以此類推,獲得高斯金字塔圖像。每一階相鄰的高斯圖像相減,就得到了高斯差分圖像,即DoG圖像。通過擬和三維二次函數以精確確定關鍵點的位置和尺度,同時去除低對比度的關鍵點和不穩定的邊緣響應點(因為DoG算子會產生較強的邊緣響應),以增強匹配穩定性、提高抗噪聲能力。利用關鍵點鄰域像素的梯度方向分布特性為每個關鍵點指定方向參數,使算子具備旋轉不變性,生成SIFT特征向量。
    接下來以關鍵點為中心取8×8的窗口。圖5(a)的中央黑點為當前關鍵點的位置,每個小格代表關鍵點鄰域所在尺度空間的一個像素,箭頭方向代表該像素的梯度方向,箭頭長度代表梯度模值,圖中圈代表高斯加權的范圍(越靠近關鍵點的像素,其梯度方向信息貢獻越大)。然后在每4×4的小塊上計算8個方向的梯度方向直方圖,繪制每個梯度方向的累加值,即可形成一個種子點,如圖5(b)所示。圖5(b)中,一個關鍵點由2×2共4個種子點組成,每個種子點有8個方向向量信息。這種鄰域方向性信息聯合的思想增強了算法抗噪聲的能力,同時對于含有定位誤差的特征匹配也提供了較好的容錯性。

4 基于CUDA的圖像拼接軟件的設計
4.1 Host端實現

    程序的Host端由C++編寫,負責控制整個程序的執行流程、所有供CPU 和GPU 所用的數據的分配管理以及Device 端模塊的調用。界面使用最基本的Windows SDK編寫。
   在數據初始化階段,包含所有之后處理步驟中所需要的圖像數據對象的生成,將輸入圖像作為高斯金字塔底層,通過系統中的PYRAMID_LEVEL宏指定金字塔的層數,在輸入圖像的尺寸基礎上循環計算各層金字塔圖像的分辨率,并對圖像進行初始化。由于所有的圖像數據需要在設備端處理,使用cudaMallocPitch函數分配數據地址空間,數據結構不再是OpenCV中的IplImage,而是GPU可以識別的uchar數組類型。
    使用cudaMemcpy2D函數將IplImage結構中的原始數據復制到相應高斯金字塔的最底層,也就是uchar數組的第一個元素,供Device端函數使用。隨后進行Kernel函數調用,對于每一個需要處理的金字塔層,Host端發起一次Kernel調用。例如:
    reduce<<<(int)ceil((float)(imageSize[l]/THREAD_NUM)), HREAD_NUM>>>(lGaussianData[l+1],rGaussianData[l+1],lGaussianData[l], rGaussianData[l], lLaplacianData[l], rLaplacianData[l],    stride[l+1], stride[l], width[l+1], height[l+1], width[l], height[l]);
4.2 Device端實現
   主要Device函數如下:
    (1) reduce()函數對左、右圖和掩碼圖像各完成一次reduce操作,生成下一層高斯金字塔圖像。reduce變換按照前文所述的方法對目標層的金字塔圖像進行逐像素處理,每一個目標像素的顏色值按一定的權重值對原始圖像中的一個5×5子塊進行計算求得。
       _global_ void reduce(uchar* lGaussianDataSrc, uchar* rGaussianDataSrc, uchar* mGaussianDataSrc,uchar* lGaussianDataDst, uchar* rGaussianDataDst, uchar* mGaussianDataDst, size_t strideSrc, size_t strideDst, int srcWidth, int srcHeight, int gauWidth, int gauHeight);
       (2) expand_and_minus()函數對左、右圖像各完成一次expand操作和減法操作,生成下一層拉普拉斯金字塔圖像。expand變換相當于reduce變換的逆過程,它對目標層的金字塔圖像進行逐像素處理,每一個目標像素的顏色值也是按reduce變換中所使用的權重值對原始圖像中的一個5×5子塊進行計算求得的。
    _global_ void expand_and_minus(uchar*lGaussianDataSrcH,uchar* rGaussianDataSrcH,uchar*lGaussianDataSrcL,uchar*rGau-
ssianDataSrcL,uchar* lLaplacianDataDst,uchar*rLaplacianDataDst,size_t strideSrc,size_t strideDst,int srcWidth, int srcHeight, int expWidth, int expHeight);
    (3) blend( )函數根據掩碼圖像的高斯金字塔以及左、右圖像的拉普拉斯金字塔合成當前層的目標圖像的拉普拉斯金字塔,所有像素值均以掩碼圖像的高斯金字塔為權重而求得。
    _global_ void blend(uchar* mGaussianDataSrc, uchar* lLaplacianDataSrc, uchar* rLaplacianDataSrc,uchar* sLaplacianDataDst, size_t stride,int lapWidth, int lapHeight)
    (4) collapse()函數對圖像的拉普拉斯金字塔分別完成一次expand操作和累加操作,本質上等同于expand操作,兩者的基本算法是相同的。不同點在于expand模塊用于各層高斯金字塔的expand操作,從而生成各層拉普拉斯金字塔,而collapse函數則用于整個融合過程最后的圖像重構步驟,將各層已經求得的拉普拉斯金字塔作擴展和累加操作,生成最后的拼接圖像。
    _global_ void collapse(uchar* sLaplacianDataSrc, uchar* sExpandDataSrc, uchar* sExpandDataDst, size_t strideSrc, size_t strideDst, int srcWidth, int srcHeight, int expWidth, int expHeight)
    本文借助于SIFT特征對于旋轉和尺度的不變性以及對于噪聲干擾良好的魯棒性進行圖像拼接與匹配,使用CUDA技術簡單地對多分辨率融合算法進行了優化,提高了其執行效率和速度。編寫了界面化的Demo程序,實現了基本的圖像拼接功能。
參考文獻
[1] 譚康.圖像拼接技術與實現[D].南京:南京理工大學, 2006.
[2] HARRIS C, STEPHENS M. A combined corner and edge detector[C]. Proceedings of the 4th Alvey Vision Conference, 1988:147-151.
[3] 張小洪,李博,楊丹.一種新的Harris多尺度角點檢測[J]. 電子與信息學報,2007(7):1735-1738.
[4] LOWE D G. Object recognition from local scale-invariant features[C]. The Proceedings of the Seventh IEEE International Conference on Computer Vision, 1999(2):1150-1157.
[5] 騫森,朱劍英. 基于改進的SIFT特征的圖像雙向匹配算法[J]. 機械科學與技術,2007(9):1179-1182.
[6] Peng Xiaoming, Ding Mingyue, Zhou Chengping, et al. Improved approach for object location under affine transformation using the Hausdorff distance[J].Optical Engineering, 2003,42(10):2794-2795.
[7] 張毓晉. 圖像工程(上冊)圖像處理(第2版)[M]. 北京:清華大學出版社,2006.
[8] LINDEBERG T. Detecting salient blob. like image structures and their scales with a scale-space primal sketch[J].International Journal of Computer Vision,1993,11(3):283-318.
[9] NVIDIA. NVIDIA CUDA Programming Guide[Z].
[10] 張舒, 褚艷利. GPU高性能運算之CUDA[M]. 北京:中國水利水電出版社,2009.

此內容為AET網站原創,未經授權禁止轉載。
热re99久久精品国产66热_欧美小视频在线观看_日韩成人激情影院_庆余年2免费日韩剧观看大牛_91久久久久久国产精品_国产原创欧美精品_美女999久久久精品视频_欧美大成色www永久网站婷_国产色婷婷国产综合在线理论片a_国产精品电影在线观看_日韩精品视频在线观看网址_97在线观看免费_性欧美亚洲xxxx乳在线观看_久久精品美女视频网站_777国产偷窥盗摄精品视频_在线日韩第一页
  • <strike id="ygamy"></strike>
  • 
    
      • <del id="ygamy"></del>
        <tfoot id="ygamy"></tfoot>
          <strike id="ygamy"></strike>
          狠狠色噜噜狠狠色综合久| 午夜精品网站| 有码中文亚洲精品| 影音国产精品| 欧美色欧美亚洲另类二区| 欧美顶级艳妇交换群宴| 欧美日韩一区二区三| 久久精品国产清高在天天线| 在线观看视频欧美| 欧美伊人久久久久久午夜久久久久| 欧美激情视频免费观看| 欧美激情第3页| 麻豆精品视频在线观看视频| 欧美日韩一区二区在线观看| 亚洲一区二区三区涩| 国产精品久久久久婷婷| 欧美日韩成人综合| 欧美精品少妇一区二区三区| 国产精品一区二区三区四区五区| 国产精品一区二区女厕厕| 国产精品美女久久久久av超清| 欧美伦理影院| 亚洲精品一区二区三区99| 午夜在线电影亚洲一区| 99re6这里只有精品| 亚洲欧美一区二区三区极速播放| 欧美日韩精品中文字幕| 午夜精品久久久久久久99樱桃| 亚洲欧美日韩国产成人精品影院| 黑人中文字幕一区二区三区| 在线观看欧美亚洲| 国产欧美在线看| 国产精品高潮呻吟久久av无限| 国产一区二区毛片| 欧美国产一区二区三区激情无套| 国产在线国偷精品产拍免费yy| 欧美天堂亚洲电影院在线观看| 99视频精品在线| 欧美日韩综合久久| 欧美亚洲视频在线看网址| 亚洲精品美女在线| 一区二区三区久久久| 亚洲精品久久久久久久久久久| 欧美日韩一区综合| 亚洲一区二区三区视频| 香蕉久久夜色精品国产使用方法| 亚洲女人小视频在线观看| 欧美一二三视频| 久久成人18免费网站| 久久亚洲综合网| 国产精品久久久久久久久婷婷| 久久超碰97人人做人人爱| 国产一区视频网站| 国产欧美日韩一区二区三区在线观看| 久久综合电影| 男女激情视频一区| 欧美在线free| 日韩亚洲精品视频| 欧美一二区视频| 欧美性片在线观看| 欧美日韩国产限制| 欧美88av| 国产精品欧美一区喷水| 欧美精品videossex性护士| 香蕉免费一区二区三区在线观看| 国模精品一区二区三区| 国产视频观看一区| 亚洲免费综合| 国产精品人人做人人爽| 欧美日韩一区免费| 美日韩精品免费观看视频| 亚洲欧美日韩精品久久久久| 精品不卡一区二区三区| 亚洲视频网站在线观看| 国产伦精品一区二区三区在线观看| 一区二区三区四区在线| 在线观看欧美成人| 免费视频一区二区三区在线观看| 夜夜爽99久久国产综合精品女不卡| 一本久久a久久免费精品不卡| 国产精品免费一区二区三区在线观看| 国产亚洲精品福利| 国产一区二区三区丝袜| 欧美午夜激情小视频| 欧美日韩在线第一页| 欧美激情中文字幕一区二区| 一区二区三区福利| 亚洲国产二区| 国产亚洲精品激情久久| 亚洲一区二区三区视频播放| 亚洲精品乱码久久久久久| 欧美日韩国产二区| 免费永久网站黄欧美| 久久久久久久久久看片| 美女免费视频一区| 91久久久亚洲精品| 99国产一区二区三精品乱码| 激情综合色综合久久| 国产精品区二区三区日本| 亚洲人成毛片在线播放女女| 国产精品女主播在线观看| 亚洲精品一区二区三区婷婷月| 美女脱光内衣内裤视频久久网站| 欧美精品一区二区三区在线看午夜| 久久免费国产精品1| 欧美三级资源在线| 国产自产女人91一区在线观看| 久久综合五月天婷婷伊人| 亚洲一区二区在线免费观看| 欧美在线观看视频一区二区三区| 久久人人爽人人| 久久精品国产亚洲一区二区| 欧美va亚洲va日韩∨a综合色| 国产精品久久77777| 一区二区三区国产盗摄| 国产一区在线观看视频| 国产一区二区三区日韩欧美| 国产一区二区三区av电影| 亚洲观看高清完整版在线观看| 国产精品美女xx| 在线国产精品播放| 亚洲春色另类小说| 一区二区久久| 久久精品av麻豆的观看方式| 国产精品二区三区四区| 欧美日韩欧美一区二区| 亚洲尤物视频网| 久久动漫亚洲| 国产精品免费视频xxxx| 黄色成人在线网站| 久久精品91久久香蕉加勒比| 亚洲自拍都市欧美小说| 亚洲精品中文字幕有码专区| 新狼窝色av性久久久久久| 国产精品都在这里| 久久久久久久一区二区| 亚洲一区免费在线观看| 欧美高清影院| 欧美手机在线视频| 91久久久精品| 欧美亚洲日本国产| 亚洲电影在线看| 欧美精品免费在线| 136国产福利精品导航网址| 中文精品视频一区二区在线观看| 久久一区免费| 欧美理论电影在线观看| 欧美极品欧美精品欧美视频| 午夜欧美理论片| 久久免费少妇高潮久久精品99| 久久国产毛片| 午夜视频精品| 久久久99爱| 亚洲欧美成人精品| 久久一区亚洲| 欧美久久久久久| 国产精品高潮视频| 亚洲欧美日韩中文在线制服| 欧美在线一区二区| 亚洲一区中文字幕在线观看| 亚洲美女视频| 激情另类综合| 国产精品久久久久77777| 亚洲激情电影在线| 一区二区国产日产| 一区二区在线视频播放| 亚洲人成77777在线观看网| 在线日韩av永久免费观看| 小黄鸭视频精品导航| 1769国产精品| 欧美中文字幕第一页| 欧美一区二区三区视频在线观看| 一区电影在线观看| 国产在线拍揄自揄视频不卡99| 欧美高清不卡在线| 国产欧美精品| 久热综合在线亚洲精品| 欧美在线1区| 欧美精品免费看| 狠狠色丁香婷婷综合影院| 日韩视频精品在线观看| 久久久久久精| 久久国产免费| 亚洲一区二区三区四区在线观看| 国产精品区免费视频| 一区二区日本视频| 久久精品人人做人人综合| 亚洲精品视频啊美女在线直播| 葵司免费一区二区三区四区五区| 亚洲国产精品久久久| 亚洲乱码日产精品bd| 一本不卡影院| 欧美天天影院| 亚洲欧美色婷婷| 久久久av水蜜桃| 亚洲欧洲免费视频| 欧美色欧美亚洲另类七区| 欧美h视频在线| 欧美一级艳片视频免费观看| 久久国产精品毛片| 亚洲日本欧美| 欧美久久综合| 欧美精品一二三| 久久午夜视频| 日韩午夜在线播放| 一区在线免费观看| 黑人中文字幕一区二区三区| 欧美日韩一区二区在线播放| 国产精品v欧美精品v日韩| 国产精品网站一区| 在线日韩中文字幕| 欧美一区二区在线免费观看| 亚洲精品午夜| 亚洲一区在线直播| 亚洲免费一在线| 国产精品一卡二卡| 国产欧美在线观看| 激情久久久久久久久久久久久久久久| 亚洲欧美国产毛片在线| 亚洲区中文字幕| 亚洲国产精品第一区二区| 欧美日韩国产成人在线免费| 久久成人18免费观看| 9色porny自拍视频一区二区| 午夜视黄欧洲亚洲| 香蕉国产精品偷在线观看不卡| 一本色道综合亚洲| 欧美va亚洲va日韩∨a综合色| 欧美风情在线| 久久天天躁狠狠躁夜夜av| 国内精品久久久久久久97牛牛| 国产精品自拍三区| 亚洲欧美三级在线| 亚洲欧美中文日韩v在线观看| 欧美午夜免费影院| 国产欧美视频在线观看| 在线看日韩欧美| 亚洲美女电影在线| 欧美成人一二三| 欧美福利一区二区| 在线观看一区二区视频| 欧美在线免费视屏| 国产欧美日韩一区二区三区| 欧美成va人片在线观看| 欧美在线亚洲一区| 欧美极品一区二区三区| 亚洲国产国产亚洲一二三| 在线亚洲国产精品网站| 美女网站久久| 亚洲黄色三级| 国产色综合天天综合网| 国产精品毛片在线看| 日韩视频一区二区三区在线播放免费观看| 国产亚洲一本大道中文在线| 欧美日韩亚洲三区| 香蕉乱码成人久久天堂爱免费| 欧美日韩高清区| 亚洲精品视频免费观看| 欧美日本一道本在线视频| 国产性做久久久久久| 久久精品久久99精品久久| 亚洲一区二区三区午夜| 一区二区三区欧美日韩| 欧美91福利在线观看| 国产一区二区按摩在线观看| 欧美精品国产一区二区| 欧美日韩一区二区三区四区五区| 女人香蕉久久**毛片精品| 欧美日韩视频一区二区三区| 亚洲第一页中文字幕| 国产精品久久久久久模特| 欧美体内谢she精2性欧美| 欧美三级电影精品| 久久亚洲春色中文字幕久久久| 国产日韩精品一区二区浪潮av| 久久精品99无色码中文字幕| 国产精品国产成人国产三级| 欧美二区在线看| 欧美特黄a级高清免费大片a级| 国产日韩欧美电影在线观看| 国产精品人成在线观看免费| 欧美va亚洲va国产综合| 欧美激情第三页| 免费看亚洲片| 久久精品一级爱片| 亚洲一本视频| 久久精品伊人| 久久精品国产99国产精品| 欧美激情一区二区三区全黄| 欧美成人免费播放| 亚洲一区三区在线观看| 亚洲在线免费观看| 在线播放豆国产99亚洲| 欧美成人精品在线观看| 欧美片第一页| 亚洲免费小视频| 久热re这里精品视频在线6| 国产精品久久国产三级国电话系列| 久久久欧美精品| 国产精品乱码一区二三区小蝌蚪| 欧美中文在线观看国产| 日韩午夜黄色| 亚洲欧美福利一区二区| 亚洲欧美另类综合偷拍| 欧美日韩视频在线一区二区观看视频| 国产精品久久影院| 国产真实精品久久二三区| 这里只有视频精品| 亚洲人成绝费网站色www| 欧美日韩美女在线观看| 欧美国产日韩一区二区三区| 欧美色综合天天久久综合精品| 国内精品久久久久影院薰衣草| 欧美xxxx在线观看| 欧美中文在线字幕| 亚洲免费观看高清完整版在线观看熊| 一区二区三区在线观看国产| 欧美电影免费观看高清完整版| 午夜精品久久久久久久白皮肤| 亚洲小视频在线观看| 国产在线视频欧美| 亚洲欧美国产高清va在线播| 亚洲激情在线观看| 91久久国产综合久久91精品网站| 亚洲黑丝一区二区| 欧美日韩免费一区二区三区| 久久9热精品视频| 狠狠色综合色综合网络|