亚洲免费av电影一区二区三区,日韩爱爱视频,51精品视频一区二区三区,91视频爱爱,日韩欧美在线播放视频,中文字幕少妇AV,亚洲电影中文字幕,久久久久亚洲av成人网址,久久综合视频网站,国产在线不卡免费播放

        ?

        基于GPU的受限玻爾茲曼機(jī)并行加速

        2016-09-14 09:17:14張立民范曉磊
        電子設(shè)計(jì)工程 2016年2期
        關(guān)鍵詞:模型設(shè)計(jì)

        張立民,劉 凱,范曉磊

        (1.海軍航空工程學(xué)院 山東 煙臺 264001;2.第二炮兵工程大學(xué) 士官職業(yè)技術(shù)教育學(xué)院,山東 青州 262500)

        基于GPU的受限玻爾茲曼機(jī)并行加速

        張立民1,劉 凱1,范曉磊2

        (1.海軍航空工程學(xué)院 山東 煙臺 264001;2.第二炮兵工程大學(xué) 士官職業(yè)技術(shù)教育學(xué)院,山東 青州 262500)

        為針對受限玻爾茲曼機(jī)處理大數(shù)據(jù)時存在的訓(xùn)練緩慢、難以得到模型最優(yōu)的問題,提出了基于GPU的RBM模型訓(xùn)練并行加速方法。首先重新規(guī)劃了對比散度算法在GPU的實(shí)現(xiàn)步驟;其次結(jié)合以往GPU并行方案,提出采用CUBLAS執(zhí)行訓(xùn)練的矩陣乘加運(yùn)算,設(shè)計(jì)周期更長、代碼更為簡潔的Tausworthe113和CLCG4的組合隨機(jī)數(shù)生成器,利用CUDA拾取紋理內(nèi)存的讀取模式實(shí)現(xiàn)了Sigmoid函數(shù)值計(jì)算;最后對訓(xùn)練時間和效果進(jìn)行檢驗(yàn)。通過MNIST手寫數(shù)字識別集實(shí)驗(yàn)證明,相較于以往RBM并行代碼,新設(shè)計(jì)的GPU并行方案在處理大規(guī)模數(shù)據(jù)集訓(xùn)練上優(yōu)勢較為明顯,加速比達(dá)到25以上。

        受限玻爾茲曼機(jī);GPU;CUDA;加速比;并行加速

        基于能量模型的受限玻爾茲曼機(jī) (Restricted Boltzmann Machine,RBM)[1]以其簡單的人工神經(jīng)網(wǎng)絡(luò)形式和快速的學(xué)習(xí)算法,越來越受到機(jī)器學(xué)習(xí)的關(guān)注,目前已經(jīng)廣泛應(yīng)用于數(shù)據(jù)降維、語音識別、3D物體識別、圖像轉(zhuǎn)換以及高維時間序列建模等機(jī)器學(xué)習(xí)問題,進(jìn)而催生出機(jī)器學(xué)習(xí)一個新的領(lǐng)域,深度學(xué)習(xí)[2],并使之得到了迅猛發(fā)展。

        但隨著大數(shù)據(jù)時代的到來,模型訓(xùn)練需要的數(shù)據(jù)越來越多,單純通過CPU進(jìn)行計(jì)算已經(jīng)逐漸失去優(yōu)勢,自2008年P(guān)aprotski等人研究了通過CUDA對RBM訓(xùn)練的加速方法[3],并且取得了較好效果之后,利用GPU的高速運(yùn)算能力以加快RBM訓(xùn)練,從而構(gòu)建深度學(xué)習(xí)模型,已經(jīng)成為了RBM訓(xùn)練的研究熱點(diǎn)[4-6]。目前RBM的CUDA加速方法主要體現(xiàn)在兩種方式,一種是基于相對成熟的CUDA并行庫,例如CUBLAS庫,另一種是根據(jù)RBM結(jié)構(gòu)特點(diǎn)和訓(xùn)練步驟設(shè)計(jì)Kernel函數(shù),對于RBM訓(xùn)練過程中目標(biāo)不同的矩陣運(yùn)算定義不同的Block和編寫Kernel函數(shù)。第一種方法的優(yōu)點(diǎn)在于,加速方式較為通用,不僅適用于RBM訓(xùn)練,還可應(yīng)用于其他與矩陣運(yùn)算相關(guān)的模型訓(xùn)練;代碼運(yùn)算穩(wěn)定,較為健壯,不容易出現(xiàn)溢出問題;但是缺點(diǎn)在于針對性不強(qiáng),不能高效利用GPU計(jì)算資源;第二種方法的優(yōu)點(diǎn)是能夠針對RBM模型,加速比好,但是缺點(diǎn)為設(shè)計(jì)較為復(fù)雜,擴(kuò)展性不強(qiáng)。因此,設(shè)計(jì)一種能夠結(jié)合兩種方式的RBM并行加速方法,既保證算法的有效性,也使得擴(kuò)展性增強(qiáng)。

        1 受限玻爾茲曼機(jī)

        1.1受限玻爾茲曼機(jī)模型

        玻爾茲曼機(jī) (Boltzmann Machine,BM)是 1985年由Hinton等提出的一種起源于統(tǒng)計(jì)物理學(xué)的隨機(jī)遞歸神經(jīng)網(wǎng)絡(luò)。作為一種隨機(jī)生成的Hopfield網(wǎng)絡(luò),BM由可見層和隱單元層組成,用可見單元和隱單元來表示隨機(jī)網(wǎng)絡(luò)和隨機(jī)環(huán)境的學(xué)習(xí)模型,通過權(quán)值表示單元之間的相關(guān)性。

        為克服BM訓(xùn)練時間過長、難以確切計(jì)算模型分布等問題,研究人員設(shè)計(jì)了受限玻爾茲曼機(jī)。RBM通過限制BM中的層內(nèi)單元連接,使得在給定同層單元狀態(tài)下臨近層單元激活概率條件獨(dú)立,其結(jié)構(gòu)如圖1所示。作為無向圖模型,RBM中的V為可見單元層,表示觀測數(shù)據(jù);H表示隱單元層,表現(xiàn)為特征檢測器,層內(nèi)隱單元狀態(tài)組合等價(jià)于觀測數(shù)據(jù)的特征組合。文獻(xiàn)[7]指出RBM的隱單元和可見單元可為任意指數(shù)族,如泊松單元、高斯單元和Sofmax單元等。文獻(xiàn)[8]從理論方面對RBM的數(shù)據(jù)泛化能力進(jìn)行了證明,得出了在足夠多數(shù)目隱單元的情況下RBM可以表征任意離散分布的結(jié)論。

        圖1 RBM單元連接圖Fig.1 RBM unit connection diagram

        以包含N個二值可見單元和M個二值隱單元的RBM為例,設(shè)定vi表示第i個可見單元的狀態(tài),hj表示第j個隱單元狀態(tài)。給定狀態(tài)(v,h),模型的能量定義如式(1)所示。

        其中,Wij表示可視單元i與隱單元j之間的連接權(quán)值,bi表示可視單元i的偏置,cj表示隱單元j的偏置,則RBM處于狀態(tài)(v,h)的概率如式(2)所示。

        由于層間單元是無連接的,所有單元的激活狀態(tài)都是條件獨(dú)立的,則隱單元和可見單元的后驗(yàn)激活概率,如式(3)、(4)所示。

        其中σ(x)=1/(1+exp(-x)),在本文中也寫為Sigmoid函數(shù)。

        1.2CD算法

        RBM通過極大似然法則對數(shù)據(jù)進(jìn)行無監(jiān)督學(xué)習(xí),即最大化訓(xùn)練數(shù)據(jù)出現(xiàn)概率。假設(shè)訓(xùn)練樣本個數(shù)為K,則RBM對數(shù)據(jù)的學(xué)習(xí)即為最大化以下目標(biāo)函數(shù),如式(5)所示。

        利用隨機(jī)梯度下降法,目標(biāo)函數(shù)式(5)對于各個參數(shù)的偏導(dǎo)如式(6)所示(以W為例)。

        對于式(6)等號右側(cè)多項(xiàng)式中的第一個項(xiàng)(Positive Phase),可以通過訓(xùn)練數(shù)據(jù)與式(4)直接計(jì)算出來;但是對于第二個項(xiàng)(Negative Phase),由于p(v.h)是無法直接通過數(shù)學(xué)推導(dǎo)出來的,故難以直接計(jì)算。

        Hinton于2002年提出了對比散度(Contrastive Divergence,CD)算法[9],通過執(zhí)行吉布斯塊采樣(block Gibbs Sampling)—分別以各個訓(xùn)練數(shù)據(jù)作為初始狀態(tài),進(jìn)行k次狀態(tài)轉(zhuǎn)移(k次Step),然后將轉(zhuǎn)移后的數(shù)據(jù)作為樣本對Negative Phase估算均值實(shí)現(xiàn)RBM的訓(xùn)練。Hinton通過實(shí)驗(yàn)證明,在實(shí)際應(yīng)用中,甚至只需要一次狀態(tài)轉(zhuǎn)移就能保證良好的學(xué)習(xí)效果。由此,在給定訓(xùn)練數(shù)據(jù) v(n)下,隱單元 j的特征 wj的迭代如式(7)所示。

        2 CUDA實(shí)現(xiàn)

        隨著顯卡技術(shù)的發(fā)展,最初用于計(jì)算機(jī)圖形渲染的圖像處理器 (Graphics Processing Unit,GPU)計(jì)算功能越來越強(qiáng)大,NVIDIA公司推出了面向通用計(jì)算的 CUDA(Compute Unified Device Architecture)運(yùn)算架構(gòu)[10-14],已經(jīng)成功地應(yīng)用到了多種計(jì)算領(lǐng)域。

        2.1算法分析

        以Step=1的CD算法為例,對RBM在訓(xùn)練過程中的CUDA優(yōu)化階段進(jìn)行分析,設(shè)定可見單元和隱單元維數(shù)分別為I和J,N為訓(xùn)練樣本數(shù)目,以箭頭表示RBM中參數(shù)參與計(jì)算的方向,則RBM的訓(xùn)練步驟如圖2所示。

        圖2 RBM訓(xùn)練算法的步驟表示Fig.2 RBM algorithm training Steps

        從圖中可以看出,RBM的訓(xùn)練主要體現(xiàn)在步驟1~6中。這6個步驟分為矩陣乘加(步驟1、3、5、6)、0~1均勻分布隨機(jī)數(shù)產(chǎn)生(步驟2、4)和Sigmoid函數(shù)計(jì)算(步驟1、3、5)3個方面。這3個方面計(jì)算要求不同,因此在設(shè)計(jì)CUDA并行化時需要單獨(dú)對待。

        2.2并行設(shè)計(jì)

        1)矩陣乘加

        鑒于矩陣運(yùn)行庫CUBLAS在處理矩陣運(yùn)算的高效性和通用性,因此這一部分計(jì)算可以直接通過該庫實(shí)現(xiàn),主要使用庫中處理雙精度矩陣乘法運(yùn)算的cublasDgemm函數(shù)。函數(shù)說明如下所示,并且所代表的運(yùn)算如式(8)所示。

        cublasDgemm(handle,CUBLAS_OP_N,CUBLAS_OP_N,mat1->_rows,mat2->_cols,mat2->_rows,&alpha,d_a,mat1->_rows,d_b,mat2->_rows,&beta,d_c,mat1->_rows)

        其中 d_a為矩陣 a(mat1)在顯存的地址,d_b為矩陣b(mat2)在顯存的地址,d_c為計(jì)算結(jié)果c的地址,需要在計(jì)算完成以后復(fù)制到計(jì)算機(jī)內(nèi)存中。

        2)0~1均勻分布隨機(jī)數(shù)的產(chǎn)生

        0~1均勻分布隨機(jī)數(shù)的產(chǎn)生需要CUDA隨機(jī)數(shù)生成器參與。所謂隨機(jī)數(shù)生成器(Random Number Generator,RNG),是指能夠產(chǎn)生隨機(jī)數(shù)序列的軟件、硬件或者二者的結(jié)合。CUDA 的RNG實(shí)質(zhì)上是偽隨機(jī)數(shù)發(fā)生器 (Pseudo Random Number Generator,PRNG),下面對常用的PRNG進(jìn)行簡單介紹。

        1)線性同余生成器(Linear Congruence Generator,LCG)

        線性同余生成器是應(yīng)用較為廣泛的一種偽隨機(jī)數(shù)生成器,優(yōu)點(diǎn)在于計(jì)算速度快、容易實(shí)現(xiàn)、便于移植等,它是利用數(shù)論中的同余運(yùn)算來產(chǎn)生偽隨機(jī)數(shù)的,其隨機(jī)數(shù)遞推公式如式(9)所示。

        其中a為乘子,c為增量,M為模數(shù),Xi為序列中的第i個數(shù),Xo稱為種子(Seed)。

        2)Tausworthe序列發(fā)生器

        Tausworthe序列發(fā)生器又稱為移位寄存器序列發(fā)生器,其工作原理是通過寄存器進(jìn)行位移操作,直接在存儲單元中生成隨機(jī)數(shù),遞推運(yùn)算公式如式(10)所示。

        其中(x0,X1,…,Xp-1)為初始值,p>q>0為整數(shù),Xi是二進(jìn)制行向量,⊕是位運(yùn)算符。但是由于Tausworthe序列存在較為嚴(yán)重的相關(guān)性,很少被單獨(dú)使用,通常用于構(gòu)建組合Tausworthe生成器。

        3)組合發(fā)生器

        組合發(fā)生器指將多個獨(dú)立的發(fā)生器按照某種方式組合在一起以產(chǎn)生偽隨機(jī)數(shù),從而得到統(tǒng)計(jì)性能更好和周期更長的PRNG。目前CUDA中的隨機(jī)數(shù)生成器通常是幾種PRNA的組合,從而達(dá)到擴(kuò)大生成器周期和摒棄單獨(dú)PRNA中存在的統(tǒng)計(jì)學(xué)缺陷的目的。文獻(xiàn) [11]提出一種混合LCG和組合Tausworthe的隨機(jī)數(shù)產(chǎn)生器,周期為 2121,文獻(xiàn)[6]進(jìn)一步對RNG進(jìn)行了改良,實(shí)現(xiàn)Tausworthe113和CLCG4的組合發(fā)生器,其周期達(dá)到2234。

        鑒于 Tausworthe113較為復(fù)雜,采用 CLCG4和組合Tausworthe共同作為隨機(jī)數(shù)生成器,周期約為 2210=121+89。其Kernel端核心代碼如下。

        3)Sigmoid函數(shù)計(jì)算

        文獻(xiàn)[5]指出,CUDA能夠?qū)崿F(xiàn)Sigmoid函數(shù)的方式有兩種,第一種是通過CUDA提供的-expf(x)函數(shù),定義Sigmoid函數(shù);第二種是利用CUDA拾取紋理內(nèi)存的讀取模式為cudaReadModeNormalizedFloat時數(shù)值被映射到[0,1]區(qū)間的能力,選用第二種方式。

        綜上,利用CUDA實(shí)現(xiàn)RBM訓(xùn)練的過程如圖3所示。

        圖3 基于CUDA的RBM訓(xùn)練并行設(shè)計(jì)Fig.3 RBM training CUDA-based parallel design

        3 實(shí)驗(yàn)分析

        為驗(yàn)證CUDA對RBM訓(xùn)練加速的效果,選取MNIST手寫數(shù)字識別集為測試對象。通過對MNIST數(shù)據(jù)集進(jìn)行訓(xùn)練,對比RBM的訓(xùn)練時間和對數(shù)據(jù)的擬合程度。實(shí)驗(yàn)平臺的CPU為Intel Core i5 760@2.8GHz,主機(jī)內(nèi)存4 GB,GPU采用GTX460(這與文獻(xiàn)[5]的GPU平臺一致,便于進(jìn)行數(shù)據(jù)對比)。

        CUDA優(yōu)化性能主要參考的指標(biāo)是加速比 (Speed-up)。加速比為程序串行執(zhí)行時間與并行執(zhí)行時間比值,數(shù)值越大,說明程序的并行化程度越高,CUDA優(yōu)化效果越好,定義如式(1)所示。

        實(shí)驗(yàn)設(shè)置RBM的結(jié)構(gòu)為784-Nhid,其中Nhid代表隱單元個數(shù)分別取100至800的8個數(shù)值,RBM參數(shù)的學(xué)習(xí)速率統(tǒng)一設(shè)置為0.01,CD訓(xùn)練循環(huán)次數(shù)不超過1000。對比方法選用文獻(xiàn)[5]提出的CUDA設(shè)計(jì)。圖4為在不同隱單元個數(shù)下的兩種CUDA設(shè)計(jì)的加速比值。

        圖4 不同CUDA方法加速比Fig.4 Speedup by different CUDA methods

        從圖中可以看出,經(jīng)過CUDA加速,每個RBM的訓(xùn)練加速比都在25以上,表現(xiàn)出了較好的并行優(yōu)化效率。當(dāng)隱單元較少,即RBM結(jié)構(gòu)小的情況下,所設(shè)計(jì)的CUDA加速與文獻(xiàn)[118]相比存在差距,原因是Sigmoid函數(shù)計(jì)算中發(fā)生的對紋理內(nèi)存的讀取增加了顯存讀取,影響了計(jì)算效率。但隨著隱單元個數(shù)的增大,所設(shè)計(jì)的CUDA程序能夠得到更大的加速比,原因在于在CUDA框架下,所設(shè)計(jì)的CUDA程序?qū)﹄[單元和可見單元隨機(jī)狀態(tài)的選取進(jìn)行了硬件加速,相當(dāng)于對RBM結(jié)構(gòu)的參數(shù)更新進(jìn)行了進(jìn)一步的并行化設(shè)計(jì)。因此,隨著隱單元個數(shù)的增大,其優(yōu)越性將得到體現(xiàn)。

        表1為當(dāng)隱單元個數(shù)設(shè)置為800時,2種不同運(yùn)算方式下的RBM訓(xùn)練時間。從表中可以看出,利用CUDA對RBM訓(xùn)練進(jìn)行并行化計(jì)算,加速比達(dá)到41.32,說明所設(shè)計(jì)實(shí)現(xiàn)的CUDA并行化對RBM的加速較為有效。

        選取不同訓(xùn)練樣本個數(shù)下的CUDA對RBM的加速比如圖5所示。

        表1 RBM訓(xùn)練時間Tab.1 RBM training time

        圖5 不同樣本數(shù)目下的CDBM加速比Fig.5 Speedup under different samples number

        從圖中可以看出,加速比基本與訓(xùn)練樣本數(shù)目成正比,同時隨著訓(xùn)練樣本數(shù)目的增多呈現(xiàn)出更為顯著的增長趨勢。

        4 結(jié) 論

        本文設(shè)計(jì)了基于CUDA的RBM并行訓(xùn)練方法,實(shí)現(xiàn)了在GPU技術(shù)下RBM的加速訓(xùn)練。相對于以往RBM的并行加速訓(xùn)練,該方法能夠更好地處理大規(guī)模數(shù)據(jù),促進(jìn)了RBM模型在未來大數(shù)據(jù)環(huán)境下的實(shí)用。從實(shí)驗(yàn)結(jié)果來看,本文設(shè)計(jì)實(shí)現(xiàn)的基于GPU的RBM并行加速實(shí)現(xiàn),具有較廣的應(yīng)用范圍和較高的工程價(jià)值。

        [1]Becker S.An information-theoretic unsupervised learning algorithm for neural networks[D].University of Toronto,1992.

        [2]Srivastava N,Salakhutdinov R R,Hinton G E.Modeling documents with deep boltzmann machines[C]//29th Conference on Uncertainty in Artificial Intelligence.Bellevue:IEEE,2013:222-227.

        [3]Daniel L.Ly,Paprotski V,Danny Y.Neural Networks on GPUs:Restricted Boltzmann Machines[C]//IEEE Conference on Machine Learning and Applications.Bellevue:IEEE,2010:307-312.

        [4]CaiX,XuZ,LaiG,etal.GPU-acceleratedrestricted boltzmann machine for collaborative filtering[M].Algorithms and Architectures for Parallel Processing.Springer Berlin Heidelberg,2012:303-316.

        [5]Lopes N,Ribeiro B,Goncalves J.Restricted Boltzmann machines and deep belief networks on multi-core processors [C]//Neural Networks(IJCNN),NewYork:Spring,2012:1-7.

        [6]薛少飛,宋彥,戴禮榮.基于多GPU的深層神經(jīng)網(wǎng)絡(luò)快速訓(xùn)練方法[J].清華大學(xué)學(xué)報(bào):自然科學(xué)版,2013,53(6):745-748.

        [7]Welling M,Rosen-Zvi M,Hinton G E.Exponential family harmoniums with an application to information retrieval[C].Advances in neural information processing systems,2004: 1481-1488.

        [8]Le Roux N,Bengio Y.Representational power of restrictedBoltzmann machines and deep belief networks[J].Neural Computation,2008,20(6):1631-1649.

        [9]G.E.Hinton.Training products of experts by minimizing contrastive divergence[J].Neural Computation,2002,14 (8):1711-1800.

        [10]王坤.基于GPU的分類并行算法的研究與實(shí)現(xiàn)[J].電子設(shè)計(jì)工程,2014,(18):39-41.

        [11]張聰,邢同舉,羅穎,等.基于GPU的數(shù)學(xué)形態(tài)學(xué)運(yùn)算并行加速研究[J].電子設(shè)計(jì)工程,2011(19):141-143.

        [12]王應(yīng)戰(zhàn),魏衍君.基于并行計(jì)算的木馬免疫算法研究[J].電子設(shè)計(jì)工程,2012(16):45-47.

        [13]榮少?。诟倪M(jìn)A*算法的水下航行器自主搜索航跡規(guī)劃[J].電子科技,2015(4):17-19,22.

        [14]蔣磊,鄒兵,吳明.基于改進(jìn)免疫遺傳算法的含分布式電源配電網(wǎng)規(guī)劃[J].陜西電力,2012(10):26-30.

        Realization of RBM parallel training based on GPU

        ZHANG Li-min1,LIU Kai1,F(xiàn)AN Xiao-lei2
        (1.Naval Aeronautical and Astronautical University,Yantai 264001,China;2.Noncommissioned Officers Vocational and Technical Education College,The Second Artillery Engineering University,Qingzhou 262500,China)

        In order to overcome the low efficiency of Restricted Boltzmann Machine handle large data,on the basic of parallel computing platform GPU,the realization of RBM training based on GPU is designed.By researching the training steps of RBM,the contrast divergence algorithm is redesigned to implement on GPU.Combined with previous GPU parallel solutions,matrix multiply-add operations are implemented by CUBLAS libraries.The combination of Tausworthe113 and CLCG4 is used as random number generation to get longer cycle and more concise random number.The CUDA pickup texture memory read mode is used to achieve sigmoid function value,and finally The MNIST handwriting digit database is conduct on the test of this new realization.The MNIST experiment results illustrated that the novel algorithm has good feasibility and is advantageous for hug amount of data.Compared to the previous RBM parallel code,this new GPU parallel processing have more obvious advantages on large data sets and the speedup rate reach at least 25.

        RBM;GPU;CUDA;speed-up;parallel programming

        TN99

        A

        1674-6236(2016)02-0028-04

        2015-03-21稿件編號:201503294

        國家自然科學(xué)基金(61032001)

        張立民(1976—),男,吉林開源人,博士,教授。研究方向:信號與信息處理、軍用仿真技術(shù)。

        猜你喜歡
        模型設(shè)計(jì)
        一半模型
        重要模型『一線三等角』
        何為設(shè)計(jì)的守護(hù)之道?
        重尾非線性自回歸模型自加權(quán)M-估計(jì)的漸近分布
        《豐收的喜悅展示設(shè)計(jì)》
        流行色(2020年1期)2020-04-28 11:16:38
        瞞天過?!律O(shè)計(jì)萌到家
        設(shè)計(jì)秀
        海峽姐妹(2017年7期)2017-07-31 19:08:17
        有種設(shè)計(jì)叫而專
        Coco薇(2017年5期)2017-06-05 08:53:16
        3D打印中的模型分割與打包
        FLUKA幾何模型到CAD幾何模型轉(zhuǎn)換方法初步研究
        热门精品一区二区三区| 少妇被躁爽到高潮无码文| 欧美中文字幕在线| 日本一区二区国产高清在线播放 | 亚洲av无码国产精品永久一区| 人妻丰满熟妇av无码处处不卡| 久久男人av资源网站无码| 久久久精品人妻一区二区三区免费 | 熟妇人妻无乱码中文字幕| 国产精品18久久久久久不卡中国| 日韩熟女精品一区二区三区视频| 99噜噜噜在线播放| 中国国语毛片免费观看视频| 久久国产亚洲精品超碰热| 狼人狠狠干首页综合网| 亚洲最大成人综合网720p| 黑人巨茎大战欧美白妇| 调教在线播放黄| av一区二区三区有码| 亚洲av无码专区在线| 国产无套护士在线观看| 亚洲无码vr| 一级黄色一区二区三区| 国产成人喷潮在线观看| 丁香综合网| 久久久亚洲精品蜜臀av| 人妻久久久一区二区三区蜜臀| 东京无码熟妇人妻av在线网址| 国产亚洲精品国产福利在线观看| 亚洲大胆美女人体一二三区| 日韩久久无码免费毛片软件| 国产香蕉97碰碰视频va碰碰看| 欧美亚洲国产精品久久久久| 精品久久综合日本久久综合网| 99精品国产在热久久无码| 久久精品片| 区二区三区亚洲精品无| 丰满熟妇乱又伦精品| 亚洲综合无码一区二区三区| 久久精品中文字幕亚洲| 色综合久久中文综合网亚洲|