色哟哟视频在线观看-色哟哟视频在线-色哟哟欧美15最新在线-色哟哟免费在线观看-国产l精品国产亚洲区在线观看-国产l精品国产亚洲区久久

0
  • 聊天消息
  • 系統(tǒng)消息
  • 評論與回復
登錄后你可以
  • 下載海量資料
  • 學習在線課程
  • 觀看技術視頻
  • 寫文章/發(fā)帖/加入社區(qū)
會員中心
創(chuàng)作中心

完善資料讓更多小伙伴認識你,還能領取20積分哦,立即完善>

3天內(nèi)不再提示

解析OneFlow BatchNorm相關算子實現(xiàn)

jf_pmFSk4VX ? 來源:GiantPandaCV ? 2022-12-23 15:08 ? 次閱讀

0x1. 前言

在ResNet中(https://github.com/pytorch/vision/blob/main/torchvision/models/resnet.py),關于BatchNorm的調(diào)用一共有兩種模式,第一種是ReLU接在BN之后:

out=self.bn1(out)
out=self.relu(out)

另外一種模式是殘差結(jié)構(gòu)引入的 BNAddReLU 的模式:

out=self.bn2(out)

ifself.downsampleisnotNone:
identity=self.downsample(x)

out+=identity
out=self.relu(out)

我們知道在 CUDA 優(yōu)化中常見的一個技巧是將一些ElementWise的算子融合到之前的計算密集型算子如卷積,矩陣乘等。在OneFlow中針對上述兩種情況并且cudnn無法fuse時分別進行了fuse和優(yōu)化,本篇文章就來解析一下這里的代碼實現(xiàn),體會其中的CUDA優(yōu)化技巧。這里的源碼開源在OneFlow的github倉庫:「https://github.com/Oneflow-Inc/oneflow/blob/master/oneflow/user/kernels/normalization_kernel.cu」 。如果本文對你產(chǎn)生了啟發(fā),不妨為OneFlow投個star。

0x2. 代碼解析

0x2.1 CUDNN BatchNorm算子的實現(xiàn)和局限

我們先來看一下OneFlow中是如何使用CUDNN庫實現(xiàn)BatchNorm算子的。代碼見:https://github.com/Oneflow-Inc/oneflow/blob/master/oneflow/user/kernels/normalization_kernel.cu#L31-L244 。這段代碼中首先實現(xiàn)了一個getCudnnBatchNormMode工具函數(shù):

cudnnBatchNormMode_tgetCudnnBatchNormMode(constint64_tdim){
if(dim==2){
returnCUDNN_BATCHNORM_PER_ACTIVATION;
}elseif(ParseBooleanFromEnv("ONEFLOW_ENABLE_NHWC",false)){
returnCUDNN_BATCHNORM_SPATIAL_PERSISTENT;
}else{
//NOTE(LiangDepeng):ThenewCUDNN_BATCHNORM_SPATIAL_PERSISTENTmodewas
//introducedinCuDNN7forperformanceoptimization,butitresultsin
//accuracylossesinconvolutionmodelssuchasResNeXt-101and
//videoR(2+1)D.WewillfallbacktothenormalCUDNN_BATCHNORM_SPATIAL
returnCUDNN_BATCHNORM_SPATIAL;
}
}

這里的dim表示輸入Tensor的維度,比如形狀為的輸入Tensor,這里的維度就是4。然后這里涉及到三種不同的cudnnBatchNormMode_t,我們看一下CUDNN的文檔(https://docs.nvidia.com/deeplearning/cudnn/api/index.html#cudnnBatchNormMode_t):

bb6519d8-828a-11ed-8abf-dac502259ad0.png

可以看到 CUDNN_BATCHNORM_PER_ACTIVATION 被用于非卷積層,在OneFlow中只有當輸入Tensor的維度為2時才選取這種模式。而CUDNN_BATCHNORM_SPATIAL_PERSISTENT這種模式只有當輸入Tensor的數(shù)據(jù)排布為NHWC方式時才會啟用。而對于其它的模式,在OneFlow中一律選取CUDNN_BATCHNORM_SPATIAL模式。

接下來閱讀一下 InferDimSizeAndDataFormat 函數(shù):

voidInferDimSizeAndDataFormat(constShapeView&x_shape,constint32_taxis,int32_t*n,int32_t*c,
int32_t*h,int32_t*w,cudnnTensorFormat_t*format){
if(x_shape.Count(axis+1)==1){
if(axis==0){
*n=1;
*h=1;
}else{
*n=x_shape.At(0);
*h=x_shape.Count(1,axis);
}
*w=1;
*c=x_shape.At(axis);
*format=CUDNN_TENSOR_NHWC;
}else{
*n=x_shape.Count(0,axis);
*c=x_shape.At(axis);
*h=x_shape.Count(axis+1);
*w=1;
*format=CUDNN_TENSOR_NCHW;
}
}

這個函數(shù)會根據(jù)輸入Tensor的shape以及axis推斷這個Tensor的內(nèi)存排布是NCHW還是NHWC模式,并設置對應的n, c, h, w變量。

//推斷和設置cudnn中的Tensor描述符
voidInferXYCudnnTensorDesc(constShapeView&xy_shape,constDataType&data_type,
constint32_taxis,cudnnTensorDescriptor_txy_desc){
int32_tn,c,h,w;
cudnnTensorFormat_tformat;
//根據(jù)輸入Tensor的shape推斷format和n,c,h,w
InferDimSizeAndDataFormat(xy_shape,axis,&n,&c,&h,&w,&format);
//根據(jù)上述的推斷結(jié)果,設置Tensor的描述符
OF_CUDNN_CHECK(
cudnnSetTensor4dDescriptor(xy_desc,format,GetCudnnDataType(data_type),n,c,h,w));
}
//根據(jù)輸入Tensor的描述符xy_desc和cudnnBatchNormMode_t模式設置參數(shù)的描述符param_desc
voidInferParamCudnnTensorDesc(constcudnnTensorDescriptor_txy_desc,cudnnBatchNormMode_tmode,
cudnnTensorDescriptor_tparam_desc){
OF_CUDNN_CHECK(cudnnDeriveBNTensorDescriptor(param_desc,xy_desc,mode));
}
//這個類就是完整使用上述的工具函數(shù)的工具類,負責推斷cudnnBatchNorm接口需要的各種描述信息比如這里的xy_desc_,param_desc_,param_data_type_和param_size_
classCudnnTensorDescHelperfinal{
public:
OF_DISALLOW_COPY_AND_MOVE(CudnnTensorDescHelper);
CudnnTensorDescHelper(constShapeView&xy_shape,constDataType&data_type,constint32_taxis,
cudnnBatchNormMode_tmode){
OF_CUDNN_CHECK(cudnnCreateTensorDescriptor(&xy_desc_));
InferXYCudnnTensorDesc(xy_shape,data_type,axis,xy_desc_);
OF_CUDNN_CHECK(cudnnCreateTensorDescriptor(¶m_desc_));
InferParamCudnnTensorDesc(xy_desc_,mode,param_desc_);
intn,c,h,w,n_stride,c_stride,h_stride,w_stride;
OF_CUDNN_CHECK(cudnnGetTensor4dDescriptor(param_desc_,¶m_data_type_,&n,&c,&h,&w,
&n_stride,&c_stride,&h_stride,&w_stride));
param_size_=c;
}
~CudnnTensorDescHelper(){
OF_CUDNN_CHECK(cudnnDestroyTensorDescriptor(param_desc_));
OF_CUDNN_CHECK(cudnnDestroyTensorDescriptor(xy_desc_));
}

cudnnTensorDescriptor_txy_desc()const{returnxy_desc_;}

cudnnTensorDescriptor_tparam_desc()const{returnparam_desc_;}

voidCheckParamTensor(constuser_op::Tensor*tensor)const{
CHECK_NOTNULL(tensor);
CHECK_EQ(tensor->shape_view().NumAxes(),1);
CHECK_EQ(tensor->shape_view().At(0),param_size_);
CHECK_EQ(GetCudnnDataType(tensor->data_type()),param_data_type_);
}

private:
cudnnTensorDescriptor_txy_desc_=nullptr;
cudnnTensorDescriptor_tparam_desc_=nullptr;
cudnnDataType_tparam_data_type_;
int32_tparam_size_=0;
};

除了這些描述信息之外,我們還可以在cudnn提供的文檔中查看BatchNorm相關的算子一般還需要什么特殊的輸入信息。我們來看 cudnnBatchNormalizationForwardTrainingEx() 這個API :https://docs.nvidia.com/deeplearning/cudnn/api/index.html#cudnnBatchNormalizationForwardTrainingEx 。

bb7f7b34-828a-11ed-8abf-dac502259ad0.png

可以看到這個算子是 cudnnBatchNormalizationForwardTraining() 這個算子的擴展,擴展的內(nèi)容就是可以我們可以傳入額外的一個Activation的算子比如ReLU以及一個Add算子分別對應我們在前言中介紹的 ResNet 中的 BNReLU 和 BNAddReLU 模式。可以看到在這個算子接口中除了對輸入Tensor x,BN后需要add的輸入Tensor z以及輸出Tensor y的描述信息外,還需要指定workspace和reserveSpace,這個workspace是cudnn的BatchNorm以NHWC模式計算時需要的GPU內(nèi)存buffer,而reserveSpace則表示當前這個配置的BN算子至少還需要多少可以申請的GPU顯存(從文檔猜測應該是和BNReLU/BNAddReLU這倆Pattern相關)。

在OneFlow中, https://github.com/Oneflow-Inc/oneflow/blob/master/oneflow/user/kernels/normalization_kernel.cu#L126-L175 以及 https://github.com/Oneflow-Inc/oneflow/blob/master/oneflow/user/kernels/normalization_kernel.cu#L637-L684 就是為了推斷BN算子以及BN擴展的算子需要的額外GPU內(nèi)存大小,然后在OneFlow的內(nèi)存池中開辟一塊顯存供調(diào)用cudnn的 cudnnBatchNormalizationForwardTrainingEx() 和 cudnnBatchNormalizationBackwardEx() 接口時使用。

關于調(diào)用cudnn的BatchNorm相關的算子api,我們還需要注意一點,那就是要使用cudnn提供的擴展接口cudnnBatchNormalizationForwardTrainingEx() 和 cudnnBatchNormalizationBackwardEx() 還存在一些限制:

bbbeb740-828a-11ed-8abf-dac502259ad0.png

首先是cudnn版本的限制,然后對于CUDNN_BATCHNORM_OPS_BN_ADD_ACTIVATION的Op模式,輸入Tensor的通道數(shù)必須是4的倍數(shù),最后這個擴展Op必須在輸入Tensor的數(shù)據(jù)排布模式是NHWC時才能啟動。這些限制對應到OneFlow的代碼在:https://github.com/Oneflow-Inc/oneflow/blob/master/oneflow/core/job_rewriter/cudnn_fused_normalization_add_relu_pass.cpp#L79-L86 。

0x2.2 善用CUDA優(yōu)化打破cudnn的限制

上面提到要使用CUDNN的擴展算子有一系列限制,我們有沒有辦法打破這限制呢?有的。以ResNet為例,針對BNReLu和BNAddReLU這兩種Pattern,我們可以分別針對ReLU和AddReLU實現(xiàn)一個CUDA Kernel,相信入門CUDA的小伙伴寫這兩個算子都沒什么問題。但如何在考慮到Backward的時候把這兩個算子優(yōu)化到位呢?OneFlow給出了一個解決方案。

前向的CUDA實現(xiàn):https://github.com/Oneflow-Inc/oneflow/blob/master/oneflow/user/kernels/normalization_kernel.cu#L246-L272bcbd3036-828a-11ed-8abf-dac502259ad0.png

反向的CUDA實現(xiàn):https://github.com/Oneflow-Inc/oneflow/blob/master/oneflow/user/kernels/normalization_kernel.cu#L246-L272

bdaaa654-828a-11ed-8abf-dac502259ad0.png

以 ReLU 算子為例,前向的輸入為x,輸出為y,后向的輸入為dy和y,輸出dx。后向計算中的y僅用來判斷對應元素是否大于0,因此可以將y替換為由前向生成的bitset(對應上述代碼中的mask),理論上可以省掉ReLU的后向算子對冗余的y的訪問操作,減少約y大小的讀取,也對應約1/3的global memory訪問。對于ReLU/ReLUAdd這種ElementWise算子來說,GPU的帶寬是極容易成為瓶頸的,通過這種優(yōu)化可以大大提升ReLU和ReLUAdd算子的帶寬。

在 《OneFlow是如何做到世界上最快的深度學習框架》(https://zhuanlan.zhihu.com/p/271740706) 文章中已經(jīng)介紹到了這種基于bitmask優(yōu)化后向算子的方案。并且文章中給出了3種方案,但沒有給出對應的代碼實現(xiàn),實際上我只讀懂了第一種和第三種方案,接下來我們描述一下這兩種方案。

Bitset mask生成方案一:順序遍歷法

這種方法是讓每個CUDA線程連續(xù)讀取內(nèi)存中的8個元素,并根據(jù)每個元素是否大于0生成一個int8類型的mask,并寫入到最終的bitset mask中。這種訪問對于寫全局內(nèi)存是連續(xù)訪問的,但對于讀(Read)全局內(nèi)存,線程間內(nèi)存訪問不連續(xù),所以沒有充分合并內(nèi)存事務。下圖展示了這種方案讀寫內(nèi)存的示例:

bdc42034-828a-11ed-8abf-dac502259ad0.png

以ReLU為例子,這種方案的代碼實現(xiàn)如下:

template
__global__voidReluGpu(int64_tn,constT*x,T*y,int8_t*mask){
CUDA_1D_KERNEL_LOOP(i,n){
int8_tmask_val=0;
for(int32_tj=0;j0);
if(is_positive){
y[offset]=sum;
mask_val|=(1<

在這種方案中,每個thread需要連續(xù)讀的8個float32數(shù)據(jù),則相鄰線程每次加載數(shù)據(jù)的間隔為32 bytes = 4 bytes * 8。所以每個線程一次加載指令就要執(zhí)行一個32字節(jié)的內(nèi)存事務。故warp內(nèi)的線程間全局內(nèi)存訪問完全沒有合并,實際有效訪存帶寬僅為 1/8,訪存效率十分低下,性能很差。

Bitset mask生成方案三:warp同步法

我們可以采用warp級別的同步原語:__ballot_sync(unsigned mask, predicate),這個函數(shù)接收兩個參數(shù),第一個參數(shù)是warp中參與計算的線程掩碼,第二個參數(shù)是要參與判斷的bool值,返回一個32bit的mask,每個bit代表warp中各個線程傳入的元素是否大于0,最后由每個warp中的0號線程將生成的mask寫入global memory中。(idea可以參考NVIDIA的性能優(yōu)化博客:https://developer.nvidia.com/blog/using-cuda-warp-level-primitives/)

這種方案的示意圖如下:

poYBAGOlVFuAWMhwAACXY46xsP8359.jpg

以ReLU為例,代碼實現(xiàn)如下:

template
__global__voidReluGpu(int64_tn,constT*x,T*y,int32_t*mask){
constint32_tlane_id=threadIdx.x%kCudaWarpSize;//如果lane_id=0,表示當前線程是一個warp的0號線程
CUDA_1D_KERNEL_LOOP(i,n){
constboolis_positive=(x[i]>0);
int32_twarp_mask=__ballot_sync(__activemask(),static_cast(is_positive));
if(lane_id==0){mask[i/kCudaWarpSize]=warp_mask;}//0號線程將生成的mask寫入globalmemory
y[i]=is_positive?sum:0;
}
}

0x3. 性能

我們這里對比一下BNReLU這個Pattern在優(yōu)化前后的后向Kernel(也就是ReLU Grad Kernel)的性能和帶寬表現(xiàn),本次測試的環(huán)境為A100 PCIE 40G,使用Nsight Compute工具進行Profile。Profile的腳本為:

importoneflowasflow
bn=flow.nn.BatchNorm2d(num_features=32,eps=1e-5,momentum=0.1).to("cuda")
fused_bn=flow.nn.FusedBatchNorm2d(32).to("cuda")
bn.train()
fused_bn.train()

x=flow.randn(16,32,112,112).to("cuda").requires_grad_()

y=flow.relu(bn(x))#這個是未優(yōu)化的實現(xiàn)
#y=fused_bn(x)#打開這行代表啟用上述介紹的優(yōu)化
res=y.sum()
res.backward()
res_scalar=res.detach().cpu().numpy()

經(jīng)過多次測試,flow.relu(bn(x))中對應的ReLU的反向Kernel耗時大概為 「48.3us」,而fused_bn(x)中對應的ReLU的反向Kernel耗時大概為 「42.8us」 ,可以說明上述基于mask掩碼降低全局內(nèi)存訪問的優(yōu)化方法是有效的。而對于BNAddReLU的Pattern來說,則可以獲得更好的性能提升,因為ReluBackward相當于將這兩個ElementWise操作給fuse了。

0x4. 總結(jié)

這里暫時寫了一下個人看OneFlow Normalization 系列算子實現(xiàn)的理解。實際上我個人還是有一些疑問在,如果后續(xù)能搞清楚的話,會繼續(xù)補充和修改。

0x5. 相關鏈接

cudnn文檔:https://docs.nvidia.com/deeplearning/cudnn/api/index.html

oneflow代碼實現(xiàn):https://github.com/Oneflow-Inc/oneflow/blob/master/oneflow/user/kernels/normalization_kernel.cu

審核編輯:湯梓紅

聲明:本文內(nèi)容及配圖由入駐作者撰寫或者入駐合作網(wǎng)站授權(quán)轉(zhuǎn)載。文章觀點僅代表作者本人,不代表電子發(fā)燒友網(wǎng)立場。文章及其配圖僅供工程師學習之用,如有內(nèi)容侵權(quán)或者其他違規(guī)問題,請聯(lián)系本站處理。 舉報投訴
  • 代碼
    +關注

    關注

    30

    文章

    4779

    瀏覽量

    68526
  • 算子
    +關注

    關注

    0

    文章

    16

    瀏覽量

    7253
  • CUDA
    +關注

    關注

    0

    文章

    121

    瀏覽量

    13620
  • OneFlow
    +關注

    關注

    0

    文章

    9

    瀏覽量

    8802

原文標題:【BBuf的CUDA筆記】二,解析 OneFlow BatchNorm 相關算子實現(xiàn)

文章出處:【微信號:GiantPandaCV,微信公眾號:GiantPandaCV】歡迎添加關注!文章轉(zhuǎn)載請注明出處。

收藏 人收藏

    評論

    相關推薦

    OneFlow Softmax算子源碼解讀之WarpSoftmax

    寫在前面:近來筆者偶然間接觸了一個深度學習框架 OneFlow,所以這段時間主要在閱讀 OneFlow 框架的 cuda 源碼。官方源碼基于不同場景分三種方式實現(xiàn) Softmax,本文主要介紹其中一種的
    的頭像 發(fā)表于 01-08 09:24 ?842次閱讀
    <b class='flag-5'>OneFlow</b> Softmax<b class='flag-5'>算子</b>源碼解讀之WarpSoftmax

    OneFlow Softmax算子源碼解讀之BlockSoftmax

    寫在前面:筆者這段時間工作太忙,身心俱疲,博客停更了一段時間,現(xiàn)在重新?lián)炱饋怼1疚闹饕庾x OneFlow 框架的第二種 Softmax 源碼實現(xiàn)細節(jié),即 block 級別的 Softmax。
    的頭像 發(fā)表于 01-08 09:26 ?707次閱讀
    <b class='flag-5'>OneFlow</b> Softmax<b class='flag-5'>算子</b>源碼解讀之BlockSoftmax

    基于GFO算子的圖像增強算法如何去實現(xiàn)

    基于GFO算子(廣義模糊算子)的圖像增強算法如何去實現(xiàn)?怎樣對圖像增強算法進行分析?
    發(fā)表于 06-04 06:24

    TensorFlow、PyTorch,“后浪”OneFlow 有沒有機會

    TensorFlow、PyTorch,“后浪”OneFlow 有沒有機會 | 一流科技工程師成誠編者按:7月31日,一流科技在創(chuàng)業(yè)1300天后,他們宣布開源自研的深度學習框架OneFlow,此前,CSDN對CEO袁進輝進行了專訪。本文中,一流科技工程師成...
    發(fā)表于 07-27 08:24

    拉普拉斯算子的FPGA實現(xiàn)方法

    拉普拉斯算子的FPGA實現(xiàn)方法  引 言   在圖像處理系統(tǒng)中常需要對圖像進行預處理。由于圖像處理的數(shù)據(jù)量大,對于實時性要求高的系統(tǒng),采用軟件實現(xiàn)通常
    發(fā)表于 02-11 11:01 ?1527次閱讀
    拉普拉斯<b class='flag-5'>算子</b>的FPGA<b class='flag-5'>實現(xiàn)</b>方法

    LOG算子在FPGA中的實現(xiàn)

    介紹了一種高斯拉普拉斯LOG算子在FPGA中的實現(xiàn)方案!并通過對一幅BMP圖像的處理!論證了在FPGA中實現(xiàn)的LOG算子的圖像增強效果
    發(fā)表于 05-16 17:12 ?50次下載
    LOG<b class='flag-5'>算子</b>在FPGA中的<b class='flag-5'>實現(xiàn)</b>

    BatchNorm是一種旨在通過固定層輸入的分布來改善神經(jīng)網(wǎng)絡訓練的技術

    作者探討了BatchNorm,優(yōu)化和Internal Covariate Shift三者之間的關系。作者在CIFAR-10數(shù)據(jù)集上分別使用和不使用BatchNorm來訓練標準的VGG網(wǎng)絡,如上圖顯示
    的頭像 發(fā)表于 07-03 14:37 ?7323次閱讀
    <b class='flag-5'>BatchNorm</b>是一種旨在通過固定層輸入的分布來改善神經(jīng)網(wǎng)絡訓練的技術

    Laplacian算子的FPGA實現(xiàn)方法

    拉普拉斯算子是一種重要的圖像增強算子,它是一種各向同性濾波器,即濾波器的響應與濾波器作用圖像的突變方向無關,而且實現(xiàn)簡單,被廣泛用于圖像銳化和高頻增強等算法中。在此,提出一種使用QuartusⅡ開發(fā)環(huán)境的Megafunction
    的頭像 發(fā)表于 06-16 17:47 ?3263次閱讀
    Laplacian<b class='flag-5'>算子</b>的FPGA<b class='flag-5'>實現(xiàn)</b>方法

    鍵盤電子實現(xiàn)

    鍵盤電子實現(xiàn)
    發(fā)表于 05-26 15:32 ?1次下載

    開源軟件-OneFlow通用深度學習框架

    ./oschina_soft/oneflow.zip
    發(fā)表于 06-20 09:26 ?2次下載
    開源軟件-<b class='flag-5'>OneFlow</b>通用深度學習框架

    Laplacian算子的硬件實現(xiàn)及結(jié)果

    使用Laplacian算子濾波是將模板與圖像做卷積運算,然后將得到的結(jié)果取絕對值后,再進行防治溢出(灰度值大于255)處理。所以在用硬件實現(xiàn)Laplacian算子時可分成三個步驟:構(gòu)造模板;使用模板對圖像進行卷積運算;對卷積后的
    發(fā)表于 07-21 09:27 ?1084次閱讀

    Sobel算子原理介紹與實現(xiàn)方法

    索貝爾算子(Sobel operator)主要用作邊緣檢測,在技術上,它是一離散性差分算子,用來運算圖像亮度函數(shù)的灰度之近似值。在圖像的任何一點使用此算子,將會產(chǎn)生對應的灰度矢量或是其法矢量Sobel 卷積因子為:
    的頭像 發(fā)表于 07-21 17:27 ?1.3w次閱讀

    flowflops:OneFlow模型的Flops計算

    用于計算 OneFlow 模型的 FLOPs 和 Parameters 的第三方庫。
    的頭像 發(fā)表于 11-16 10:04 ?1194次閱讀

    解析OneFlow Element-Wise算子實現(xiàn)方法

    雖然這種寫法非常簡單明了,但卻存在明顯的性能問題。所以這篇文章將基于OneFlow開源的Element-Wise CUDA算子方案來解釋如何寫一個高性能的Element-Wise CUDA算子
    的頭像 發(fā)表于 12-12 10:54 ?1552次閱讀

    自定義算子開發(fā)

    一個完整的自定義算子應用過程包括注冊算子算子實現(xiàn)、含自定義算子模型轉(zhuǎn)換和運行含自定義op模型四個階段。在大多數(shù)情況下,您的模型應該可以通過使用hb_mapper工具完成轉(zhuǎn)換并順利部署
    的頭像 發(fā)表于 04-07 16:11 ?2794次閱讀
    自定義<b class='flag-5'>算子</b>開發(fā)
    主站蜘蛛池模板: 免费国产成人高清在线看软件| 欧美中文字幕一区二区三区| 色播成人影院| 在线高清无码欧美久章草| 国产成人精品男人免费| 免费A级毛片无码鲁大师| 亚洲另类欧美综合在线| 俄罗斯18xv在线观看| 蜜芽一区二区国产精品| 亚洲欧美一区二区三区导航| 俄罗斯6一12呦女精品| 男人插女人逼逼| 在线播放免费人成毛片视频| 国产欧美一本道无码| 日本护士在线观看| 99爱在线观看精品视频| 久久精品人人做人人爽97| 小货SAO边洗澡边CAO你动漫| 成人小视频在线观看免费| 蜜臀AV99无码精品国产专区| 伊人久久综合热青草| 国产一区二区精品视频| 翁止熄痒禁伦短文合集免费视频| sihu国产精品永久免费| 伦理片飘花免费影院| 一级毛片视频免费| 黑人 尺寸 强行害怕 痛哭| 性色欲情网站IWWW| 国产白丝JK被疯狂输出视频| 日本三区四区免费高清不卡| 99热在线精品视频| 蜜芽无码亚洲资源网站| 91九色视频无限观看免费| 久久久精品国产免费A片胖妇女| 性奴公司 警花| chinesevideos原创麻豆| 国产一级做a爰片久久毛片男| 青青青青草原国产免费| 亚洲爆乳无码精品AAA片蜜桃| 爱看吧孕妇网| 可以看的黄页的网站|