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

    CUDA加速的DNA-蛋白質(zhì)匹配及其優(yōu)化

    2013-09-07 02:52:46陳春雷慕德俊張慧翔
    電子技術(shù)應(yīng)用 2013年9期
    關(guān)鍵詞:方法

    陳春雷,慕德俊,張慧翔,胡 偉

    (西北工業(yè)大學(xué) 自動(dòng)化學(xué)院,陜西 西安 710072)

    結(jié)構(gòu)生物信息學(xué)已成為當(dāng)前計(jì)算機(jī)科學(xué)領(lǐng)域的一個(gè)研究熱點(diǎn)。其主要研究目標(biāo)是獲取生物系統(tǒng)的高分辨率結(jié)構(gòu)信息,從而精確地推理系統(tǒng)的功能、預(yù)測(cè)某些修改或擾動(dòng)對(duì)系統(tǒng)造成的影響。與生物信息學(xué)的其他領(lǐng)域相比,結(jié)構(gòu)生物信息學(xué)面臨眾多新的挑戰(zhàn)。其中之一是多數(shù)結(jié)構(gòu)生物信息學(xué)問(wèn)題的搜索空間是連續(xù)的;換言之,生物系統(tǒng)的微觀結(jié)構(gòu)一般通過(guò)原子在空間坐標(biāo)系中的位置來(lái)表示,而坐標(biāo)通常是連續(xù)變量,因此搜索空間是無(wú)窮的。雖然某些近似方法可以在一定程度上應(yīng)對(duì)結(jié)構(gòu)生物信息學(xué)問(wèn)題的這種連續(xù)特性,但是從搜索空間中求得最終解仍然需要進(jìn)行大量的運(yùn)算[1]。

    DNA-蛋白質(zhì)匹配是生物信息學(xué)的一個(gè)重要問(wèn)題。傳統(tǒng)的序列分析方法往往會(huì)導(dǎo)致較高的多重檢驗(yàn)錯(cuò)誤率 (false-positive rate)[2]。多數(shù)基于結(jié)構(gòu)生物信息學(xué)的DNA-蛋白質(zhì)匹配方法把獲得轉(zhuǎn)錄因子-DNA組合體的信息作為一個(gè)必備條件,而轉(zhuǎn)錄因子-DNA組合體的信息通常需要通過(guò)實(shí)驗(yàn)獲得,這種實(shí)驗(yàn)通常需要耗費(fèi)大量的時(shí)間。參考文獻(xiàn)[2]提出的基于退火算法的匹配算法避開(kāi)了這個(gè)問(wèn)題,但是使用該方法進(jìn)行DNA-蛋白質(zhì)匹配仍需要較大的運(yùn)算量。參考文獻(xiàn)[3]基于參考集索引技術(shù)提出了一種快速的序列分析算法,并將其應(yīng)用于DNA-蛋白質(zhì)匹配。參考文獻(xiàn)[4]采用多CPU的方式設(shè)計(jì)實(shí)現(xiàn)了并行化的DNA-蛋白質(zhì)序列分析方法,并取得了較高的加速比。但上述工作均未涉及基于結(jié)構(gòu)生物信息學(xué)的匹配方法的加速問(wèn)題。

    應(yīng)對(duì)大運(yùn)算量問(wèn)題的一個(gè)常用方法是并行計(jì)算。NVIDIA公司統(tǒng)一計(jì)算設(shè)備架構(gòu)CUDA(Compute Unified Device Architecture)是一種全新的并行計(jì)算架構(gòu)。在CUDA的支持下,通用計(jì)算圖形處理單元GPGPU(General Purpose Graphic Processing Unit)可以作為一種通用的效率、硬件加速器,發(fā)揮出強(qiáng)大的運(yùn)算能力[5]。

    本文針對(duì)參考文獻(xiàn)[2]的DNA-蛋白質(zhì)匹配方法,從計(jì)算的角度對(duì)其進(jìn)行分析,設(shè)計(jì)實(shí)現(xiàn)并優(yōu)化了基于CUDA的加速方法,通過(guò)實(shí)驗(yàn)驗(yàn)證了其加速性能。

    1 相關(guān)知識(shí)介紹

    1.1 基于退火算法的DNA-蛋白質(zhì)匹配算法

    給定一條DNA鏈和一條蛋白質(zhì)鏈,兩者之間的相對(duì)位置存在很多種可能。假定兩者均為剛體,DNA-蛋白質(zhì)組合體的物理能量主要受兩者相對(duì)位置的影響,組合體的物理能量越低,其穩(wěn)定性越強(qiáng)。該方法試圖尋找使得組合體最穩(wěn)定的相對(duì)位置,這一結(jié)果對(duì)生物信息學(xué)的研究具有重要意義。

    從計(jì)算的角度看,該方法的基礎(chǔ)是退火算法。為了提高搜索精度,該匹配方法通常需要隨機(jī)產(chǎn)生大量初始“種子”,為每個(gè)“種子”啟動(dòng)一個(gè)基于退火算法的搜索過(guò)程,在整個(gè)解空間中尋找最穩(wěn)定的相對(duì)位置。這種匹配方法的流程如圖1所示。

    圖1 基于退火算法的DNA-蛋白質(zhì)匹配方法流程圖

    圖1中的初始值為T0,溫度T及溫度衰減周期interval為正整數(shù),溫度衰減系數(shù)為a。每當(dāng)算法所經(jīng)歷的平移和旋轉(zhuǎn)次數(shù)達(dá)到interval的非零整數(shù)倍時(shí),T衰減為原來(lái)的a倍(0<a<1),Tmin為溫度的最小值。

    DNA-蛋白質(zhì)組合體的物理能量E1、E2分別代表第n次(1≤n≤Smax,Smax為平移和旋轉(zhuǎn)次數(shù)的最大值)平移和旋轉(zhuǎn)前后組合體的物理能量。Emin代表通過(guò)算法得到的組合體物理能量的最小值。如果E2<E1,則接受第n次平移和旋轉(zhuǎn)的結(jié)果;否則,計(jì)算指數(shù)式exp(-(E2-E1)/T)的值,并使用C++庫(kù)函數(shù)drand48()產(chǎn)生一個(gè)在區(qū)間[0,1]上均勻分布的隨機(jī)數(shù)。如果指數(shù)式的值小于隨機(jī)數(shù)的值,則接受第n次平移旋轉(zhuǎn)的結(jié)果,反之不接受。

    step為當(dāng)前進(jìn)行到了第幾次平移和旋轉(zhuǎn)。

    1.2 CUDA編程模型及線程調(diào)度方式的特點(diǎn)

    CUDA程序通常包含CPU串行代碼和GPU并行代碼,后者被稱作CUDA程序的內(nèi)核。在執(zhí)行內(nèi)核時(shí),CUDA會(huì)產(chǎn)生大量并行工作的線程,以單指令多數(shù)據(jù)SIMD(Single Instruction Multiple Data)方式完成整個(gè)內(nèi)核的計(jì)算任務(wù)。CUDA采用網(wǎng)格(grid)和線程塊(block)來(lái)組織線程[6]。一個(gè)block的線程又被劃分為一個(gè)或多個(gè)線程組(warp)。warp是CUDA程序最基本的調(diào)度單位,屬于同一個(gè)warp的各個(gè)線程會(huì)從CUDA程序代碼段中相同的程序地址同時(shí)開(kāi)始執(zhí)行,但是各線程具有獨(dú)立的當(dāng)前指令指針和寄存器狀態(tài)。因此,各線程可能會(huì)有不同的執(zhí)行路徑,例如執(zhí)行不同的分支選擇結(jié)構(gòu)代碼或循環(huán)結(jié)構(gòu)代碼[7]。warp執(zhí)行特點(diǎn)如圖2所示。

    圖2 warp的執(zhí)行特點(diǎn)示意圖

    假設(shè)一個(gè)warp中共有4個(gè)線程:線程1~線程4,它們的執(zhí)行時(shí)間分別是t1~t4,并且t3<t1<t2<t4。因?yàn)镃UDA的基本調(diào)度單位是一個(gè)完整的warp而非單個(gè)線程,所以整個(gè)warp的執(zhí)行時(shí)間取決于執(zhí)行時(shí)間最長(zhǎng)的線程t4。其他3個(gè)線程在執(zhí)行完畢后必須等待還沒(méi)有執(zhí)行完畢的線程,因此有一段時(shí)間處于空閑狀態(tài),相應(yīng)的計(jì)算資源也就被閑置了。例如,線程3閑置的計(jì)算資源最多,其空閑時(shí)長(zhǎng)為(t4-t3)。造成計(jì)算資源閑置的原因是同一warp中各個(gè)線程的執(zhí)行路徑不同;而CUDA采用的是SIMD并行方式,執(zhí)行路徑的不同是由每個(gè)線程所處理的數(shù)據(jù)差異造成的。因此,依照一定規(guī)則對(duì)數(shù)據(jù)進(jìn)行重新排列是減少這種計(jì)算資源閑置狀況的常用方法[8]。重排規(guī)則通常視具體應(yīng)用而定。

    2 使用CUDA加速DNA-蛋白質(zhì)匹配

    2.1 從計(jì)算角度分析匹配算法

    匹配方法的流程已由圖1給出,除參數(shù)初始化外,該方法可分為三部分:(1)平移、旋轉(zhuǎn)DNA和蛋白質(zhì);(2)計(jì)算DNA-蛋白質(zhì)組合體的能量;(3)后續(xù)處理。這3部分在普通CPU平臺(tái)上所耗時(shí)間依次為∶73.624 s、1138.945 s、110.809 s(僅使用一個(gè)隨機(jī)初始種子,退火算法參數(shù)的預(yù)設(shè)值及軟硬件平臺(tái)參數(shù)指標(biāo)見(jiàn)本文第3部分),實(shí)驗(yàn)使用的DNA、蛋白質(zhì)結(jié)構(gòu)數(shù)據(jù)來(lái)自參考文獻(xiàn)[9]編號(hào)為2GLI的文件。其他的DNA、蛋白質(zhì)結(jié)構(gòu)數(shù)據(jù)的實(shí)驗(yàn)結(jié)果也顯示了計(jì)算能量所耗費(fèi)的時(shí)間遠(yuǎn)超過(guò)其他兩個(gè)部分。計(jì)算能量時(shí),DNA-蛋白質(zhì)組合體所處的三維空間被劃分成若干個(gè)棱長(zhǎng)為埃米(10-10m)級(jí)的晶格。一條DNA鏈由若干個(gè)DNA原子構(gòu)成,一條蛋白質(zhì)鏈由若干個(gè)蛋白質(zhì)原子構(gòu)成;以2GLI組合體為例,共有855個(gè)DNA原子和1 270個(gè)蛋白質(zhì)原子。每次平移和旋轉(zhuǎn)前后,每個(gè)原子所屬的晶格都可能發(fā)生改變,每個(gè)晶格所包含的原子個(gè)數(shù)也可能發(fā)生改變。Di(i=1,2,…,ND),Pj(j=1,2,…,NP)分別代表組合體中的DNA原子數(shù)量和蛋白質(zhì)原子數(shù)量。DNA原子i(i=1,2,…,ND)在三維空間中所處的晶格和相鄰的晶格共有27個(gè),統(tǒng)稱為原子i的臨近晶格(Neighboring Lattice),以pi,js(x,1,2,…,ni)代表這27個(gè)晶格中的所有蛋白質(zhì)原子,這些原子即原子i的相鄰蛋白質(zhì)原子。以di,jx代表DNA原子i與蛋白質(zhì)原子jx之間的歐式距離,兩個(gè)原子之間的能量是di,jx的函數(shù):E(di,jx)。則組合體的總能量為:

    根據(jù)以上分析可知:(1)計(jì)算能量是匹配方法中最耗時(shí)部分;(2)E(di,jx)(i=1,2,…,ND;x=1,2,…,ni)的計(jì)算是相互獨(dú)立的,無(wú)依賴關(guān)系,適于并行。根據(jù)阿姆達(dá)爾定律[10],選擇計(jì)算能量部分進(jìn)行并行化。

    2.2 基于CUDA的并行匹配算法

    以下偽代碼給出了串行方式計(jì)算能量的步驟:

    算法1:串行方式計(jì)算能量

    (1)E←0

    (2)for i=1 to ND

    (3) for all neighboring lattice of atomi∶la=1 to 27

    (4) for all protein atoms in lattice la

    (5) compute E(di,jx)and E←E+E(di,jx)

    (6) end for

    (7) end for

    (8)end for

    令CUDA block為一維結(jié)構(gòu),一個(gè)block中包含B個(gè)線程,一個(gè)初始種子的搜索過(guò)程僅由一個(gè)block完成?;谝陨洗写a,給出CUDA內(nèi)核的偽代碼:

    算法2:并行方式計(jì)算能量

    (1)tid←threadIdx.x

    (2)Ethis_thread←0,stride←「ND/B?

    (3)for s=0 to stride

    (4)i←s.B+tid

    (5)if i<ND

    (6) for all neighboring lattice of atom i∶la=1 to 27

    (7) for all protein atoms in lattice la

    (8) compute and E(di,jx)and Ethis_thread←Ethis_thread+E(di,jx)

    (9) end for

    (10) end for

    (11)end for

    算法2實(shí)現(xiàn)了算法1第2行循環(huán)操作的并行化。每個(gè)CUDA線程對(duì)應(yīng)一個(gè)DNA原子,依次進(jìn)入該原子的27個(gè)相鄰晶格,找到各個(gè)相鄰蛋白質(zhì)原子,計(jì)算E(di,jx),并累加到能量的部分和:Ethis_thread。在CUDA內(nèi)核返回后,通過(guò)累加各線程的能量部分和可以得到組合體的總能量。另外,如果B<ND,則需要把ND個(gè)DNA原子平均分成「ND/B?塊(最后一塊可能不足B個(gè)原子),然后,由整個(gè)block的線程依次處理各塊,算法2第3行的循環(huán)結(jié)構(gòu)即為了達(dá)到這個(gè)目的。

    2.3 優(yōu)化并行算法

    考慮到本文1.2節(jié)所述warp的執(zhí)行特點(diǎn),上述并行方式存在一定的不足。算法2中第7行的循環(huán)次數(shù)取決于當(dāng)前晶格中蛋白質(zhì)原子的個(gè)數(shù),同一warp中各線程的循環(huán)次數(shù)可能會(huì)不同,如果差異過(guò)大,會(huì)造成嚴(yán)重的計(jì)算資源閑置;對(duì)所有線程而言,算法2第3行、第6行循環(huán)結(jié)構(gòu)的循環(huán)次數(shù)是確定的。

    假設(shè)DNA存儲(chǔ)在一個(gè)數(shù)組中,且該數(shù)組位于GPU主存(global memory)中。為了盡量減少計(jì)算資源閑置,重排DNA原子的原有次序(DNA鏈中的次序),處于同一個(gè)晶格中的DNA原子被存儲(chǔ)在數(shù)組的某一段連續(xù)區(qū)域內(nèi)。采用這種優(yōu)化措施是基于以下原因:(1)由式(1)可知,總能量由E(di,jx)做算術(shù)加法得到,與E(di,jx)的計(jì)算次序無(wú)關(guān);(2)為了提高GPU主存的帶寬利用率,同一warp中的線程通常從主存中的臨近區(qū)域讀取數(shù)據(jù),即內(nèi)存訪問(wèn)對(duì)齊 (coalescing)[6];(3)DNA的雙螺旋結(jié)構(gòu)是非線性的,DNA在鏈中相鄰的原子未必處于同一晶格中;(4)同一晶格中的DNA原子的臨近蛋白質(zhì)原子數(shù)量相同,重排可以減少因線程執(zhí)行路徑差異造成的計(jì)算資源閑置。重排的效果如圖3所示。

    圖3 DNA原子存儲(chǔ)次序重排示意圖

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

    3.1 實(shí)驗(yàn)平臺(tái)

    硬 件 平 臺(tái):Intel?CoreTM2 E7500 CPU,2.93 GHz,NVIDIA GTX 660圖形處理器(單個(gè)warp包含32個(gè)線程),CPU主存4 GB。

    軟件平臺(tái):Linux操作系統(tǒng)內(nèi)核版本2.6.18-194.el5,gcc版本4.1.2,CUDA版本5.0。

    3.2 實(shí)驗(yàn)參數(shù)設(shè)置與結(jié)果

    DNA-蛋白質(zhì)結(jié)構(gòu)數(shù)據(jù)編號(hào)為2GLI。CPU版程序以單線程串行方式執(zhí)行;GPU版程序block size固定為512;各block從不同的初始種子出發(fā),分別基于退火算法進(jìn)行搜索,每個(gè)block對(duì)應(yīng)一個(gè)初始種子。退火算法參數(shù):初始溫度T0=120℃,最低溫度Tmin=0.001℃,溫度衰減周期interval=800,溫度衰減系數(shù)a=0.99,最大平移旋轉(zhuǎn)次數(shù)Smax=106。

    實(shí)驗(yàn)結(jié)果如表1所示。與單線程CPU程序相比,未經(jīng)優(yōu)化的GPU程序?qū)@得最高可達(dá)8倍左右的加速比;而經(jīng)過(guò)重排優(yōu)化后,加速比在此基礎(chǔ)上進(jìn)一步顯著提升,最高可達(dá)15倍左右。

    表1 實(shí)驗(yàn)結(jié)果

    本文針對(duì)一種典型的DNA-蛋白質(zhì)匹配算法,設(shè)計(jì)實(shí)現(xiàn)了基于CUDA的并行化方法,從線程調(diào)度的角度對(duì)該方法進(jìn)行優(yōu)化,并通過(guò)實(shí)驗(yàn)驗(yàn)證了加速性能。實(shí)際應(yīng)用往往需要為一個(gè)組合體產(chǎn)生大量的初始種子(數(shù)百個(gè)),并為每個(gè)種子開(kāi)啟一個(gè)基于退火算法的搜索過(guò)程;其目的是達(dá)到較高的匹配精度。實(shí)驗(yàn)顯示,使用單個(gè)GPGPU加速,在單個(gè)block包含的線程數(shù)不變的前提下,隨著初始種子數(shù)量的增加,加速比逐漸趨于穩(wěn)定。例如,當(dāng)初始種子個(gè)數(shù)超過(guò)40后,加速比基本穩(wěn)定在15倍左右。其原因在于單個(gè)GPGPU的計(jì)算能力存在上限,當(dāng)種子足夠多時(shí),其計(jì)算能力已得到較充分利用,無(wú)法繼續(xù)提高加速比。為了滿足實(shí)際應(yīng)用的需求,下一步的工作將考慮使用基于GPGPU的集群來(lái)加速匹配算法,以進(jìn)一步提高加速比。

    [1]BOURNE P E,WEISSIG H.Structural bioinformatics[M].Hoboken∶Wiley-Liss Inc,2003.

    [2]Liu Zhijie,Guo Juntao,Li Ting,et al.Structure-based prediction of transcription factor binding sites using a protein-DNA docking approach[J].Proteins,2008,72(4)∶1114-1124.

    [3]戴東波,熊赟,朱揚(yáng)勇.基于參考集索引的高效序列相似性查找算法[J].軟件學(xué)報(bào),2011(4)∶718-731.

    [4]Zhang Zhang,Xiao Jingfa,Wu Jiayan,et al.ParaAT∶A parallel tool for constructing multiple protein-coding DNA alignments[J].Biochemical Biophysical Research Communications,2012,419(4)∶779-781.

    [5]徐新海,楊學(xué)軍,林宇斐,等.一種面向CPU-GPU異構(gòu)系統(tǒng)的容錯(cuò)方法[J].軟件學(xué)報(bào),2011,22(10):2538-2552.

    [6]NVIDIA Cooperation.CUDA programming guide version 5.0[EB/OL].[2013-05-15].http://docs.nvidia.com/cuda/cudac-programming-guide/.

    [7]TIANYI D H,TAREK S A.Reducing branch divergence in GPU programs[A].In Proceedings of the Fourth Workshop on General Purpose Processing on Graphics Processing Units[C].London∶ACM.2012.

    [8]IMEN C,AHCENE B,NOUREDINE M.Reducing thread divergence in GPU-based b&b applied to the flow-shop problem[A].In Proceedings of the 9th International Conference on Parallel[C].Berlin∶Springer-Verlag.2011.

    [9]Rutgers and UCSD.Protein Data Bank[DB/OL].[2009-02-24].http://www.rcsb.org/pdb/explore/explore.do?structureId=2gli.

    [10]GENE A.Validity of the single processor approach to achieving large-scale computing capabilities[A].In Proceedings of the April 18-20(AFIPS′67)[C].New York∶ACM.1967.

    猜你喜歡
    方法
    中醫(yī)特有的急救方法
    中老年保健(2021年9期)2021-08-24 03:52:04
    高中數(shù)學(xué)教學(xué)改革的方法
    化學(xué)反應(yīng)多變幻 “虛擬”方法幫大忙
    變快的方法
    兒童繪本(2020年5期)2020-04-07 17:46:30
    學(xué)習(xí)方法
    可能是方法不對(duì)
    用對(duì)方法才能瘦
    Coco薇(2016年2期)2016-03-22 02:42:52
    最有效的簡(jiǎn)單方法
    山東青年(2016年1期)2016-02-28 14:25:23
    四大方法 教你不再“坐以待病”!
    Coco薇(2015年1期)2015-08-13 02:47:34
    賺錢方法
    丁香六月天网| 街头女战士在线观看网站| 精品久久久久久电影网| 久久综合国产亚洲精品| 免费黄色在线免费观看| 一区二区日韩欧美中文字幕| 国产免费福利视频在线观看| 国产一区二区在线观看av| 综合色丁香网| 久久精品久久久久久久性| 国产又爽黄色视频| 满18在线观看网站| 最新在线观看一区二区三区 | 一级片'在线观看视频| 大片免费播放器 马上看| 啦啦啦啦在线视频资源| 免费看不卡的av| 午夜福利网站1000一区二区三区| 亚洲国产欧美在线一区| 色婷婷久久久亚洲欧美| 一区福利在线观看| 香蕉丝袜av| 哪个播放器可以免费观看大片| 国产人伦9x9x在线观看| 日韩中文字幕欧美一区二区 | 一级片免费观看大全| 久久久久视频综合| 免费在线观看视频国产中文字幕亚洲 | 国产av码专区亚洲av| 亚洲精品,欧美精品| 高清av免费在线| 国产在视频线精品| 嫩草影视91久久| 在线观看www视频免费| 欧美在线黄色| 亚洲av欧美aⅴ国产| 成年女人毛片免费观看观看9 | 亚洲国产精品成人久久小说| 一级爰片在线观看| av有码第一页| 大香蕉久久网| 丝袜人妻中文字幕| 少妇的丰满在线观看| 久久99热这里只频精品6学生| 韩国av在线不卡| 亚洲精品久久成人aⅴ小说| 亚洲第一区二区三区不卡| 在线观看免费午夜福利视频| 亚洲精品久久午夜乱码| 如何舔出高潮| 久久女婷五月综合色啪小说| 色婷婷av一区二区三区视频| 国产成人a∨麻豆精品| 亚洲成人一二三区av| 久久亚洲国产成人精品v| 久久国产精品大桥未久av| 亚洲欧洲国产日韩| 午夜免费鲁丝| 在线亚洲精品国产二区图片欧美| 19禁男女啪啪无遮挡网站| 国产精品无大码| 精品第一国产精品| 欧美 日韩 精品 国产| 国产精品二区激情视频| 久久久久精品久久久久真实原创| 天天躁日日躁夜夜躁夜夜| 国产深夜福利视频在线观看| av.在线天堂| 女人精品久久久久毛片| a级毛片黄视频| 免费少妇av软件| 免费观看人在逋| 性少妇av在线| 亚洲欧美清纯卡通| 搡老岳熟女国产| 亚洲国产精品一区二区三区在线| 欧美少妇被猛烈插入视频| 欧美日韩av久久| 久久狼人影院| 久久99精品国语久久久| 啦啦啦在线免费观看视频4| 丝袜在线中文字幕| 久久精品人人爽人人爽视色| 爱豆传媒免费全集在线观看| 老司机在亚洲福利影院| 国产伦理片在线播放av一区| 国产欧美日韩综合在线一区二区| 丝袜美足系列| 巨乳人妻的诱惑在线观看| 天天添夜夜摸| 亚洲综合色网址| 成年人午夜在线观看视频| 美女中出高潮动态图| 亚洲一码二码三码区别大吗| 91老司机精品| 人人妻,人人澡人人爽秒播 | 欧美精品一区二区大全| 亚洲国产精品一区三区| 男人舔女人的私密视频| 美女脱内裤让男人舔精品视频| 极品人妻少妇av视频| 亚洲婷婷狠狠爱综合网| 男女无遮挡免费网站观看| 电影成人av| 多毛熟女@视频| 国产精品国产三级国产专区5o| 久久亚洲国产成人精品v| 波多野结衣一区麻豆| 亚洲色图综合在线观看| 熟妇人妻不卡中文字幕| 永久免费av网站大全| 自拍欧美九色日韩亚洲蝌蚪91| www日本在线高清视频| 亚洲国产毛片av蜜桃av| 视频在线观看一区二区三区| 久久 成人 亚洲| 国产精品蜜桃在线观看| 亚洲成人av在线免费| 日韩大片免费观看网站| 97精品久久久久久久久久精品| 亚洲视频免费观看视频| 视频区图区小说| 热re99久久国产66热| 丝袜喷水一区| 久久久亚洲精品成人影院| 制服人妻中文乱码| 极品少妇高潮喷水抽搐| 日韩大码丰满熟妇| 色播在线永久视频| 热99久久久久精品小说推荐| 日本av免费视频播放| 久久99热这里只频精品6学生| 久久精品国产a三级三级三级| 99精国产麻豆久久婷婷| 黑人猛操日本美女一级片| 青春草国产在线视频| 不卡视频在线观看欧美| 一本色道久久久久久精品综合| 成年人午夜在线观看视频| 蜜桃在线观看..| 亚洲美女黄色视频免费看| xxxhd国产人妻xxx| 国产男女超爽视频在线观看| 少妇人妻久久综合中文| 精品少妇一区二区三区视频日本电影 | 大香蕉久久网| 欧美中文综合在线视频| 三上悠亚av全集在线观看| 看十八女毛片水多多多| 一级片'在线观看视频| 亚洲欧美激情在线| 黑丝袜美女国产一区| 在线观看三级黄色| 悠悠久久av| 午夜福利在线免费观看网站| www.熟女人妻精品国产| 女人被躁到高潮嗷嗷叫费观| 精品福利永久在线观看| 精品少妇久久久久久888优播| 熟妇人妻不卡中文字幕| 1024视频免费在线观看| 丝袜美足系列| a级片在线免费高清观看视频| av福利片在线| 曰老女人黄片| 夫妻性生交免费视频一级片| 啦啦啦在线观看免费高清www| 少妇人妻精品综合一区二区| 成人亚洲欧美一区二区av| 亚洲成国产人片在线观看| 黄色视频在线播放观看不卡| 成人亚洲精品一区在线观看| 大片电影免费在线观看免费| 三上悠亚av全集在线观看| 欧美xxⅹ黑人| 久久 成人 亚洲| 欧美在线黄色| 亚洲av成人精品一二三区| 不卡av一区二区三区| 国产一卡二卡三卡精品 | 99热全是精品| 免费女性裸体啪啪无遮挡网站| 制服人妻中文乱码| 国产日韩欧美在线精品| 婷婷色综合大香蕉| 高清欧美精品videossex| 亚洲人成网站在线观看播放| 99久久精品国产亚洲精品| 久久久久精品性色| 伦理电影大哥的女人| 欧美日韩亚洲国产一区二区在线观看 | 不卡av一区二区三区| 国产一卡二卡三卡精品 | 婷婷成人精品国产| 午夜福利网站1000一区二区三区| 亚洲成人手机| 一本一本久久a久久精品综合妖精| 日韩不卡一区二区三区视频在线| 精品少妇一区二区三区视频日本电影 | 国产欧美日韩综合在线一区二区| 亚洲精品第二区| 人人妻人人澡人人爽人人夜夜| av网站在线播放免费| 在线观看www视频免费| 叶爱在线成人免费视频播放| 综合色丁香网| 国产精品麻豆人妻色哟哟久久| 18在线观看网站| 丝袜美足系列| 亚洲国产精品999| 菩萨蛮人人尽说江南好唐韦庄| 最近的中文字幕免费完整| 久久久久久久久久久免费av| 久久午夜综合久久蜜桃| 精品一区二区三卡| 操美女的视频在线观看| 国产精品蜜桃在线观看| 国产极品天堂在线| 一二三四中文在线观看免费高清| 青春草视频在线免费观看| 在线观看免费午夜福利视频| 日韩中文字幕欧美一区二区 | 可以免费在线观看a视频的电影网站 | 交换朋友夫妻互换小说| 精品少妇久久久久久888优播| 巨乳人妻的诱惑在线观看| 免费观看a级毛片全部| 国产老妇伦熟女老妇高清| 九草在线视频观看| 国产一区二区 视频在线| 亚洲国产毛片av蜜桃av| 亚洲国产av新网站| 国产无遮挡羞羞视频在线观看| 人体艺术视频欧美日本| 啦啦啦中文免费视频观看日本| 在线观看一区二区三区激情| 男人操女人黄网站| 国产精品免费视频内射| 免费女性裸体啪啪无遮挡网站| 日韩成人av中文字幕在线观看| 欧美97在线视频| 免费在线观看黄色视频的| 久久精品国产综合久久久| 午夜福利,免费看| 国产成人免费观看mmmm| 飞空精品影院首页| 亚洲精品,欧美精品| 超碰成人久久| 90打野战视频偷拍视频| 欧美久久黑人一区二区| 免费少妇av软件| 欧美日韩福利视频一区二区| 一级黄片播放器| 国产av精品麻豆| 午夜免费鲁丝| 亚洲第一青青草原| 日韩av免费高清视频| 亚洲精品国产av蜜桃| 亚洲一区二区三区欧美精品| 久久精品亚洲熟妇少妇任你| 午夜福利在线免费观看网站| 国产成人免费无遮挡视频| 成人午夜精彩视频在线观看| 精品久久蜜臀av无| 国产在线一区二区三区精| 汤姆久久久久久久影院中文字幕| 男人爽女人下面视频在线观看| www.av在线官网国产| 国产激情久久老熟女| 国产乱人偷精品视频| 亚洲精品,欧美精品| 成年人午夜在线观看视频| 国产精品一区二区在线观看99| 国产精品三级大全| 国产一区有黄有色的免费视频| 久久精品久久久久久噜噜老黄| 日韩 亚洲 欧美在线| 久久国产精品大桥未久av| 欧美日韩视频高清一区二区三区二| avwww免费| 精品国产一区二区久久| 丝袜在线中文字幕| 国产极品天堂在线| 免费黄色在线免费观看| 制服丝袜香蕉在线| 欧美 亚洲 国产 日韩一| 啦啦啦 在线观看视频| 午夜久久久在线观看| 黑人巨大精品欧美一区二区蜜桃| 亚洲天堂av无毛| 97精品久久久久久久久久精品| 天堂中文最新版在线下载| 中文字幕高清在线视频| 亚洲,欧美,日韩| 久久久精品94久久精品| 日韩欧美一区视频在线观看| 777米奇影视久久| 亚洲av成人不卡在线观看播放网 | 99国产精品免费福利视频| 欧美日韩亚洲国产一区二区在线观看 | 在线观看www视频免费| 国产精品久久久久久人妻精品电影 | 熟女av电影| 18禁裸乳无遮挡动漫免费视频| 男人舔女人的私密视频| 日韩伦理黄色片| 久久久久久免费高清国产稀缺| 欧美精品一区二区大全| 国产精品久久久久久精品电影小说| videosex国产| 日韩成人av中文字幕在线观看| 热99久久久久精品小说推荐| 男女国产视频网站| 亚洲精品久久成人aⅴ小说| 国产成人欧美在线观看 | 国产精品.久久久| 免费av中文字幕在线| 国产亚洲av片在线观看秒播厂| 69精品国产乱码久久久| 亚洲成人一二三区av| 亚洲精品日韩在线中文字幕| 欧美精品亚洲一区二区| 交换朋友夫妻互换小说| 在现免费观看毛片| 亚洲欧美一区二区三区黑人| 免费日韩欧美在线观看| 天堂中文最新版在线下载| 丰满少妇做爰视频| 国产熟女欧美一区二区| 欧美人与性动交α欧美软件| 一级毛片黄色毛片免费观看视频| 最近手机中文字幕大全| 丝袜美腿诱惑在线| 天天躁日日躁夜夜躁夜夜| 国产免费现黄频在线看| 91精品国产国语对白视频| 中文字幕人妻丝袜一区二区 | 久久久国产精品麻豆| xxxhd国产人妻xxx| 国产欧美亚洲国产| 久久青草综合色| 亚洲久久久国产精品| 久久久久精品国产欧美久久久 | 欧美日韩精品网址| 国产精品 国内视频| 亚洲国产精品成人久久小说| 老鸭窝网址在线观看| 欧美日韩av久久| 国产福利在线免费观看视频| 老司机在亚洲福利影院| 亚洲激情五月婷婷啪啪| 下体分泌物呈黄色| 日韩av不卡免费在线播放| 国产老妇伦熟女老妇高清| 欧美xxⅹ黑人| 曰老女人黄片| 麻豆av在线久日| 亚洲欧美中文字幕日韩二区| 欧美xxⅹ黑人| 下体分泌物呈黄色| 亚洲精品视频女| a 毛片基地| 亚洲欧美精品自产自拍| 国产精品嫩草影院av在线观看| 夫妻性生交免费视频一级片| 亚洲欧美清纯卡通| 高清黄色对白视频在线免费看| 国语对白做爰xxxⅹ性视频网站| 色网站视频免费| 久久99热这里只频精品6学生| 亚洲在久久综合| 久久99热这里只频精品6学生| 亚洲第一青青草原| 天堂8中文在线网| 亚洲,欧美,日韩| 午夜免费鲁丝| 伦理电影大哥的女人| 午夜福利视频精品| 亚洲人成77777在线视频| 少妇猛男粗大的猛烈进出视频| 女人久久www免费人成看片| 久久99热这里只频精品6学生| 宅男免费午夜| 最近中文字幕高清免费大全6| 国产一区二区三区综合在线观看| 亚洲 欧美一区二区三区| 欧美黄色片欧美黄色片| 亚洲 欧美一区二区三区| 制服人妻中文乱码| av网站在线播放免费| 欧美97在线视频| 高清av免费在线| 午夜免费观看性视频| 亚洲av福利一区| 国产一区有黄有色的免费视频| 日韩中文字幕视频在线看片| 欧美av亚洲av综合av国产av | 精品亚洲成国产av| 一级片'在线观看视频| 日韩精品免费视频一区二区三区| 亚洲欧洲日产国产| 尾随美女入室| 高清不卡的av网站| 精品酒店卫生间| 久久97久久精品| 久久久久久久久久久免费av| 一本大道久久a久久精品| 国产免费现黄频在线看| 精品午夜福利在线看| www.自偷自拍.com| 亚洲免费av在线视频| 国产乱人偷精品视频| 一本—道久久a久久精品蜜桃钙片| 天天影视国产精品| 午夜久久久在线观看| 国产午夜精品一二区理论片| 成人18禁高潮啪啪吃奶动态图| a级片在线免费高清观看视频| 久久人人爽av亚洲精品天堂| 欧美人与善性xxx| 观看美女的网站| 深夜精品福利| 黄色毛片三级朝国网站| 国产在线视频一区二区| 久久精品国产综合久久久| 久久人妻熟女aⅴ| 国产精品久久久久久精品电影小说| 一级,二级,三级黄色视频| 亚洲精品第二区| 日本av免费视频播放| 亚洲精品国产av蜜桃| 日本vs欧美在线观看视频| 国产精品久久久av美女十八| 日本wwww免费看| 中文字幕制服av| 日本av免费视频播放| 熟妇人妻不卡中文字幕| 日韩中文字幕视频在线看片| 久久毛片免费看一区二区三区| 久久人人爽av亚洲精品天堂| 成人亚洲精品一区在线观看| 老司机亚洲免费影院| 永久免费av网站大全| 韩国av在线不卡| 婷婷色综合大香蕉| 成人三级做爰电影| 新久久久久国产一级毛片| av片东京热男人的天堂| 日韩精品免费视频一区二区三区| 黑人欧美特级aaaaaa片| 90打野战视频偷拍视频| 久久久久网色| 黄片无遮挡物在线观看| 男人舔女人的私密视频| 一本—道久久a久久精品蜜桃钙片| 久久精品熟女亚洲av麻豆精品| 2018国产大陆天天弄谢| 赤兔流量卡办理| 亚洲精品久久午夜乱码| 国产熟女欧美一区二区| 97精品久久久久久久久久精品| 男人操女人黄网站| 亚洲伊人色综图| 少妇的丰满在线观看| 赤兔流量卡办理| 久久青草综合色| 99精国产麻豆久久婷婷| 国产黄色视频一区二区在线观看| 一区在线观看完整版| 日日啪夜夜爽| 久久久久国产一级毛片高清牌| 国产精品久久久av美女十八| 国产片内射在线| 19禁男女啪啪无遮挡网站| 日韩免费高清中文字幕av| 国产av码专区亚洲av| 下体分泌物呈黄色| 一区在线观看完整版| 女人高潮潮喷娇喘18禁视频| 亚洲精品国产一区二区精华液| 亚洲人成网站在线观看播放| 一二三四中文在线观看免费高清| 男女免费视频国产| 一级a爱视频在线免费观看| 国产黄频视频在线观看| 韩国精品一区二区三区| 一级片'在线观看视频| 赤兔流量卡办理| 中文字幕av电影在线播放| 黄色视频不卡| 王馨瑶露胸无遮挡在线观看| 亚洲欧美精品自产自拍| 精品国产超薄肉色丝袜足j| 欧美日韩一级在线毛片| 乱人伦中国视频| 国产麻豆69| 在线观看免费视频网站a站| 午夜免费鲁丝| 黄频高清免费视频| 女性被躁到高潮视频| 电影成人av| 女的被弄到高潮叫床怎么办| 如日韩欧美国产精品一区二区三区| 在线亚洲精品国产二区图片欧美| 中文精品一卡2卡3卡4更新| 热99久久久久精品小说推荐| 国产成人免费无遮挡视频| 在线看a的网站| 亚洲成人一二三区av| 十分钟在线观看高清视频www| 国产精品国产av在线观看| 91精品国产国语对白视频| 国产成人精品久久二区二区91 | 免费女性裸体啪啪无遮挡网站| 亚洲精品成人av观看孕妇| 少妇人妻 视频| 亚洲欧美色中文字幕在线| 成人影院久久| 日韩,欧美,国产一区二区三区| 不卡av一区二区三区| 街头女战士在线观看网站| 精品亚洲成a人片在线观看| 久久久久国产一级毛片高清牌| 一本一本久久a久久精品综合妖精| 咕卡用的链子| 亚洲av国产av综合av卡| 久久久久久久国产电影| 国产精品三级大全| 欧美激情极品国产一区二区三区| 亚洲国产精品成人久久小说| 少妇被粗大猛烈的视频| 在线观看免费视频网站a站| 日日爽夜夜爽网站| av网站免费在线观看视频| 三上悠亚av全集在线观看| 亚洲精品国产一区二区精华液| 亚洲国产日韩一区二区| 美女国产高潮福利片在线看| 亚洲人成电影观看| www.自偷自拍.com| 亚洲精品国产色婷婷电影| 看免费成人av毛片| 色婷婷久久久亚洲欧美| 中文字幕人妻熟女乱码| 欧美97在线视频| 国产av国产精品国产| 在线天堂最新版资源| 亚洲精品第二区| 99久久精品国产亚洲精品| a级片在线免费高清观看视频| 人妻人人澡人人爽人人| 街头女战士在线观看网站| 国产精品麻豆人妻色哟哟久久| 欧美日韩视频高清一区二区三区二| 久久人妻熟女aⅴ| 老司机靠b影院| 精品久久久久久电影网| 久久亚洲国产成人精品v| 亚洲精品国产av成人精品| 天堂中文最新版在线下载| 飞空精品影院首页| 国产精品三级大全| 欧美最新免费一区二区三区| 青春草国产在线视频| 18禁裸乳无遮挡动漫免费视频| 亚洲在久久综合| 一边摸一边抽搐一进一出视频| 日本一区二区免费在线视频| 国产午夜精品一二区理论片| 人妻人人澡人人爽人人| 欧美日韩视频精品一区| 日韩精品免费视频一区二区三区| av卡一久久| 欧美日韩成人在线一区二区| 一级爰片在线观看| 一二三四中文在线观看免费高清| 视频在线观看一区二区三区| 日本91视频免费播放| 看免费av毛片| 日韩一本色道免费dvd| 极品少妇高潮喷水抽搐| 欧美日韩视频高清一区二区三区二| 国产成人精品久久久久久| 国产欧美日韩一区二区三区在线| 国产精品亚洲av一区麻豆 | 美国免费a级毛片| 亚洲人成电影观看| 天天影视国产精品| 久久韩国三级中文字幕| 麻豆av在线久日| 波多野结衣av一区二区av| 精品一区在线观看国产| av.在线天堂| 午夜精品国产一区二区电影| 一边亲一边摸免费视频| 伦理电影大哥的女人| 麻豆乱淫一区二区| 最近最新中文字幕免费大全7| 亚洲精品美女久久av网站| 国产精品免费视频内射| 最新在线观看一区二区三区 | 制服丝袜香蕉在线| 成年av动漫网址| 成人漫画全彩无遮挡| 成年美女黄网站色视频大全免费| 国产精品久久久久久久久免| 精品一区在线观看国产| 欧美xxⅹ黑人| 黑人巨大精品欧美一区二区蜜桃| 黑人欧美特级aaaaaa片| 国产精品久久久久久人妻精品电影 | 亚洲精品,欧美精品| 精品国产一区二区三区四区第35| 99久久综合免费| 亚洲精品久久久久久婷婷小说|