《電子技術(shù)應(yīng)用》
您所在的位置:首頁 > 嵌入式技術(shù) > 設(shè)計(jì)應(yīng)用 > CUDA加速分形火焰繪制
CUDA加速分形火焰繪制
來源:微型機(jī)與應(yīng)用2013年第23期
劉進(jìn)鋒
(寧夏大學(xué) 數(shù)學(xué)計(jì)算機(jī)學(xué)院,寧夏 銀川 750021)
摘要: 提出了一種基于CUDA的并行分形火焰繪制算法,該算法利用了GPU的單指令多線程的特點(diǎn),將常用于迭代函數(shù)系統(tǒng)(IFS)的傳統(tǒng)的混沌游戲(Chaos Game)算法作了并行化修改,并在圖形輸出時(shí)利用了CUDA與OpenGL互操作加速分形火焰繪制。實(shí)驗(yàn)證明,該并行方法比CPU上運(yùn)行的普通算法快了15倍左右,能夠?qū)崟r(shí)繪制分形火焰圖形。在上述基本算法的基礎(chǔ)上,又進(jìn)一步研究了消除分支分歧的改進(jìn)算法,改進(jìn)算法的運(yùn)行時(shí)間具有相對(duì)于變換函數(shù)數(shù)量的恒定性,多數(shù)情況下比基本算法性能更優(yōu)越。
Abstract:
Key words :

摘  要: 提出了一種基于CUDA的并行分形火焰繪制算法,該算法利用了GPU的單指令多線程的特點(diǎn),將常用于迭代函數(shù)系統(tǒng)(IFS)的傳統(tǒng)的混沌游戲(Chaos Game)算法作了并行化修改,并在圖形輸出時(shí)利用了CUDA與OpenGL互操作加速分形火焰繪制。實(shí)驗(yàn)證明,該并行方法比CPU上運(yùn)行的普通算法快了15倍左右,能夠?qū)崟r(shí)繪制分形火焰圖形。在上述基本算法的基礎(chǔ)上,又進(jìn)一步研究了消除分支分歧的改進(jìn)算法,改進(jìn)算法的運(yùn)行時(shí)間具有相對(duì)于變換函數(shù)數(shù)量的恒定性,多數(shù)情況下比基本算法性能更優(yōu)越。
關(guān)鍵詞: 分形火焰;迭代函數(shù)系統(tǒng);統(tǒng)一計(jì)算設(shè)備架構(gòu);圖形處理器

 分形通常被定義為“一個(gè)粗糙或零碎的幾何形狀,可以分成數(shù)個(gè)部分,且每一部分都(至少近似地)是整體縮小后的形狀”,即具有自相似的性質(zhì)。分形一詞于1975年由曼德博創(chuàng)造出,來自拉丁文“frāctus”,有“零碎”、“破裂”之意。分形有幾種類型,可以分別依據(jù)表現(xiàn)出的精確自相似性、半自相似性和統(tǒng)計(jì)自相似性來定義。雖然分形是一個(gè)數(shù)學(xué)構(gòu)造,但也可以在自然界中找到。分形在藝術(shù)作品、醫(yī)學(xué)、土力學(xué)、地震學(xué)和技術(shù)分析中都有應(yīng)用。
 分形火焰是迭代函數(shù)系統(tǒng)IFS(Iterated Function System)分形的一種更復(fù)雜的變種形式,與普通的IFS相比,分形火焰能產(chǎn)生變化多端、引人入勝的圖形,但與此同時(shí),其計(jì)算復(fù)雜度很高。
 圖形處理器(GPU)原本是處理計(jì)算機(jī)圖形的專用設(shè)備,近十年來,由于高清晰度復(fù)雜圖形實(shí)時(shí)處理的需求,GPU發(fā)展成為高并行度、多線程、多核的處理器。目前,主流GPU的運(yùn)算能力已超過主流通用CPU,從發(fā)展趨勢(shì)上來看,將來差距會(huì)越拉越大。GPU卓越的性能對(duì)開發(fā)GPGPU(使用GPU進(jìn)行通用計(jì)算)非常具有吸引力。統(tǒng)一計(jì)算設(shè)備架構(gòu)CUDA(Compute Unified Device Architecture)是NVIDIA公司伴隨著統(tǒng)一渲染架構(gòu)而推出的一種通用的GPU編程模型,可以將GPU視為一個(gè)并行數(shù)據(jù)計(jì)算的設(shè)備,對(duì)所進(jìn)行的計(jì)算進(jìn)行分配和管理[1]。在CUDA的架構(gòu)中,通用計(jì)算不再像過去的GPGPU那樣必須將計(jì)算映射到圖形API中,開發(fā)者無需學(xué)習(xí)復(fù)雜的顯示芯片的指令或是特殊的結(jié)構(gòu),CUDA編程語言只是對(duì)標(biāo)準(zhǔn)的C語言作了少量擴(kuò)展,因此開發(fā)門檻大大降低了。目前有很多基于CUDA的GPU計(jì)算的研究成果,有不少計(jì)算問題特別適合這種計(jì)算模式。關(guān)于CUDA的相關(guān)應(yīng)用可參閱參考文獻(xiàn)[2]和參考文獻(xiàn)[3]。
 CUDA加速的分形繪制的研究并不多,其中NVIDIA的SDK提供了基于CUDA的加速M(fèi)andelbrot、Julia集生成的實(shí)例[4]。其算法的基本過程是平面圖的每一個(gè)像素對(duì)應(yīng)CUDA并行計(jì)算中的一個(gè)線程,該并行方法能比CPU上的普通方法加速幾十倍。參考文獻(xiàn)[5]描述了基于CUDA的并行分形IFS點(diǎn)遞歸繪制算法。
本文的研究是利用GPU的計(jì)算能力加速分形火焰繪制,使之能更加實(shí)用。
1 背景及相關(guān)知識(shí)
1.1 經(jīng)典的迭代函數(shù)系統(tǒng)

 數(shù)學(xué)中,迭代函數(shù)系統(tǒng)是構(gòu)造分形的一種方法,由此產(chǎn)生的結(jié)構(gòu)是自相似的。IFS分形可以是任何維度[6],但通常在二維空間中計(jì)算和繪制。分形由一些自身的拷貝聯(lián)合構(gòu)成,每個(gè)拷貝根據(jù)一個(gè)函數(shù)來作變換。函數(shù)通常是收縮的,意思是它們會(huì)使圖像點(diǎn)更靠近,使形狀更小。因此,IFS分形的形狀由一些可能重疊的自身拷貝組成,其中每個(gè)拷貝也由自身的一些拷貝組成,如此無限下去。這也是自相似分形特性的來源。

1.2 分形火焰簡(jiǎn)介
 分形火焰是1992年由DRAVES S提出的[7],它可以說是分形迭代函數(shù)的一種形式,但有以下3方面不同:
 (1)迭代時(shí)使用非線性函數(shù)而不是仿射變換;
 (2)色調(diào)映射顯示是對(duì)數(shù)密度而不是線性的;
 (3)顏色是根據(jù)結(jié)構(gòu)(即通過采取的遞歸路徑)決定的,而不是采用單色的或根據(jù)點(diǎn)的密度決定。
色調(diào)映射和著色旨在顯示盡可能多的分形的詳細(xì)信息,這樣通常會(huì)產(chǎn)生更美觀、更絢麗的圖像。
分形火焰的每個(gè)函數(shù)具有如下的形式:

3 分形火焰并行算法詳細(xì)描述
 在分形火焰直方圖生成階段,使用串行混沌游戲算法的過程是:選擇一個(gè)點(diǎn)開始,并根據(jù)概率挑選一個(gè)函數(shù)將該點(diǎn)代入來計(jì)算,得到下一個(gè)點(diǎn),然后用新點(diǎn)繼續(xù)進(jìn)行下一次迭代,如此不斷迭代下去。
 本文提出的并行算法是傳統(tǒng)混沌游戲算法的擴(kuò)展:算法開始時(shí)選擇n個(gè)而不是一個(gè)起始點(diǎn),參考文獻(xiàn)[9]說明了選擇n個(gè)起始點(diǎn)的方法開始迭代和選擇1個(gè)起始點(diǎn)開始迭代得到的結(jié)果一致。并行的實(shí)現(xiàn)是通過將一個(gè)線程分配給某一個(gè)起始點(diǎn),n個(gè)線程同時(shí)分別獨(dú)立執(zhí)行混沌游戲。通過這種方式,并行算法提高了速度。
 當(dāng)然,上述只是并行算法的主要思想,因?yàn)榉中位鹧姹容^復(fù)雜,具體的算法還有很多細(xì)節(jié)問題需要考慮。實(shí)現(xiàn)主要包括3個(gè)在GPU上并行執(zhí)行的內(nèi)核函數(shù):warm_up、iterate_batch和output_for_rending。與常見的串行混沌游戲一樣,在warm_up函數(shù)中的每個(gè)線程選擇一個(gè)起始點(diǎn),迭代十幾次,但只保持最后的結(jié)果,放棄前面迭代得到的值。warm_up函數(shù)計(jì)較簡(jiǎn)單,沒必要描述實(shí)現(xiàn)細(xì)節(jié)。Iterate_batch函數(shù)接收warm_up函數(shù)生成的點(diǎn)并迭代數(shù)十或數(shù)百遍,然后計(jì)算得到的每個(gè)點(diǎn)的直方圖。Iterate_batch函數(shù)的核心思想就是n個(gè)點(diǎn)同時(shí)迭代的并行混沌游戲算法,也是生成分形火焰圖形的核心。其偽代碼如下:
//輸入:ractalInfo存儲(chǔ)該fractal flame信息的結(jié)構(gòu)體;iterPosStateBuffer是warm_up后存儲(chǔ)點(diǎn)位置的數(shù)組;iterColorStateBuffer是warm_up后存儲(chǔ)顏色的數(shù)組;randBuffer用戶指定的每個(gè)函數(shù)選擇的概率
//輸出:accumBuffer存儲(chǔ)累計(jì)直方圖信息,該信息在繪制階段會(huì)用到
Iterate_batch()
{lid=threadIdx.x;
gid=(blockIdx.x*blockDim.x+threadIdx.x);
pos=iterPosStateBuffer[gid];
color=iterColorStateBuffer[gid];
for(iter=0;iter<ITERCOUNT;iter++)
{fIndex=chooseRandomBranch(randBuffer+lid,fractalInfo);//根據(jù)randBuffer選擇一個(gè)函數(shù)
iterate(&pos,&color,fIndex,fractalInfo);
//迭代生成新的點(diǎn)和顏色
screenPos=getPos(fractalInfo,pos);
//將計(jì)算得到的點(diǎn)的位置轉(zhuǎn)換為屏幕上點(diǎn)的位置
calhistogram(color,screenPos,accumBuffer);
//計(jì)算每個(gè)屏幕點(diǎn)的統(tǒng)計(jì)直方圖和顏色,保存到accumBuffer
}}
Output_for_rending函數(shù)主要完成圖形繪制功能,它接收從iterate_batch傳遞來的直方圖信息,作色調(diào)映射和伽瑪校正,并將每個(gè)像素的最終顯示顏色放在outputBuffer。其偽代碼如下:
//輸入:ractalInfo存儲(chǔ)了該fractal flame信息的結(jié)構(gòu)體;
accumBuffer是從Iterate_batch函數(shù)傳遞過來的,保存了直方圖信息
//輸出:outputBuffer保存了最終數(shù)據(jù),用于通過OpenGL繪制最終圖形
output_for_rending(ractalInfo,accumBuffer)
{x=(blockIdx.x*blockDim.x+threadIdx.x);
y=(blockIdx.y*blockDim.y+threadIdx.y);
pix=tonemap(fractalInfo,accumBuffer,x,y);
//根據(jù)直方圖信息,使用色調(diào)映射和伽瑪校正產(chǎn)生實(shí)際
//顯示的顏色
setoutput(outputBuffer,x,y,pix);
//將每個(gè)像素的顯示顏色保存到outputBuffer
}
 根據(jù)實(shí)驗(yàn)觀察,圖像繪制階段大約消耗總時(shí)間的70%,為了加速繪制過程,需要利用CUDA和Open GL互操作[10]。互操作的基本方法是將Open GL緩沖區(qū)映射到CUDA的內(nèi)存空間。CUDA用于計(jì)算和數(shù)據(jù)生成,Open GL用來繪制像素或頂點(diǎn),因?yàn)樗鼈児蚕硗粌?nèi)存空間,無需在CPU內(nèi)存和GPU內(nèi)存間移動(dòng)數(shù)據(jù),所以速度非常快。調(diào)用output_for_rending函數(shù)之前,需要做一些工作使CUDA和Open GL相關(guān)聯(lián),并設(shè)置outputBuffer與CUDA和Open GL共享。
4 對(duì)基本并行算法的改進(jìn)
 分形火焰的基本并行實(shí)現(xiàn)是每個(gè)線程對(duì)應(yīng)一個(gè)數(shù)據(jù)點(diǎn),各線程隨機(jī)選擇一個(gè)變換函數(shù)計(jì)算,用計(jì)算得到的數(shù)據(jù)點(diǎn)替換原數(shù)據(jù)點(diǎn),這個(gè)過程重復(fù)m次。該并行算法需要每個(gè)線程根據(jù)隨機(jī)數(shù)選擇變換函數(shù),因而會(huì)導(dǎo)致CUDA嚴(yán)重的分支分歧[1]。變換函數(shù)越多,越有可能繪制出更有趣的圖像,因此,應(yīng)該盡可能允許有更多的變換函數(shù)。然而用到的變換函數(shù)越多,分支分歧問題就越嚴(yán)重。
 本文提出一種改進(jìn)的算法,該算法通過預(yù)分類能移除隨機(jī)選擇函數(shù),可以消除分支分歧。具體方法是通過隨機(jī)數(shù)據(jù)訪問來替代隨機(jī)選定函數(shù),即每個(gè)線程分配一個(gè)固定的函數(shù),在每次迭代中隨機(jī)選擇一個(gè)數(shù)據(jù)點(diǎn)。這種選擇是通過數(shù)據(jù)和線程索引之間的一個(gè)隨機(jī)雙射函數(shù)映射實(shí)現(xiàn)的。通過這種方法,指令能夠以最佳方式靜態(tài)地分配給線程,預(yù)先計(jì)算好固定的置換,它們不依賴于動(dòng)態(tài)數(shù)據(jù)。每個(gè)線程使用分配的函數(shù),通過置換索引間接地訪問數(shù)據(jù)數(shù)組,然后將該函數(shù)的計(jì)算結(jié)果寫回。為了避免寫訪問時(shí)的條件競(jìng)爭(zhēng),要么結(jié)果數(shù)據(jù)寫回讀取的位置,要么需要第二個(gè)數(shù)組來存儲(chǔ)數(shù)據(jù)點(diǎn),每次迭代挑選一個(gè)新的置換。
 上述的過程需要通過創(chuàng)建多個(gè)包含數(shù)據(jù)點(diǎn)索引及隨機(jī)數(shù)的數(shù)組來進(jìn)行置換,隨機(jī)數(shù)可由Mersenne Twister方法生成[11],數(shù)組根據(jù)隨機(jī)數(shù)排序。所有的變換函數(shù)是在一個(gè)大的switch語句中實(shí)現(xiàn)的。用一個(gè)結(jié)構(gòu)描述變化索引及縮放因子,該結(jié)構(gòu)存儲(chǔ)在常量?jī)?nèi)存中。
5 實(shí)驗(yàn)結(jié)果
 上述并行分形火焰算法是在NVIDIA GeForce GTX9800+上實(shí)現(xiàn)的,該GPU有128個(gè)頻率為1.836 GHz的流處理器,分為16個(gè)SM,896 MB顯存,裝配在CPU為Intel Core 2 Duo E7400 2.8 GHz,內(nèi)存為1 GB的計(jì)算機(jī)上。為了作對(duì)比,同時(shí)實(shí)現(xiàn)了CPU上運(yùn)行的串行版本。實(shí)驗(yàn)結(jié)果表明,本文提出的基于CUDA的基本并行算法在大多數(shù)情況下比CPU上的串行算法快了約15倍,能做到分形火焰實(shí)時(shí)繪制。
 實(shí)驗(yàn)表明,消除分支分歧的算法與基本并行算法相比,基本并行算法的運(yùn)行時(shí)間與變換函數(shù)的數(shù)目呈現(xiàn)線性增長(zhǎng)關(guān)系,而消除分支分歧的改進(jìn)算法運(yùn)行時(shí)間恒定,多數(shù)情況下性能都比基本算法高。在變換函數(shù)數(shù)目為20個(gè)時(shí),消除了分支分歧的算法比基本算法快了約兩倍。
 本文提出了一種并行的分形火焰算法,該方法利用了GPU的計(jì)算能力及CUDA和Open GL互操作方式,速度快到足夠?qū)崟r(shí)高幀速率繪制。該算法使得在生成分形火焰圖形時(shí),能實(shí)時(shí)更改參數(shù)并觀察繪制效果,可以有效地幫助加深對(duì)迭代函數(shù)系統(tǒng)的認(rèn)識(shí)。本文方法對(duì)其他圖像實(shí)時(shí)繪制的應(yīng)用也有很高的參考價(jià)值。
參考文獻(xiàn)
[1] NVIDIA Corporation. NVIDIA CUDA programming guide version 3.2[EB/OL]. (2011-03). http://developer.nvidia.com/cuda.
[2] Hwu Wenmei, RODRIGUES  C, RYOO S, et al. Compute unified device architecture application suitability[J]. Computing in Science and Engineering, 2009,11(3): 16-26.
[3] Hwu Wenmei.  GPU computing gems[M]. New York: Morgan Kaufmann,2011.
[4] GRANGER M. Mandelbrot CUDA SDK code samples[EB/OL].http://developer.nvidia.com/cuda.2011-03.
[5] KWANIEWSKI B. CUDA based parallel version of point recursive rendering algorithm of IFS attractor[C]. 12th International PhD Workshop OWD, 2010: 23-26.
[6] WHITELAW M. Metacreation: art and artificial life[M]. MIT Press, 2004.
[7] DRAVES S, RECKASE E. The fractal flame algorithm[EB/OL].http://flam3.com/flame draves.pdf. 2008-11.
[8] NIKIEL S. Iterated function systems for real-time image synthesis[M]. London: Springer, 2007.
[9] DUTIL N. Construction of fractal objects with iterated function systems[EB/OL].(2000-10).http://www.cs.mcgill.ca/~ndutil/project.pdf .
[10] 劉進(jìn)鋒,郭雷.CUDA和OpenGL互操作的實(shí)現(xiàn)及分析[J].微型機(jī)與應(yīng)用,2011(23):40-43.
[11] MATSUMOTO M, NISHIMURA T. Mersenne twister: a 623-dimensionally equidistributed uniform pseudorandom number generator[J]. ACM Transactions on Modeling and Computer Simulation, 1998,8(1):3-30.

此內(nèi)容為AET網(wǎng)站原創(chuàng),未經(jīng)授權(quán)禁止轉(zhuǎn)載。
主站蜘蛛池模板: 99久热re在线精品99 6热视频 | 日韩91在线 | 欧美又大又色又爽bbbbb片 | 欧美视频在线观看一区 | 日本网站免费 | 无码熟妇αⅴ人妻又粗又大 | 午夜精品成人 | 色偷偷导航 | 用力使劲高潮了888av | av手机免费在线观看 | 日韩欧美高清片 | 国产婷婷色一区二区三区在线 | 久草天堂 | 国产乱人伦中文无无码视频试看 | 两个黑人大战嫩白金发美女 | 国产精品乱码一区二区三区视频 | 久久精品女人的天堂av | 777奇米888色狠狠俺也去 | 能免费看av的网站 | 91午夜少妇三级全黄 | 久久黄色一级片 | 自拍偷拍第3页 | 肉视频在线观看 | 免费毛片网 | 亚洲综合国产一区二区三区 | 亚洲国产成人丁香五月激情 | 日本一区二区三区免费视频 | 国产日韩一区二区在线 | 亚洲精品久久久久久久久久吃药 | 果冻传媒色av国产在线播放 | 男人边做边吃奶头视频 | 国产精品96久久久久久久 | 天天色综合久久 | 亚洲国产无线乱码在线观看 | 日韩在线视频一区二区三区 | 精品国产不卡一区二区三区 | 日本高清视频免费看 | 国产精品久久国产精品99 gif | 欧美韩日精品 | 美女又爽又黄视频毛茸茸 | 欧美片一区二区三区 | 成人妇女免费播放久久久 | 国产精品乱 | 香蕉黄色网 | 性无码一区二区三区在线观看 | 超碰caopeng| 色婷婷精品视频 | 亚洲福利视频在线 | 久久99久久99精品免观看粉嫩 | 国产精品18久久久久久麻辣 | 香蕉成人在线视频 | 亚洲视频三区 | 黑人巨茎美女高潮视频 | 忍不住的亲子中文字幕 | 国产娇小性色xxxxx视频 | 精品蜜桃一区二区三区 | 大桥未久亚洲精品久久久强制中出 | 成人免费高清 | 日韩六九视频 | 中文字幕av手机版 | 人人玩人人添人人澡 | 国产精品国产三级在线专区 | 国产精品久久久久久久午夜片 | 少妇性荡欲视频 | 正在播放国产真实哭都没用 | 免费在线观看视频a | 亚洲v国产v欧美v久久久久久 | 亚洲一区二区三区四区五区六 | 色一情一乱一乱一区免费网站 | 日本www一道久久久免费 | 国产精品18hdxxxⅹ在线 | 国产无遮无挡120秒 国产无遮掩 | 女女女女女裸体处开bbb | 国产在线激情视频 | 97超碰超碰 | 成人妇女淫片aaaa视频 | 黄色三级在线观看 | 人人干在线视频 | 日本a级片网站 | 成人做爰69片免费观看 | 欧美福利在线视频 | 国产三级播放 | 亚洲自拍网站 | 五月香蕉网| 强行处破女系列中文字幕 | 久热网站 | 特黄三级男人添女人下面 | 日本不卡视频在线 | 天干啦夜天干天干在线线 | 亚洲天堂网站在线 | 免费精品一区 | 国产精品久久久久久久 | 免费国精产品自偷自偷免费看 | 日韩欧美一区在线 | 韩日精品视频在线观看 | 黑白配在线观看免费观看 | 97丨九色丨蜜臀 | 深夜视频在线观看免费 | 天堂va久久久噜噜噜久久va | 成人在线视频免费观看 | 蜜臀av无码人妻精品 | 日本三级韩国三级三级a级按摩 | 精品人人妻人人澡人人爽人人 | 性一交一乱一色一免费无遮挡 | 免费日韩av | 国产精品免费无遮挡无码永久视频 | 视频国产精品 | 久久久不卡国产精品一区二区 | 欧美午夜精品理论片 | 99热这里只有精品9 99热这里只有精品99 | 爱情岛免费永久网站 | 国产女同疯狂作爱系列 | 亚洲女同一区二区 | 日韩欧美激情兽交 | 精品探花 | 久久九九网站 | 欧美视频免费在线观看 | 91禁在线动漫 | www一区| 欧美第一页浮力影院 | 三级三级三级三级 | 国产情侣激情在线对白 | 亚洲欧美日韩系列 | 日本真人做爰免费的视频 | 国产不卡一区 | 今夜无人入睡在线观看 | 九色中文 | 国产成人精品日本亚洲专区 | 久久久久久久久精 | 尤物在线视频观看 | 麻豆国产人妻欲求不满 | 含紧一点h边做边走动免费视频 | 九月婷婷人人澡人人添人人爽 | 欧美群妇大交群 | 免费看欧美中韩毛片影院 | a v免费视频| 天天天天天干 | 久久国产劲爆∧v内射 | 欧美日韩一区二区视频在线观看 | 极品少妇xxxx精品少妇偷拍 | 成熟老妇女视频 | 又色又爽又黄的美女裸体网站 | 在线超碰91 | 一级片观看 | 一本色道久久hezyo无码 | 午夜剧院免费观看 | 乱淫交换粗大多p | 国产女女| 日韩精品一区二区三区蜜臀 | 久久精品人人做人人爽97 | 一本大道伊人av久久综合 | 国产精品国产成人国产三级 | 51久久久 | 暖暖日本在线观看免费 | 手机看片1024国产 | 欧美亚洲国产一区二区三区 | 97在线观看| 久热这里只有精品6 | 国产资源在线视频 | 久久精品成人免费国产片桃视频 | 毛片福利视频 | 国产精品妇女一二三区 | 欧美人与性动交xxⅹxx | 99热这里只有精品99 | 不卡av在线播放 | 牲交欧美兽交欧美 | 成人aaaa | 精产国品一二三区 | 天天躁夜夜躁很很躁麻豆 | 日韩av免费在线观看 | 日韩在线一区二区三区 | 国产成人自拍网站 | 91一区二区视频 | 国内精品小视频 | 国产天美传媒性色av | 性仑少妇av啪啪a毛片 | 亚洲色成人一区二区三区小说 | 亚洲欧美日韩中文在线制服 | 风间由美不戴奶罩邻居勃起av | 超碰97人人爱 | 国产人妻精品无码av在线 | 免费看黄色的视频 | 亚洲第一成人区av桥本有菜 | 国产精品久久久久久久久久综合 | 98色| 精品aⅴ一区二区三区 | 国产对白叫床清晰在线播放图片 | 午夜精品一区二区三区aa毛片 | 国产中文字幕一区二区 | 欧美一级在线播放 | av不卡中文字幕 | 国产亚洲精品久久久久四川人 | 愉拍自拍第43页免费 | 巨乳动漫美女 | 亚洲福利一区二区三区 | 亚洲欧美不卡 | 日本欧美一级片 | 日本a在线天堂 | 国产在线精品成人一区二区 | 91成人免费在线观看 | 在线亚洲人成电影网站色www | 欧美日韩一区二区区别是什么 | www.黄色av | 黄色链接视频 | 国产成人二区 | 国语对白乱妇激情视频 | 日本无遮挡真人祼交视频 | 嘿嘿射在线| 午夜毛片在线 | 国产亚洲二区 | 自拍偷拍亚洲欧洲 | 99久热在线精品996热是什么 | 亚洲福利小视频 | 国产亚洲精品日韩在线tv黄 | 亚洲a∨国产av综合av下载 | 日本免费www| 亚洲精品一区二区 | 一本一道久久久a久久久精品蜜臀 | av最新版天堂资源在线 | av在线免费观看网址 | 69视频在线看 | 国产内射999视频一区 | 欧美日比视频 | 中文在线а√天堂 | 中文久久乱码一区二区 | 黑人玩弄出轨人妻松雪 | 久久久久久三区 | 欧美日韩视频免费观看 | 日本色中色 | 免费看黄色一级毛片 | 女女百合高h喷汁呻吟视频 女女百合国产免费网站 | 成人特级片 | a一区二区三区乱码在线 | 欧洲 | 久久亚洲精品成人av无码网站 | 色七七久久| 性感av在线 | 看免费的毛片 | 国产精品精品视频一区二区三区 | 蜜桃视频黄色 | 天天躁日日躁狠狠躁超碰97 | 欧美在线xxxx | 久久久久青草大香综合精品 | 最近2019年好看中文字幕视频 | 国产在线视频网站 | 91久久一区 | 性一交一乱一乱一视频 | 精品乱码久久久久久中文字幕 | 日韩免费黄色 | 少妇自摸视频 | 国产999在线观看 | 超碰人人91 | 美女18免费视频 | 四虎激情| 中文字幕av无码不卡免费 | 日韩av官网| 国产精品久久免费视频 | 亚洲免费一区二区 | 午夜激情综合网 | 久久人人爽爽爽人久久久 | 亚洲欧美日本另类 | 日日噜噜夜夜狠狠va视频v | 亚洲久草视频 | 爽爽影院在线 | 亚洲成av人的天堂在线观看 | 亚洲自拍偷拍图 | 韩国午夜激情 | 国产视频麻豆 | 综合久久中文字幕 | 凹凸av在线 | 一边吃奶一边做动态图 | 亚洲成色在线 | 浓毛老太交欧美老妇热爱乱 | 国产免费拔擦拔擦8x网址 | 中文字幕一区二区三区人妻少妇 | 国产xxxx做受性欧美88 | ass亚洲曰本人体私拍ass | 一个人看的免费高清www视频 | 久久久中文字幕日本无吗 | 欧洲成人综合 | 狠狠爱综合 | 国产嫩草影院在线观看88 | 国产精品一二三区成毛片视频 | 三级av在线免费观看 | 桃色视频.m3u8 | 日本韩国在线播放 | 成人免费看片又大又黄 | 欧美成人一区二免费视频软件 | 91蜜桃婷婷狠狠久久综合9色 | aav在线| 天堂va在我观看 | 在线中文视频 | 婷婷国产一区二区三区 | 国产成人综合自拍 | 在线黄色免费 | 1000部拍拍拍18勿入在线看 | 黄色美女毛片 | 日本在线激情 | 免费观看又污又黄的网站 | 91tv亚洲精品香蕉国产一区 | 嘿咻视频在线观看 | 国产精品自在线拍国产手青青机版 | 国产91清纯白嫩初高中在线观看 | 成人久久久 | 免费看片在线观看www | 女人黄色片 | 国产成人久久精品麻豆二区 | 国产丰满精品伦一区二区三级视频 | 天堂在线一区二区 | 欧美 国产 亚洲 卡通 综合 | 国产在线视频99 | 欧美亚洲综合另类 | 消息称老熟妇乱视频一区二区 | 久久久久国产免费 | 亚洲国产精品成人一区二区在线 | 136av导航| 处破痛哭a√18成年片免费 | 亚洲一区二区三区四区五区六 | 天天av天天翘天天综合网 | 国产又黄又硬又湿又黄 | japanese丰满少妇最高潮 | 亚洲不卡中文字幕 | 狠狠亚洲婷婷综合色香五月 | 情一色一乱一欲一区二区 | 色噜噜亚洲男人的天堂 | 一级少妇精品久久久久久久 | 欧美爱爱网址 | 外国a级片 | 免费污视频在线观看 | 高h肉放荡爽全文寂寞少妇 高h肉各种姿势g短篇np视频 | 久久久精品视频免费看 | www.com久久| 欧美日韩亚洲精品瑜伽裤 | 老司机在线精品视频 | 国产精品aaaa | 91p九色 | 久久99精品久久久久麻豆 | 国产精品免费一区二区三区都可以 | 国产欧美一区二区在线观看 | 看国产一级毛片 | 国产黄色网络 | 久久久精品小视频 | 成人网久久 | 国产老少配bbbb搡bbbb | 日韩一级二级 | 91精品国产91久久久久久久久久久久 | 91人人草| 欧美视频黄| 麻豆精品国产入口 | 亚洲欧美久久 | 亚洲人成影院在线无码按摩店 | 国产精品精品 | 六姐妹在线观看 | 亚洲热在线观看 | 午夜在线观看视频网站 | 少妇尝试黑人粗吊受不了 | 99久久精品无码一区二区毛片 | 欧美乱妇无乱码大黄a片 | 久久精品99久久 | 亚洲videos| 精品成人在线视频 | 国产三级精品一区二区三区视频 | 国产 日韩 欧美 成人 | 久久久久久久香蕉 | 亚欧洲精品在线 | 国产大片内射1区2区 | 国产成人无码一区二区在线播放 | 国产精品久久久久久亚洲伦 | av无码精品一区二区三区 | 中国中文字幕伦av在线看片 | 污视频在线观看免费网站 | 俄罗斯丰满熟妇hd | 麻豆三级 | 亚洲国产一区视频 | 国产午夜精品福利视频 | 欧美成人午夜精品 | 在线视频观看免费视频18 | 国产精品午夜无码av体验区 | 国产主播福利在线 | 久久久综合视频 | 日产av在线播放 | 欧美精品一级 | 麻豆精品a∨在线观看 | 久久在线精品 | 亚洲国产欧美日本视频 | 日日射日日干 | 亚洲午夜免费视频 | 欧美性猛交7777777 | 激情总合网 | 欧美一级爽aaaaa大片 | 国产youjizz| 亚洲国产成人005 | 欧美一级欧美三级在线观看 | 妩媚尤物娇喘无力呻吟在线视频 | 国产性天天综合网 | 国产免费无遮挡吸乳视频 | 久久久久久久久久久久网站 | 天天插夜夜爽 | 国产毛片久久久久久美女视频 | 亚洲h视频在线 | 国产乱色国产精品播放视频 | 国产九九九 | h视频亚洲 | 中文字幕乱码在线 | 中文字幕av一区二区三区高 | 国产成人福利av综合导航 | 狠狠色婷婷狠狠狠亚洲综合 | 日韩精品视频在线一区 | 日韩精品一区二区三区中文 | 国产精品免费福利久久 | 亚洲va在线va天堂xxxx中文 | 午夜精品久久久久久久久 | 男人下部进女人下部视频 | 狠狠精品久久久无码中文字幕 | 成人网页在线观看 | 91精品国产亚一区二区三区老牛 | 999精品视频在线观看 | av夜夜| 亚洲第一无码av无码专区 | 国产三级久久久精品麻豆三级 | 丝袜 亚洲 另类 欧美 重口 | 久久久久久人妻一区精品 | av在线不卡免费 | 国产片网址 | 国产内射999视频一区 | 特级淫片aaaaaaa级附近的 | h网站免费在线观看 | 一级做a爰片性色毛片精油 一级做a爰片性色毛片视频停止 | 成人手机在线免费视频 | 亚洲天堂区 | 一级片在线免费观看 | 欧美激情视频一区二区 | 欧美做受69 | 国产成人精品亚洲男人的天堂 | 亚洲精品一区二区三区在线观看 | 日本亚洲一区 | 免费在线黄色av | 亚洲精品无码久久久久久 | 国产熟睡乱子伦视频 | 人人干超碰| 已婚少妇露脸日出白浆 | 亚洲天堂美女视频 | 国产a三级久久精品 | 国产极品美女在线精品图片 | 国产成人久久婷婷精品流白浆 | 日韩高清黄色 | 亚洲玉足av久久影视 | 99久久伊人精品综合观看 | 欧美在线视频免费播放 | 女人高潮流白浆视频 | 国产一区二区精品久久 | 贱奴的sm(高h调教) | 亚洲综合av在线在线播放 | 亚洲一区精品无码 | 久久99热久久99精品 | 亚洲乱码国产乱码精品精大量 | 国产第三页 | 91色区 | 欧美一区二区视频在线观看 | 国产一级久久 | 国产情侣激情在线对白 | 又黄又爽又色qq群 | 国产美女免费看 | 国产无遮挡猛进猛出免费软件 | 日韩少妇中文字幕 | 色综合中文 | 美女胸18大禁视频网站 | 欧美在线www | 潮喷无码正在播放 | 四虎永久在线精品免费播放 | 亚洲美女爱爱 | 欧美一级激情 | 欧美乱大交xxxxx潮喷 | 欧美色涩| yy6080亚洲精品一区 | 天天干天天色综合 | 日韩精品视频在线观看一区二区 | 久久久久久www | 日本又色又爽又黄的a片吻戏 | 亚洲高清网 | 日韩大片在线观看 | 国精产品一区一区三区mba下载 | 欧美一a一片一级一片 | 九九久久免费视频 | 中文字幕在线免费看线人 | 欧美爱爱视频网站 | 色男人av| 清纯粉嫩极品夜夜嗨av | 久久久久久久久免费视频 | 日本一区二区不卡在线 | 欧美老肥婆性猛交视频 | 国产馆在线观看 | 国产91玉足脚交在线播放 | 亚洲男同视频网站 | 亚洲最新网址 | 国产无遮挡aaa片爽爽 | 在线观看aa| 日韩综合中文字幕 | 中文在线天堂网 | 手机在线看片 | 久久久黄色一级片 | 一 级 黄 色 片免费网站 | 2019精品手机国产品在线 | 人妻激情偷乱视频一区二区三区 | 亚洲精品中文字幕 | 男人巨茎大战欧美白妇 | 国产手机在线精品 | 欧美成人看片一区二区三区尤物 | 日本男女激情视频 | 深夜在线免费观看 | 免费在线观看黄色av | 日本免费福利视频 | 国产无遮挡一区二区三区毛片日本 | 国产嫩草在线观看 | 久久国产精品综合 | 波多野结衣福利视频 | 国产裸体丰满白嫩大尺度尤物可乐 | 亚洲精品久久久蜜桃 | 欧美不卡一区二区三区 | 欧美做爰爽爽爽爽爽爽 | 国产cd人妖ts在线观看 | 午夜爱爱免费视频体验区 | 性xxxxx大片免费视频 | 亚洲xx在线 | 国产69久久精品成人看 | 国产韩国精品一区二区三区 | 亚洲精品在线不卡 | 国产色视频一区二区三区qq号 | av天堂久久天堂色综合 | 国产乱子经典视频在线观看 | 欧美三级韩国三级日本三斤 | 狠狠色狠狠色很很综合很久久 | 91无人区乱码卡一卡二卡 | 五月天激情小说 | 国产精品成人va在线播放 | 色婷婷狠狠干 | 久久福利影视 | 欧美xxxx888| 亚洲 日本 欧美 中文幕 | 黄色成人免费网站 | 精品久久久久久亚洲精品 | 人人爽久久久噜噜噜婷婷 | 国产浮力视频 | 亚洲网站在线 | 亚洲一卡二卡 | 国产亚洲日韩一区二区三区 | 91亚洲国产成人精品一区二三 | 午夜精品在线播放 | 精品国产sm最大网站 | 深夜老司机福利 | 无码人妻品一区二区三区精99 | 久久一卡二卡三卡四卡 | 美女一区 | 亚洲制服丝袜精品久久 | 亚洲图色视频 | 91tv亚洲精品香蕉国产一区 | 国产精品久久久久久久久久久久冷 | 国产一区二区三区四区精 | 四虎国产精品永久在线国在线 | av一二三四 | 国产成人一区二区三区别 | 欧美巨猛xxxx猛交黑人97人 | 97色伦图片 | 红桃视频 国产 | 国产情趣视频 | 午夜少妇性高湖久久久久 | 亚洲蜜芽在线精品一区 | 国产女人与拘做视频免费 | 99国产精品欲 | 国产美女精品视频国产 | 欧美性高潮| 天天玩天天干天天操 | 91国偷自产一区二区使用方法 | 黄色成年网站 | 亚洲成人综合视频 | 精品一卡二卡三卡四卡 | 男女午夜激情视频 | 玩弄japan白嫩少妇hd | 日本又黄又硬又爽的大片 | 午夜艹逼 | 久久99精品久久久久久 | 国产精品777 | 国产亚洲精品一区二区三区 | 在线视频你懂得 | 国产精品一二三四五区 | 亚洲美女综合网 | 伊人久久大香线蕉无码 | 日韩av在线播 | 欧美成人aaaaaaaa免费 | 成年人视频网 | a视频| 久久老司机 | 综合欧美一区二区三区 | 青青草黄色| 久久艹逼视频 | 成人爽a毛片在线视频 | 国产精品主播一区二区 | 成人不卡 | 俺也来俺也去俺也射 | 国产露脸150部国语对白 | 国产精品国产三级国产有见不卡 | 黄色激情在线观看 | 国产精品夜夜春夜夜爽久久小 | 中文字幕一区三级久久日本 | 浪潮av激情高潮国产精品香港 | 国产剧情一区在线 | 久久996re热这里只有精品无码 | 亚洲第一免费网站 | 麻豆精品影院 | 毛片毛片毛片毛片毛片毛片毛片毛片毛片毛片 | 在线观看视频国产 | 久久国产资源 | 中国黄色网页 | 无码国产一区二区三区四区 | 森泽佳奈在线播放 | 国产女主播高潮在线播放 |