999精品在线视频,手机成人午夜在线视频,久久不卡国产精品无码,中日无码在线观看,成人av手机在线观看,日韩精品亚洲一区中文字幕,亚洲av无码人妻,四虎国产在线观看 ?

一種基于GPU的高性能稀疏卷積神經網絡優化*

2018-02-26 10:12:52邢座程陳頊顥
計算機工程與科學 2018年12期
關鍵詞:優化

方 程,邢座程,陳頊顥,張 洋

(國防科技大學計算機學院,湖南長沙410073)

1 引言

CNN(Convolutional Neural Network)目前作為深度學習領域中的一個重要模型,在計算機圖像[1]、語音識別[2]、游戲比賽[3]以及機器人[4]等方面都扮演著越來越重要的角色。但是,隨著CNN的發展,CNN網絡規模和網絡層數不斷增加,參數規模也變得越來越龐大。1990年,早期的卷積神經網絡模型用于手寫識別使用的參數數量不到100M[5]。20 年后,AlexNet[1]和 VGG[6]分別使用了61M和138M個參數來對1 000個圖像進行分類。顯而易見,這對CNN實現過程中的硬件資源、網絡結構、算法優化等各方面都會產生諸多的挑戰。CNN 加速[7]、CNN 參數的量化分析與研究[8]、參數規模的縮小與權重刪減[9-12]都將成為熱門的研究方向。

為了應對CNN對計算需求量的不斷增加,采用高性能的 GPU 已經成為加速 CNN[13,14]的一項重要措施。此外還有研究給出了壓縮CNN的解決方案。壓縮CNN方法主要分為兩類:一類是基于分解[15,16],另一類是基于刪減[17,18]。基于刪減的方法是在保證網絡訓練和測試結果精確度[8]沒有損失的前提下減少參數數量。權重刪減下的深度壓縮[9,10]方式能夠使 AlexNet和 VGG-16 的參數規模分別縮小9倍和13倍,在CPU和GPU的架構下實現了3~4倍的加速。但是,權重刪減的性能提升遠遠落后于實際減少乘累加操作的性能提升,特別是在GPU的硬件設備上,這一類似的性能損失經常發生。同時,有研究提出了全新的直接稀疏卷積算法[19],它在CPU的架構下相比原有的稠密算法在AlexNet卷積層上實現了3.1×~7.3×的加速。對于訓練好的CNN來說,卷積層的卷積運算是測試過程運行時間最主要的部分。所以,對卷積層的卷積運算優化成為解決該加速優化問題中的關鍵路徑。

由于權重刪減的方式犧牲了數據的規則性,SCNN(Sparse Convolutional Neural Network)內部產生了大量的稀疏計算成分。稀疏數據處理與GPU體系結構特性不匹配[20]。原有GPU架構下對卷積核提供卷積數學運算的cuBLAS和cuSPARSE并不能很好地應對這種不匹配。同時,GPU和CPU在體系結構上存在的差異使得很多針對CPU優化的稀疏卷積算法并不能在GPU上適用。我們采用了一個高效的直接稀疏卷積算法,對其在GPU的平臺上進行優化,從而解決權重刪減產生稀疏數據所帶來的性能損失。

本篇論文主要貢獻在以下幾點:

(1)針對卷積層關鍵優化路徑上完成直接稀疏卷積算法[19]在GPU架構上的并行化實現,打破GPU上采用傳統稠密算法的局限性,給出了一種可行且高效的GPU加速SCNN方案。

(2)采用最大限度的線程映射,充分利用GPU的硬件計算資源,防止產生稀疏結構運算對GPU計算資源的浪費。

(3)采用最優的任務調度,合理安排每個單線程的任務工作,減少線程同步過程中某一部分線程等待時間,提高資源利用率。

(4)充分利用直接稀疏卷積算法數據處理過程中的數據局部性,增加數據復用,對于同一block下的所有線程,采用共享內存來減少數據訪存時間。

最終我們在CAFFE(Convolutional Architecture for Fast Feature Embedding)架構下所實現的稀疏卷積神經網絡,對同一訓練好的 AlexNet、GoogleNet、ResNet,在 GPU GTX1060 上,與 CAFFE本身搭建的由cuBLAS和cuSPARSE所提供的數學庫支持的卷積層進行測試對比。相比cuBLAS的實現,我們在 AlexNet、GoogleNet、ResNet上性能提升分別達到1.07 × ~1.23 ×、1.17 × ~3.51 ×、1.32×~5.00×的加速比。相比cuSPARSE的實現,在 AlexNet、GoogleNet、ResNet上性能提升分別達到1.31 × ~1.42 ×、1.09 × ~2.00 ×、1.07 × ~3.22×的加速比。

2 背景

這一節介紹實現卷積運算的幾種方式,并說明了它們各自的局限性,從而闡述本文的研究意義和研究背景。

2.1 降維方式

目前很多CNN卷積層的卷積操作都是通過降維方式實現的[18]。圖1所示是一個簡單的用降維方式實現卷積的例子,圖中參數可參考表1。

Table 1 Description of convolution parameters表1 卷積參數描述

假設輸入特征矩陣的batchsize為1,其輸入通道數為C,輸入特征矩陣大小為H×W,預輸出通道數為M,每一個卷積核的實際大小為R×S(在實際應用中,可以通過設置步長U來控制卷積核在輸入特征矩陣上的局部感知區域的位置,后文我們假設U默認為1)。那么一共有M個卷積核,每個卷積核包含C個通道。降維方式通過將輸入特征矩陣和卷積核分別以行展開的方式生成新的特征矩陣Ilowering和卷積核矩陣Wlowering。那么,最終卷積的計算過程可以表示為:

降維方式將卷積運算轉換為矩陣乘法。在基礎線性代數子程序庫BLAS(Basic Linear Algebra Subprograms)中,GEMM(GEneralized Matrix Multiplication)的函數接口實現了兩個稠密矩陣的乘法運算。在CAFFE的框架下,CNN卷積層中卷積運算所采用的方式也是降維方式,具體是通過im2col函數和GEMM函數實現。此外,CAFFE還支持了CUDA版本下由cuBLAS所提供的GPU架構下并行實現的降維方式。

在降維方式展開生成新的特征矩陣Ilowering的過程中,卷積核所感知的局部區域重疊部分的元素都進行了多次重復復制,增加了存儲開銷。特別是在SCNN中,這種大量的數據重復復制浪費了大量的存儲資源。此外,GEMM是針對稠密矩陣實現的矩陣乘法,對于處理稀疏矩陣浪費了GPU大量的計算資源。所以,我們需要一個針對GPU實現的稀疏卷積運算。

2.2 直接稀疏卷積

直接稀疏卷積(Direct Sparse Convolutions)作為一種全新的卷積方式在2017年的ICLR會議上被首次提出[19]。該算法在CPU的架構下相比原有的算法在AlexNet卷積層上實現了3.1×~7.3×的加速。

相比降維方式,直接稀疏卷積去除了輸入特征矩陣中的數據重復復制。該算法將卷積核矩陣的規模擴展到輸入矩陣的相同大小。對于延展后的卷積核行展開生成向量Wm,其長度為C×H×W。由于有M個卷積核,對每一個卷積核進行延展后得到了M×(C×H×W)的權重矩陣。對于該批次任務下的輸入矩陣以行展開的方式形成列向量I,其長度為C×H×W。那么,在計算卷積的過程中,對于不同感知區域的元素可以通過調整向量I的起始指針,使得卷積核映射到正確的局部區域。其具體算法如圖2所示。

該批次任務下,直接稀疏卷積結果可以表示為:Om=Wm·Ivirtual。其中矩陣Ivirtual是由列向量I調整起始指針所得到的。那么,我們可以進一步簡化結果為:Om,y,x=Wm·Iy·W+x。所有輸出通道下的稀疏向量Wm構成稀疏矩陣WSparse,采用行壓縮存儲CSR(Compressed Spares Row)格式,存儲如圖3所示。數組value記錄矩陣Wsparse中的非零元素。數組colidx記錄每個非零元素在矩陣Wsparse中的列指針。數組rowptr記錄矩陣Wsparse中每一行起始元素在value中的指針。

直接稀疏卷積將卷積運算抽象成稀疏向量Wm對稠密向量Iy·W+x的內積。此外,由于 SCNN采用CSR或CSC(Compressed Sparse Column)的稀疏數據存儲格式,對于運算過程中的延展實際上并沒有增加存儲開銷,只是調整了矩陣中非零元素的行列指針。相比降維方式,直接稀疏卷積更適合在GPU上實現SCNN。

3 設計與實現

本節介紹本文所提方法的具體實現和優化。由于權重刪減后SCNN產生了大量稀疏數據結構,而傳統的降維方式并不能保證稀疏矩陣卷積的計算性能,本文采用全新的直接稀疏卷積來替代降維方式,彌補性能損失。除此以外,GPU的體系結構特征需要在實現過程中對線程映射、任務分配以及內存管理進行更多的考慮和優化。

3.1 概述

直接稀疏卷積的實現主要由兩部分組成:(1)數據預處理,主要完成對卷積核矩陣的延展,生成稀疏向量Wm和稠密向量I;(2)計算過程,主要完成所有的MAC操作,并準確更新計算過程中的指針。

第(1)部分如圖4所示。在這里,權重矩陣為M×(C×R×S)的稀疏矩陣,按照CSR格式存儲于物理內存中。對于輸出通道m中的第j個非零元素(c,y,x)有:

其中col=colidx[j]。那么,延展后的權重矩陣大小為M ×(C ×H ×W),同一個非零元素(c,y,x)的CSR格式存儲下的列指針更新為:colidx[j]=(c*H+y)*W+x。

直接稀疏卷積的計算過程可以表示為:Om,y,x=Wm·Iy·W+x。其核心在于實現稀疏向量Wm與稠密向量Iy·W+x的內積運算。對于計算輸出矩陣中的點(m,y,x),需要完成的MAC操作數取決于稀疏向量Wm的非零元素數目。由于對同一輸出通道m中的所有點,稀疏向量Wm是恒定不變的,所以計算這些輸出節點所需要的MAC操作數相等。在直接稀疏卷積算法中矩陣Ivirtual是由向量I生成,其中每一個列向量Iy·W+x的起始指針所指向的元素為I[y·W+x]。根據這一特點,我們僅將向量I的元素常駐內存,而不是存儲整個稠密矩陣Ivirtual。

考慮到實際的CNN模型中,所有卷積層經過權重刪減后的稀疏度存在差異,我們通過下列方式來計算當前卷積層的稀疏度:

其中,Nnonzero為當前卷積層的所有非零元素數目,M為當前卷積層輸出通道數,kernel_size為卷積核規模大小。

對于不同稀疏度的卷積層,我們設置一個閾值。稀疏度大于該閾值的卷積層采用優化后的直接稀疏卷積方式,小于該閾值的卷積層則仍采用原有的降維方式。對于稠密數據和稀疏數據的分別處理,使得對于任意稀疏度的卷積層都能夠實現最佳的計算性能,可以最大限度提高整個網絡的運行性能。由于在最終實驗過程中采用了IntelSkimcaffe開源項目(https://github.com/IntelLabs/Skim-Caffe)中的稀疏CNN網絡結構,CNN中的卷積層的稀疏度集中在0和0.7~0.96這兩個區域,所以設置閾值僅僅是排除了稀疏度為0的稠密層。

3.2 并行策略

相比CPU,GPU擁有更多的處理核心,如何合理分配和充分利用這些處理核心是本文設計的關鍵。接下來我們將分別介紹直接稀疏卷積兩個過程的并行策略。

對于過程一,即圖5中所示oc=m時所有非零元素的列指針更新。

將整個權重矩陣進行延展就是更新權重矩陣內所有非零元素的列指針colidx。那么,我們設置線程Threadm完成稀疏向量Wm內所有非零元素的列指針更新。

對于過程二,每一個線程計算輸出特征矩陣中的一個點(m,y,x),如圖5所示。由于輸入特征矩陣Ivirtual中每一列向量是由稠密列向量I移動初始指針得到的,那么我們將稠密列向量I的所有數據常駐內存。當需要計算不同的輸出點(m,y,x)時,計算其對應列向量 Iy·W+x相對向量 I的偏移量pos,其計算公式為:pos=y·W+x。通過對向量I的初始指針增加pos偏移量得到對應向量Iy·W+x:*inputptr=input+pos。該過程避免了數據的重復復制,僅通過調整指針來完成當前輸出通道的全部計算。由于Wsparse作為稀疏矩陣采用CSR的格式存儲在物理內存中,第m個卷積核下所有非零元素對應存儲在rowptr[m]行。該行元素在物理內存中存儲非零元素的一維數組value中的起始位置為 row_strat(row_start=rowptr[m]),結束位置為row_end(row_end=rowptr[m+1])。那么,對于線程 Thread(z,y,x),需要完成下列計算:

輸出點(m,y,x)與線程 Thread(z,y,x)一一對應。

通過分別對過程一和過程二實現并行化,我們在GPU的架構下實現了直接稀疏卷積。在實際的測試中,這一實現的具體性能并沒有達到預期效果(這一點將在第4節具體說明)。所以,接下來增加了對數據局部性的考慮,對實現的并行策略進行了進一步優化。

3.3 局部性優化

由于輸入特征向量的數據復用,我們采用了Ivirtual的方式來減小帶寬需求。通過更改訪存指針來讀取向量I中的值。同樣地,為了增加Cache塊的命中率,希望優先計算同一輸出通道的值。由于實際測試性能達不到預期效果,我們增加了共享內存優化的版本,其具體映射規則如圖6所示。

對于輸出通道m,需要E×F個線程來完成計算任務。但是在實際情況中,GPU所能設置的最大block_size小于E×F,所以對于同一個block內的所有線程會在短時間內經常訪問向量Wm,直到該block內的所有線程完成計算。此時內存常駐的數據僅僅只有向量I和向量Wm。

由于同一個block下的所有線程都要求對Wm進行數據訪問,我們將Wm放入共享內存中,以減少Wm的數據訪存時間。共享內存對于同一block塊下的線程是共同可見的。考慮到GPU內共享內存大小的限制,將Wm分塊化,塊Tilei為特定長度的一維數組。將Tilei的長度設定為block下的線程總數,并使Tilei包含的數據能夠常駐共享內存。由于Wm采用CSR格式存儲,那么僅需將對應數組value和數組colidx的值存入共享內存。在計算輸出結果前,需要將Tilei對應的value和colidx寫入共享內存中的數組valueshared和數組colidxshared。由于 Tilei長度與線程數相等,那么對于線程Thread(z,y,x)需要完成的讀寫工作如下所示:

每個線程只需要將Tilei塊內一個元素的value和colidx數組值寫入共享內存。其中,blockDim.y為GPU線程設置中block塊在y方向上的維度大小,Tilesize為設置的Tile塊的長度。

為了防止讀后寫,為同一block下的所有線程增加同步操作。線程 Thread(z,y,x)將 Tilei數據寫入共享內存后進行等待,直到所有線程完成操作。當block 塊內所有線程完成同步后,線程 Thread(z,y,x)需要完成共享內存內向量Tilei與向量Iy·W+x的內積運算。其具體計算如下所示:

每一次累加操作后同步線程,當訪問共享內存未命中時,替換下一個Tile塊到共享內存。將每個塊替換下來的部分和保存在寄存器sum中,這樣當下一個塊被換進共享內存時,線程能夠正常工作。當輸出通道m所有Tile塊都被替換過后,將部分和sum輸出:output[(m·E+y)·F+x]=sum。輸出點(m,y,x)與線程 Thread(z,y,x)的映射關系與之前的一樣。

通過增加共享內存以及對數據局部性的考慮,實驗結果最終達到了預期性能。相比未優化的直接稀疏卷積,本文在GPU上實現了更為高效的性能。

4 性能評估

4.1 總體性能

實驗采用的GPU型號為GTX 1060 3 GB。設置稀疏度閾值為0.6,batchsize為128。訓練好的AlexNet模型包含5層卷積層,每層的稀疏度根據公式計算的結果如表2所示。

Table 2 Parameters of AlexNet convolution layers表2 AlexNet卷積層參數

對稀疏度大于0.6 的 CONV2、CONV3、CONV4和CONV5四個卷積層采用直接稀疏卷積的方式。設置一個block塊下的總線程數為1 024,那么每次替換進共享內存的Tile塊長度為1 024。實驗結果記錄了50 000次迭代中每一次迭代完成Forward過程所需的時間,每100次迭代的Forward執行時間取平均值,具體結果如圖7所示。其中Base為未優化初始版本的執行時間曲線,Tiled為增加共享內存優化后版本的執行時間曲線。

Tiled版本相比Base版本在各層上都有較大的性能提升。在各層上Tiled版本的性能分別提升了 46.7%、41.1%、41.5%、42.6%。這說明本文的優化在GPU架構上起到了實質性作用。通過增加共享內存,合理分配線程,增加數據復用,在GPU架構上實現直接稀疏卷積,實現了高效的稀疏卷積神經網絡優化。本文采用的直接稀疏卷積并行方式適應了GPU的體系結構特征,充分利用了硬件計算資源。為了進一步說明本文設計性能的優越性,將在4.2小節與現有的CNN卷積層實現進行性能對比。

4.2 執行時間分析

為了進一步說明優化后的性能提升,在Alex-Net模型基礎上對比本文的設計與原有CAFFE框架下所實現的卷積神經網絡。CAFFE通過cu-BLAS提供的函數接口在GPU上主要實現了降維。cuBLAS是在GPU上實現的CUDA數學函數庫。此外,CAFFE還采用了cuSPARSE庫來優化處理稀疏卷積。給出了batchsize為64時的AlexNet模型各層執行時間對比,如圖8所示。

同樣地,只列出了AlexNet中稀疏度大于0.6的四個卷積層以及它們的總執行時間。從圖8可以看到,相比cuBLAS實現方法,本文的優化方法僅在CONV2上有略微的性能損失,而在CONV3、CONV4、CONV5上的性能分別提升了41.1%、26.2%、40.1%,且總體性能提升了10%;相比 cuSPARSE實現方法,本文的優化方法在 CONV2、CONV3、CONV4、CONV5上的性能分別提升了29.6%、39.2%、47.1%、67.4%,且總體性能提升了41.1%。通過分析表2給出了未刪減前AlexNet各層結構參數,包括輸入特征矩陣大小、卷積核大小以及稀疏度。而對于 CONV3、CONV4、CONV5這三層來說,其稀疏度均高于CONV2的稀疏度,從而證明了本文的設計針對大規模高稀疏度數據有顯著優化效果。此外,圖9給出了不同batchsize下AlexNet各層的加速比。相比cuBLAS,本文的優化方法在batchsize為192時得到了最佳的加速比;相比cuSPARSE,當batchsize在32~64時性能最佳,從網絡整體性能來看,batchsize為64時加速性能最佳。這是由于batchsize過小或過大都會使在線負載任務過輕或過重,不能合理利用硬件計算資源。

本文還對GoogleNet和ResNet模型進行了測試,同樣也只給出了稀疏度大于0.6的卷積層的加速比,如圖10所示。

對于 GoogleNet,相比 cuBLAS,本文優化方法僅在低維度有1層性能有略微的損失,其余高維度稀疏層實現了 1.17× ~3.51×加速;相比 cuSPARSE,本文優化方法僅在高維度和低維度各出現了1層性能損失,其余各層實現了1.09×~2.00×加速;在總體性能上,相比cuBLAS和cuSPARSE的方式分別實現了1.34×和1.21×加速。對于ResNet,相比cuBLAS,本文優化方法在所有稀疏層實現了1.32× ~5.00×加速;相比 cuSPARSE,本文優化方法僅在高維度出現了2層性能損失,其余各層實現了1.07×~3.22×加速;在總體性能上,相比cuBLAS和cuSPARSE的方式分別實現了2.43×和1.97×加速。

移動學習中學習評價是在網絡課程學習的過程中對學生的學習過程和學習結果進行價值判斷的過程[5]。移動學習評價設計的缺失和無效已經成為制約網絡課程發揮實際效力的關鍵因素。54.5%的學生希望針對移動學習進行學習評價。移動學習可以通過在線反饋獲取學生的學習評價,Bb平臺提供給學生多種學習反饋方式,學生反映不一。58.6%的學生希望可以看到每題得分和總成績,通過每題的得分狀況,對自己掌握的知識進行針對性的學習;33.8%的學生需要教師評語,教師評語可以更加直觀、深入的評價學生測試中的問題,便于學生理解與反思。

由于各層刪減后的數據規則性也會對實驗結果產生一定的影響,所以在本文優化方案的測試結果中也出現了某些層的性能損失。但是,相比cu-BLAS和cuSPARSE,本文方法對高稀疏度層的優化加速效果顯著。總體來說,本文優化方法實現了基于GPU架構的SCNN加速優化。

5 相關工作

相比傳統意義上GPU加速CNN的實現方案[13,14],本文采用更優的數學內核和卷積運算算法,提高了整個系統的可優化程度。相比其他采用更合理的刪減方式來切合GPU硬件特性[20]的加速方案,本文所提供的加速方案具有更好的可移植性和可靠性,對數據預處理的消耗小。此外,在文獻[21]所提出的Escort優化版本上,本文改進了并行策略和映射規則,取得了更高的加速比。

6 結束語

本文通過在GPU上實現直接稀疏卷積算法,打破了GPU架構下傳統稠密算法對于稀疏結構處理的局限性,有效解決了權重刪減后SCNN在GPU上運行出現性能損失的問題。對于高稀疏度,甚至是GPU所不擅長處理的不規則數據,本文的設計仍然有著極大的優勢。相比CAFEE下cuBLAS的實現,本文方法在 AlexNet、GoogleNet、ResNet上的性能提升分別達到 1.07× ~1.23 ×、1.17× ~3.51×、1.32 × ~ 5.00 ×。相比 cuSPARSE 的實現,本文方法在 AlexNet、GoogleNet、ResNet上的性能提升分別達到1.31× ~1.42×、1.09 × ~2.00 ×、1.07 × ~3.22 ×。

猜你喜歡
優化
超限高層建筑結構設計與優化思考
房地產導刊(2022年5期)2022-06-01 06:20:14
PEMFC流道的多目標優化
能源工程(2022年1期)2022-03-29 01:06:28
民用建筑防煙排煙設計優化探討
關于優化消防安全告知承諾的一些思考
一道優化題的幾何解法
由“形”啟“數”優化運算——以2021年解析幾何高考題為例
圍繞“地、業、人”優化產業扶貧
今日農業(2020年16期)2020-12-14 15:04:59
事業單位中固定資產會計處理的優化
消費導刊(2018年8期)2018-05-25 13:20:08
4K HDR性能大幅度優化 JVC DLA-X8 18 BC
幾種常見的負載均衡算法的優化
電子制作(2017年20期)2017-04-26 06:57:45
主站蜘蛛池模板: 91人妻日韩人妻无码专区精品| 老司机午夜精品视频你懂的| 国产美女人喷水在线观看| 在线观看91香蕉国产免费| 91免费在线看| AV无码一区二区三区四区| 国产在线啪| Jizz国产色系免费| 亚洲国产黄色| 九九热这里只有国产精品| 国产亚洲高清在线精品99| av在线手机播放| 五月六月伊人狠狠丁香网| 老司机久久99久久精品播放| 国产精品尤物在线| 精品夜恋影院亚洲欧洲| 日本国产精品| 国产99视频在线| 亚洲高清资源| 狠狠综合久久久久综| 日韩无码白| 日韩经典精品无码一区二区| 不卡网亚洲无码| 久久久久久国产精品mv| 日韩精品成人在线| 亚洲综合第一区| 亚洲日韩AV无码一区二区三区人| 999国产精品| 九色免费视频| 99在线免费播放| 亚洲成aⅴ人片在线影院八| 99视频在线观看免费| 亚洲视频四区| 91午夜福利在线观看| 91久久偷偷做嫩草影院电| 视频一本大道香蕉久在线播放| 色婷婷啪啪| 在线色综合| 国产在线一区二区视频| 女人18毛片水真多国产| 国产欧美视频综合二区 | 国产在线观看91精品亚瑟| 最新亚洲av女人的天堂| 岛国精品一区免费视频在线观看| 永久成人无码激情视频免费| 黄色网址免费在线| 国产乱人伦AV在线A| 日韩精品一区二区深田咏美| 特级欧美视频aaaaaa| 国产精品无码影视久久久久久久| 日韩专区欧美| 97青青青国产在线播放| 亚洲IV视频免费在线光看| 一级高清毛片免费a级高清毛片| 国产高清在线丝袜精品一区| 国产午夜精品一区二区三| 亚洲精品福利视频| 91精品国产自产在线老师啪l| 在线永久免费观看的毛片| 中文字幕永久视频| 国产在线观看人成激情视频| 免费又爽又刺激高潮网址 | 任我操在线视频| 免费国产福利| 久青草免费视频| 51国产偷自视频区视频手机观看| 日本中文字幕久久网站| 9久久伊人精品综合| 男人天堂亚洲天堂| 成色7777精品在线| 波多野结衣一级毛片| 亚洲视频免费在线看| 久久综合AV免费观看| 91久久偷偷做嫩草影院电| 免费看a级毛片| 免费人成又黄又爽的视频网站| 性69交片免费看| 无码高潮喷水专区久久| 久久窝窝国产精品午夜看片| 国内老司机精品视频在线播出| 精品欧美一区二区三区在线| 精品久久777|