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

    一種基于GPU的高性能稀疏卷積神經(jīng)網(wǎng)絡(luò)優(yōu)化*

    2018-02-26 10:12:52邢座程陳頊顥
    計算機(jī)工程與科學(xué) 2018年12期
    關(guān)鍵詞:共享內(nèi)存指針線程

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

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

    1 引言

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

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

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

    本篇論文主要貢獻(xiàn)在以下幾點(diǎn):

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

    (2)采用最大限度的線程映射,充分利用GPU的硬件計算資源,防止產(chǎn)生稀疏結(jié)構(gòu)運(yùn)算對GPU計算資源的浪費(fèi)。

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

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

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

    2 背景

    這一節(jié)介紹實(shí)現(xiàn)卷積運(yùn)算的幾種方式,并說明了它們各自的局限性,從而闡述本文的研究意義和研究背景。

    2.1 降維方式

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

    Table 1 Description of convolution parameters表1 卷積參數(shù)描述

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

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

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

    2.2 直接稀疏卷積

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

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

    該批次任務(wù)下,直接稀疏卷積結(jié)果可以表示為:Om=Wm·Ivirtual。其中矩陣Ivirtual是由列向量I調(diào)整起始指針?biāo)玫降摹D敲?,我們可以進(jìn)一步簡化結(jié)果為:Om,y,x=Wm·Iy·W+x。所有輸出通道下的稀疏向量Wm構(gòu)成稀疏矩陣WSparse,采用行壓縮存儲CSR(Compressed Spares Row)格式,存儲如圖3所示。數(shù)組value記錄矩陣Wsparse中的非零元素。數(shù)組colidx記錄每個非零元素在矩陣Wsparse中的列指針。數(shù)組rowptr記錄矩陣Wsparse中每一行起始元素在value中的指針。

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

    3 設(shè)計與實(shí)現(xiàn)

    本節(jié)介紹本文所提方法的具體實(shí)現(xiàn)和優(yōu)化。由于權(quán)重刪減后SCNN產(chǎn)生了大量稀疏數(shù)據(jù)結(jié)構(gòu),而傳統(tǒng)的降維方式并不能保證稀疏矩陣卷積的計算性能,本文采用全新的直接稀疏卷積來替代降維方式,彌補(bǔ)性能損失。除此以外,GPU的體系結(jié)構(gòu)特征需要在實(shí)現(xiàn)過程中對線程映射、任務(wù)分配以及內(nèi)存管理進(jìn)行更多的考慮和優(yōu)化。

    3.1 概述

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

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

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

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

    考慮到實(shí)際的CNN模型中,所有卷積層經(jīng)過權(quán)重刪減后的稀疏度存在差異,我們通過下列方式來計算當(dāng)前卷積層的稀疏度:

    其中,Nnonzero為當(dāng)前卷積層的所有非零元素數(shù)目,M為當(dāng)前卷積層輸出通道數(shù),kernel_size為卷積核規(guī)模大小。

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

    3.2 并行策略

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

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

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

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

    輸出點(diǎn)(m,y,x)與線程 Thread(z,y,x)一一對應(yīng)。

    通過分別對過程一和過程二實(shí)現(xiàn)并行化,我們在GPU的架構(gòu)下實(shí)現(xiàn)了直接稀疏卷積。在實(shí)際的測試中,這一實(shí)現(xiàn)的具體性能并沒有達(dá)到預(yù)期效果(這一點(diǎn)將在第4節(jié)具體說明)。所以,接下來增加了對數(shù)據(jù)局部性的考慮,對實(shí)現(xiàn)的并行策略進(jìn)行了進(jìn)一步優(yōu)化。

    3.3 局部性優(yōu)化

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

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

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

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

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

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

    通過增加共享內(nèi)存以及對數(shù)據(jù)局部性的考慮,實(shí)驗(yàn)結(jié)果最終達(dá)到了預(yù)期性能。相比未優(yōu)化的直接稀疏卷積,本文在GPU上實(shí)現(xiàn)了更為高效的性能。

    4 性能評估

    4.1 總體性能

    實(shí)驗(yàn)采用的GPU型號為GTX 1060 3 GB。設(shè)置稀疏度閾值為0.6,batchsize為128。訓(xùn)練好的AlexNet模型包含5層卷積層,每層的稀疏度根據(jù)公式計算的結(jié)果如表2所示。

    Table 2 Parameters of AlexNet convolution layers表2 AlexNet卷積層參數(shù)

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

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

    4.2 執(zhí)行時間分析

    為了進(jìn)一步說明優(yōu)化后的性能提升,在Alex-Net模型基礎(chǔ)上對比本文的設(shè)計與原有CAFFE框架下所實(shí)現(xiàn)的卷積神經(jīng)網(wǎng)絡(luò)。CAFFE通過cu-BLAS提供的函數(shù)接口在GPU上主要實(shí)現(xiàn)了降維。cuBLAS是在GPU上實(shí)現(xiàn)的CUDA數(shù)學(xué)函數(shù)庫。此外,CAFFE還采用了cuSPARSE庫來優(yōu)化處理稀疏卷積。給出了batchsize為64時的AlexNet模型各層執(zhí)行時間對比,如圖8所示。

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

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

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

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

    由于各層刪減后的數(shù)據(jù)規(guī)則性也會對實(shí)驗(yàn)結(jié)果產(chǎn)生一定的影響,所以在本文優(yōu)化方案的測試結(jié)果中也出現(xiàn)了某些層的性能損失。但是,相比cu-BLAS和cuSPARSE,本文方法對高稀疏度層的優(yōu)化加速效果顯著??傮w來說,本文優(yōu)化方法實(shí)現(xiàn)了基于GPU架構(gòu)的SCNN加速優(yōu)化。

    5 相關(guān)工作

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

    6 結(jié)束語

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

    猜你喜歡
    共享內(nèi)存指針線程
    通過QT實(shí)現(xiàn)進(jìn)程間的通信
    偷指針的人
    娃娃畫報(2019年5期)2019-06-17 16:58:10
    為什么表的指針都按照順時針方向轉(zhuǎn)動
    基于PCI總線的多處理器協(xié)同機(jī)制研究
    淺談linux多線程協(xié)作
    基于改進(jìn)Hough變換和BP網(wǎng)絡(luò)的指針儀表識別
    電測與儀表(2015年5期)2015-04-09 11:30:42
    QNX下PEX8311多路實(shí)時數(shù)據(jù)采集的驅(qū)動設(shè)計
    電子世界(2014年21期)2014-04-29 06:41:36
    ARM Cortex—MO/MO+單片機(jī)的指針變量替換方法
    一種高效RTAI 共享內(nèi)存管理層的研究與實(shí)現(xiàn)*
    Linux線程實(shí)現(xiàn)技術(shù)研究
    少妇粗大呻吟视频| 日本 av在线| 国产精品香港三级国产av潘金莲| АⅤ资源中文在线天堂| 在线观看午夜福利视频| 狂野欧美激情性xxxx| 国产成人av激情在线播放| 91九色精品人成在线观看| 亚洲三区欧美一区| 国产男靠女视频免费网站| 亚洲成人国产一区在线观看| 亚洲狠狠婷婷综合久久图片| 波多野结衣一区麻豆| 99久久综合精品五月天人人| 91九色精品人成在线观看| 亚洲精品中文字幕在线视频| 久久天堂一区二区三区四区| 免费搜索国产男女视频| 成熟少妇高潮喷水视频| 久久精品亚洲精品国产色婷小说| 嫩草影视91久久| 午夜免费观看网址| 夜夜爽天天搞| 老司机福利观看| 久久国产亚洲av麻豆专区| 久99久视频精品免费| 禁无遮挡网站| e午夜精品久久久久久久| 人人澡人人妻人| 国产成人精品久久二区二区免费| 久久久久久免费高清国产稀缺| 中文字幕人成人乱码亚洲影| 极品教师在线免费播放| 精品国产一区二区久久| 一级毛片精品| 女人被躁到高潮嗷嗷叫费观| 91麻豆精品激情在线观看国产| 叶爱在线成人免费视频播放| 国产麻豆成人av免费视频| 国内久久婷婷六月综合欲色啪| 一本大道久久a久久精品| 18禁观看日本| 免费高清视频大片| 欧美色视频一区免费| 国产激情久久老熟女| 国产熟女xx| 欧美一区二区精品小视频在线| 婷婷六月久久综合丁香| 老鸭窝网址在线观看| 亚洲免费av在线视频| 啦啦啦韩国在线观看视频| 九色亚洲精品在线播放| 在线观看免费视频网站a站| 国产精品自产拍在线观看55亚洲| 久久这里只有精品19| 国产日韩一区二区三区精品不卡| 黑人巨大精品欧美一区二区mp4| 国产av在哪里看| 久久香蕉国产精品| 欧美日韩福利视频一区二区| 久久 成人 亚洲| 亚洲中文字幕日韩| 黄片小视频在线播放| 少妇熟女aⅴ在线视频| 不卡一级毛片| 亚洲第一青青草原| 1024视频免费在线观看| 狂野欧美激情性xxxx| 人妻丰满熟妇av一区二区三区| 变态另类成人亚洲欧美熟女 | 99riav亚洲国产免费| 91麻豆av在线| 丁香六月欧美| 亚洲 欧美 日韩 在线 免费| 国产私拍福利视频在线观看| 高清毛片免费观看视频网站| xxx96com| 不卡一级毛片| 国产一区二区三区综合在线观看| 国产熟女午夜一区二区三区| 免费在线观看日本一区| 一区二区三区高清视频在线| av有码第一页| 国产成人一区二区三区免费视频网站| 亚洲av成人一区二区三| 日本五十路高清| 欧美日本中文国产一区发布| bbb黄色大片| 夜夜躁狠狠躁天天躁| 又紧又爽又黄一区二区| 国产精品,欧美在线| 中国美女看黄片| 亚洲欧美精品综合久久99| 欧美乱码精品一区二区三区| 91成人精品电影| 成人免费观看视频高清| 国产麻豆69| 精品一区二区三区av网在线观看| 妹子高潮喷水视频| 国产私拍福利视频在线观看| 在线观看免费日韩欧美大片| 50天的宝宝边吃奶边哭怎么回事| 亚洲专区中文字幕在线| 国内毛片毛片毛片毛片毛片| 老司机午夜福利在线观看视频| 91九色精品人成在线观看| av天堂在线播放| 变态另类丝袜制服| 国产精品美女特级片免费视频播放器 | 91精品国产国语对白视频| 丝袜在线中文字幕| 一个人免费在线观看的高清视频| 国产97色在线日韩免费| 一边摸一边抽搐一进一出视频| 国产xxxxx性猛交| av视频免费观看在线观看| 老熟妇仑乱视频hdxx| 欧美日韩乱码在线| 99久久久亚洲精品蜜臀av| 在线十欧美十亚洲十日本专区| √禁漫天堂资源中文www| 在线观看免费午夜福利视频| 高清毛片免费观看视频网站| 中国美女看黄片| 欧美成人一区二区免费高清观看 | 久久久久久人人人人人| 嫁个100分男人电影在线观看| 男女下面进入的视频免费午夜 | 制服丝袜大香蕉在线| 亚洲av电影在线进入| 最近最新免费中文字幕在线| bbb黄色大片| 99riav亚洲国产免费| 成人免费观看视频高清| 亚洲人成电影免费在线| 波多野结衣av一区二区av| 精品一区二区三区视频在线观看免费| 亚洲一区中文字幕在线| 不卡av一区二区三区| 免费在线观看视频国产中文字幕亚洲| 香蕉久久夜色| 熟妇人妻久久中文字幕3abv| 女人爽到高潮嗷嗷叫在线视频| 好看av亚洲va欧美ⅴa在| 国产亚洲精品av在线| 老司机深夜福利视频在线观看| 婷婷精品国产亚洲av在线| 亚洲中文av在线| 18禁黄网站禁片午夜丰满| 国产成人精品久久二区二区免费| 丁香欧美五月| 久久精品91无色码中文字幕| 日日爽夜夜爽网站| 国产亚洲精品第一综合不卡| 夜夜爽天天搞| 最近最新中文字幕大全电影3 | 国产一级毛片七仙女欲春2 | 免费在线观看亚洲国产| 色婷婷久久久亚洲欧美| 国产一卡二卡三卡精品| 别揉我奶头~嗯~啊~动态视频| 国产精品电影一区二区三区| 久久久国产欧美日韩av| 中亚洲国语对白在线视频| 国产片内射在线| 亚洲中文字幕日韩| 久久久精品欧美日韩精品| 桃色一区二区三区在线观看| 欧美日韩精品网址| 中文字幕人妻丝袜一区二区| 极品人妻少妇av视频| 亚洲av片天天在线观看| 精品国产美女av久久久久小说| a在线观看视频网站| 91老司机精品| 最好的美女福利视频网| 电影成人av| 欧美大码av| 亚洲精品国产区一区二| 亚洲性夜色夜夜综合| 可以免费在线观看a视频的电影网站| 欧美乱色亚洲激情| 国产一卡二卡三卡精品| 免费不卡黄色视频| 午夜两性在线视频| 亚洲七黄色美女视频| 成人18禁在线播放| 亚洲全国av大片| 午夜亚洲福利在线播放| 脱女人内裤的视频| 美女午夜性视频免费| 中国美女看黄片| 可以在线观看毛片的网站| 午夜免费成人在线视频| 日韩高清综合在线| 欧美乱码精品一区二区三区| 88av欧美| 亚洲第一欧美日韩一区二区三区| 国产99白浆流出| 淫秽高清视频在线观看| 级片在线观看| 国产又爽黄色视频| 日本一区二区免费在线视频| 国产精品免费视频内射| 精品久久久久久久毛片微露脸| 中出人妻视频一区二区| 妹子高潮喷水视频| 69av精品久久久久久| 国产aⅴ精品一区二区三区波| 亚洲全国av大片| 亚洲av熟女| 两个人看的免费小视频| 成人三级黄色视频| 日韩大码丰满熟妇| 一本大道久久a久久精品| 91麻豆精品激情在线观看国产| 无人区码免费观看不卡| 国语自产精品视频在线第100页| 黄片小视频在线播放| 首页视频小说图片口味搜索| 女人被躁到高潮嗷嗷叫费观| 久久精品91无色码中文字幕| 亚洲av片天天在线观看| 国产精品一区二区三区四区久久 | 国产精品秋霞免费鲁丝片| 亚洲va日本ⅴa欧美va伊人久久| 国产精品久久久久久精品电影 | 黑人欧美特级aaaaaa片| 亚洲成人精品中文字幕电影| 后天国语完整版免费观看| 搡老熟女国产l中国老女人| 中文字幕人成人乱码亚洲影| 天堂动漫精品| 丝袜美腿诱惑在线| 久久久精品国产亚洲av高清涩受| 国产亚洲欧美精品永久| 青草久久国产| av片东京热男人的天堂| 麻豆国产av国片精品| www.熟女人妻精品国产| 女同久久另类99精品国产91| 亚洲精品一区av在线观看| 男女之事视频高清在线观看| 99在线视频只有这里精品首页| 成人欧美大片| 日韩中文字幕欧美一区二区| 日韩欧美一区视频在线观看| 又大又爽又粗| 亚洲 欧美一区二区三区| 一区二区三区高清视频在线| 啦啦啦 在线观看视频| 国产麻豆成人av免费视频| 日日摸夜夜添夜夜添小说| 99国产极品粉嫩在线观看| 国产精品日韩av在线免费观看 | 精品国产超薄肉色丝袜足j| 国产精品爽爽va在线观看网站 | 性少妇av在线| 免费无遮挡裸体视频| 人人妻人人爽人人添夜夜欢视频| 国产成人一区二区三区免费视频网站| 成人国产综合亚洲| 91精品三级在线观看| 久久草成人影院| 日韩免费av在线播放| 亚洲熟妇中文字幕五十中出| 性少妇av在线| 亚洲成人免费电影在线观看| 久久精品人人爽人人爽视色| 亚洲欧美精品综合久久99| 老汉色∧v一级毛片| 亚洲狠狠婷婷综合久久图片| 亚洲第一欧美日韩一区二区三区| www国产在线视频色| 久久久久精品国产欧美久久久| 成人国产一区最新在线观看| videosex国产| 国内精品久久久久久久电影| 在线十欧美十亚洲十日本专区| 国产一区在线观看成人免费| 天天添夜夜摸| 国内精品久久久久久久电影| 久久 成人 亚洲| 老熟妇仑乱视频hdxx| 少妇裸体淫交视频免费看高清 | 如日韩欧美国产精品一区二区三区| 成熟少妇高潮喷水视频| 欧美日韩瑟瑟在线播放| 国产野战对白在线观看| 国产精品精品国产色婷婷| 别揉我奶头~嗯~啊~动态视频| 精品久久久久久久久久免费视频| 亚洲精品一卡2卡三卡4卡5卡| 国产蜜桃级精品一区二区三区| av网站免费在线观看视频| 精品日产1卡2卡| 国产视频一区二区在线看| 精品国产一区二区三区四区第35| 久久中文字幕一级| 精品国产一区二区久久| 91老司机精品| 999久久久精品免费观看国产| 亚洲,欧美精品.| 满18在线观看网站| 在线观看免费午夜福利视频| 久久久久精品国产欧美久久久| 国产亚洲精品av在线| 韩国精品一区二区三区| 91精品三级在线观看| 亚洲国产精品成人综合色| 在线观看免费午夜福利视频| 中文字幕最新亚洲高清| 国语自产精品视频在线第100页| 伊人久久大香线蕉亚洲五| 一级a爱视频在线免费观看| 午夜免费成人在线视频| 涩涩av久久男人的天堂| 欧美在线黄色| 亚洲精品久久成人aⅴ小说| 亚洲国产精品999在线| 91麻豆精品激情在线观看国产| 国产成人精品无人区| 99久久精品国产亚洲精品| 色综合婷婷激情| 12—13女人毛片做爰片一| 一进一出好大好爽视频| 两个人视频免费观看高清| 国产极品粉嫩免费观看在线| 亚洲欧美精品综合久久99| 一级毛片精品| 一级毛片高清免费大全| 国产精品99久久99久久久不卡| 一级a爱片免费观看的视频| 一个人免费在线观看的高清视频| 色综合欧美亚洲国产小说| 免费在线观看视频国产中文字幕亚洲| 亚洲精品在线观看二区| АⅤ资源中文在线天堂| 久久久久国内视频| 国产成人欧美在线观看| 欧美激情 高清一区二区三区| 国产亚洲欧美在线一区二区| 大香蕉久久成人网| 亚洲欧美精品综合久久99| 久久国产亚洲av麻豆专区| 精品久久久久久久人妻蜜臀av | 欧美国产精品va在线观看不卡| av天堂在线播放| 露出奶头的视频| 国产精品秋霞免费鲁丝片| 国产亚洲精品av在线| 涩涩av久久男人的天堂| 免费观看人在逋| 99久久精品国产亚洲精品| 一本大道久久a久久精品| 成人三级黄色视频| 免费女性裸体啪啪无遮挡网站| 非洲黑人性xxxx精品又粗又长| 国产亚洲精品av在线| 99热只有精品国产| 欧美另类亚洲清纯唯美| 日韩 欧美 亚洲 中文字幕| 99在线视频只有这里精品首页| 国语自产精品视频在线第100页| 最近最新中文字幕大全免费视频| 精品高清国产在线一区| 亚洲成人国产一区在线观看| 亚洲欧美精品综合久久99| 免费少妇av软件| 一级作爱视频免费观看| 亚洲电影在线观看av| 久久久久国内视频| 欧美乱色亚洲激情| 午夜福利免费观看在线| 热re99久久国产66热| 久久精品国产亚洲av香蕉五月| 两人在一起打扑克的视频| 亚洲精品国产区一区二| tocl精华| 亚洲中文字幕日韩| 男人操女人黄网站| 久久久久久免费高清国产稀缺| 久久九九热精品免费| 一边摸一边抽搐一进一出视频| 亚洲黑人精品在线| 欧美绝顶高潮抽搐喷水| 少妇裸体淫交视频免费看高清 | 亚洲精华国产精华精| 欧美日本视频| 狠狠狠狠99中文字幕| 色播在线永久视频| 日本 av在线| 午夜精品久久久久久毛片777| 国产亚洲欧美98| 在线国产一区二区在线| 一区二区三区高清视频在线| 19禁男女啪啪无遮挡网站| 可以在线观看的亚洲视频| 波多野结衣巨乳人妻| 午夜影院日韩av| 日本 欧美在线| 午夜久久久久精精品| 日韩中文字幕欧美一区二区| 美女免费视频网站| 两个人视频免费观看高清| 久久热在线av| 亚洲一区二区三区色噜噜| 久久久久久久午夜电影| 亚洲一卡2卡3卡4卡5卡精品中文| 日日夜夜操网爽| 亚洲av成人不卡在线观看播放网| 亚洲国产欧美一区二区综合| 黄色a级毛片大全视频| 涩涩av久久男人的天堂| 色播亚洲综合网| 99国产精品99久久久久| 波多野结衣av一区二区av| 天天添夜夜摸| 国内毛片毛片毛片毛片毛片| 国产精品日韩av在线免费观看 | 国产亚洲精品av在线| 精品国产美女av久久久久小说| 国产97色在线日韩免费| 亚洲国产高清在线一区二区三 | 999精品在线视频| 午夜福利在线观看吧| 亚洲情色 制服丝袜| 在线观看免费视频网站a站| 久久久国产成人免费| 麻豆久久精品国产亚洲av| 人成视频在线观看免费观看| 国产1区2区3区精品| 国产一区二区在线av高清观看| 亚洲男人的天堂狠狠| 搡老岳熟女国产| 丝袜美腿诱惑在线| svipshipincom国产片| 国语自产精品视频在线第100页| 亚洲,欧美精品.| 久久香蕉激情| www日本在线高清视频| 精品乱码久久久久久99久播| bbb黄色大片| 成人国产综合亚洲| 99精品欧美一区二区三区四区| 老司机午夜十八禁免费视频| 欧美精品啪啪一区二区三区| 亚洲av成人av| 日日爽夜夜爽网站| 日本精品一区二区三区蜜桃| 最新美女视频免费是黄的| 19禁男女啪啪无遮挡网站| 人人澡人人妻人| 午夜福利,免费看| 国内精品久久久久精免费| 在线免费观看的www视频| 多毛熟女@视频| 村上凉子中文字幕在线| 亚洲av电影不卡..在线观看| 国产成人系列免费观看| 老司机福利观看| 亚洲欧美日韩无卡精品| 精品无人区乱码1区二区| 999久久久国产精品视频| 真人一进一出gif抽搐免费| 午夜视频精品福利| 女警被强在线播放| 黑人巨大精品欧美一区二区蜜桃| 久久久久国产一级毛片高清牌| 久久国产乱子伦精品免费另类| 可以免费在线观看a视频的电影网站| 久久婷婷人人爽人人干人人爱 | 91成人精品电影| 国产成人影院久久av| 丝袜在线中文字幕| 女警被强在线播放| 午夜免费成人在线视频| 国产精品 欧美亚洲| 国产精品香港三级国产av潘金莲| 午夜免费鲁丝| 久久精品影院6| 日日夜夜操网爽| 久久人妻熟女aⅴ| 好看av亚洲va欧美ⅴa在| 亚洲性夜色夜夜综合| 777久久人妻少妇嫩草av网站| 亚洲黑人精品在线| 精品国产超薄肉色丝袜足j| 一区二区三区精品91| 一边摸一边做爽爽视频免费| 国产国语露脸激情在线看| 波多野结衣高清无吗| 亚洲人成伊人成综合网2020| 日本 av在线| 免费久久久久久久精品成人欧美视频| 男人舔女人下体高潮全视频| 50天的宝宝边吃奶边哭怎么回事| 老汉色av国产亚洲站长工具| 我的亚洲天堂| 男女之事视频高清在线观看| 国产精品国产高清国产av| 男人操女人黄网站| 日本免费一区二区三区高清不卡 | av中文乱码字幕在线| 免费在线观看黄色视频的| 久久人人97超碰香蕉20202| 欧美成人免费av一区二区三区| 亚洲国产精品999在线| 97超级碰碰碰精品色视频在线观看| 午夜a级毛片| 丰满人妻熟妇乱又伦精品不卡| 两个人免费观看高清视频| 一级毛片精品| 国产精品自产拍在线观看55亚洲| 亚洲五月色婷婷综合| 大香蕉久久成人网| 身体一侧抽搐| 久久草成人影院| av在线播放免费不卡| 成熟少妇高潮喷水视频| 一进一出好大好爽视频| 久久久久久久久免费视频了| 免费看美女性在线毛片视频| 女性被躁到高潮视频| 51午夜福利影视在线观看| 女性被躁到高潮视频| 中国美女看黄片| 日韩欧美国产一区二区入口| 亚洲一区高清亚洲精品| 丰满的人妻完整版| 久久人妻av系列| 黄色视频不卡| 免费一级毛片在线播放高清视频 | 多毛熟女@视频| 国产精品日韩av在线免费观看 | 日韩欧美三级三区| 好看av亚洲va欧美ⅴa在| 久久中文字幕人妻熟女| 激情在线观看视频在线高清| 日韩有码中文字幕| 此物有八面人人有两片| 亚洲免费av在线视频| 一级a爱片免费观看的视频| 亚洲avbb在线观看| 视频区欧美日本亚洲| 一级毛片女人18水好多| 久久久国产精品麻豆| 在线观看66精品国产| 18禁裸乳无遮挡免费网站照片 | 精品久久久精品久久久| 国内毛片毛片毛片毛片毛片| 999久久久精品免费观看国产| 人人妻人人澡人人看| 男女做爰动态图高潮gif福利片 | 日韩欧美国产一区二区入口| 亚洲人成伊人成综合网2020| 国产私拍福利视频在线观看| 亚洲欧美日韩高清在线视频| 亚洲熟妇熟女久久| 午夜激情av网站| 黄片大片在线免费观看| 丁香欧美五月| www.精华液| 久久精品国产亚洲av香蕉五月| 亚洲情色 制服丝袜| 久久精品aⅴ一区二区三区四区| 免费不卡黄色视频| 韩国精品一区二区三区| 国内久久婷婷六月综合欲色啪| 人人妻人人爽人人添夜夜欢视频| 波多野结衣巨乳人妻| 丝袜人妻中文字幕| 一区二区三区高清视频在线| svipshipincom国产片| 亚洲久久久国产精品| 热re99久久国产66热| 高清毛片免费观看视频网站| 好男人电影高清在线观看| 久久久久精品国产欧美久久久| 99精品欧美一区二区三区四区| 午夜久久久久精精品| 免费女性裸体啪啪无遮挡网站| 国产精品1区2区在线观看.| 亚洲成a人片在线一区二区| 97人妻精品一区二区三区麻豆 | 91国产中文字幕| 精品电影一区二区在线| 天堂影院成人在线观看| 国产午夜精品久久久久久| 国产乱人伦免费视频| 午夜福利视频1000在线观看 | 亚洲成国产人片在线观看| 久久国产精品男人的天堂亚洲| 麻豆一二三区av精品| 亚洲九九香蕉| 亚洲欧美精品综合一区二区三区| 国产亚洲欧美98| 久久精品国产亚洲av香蕉五月| av电影中文网址| 一级片免费观看大全| 村上凉子中文字幕在线| 校园春色视频在线观看| 日本精品一区二区三区蜜桃| 免费无遮挡裸体视频| 一区二区三区精品91| 国产不卡一卡二| 两个人看的免费小视频| 琪琪午夜伦伦电影理论片6080| 精品少妇一区二区三区视频日本电影| 久久香蕉激情| 久久精品国产清高在天天线| 欧美日韩亚洲国产一区二区在线观看| 欧美日本亚洲视频在线播放| 99在线人妻在线中文字幕| 给我免费播放毛片高清在线观看|