From 81d726030b35cff1b726f20a166ffff45c7979f2 Mon Sep 17 00:00:00 2001 From: James McClure Date: Fri, 12 Feb 2021 09:36:50 -0500 Subject: [PATCH] add freelee gpu skeleton --- cuda/FreeLee.cu | 2013 +++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 2013 insertions(+) create mode 100644 cuda/FreeLee.cu diff --git a/cuda/FreeLee.cu b/cuda/FreeLee.cu new file mode 100644 index 00000000..bc641ed9 --- /dev/null +++ b/cuda/FreeLee.cu @@ -0,0 +1,2013 @@ +#include + +#define STOKES + +__global__ void dvc_ScaLBL_D3Q19_FreeLeeModel_TwoFluid_Init(double *gqbar, double *mu_phi, double *ColorGrad, double Fx, double Fy, double Fz, int Np) +{ + int n; + double p = 1.0;//NOTE: take initial pressure p=1.0 + double chem; + double cg_x,cg_y,cg_z; + + //for (n=0; n 1.f) phi = 1.0; + if (phi < -1.f) phi = -1.0; + Den[idx] = rhoA + 0.5*(1.0-phi)*(rhoB-rhoA); + + //compute unit normal of color gradient + nx = ColorGrad[idx+0*Np]; + ny = ColorGrad[idx+1*Np]; + nz = ColorGrad[idx+2*Np]; + cg_mag = sqrt(nx*nx+ny*ny+nz*nz); + double ColorMag_temp = cg_mag; + if (cg_mag==0.0) ColorMag_temp=1.0; + nx = nx/ColorMag_temp; + ny = ny/ColorMag_temp; + nz = nz/ColorMag_temp; + + theta = M*cs2_inv*(1-4.0*phi*phi)/W; + + hq[0*Np+idx]=0.3333333333333333*(phi); + hq[1*Np+idx]=0.1111111111111111*(phi+theta*nx); + hq[2*Np+idx]=0.1111111111111111*(phi-theta*nx); + hq[3*Np+idx]=0.1111111111111111*(phi+theta*ny); + hq[4*Np+idx]=0.1111111111111111*(phi-theta*ny); + hq[5*Np+idx]=0.1111111111111111*(phi+theta*nz); + hq[6*Np+idx]=0.1111111111111111*(phi-theta*nz); + + } +} + +__global__ void dvc_ScaLBL_D3Q7_AAodd_FreeLeeModel_PhaseField(int *neighborList, int *Map, double *hq, double *Den, double *Phi, + double rhoA, double rhoB, int start, int finish, int Np){ + + int idx,n,nread; + double fq,phi; + int S = Np/NBLOCKS/NTHREADS + 1; + for (int s=0; s 10Np => odd part of dist) + m1 = dist[nr1]; // reading the f1 data into register fq + + nr2 = neighborList[n+Np]; // neighbor 1 ( < 10Np => even part of dist) + m2 = dist[nr2]; // reading the f2 data into register fq + + // q=3 + nr3 = neighborList[n+2*Np]; // neighbor 4 + m3 = dist[nr3]; + + // q = 4 + nr4 = neighborList[n+3*Np]; // neighbor 3 + m4 = dist[nr4]; + + // q=5 + nr5 = neighborList[n+4*Np]; + m5 = dist[nr5]; + + // q = 6 + nr6 = neighborList[n+5*Np]; + m6 = dist[nr6]; + + // q=7 + nr7 = neighborList[n+6*Np]; + m7 = dist[nr7]; + + // q = 8 + nr8 = neighborList[n+7*Np]; + m8 = dist[nr8]; + + // q=9 + nr9 = neighborList[n+8*Np]; + m9 = dist[nr9]; + + // q = 10 + nr10 = neighborList[n+9*Np]; + m10 = dist[nr10]; + + // q=11 + nr11 = neighborList[n+10*Np]; + m11 = dist[nr11]; + + // q=12 + nr12 = neighborList[n+11*Np]; + m12 = dist[nr12]; + + // q=13 + nr13 = neighborList[n+12*Np]; + m13 = dist[nr13]; + + // q=14 + nr14 = neighborList[n+13*Np]; + m14 = dist[nr14]; + + // q=15 + nr15 = neighborList[n+14*Np]; + m15 = dist[nr15]; + + // q=16 + nr16 = neighborList[n+15*Np]; + m16 = dist[nr16]; + + // q=17 + nr17 = neighborList[n+16*Np]; + m17 = dist[nr17]; + + // q=18 + nr18 = neighborList[n+17*Np]; + m18 = dist[nr18]; + + //compute fluid velocity + ux = 3.0/rho0*(m1-m2+m7-m8+m9-m10+m11-m12+m13-m14+0.5*(chem*nx+Fx)/3.0); + uy = 3.0/rho0*(m3-m4+m7-m8-m9+m10+m15-m16+m17-m18+0.5*(chem*ny+Fy)/3.0); + uz = 3.0/rho0*(m5-m6+m11-m12-m13+m14+m15-m16-m17+m18+0.5*(chem*nz+Fz)/3.0); + //compute pressure + p = (m0+m2+m1+m4+m3+m6+m5+m8+m7+m10+m9+m12+m11+m14+m13+m16+m15+m18+m17) + +0.5*(rhoA-rhoB)/2.0/3.0*(ux*nx+uy*ny+uz*nz); + + //compute equilibrium distributions + feq0 = 0.3333333333333333*p - 0.25*(Fx*ux + Fy*uy + Fz*uz)*(-0.6666666666666666 + ux*ux + uy*uy + uz*uz) - + 0.16666666666666666*rho0*(ux*ux + uy*uy + uz*uz) - 0.5*(-(nx*ux) - ny*uy - nz*uz)* + (-0.08333333333333333*(rhoA - rhoB)*(ux*ux + uy*uy + uz*uz) + chem*(0.3333333333333333 - 0.5*(ux*ux + uy*uy + uz*uz))); + feq1 = 0.05555555555555555*p - 0.08333333333333333*rho0*(-ux*ux + 0.3333333333333333*(-2*ux + ux*ux + uy*uy + uz*uz)) - + 0.125*(Fx*(-1. + ux) + Fy*uy + Fz*uz)*(-0.2222222222222222 - 1.*ux*ux + + 0.3333333333333333*(-2.*ux + ux*ux + uy*uy + uz*uz)) - 0.0625*(nx - nx*ux - ny*uy - nz*uz)* + (2*chem*ux*ux - 0.3333333333333333*((-rhoA + rhoB)*ux*ux + 2*chem*(-2*ux + ux*ux + uy*uy + uz*uz)) + + 0.1111111111111111*(4*chem - (rhoA - rhoB)*(-2*ux + ux*ux + uy*uy + uz*uz))); + feq2 = 0.05555555555555555*p - 0.08333333333333333*rho0*(-ux*ux + 0.3333333333333333*(2*ux + ux*ux + uy*uy + uz*uz)) - + 0.125*(Fx + Fx*ux + Fy*uy + Fz*uz)*(-0.2222222222222222 - 1.*ux*ux + + 0.3333333333333333*(2.*ux + ux*ux + uy*uy + uz*uz)) - 0.0625*(nx + nx*ux + ny*uy + nz*uz)* + (-2.*chem*ux*ux + 0.1111111111111111*(-4.*chem + rhoB*(-2.*ux - 1.*ux*ux - 1.*uy*uy - 1.*uz*uz) + + rhoA*(2.*ux + ux*ux + uy*uy + uz*uz)) + 0.3333333333333333*((-1.*rhoA + rhoB)*ux*ux + + chem*(4.*ux + 2.*ux*ux + 2.*uy*uy + 2.*uz*uz))); + feq3 = 0.05555555555555555*p - 0.08333333333333333*rho0*(-uy*uy + 0.3333333333333333*(ux*ux - 2*uy + uy*uy + uz*uz)) - + 0.125*(Fx*ux + Fy*(-1. + uy) + Fz*uz)*(-0.2222222222222222 - 1.*uy*uy + + 0.3333333333333333*(ux*ux - 2.*uy + uy*uy + uz*uz)) - 0.0625*(ny - nx*ux - ny*uy - nz*uz)* + (2*chem*uy*uy - 0.3333333333333333*((-rhoA + rhoB)*uy*uy + 2*chem*(ux*ux - 2*uy + uy*uy + uz*uz)) + + 0.1111111111111111*(4*chem - (rhoA - rhoB)*(ux*ux - 2*uy + uy*uy + uz*uz))); + feq4 = 0.05555555555555555*p - 0.08333333333333333*rho0*(-uy*uy + 0.3333333333333333*(ux*ux + 2*uy + uy*uy + uz*uz)) - + 0.125*(Fy + Fx*ux + Fy*uy + Fz*uz)*(-0.2222222222222222 - 1.*uy*uy + + 0.3333333333333333*(ux*ux + 2.*uy + uy*uy + uz*uz)) - 0.0625*(ny + nx*ux + ny*uy + nz*uz)* + (-2.*chem*uy*uy + 0.1111111111111111*(-4.*chem + rhoB*(-1.*ux*ux - 2.*uy - 1.*uy*uy - 1.*uz*uz) + + rhoA*(ux*ux + 2.*uy + uy*uy + uz*uz)) + 0.3333333333333333*((-1.*rhoA + rhoB)*uy*uy + + chem*(2.*ux*ux + 4.*uy + 2.*uy*uy + 2.*uz*uz))); + feq5 = 0.05555555555555555*p - 0.08333333333333333*rho0*(-uz*uz + 0.3333333333333333*(ux*ux + uy*uy + (-2 + uz)*uz)) - + 0.125*(Fx*ux + Fy*uy + Fz*(-1. + uz))*(-0.2222222222222222 - 1.*uz*uz + + 0.3333333333333333*(ux*ux + uy*uy + (-2. + uz)*uz)) - 0.0625*(nx*ux + ny*uy + nz*(-1. + uz))* + (-2.*chem*uz*uz + 0.1111111111111111*(-4.*chem + rhoB*(-1.*ux*ux - 1.*uy*uy + (2. - 1.*uz)*uz) + + rhoA*(ux*ux + uy*uy + (-2. + uz)*uz)) + 0.3333333333333333*((-1.*rhoA + rhoB)*uz*uz + + chem*(2.*ux*ux + 2.*uy*uy + uz*(-4. + 2.*uz)))); + feq6 = 0.05555555555555555*p - 0.08333333333333333*rho0*(-uz*uz + 0.3333333333333333*(ux*ux + uy*uy + uz*(2 + uz))) - + 0.125*(Fz + Fx*ux + Fy*uy + Fz*uz)*(-0.2222222222222222 - 1.*uz*uz + + 0.3333333333333333*(ux*ux + uy*uy + uz*(2. + uz))) - 0.0625*(nz + nx*ux + ny*uy + nz*uz)* + (-2.*chem*uz*uz + 0.1111111111111111*(-4.*chem + rhoB*(-1.*ux*ux - 1.*uy*uy + (-2. - 1.*uz)*uz) + + rhoA*(ux*ux + uy*uy + uz*(2. + uz))) + 0.3333333333333333*((-1.*rhoA + rhoB)*uz*uz + + chem*(2.*ux*ux + 2.*uy*uy + uz*(4. + 2.*uz)))); + feq7 = 0.027777777777777776*p - 0.041666666666666664*rho0* + (-(ux + uy)*(ux + uy) + 0.3333333333333333*(-2*ux + ux*ux - 2*uy + uy*uy + uz*uz)) - + 0.0625*(Fx*(-1. + ux) + Fy*(-1. + uy) + Fz*uz)*(-0.2222222222222222 - 1.*ux*ux - 2.*ux*uy - 1.*uy*uy + + 0.3333333333333333*(-2.*ux + ux*ux - 2.*uy + uy*uy + uz*uz)) - 0.03125*(nx + ny - nx*ux - ny*uy - nz*uz)* + (2*chem*(ux + uy)*(ux + uy) + 0.3333333333333333*((rhoA - rhoB)*(ux + uy)*(ux + uy) - 2*chem*(-2*ux + ux*ux - 2*uy + uy*uy + uz*uz)) + + 0.1111111111111111*(4*chem - (rhoA - rhoB)*(-2*ux + ux*ux - 2*uy + uy*uy + uz*uz))); + feq8 = 0.027777777777777776*p - 0.041666666666666664*rho0* + (-(ux + uy)*(ux + uy) + 0.3333333333333333*(2*ux + ux*ux + 2*uy + uy*uy + uz*uz)) - + 0.0625*(Fx + Fy + Fx*ux + Fy*uy + Fz*uz)*(-0.2222222222222222 - 1.*ux*ux - 2.*ux*uy - 1.*uy*uy + + 0.3333333333333333*(2.*ux + ux*ux + 2.*uy + uy*uy + uz*uz)) - 0.03125*(-(nx*(1 + ux)) - ny*(1 + uy) - nz*uz)* + (2*chem*(ux + uy)*(ux + uy) - 0.3333333333333333*(-((rhoA - rhoB)*(ux + uy)*(ux + uy)) + + 2*chem*(2*ux + ux*ux + 2*uy + uy*uy + uz*uz)) + 0.1111111111111111* + (4*chem - (rhoA - rhoB)*(2*ux + ux*ux + 2*uy + uy*uy + uz*uz))); + feq9 = 0.027777777777777776*p - 0.041666666666666664*rho0* + (-(ux - uy)*(ux - uy) + 0.3333333333333333*(-2*ux + ux*ux + 2*uy + uy*uy + uz*uz)) - + 0.0625*(Fy + Fx*(-1. + ux) + Fy*uy + Fz*uz)*(-0.2222222222222222 - 1.*ux*ux + 2.*ux*uy - 1.*uy*uy + + 0.3333333333333333*(-2.*ux + ux*ux + 2.*uy + uy*uy + uz*uz)) - 0.03125*(nx - nx*ux - ny*(1 + uy) - nz*uz)* + (2*chem*(ux - uy)*(ux - uy) - 0.3333333333333333*(-((rhoA - rhoB)*(ux - uy)*(ux - uy)) + + 2*chem*(-2*ux + ux*ux + 2*uy + uy*uy + uz*uz)) + 0.1111111111111111* + (4*chem - (rhoA - rhoB)*(-2*ux + ux*ux + 2*uy + uy*uy + uz*uz))); + feq10 = 0.027777777777777776*p - 0.041666666666666664*rho0* + (-(ux - uy)*(ux - uy) + 0.3333333333333333*(2*ux + ux*ux - 2*uy + uy*uy + uz*uz)) - + 0.0625*(Fx*(1 + ux) + Fy*(-1. + uy) + Fz*uz)*(-0.2222222222222222 - 1.*ux*ux + 2.*ux*uy - 1.*uy*uy + + 0.3333333333333333*(2.*ux + ux*ux - 2.*uy + uy*uy + uz*uz)) - 0.03125*(ny - nx*(1 + ux) - ny*uy - nz*uz)* + (2*chem*(ux - uy)*(ux - uy) - 0.3333333333333333*(-((rhoA - rhoB)*(ux - uy)*(ux - uy)) + + 2*chem*(2*ux + ux*ux - 2*uy + uy*uy + uz*uz)) + 0.1111111111111111* + (4*chem - (rhoA - rhoB)*(2*ux + ux*ux - 2*uy + uy*uy + uz*uz))); + feq11 = 0.027777777777777776*p - 0.041666666666666664*rho0* + (-(ux + uz)*(ux + uz) + 0.3333333333333333*(-2*ux + ux*ux + uy*uy + (-2 + uz)*uz)) - + 0.0625*(Fx*(-1. + ux) + Fy*uy + Fz*(-1. + uz))*(-0.2222222222222222 - 1.*ux*ux - 2.*ux*uz - 1.*uz*uz + + 0.3333333333333333*(-2.*ux + ux*ux + uy*uy + (-2. + uz)*uz)) - 0.03125*(nx + nz - nx*ux - ny*uy - nz*uz)* + (2*chem*(ux + uz)*(ux + uz) + 0.3333333333333333*((rhoA - rhoB)*(ux + uz)*(ux + uz) - 2*chem*(-2*ux + ux*ux + uy*uy + (-2 + uz)*uz)) + + 0.1111111111111111*(4*chem - (rhoA - rhoB)*(-2*ux + ux*ux + uy*uy + (-2 + uz)*uz))); + feq12 = 0.027777777777777776*p - 0.041666666666666664*rho0* + (-(ux + uz)*(ux + uz) + 0.3333333333333333*(2*ux + ux*ux + uy*uy + uz*(2 + uz))) - + 0.0625*(Fx + Fz + Fx*ux + Fy*uy + Fz*uz)*(-0.2222222222222222 - 1.*ux*ux - 2.*ux*uz - 1.*uz*uz + + 0.3333333333333333*(2.*ux + ux*ux + uy*uy + uz*(2. + uz))) - 0.03125*(-(nx*(1 + ux)) - ny*uy - nz*(1 + uz))* + (2*chem*(ux + uz)*(ux + uz) - 0.3333333333333333*(-((rhoA - rhoB)*(ux + uz)*(ux + uz)) + + 2*chem*(2*ux + ux*ux + uy*uy + uz*(2 + uz))) + 0.1111111111111111* + (4*chem - (rhoA - rhoB)*(2*ux + ux*ux + uy*uy + uz*(2 + uz)))); + feq13 = 0.027777777777777776*p - 0.041666666666666664*rho0* + (-(ux - uz)*(ux - uz) + 0.3333333333333333*(-2*ux + ux*ux + uy*uy + uz*(2 + uz))) - + 0.0625*(Fz + Fx*(-1. + ux) + Fy*uy + Fz*uz)*(-0.2222222222222222 - 1.*ux*ux + 2.*ux*uz - 1.*uz*uz + + 0.3333333333333333*(-2.*ux + ux*ux + uy*uy + uz*(2. + uz))) - 0.03125*(nx - nx*ux - ny*uy - nz*(1 + uz))* + (2*chem*(ux - uz)*(ux - uz) - 0.3333333333333333*(-((rhoA - rhoB)*(ux - uz)*(ux - uz)) + + 2*chem*(-2*ux + ux*ux + uy*uy + uz*(2 + uz))) + 0.1111111111111111* + (4*chem - (rhoA - rhoB)*(-2*ux + ux*ux + uy*uy + uz*(2 + uz)))); + feq14 = 0.027777777777777776*p - 0.041666666666666664*rho0* + (-(ux - uz)*(ux - uz) + 0.3333333333333333*(2*ux + ux*ux + uy*uy + (-2 + uz)*uz)) - + 0.0625*(Fx*(1 + ux) + Fy*uy + Fz*(-1. + uz))*(-0.2222222222222222 - 1.*ux*ux + 2.*ux*uz - 1.*uz*uz + + 0.3333333333333333*(2.*ux + ux*ux + uy*uy + (-2. + uz)*uz)) - 0.03125*(nz - nx*(1 + ux) - ny*uy - nz*uz)* + (2*chem*(ux - uz)*(ux - uz) - 0.3333333333333333*(-((rhoA - rhoB)*(ux - uz)*(ux - uz)) + + 2*chem*(2*ux + ux*ux + uy*uy + (-2 + uz)*uz)) + 0.1111111111111111* + (4*chem - (rhoA - rhoB)*(2*ux + ux*ux + uy*uy + (-2 + uz)*uz))); + feq15 = 0.027777777777777776*p - 0.041666666666666664*rho0* + (-(uy + uz)*(uy + uz) + 0.3333333333333333*(ux*ux - 2*uy + uy*uy + (-2 + uz)*uz)) - + 0.0625*(Fx*ux + Fy*(-1. + uy) + Fz*(-1. + uz))*(-0.2222222222222222 - 1.*uy*uy - 2.*uy*uz - 1.*uz*uz + + 0.3333333333333333*(ux*ux - 2.*uy + uy*uy + (-2. + uz)*uz)) - 0.03125*(ny + nz - nx*ux - ny*uy - nz*uz)* + (2*chem*(uy + uz)*(uy + uz) + 0.3333333333333333*((rhoA - rhoB)*(uy + uz)*(uy + uz) - 2*chem*(ux*ux - 2*uy + uy*uy + (-2 + uz)*uz)) + + 0.1111111111111111*(4*chem - (rhoA - rhoB)*(ux*ux - 2*uy + uy*uy + (-2 + uz)*uz))); + feq16 = 0.027777777777777776*p - 0.041666666666666664*rho0* + (-(uy + uz)*(uy + uz) + 0.3333333333333333*(ux*ux + 2*uy + uy*uy + uz*(2 + uz))) - + 0.0625*(Fy + Fz + Fx*ux + Fy*uy + Fz*uz)*(-0.2222222222222222 - 1.*uy*uy - 2.*uy*uz - 1.*uz*uz + + 0.3333333333333333*(ux*ux + 2.*uy + uy*uy + uz*(2. + uz))) - 0.03125*(-(nx*ux) - ny*(1 + uy) - nz*(1 + uz))* + (2*chem*(uy + uz)*(uy + uz) - 0.3333333333333333*(-((rhoA - rhoB)*(uy + uz)*(uy + uz)) + + 2*chem*(ux*ux + 2*uy + uy*uy + uz*(2 + uz))) + 0.1111111111111111* + (4*chem - (rhoA - rhoB)*(ux*ux + 2*uy + uy*uy + uz*(2 + uz)))); + feq17 = 0.027777777777777776*p - 0.041666666666666664*rho0* + (-(uy - uz)*(uy - uz) + 0.3333333333333333*(ux*ux - 2*uy + uy*uy + uz*(2 + uz))) - + 0.0625*(Fz + Fx*ux + Fy*(-1. + uy) + Fz*uz)*(-0.2222222222222222 - 1.*uy*uy + 2.*uy*uz - 1.*uz*uz + + 0.3333333333333333*(ux*ux - 2.*uy + uy*uy + uz*(2. + uz))) - 0.03125*(ny - nx*ux - ny*uy - nz*(1 + uz))* + (2*chem*(uy - uz)*(uy - uz) - 0.3333333333333333*(-((rhoA - rhoB)*(uy - uz)*(uy - uz)) + + 2*chem*(ux*ux - 2*uy + uy*uy + uz*(2 + uz))) + 0.1111111111111111* + (4*chem - (rhoA - rhoB)*(ux*ux - 2*uy + uy*uy + uz*(2 + uz)))); + feq18 = 0.027777777777777776*p - 0.041666666666666664*rho0* + (-(uy - uz)*(uy - uz) + 0.3333333333333333*(ux*ux + 2*uy + uy*uy + (-2 + uz)*uz)) - + 0.0625*(Fx*ux + Fy*(1 + uy) + Fz*(-1. + uz))*(-0.2222222222222222 - 1.*uy*uy + 2.*uy*uz - 1.*uz*uz + + 0.3333333333333333*(ux*ux + 2.*uy + uy*uy + (-2. + uz)*uz)) - 0.03125*(nz - nx*ux - ny*(1 + uy) - nz*uz)* + (2*chem*(uy - uz)*(uy - uz) - 0.3333333333333333*(-((rhoA - rhoB)*(uy - uz)*(uy - uz)) + + 2*chem*(ux*ux + 2*uy + uy*uy + (-2 + uz)*uz)) + 0.1111111111111111* + (4*chem - (rhoA - rhoB)*(ux*ux + 2*uy + uy*uy + (-2 + uz)*uz))); + + //------------------------------------------------- BCK collison ------------------------------------------------------------// + // q=0 + dist[n] = m0 - (m0-feq0)/tau + 0.25*(2*(Fx*ux + Fy*uy + Fz*uz)*(-0.6666666666666666 + ux*ux + uy*uy + uz*uz) + + (mgx*ux + mgy*uy + mgz*uz)*(2*chem*(ux*ux + uy*uy + uz*uz) + + 0.3333333333333333*(-4*chem + (rhoA - rhoB)*(ux*ux + uy*uy + uz*uz)))); + + // q = 1 + dist[nr2] = m1 - (m1-feq1)/tau + 0.125*(2*(Fx*(-1 + ux) + Fy*uy + Fz*uz)*(-0.2222222222222222 - ux*ux + + 0.3333333333333333*(-2*ux + ux*ux + uy*uy + uz*uz)) + + (mgx*(-1 + ux) + mgy*uy + mgz*uz)*(-2*chem*(ux*ux) + + 0.3333333333333333*((-rhoA + rhoB)*(ux*ux) + 2*chem*(-2*ux + ux*ux + uy*uy + uz*uz)) + + 0.1111111111111111*(-4*chem + (rhoA - rhoB)*(-2*ux + ux*ux + uy*uy + uz*uz)))); + + // q=2 + dist[nr1] = m2 - (m2-feq2)/tau + 0.125*(2*(Fx + Fx*ux + Fy*uy + Fz*uz)*(-0.2222222222222222 - ux*ux + + 0.3333333333333333*(2*ux + ux*ux + uy*uy + uz*uz)) + + (mgx + mgx*ux + mgy*uy + mgz*uz)*(-2*chem*(ux*ux) + + 0.3333333333333333*((-rhoA + rhoB)*(ux*ux) + 2*chem*(2*ux + ux*ux + uy*uy + uz*uz)) + + 0.1111111111111111*(-4*chem + (rhoA - rhoB)*(2*ux + ux*ux + uy*uy + uz*uz)))); + + // q = 3 + dist[nr4] = m3 - (m3-feq3)/tau + 0.125*(2*(Fx*ux + Fy*(-1 + uy) + Fz*uz)*(-0.2222222222222222 - uy*uy + + 0.3333333333333333*(ux*ux - 2*uy + uy*uy + uz*uz)) + + (mgx*ux + mgy*(-1 + uy) + mgz*uz)*(-2*chem*(uy*uy) + + 0.3333333333333333*((-rhoA + rhoB)*(uy*uy) + 2*chem*(ux*ux - 2*uy + uy*uy + uz*uz)) + + 0.1111111111111111*(-4*chem + (rhoA - rhoB)*(ux*ux - 2*uy + uy*uy + uz*uz)))); + + // q = 4 + dist[nr3] = m4 - (m4-feq4)/tau + 0.125*(2*(Fy + Fx*ux + Fy*uy + Fz*uz)*(-0.2222222222222222 - uy*uy + + 0.3333333333333333*(ux*ux + 2*uy + uy*uy + uz*uz)) + + (mgy + mgx*ux + mgy*uy + mgz*uz)*(-2*chem*(uy*uy) + + 0.3333333333333333*((-rhoA + rhoB)*(uy*uy) + 2*chem*(ux*ux + 2*uy + uy*uy + uz*uz)) + + 0.1111111111111111*(-4*chem + (rhoA - rhoB)*(ux*ux + 2*uy + uy*uy + uz*uz)))); + + // q = 5 + dist[nr6] = m5 - (m5-feq5)/tau + 0.125*(2*(Fx*ux + Fy*uy + Fz*(-1 + uz))*(-0.2222222222222222 - uz*uz + + 0.3333333333333333*(ux*ux + uy*uy + (-2 + uz)*uz)) + + (mgx*ux + mgy*uy + mgz*(-1 + uz))*(-2*chem*(uz*uz) + + 0.3333333333333333*((-rhoA + rhoB)*(uz*uz) + 2*chem*(ux*ux + uy*uy + (-2 + uz)*uz)) + + 0.1111111111111111*(-4*chem + (rhoA - rhoB)*(ux*ux + uy*uy + (-2 + uz)*uz)))); + + // q = 6 + dist[nr5] = m6 - (m6-feq6)/tau + 0.125*(2*(Fz + Fx*ux + Fy*uy + Fz*uz)*(-0.2222222222222222 - uz*uz + + 0.3333333333333333*(ux*ux + uy*uy + uz*(2 + uz))) + + (mgz + mgx*ux + mgy*uy + mgz*uz)*(-2*chem*(uz*uz) + + 0.3333333333333333*((-rhoA + rhoB)*(uz*uz) + 2*chem*(ux*ux + uy*uy + uz*(2 + uz))) + + 0.1111111111111111*(-4*chem + (rhoA - rhoB)*(ux*ux + uy*uy + uz*(2 + uz))))); + + // q = 7 + dist[nr8] = m7 - (m7-feq7)/tau + 0.0625*(-2*(Fx*(-1 + ux) + Fy*(-1 + uy) + Fz*uz)* + (0.2222222222222222 + (ux + uy)*(ux + uy) - + 0.3333333333333333*(-2*ux + ux*ux - 2*uy + uy*uy + uz*uz)) + + (mgx*(-1 + ux) + mgy*(-1 + uy) + mgz*uz)* + (-2*chem*((ux + uy)*(ux + uy)) + 0.3333333333333333* + (-((rhoA - rhoB)*((ux + uy)*(ux + uy))) + 2*chem*(-2*ux + ux*ux - 2*uy + uy*uy + uz*uz)) + + 0.1111111111111111*(-4*chem + (rhoA - rhoB)*(-2*ux + ux*ux - 2*uy + uy*uy + uz*uz)))); + + // q = 8 + dist[nr7] = m8 - (m8-feq8)/tau + 0.0625*(2*(Fx + Fy + Fx*ux + Fy*uy + Fz*uz)* + (-0.2222222222222222 - (ux + uy)*(ux + uy) + + 0.3333333333333333*(2*ux + ux*ux + 2*uy + uy*uy + uz*uz)) + + (mgx + mgy + mgx*ux + mgy*uy + mgz*uz)* + (-2*chem*((ux + uy)*(ux + uy)) + 0.3333333333333333* + (-((rhoA - rhoB)*((ux + uy)*(ux + uy))) + 2*chem*(2*ux + ux*ux + 2*uy + uy*uy + uz*uz)) + + 0.1111111111111111*(-4*chem + (rhoA - rhoB)*(2*ux + ux*ux + 2*uy + uy*uy + uz*uz)))); + + // q = 9 + dist[nr10] = m9 - (m9-feq9)/tau + 0.0625*(2*(Fy + Fx*(-1 + ux) + Fy*uy + Fz*uz)* + (-0.2222222222222222 - (ux - uy)*(ux - uy) + + 0.3333333333333333*(-2*ux + ux*ux + 2*uy + uy*uy + uz*uz)) + + (mgy + mgx*(-1 + ux) + mgy*uy + mgz*uz)* + (-2*chem*((ux - uy)*(ux - uy)) + 0.3333333333333333* + (-((rhoA - rhoB)*((ux - uy)*(ux - uy))) + 2*chem*(-2*ux + ux*ux + 2*uy + uy*uy + uz*uz)) + + 0.1111111111111111*(-4*chem + (rhoA - rhoB)*(-2*ux + ux*ux + 2*uy + uy*uy + uz*uz)))); + + // q = 10 + dist[nr9] = m10 - (m10-feq10)/tau + 0.0625*(2*(Fx*(1 + ux) + Fy*(-1 + uy) + Fz*uz)* + (-0.2222222222222222 - (ux - uy)*(ux - uy) + + 0.3333333333333333*(2*ux + ux*ux - 2*uy + uy*uy + uz*uz)) + + (mgx*(1 + ux) + mgy*(-1 + uy) + mgz*uz)* + (-2*chem*((ux - uy)*(ux - uy)) + 0.3333333333333333* + (-((rhoA - rhoB)*((ux - uy)*(ux - uy))) + 2*chem*(2*ux + ux*ux - 2*uy + uy*uy + uz*uz)) + + 0.1111111111111111*(-4*chem + (rhoA - rhoB)*(2*ux + ux*ux - 2*uy + uy*uy + uz*uz)))); + + // q = 11 + dist[nr12] = m11 - (m11-feq11)/tau + 0.0625*(-2*(Fx*(-1 + ux) + Fy*uy + Fz*(-1 + uz))* + (0.2222222222222222 + (ux + uz)*(ux + uz) - + 0.3333333333333333*(-2*ux + ux*ux + uy*uy + (-2 + uz)*uz)) + + (mgx*(-1 + ux) + mgy*uy + mgz*(-1 + uz))* + (-2*chem*((ux + uz)*(ux + uz)) + 0.3333333333333333* + (-((rhoA - rhoB)*((ux + uz)*(ux + uz))) + 2*chem*(-2*ux + ux*ux + uy*uy + (-2 + uz)*uz)) + + 0.1111111111111111*(-4*chem + (rhoA - rhoB)*(-2*ux + ux*ux + uy*uy + (-2 + uz)*uz)))); + + // q = 12 + dist[nr11] = m12 - (m12-feq12)/tau + 0.0625*(2*(Fx + Fz + Fx*ux + Fy*uy + Fz*uz)* + (-0.2222222222222222 - (ux + uz)*(ux + uz) + 0.3333333333333333*(2*ux + ux*ux + uy*uy + uz*(2 + uz))) + + (mgx + mgz + mgx*ux + mgy*uy + mgz*uz)* + (-2*chem*((ux + uz)*(ux + uz)) + 0.3333333333333333* + (-((rhoA - rhoB)*((ux + uz)*(ux + uz))) + 2*chem*(2*ux + ux*ux + uy*uy + uz*(2 + uz))) + + 0.1111111111111111*(-4*chem + (rhoA - rhoB)*(2*ux + ux*ux + uy*uy + uz*(2 + uz))))); + + // q = 13 + dist[nr14] = m13 - (m13-feq13)/tau + 0.0625*(2*(Fz + Fx*(-1 + ux) + Fy*uy + Fz*uz)* + (-0.2222222222222222 - (ux - uz)*(ux - uz) + + 0.3333333333333333*(-2*ux + ux*ux + uy*uy + uz*(2 + uz))) + + (mgz + mgx*(-1 + ux) + mgy*uy + mgz*uz)* + (-2*chem*((ux - uz)*(ux - uz)) + 0.3333333333333333* + (-((rhoA - rhoB)*((ux - uz)*(ux - uz))) + 2*chem*(-2*ux + ux*ux + uy*uy + uz*(2 + uz))) + + 0.1111111111111111*(-4*chem + (rhoA - rhoB)*(-2*ux + ux*ux + uy*uy + uz*(2 + uz))))); + + // q= 14 + dist[nr13] = m14 - (m14-feq14)/tau + 0.0625*(2*(Fx*(1 + ux) + Fy*uy + Fz*(-1 + uz))* + (-0.2222222222222222 - (ux - uz)*(ux - uz) + + 0.3333333333333333*(2*ux + ux*ux + uy*uy + (-2 + uz)*uz)) + + (mgx*(1 + ux) + mgy*uy + mgz*(-1 + uz))* + (-2*chem*((ux - uz)*(ux - uz)) + 0.3333333333333333* + (-((rhoA - rhoB)*((ux - uz)*(ux - uz))) + 2*chem*(2*ux + ux*ux + uy*uy + (-2 + uz)*uz)) + + 0.1111111111111111*(-4*chem + (rhoA - rhoB)*(2*ux + ux*ux + uy*uy + (-2 + uz)*uz)))); + + // q = 15 + dist[nr16] = m15 - (m15-feq15)/tau + 0.0625*(-2*(Fx*ux + Fy*(-1 + uy) + Fz*(-1 + uz))* + (0.2222222222222222 + (uy + uz)*(uy + uz) - 0.3333333333333333*(ux*ux - 2*uy + uy*uy + (-2 + uz)*uz)) + + (mgx*ux + mgy*(-1 + uy) + mgz*(-1 + uz))* + (-2*chem*((uy + uz)*(uy + uz)) + 0.3333333333333333* + (-((rhoA - rhoB)*((uy + uz)*(uy + uz))) + 2*chem*(ux*ux - 2*uy + uy*uy + (-2 + uz)*uz)) + + 0.1111111111111111*(-4*chem + (rhoA - rhoB)*(ux*ux - 2*uy + uy*uy + (-2 + uz)*uz)))); + + // q = 16 + dist[nr15] = m16 - (m16-feq16)/tau + 0.0625*(2*(Fy + Fz + Fx*ux + Fy*uy + Fz*uz)* + (-0.2222222222222222 - (uy + uz)*(uy + uz) + 0.3333333333333333*(ux*ux + 2*uy + uy*uy + uz*(2 + uz))) + + (mgy + mgz + mgx*ux + mgy*uy + mgz*uz)* + (-2*chem*((uy + uz)*(uy + uz)) + 0.3333333333333333* + (-((rhoA - rhoB)*((uy + uz)*(uy + uz))) + 2*chem*(ux*ux + 2*uy + uy*uy + uz*(2 + uz))) + + 0.1111111111111111*(-4*chem + (rhoA - rhoB)*(ux*ux + 2*uy + uy*uy + uz*(2 + uz))))); + + // q = 17 + dist[nr18] = m17 - (m17-feq17)/tau + 0.0625*(2*(Fz + Fx*ux + Fy*(-1 + uy) + Fz*uz)* + (-0.2222222222222222 - (uy - uz)*(uy - uz) + 0.3333333333333333*(ux*ux - 2*uy + uy*uy + uz*(2 + uz))) + + (mgz + mgx*ux + mgy*(-1 + uy) + mgz*uz)* + (-2*chem*((uy - uz)*(uy - uz)) + 0.3333333333333333* + (-((rhoA - rhoB)*((uy - uz)*(uy - uz))) + 2*chem*(ux*ux - 2*uy + uy*uy + uz*(2 + uz))) + + 0.1111111111111111*(-4*chem + (rhoA - rhoB)*(ux*ux - 2*uy + uy*uy + uz*(2 + uz))))); + + // q = 18 + dist[nr17] = m18 - (m18-feq18)/tau + 0.0625*(2*(Fx*ux + Fy*(1 + uy) + Fz*(-1 + uz))* + (-0.2222222222222222 - (uy - uz)*(uy - uz) + + 0.3333333333333333*(ux*ux + 2*uy + uy*uy + (-2 + uz)*uz)) + + (mgx*ux + mgy*(1 + uy) + mgz*(-1 + uz))* + (-2*chem*((uy - uz)*(uy - uz)) + 0.3333333333333333* + (-((rhoA - rhoB)*((uy - uz)*(uy - uz))) + 2*chem*(ux*ux + 2*uy + uy*uy + (-2 + uz)*uz)) + + 0.1111111111111111*(-4*chem + (rhoA - rhoB)*(ux*ux + 2*uy + uy*uy + (-2 + uz)*uz)))); + //----------------------------------------------------------------------------------------------------------------------------------------// + + + // ----------------------------- compute phase field evolution ---------------------------------------- + //Normalize the Color Gradient + C = sqrt(nx*nx+ny*ny+nz*nz); + double ColorMag = C; + if (C==0.0) ColorMag=1.0; + nx = nx/ColorMag; + ny = ny/ColorMag; + nz = nz/ColorMag; + //compute surface tension-related parameter + theta = M*4.5*(1-4.0*phi*phi)/W; + + //load distributions of phase field + //q=0 + h0 = hq[n]; + //q=1 + h1 = hq[nr1]; + + //q=2 + h2 = hq[nr2]; + + //q=3 + h3 = hq[nr3]; + + //q=4 + h4 = hq[nr4]; + + //q=5 + h5 = hq[nr5]; + + //q=6 + h6 = hq[nr6]; + + //-------------------------------- BGK collison for phase field ---------------------------------// + // q = 0 + hq[n] = h0 - (h0 - 0.3333333333333333*phi)/tauM; + + // q = 1 + hq[nr2] = h1 - (h1 - phi*(0.1111111111111111 + 0.5*ux) - (0.5*M*nx*(1 - 4*phi*phi))/W)/tauM; + + // q = 2 + hq[nr1] = h2 - (h2 - phi*(0.1111111111111111 - 0.5*ux) + (0.5*M*nx*(1 - 4*phi*phi))/W)/tauM; + + // q = 3 + hq[nr4] = h3 - (h3 - phi*(0.1111111111111111 + 0.5*uy) - (0.5*M*ny*(1 - 4*phi*phi))/W)/tauM; + + // q = 4 + hq[nr3] = h4 - (h4 - phi*(0.1111111111111111 - 0.5*uy) + (0.5*M*ny*(1 - 4*phi*phi))/W)/tauM; + + // q = 5 + hq[nr6] = h5 - (h5 - phi*(0.1111111111111111 + 0.5*uz) - (0.5*M*nz*(1 - 4*phi*phi))/W)/tauM; + + // q = 6 + hq[nr5] = h6 - (h6 - phi*(0.1111111111111111 - 0.5*uz) + (0.5*M*nz*(1 - 4*phi*phi))/W)/tauM; + //........................................................................ + + //Update velocity on device + Vel[0*Np+n] = ux; + Vel[1*Np+n] = uy; + Vel[2*Np+n] = uz; + //Update pressure on device + Pressure[n] = p; + //Update chemical potential on device + mu_phi[n] = chem; + //Update color gradient on device + ColorGrad[0*Np+n] = nx; + ColorGrad[1*Np+n] = ny; + ColorGrad[2*Np+n] = nz; + } + } +} + +__global__ void dvc_ScaLBL_D3Q19_AAeven_FreeLeeModel(int *Map, double *dist, double *hq, double *Den, double *Phi, double *mu_phi, double *Vel, double *Pressure, double *ColorGrad, + double rhoA, double rhoB, double tauA, double tauB, double tauM, double kappa, double beta, double W, double Fx, double Fy, double Fz, + int strideY, int strideZ, int start, int finish, int Np){ + + int n,nn,nn2x,ijk; + //int nr1,nr2,nr3,nr4,nr5,nr6,nr7,nr8,nr9,nr10,nr11,nr12,nr13,nr14,nr15,nr16,nr17,nr18; + double ux,uy,uz;//fluid velocity + double p;//pressure + double chem;//chemical potential + double phi; //phase field + double rho0;//fluid density + // distribution functions + double m1,m2,m4,m6,m8,m9,m10,m11,m12,m13,m14,m15,m16,m17,m18; + double m0,m3,m5,m7; + double mm1,mm2,mm4,mm6,mm8,mm9,mm10,mm11,mm12,mm13,mm14,mm15,mm16,mm17,mm18; + double mm3,mm5,mm7; + double feq0,feq1,feq2,feq3,feq4,feq5,feq6,feq7,feq8,feq9,feq10,feq11,feq12,feq13,feq14,feq15,feq16,feq17,feq18; + double nx,ny,nz;//normal color gradient + double mgx,mgy,mgz;//mixed gradient reaching secondary neighbor + + //double f0,f1,f2,f3,f4,f5,f6,f7,f8,f9,f10,f11,f12,f13,f14,f15,f16,f17,f18; + double h0,h1,h2,h3,h4,h5,h6;//distributions for LB phase field + double tau;//position dependent LB relaxation time for fluid + double C,theta; + double M = 2.0/9.0*(tauM-0.5);//diffusivity (or mobility) for the phase field D3Q7 + + // for (int n=start; n 10Np => odd part of dist) + m1 = dist[nr1]; // reading the f1 data into register fq + + nr2 = neighborList[n+Np]; // neighbor 1 ( < 10Np => even part of dist) + m2 = dist[nr2]; // reading the f2 data into register fq + + // q=3 + nr3 = neighborList[n+2*Np]; // neighbor 4 + m3 = dist[nr3]; + + // q = 4 + nr4 = neighborList[n+3*Np]; // neighbor 3 + m4 = dist[nr4]; + + // q=5 + nr5 = neighborList[n+4*Np]; + m5 = dist[nr5]; + + // q = 6 + nr6 = neighborList[n+5*Np]; + m6 = dist[nr6]; + + // q=7 + nr7 = neighborList[n+6*Np]; + m7 = dist[nr7]; + + // q = 8 + nr8 = neighborList[n+7*Np]; + m8 = dist[nr8]; + + // q=9 + nr9 = neighborList[n+8*Np]; + m9 = dist[nr9]; + + // q = 10 + nr10 = neighborList[n+9*Np]; + m10 = dist[nr10]; + + // q=11 + nr11 = neighborList[n+10*Np]; + m11 = dist[nr11]; + + // q=12 + nr12 = neighborList[n+11*Np]; + m12 = dist[nr12]; + + // q=13 + nr13 = neighborList[n+12*Np]; + m13 = dist[nr13]; + + // q=14 + nr14 = neighborList[n+13*Np]; + m14 = dist[nr14]; + + // q=15 + nr15 = neighborList[n+14*Np]; + m15 = dist[nr15]; + + // q=16 + nr16 = neighborList[n+15*Np]; + m16 = dist[nr16]; + + // q=17 + nr17 = neighborList[n+16*Np]; + m17 = dist[nr17]; + + // q=18 + nr18 = neighborList[n+17*Np]; + m18 = dist[nr18]; + + //compute fluid velocity + ux = 3.0/rho0*(m1-m2+m7-m8+m9-m10+m11-m12+m13-m14+0.5*(Fx)/3.0); + uy = 3.0/rho0*(m3-m4+m7-m8-m9+m10+m15-m16+m17-m18+0.5*(Fy)/3.0); + uz = 3.0/rho0*(m5-m6+m11-m12-m13+m14+m15-m16-m17+m18+0.5*(Fz)/3.0); + //compute pressure + p = (m0+m2+m1+m4+m3+m6+m5+m8+m7+m10+m9+m12+m11+m14+m13+m16+m15+m18+m17); + + //------------------------------------------------- BCK collison ------------------------------------------------------------// + // q=0 + dist[n] = m0 + 0.5*(Fx*ux + Fy*uy + Fz*uz)*(-0.6666666666666666 + ux*ux + uy*uy + uz*uz) - + (m0 - 0.3333333333333333*p + 0.25*(Fx*ux + Fy*uy + Fz*uz)* + (-0.6666666666666666 + ux*ux + uy*uy + uz*uz) + 0.16666666666666666*rho0*(ux*ux + uy*uy + uz*uz))/ + tau; + + // q = 1 + dist[nr2] = m1 + 0.25*(Fx*(-1 + ux) + Fy*uy + Fz*uz)*(-0.2222222222222222 - ux*ux + + 0.3333333333333333*(-2*ux + ux*ux + uy*uy + uz*uz)) - + (m1 - 0.05555555555555555*p + 0.08333333333333333*rho0* + (-(ux*ux) + 0.3333333333333333*(-2*ux + ux*ux + uy*uy + uz*uz)) + + 0.125*(Fx*(-1. + ux) + Fy*uy + Fz*uz)* + (-0.2222222222222222 - 1.*(ux*ux) + 0.3333333333333333*(-2.*ux + ux*ux + uy*uy + uz*uz)))/tau; + + // q=2 + dist[nr1] = m2 + 0.25*(Fx + Fx*ux + Fy*uy + Fz*uz)*(-0.2222222222222222 - ux*ux + + 0.3333333333333333*(2*ux + ux*ux + uy*uy + uz*uz)) - + (m2 - 0.05555555555555555*p + 0.08333333333333333*rho0* + (-(ux*ux) + 0.3333333333333333*(2*ux + ux*ux + uy*uy + uz*uz)) + + 0.125*(Fx + Fx*ux + Fy*uy + Fz*uz)*(-0.2222222222222222 - 1.*(ux*ux) + + 0.3333333333333333*(2.*ux + ux*ux + uy*uy + uz*uz)))/tau; + + // q = 3 + dist[nr4] = m3 + 0.25*(Fx*ux + Fy*(-1 + uy) + Fz*uz)*(-0.2222222222222222 - uy*uy + + 0.3333333333333333*(ux*ux - 2*uy + uy*uy + uz*uz)) - + (m3 - 0.05555555555555555*p + 0.08333333333333333*rho0* + (-(uy*uy) + 0.3333333333333333*(ux*ux - 2*uy + uy*uy + uz*uz)) + + 0.125*(Fx*ux + Fy*(-1. + uy) + Fz*uz)* + (-0.2222222222222222 - 1.*(uy*uy) + 0.3333333333333333*(ux*ux - 2.*uy + uy*uy + uz*uz)))/tau; + + // q = 4 + dist[nr3] = m4 + 0.25*(Fy + Fx*ux + Fy*uy + Fz*uz)*(-0.2222222222222222 - uy*uy + + 0.3333333333333333*(ux*ux + 2*uy + uy*uy + uz*uz)) - + (m4 - 0.05555555555555555*p + 0.08333333333333333*rho0* + (-(uy*uy) + 0.3333333333333333*(ux*ux + 2*uy + uy*uy + uz*uz)) + + 0.125*(Fy + Fx*ux + Fy*uy + Fz*uz)*(-0.2222222222222222 - 1.*(uy*uy) + + 0.3333333333333333*(ux*ux + 2.*uy + uy*uy + uz*uz)))/tau; + + // q = 5 + dist[nr6] = m5 + 0.25*(Fx*ux + Fy*uy + Fz*(-1 + uz))*(-0.2222222222222222 - uz*uz + + 0.3333333333333333*(ux*ux + uy*uy + (-2 + uz)*uz)) - + (m5 - 0.05555555555555555*p + 0.08333333333333333*rho0* + (-(uz*uz) + 0.3333333333333333*(ux*ux + uy*uy + (-2 + uz)*uz)) + + 0.125*(Fx*ux + Fy*uy + Fz*(-1. + uz))* + (-0.2222222222222222 - 1.*(uz*uz) + 0.3333333333333333*(ux*ux + uy*uy + (-2. + uz)*uz)))/tau; + + // q = 6 + dist[nr5] = m6 + 0.25*(Fz + Fx*ux + Fy*uy + Fz*uz)*(-0.2222222222222222 - uz*uz + + 0.3333333333333333*(ux*ux + uy*uy + uz*(2 + uz))) - + (m6 - 0.05555555555555555*p + 0.08333333333333333*rho0* + (-(uz*uz) + 0.3333333333333333*(ux*ux + uy*uy + uz*(2 + uz))) + + 0.125*(Fz + Fx*ux + Fy*uy + Fz*uz)*(-0.2222222222222222 - 1.*(uz*uz) + + 0.3333333333333333*(ux*ux + uy*uy + uz*(2. + uz))))/tau; + + // q = 7 + dist[nr8] = m7 - 0.125*(Fx*(-1 + ux) + Fy*(-1 + uy) + Fz*uz)* + (0.2222222222222222 + (ux + uy)*(ux + uy) - 0.3333333333333333*(-2*ux + ux*ux - 2*uy + uy*uy + uz*uz))\ + - (m7 - 0.027777777777777776*p + 0.041666666666666664*rho0* + (-((ux + uy)*(ux + uy)) + 0.3333333333333333*(-2*ux + ux*ux - 2*uy + uy*uy + uz*uz)) + + 0.0625*(Fx*(-1. + ux) + Fy*(-1. + uy) + Fz*uz)* + (-0.2222222222222222 - 1.*(ux*ux) - 2.*ux*uy - 1.*(uy*uy) + + 0.3333333333333333*(-2.*ux + ux*ux - 2.*uy + uy*uy + uz*uz)))/tau; + + // q = 8 + dist[nr7] = m8 + 0.125*(Fx + Fy + Fx*ux + Fy*uy + Fz*uz)* + (-0.2222222222222222 - (ux + uy)*(ux + uy) + 0.3333333333333333*(2*ux + ux*ux + 2*uy + uy*uy + uz*uz))\ + - (m8 - 0.027777777777777776*p + 0.041666666666666664*rho0* + (-((ux + uy)*(ux + uy)) + 0.3333333333333333*(2*ux + ux*ux + 2*uy + uy*uy + uz*uz)) + + 0.0625*(Fx + Fy + Fx*ux + Fy*uy + Fz*uz)* + (-0.2222222222222222 - 1.*(ux*ux) - 2.*ux*uy - 1.*(uy*uy) + + 0.3333333333333333*(2.*ux + ux*ux + 2.*uy + uy*uy + uz*uz)))/tau; + + // q = 9 + dist[nr10] = m9 + 0.125*(Fy + Fx*(-1 + ux) + Fy*uy + Fz*uz)* + (-0.2222222222222222 - (ux - uy)*(ux - uy) + 0.3333333333333333*(-2*ux + ux*ux + 2*uy + uy*uy + uz*uz)) + - (m9 - 0.027777777777777776*p + 0.041666666666666664*rho0* + (-((ux - uy)*(ux - uy)) + 0.3333333333333333*(-2*ux + ux*ux + 2*uy + uy*uy + uz*uz)) + + 0.0625*(Fy + Fx*(-1. + ux) + Fy*uy + Fz*uz)* + (-0.2222222222222222 - 1.*(ux*ux) + 2.*ux*uy - 1.*(uy*uy) + + 0.3333333333333333*(-2.*ux + ux*ux + 2.*uy + uy*uy + uz*uz)))/tau; + + // q = 10 + dist[nr9] = m10 + 0.125*(Fx*(1 + ux) + Fy*(-1 + uy) + Fz*uz)* + (-0.2222222222222222 - (ux - uy)*(ux - uy) + 0.3333333333333333*(2*ux + ux*ux - 2*uy + uy*uy + uz*uz))\ + - (m10 - 0.027777777777777776*p + 0.041666666666666664*rho0* + (-((ux - uy)*(ux - uy)) + 0.3333333333333333*(2*ux + ux*ux - 2*uy + uy*uy + uz*uz)) + + 0.0625*(Fx*(1 + ux) + Fy*(-1. + uy) + Fz*uz)* + (-0.2222222222222222 - 1.*(ux*ux) + 2.*ux*uy - 1.*(uy*uy) + + 0.3333333333333333*(2.*ux + ux*ux - 2.*uy + uy*uy + uz*uz)))/tau; + + // q = 11 + dist[nr12] = m11 - 0.125*(Fx*(-1 + ux) + Fy*uy + Fz*(-1 + uz))* + (0.2222222222222222 + (ux + uz)*(ux + uz) - 0.3333333333333333*(-2*ux + ux*ux + uy*uy + (-2 + uz)*uz))\ + - (m11 - 0.027777777777777776*p + 0.041666666666666664*rho0* + (-((ux + uz)*(ux + uz)) + 0.3333333333333333*(-2*ux + ux*ux + uy*uy + (-2 + uz)*uz)) + + 0.0625*(Fx*(-1. + ux) + Fy*uy + Fz*(-1. + uz))* + (-0.2222222222222222 - 1.*(ux*ux) - 2.*ux*uz - 1.*(uz*uz) + + 0.3333333333333333*(-2.*ux + ux*ux + uy*uy + (-2. + uz)*uz)))/tau; + + // q = 12 + dist[nr11] = m12 + 0.125*(Fx + Fz + Fx*ux + Fy*uy + Fz*uz)* + (-0.2222222222222222 - (ux + uz)*(ux + uz) + 0.3333333333333333*(2*ux + ux*ux + uy*uy + uz*(2 + uz)))\ + - (m12 - 0.027777777777777776*p + 0.041666666666666664*rho0* + (-((ux + uz)*(ux + uz)) + 0.3333333333333333*(2*ux + ux*ux + uy*uy + uz*(2 + uz))) + + 0.0625*(Fx + Fz + Fx*ux + Fy*uy + Fz*uz)* + (-0.2222222222222222 - 1.*(ux*ux) - 2.*ux*uz - 1.*(uz*uz) + + 0.3333333333333333*(2.*ux + ux*ux + uy*uy + uz*(2. + uz))))/tau; + + // q = 13 + dist[nr14] = m13 + 0.125*(Fz + Fx*(-1 + ux) + Fy*uy + Fz*uz)* + (-0.2222222222222222 - (ux - uz)*(ux - uz) + 0.3333333333333333*(-2*ux + ux*ux + uy*uy + uz*(2 + uz)))\ + - (m13 - 0.027777777777777776*p + 0.041666666666666664*rho0* + (-((ux - uz)*(ux - uz)) + 0.3333333333333333*(-2*ux + ux*ux + uy*uy + uz*(2 + uz))) + + 0.0625*(Fz + Fx*(-1. + ux) + Fy*uy + Fz*uz)* + (-0.2222222222222222 - 1.*(ux*ux) + 2.*ux*uz - 1.*(uz*uz) + + 0.3333333333333333*(-2.*ux + ux*ux + uy*uy + uz*(2. + uz))))/tau; + + // q= 14 + dist[nr13] = m14 + 0.125*(Fx*(1 + ux) + Fy*uy + Fz*(-1 + uz))* + (-0.2222222222222222 - (ux - uz)*(ux - uz) + 0.3333333333333333*(2*ux + ux*ux + uy*uy + (-2 + uz)*uz))\ + - (m14 - 0.027777777777777776*p + 0.041666666666666664*rho0* + (-((ux - uz)*(ux - uz)) + 0.3333333333333333*(2*ux + ux*ux + uy*uy + (-2 + uz)*uz)) + + 0.0625*(Fx*(1 + ux) + Fy*uy + Fz*(-1. + uz))* + (-0.2222222222222222 - 1.*(ux*ux) + 2.*ux*uz - 1.*(uz*uz) + + 0.3333333333333333*(2.*ux + ux*ux + uy*uy + (-2. + uz)*uz)))/tau; + + // q = 15 + dist[nr16] = m15 - 0.125*(Fx*ux + Fy*(-1 + uy) + Fz*(-1 + uz))* + (0.2222222222222222 + (uy + uz)*(uy + uz) - 0.3333333333333333*(ux*ux - 2*uy + uy*uy + (-2 + uz)*uz))\ + - (m15 - 0.027777777777777776*p + 0.041666666666666664*rho0* + (-((uy + uz)*(uy + uz)) + 0.3333333333333333*(ux*ux - 2*uy + uy*uy + (-2 + uz)*uz)) + + 0.0625*(Fx*ux + Fy*(-1. + uy) + Fz*(-1. + uz))* + (-0.2222222222222222 - 1.*(uy*uy) - 2.*uy*uz - 1.*(uz*uz) + + 0.3333333333333333*(ux*ux - 2.*uy + uy*uy + (-2. + uz)*uz)))/tau; + + // q = 16 + dist[nr15] = m16 + 0.125*(Fy + Fz + Fx*ux + Fy*uy + Fz*uz)* + (-0.2222222222222222 - (uy + uz)*(uy + uz) + 0.3333333333333333*(ux*ux + 2*uy + uy*uy + uz*(2 + uz)))\ + - (m16 - 0.027777777777777776*p + 0.041666666666666664*rho0* + (-((uy + uz)*(uy + uz)) + 0.3333333333333333*(ux*ux + 2*uy + uy*uy + uz*(2 + uz))) + + 0.0625*(Fy + Fz + Fx*ux + Fy*uy + Fz*uz)* + (-0.2222222222222222 - 1.*(uy*uy) - 2.*uy*uz - 1.*(uz*uz) + + 0.3333333333333333*(ux*ux + 2.*uy + uy*uy + uz*(2. + uz))))/tau; + + // q = 17 + dist[nr18] = m17 + 0.125*(Fz + Fx*ux + Fy*(-1 + uy) + Fz*uz)* + (-0.2222222222222222 - (uy - uz)*(uy - uz) + 0.3333333333333333*(ux*ux - 2*uy + uy*uy + uz*(2 + uz)))\ + - (m17 - 0.027777777777777776*p + 0.041666666666666664*rho0* + (-((uy - uz)*(uy - uz)) + 0.3333333333333333*(ux*ux - 2*uy + uy*uy + uz*(2 + uz))) + + 0.0625*(Fz + Fx*ux + Fy*(-1. + uy) + Fz*uz)* + (-0.2222222222222222 - 1.*(uy*uy) + 2.*uy*uz - 1.*(uz*uz) + + 0.3333333333333333*(ux*ux - 2.*uy + uy*uy + uz*(2. + uz))))/tau; + + // q = 18 + dist[nr17] = m18 + 0.125*(Fx*ux + Fy*(1 + uy) + Fz*(-1 + uz))* + (-0.2222222222222222 - (uy - uz)*(uy - uz) + 0.3333333333333333*(ux*ux + 2*uy + uy*uy + (-2 + uz)*uz))\ + - (m18 - 0.027777777777777776*p + 0.041666666666666664*rho0* + (-((uy - uz)*(uy - uz)) + 0.3333333333333333*(ux*ux + 2*uy + uy*uy + (-2 + uz)*uz)) + + 0.0625*(Fx*ux + Fy*(1 + uy) + Fz*(-1. + uz))* + (-0.2222222222222222 - 1.*(uy*uy) + 2.*uy*uz - 1.*(uz*uz) + + 0.3333333333333333*(ux*ux + 2.*uy + uy*uy + (-2. + uz)*uz)))/tau; + //----------------------------------------------------------------------------------------------------------------------------------------// + + + //Update velocity on device + Vel[0*Np+n] = ux; + Vel[1*Np+n] = uy; + Vel[2*Np+n] = uz; + //Update pressure on device + Pressure[n] = p; + } + } +} + +__global__ void dvc_ScaLBL_D3Q19_AAeven_FreeLeeModel_SingleFluid_BGK(double *dist, double *Vel, double *Pressure, + double tau, double rho0, double Fx, double Fy, double Fz, int start, int finish, int Np){ + + int n; + double ux,uy,uz;//fluid velocity + double p;//pressure + // distribution functions + double m1,m2,m4,m6,m8,m9,m10,m11,m12,m13,m14,m15,m16,m17,m18; + double m0,m3,m5,m7; + + // for (int n=start; n