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

    Compute Unified Device Architecture Implementation of Euler/Navier-Stokes Solver on Graphics Processing Unit Desktop Platform for 2-D Compressible Flows

    2016-12-01 03:18:41ZhangJialeChenHongquan

    Zhang Jiale,Chen Hongquan

    College of Aerospace Engineering,Nanjing University of Aeronautics and Astronautics,Nanjing 210016,P.R.China

    Compute Unified Device Architecture Implementation of Euler/Navier-Stokes Solver on Graphics Processing Unit Desktop Platform for 2-D Compressible Flows

    Zhang Jiale,Chen Hongquan*

    College of Aerospace Engineering,Nanjing University of Aeronautics and Astronautics,Nanjing 210016,P.R.China

    Personal desktop platform with teraflops peak performance of thousands of cores is realized at the price of conventional workstations using the programmable graphics processing units(GPUs).A GPU-based parallel Euler/Navier-Stokes solver is developed for 2-D compressible flows by using NVIDIA′s Compute Unified Device Architecture(CUDA)programming model in CUDA Eortran programming language.The techniques of implementation of CUDA kernels,double-layered thread hierarchy and variety memory hierarchy are presented to form the GPU-based algorithm of Euler/Navier-Stokes equations.The resulting parallel solver is validated by a set of typical test flow cases.The numerical results show that dozens of times speedup relative to a serial CPU implementation can be achieved using a single GPU desktop platform,which demonstrates that a GPU desktop can serve as a cost-effective parallel computing platform to accelerate computational fluid dynamics(CED)simulations substantially.

    graphics processing unit(GPU);GPU parallel computing;compute unified device architecture(CUDA)Eortran;finite volume method(EVM);acceleration

    0 Introduction

    Developing computing codes with efficient parallelization is greatly demanded for the real applications occurred in computational fluid dynamics(CED)like simulations of flows over aerodynamic bodies of considering complex three-dimensional configurations.The efficiency of traditional CED solvers based on CPU architecture meets the bottleneck because limited by the clock frequency and data transmission bandwidth of CPUs[1].In recent years,a new streaming processors,often called graphics processing unit(GPU),has become available for compute-intensive parallel tasks[2].The memory bandwidth and floatingpoint performances of modern GPUs are orders of magnitude faster than a standard CPU.And the growing gap in peak performance,measured in floating point operations per second(ELOPS)between GPU and CPU over the last decade,is illustrated by Fig.1[2].Currently,NVIDIA GPUs outperform Intel CPUs on floating point performance and memory bandwidth,both by a factor of roughly ten[2].

    At the beginning,the GPU′s old fixed-function pipeline did not allow complex operations, thus using GPUs for general-purpose computation was a complicated exercise[3].As the formation and development of new efficient programming model,the genera-purpose computation programmers can perform the parallel computation on GPU,and transfer date between CPU and GPU more convenient and effective than before. NVIDIA′s compute unified device architecture(CUDA)[2]is one such model that supports native high-level programming language on its own line of GPUs.The programmers can use CUDA C/C ++,CUDA Eortran or other programming languages in the GPU codes.

    *Corresponding author,E-mail address:hqchenam@nuaa.edu.cn.

    How to cite this article:Zhang Jiale,Chen Hongquan.Compute unified device architecture implementation of Euler/Navier-Stokes Solver on graphics processing unit desktop platform for 2-D compressible flows[J].Trans.Nanjing Univ.Aero. Astro.,2016,33(5):536-545.

    http://dx.doi.org/10.16356/j.1005-1120.2016.05.536

    Fig.1 Eloating-point performance and memory bandwidth for Intel CPU and NVIDIA GPU[2]

    High performance parallel computing with CUDA has already attracted the CED researchers. Brandvik and Pullan[4]completed porting a twoand three-dimensional Euler code for modeling inviscid flow onto GPU in 2008.The resultant solver ran 30 times faster for a two-dimensional case and 15 times faster for a three-dimensional case. At around the same time,Elsen,et al.[5]rewrote part of the Navier-Stokes Stanford University solver(NSSUS)to model hypersonic flow on GPU. The performance of their code ranges from 15 to 40 times speedup compared with the original solver.Molemaker,et al.[6]developed a multi-grid method to solve the pressure Poisson equation. The CUDA implementation of the multi-grid pressure Poisson solver produced a speedup of 55 times relative to a 2.2 MHz AMD Opteron processor[6].In addition to the finite volume method, the discontinuous Galerkin(DG)method has also been implemented on the GPU by Klockner,et al.[7]in 2009.Zhang and Han[8]developed an implicit data-parallel scheme to solve computational fluid dynamics problems on GPU platform in 2010,and obtained 28 times speedup for an implicit calculation of ONERA M6 wing.In 2011, Liu,et al.[9]acceletated the simulation of large sparse linear and nonlinear equations on multiple pieces of GPU.Gu and Xu[10]achieved a real-time liquid payload hydrodynamics simulation method on GPU using smoothed particle hydrodynamics (SPH)method in 2013,and obtained a performance of 12 times speedup of a aerial CPU implementation.All the past and recent works in writing CED solvers on GPU have shown encouraging results.

    In the paper,we focus on the developing of personal desktop platform with the programmable graphics processing units of thousands of cores. The techniques of implementation of CUDA kernels,double-layered thread hierarchy and variety memory hierarchy are discussed to develop a GPU-based parallel Euler/Navier-Stokes solver in CUDA Eortran programming language.A set of typical test flow cases are selected for both validation and performance analysis.Dozens of times speedup(about 38 times upmost for present test cases)relative to a serial CPU implementation can be achieved using GPU desktop platform developed for the present solver.

    1 Governing Equations and Numerical Approach

    1.1 Governing equations

    The 2-D Navier-Stokes equations governing compressible fluid flows can be expressed in conservative form[11]as

    where W is the vector of conservative variables.Fcand Gcare the convective flux terms;Fvand Gvthe viscous flux terms.They are defined aswhereρ,p,T,E and H denote the density,pressure,temperature,total energy per unit mass, and total enthalpy per unit mass,respectively.u and v are the cartesian components of the velocity vector.Eor a perfect gas,these quantities satisfy

    In these equations,μand k denote the dynamic viscosity coefficient and thermal conductivity coefficient,respectively.Eor perfect gas,the dynamic viscosity coefficientμcan be calculated by the Sutherland formula[11].The thermal conductivity coefficient k has relationship toμas

    an d

    The components of the viscous stress tensorτijis defined as

    Eor perfect gas,the ratio of specific heats of fluid γ=1.4,and the Prandtl number Pr=0.72.

    1.2 Numerical approach

    Once a mathematical model is defined,we need a method for approximating the differential equations by obtaining a system of equations at a set of discrete elements in space and time.In our application of fluid dynamics,the finite volume method(EVM)in cell-centered scheme is adopted.By divided the flow field into a set of discrete control volumes,the integral form of the equation can be applied over each control volume.Eor control volume i,the discretization equations is as follows

    where the convective flux and viscous flux are defined as

    where n andΔSijdenote the edge number and the area of edge j of the control volume i,respectively.

    The convective flux term is evaluated by central scheme with artificial dissipation,as

    where the construction of the artificial dissipative term is given by

    whereε(2)andε(4)are adaptive coefficients.λ=is spectral radius of the Jacobian matrix, in which U=unx+vnyis the normal velocity of the edge,andthe speed of sound.The detailed description of the construction of the dissipative term can be found in Ref.[11].

    The viscous flux term is evaluated by central scheme,as

    In order to obtain the steady solution,an explicit four-stage Runge-Kutta scheme is adopted for time integration

    where the superscript n and n+1 denote the current and the next time level.The coefficientαmcan be found in Ref.[11].To accelerate the convergence,Δt is taken as the local time step[11].

    In the case of inviscid flow governed by the Euler equations,the fluid slips over the solid wall and there is no flow normal to the surface,that is V·n=0.And in the case of viscous flow governed by the Navier-Stokes equations,the relative velocity between the surface and the fluid directly at the surface is zero,that is u=v=0.In the far field,the Riemann boundary condition is adopted,and the details can be found in Ref.[11].

    2 Programming Model and GPU Implementation

    NVIDIA′s CUDA[2]is a parallel computing platform and programming model that leveragesthe powerful compute engine of their GPUs.CUDA provides the developers with a software environment,an extension to C,Eortran and some other programming languages,to launch and manage massively parallel computations on GPUs.The reader can refer to the CUDA programming guide for more details[2].In the section,we summarize the GPU hardware architecture and the programming model with CUDA Eortran.Hereinafter,we use the term″host″to refer to CPU and the term″device″to refer to GPU.

    2.1 GPU architecture

    GPU is specialized for compute-intensive, highly parallel computation,and designed such that more transistors are devoted to data processing rather than data caching and flow control in CPU,as schematically illustrated by Fig.2.

    Fig.2 Different architecture of CPU and GPU

    GPU is a set of stream multiprocessors, which is an extension of single instruction multiple data(SIMD)paradigm architecture.This design architecture,as the name suggests,makes CUDA optimal for performing a single instruction in parallel on different sets of data on NVIDIA GPUs.A full Kepler GK110(the latest generation GPU architecture)is composed of fifteen streaming multiprocessors(SM)and six 64-bit memory controllers.As shown in Fig.3(a),each SM contains a collection of processing cores(192 in Kepler architecture),shared memory,Register Eile,L1 cache memory,read-only data cache,L2 cache memory,Warp scheduler,and other processing or control units.Each green border square represents a CUDA core which is then mapped to a thread by the runtime system.Closer to the core,the local registers allow fast ALU operations.The shared memory which is also closer to the processors will be used to store data shared by all the cores of a SM.The global memory is used to store the main computing data. And different kinds of cache,including L1 cache, read-only data cache and L2 cache,are used to provide efficient,high speed data access of global memory by processors.Depending on the GPU architecture,a wide variety of combinations of cache memory,global memory,number of streaming multiprocessors,and thread count are available.CUDA-enabled GPUs with the Kepler architecture are powerful processors that contain thousands of cores which are able of converting a simple computer into a high-performance computer.

    2.2 Programming model

    In the CUDA programming model,the parallel code executed on GPU is called the kernel. Before executing a kernel,the processing data must be transferred to memory on GPU.Next, CPU initiates kernel execution on GPU.After the kernel is completed executed,CPU retrieves the processed data from GPU.This process is illustrated in Fig.4.

    The kernel is launched from the host side (CPU),and it is organized to a thread grid,as illustrated in Fig.3(b).Each grid contains a set of blocks,and each block contains a set of threads. The number of threads per block and the grid size (number of blocks)need to be defined before launching the kernel.All the threads from a specific block have to execute on the same SMX, they can access to the same shared memory and can be synchronized.On the other hand,threads from different blocks cannot synchronize and can exchange data only through the global memory.

    CUDA Eortran API is an extension to the Eortran programming language.It provides functions and keywords to manage the computations on GPU.The reader can refer to Ref.[12]for more details.In CUDA Eortran API,the qualifiers attributes(host),attributes(global)and at-tributes(device)specify whether CPU or GPU should execute and call the qualified subroutine or function.The qualifiers device,shared and constant are introduced to define the type of memory a variable should use,and the function CUDAMalloc(),CUDAMemcpy()and CUDAFree() are used to allocate memory on GPU,copy data between CPU memory and device memory,and free memory on GPU,respectively.

    Fig.3 Kepler GPU hardware architecture and CUDA thread organization

    Fig.4 CUDA process flow

    In addition,the kernel is launched by specifying the size of the grid(number of blocks)and the size of the block(number of threads per block)using the flowing prototype:kernel_function<<<grid Size,blockSize>>>().The function syncthreads()is used inside a kernel to synchronize all the threads of a same block,and the function CUDADeviceSynchronize()is used inside a host function to synchronize the host (CPU)and device(GPU).

    2.3 CUDA implementation

    Explicit Runge-Kutta scheme is selected for present CUDA implementation.Since the algorithm is explicit in both space and time,threads can run independently of each other for a majority of the time-step.This ensures the thread processors are fully used and none is idle.There are only a few thread synchronization must occurred, such as after updating the flow variables in grid elements,or when judging the convergence.In Additional,the algorithm is designed in unstructured model,that suitable for both structured and unstructured grid.Because of that,the grid cells are arranged into one dimensional form,and the size of the thread block and thread grid are both in one-dimensional form,as shown in Fig.5. Each thread corresponds to a grid element,and takes charge of the computation of the specified element.Every n(thread number per block) counts of threads are combined into a thread block,and the size of thread grid is defined to m=(N+n-1)/n.Thus,all the thread counts is m*n,a little larger than the element count n N(number of elements),and the threads can cover all the element that needing be solved.

    Fig.5 1-D permutation and correspondence between grid elements and GPU threads

    Listing 1 shows the simplified host side code of the time iteration procedure.In the code snippet,two nested loops are composed.The outer loop is used for advancing the solution in time, and the inner loop is used for four-stage Runge-Kutta iteration.As shown in Listing 1,almost all of the computing tasks are composed into several kernel functions that could run on the device (GPU),and the host(CPU)is just used to organize and control the computational flow.The memory allocation of the arrays on the device is only one before starting the time stepping.

    Listing 2 shows the device code that compute the primitive flow variables by the conserved flow variables.In the beginning,the qualifier″attributes(global)″is used to define kernel_primvar as a kernel subroutine.The qualifier″device″in variable define line is used to express the val_consvar and val_primvar are GPU arrays,and the qualifier″value″is used to express the constant val_ nElem is a variable that define in host and used in kernel.The code in line 10 is used to obtain the thread ID,and the″if″structure in line 12 is used to limit the thread ID in case computing beyond the upper bound.

    Listing 1 Simplified host side code of time iteration procedure

    Listing 2 Simplified device side code for computing primitive variables from conservative variables

    3 Results and Discussions

    Based on the above method and the hardware architecture,we develop both CPU serial code and GPU parallel code for solving the 2-D Euler/Navier-Stokes equations.All the codes are 32 bit floating-point precision since GPU we used operates in single precision.In all our cases,CPU used is a single core of an Intel Core i5-3450(3.1 GHz,6 MB L3 cache,can turbo to 3.5 GHz)and GPU used is an NVIDIA GTX Titan(2688 scalar processor cores at 837 MHz,6 GB Memory). Both CPU and GPU can represent their current advanced level of computing.

    Eirstly,the GPU solver is tested on a transonic inviscid flow around the RAE2822 airfoil, and a subsonic viscous flow around the NACA0012 airfoil.The results of these two-dimensional cases are given below and compared with experimental data and reference results.Then the resulting GPU solver is tested with a set of cases for the steady-state flow around the NACA0012 airfoil on a series of structure grids.The results and the performance are given and analyzed below.

    3.1 Validation of GPU implementation of Euler/ Navier-Stokes solver

    Eirstly,the GPU solver is tested with a case of a transonic inviscid flow around the RAE2822 airfoil.This case is run with four stages Runge-Kutta iteration,with a CEL number of 5 and implicit residual smoothing method.We assume the free stream flow conditions with a Mach number of 0.729 and a 2.31°angle of attack in this case. Fig.6(a)shows the contours of Mach number near the airfoil,and Fig.6(b)presents the comparison of pressure coefficient data among the experiment data,the reference data and our result. The reference data NPARC and WIND are provided by NASA.Our result produces reasonably accurate results on the leading edge of the airfoil and a very crisp,clean shock.The position of shock wave on the upper surface of the airfoil of our result is a little later than the reference results,but closer to the experimental data.The pressure coefficient on the trailing edge of the airfoil of our result is a little lower than the experimental data and viscous reference results,but closer to the inviscid results from Eluent and Ref.[13].These behaviors are fairly typical of transonic Euler CED solutions.

    Fig.6 Euler results for RAE2822 airfoil at Ma∞=0. 729 and 2.31°of Ao A

    Then,the GPU solver is tested with a case of a subsonic viscous flow around the NACA0012 airfoil.The case is running with same temporal discretization scheme and same CEL number of the front case.We assume free stream flow conditions with a Mach number of 0.5,a zero angle of attack and Reynolds number at 5 000 in this case.Fig.7(a)shows the contours of Mach number near the airfoil,and Fig.7(b)presents the comparison of pressure coefficient data between the reference data[14-17]and our result.As we can see in Figs.7(a,b),our GPU solver produces reasonably accurate results in this case of a viscous flow.

    The results of our GPU solver for both inviscid and viscous flows seem fairly reasonable with experimental data and refer results.

    Fig.7 Laminar results for NACA0012 airfoil at Ma∞= 0.5,zero of Ao A and Rec=5 000

    3.2 Performance results and evaluation

    The following computing hardware is utilized in the paper.A dual-CPU/single-GPU platform is built with an Intel Core i5-3450 Quad 3.1 GHz CPU,16 GB of memory and a NVIDIA GTX Titan board.The Titan board provides 2 688 streaming processor cores and 6 GB of global device memory.And the Titan GPU used in the paper can deliver a theoretical peak performance of 4 TELOPS(Trillion floating-point operations per second).

    To assess the performance of the Euler/ Navier-Stokes solver,we use both the inviscid and viscous conditions to solve for the steadystate flows around the NACA0012 airfoil on varying sizes of structure grids.The grid sizes selected for comparison are 192×64,384×128, 768×256,and 1 536×512.The performance is measured by dividing the total number of grid elements by the time for a single iteration of the solver,yielding the throughput in millions of cellstep per second.

    Using a single CPU core,the performance of our CED code is approximately 0.6 million cellsteps per second for Euler solver and 0.45 million cell-steps per second for NS solver,shown in Fig.8 and the third column of Table 1.We can find that the CPU performance is almost no growth with the increase of the gird size.That is because the smallest size of grid have already made the full use of CPU.On the other hand,the performance of our GPU implementation is more efficient,as presented in Fig.8 and the fourth column in Table 1.And the computational speedup is shown in the fifth column.We can find that present GPU performs much faster than the implementation of CPU,at a radio of tens.

    Fig.8 Performance of single CPU and GPU on NACA0012 case with varying grid size

    Table 1 Performance of CPU and GPU with different size of grids

    As shown in Table 1,the speedup of GPU isgrowing with the increase of the grid size.On the smallest size of grid,the level of parallelism does not fully utilize the massively parallel GPU architecture,and performance is only 11.9 for inviscid flow case and 10.4 for laminar flow case above that of CPU(Table 1).As the mesh size increases,GPU becomes more efficient and perform much faster.The speedup grows up to 25.0, 38.1 and 38.8 times in the inviscid flow cases and 19.9,26.9 and 29.9 times in viscous flow cases on the increasing grid size of 384×128,768×256 and 1 536×512,respectively.The speedup numbers are impressive for large problem size,because the arithmetic intensity on the GPU increases with problem size.

    4 Conclusions

    The implementation of Euler/Navier-Stokes equations for 2-D compressible flows on personal desktop platform with a GPU is presented.With CUDA Eortran programming language, NVIDIA′s CUDA programming model is used to implement the discretized form of the governing equations.In the numerical solution of 2-D compressible flows,the GPU code has perform dozens of speedup compared with the serial CPU code(about 38 times upmost for present test cases).In addition,we observe that the performance of speedup becomes better with increasing the size of corresponding computational problems,which suggests that the GPU code can be well suited to the complex engineering problems.

    Acknowledgements

    This work was supported by the National Natural Science Eoundation of China(No.11172134);the Eunding of Jiangsu Innovation Program for Graduate Education(No. CXLX13_132).

    [1] PATTERSON D A,HENNESSY J L.Computer architecture:A quantitative approach[M].4th ed.[S. l.]:Morgan Kaufmann Publishers Inc,2007.

    [2] NVIDIA.NVIDIA CUDA C programming guide (version 7.0)[EB/OL].(2015-03-05)[2015-09-01].http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html.

    [3] OWENS J D,HOUSTON M,LUEBKE D,et al. GPU computing[J].Proceedings of the IEEE,2008, 96(5):879-899.

    [4] BRANDVIK T,PULLAN G.Acceleration of a 3D Euler solver using commodity graphics hardware: AIAA 2008-607[R].2008.

    [5] ELSEN E,LEGRESLEY P,DARVE E.Large calculation of the flow over a hypersonic vehicle using a GPU[J].Journal of Computational Physics,2008, 227(24):10148-10161.

    [6] MOLEMARKER J,COHEN J M,PATEL S,et al. Low viscosity flow simulations for animation[C]// Eurographics/ACM SIGGRAPH Symposium on Computer Animation.Dublin,Ireland:SCA,2008: 115-116.

    [7] KL?CKNER A,WARBURTON T,BRIDGE J,et al.Nodal discontinuous Galerkin methods on graphics processors[J].Journal of Computational Physics, 2009,228(21):7863-7882.

    [8] ZHANG B,HAN J.Parallel computing methods for CED using a GPU and implicit scheme[J].Acta Aeronautica et Astronautica Sinica,2010,21(2): 249-256.

    [9] LIU S,ZHONG C,CHEN X.Solvers for systems of large sparse linear and nonlinear equations based on multi-GPUs[J].Transactions of Nanjing University of Aeronautics and Astronautics,2011,28(3): 300-308.

    [10]GU A,XU J.Liquid payload hydrodynamics simulation method for real-time aircraft simulation system [J].Journal of Nanjing University of Aeronautics and Astronautics,2013,45(4):491-496.(in Chinese)

    [11]BLAZEK J.Computational fluid dynamics:Principles and applications[M].2nd ed.[S.l.]:Elsevier, 2001.

    [12]The Portland Group.CUDA Eortran programming guide and reference,Release 2014[EB/OL].(2014-01-01)[2015-09-01].http://www.pgroup.com/ doc/pgicudaforug.pdf.

    [13]AHMED E,AHMED E,AEEAN M,et al.Numerical solution of 2D Euler equations for the transonic flow past over NACA0012 and RAE2822 airfoils using high order accurate Runge-Kutta discontinuous Galerkin method[C]//Proceeding of 2014 11th International Bhurban Conference on Applied Sciences &Technology(IBCAST).[S.l.]:IBCAST,2014: 218-213.

    [14]LARA P R M,MORGAN K.A review and comparative study of upwind based schemes for compressible flow computation.Part III:Multidimensional extension on unstructured grids[J].Archives of Computational Methods in Engineering,2002,9(3):207-256.

    [15]CARAENI D,EUCHS L.Compact third-order multidimensional upwind scheme for Navier-Stokes simulations[J].Theoretical and Computational Eluid Dynamics,2002,15(6):373-401.

    [16]CHASSAING J C,KHELLADI S,NOGUEIRA X. Accuracy assessment of a high-order moving least squares finite volume method for compressible flows [J].Computers&Eluid,2013,71(2):41-53.

    [17]LIU Xueqiang.The research of N-S equations′solution using hybrid grids and multi-grid methods and its applications[D].Nanjing:Nanjing University of Aeronautics and Astronautics,2001.(in Chinese)

    Ms.Zhang Jiale received B.S.degree in flight vehicle design and engineering from Nanjing University of Aeronautics and Astronautics(NUAA)in 2009,and M.S.degree in aerodynamics also from NUAA in 2012,respectively. Erom 2012 to present,he is a Ph.D.candidate in College of Aerospace Engineering,Nanjing University of Aeronautics and Astronautics.His research is focused on computational fluid dynamics and high performance computing.

    Prof.Chen Hongquan received B.S.and Ph.D.degrees in aerodynamics from Nanjing University of Aeronautics and Astronautics(NUAA)in 1984 and 1990,respectively. Erom 1993 to 1994,he was a Postdoctoral Research Eellow in University of Paris VI and Dassault Aircraft Company, Erance.Erom 1990 to present,he has been a full professor in College of Aerospace Engineering,Nanjing University of Aeronautics and Astronautics.His research is focused on aerodynamics,which includes computational fluid dynamics,computational electromagnetic,genetic algorithm and multidiscipline design optimization(MDO).

    (Executive Editor:Xu Chengting)

    V211.3 Document code:A Article ID:1005-1120(2016)05-0536-10

    (Received 16 December 2015;revised 30 March 2016;accepted 28 April 2016)

    国产女主播在线喷水免费视频网站| 亚洲最大成人手机在线| 边亲边吃奶的免费视频| 欧美3d第一页| 一本色道久久久久久精品综合| www.av在线官网国产| 国产男人的电影天堂91| 久久久久久久久久久免费av| 男人狂女人下面高潮的视频| 成人国产av品久久久| 中国三级夫妇交换| 久久精品国产亚洲av涩爱| 啦啦啦中文免费视频观看日本| 亚洲成色77777| 如何舔出高潮| 激情五月婷婷亚洲| 精品人妻偷拍中文字幕| 秋霞伦理黄片| 老司机影院成人| 99久久中文字幕三级久久日本| 麻豆久久精品国产亚洲av| 狠狠精品人妻久久久久久综合| a级毛片免费高清观看在线播放| 日韩欧美精品v在线| 91精品一卡2卡3卡4卡| 秋霞在线观看毛片| 国产成人a∨麻豆精品| 看非洲黑人一级黄片| 女人久久www免费人成看片| 久久精品国产亚洲网站| 搡女人真爽免费视频火全软件| 赤兔流量卡办理| 在线免费观看不下载黄p国产| 国产精品精品国产色婷婷| 一区二区av电影网| 肉色欧美久久久久久久蜜桃 | 久久久亚洲精品成人影院| 九草在线视频观看| 日韩欧美一区视频在线观看 | 欧美97在线视频| 好男人在线观看高清免费视频| 国产精品国产三级国产av玫瑰| 我的老师免费观看完整版| 人妻一区二区av| 欧美日韩亚洲高清精品| 嫩草影院入口| 一级片'在线观看视频| 美女主播在线视频| av网站免费在线观看视频| 美女被艹到高潮喷水动态| 亚洲国产精品成人久久小说| 国产精品久久久久久av不卡| 麻豆国产97在线/欧美| 日韩精品有码人妻一区| 一级毛片黄色毛片免费观看视频| 80岁老熟妇乱子伦牲交| 亚洲精品乱码久久久v下载方式| 乱系列少妇在线播放| 卡戴珊不雅视频在线播放| 免费av毛片视频| 亚洲熟女精品中文字幕| 国产精品一区二区三区四区免费观看| 夜夜爽夜夜爽视频| 日本三级黄在线观看| 久久久久久久国产电影| 97热精品久久久久久| 久久影院123| a级毛片免费高清观看在线播放| 97超视频在线观看视频| 亚洲欧美成人综合另类久久久| 日韩一本色道免费dvd| 两个人的视频大全免费| 国产高清有码在线观看视频| 亚洲国产精品999| 涩涩av久久男人的天堂| 美女高潮的动态| 乱系列少妇在线播放| 热99国产精品久久久久久7| 婷婷色综合www| 午夜亚洲福利在线播放| 国产精品人妻久久久影院| 一级毛片 在线播放| av免费观看日本| 成人毛片60女人毛片免费| .国产精品久久| 搞女人的毛片| 国产精品99久久久久久久久| 日韩在线高清观看一区二区三区| 青春草视频在线免费观看| 国产91av在线免费观看| 天堂俺去俺来也www色官网| 亚洲精品日本国产第一区| 丰满少妇做爰视频| 好男人视频免费观看在线| 日日摸夜夜添夜夜爱| 欧美一区二区亚洲| 搡老乐熟女国产| 国国产精品蜜臀av免费| 免费不卡的大黄色大毛片视频在线观看| 日本wwww免费看| 亚洲天堂av无毛| 亚洲精品日韩av片在线观看| 亚洲一区二区三区欧美精品 | 亚洲,一卡二卡三卡| 亚洲,一卡二卡三卡| 男男h啪啪无遮挡| 一区二区三区四区激情视频| 在线观看美女被高潮喷水网站| 亚洲av福利一区| 日本wwww免费看| 99久久精品热视频| 日韩精品有码人妻一区| 少妇丰满av| 深夜a级毛片| freevideosex欧美| 久久久午夜欧美精品| 天天躁夜夜躁狠狠久久av| 男的添女的下面高潮视频| 成人毛片60女人毛片免费| 亚洲av中文av极速乱| 男女啪啪激烈高潮av片| 人妻夜夜爽99麻豆av| 精品人妻偷拍中文字幕| 日日啪夜夜撸| av在线观看视频网站免费| 国产精品福利在线免费观看| 欧美日韩在线观看h| 免费高清在线观看视频在线观看| 蜜臀久久99精品久久宅男| 欧美日韩在线观看h| 少妇人妻久久综合中文| 日韩三级伦理在线观看| 国产国拍精品亚洲av在线观看| 欧美国产精品一级二级三级 | 蜜桃久久精品国产亚洲av| 高清在线视频一区二区三区| tube8黄色片| 成人一区二区视频在线观看| 国产高清三级在线| 久久精品国产鲁丝片午夜精品| 直男gayav资源| 国产亚洲av片在线观看秒播厂| 在线观看一区二区三区激情| 黄色欧美视频在线观看| 少妇的逼水好多| 久久精品国产亚洲av涩爱| 午夜福利在线观看免费完整高清在| 国产女主播在线喷水免费视频网站| 国产成人精品久久久久久| 亚洲成色77777| 夫妻性生交免费视频一级片| 免费观看av网站的网址| 交换朋友夫妻互换小说| 国产黄频视频在线观看| 26uuu在线亚洲综合色| 亚洲一级一片aⅴ在线观看| 七月丁香在线播放| 99久久中文字幕三级久久日本| 精品午夜福利在线看| xxx大片免费视频| 久久久精品94久久精品| 久久综合国产亚洲精品| 综合色av麻豆| 一二三四中文在线观看免费高清| 国语对白做爰xxxⅹ性视频网站| 精品一区二区三区视频在线| 97精品久久久久久久久久精品| 国产欧美日韩一区二区三区在线 | 日韩 亚洲 欧美在线| 人妻少妇偷人精品九色| 亚洲欧美日韩另类电影网站 | 亚洲av男天堂| 九九久久精品国产亚洲av麻豆| 少妇丰满av| 3wmmmm亚洲av在线观看| 最近最新中文字幕免费大全7| 国产精品爽爽va在线观看网站| 久久影院123| 噜噜噜噜噜久久久久久91| 亚洲欧洲日产国产| 国产欧美另类精品又又久久亚洲欧美| av又黄又爽大尺度在线免费看| 国产男女内射视频| 九色成人免费人妻av| 午夜免费鲁丝| 日韩亚洲欧美综合| 男人狂女人下面高潮的视频| 直男gayav资源| 免费看不卡的av| 久久女婷五月综合色啪小说 | 国产 一区精品| 丰满人妻一区二区三区视频av| av在线天堂中文字幕| 国产av国产精品国产| 在线 av 中文字幕| 九九爱精品视频在线观看| 久久久a久久爽久久v久久| 亚洲国产精品999| 日本-黄色视频高清免费观看| 国产精品无大码| 日韩制服骚丝袜av| a级毛片免费高清观看在线播放| 69av精品久久久久久| 久久韩国三级中文字幕| 国产人妻一区二区三区在| 精品99又大又爽又粗少妇毛片| tube8黄色片| 国产伦在线观看视频一区| 免费观看av网站的网址| 久久国产乱子免费精品| 亚洲av福利一区| 一级毛片我不卡| 亚洲最大成人手机在线| 国产乱来视频区| 三级国产精品片| 内地一区二区视频在线| 国产久久久一区二区三区| 99热这里只有是精品50| 白带黄色成豆腐渣| 男人舔奶头视频| 国产精品成人在线| 成人黄色视频免费在线看| xxx大片免费视频| av国产精品久久久久影院| 简卡轻食公司| av黄色大香蕉| 国产毛片在线视频| 国产免费一级a男人的天堂| eeuss影院久久| 97在线视频观看| 久久午夜福利片| 建设人人有责人人尽责人人享有的 | 乱码一卡2卡4卡精品| 欧美+日韩+精品| 两个人的视频大全免费| 七月丁香在线播放| 成人漫画全彩无遮挡| 最近手机中文字幕大全| 欧美极品一区二区三区四区| 国产精品99久久久久久久久| 亚洲一区二区三区欧美精品 | 国产91av在线免费观看| 久久久精品免费免费高清| 99视频精品全部免费 在线| 欧美bdsm另类| 日本猛色少妇xxxxx猛交久久| 欧美日韩国产mv在线观看视频 | 极品教师在线视频| 午夜免费男女啪啪视频观看| 交换朋友夫妻互换小说| 在线观看三级黄色| av卡一久久| 一本久久精品| 狂野欧美激情性xxxx在线观看| 欧美xxⅹ黑人| 乱系列少妇在线播放| 赤兔流量卡办理| 成人特级av手机在线观看| 18禁裸乳无遮挡动漫免费视频 | 中文字幕人妻熟人妻熟丝袜美| 日本熟妇午夜| av免费观看日本| 国产综合精华液| 欧美xxxx性猛交bbbb| 国产日韩欧美在线精品| 一级毛片黄色毛片免费观看视频| 国产一区二区三区av在线| 男女国产视频网站| 大香蕉97超碰在线| 狠狠精品人妻久久久久久综合| 国产成人freesex在线| 熟女人妻精品中文字幕| 国产探花在线观看一区二区| 欧美xxⅹ黑人| 高清日韩中文字幕在线| 日本爱情动作片www.在线观看| 简卡轻食公司| 国产日韩欧美在线精品| 在线天堂最新版资源| 国产成人免费观看mmmm| 欧美成人a在线观看| 蜜桃亚洲精品一区二区三区| 午夜免费观看性视频| 婷婷色麻豆天堂久久| 国产精品人妻久久久影院| av卡一久久| 亚洲精品一区蜜桃| 精品酒店卫生间| 国内揄拍国产精品人妻在线| 青春草亚洲视频在线观看| 性色avwww在线观看| 国产大屁股一区二区在线视频| 亚洲av二区三区四区| 亚洲精品日韩av片在线观看| 日本av手机在线免费观看| 欧美3d第一页| 伦精品一区二区三区| 亚洲av国产av综合av卡| 97在线视频观看| 乱系列少妇在线播放| 欧美一级a爱片免费观看看| 老司机影院毛片| 久久久成人免费电影| 六月丁香七月| 亚洲国产成人一精品久久久| 看免费成人av毛片| 少妇被粗大猛烈的视频| 亚洲,一卡二卡三卡| 国产精品伦人一区二区| 亚洲在久久综合| 中文欧美无线码| 久久综合国产亚洲精品| 女人被狂操c到高潮| 看免费成人av毛片| 亚洲美女视频黄频| 成人无遮挡网站| 在线观看人妻少妇| 美女脱内裤让男人舔精品视频| 亚洲自偷自拍三级| 国产老妇伦熟女老妇高清| 波多野结衣巨乳人妻| 国产极品天堂在线| 亚洲自偷自拍三级| 国产亚洲精品久久久com| 99热这里只有是精品50| 熟女电影av网| 中国三级夫妇交换| 99久国产av精品国产电影| 欧美一级a爱片免费观看看| 男女边摸边吃奶| 美女xxoo啪啪120秒动态图| 干丝袜人妻中文字幕| 熟女av电影| 日韩一本色道免费dvd| 久久综合国产亚洲精品| 日韩一区二区三区影片| av线在线观看网站| av.在线天堂| 国产欧美另类精品又又久久亚洲欧美| 亚洲最大成人手机在线| 日韩一区二区视频免费看| 日韩一本色道免费dvd| 我要看日韩黄色一级片| 女人被狂操c到高潮| 少妇高潮的动态图| 一区二区av电影网| 日本欧美国产在线视频| 久久久精品免费免费高清| 欧美亚洲 丝袜 人妻 在线| 免费观看性生交大片5| 精品99又大又爽又粗少妇毛片| av福利片在线观看| 国产亚洲91精品色在线| 亚洲成人久久爱视频| 亚洲精华国产精华液的使用体验| 97人妻精品一区二区三区麻豆| 精品久久久精品久久久| 亚洲人成网站高清观看| 成人毛片60女人毛片免费| 久久久精品94久久精品| 韩国高清视频一区二区三区| av在线亚洲专区| eeuss影院久久| 97热精品久久久久久| 男女那种视频在线观看| 高清午夜精品一区二区三区| 国产成人免费观看mmmm| 国产黄色视频一区二区在线观看| 十八禁网站网址无遮挡 | 日本黄色片子视频| 亚洲国产日韩一区二区| 一级毛片 在线播放| 简卡轻食公司| 伦精品一区二区三区| 久久99热这里只有精品18| 国产片特级美女逼逼视频| 午夜免费男女啪啪视频观看| 日韩中字成人| 国产免费一区二区三区四区乱码| 观看免费一级毛片| 日韩一本色道免费dvd| 校园人妻丝袜中文字幕| 久久精品国产鲁丝片午夜精品| 三级男女做爰猛烈吃奶摸视频| 亚洲内射少妇av| 亚洲精品456在线播放app| 国产精品一区二区三区四区免费观看| 亚洲av男天堂| 99热这里只有是精品在线观看| 亚洲成色77777| 国产国拍精品亚洲av在线观看| 麻豆精品久久久久久蜜桃| 一个人看的www免费观看视频| 国产男女内射视频| 晚上一个人看的免费电影| 亚洲av日韩在线播放| 午夜亚洲福利在线播放| 亚洲精品乱码久久久v下载方式| 熟女电影av网| 亚洲最大成人av| 亚洲精品一区蜜桃| 日韩av不卡免费在线播放| 日本爱情动作片www.在线观看| 中文在线观看免费www的网站| 久久精品人妻少妇| 久久97久久精品| 精品一区二区三卡| 日本熟妇午夜| 少妇被粗大猛烈的视频| 久久久久网色| 欧美人与善性xxx| 免费不卡的大黄色大毛片视频在线观看| 少妇熟女欧美另类| 三级经典国产精品| 亚州av有码| a级毛色黄片| 高清午夜精品一区二区三区| 舔av片在线| 中文字幕人妻熟人妻熟丝袜美| 亚洲精品日韩在线中文字幕| 少妇猛男粗大的猛烈进出视频 | 国产精品久久久久久精品电影| 干丝袜人妻中文字幕| 免费大片黄手机在线观看| 亚洲综合色惰| 人妻系列 视频| 国产老妇女一区| 中国国产av一级| 99久久人妻综合| 日韩人妻高清精品专区| 国产成人免费无遮挡视频| 免费人成在线观看视频色| 在线观看一区二区三区| 黄色视频在线播放观看不卡| 国产毛片a区久久久久| 久久久久久久午夜电影| 国产乱人偷精品视频| 亚洲精品视频女| 亚洲婷婷狠狠爱综合网| 欧美高清成人免费视频www| 99久久精品热视频| 青春草国产在线视频| 免费观看av网站的网址| av国产精品久久久久影院| 国产成人a区在线观看| 免费看a级黄色片| 毛片女人毛片| 国产毛片a区久久久久| 2018国产大陆天天弄谢| 熟妇人妻不卡中文字幕| 汤姆久久久久久久影院中文字幕| 97热精品久久久久久| 国产精品不卡视频一区二区| 国产精品一及| 99热6这里只有精品| 五月开心婷婷网| 男女国产视频网站| 国产亚洲av片在线观看秒播厂| 欧美日韩在线观看h| 2021少妇久久久久久久久久久| 日本熟妇午夜| 男女国产视频网站| 亚洲av在线观看美女高潮| 久久99蜜桃精品久久| 18禁裸乳无遮挡免费网站照片| 看黄色毛片网站| 国产亚洲5aaaaa淫片| 婷婷色综合大香蕉| av国产久精品久网站免费入址| 波多野结衣巨乳人妻| 亚洲国产av新网站| 97精品久久久久久久久久精品| 久久女婷五月综合色啪小说 | 一区二区三区乱码不卡18| 十八禁网站网址无遮挡 | 中文资源天堂在线| 国产精品人妻久久久影院| 亚洲精品中文字幕在线视频 | 97超视频在线观看视频| 美女高潮的动态| 久久鲁丝午夜福利片| 噜噜噜噜噜久久久久久91| 青春草视频在线免费观看| 精华霜和精华液先用哪个| 午夜爱爱视频在线播放| 岛国毛片在线播放| 男女国产视频网站| 一级爰片在线观看| 国产一区二区在线观看日韩| 欧美xxxx黑人xx丫x性爽| 亚洲在久久综合| 久久久精品94久久精品| 精品亚洲乱码少妇综合久久| 91aial.com中文字幕在线观看| 国产 精品1| 波野结衣二区三区在线| 国产亚洲最大av| 久久久a久久爽久久v久久| 街头女战士在线观看网站| 国产精品嫩草影院av在线观看| 精品酒店卫生间| 国国产精品蜜臀av免费| 国产高清三级在线| 日韩中字成人| 99热6这里只有精品| 九九在线视频观看精品| 天美传媒精品一区二区| 国产色爽女视频免费观看| 99热国产这里只有精品6| 99热全是精品| 亚洲精品aⅴ在线观看| videossex国产| 另类亚洲欧美激情| 亚洲精品色激情综合| 乱码一卡2卡4卡精品| 日本wwww免费看| 久久久久久久大尺度免费视频| 少妇人妻一区二区三区视频| 久久久久精品性色| 夜夜看夜夜爽夜夜摸| 好男人视频免费观看在线| 搞女人的毛片| 亚洲最大成人av| 免费少妇av软件| 亚洲av成人精品一二三区| 国产精品一二三区在线看| 国产爱豆传媒在线观看| 男人爽女人下面视频在线观看| 全区人妻精品视频| 欧美zozozo另类| 国产精品99久久久久久久久| 久久久精品94久久精品| 观看免费一级毛片| 精品久久久久久久久亚洲| 成人鲁丝片一二三区免费| 国产爱豆传媒在线观看| 一级毛片我不卡| 成人免费观看视频高清| 中国国产av一级| 18+在线观看网站| 久久精品综合一区二区三区| 国产 一区 欧美 日韩| 久久久久久久久大av| 亚洲精品乱码久久久v下载方式| 2022亚洲国产成人精品| 免费黄色在线免费观看| 最近中文字幕高清免费大全6| 国产 一区 欧美 日韩| 亚洲av男天堂| 国产久久久一区二区三区| 少妇人妻一区二区三区视频| 国产亚洲av嫩草精品影院| 夜夜看夜夜爽夜夜摸| videossex国产| 大陆偷拍与自拍| av国产久精品久网站免费入址| 人体艺术视频欧美日本| 欧美xxxx黑人xx丫x性爽| 波野结衣二区三区在线| 久久97久久精品| 老师上课跳d突然被开到最大视频| 国产高潮美女av| 国产精品三级大全| 联通29元200g的流量卡| 欧美亚洲 丝袜 人妻 在线| 色吧在线观看| 亚洲精品一区蜜桃| 国产伦精品一区二区三区视频9| 久久久久久久国产电影| 日产精品乱码卡一卡2卡三| 超碰97精品在线观看| freevideosex欧美| 亚洲精品第二区| 成人毛片a级毛片在线播放| 91在线精品国自产拍蜜月| 色5月婷婷丁香| 久久久久久久午夜电影| 欧美老熟妇乱子伦牲交| 国产精品精品国产色婷婷| 久久久精品免费免费高清| 久久久久久久午夜电影| 国产午夜精品一二区理论片| 色哟哟·www| 看黄色毛片网站| 我要看日韩黄色一级片| 国产一区二区三区综合在线观看 | 日韩一区二区视频免费看| 国产一区亚洲一区在线观看| 新久久久久国产一级毛片| 免费av毛片视频| 久久精品国产a三级三级三级| 最新中文字幕久久久久| 国产欧美日韩一区二区三区在线 | 色哟哟·www| a级一级毛片免费在线观看| 国产成人精品婷婷| 国产大屁股一区二区在线视频| 久久久久久伊人网av| 人妻 亚洲 视频| 一区二区三区乱码不卡18| 国产精品无大码| 中文字幕久久专区| 国产精品久久久久久av不卡| 婷婷色麻豆天堂久久| 真实男女啪啪啪动态图| av国产精品久久久久影院| 91午夜精品亚洲一区二区三区| 极品少妇高潮喷水抽搐| 欧美性感艳星| 麻豆久久精品国产亚洲av| 欧美高清性xxxxhd video| 在线观看国产h片| 日本免费在线观看一区| 久久久久久久大尺度免费视频| 久久精品综合一区二区三区|