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

    基于CUDA技術(shù)的離散小波變換算法研究與實現(xiàn)

    2020-02-22 03:10:52張金霜
    現(xiàn)代信息科技 2020年17期

    摘? 要:針對離散小波變換過程比較耗時、不利于實際工程應(yīng)用的問題,提出利用基于GPU平臺的CUDA技術(shù)對小波變換算法做并行化改造,從而提高算法執(zhí)行效率。該文分析了小波Mallat算法并行化的可行性,并詳細介紹了算法的改造過程。實驗表明,基于GPU/CUDA技術(shù)的并行小波Mallat算法,相較于串行小波變換算法,執(zhí)行速度最高提升了50余倍,且算法效率與計算量成正向關(guān)系。

    關(guān)鍵詞:CUDA;并行程序設(shè)計;離散小波變換;圖像壓縮

    中圖分類號:TN911;TP399? ? ? ?文獻標識碼:A 文章編號:2096-4706(2020)17-0072-05

    Abstract:In view of the time-consuming process of discrete wavelet transform and not conducive to practical engineering applications,it is proposed to use GPU-based CUDA technology to parallelize the wavelet transform algorithm,thereby improving the efficiency of algorithm execution. This article analyzes the feasibility of the parallelization of the wavelet Mallat algorithm,and introduces the algorithm transformation process in detail. Experiments show that the parallel wavelet Mallat algorithm based on GPU/CUDA technology,compared with the serial wavelet transform algorithm,the execution speed is up to 50 times faster,and the algorithm efficiency has a positive relationship with the amount of calculation.

    Keywords:CUDA;parallel programming;discrete wavelet transform;image compression

    0? 引? 言

    小波分析廣泛應(yīng)用于信號分析、圖像識別、計算機視覺和數(shù)據(jù)壓縮等,JPEG2000標準的出現(xiàn)則標志著小波變換在圖像壓縮方面走向成熟,Mallat算法將小波變換真正應(yīng)用于實際,該算法采用濾波器執(zhí)行離散小波變換,但對于數(shù)據(jù)量較大的圖像信息,小波變換的過程耗時較長。為了提升小波變換算法的執(zhí)行效率,可考慮將傳統(tǒng)CPU上的串行結(jié)構(gòu)算法改成基于CPU/GPU異構(gòu)算法,通過GPU內(nèi)部大量獨立的計算單元,實現(xiàn)計算多線程并發(fā)執(zhí)行[1-3]。

    為方便作者在后續(xù)的圖像分析研究工作中,使用高效的離散小波變換技術(shù),及在自研的工程項目中便捷地調(diào)用小波變換算法,因此需要自己重新編寫小波變換程序。本文將結(jié)合小波Mallat算法的數(shù)學理論與程序代碼邏輯,論證該算法具備并行化改造的可行性;再從程序設(shè)計優(yōu)化的角度,利用CUDA技術(shù)改寫小波Mallat算法中的耗時部分,通過利用GPU強大的計算資源來提升算法效率。

    1? CUDA技術(shù)

    CUDA是由顯卡廠商NVIDIA推出的通用并行計算框架,允許設(shè)計人員使用C/C++語言和CUDA擴展庫的形式編寫程序,避免了傳統(tǒng)方法中將通用計算轉(zhuǎn)換為圖形處理的過程,大大降低程序設(shè)計的難度和系統(tǒng)維護成本。

    在CUDA架構(gòu)下,程序被分成兩個部分:host端程序和device端程序。host端程序運行在CPU上,device端程序又叫kernel,運行在GPU上。通常先由host端程序準備數(shù)據(jù),對設(shè)備進行必要的初始化,把數(shù)據(jù)從內(nèi)存?zhèn)魉偷斤@存中,執(zhí)行device端程序,待GPU運行完后,host端程序再從顯存中把結(jié)果取回內(nèi)存中。

    CUDA的線程模型分三個層次:線程(thread)、線程塊(block)、線程網(wǎng)格(grid)。其中,kernel以grid的形式組織,每個線程網(wǎng)格由若干個block組成,而每個線程塊又由若干個thread組成。block內(nèi)的thread之間是并行執(zhí)行的,而grid內(nèi)的block之間也是并行執(zhí)行的,grid與grid之間是串行執(zhí)行的,即GPU同時只能處理一個grid[4]。

    2? 離散小波變換算法

    離散小波及離散小波變換的表達式為[5]:

    其中a為尺度因子,b為平移因子,且a>1,b∈R;m,n∈Z2,m為頻率范圍指數(shù),n為時間步長變化指數(shù);ψm,n(t)為基本小波或母小波;f(t)為原函數(shù),〈〉為平方可積函數(shù)空間L2(R),而〈f,ψm,n〉為小波變換。

    離散小波變換是一種時頻分析,它從集中在某個區(qū)間上基本函數(shù)開始,以規(guī)定步長向左或者向右移動基本波形,并用尺度因子a擴張或收縮來構(gòu)造函數(shù)的系數(shù)。

    如果尺度以二進方式離散變化,可得小波函數(shù)為:

    其中參數(shù)j決定伸縮,而k確定平移幅度,且j,k∈Z。

    由于圖像信息一般是二維的,因此需要將小波理論從一維推廣至二維,二維基本小波如式(4)所示。

    假設(shè)一幅M×M像素的圖像f1(x,y),根據(jù)Mallat算法,對于j=0,則源圖像的尺度為20=1。j值每增大1都使尺度加倍,分辨率減半。每次隔行隔列采樣,圖像被分解為4個大小為原來尺寸四分之一的子頻帶區(qū)域,這4個子頻帶區(qū)域的每一個都是由源圖像與一個小波基函數(shù)內(nèi)積后,再經(jīng)過x和y方向都進行2倍的間隔抽樣而產(chǎn)生的。因此,對于第一個層次j=1,可寫成:

    其中f20(m,n)為頻帶區(qū)域,其下標為尺度,上標為索引而非指數(shù)。

    對于后繼的層次(j>1),(x,y)都以完全相同的方式分解而構(gòu)成4個尺度上的更小的圖像。由此可知,執(zhí)行小波變換的有效方法是使用濾波器,將內(nèi)積寫成卷積的形式:

    因為尺度函數(shù)和小波函數(shù)都是可分離的,所以每個卷積都可以分解成在 (x,y)的行和列上的一維卷積。由式

    (6)可知,實現(xiàn)Mallat小波分解的核心是將待處理數(shù)據(jù)與濾波器進行卷積,最后得到一系列小波系數(shù)[6]。

    3? 基于CUDA實現(xiàn)Mallat算法加速

    3.1? 算法整體設(shè)計

    要利用GPU來加速變換過程,則必須結(jié)合GPU的特性來優(yōu)化算法的效率。算法實現(xiàn)的總體流程如圖1所示。

    算法從host端開始,首先在內(nèi)存中分配兩個具有與被處理圖像像素數(shù)據(jù)相同大小的空間,分別用于存放被處理圖像的源數(shù)據(jù)和處理結(jié)果,并把圖像像素數(shù)據(jù)讀入到內(nèi)存。圖像要在GPU中進行處理,因此需要把圖像像素數(shù)據(jù)復(fù)制到顯存中。在顯存中申請兩個具有與被處理圖像像素數(shù)據(jù)相同大小的空間,分別用于存放圖像源數(shù)據(jù)和處理結(jié)果。其中,存放圖像源數(shù)據(jù)的空間采用紋理存儲器,存放處理結(jié)果的空間采用全局存儲器。

    每一個線程處理一個像素點的卷積運算。小波變換完成后,將結(jié)果寫入到全局存儲器中,同樣是每個線程寫入一個系數(shù),至此device端的工作完成。最后,host端進行數(shù)據(jù)量化和小波重構(gòu),并顯示最終結(jié)果,整個算法完成。

    3.2? 源程序分析

    本文采用的小波Mallat算法的偽代碼為:

    //卷積操作

    void Convolution {

    for(..;..;..) {

    for(..;..;..) {

    行(列)前半部分時域數(shù)據(jù) += 時域數(shù)據(jù)*分解尺度函數(shù);

    行(列)后半部分時域數(shù)據(jù) += 時域數(shù)據(jù)*分解母函數(shù);

    }

    }

    }

    //離散小波變換過程

    void Wavelet() {

    for(n=0;n

    Height=lHeight>>2;? ? ?//圖像低頻信息的高度

    Width=lWidth>>2;? ? ? ?//圖像低頻信息的寬度

    for(i = 0; i < Height; i++) {

    for(j=0 ; j< Width ; j++)

    { 獲取行數(shù)據(jù) }

    Convolution(…);? ?// 對行方向進行卷積運算

    }

    for(i = 0; i < Width; i++) {

    for(j=0;j< Height;j++)

    { 獲取列數(shù)據(jù) }

    Convolution(…);? ?// 對列方向進行卷積運算

    }

    } //小波變換結(jié)束

    }

    循環(huán)控制小波變換的層數(shù),每循環(huán)一次,則進行一次完整的行卷積運算和列卷積運算,同時小波變換處理的范圍縮小為原來的1/4。此時圖像的低頻有效信息集中在圖像的左上角,而其他區(qū)域為高頻細節(jié)信息,不屬于下一次循環(huán)小波處理的區(qū)域。

    由上述代碼可知,該算法耗時的部分出現(xiàn)在圖像的遍歷和卷積過程,改進的關(guān)鍵在于如何將該部分串行代碼變成運行在GPU上并行代碼。

    并行算法設(shè)計思路為:

    (1)串行算法中,通過兩重循環(huán),實現(xiàn)圖像的遍歷。針對這個操作,在CUDA中,可以將源圖像數(shù)據(jù)映射到紋理存儲器(這過程稱為紋理映射),利用紋理存儲器支持二維尋址的特性,實現(xiàn)對像元的快速訪問。

    (2)串行算法中,每個像素點要完成卷積運算,即使每個像素點的計算相對獨立,也必須依次進行。針對這個操作,在CUDA中,可以開辟多個線程,每個線程處理一個像素點,獨立完成卷積。

    3.3? 并行算法設(shè)計

    3.3.1? 線程劃分

    在GPU的SM(Streaming Multiprocessor,多處理器)中,線程被組織成warp(線程束)執(zhí)行,一個warp中包含32個線程,所以一個block中的線程數(shù)最好是32的倍數(shù),這里選擇一個block中申請16×16個thread,即定義dim3 threadBlock(16,16)。把整個圖像作為一個grid,整個grid包含和圖像像素點相等的線程數(shù),則block的數(shù)量為(width/threadBlock.x)×(height/threadBlock.y),其中,width和height分別表示圖像的寬和高。線程劃分如圖2所示。

    3.3.2? 存儲器分配

    為了存放源圖像數(shù)據(jù)和處理后的圖像數(shù)據(jù),必須開辟兩個和源圖像數(shù)據(jù)相同大小的存儲空間。

    首先,存放源圖像數(shù)據(jù)可采用CUDA數(shù)組,這可方便地將數(shù)據(jù)映射到紋理存儲器。紋理數(shù)據(jù)按16的倍數(shù)對齊,符合“合并訪問”的要求,能提高訪存效率。由于在做列卷積的過程中,需要用到行卷積后的數(shù)據(jù),為此要保存原來的行卷積結(jié)果,并更新紋理存儲器中數(shù)據(jù)。紋理存儲器是個只讀存儲器,不支持數(shù)據(jù)寫入,但可以通過更新與之綁定的CUDA數(shù)組中的值,來達到更行紋理存儲器中的數(shù)據(jù)的目的。這個過程中,因為數(shù)據(jù)仍然在device端進行傳遞,耗時不長。

    其次,存放小波變換處理后的數(shù)據(jù),可采用全局存儲器。全局存儲器使用的是普通顯存,整個grid的所有線程都能訪問,由于沒有片內(nèi)緩存,訪問速度慢,但通過合并訪問,可以隱藏延時。

    此外,卷積過程中用到的“分解尺度函數(shù)”和“分解母函數(shù)”初始化后不再發(fā)生改變,故可將它們存放在常數(shù)存儲器中。常數(shù)存儲器大小為64 KB,是只讀的地址空間,其數(shù)據(jù)位于顯存,但擁有緩存加速。在CUDA程序中常數(shù)存儲器一般用于存儲需要頻繁訪問的只讀參數(shù)[7]。

    3.3.3? 并行算法實現(xiàn)

    host端控制小波變換的層數(shù),device端執(zhí)行一次小波變換。host端關(guān)鍵偽代碼為:

    for(int n=0;n

    width=lWidth>>n? ? ;? //通過移位操作,控制小波變化的范圍

    height=lHeight>>n ;? //通過移位操作,控制小波變化的范圍

    //根據(jù)小波變換的層次,不斷調(diào)整block的數(shù)量

    dim3 blockGrid(width/threadBlock.x, height/threadBlock.y);

    convolutionRow<<>>(.. ..); //行方向卷積

    cudaMemcpyToArray(cu_array,0,0,d_Result,size,cudaMemcpyDeviceToDevice); //更新紋理數(shù)據(jù)

    convolutionCol<<>>(.. ..); //列方向卷積

    cudaMemcpyToArray(cu_array,0,0,d_Result,size,cudaMemcpyDeviceToDevice); //更新紋理數(shù)據(jù)

    }

    因為每次小波變換處理的區(qū)域都在發(fā)生改變,為了在device端更方便地控制線程操作,故每次循環(huán)都重新定義grid,修改grid中block的數(shù)量。

    接下來,調(diào)用device端的行(列)卷積運算,進行小波變化,每完成一次行(列)卷積,都必須更新紋理數(shù)據(jù),以供下一次循環(huán)數(shù)據(jù)調(diào)用。device端小波變換偽代碼為:

    //行方向上卷積

    if(x

    for(int j=0; j< 某參數(shù); j++){

    m=(某計算)% width;? // width是小波變換作用的區(qū)域

    tmp += tex2D(texData, ((float)m + 0.5f) , iy) *LF[j]; //時域與分解尺度函數(shù)相乘

    }

    f[y*lwidth+x] = tmp; // lwidth是原圖像的寬度,f存放結(jié)果

    }else if(x

    for(int j=0; j< 某參數(shù); j++){

    m=(某計算)% Width; // width是小波變換作用的區(qū)域

    tmp += tex2D(texData, ((float)m + 0.5f) , iy) *HF[j]; //時域與分解母函數(shù)相乘

    }

    f[y*lwidth+x] = tmp; // lwidth是原圖像的寬度,f存放結(jié)果

    }

    上述代碼的執(zhí)行過程是并行的,即每個線程會并發(fā)地去執(zhí)行相應(yīng)的計算,然后將計算的結(jié)果保存在全局存儲器中。列方向卷積同理,不再列舉。

    由上述分析可見,算法的時間復(fù)雜度由原來的四重循環(huán)變?yōu)閮芍匮h(huán),而兩重循環(huán)的基數(shù)非常小,所以大大地縮短了算法運行時間。

    4? 實驗結(jié)果分析

    本程序目的在于對比小波變換在CPU和GPU上的運行效率,通過調(diào)用同一幅圖片進行小波變換,取得各自的執(zhí)行時間,最后算出加速比。

    硬件上,CPU采用IntelCore i5-7200,主頻為2.5 GHz;GPU采用NVIDIA公司的GeForce GTX660,支持CUDA 4.0以上驅(qū)動版本,顯存容量2 GB。

    實驗程序采用Daubechies D4小波,變換層次為3。所有數(shù)據(jù)為5次計算時間平均值,CUDA計算時間包括了數(shù)據(jù)在host端(主機)與device端(顯卡存儲器)中交換的時間。

    經(jīng)過測試,小波變換在CPU和GPU上執(zhí)行的效率對比如表1所示。從表中數(shù)據(jù)可以看出,應(yīng)用CUDA技術(shù)加速小波變換,均達到兩個數(shù)量級的加速效果。

    由圖3可知,隨著圖像大小的增加,CPU上小波變換所需的時間急劇上升,而GPU上執(zhí)行所需時間則增長緩慢。由圖4可見,圖像大小與加速比總體呈正向關(guān)系。

    5? 結(jié)? 論

    本文重點分析了小波Mallat算法并行化的可行性,并成功利用CUDA技術(shù)實現(xiàn)了高速并行小波變換算法。實驗表明,相較于CPU上的串行小波變換算法,本文算法的執(zhí)行效率提升了數(shù)十倍,最高達到了約51.6倍,且算法效率與計算量成正向關(guān)系。通過合理設(shè)計CUDA程序,將耗時的串行計算過程并行化,充分利用GPU強大的計算資源,能最大限度地縮短執(zhí)行時間,使得小波技術(shù)更好地應(yīng)用于實際工程系統(tǒng)中。

    參考文獻:

    [1] 陳慶奎,王海峰,那麗春,等.圖形處理器通用計算的研究綜述 [J].黑龍江大學自然科學學報,2012,29(5):672-679.

    [2] 陳茜.基于GPU的數(shù)字圖像處理并行算法的研究 [D].西安:中國科學院研究生院(西安光學精密機械研究所),2014.

    [3] 白兆峰.基于GPU的JPEG2000圖像壓縮編碼技術(shù)的研究 [D].西安:中國科學院大學(中國科學院西安光學精密機械研究所),2016.

    [4] 張舒,褚艷利.GPU高性能運算之CUDA [M].北京:中國水利水電出版社,2009:20-22.

    [5] 陳天華.數(shù)字圖像處理 [M].北京:清華大學出版社,2007:132-133.

    [6] 韓志偉,劉志剛,魯曉帆,等.基于CUDA的高速并行小波算法及其在電力系統(tǒng)諧波分析中的應(yīng)用 [J].電力自動化設(shè)備,2010,30(1):98-101+105.

    [7] 鄒巖,楊志義,張凱龍.CUDA并行程序的內(nèi)存訪問優(yōu)化技術(shù)研究 [J].計算機測量與控制,2009,17(12):2504-2506.

    作者簡介:張金霜(1987—),男,漢族,廣東茂名人,教師,畢業(yè)于華南農(nóng)業(yè)大學,碩士研究生,研究方向:計算機教育。

    97超级碰碰碰精品色视频在线观看| 中文字幕av成人在线电影| 久久久精品大字幕| 亚洲在线自拍视频| 精品一区二区三区av网在线观看| 怎么达到女性高潮| 男人的好看免费观看在线视频| 天堂动漫精品| 久久久成人免费电影| 国内揄拍国产精品人妻在线| 村上凉子中文字幕在线| 一区二区三区激情视频| 国内少妇人妻偷人精品xxx网站| 此物有八面人人有两片| 一区二区三区激情视频| 此物有八面人人有两片| 亚洲自拍偷在线| 久久久久久九九精品二区国产| 人人妻人人澡欧美一区二区| eeuss影院久久| 他把我摸到了高潮在线观看| 免费高清视频大片| 久久久久久大精品| 欧美成人性av电影在线观看| av视频在线观看入口| www.www免费av| 久久精品91无色码中文字幕| 日韩欧美精品免费久久 | 成年免费大片在线观看| 亚洲第一欧美日韩一区二区三区| 久久伊人香网站| 99热精品在线国产| 欧美乱码精品一区二区三区| 亚洲国产精品合色在线| 看片在线看免费视频| 99久久99久久久精品蜜桃| 亚洲色图av天堂| 成人av一区二区三区在线看| 此物有八面人人有两片| 热99在线观看视频| 午夜免费观看网址| 国产一区二区在线av高清观看| 我要搜黄色片| 91久久精品国产一区二区成人 | 国产 一区 欧美 日韩| 中文在线观看免费www的网站| 国产在线精品亚洲第一网站| 91在线观看av| 黄色成人免费大全| 国产一区二区三区视频了| 亚洲美女黄片视频| 久久久成人免费电影| 久久香蕉国产精品| 天天添夜夜摸| 久久精品91无色码中文字幕| 热99re8久久精品国产| 色av中文字幕| 久久久精品大字幕| 韩国av一区二区三区四区| 久久久久国产精品人妻aⅴ院| 麻豆成人av在线观看| 91字幕亚洲| 久久婷婷人人爽人人干人人爱| 久久香蕉国产精品| 91在线精品国自产拍蜜月 | 中文字幕久久专区| 国产色爽女视频免费观看| 成人一区二区视频在线观看| 国产精品亚洲美女久久久| 久久久久免费精品人妻一区二区| 啦啦啦观看免费观看视频高清| 免费人成在线观看视频色| 脱女人内裤的视频| 国产伦一二天堂av在线观看| 久久久久免费精品人妻一区二区| 搡老熟女国产l中国老女人| 一本综合久久免费| 亚洲成av人片在线播放无| 99热精品在线国产| 亚洲av成人精品一区久久| 亚洲在线自拍视频| 九色国产91popny在线| 一本精品99久久精品77| 亚洲久久久久久中文字幕| av专区在线播放| 脱女人内裤的视频| 日韩欧美国产在线观看| 精品久久久久久久久久久久久| 高清在线国产一区| 欧美极品一区二区三区四区| 国产单亲对白刺激| 国产精品久久久久久精品电影| 国产精品一及| 国产私拍福利视频在线观看| 国产精品一及| 午夜免费观看网址| 亚洲性夜色夜夜综合| 黄色视频,在线免费观看| 岛国在线免费视频观看| 亚洲黑人精品在线| 成人午夜高清在线视频| 欧美成人a在线观看| 亚洲黑人精品在线| 女人十人毛片免费观看3o分钟| 天美传媒精品一区二区| 亚洲av一区综合| 日本a在线网址| а√天堂www在线а√下载| 亚洲av第一区精品v没综合| 中国美女看黄片| 最近最新免费中文字幕在线| 国产成人影院久久av| 久久九九热精品免费| 亚洲欧美日韩高清专用| 国产v大片淫在线免费观看| 婷婷丁香在线五月| 日本在线视频免费播放| 一卡2卡三卡四卡精品乱码亚洲| 亚洲av电影不卡..在线观看| 国产97色在线日韩免费| 亚洲av成人精品一区久久| 久久久久久久久中文| 国产精品亚洲美女久久久| 欧美激情久久久久久爽电影| 亚洲va日本ⅴa欧美va伊人久久| 国产一级毛片七仙女欲春2| 欧美三级亚洲精品| 亚洲一区二区三区色噜噜| 午夜福利高清视频| 波多野结衣高清无吗| 女人高潮潮喷娇喘18禁视频| 国产成人系列免费观看| 男女下面进入的视频免费午夜| 国产黄a三级三级三级人| 女生性感内裤真人,穿戴方法视频| 国产精品精品国产色婷婷| 国产精品女同一区二区软件 | 午夜福利免费观看在线| 麻豆国产97在线/欧美| 久久久久久大精品| 国产精华一区二区三区| 欧美成狂野欧美在线观看| 欧美日韩福利视频一区二区| 日韩亚洲欧美综合| 少妇人妻精品综合一区二区 | 婷婷亚洲欧美| 亚洲av五月六月丁香网| 午夜福利在线观看免费完整高清在 | 久久久久国产精品人妻aⅴ院| 老师上课跳d突然被开到最大视频 久久午夜综合久久蜜桃 | 国产真人三级小视频在线观看| 久久这里只有精品中国| 日韩欧美三级三区| 久久精品影院6| 欧美黑人巨大hd| 国产精品嫩草影院av在线观看 | 内地一区二区视频在线| 在线播放国产精品三级| 精品一区二区三区av网在线观看| 欧美日韩精品网址| 在线天堂最新版资源| 中文字幕av成人在线电影| 五月伊人婷婷丁香| 国产免费av片在线观看野外av| 国产高清视频在线观看网站| av视频在线观看入口| 久久久久九九精品影院| 别揉我奶头~嗯~啊~动态视频| 亚洲成a人片在线一区二区| 色综合欧美亚洲国产小说| 色视频www国产| 99热精品在线国产| 搞女人的毛片| 亚洲av二区三区四区| 1000部很黄的大片| 亚洲乱码一区二区免费版| 久久99热这里只有精品18| 免费av毛片视频| 九九久久精品国产亚洲av麻豆| 亚洲av一区综合| 最后的刺客免费高清国语| 亚洲成人久久爱视频| 欧美极品一区二区三区四区| 黑人欧美特级aaaaaa片| 亚洲黑人精品在线| aaaaa片日本免费| 国产成人av激情在线播放| 中文字幕高清在线视频| 一个人看视频在线观看www免费 | 亚洲美女黄片视频| 五月伊人婷婷丁香| 在线免费观看的www视频| 琪琪午夜伦伦电影理论片6080| 一夜夜www| 欧美成人一区二区免费高清观看| av女优亚洲男人天堂| 精品熟女少妇八av免费久了| 啦啦啦观看免费观看视频高清| 欧美日韩国产亚洲二区| 日本成人三级电影网站| 国产一区二区三区视频了| 精品人妻1区二区| 国内揄拍国产精品人妻在线| 国产成年人精品一区二区| 亚洲性夜色夜夜综合| 18禁黄网站禁片午夜丰满| 国产精品亚洲av一区麻豆| 欧美国产日韩亚洲一区| 亚洲欧美日韩高清在线视频| 免费av观看视频| 两性午夜刺激爽爽歪歪视频在线观看| 亚洲专区中文字幕在线| 五月伊人婷婷丁香| 五月玫瑰六月丁香| 国产精品自产拍在线观看55亚洲| 久久久久久久精品吃奶| 欧美成人一区二区免费高清观看| svipshipincom国产片| 亚洲av中文字字幕乱码综合| 亚洲色图av天堂| 国内精品美女久久久久久| 欧美日本亚洲视频在线播放| 亚洲激情在线av| 国产亚洲欧美在线一区二区| 欧美乱妇无乱码| 国产97色在线日韩免费| 国产综合懂色| 在线观看舔阴道视频| 欧美日韩综合久久久久久 | 日韩欧美 国产精品| 国产欧美日韩一区二区三| 国产精品三级大全| 亚洲七黄色美女视频| 麻豆成人午夜福利视频| 99久国产av精品| 在线播放国产精品三级| av福利片在线观看| 一进一出抽搐gif免费好疼| eeuss影院久久| 九色成人免费人妻av| av视频在线观看入口| 99在线视频只有这里精品首页| 欧美一级a爱片免费观看看| 91九色精品人成在线观看| 日本成人三级电影网站| 波多野结衣高清无吗| 丰满的人妻完整版| 国产精品爽爽va在线观看网站| 久久九九热精品免费| 99精品久久久久人妻精品| 色av中文字幕| 啦啦啦观看免费观看视频高清| 内射极品少妇av片p| 在线观看66精品国产| 亚洲最大成人中文| 亚洲最大成人中文| 亚洲成人中文字幕在线播放| 色综合亚洲欧美另类图片| 欧美一区二区国产精品久久精品| 免费人成视频x8x8入口观看| 中国美女看黄片| 热99re8久久精品国产| 日韩大尺度精品在线看网址| 香蕉久久夜色| 日韩中文字幕欧美一区二区| 亚洲乱码一区二区免费版| 国产一区二区在线av高清观看| 最后的刺客免费高清国语| 夜夜躁狠狠躁天天躁| 白带黄色成豆腐渣| 精华霜和精华液先用哪个| 麻豆国产97在线/欧美| 中国美女看黄片| 九色国产91popny在线| 757午夜福利合集在线观看| 日本五十路高清| 长腿黑丝高跟| 一级作爱视频免费观看| av黄色大香蕉| 色老头精品视频在线观看| 欧美又色又爽又黄视频| 一二三四社区在线视频社区8| 色精品久久人妻99蜜桃| 国产成人a区在线观看| 国语自产精品视频在线第100页| 人人妻人人澡欧美一区二区| www.999成人在线观看| 中国美女看黄片| 高潮久久久久久久久久久不卡| 一个人看的www免费观看视频| 色在线成人网| 哪里可以看免费的av片| 一个人观看的视频www高清免费观看| 欧美日韩福利视频一区二区| 免费av观看视频| 精品一区二区三区人妻视频| 99久久精品热视频| 观看美女的网站| 国产成人av教育| 亚洲国产日韩欧美精品在线观看 | 偷拍熟女少妇极品色| 免费观看的影片在线观看| 人妻久久中文字幕网| 久久国产精品影院| www日本黄色视频网| 精品久久久久久久毛片微露脸| 国模一区二区三区四区视频| 狂野欧美激情性xxxx| 国产激情欧美一区二区| 亚洲av二区三区四区| 亚洲精品美女久久久久99蜜臀| 男女视频在线观看网站免费| 亚洲一区高清亚洲精品| 757午夜福利合集在线观看| 男人舔奶头视频| 99热只有精品国产| 欧美性感艳星| 一本一本综合久久| 法律面前人人平等表现在哪些方面| 在线天堂最新版资源| 久久精品夜夜夜夜夜久久蜜豆| 久久精品91无色码中文字幕| 又爽又黄无遮挡网站| 婷婷精品国产亚洲av| 亚洲国产欧洲综合997久久,| 18美女黄网站色大片免费观看| 久久国产精品人妻蜜桃| ponron亚洲| 五月玫瑰六月丁香| 在线看三级毛片| 99热精品在线国产| 国产野战对白在线观看| 午夜福利免费观看在线| 成人性生交大片免费视频hd| 99久久成人亚洲精品观看| 手机成人av网站| 亚洲精品成人久久久久久| 久久久精品欧美日韩精品| 真人一进一出gif抽搐免费| 级片在线观看| 国内精品一区二区在线观看| 亚洲激情在线av| 亚洲黑人精品在线| 国产精品久久久久久久久免 | 日本免费一区二区三区高清不卡| 亚洲av日韩精品久久久久久密| 99精品久久久久人妻精品| 久久久久久久亚洲中文字幕 | 国内精品久久久久精免费| 国语自产精品视频在线第100页| 人人妻人人澡欧美一区二区| 综合色av麻豆| 国产中年淑女户外野战色| 成人欧美大片| 国产真实伦视频高清在线观看 | 亚洲精品456在线播放app | 欧美日韩综合久久久久久 | 欧美三级亚洲精品| 国产伦精品一区二区三区视频9 | 国产91精品成人一区二区三区| 少妇裸体淫交视频免费看高清| 女人被狂操c到高潮| 国产毛片a区久久久久| 午夜影院日韩av| 久久精品综合一区二区三区| 欧美成狂野欧美在线观看| 国产亚洲精品综合一区在线观看| 三级国产精品欧美在线观看| 国产精品久久久久久精品电影| 欧美成人一区二区免费高清观看| 最近最新中文字幕大全电影3| 九九在线视频观看精品| 亚洲av电影在线进入| 亚洲欧美日韩高清专用| 精品国内亚洲2022精品成人| 悠悠久久av| 日韩欧美三级三区| 国产v大片淫在线免费观看| 国产精品一区二区三区四区免费观看 | 国产成人系列免费观看| 日本 欧美在线| 精品久久久久久久久久免费视频| 国内精品美女久久久久久| 深爱激情五月婷婷| 制服人妻中文乱码| 亚洲 欧美 日韩 在线 免费| 国产成年人精品一区二区| 男插女下体视频免费在线播放| 级片在线观看| 伊人久久大香线蕉亚洲五| 亚洲熟妇熟女久久| 热99re8久久精品国产| 熟妇人妻久久中文字幕3abv| 宅男免费午夜| 成年人黄色毛片网站| 日本一本二区三区精品| 黄色日韩在线| 搡女人真爽免费视频火全软件 | 波野结衣二区三区在线 | 中文字幕人妻丝袜一区二区| 亚洲精品影视一区二区三区av| 欧美3d第一页| 99国产精品一区二区三区| 久久精品夜夜夜夜夜久久蜜豆| 亚洲五月天丁香| 在线播放国产精品三级| 色哟哟哟哟哟哟| 欧美在线黄色| 亚洲va日本ⅴa欧美va伊人久久| 亚洲精品一卡2卡三卡4卡5卡| 国产精品综合久久久久久久免费| 国内揄拍国产精品人妻在线| 偷拍熟女少妇极品色| 欧美成狂野欧美在线观看| 欧美成人免费av一区二区三区| 淫妇啪啪啪对白视频| 又黄又爽又免费观看的视频| 国产高潮美女av| 九九久久精品国产亚洲av麻豆| 一个人观看的视频www高清免费观看| 法律面前人人平等表现在哪些方面| 亚洲精品一卡2卡三卡4卡5卡| 久久久久久久久大av| 美女cb高潮喷水在线观看| 特级一级黄色大片| 亚洲片人在线观看| 亚洲国产精品成人综合色| 特级一级黄色大片| 欧美乱色亚洲激情| 国产99白浆流出| 久久精品国产亚洲av香蕉五月| 亚洲av一区综合| 亚洲在线自拍视频| 国产精品综合久久久久久久免费| 亚洲狠狠婷婷综合久久图片| 午夜激情福利司机影院| 欧美日本视频| 久久国产乱子伦精品免费另类| 女生性感内裤真人,穿戴方法视频| 免费在线观看亚洲国产| 最近最新中文字幕大全电影3| 身体一侧抽搐| 久久精品综合一区二区三区| 久久久久性生活片| 亚洲精品在线美女| 亚洲av免费在线观看| 久久天躁狠狠躁夜夜2o2o| 午夜福利成人在线免费观看| 91九色精品人成在线观看| 老熟妇仑乱视频hdxx| 深夜精品福利| 日韩av在线大香蕉| 97超级碰碰碰精品色视频在线观看| 久久久久国产精品人妻aⅴ院| 国产野战对白在线观看| 欧美性猛交黑人性爽| 熟女人妻精品中文字幕| 怎么达到女性高潮| 全区人妻精品视频| 长腿黑丝高跟| aaaaa片日本免费| 色老头精品视频在线观看| 黄色女人牲交| 欧美3d第一页| 黄色成人免费大全| 免费高清视频大片| 亚洲成a人片在线一区二区| 欧美激情久久久久久爽电影| 禁无遮挡网站| 一二三四社区在线视频社区8| 天美传媒精品一区二区| 免费搜索国产男女视频| 又爽又黄无遮挡网站| 成人永久免费在线观看视频| 亚洲真实伦在线观看| 18禁国产床啪视频网站| 69av精品久久久久久| 国产v大片淫在线免费观看| 国产成+人综合+亚洲专区| 丁香欧美五月| 亚洲五月天丁香| 日韩欧美国产在线观看| 精品人妻偷拍中文字幕| 女人十人毛片免费观看3o分钟| 特大巨黑吊av在线直播| 国产精品国产高清国产av| 国产精品,欧美在线| 操出白浆在线播放| 欧美绝顶高潮抽搐喷水| 久久精品国产亚洲av香蕉五月| 午夜老司机福利剧场| 成人国产综合亚洲| 精品乱码久久久久久99久播| 日韩高清综合在线| 色尼玛亚洲综合影院| 久久久久国产精品人妻aⅴ院| 在线视频色国产色| 噜噜噜噜噜久久久久久91| 老司机午夜福利在线观看视频| 韩国av一区二区三区四区| 国产三级在线视频| 我要搜黄色片| 成人鲁丝片一二三区免费| 国产午夜精品论理片| 亚洲国产精品999在线| 两个人视频免费观看高清| 国产单亲对白刺激| 色老头精品视频在线观看| 久久精品国产亚洲av香蕉五月| 毛片女人毛片| 女同久久另类99精品国产91| 丰满乱子伦码专区| 亚洲男人的天堂狠狠| 亚洲 国产 在线| 叶爱在线成人免费视频播放| 国产av麻豆久久久久久久| 在线视频色国产色| 99riav亚洲国产免费| 看片在线看免费视频| 日本熟妇午夜| 午夜福利免费观看在线| 色尼玛亚洲综合影院| 老熟妇仑乱视频hdxx| 亚洲av熟女| 中文在线观看免费www的网站| 村上凉子中文字幕在线| 五月玫瑰六月丁香| 禁无遮挡网站| 国产伦精品一区二区三区四那| 一区二区三区高清视频在线| 亚洲精品成人久久久久久| 听说在线观看完整版免费高清| 禁无遮挡网站| 色尼玛亚洲综合影院| 精品国产亚洲在线| 最新中文字幕久久久久| 成人性生交大片免费视频hd| 中文字幕高清在线视频| 午夜免费成人在线视频| а√天堂www在线а√下载| 久久伊人香网站| 国产极品精品免费视频能看的| 亚洲无线在线观看| 色av中文字幕| 国产高潮美女av| 日韩av在线大香蕉| 99久久成人亚洲精品观看| 大型黄色视频在线免费观看| 精品久久久久久久毛片微露脸| 99久久成人亚洲精品观看| 最好的美女福利视频网| 男插女下体视频免费在线播放| 91九色精品人成在线观看| 色吧在线观看| 国产淫片久久久久久久久 | 噜噜噜噜噜久久久久久91| 精品午夜福利视频在线观看一区| 亚洲最大成人中文| 最新中文字幕久久久久| 俄罗斯特黄特色一大片| 国产单亲对白刺激| 欧美精品啪啪一区二区三区| www日本黄色视频网| 色综合站精品国产| 欧美黄色淫秽网站| 国产激情偷乱视频一区二区| ponron亚洲| 性色av乱码一区二区三区2| 国产69精品久久久久777片| 在线十欧美十亚洲十日本专区| 一进一出好大好爽视频| 欧美中文综合在线视频| 丰满人妻一区二区三区视频av | 国产亚洲精品一区二区www| 18禁在线播放成人免费| 国产精品亚洲美女久久久| 观看美女的网站| 成人国产综合亚洲| 中亚洲国语对白在线视频| 99视频精品全部免费 在线| 性色avwww在线观看| 在线天堂最新版资源| 男人和女人高潮做爰伦理| 日韩人妻高清精品专区| 99久久精品国产亚洲精品| 亚洲国产精品999在线| 中出人妻视频一区二区| 国产精品1区2区在线观看.| 99热只有精品国产| 久久久久久久精品吃奶| av黄色大香蕉| 国产爱豆传媒在线观看| 深爱激情五月婷婷| 国产精品久久电影中文字幕| 国产99白浆流出| 俄罗斯特黄特色一大片| 99热这里只有精品一区| 黄色日韩在线| 国产一区二区在线av高清观看| 成熟少妇高潮喷水视频| 日本黄色视频三级网站网址| 有码 亚洲区| 成年女人毛片免费观看观看9| 欧美不卡视频在线免费观看| 一二三四社区在线视频社区8| 成年免费大片在线观看| 亚洲成人免费电影在线观看| 在线观看美女被高潮喷水网站 | 国产亚洲精品综合一区在线观看| 国产一区二区在线观看日韩 | 亚洲av免费在线观看| 成年女人毛片免费观看观看9| 老熟妇仑乱视频hdxx| 精品免费久久久久久久清纯| 啦啦啦观看免费观看视频高清|