From 4c8ec687d23f01f1fbb429619a4aa6bb2fc093f5 Mon Sep 17 00:00:00 2001
From: Rex Zhe Li <zhe.rex.li@gmail.com>
Date: Wed, 24 Feb 2021 20:28:36 -0500
Subject: [PATCH 01/10] unfinished work: try to add capillary penalty term in
 GreyscaleColor

---
 cuda/GreyscaleColor.cu | 1451 ++++++++++++++++++++++++++++++++++++++++
 1 file changed, 1451 insertions(+)

diff --git a/cuda/GreyscaleColor.cu b/cuda/GreyscaleColor.cu
index 58dfa311..0cc49bcd 100644
--- a/cuda/GreyscaleColor.cu
+++ b/cuda/GreyscaleColor.cu
@@ -1447,6 +1447,1457 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor(int *Map, double *dist,
 	}
 }
 
+//CP: capillary penalty
+// also turn off recoloring for grey nodes
+__global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int *Map, double *dist, double *Aq, double *Bq, double *Den,
+		 double *Phi, double *GreySolidGrad, double *Poros,double *Perm, double *Velocity, double *Pressure,
+         double rhoA, double rhoB, double tauA, double tauB,double tauA_eff,double tauB_eff,double alpha, double beta,
+		double Gx, double Gy, double Gz, bool RecoloringOff, int strideY, int strideZ, int start, int finish, int Np){
+
+	int n,nn,ijk,nread;
+	int nr1,nr2,nr3,nr4,nr5,nr6;
+	int nr7,nr8,nr9,nr10;
+	int nr11,nr12,nr13,nr14;
+	//int nr15,nr16,nr17,nr18;
+	double fq;
+	// conserved momemnts
+	double rho,jx,jy,jz;
+	//double vx,vy,vz,v_mag;
+    //double ux,uy,uz,u_mag;
+    double ux,uy,uz;
+	// non-conserved moments
+	double m1,m2,m4,m6,m8,m9,m10,m11,m12,m13,m14,m15,m16,m17,m18;
+	double m3,m5,m7;
+	double nA,nB; // number density
+	double a1,b1,a2,b2,nAB,delta;
+	double C,nx,ny,nz; //color gradient magnitude and direction
+	double phi,tau,rho0,rlx_setA,rlx_setB;
+    double Fcpx,Fcpy,Fcpz;
+
+    //double GeoFun=0.0;//geometric function from Guo's PRE 66, 036304 (2002)
+    double porosity;
+    double perm;//voxel permeability
+    double c0, c1; //Guo's model parameters
+    double tau_eff;
+    double mu_eff;//kinematic viscosity
+    double nx_gs,ny_gs,nz_gs;//grey-solid color gradient
+    double nx_phase,ny_phase,nz_phase,C_phase;
+    double Fx,Fy,Fz;
+
+	const double mrt_V1=0.05263157894736842;
+	const double mrt_V2=0.012531328320802;
+	const double mrt_V3=0.04761904761904762;
+	const double mrt_V4=0.004594820384294068;
+	const double mrt_V5=0.01587301587301587;
+	const double mrt_V6=0.0555555555555555555555555;
+	const double mrt_V7=0.02777777777777778;
+	const double mrt_V8=0.08333333333333333;
+	const double mrt_V9=0.003341687552213868;
+	const double mrt_V10=0.003968253968253968;
+	const double mrt_V11=0.01388888888888889;
+	const double mrt_V12=0.04166666666666666;
+
+	int S = Np/NBLOCKS/NTHREADS + 1;
+	for (int s=0; s<S; s++){
+		//........Get 1-D index for this thread....................
+		n =  S*blockIdx.x*blockDim.x + s*blockDim.x + threadIdx.x + start;
+		if (n<finish) {
+			// read the component number densities
+			nA = Den[n];
+			nB = Den[Np + n];
+
+            porosity = Poros[n];
+            perm = Perm[n];
+            nx_gs = GreySolidGrad[n+0*Np];
+            ny_gs = GreySolidGrad[n+1*Np];
+            nz_gs = GreySolidGrad[n+2*Np];
+
+			// compute phase indicator field
+			phi=(nA-nB)/(nA+nB);
+
+			// local density
+			rho0=rhoA + 0.5*(1.0-phi)*(rhoB-rhoA);
+			// local relaxation time
+			tau=tauA + 0.5*(1.0-phi)*(tauB-tauA);
+			tau_eff=tauA_eff + 0.5*(1.0-phi)*(tauB_eff-tauA_eff);
+			rlx_setA = 1.f/tau;
+			rlx_setB = 8.f*(2.f-rlx_setA)/(8.f-rlx_setA);
+            mu_eff = (tau_eff-0.5)/3.0;//kinematic viscosity
+			
+			// Get the 1D index based on regular data layout
+			ijk = Map[n];
+			//					COMPUTE THE COLOR GRADIENT
+			//........................................................................
+			//.................Read Phase Indicator Values............................
+			//........................................................................
+			nn = ijk-1;							// neighbor index (get convention)
+			m1 = Phi[nn];						// get neighbor for phi - 1
+			//........................................................................
+			nn = ijk+1;							// neighbor index (get convention)
+			m2 = Phi[nn];						// get neighbor for phi - 2
+			//........................................................................
+			nn = ijk-strideY;							// neighbor index (get convention)
+			m3 = Phi[nn];					// get neighbor for phi - 3
+			//........................................................................
+			nn = ijk+strideY;							// neighbor index (get convention)
+			m4 = Phi[nn];					// get neighbor for phi - 4
+			//........................................................................
+			nn = ijk-strideZ;						// neighbor index (get convention)
+			m5 = Phi[nn];					// get neighbor for phi - 5
+			//........................................................................
+			nn = ijk+strideZ;						// neighbor index (get convention)
+			m6 = Phi[nn];					// get neighbor for phi - 6
+			//........................................................................
+			nn = ijk-strideY-1;						// neighbor index (get convention)
+			m7 = Phi[nn];					// get neighbor for phi - 7
+			//........................................................................
+			nn = ijk+strideY+1;						// neighbor index (get convention)
+			m8 = Phi[nn];					// get neighbor for phi - 8
+			//........................................................................
+			nn = ijk+strideY-1;						// neighbor index (get convention)
+			m9 = Phi[nn];					// get neighbor for phi - 9
+			//........................................................................
+			nn = ijk-strideY+1;						// neighbor index (get convention)
+			m10 = Phi[nn];					// get neighbor for phi - 10
+			//........................................................................
+			nn = ijk-strideZ-1;						// neighbor index (get convention)
+			m11 = Phi[nn];					// get neighbor for phi - 11
+			//........................................................................
+			nn = ijk+strideZ+1;						// neighbor index (get convention)
+			m12 = Phi[nn];					// get neighbor for phi - 12
+			//........................................................................
+			nn = ijk+strideZ-1;						// neighbor index (get convention)
+			m13 = Phi[nn];					// get neighbor for phi - 13
+			//........................................................................
+			nn = ijk-strideZ+1;						// neighbor index (get convention)
+			m14 = Phi[nn];					// get neighbor for phi - 14
+			//........................................................................
+			nn = ijk-strideZ-strideY;					// neighbor index (get convention)
+			m15 = Phi[nn];					// get neighbor for phi - 15
+			//........................................................................
+			nn = ijk+strideZ+strideY;					// neighbor index (get convention)
+			m16 = Phi[nn];					// get neighbor for phi - 16
+			//........................................................................
+			nn = ijk+strideZ-strideY;					// neighbor index (get convention)
+			m17 = Phi[nn];					// get neighbor for phi - 17
+			//........................................................................
+			nn = ijk-strideZ+strideY;					// neighbor index (get convention)
+			m18 = Phi[nn];					// get neighbor for phi - 18
+			//............Compute the Color Gradient...................................
+			nx_phase = -(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
+			ny_phase = -(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
+			nz_phase = -(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
+			C_phase = sqrt(nx_phase*nx_phase+ny_phase*ny_phase+nz_phase*nz_phase);
+
+            //correct the normal color gradient by considering the effect of grey solid
+            nx = nx_phase + (1.0-porosity)*nx_gs; 
+            ny = ny_phase + (1.0-porosity)*ny_gs; 
+            nz = nz_phase + (1.0-porosity)*nz_gs; 
+            if (C_phase==0.0){
+                nx = nx_phase; 
+                ny = ny_phase;
+                nz = nz_phase;
+            }
+
+			//...........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;		
+            //----------- Introduce capillary penalty force -------------------------
+            //NOTE: apply only to grey nodes
+            Fcpx = 0.5*alpha*C*W/sqrt(perm)*nx;
+            Fcpy = 0.5*alpha*C*W/sqrt(perm)*ny;
+            Fcpz = 0.5*alpha*C*W/sqrt(perm)*nz;
+            if (porosity==1.0){
+                Fcpx = 0.0;
+                Fcpy = 0.0;
+                Fcpz = 0.0;
+            }
+
+			// q=0
+			fq = dist[n];
+			rho = fq;
+			m1  = -30.0*fq;
+			m2  = 12.0*fq;
+
+			// q=1
+			//nread = neighborList[n]; // neighbor 2 
+			//fq = dist[nread]; // reading the f1 data into register fq		
+			nr1 = neighborList[n]; 
+			fq = dist[nr1]; // reading the f1 data into register fq
+			rho += fq;
+			m1 -= 11.0*fq;
+			m2 -= 4.0*fq;
+			jx = fq;
+			m4 = -4.0*fq;
+			m9 = 2.0*fq;
+			m10 = -4.0*fq;
+
+			// f2 = dist[10*Np+n];
+			//nread = neighborList[n+Np]; // neighbor 1 ( < 10Np => even part of dist)
+			//fq = dist[nread];  // reading the f2 data into register fq
+			nr2 = neighborList[n+Np]; // neighbor 1 ( < 10Np => even part of dist)
+			fq = dist[nr2];  // reading the f2 data into register fq
+			rho += fq;
+			m1 -= 11.0*(fq);
+			m2 -= 4.0*(fq);
+			jx -= fq;
+			m4 += 4.0*(fq);
+			m9 += 2.0*(fq);
+			m10 -= 4.0*(fq);
+
+			// q=3
+			//nread = neighborList[n+2*Np]; // neighbor 4
+			//fq = dist[nread];
+			nr3 = neighborList[n+2*Np]; // neighbor 4
+			fq = dist[nr3];
+			rho += fq;
+			m1 -= 11.0*fq;
+			m2 -= 4.0*fq;
+			jy = fq;
+			m6 = -4.0*fq;
+			m9 -= fq;
+			m10 += 2.0*fq;
+			m11 = fq;
+			m12 = -2.0*fq;
+
+			// q = 4
+			//nread = neighborList[n+3*Np]; // neighbor 3
+			//fq = dist[nread];
+			nr4 = neighborList[n+3*Np]; // neighbor 3
+			fq = dist[nr4];
+			rho+= fq;
+			m1 -= 11.0*fq;
+			m2 -= 4.0*fq;
+			jy -= fq;
+			m6 += 4.0*fq;
+			m9 -= fq;
+			m10 += 2.0*fq;
+			m11 += fq;
+			m12 -= 2.0*fq;
+
+			// q=5
+			//nread = neighborList[n+4*Np];
+			//fq = dist[nread];
+			nr5 = neighborList[n+4*Np];
+			fq = dist[nr5];
+			rho += fq;
+			m1 -= 11.0*fq;
+			m2 -= 4.0*fq;
+			jz = fq;
+			m8 = -4.0*fq;
+			m9 -= fq;
+			m10 += 2.0*fq;
+			m11 -= fq;
+			m12 += 2.0*fq;
+
+
+			// q = 6
+			//nread = neighborList[n+5*Np];
+			//fq = dist[nread];
+			nr6 = neighborList[n+5*Np];
+			fq = dist[nr6];
+			rho+= fq;
+			m1 -= 11.0*fq;
+			m2 -= 4.0*fq;
+			jz -= fq;
+			m8 += 4.0*fq;
+			m9 -= fq;
+			m10 += 2.0*fq;
+			m11 -= fq;
+			m12 += 2.0*fq;
+
+			// q=7
+			//nread = neighborList[n+6*Np];
+			//fq = dist[nread];
+			nr7 = neighborList[n+6*Np];
+			fq = dist[nr7];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jx += fq;
+			m4 += fq;
+			jy += fq;
+			m6 += fq;
+			m9  += fq;
+			m10 += fq;
+			m11 += fq;
+			m12 += fq;
+			m13 = fq;
+			m16 = fq;
+			m17 = -fq;
+
+			// q = 8
+			//nread = neighborList[n+7*Np];
+			//fq = dist[nread];
+			nr8 = neighborList[n+7*Np];
+			fq = dist[nr8];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jx -= fq;
+			m4 -= fq;
+			jy -= fq;
+			m6 -= fq;
+			m9 += fq;
+			m10 += fq;
+			m11 += fq;
+			m12 += fq;
+			m13 += fq;
+			m16 -= fq;
+			m17 += fq;
+
+			// q=9
+			//nread = neighborList[n+8*Np];
+			//fq = dist[nread];
+			nr9 = neighborList[n+8*Np];
+			fq = dist[nr9];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jx += fq;
+			m4 += fq;
+			jy -= fq;
+			m6 -= fq;
+			m9 += fq;
+			m10 += fq;
+			m11 += fq;
+			m12 += fq;
+			m13 -= fq;
+			m16 += fq;
+			m17 += fq;
+
+			// q = 10
+			//nread = neighborList[n+9*Np];
+			//fq = dist[nread];
+			nr10 = neighborList[n+9*Np];
+			fq = dist[nr10];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jx -= fq;
+			m4 -= fq;
+			jy += fq;
+			m6 += fq;
+			m9 += fq;
+			m10 += fq;
+			m11 += fq;
+			m12 += fq;
+			m13 -= fq;
+			m16 -= fq;
+			m17 -= fq;
+
+			// q=11
+			//nread = neighborList[n+10*Np];
+			//fq = dist[nread];
+			nr11 = neighborList[n+10*Np];
+			fq = dist[nr11];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jx += fq;
+			m4 += fq;
+			jz += fq;
+			m8 += fq;
+			m9 += fq;
+			m10 += fq;
+			m11 -= fq;
+			m12 -= fq;
+			m15 = fq;
+			m16 -= fq;
+			m18 = fq;
+
+			// q=12
+			//nread = neighborList[n+11*Np];
+			//fq = dist[nread];
+			nr12 = neighborList[n+11*Np];
+			fq = dist[nr12];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jx -= fq;
+			m4 -= fq;
+			jz -= fq;
+			m8 -= fq;
+			m9 += fq;
+			m10 += fq;
+			m11 -= fq;
+			m12 -= fq;
+			m15 += fq;
+			m16 += fq;
+			m18 -= fq;
+
+			// q=13
+			//nread = neighborList[n+12*Np];
+			//fq = dist[nread];
+			nr13 = neighborList[n+12*Np];
+			fq = dist[nr13];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jx += fq;
+			m4 += fq;
+			jz -= fq;
+			m8 -= fq;
+			m9 += fq;
+			m10 += fq;
+			m11 -= fq;
+			m12 -= fq;
+			m15 -= fq;
+			m16 -= fq;
+			m18 -= fq;
+
+			// q=14
+			//nread = neighborList[n+13*Np];
+			//fq = dist[nread];
+			nr14 = neighborList[n+13*Np];
+			fq = dist[nr14];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jx -= fq;
+			m4 -= fq;
+			jz += fq;
+			m8 += fq;
+			m9 += fq;
+			m10 += fq;
+			m11 -= fq;
+			m12 -= fq;
+			m15 -= fq;
+			m16 += fq;
+			m18 += fq;
+
+			// q=15
+			nread = neighborList[n+14*Np];
+			fq = dist[nread];
+			//fq = dist[17*Np+n];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jy += fq;
+			m6 += fq;
+			jz += fq;
+			m8 += fq;
+			m9 -= 2.0*fq;
+			m10 -= 2.0*fq;
+			m14 = fq;
+			m17 += fq;
+			m18 -= fq;
+
+			// q=16
+			nread = neighborList[n+15*Np];
+			fq = dist[nread];
+			//fq = dist[8*Np+n];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jy -= fq;
+			m6 -= fq;
+			jz -= fq;
+			m8 -= fq;
+			m9 -= 2.0*fq;
+			m10 -= 2.0*fq;
+			m14 += fq;
+			m17 -= fq;
+			m18 += fq;
+
+			// q=17
+			//fq = dist[18*Np+n];
+			nread = neighborList[n+16*Np];
+			fq = dist[nread];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jy += fq;
+			m6 += fq;
+			jz -= fq;
+			m8 -= fq;
+			m9 -= 2.0*fq;
+			m10 -= 2.0*fq;
+			m14 -= fq;
+			m17 += fq;
+			m18 += fq;
+
+			// q=18
+			nread = neighborList[n+17*Np];
+			fq = dist[nread];
+			//fq = dist[9*Np+n];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jy -= fq;
+			m6 -= fq;
+			jz += fq;
+			m8 += fq;
+			m9 -= 2.0*fq;
+			m10 -= 2.0*fq;
+			m14 -= fq;
+			m17 -= fq;
+			m18 -= fq;
+			
+            //// Compute greyscale related parameters
+            ux = (jx/rho0+0.5*porosity*Gx+0.5*Fcpx/rho0)/(1.0+0.5*porosity*mu_eff/perm);
+            uy = (jy/rho0+0.5*porosity*Gy+0.5*Fcpy/rho0)/(1.0+0.5*porosity*mu_eff/perm);
+            uz = (jz/rho0+0.5*porosity*Gz+0.5*Fcpz/rho0)/(1.0+0.5*porosity*mu_eff/perm);
+
+            //Update the total force to include linear (Darcy) and nonlinear (Forchheimer) drags due to the porous medium
+            Fx = rho0*(-porosity*mu_eff/perm*ux + porosity*Gx)+Fcpx;
+            Fy = rho0*(-porosity*mu_eff/perm*uy + porosity*Gy)+Fcpy;
+            Fz = rho0*(-porosity*mu_eff/perm*uz + porosity*Gz)+Fcpz;
+            if (porosity==1.0){
+                Fx=rho0*(porosity*Gx);
+                Fy=rho0*(porosity*Gy);
+                Fz=rho0*(porosity*Gz);
+            }
+
+			// write the velocity 
+			Velocity[n] = ux;
+			Velocity[Np+n] = uy;
+			Velocity[2*Np+n] = uz;
+            //Pressure[n] = rho/3.f/porosity;
+            Pressure[n] = rho/3.f;
+
+			//........................................................................
+			//..............carry out relaxation process..............................
+			//..........Toelke, Fruediger et. al. 2006................................
+            //---------------- NO higher-order force -------------------------------//
+			if (C == 0.0)	nx = ny = nz = 0.0;
+			m1 = m1 + rlx_setA*((19*(ux*ux+uy*uy+uz*uz)*rho0/porosity - 11*rho) -19*alpha*C - m1);
+			m2 = m2 + rlx_setA*((3*rho - 5.5*(ux*ux+uy*uy+uz*uz)*rho0/porosity)- m2);
+            jx = jx + Fx;
+			m4 = m4 + rlx_setB*((-0.6666666666666666*ux*rho0)- m4)
+                    + (1-0.5*rlx_setB)*(-0.6666666666666666*Fx);
+            jy = jy + Fy;
+			m6 = m6 + rlx_setB*((-0.6666666666666666*uy*rho0)- m6)
+                    + (1-0.5*rlx_setB)*(-0.6666666666666666*Fy);
+            jz = jz + Fz;
+			m8 = m8 + rlx_setB*((-0.6666666666666666*uz*rho0)- m8)
+                    + (1-0.5*rlx_setB)*(-0.6666666666666666*Fz);
+			m9 = m9 + rlx_setA*(((2*ux*ux-uy*uy-uz*uz)*rho0/porosity) + 0.5*alpha*C*(2*nx*nx-ny*ny-nz*nz) - m9);
+			m10 = m10 + rlx_setA*( - m10);
+            //m10 = m10 + rlx_setA*(-0.5*rho0*((2*ux*ux-uy*uy-uz*uz)/porosity)- m10);
+			m11 = m11 + rlx_setA*(((uy*uy-uz*uz)*rho0/porosity) + 0.5*alpha*C*(ny*ny-nz*nz)- m11);
+			m12 = m12 + rlx_setA*( - m12);
+            //m12 = m12 + rlx_setA*(-0.5*(rho0*(uy*uy-uz*uz)/porosity)- m12);
+			m13 = m13 + rlx_setA*( (ux*uy*rho0/porosity) + 0.5*alpha*C*nx*ny - m13);
+			m14 = m14 + rlx_setA*( (uy*uz*rho0/porosity) + 0.5*alpha*C*ny*nz - m14);
+			m15 = m15 + rlx_setA*( (ux*uz*rho0/porosity) + 0.5*alpha*C*nx*nz - m15);
+			m16 = m16 + rlx_setB*( - m16);
+			m17 = m17 + rlx_setB*( - m17);
+			m18 = m18 + rlx_setB*( - m18);
+            //----------------------------------------------------------------------//
+
+            //----------------With higher-order force ------------------------------//
+			//if (C == 0.0)	nx = ny = nz = 0.0;
+			//m1 = m1 + rlx_setA*((19*(ux*ux+uy*uy+uz*uz)*rho0/porosity - 11*rho) -19*alpha*C - m1)
+            //        + (1-0.5*rlx_setA)*38*(Fx*ux+Fy*uy+Fz*uz)/porosity;
+			//m2 = m2 + rlx_setA*((3*rho - 5.5*(ux*ux+uy*uy+uz*uz)*rho0/porosity)- m2)
+            //        + (1-0.5*rlx_setA)*11*(-Fx*ux-Fy*uy-Fz*uz)/porosity;
+            //jx = jx + Fx;
+			//m4 = m4 + rlx_setB*((-0.6666666666666666*ux*rho0)- m4)
+            //        + (1-0.5*rlx_setB)*(-0.6666666666666666*Fx);
+            //jy = jy + Fy;
+			//m6 = m6 + rlx_setB*((-0.6666666666666666*uy*rho0)- m6)
+            //        + (1-0.5*rlx_setB)*(-0.6666666666666666*Fy);
+            //jz = jz + Fz;
+			//m8 = m8 + rlx_setB*((-0.6666666666666666*uz*rho0)- m8)
+            //        + (1-0.5*rlx_setB)*(-0.6666666666666666*Fz);
+			//m9 = m9 + rlx_setA*(((2*ux*ux-uy*uy-uz*uz)*rho0/porosity) + 0.5*alpha*C*(2*nx*nx-ny*ny-nz*nz) - m9)
+            //        + (1-0.5*rlx_setA)*(4*Fx*ux-2*Fy*uy-2*Fz*uz)/porosity;
+			////m10 = m10 + rlx_setA*( - m10);
+            //m10 = m10 + rlx_setA*(-0.5*rho0*((2*ux*ux-uy*uy-uz*uz)/porosity)- m10)
+            //          + (1-0.5*rlx_setA)*(-2*Fx*ux+Fy*uy+Fz*uz)/porosity;
+			//m11 = m11 + rlx_setA*(((uy*uy-uz*uz)*rho0/porosity) + 0.5*alpha*C*(ny*ny-nz*nz)- m11)
+            //          + (1-0.5*rlx_setA)*(2*Fy*uy-2*Fz*uz)/porosity;
+			////m12 = m12 + rlx_setA*( - m12);
+            //m12 = m12 + rlx_setA*(-0.5*(rho0*(uy*uy-uz*uz)/porosity)- m12)
+            //          + (1-0.5*rlx_setA)*(-Fy*uy+Fz*uz)/porosity;
+			//m13 = m13 + rlx_setA*( (ux*uy*rho0/porosity) + 0.5*alpha*C*nx*ny - m13);
+            //          + (1-0.5*rlx_setA)*(Fy*ux+Fx*uy)/porosity;
+			//m14 = m14 + rlx_setA*( (uy*uz*rho0/porosity) + 0.5*alpha*C*ny*nz - m14);
+            //          + (1-0.5*rlx_setA)*(Fz*uy+Fy*uz)/porosity;
+			//m15 = m15 + rlx_setA*( (ux*uz*rho0/porosity) + 0.5*alpha*C*nx*nz - m15);
+            //          + (1-0.5*rlx_setA)*(Fz*ux+Fx*uz)/porosity;
+			//m16 = m16 + rlx_setB*( - m16);
+			//m17 = m17 + rlx_setB*( - m17);
+			//m18 = m18 + rlx_setB*( - m18);
+            //----------------------------------------------------------------------//
+
+			//.................inverse transformation......................................................
+			// q=0
+			fq = mrt_V1*rho-mrt_V2*m1+mrt_V3*m2;
+			dist[n] = fq;
+
+			// q = 1
+			fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(jx-m4)+mrt_V6*(m9-m10);
+			//nread = neighborList[n+Np];
+			dist[nr2] = fq;
+
+			// q=2
+			fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(m4-jx)+mrt_V6*(m9-m10);
+			//nread = neighborList[n];
+			dist[nr1] = fq;
+
+			// q = 3
+			fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(jy-m6)+mrt_V7*(m10-m9)+mrt_V8*(m11-m12);
+			//nread = neighborList[n+3*Np];
+			dist[nr4] = fq;
+
+			// q = 4
+			fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(m6-jy)+mrt_V7*(m10-m9)+mrt_V8*(m11-m12);
+			//nread = neighborList[n+2*Np];
+			dist[nr3] = fq;
+
+			// q = 5
+			fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(jz-m8)+mrt_V7*(m10-m9)+mrt_V8*(m12-m11);
+			//nread = neighborList[n+5*Np];
+			dist[nr6] = fq;
+
+			// q = 6
+			fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(m8-jz)+mrt_V7*(m10-m9)+mrt_V8*(m12-m11);
+			//nread = neighborList[n+4*Np];
+			dist[nr5] = fq;
+
+			// q = 7
+			fq = mrt_V1*rho+mrt_V9*m1+mrt_V10*m2+0.1*(jx+jy)+0.025*(m4+m6)+
+					mrt_V7*m9+mrt_V11*m10+mrt_V8*m11+mrt_V12*m12+0.25*m13+0.125*(m16-m17);
+			//nread = neighborList[n+7*Np];
+			dist[nr8] = fq;
+
+			// q = 8
+			fq = mrt_V1*rho+mrt_V9*m1+mrt_V10*m2-0.1*(jx+jy)-0.025*(m4+m6) +mrt_V7*m9+mrt_V11*m10+mrt_V8*m11
+					+mrt_V12*m12+0.25*m13+0.125*(m17-m16);
+			//nread = neighborList[n+6*Np];
+			dist[nr7] = fq;
+
+			// q = 9
+			fq = mrt_V1*rho+mrt_V9*m1+mrt_V10*m2+0.1*(jx-jy)+0.025*(m4-m6)+
+					mrt_V7*m9+mrt_V11*m10+mrt_V8*m11+mrt_V12*m12-0.25*m13+0.125*(m16+m17);
+			//nread = neighborList[n+9*Np];
+			dist[nr10] = fq;
+
+			// q = 10
+			fq = mrt_V1*rho+mrt_V9*m1+mrt_V10*m2+0.1*(jy-jx)+0.025*(m6-m4)+
+					mrt_V7*m9+mrt_V11*m10+mrt_V8*m11+mrt_V12*m12-0.25*m13-0.125*(m16+m17);
+			//nread = neighborList[n+8*Np];
+			dist[nr9] = fq;
+
+			// q = 11
+			fq = mrt_V1*rho+mrt_V9*m1
+					+mrt_V10*m2+0.1*(jx+jz)+0.025*(m4+m8)
+					+mrt_V7*m9+mrt_V11*m10-mrt_V8*m11
+					-mrt_V12*m12+0.25*m15+0.125*(m18-m16);
+			//nread = neighborList[n+11*Np];
+			dist[nr12] = fq;
+
+			// q = 12
+			fq = mrt_V1*rho+mrt_V9*m1+mrt_V10*m2-0.1*(jx+jz)-0.025*(m4+m8)+
+					mrt_V7*m9+mrt_V11*m10-mrt_V8*m11-mrt_V12*m12+0.25*m15+0.125*(m16-m18);
+			//nread = neighborList[n+10*Np];
+			dist[nr11]= fq;
+
+			// q = 13
+			fq = mrt_V1*rho+mrt_V9*m1
+					+mrt_V10*m2+0.1*(jx-jz)+0.025*(m4-m8)
+					+mrt_V7*m9+mrt_V11*m10-mrt_V8*m11
+					-mrt_V12*m12-0.25*m15-0.125*(m16+m18);
+			//nread = neighborList[n+13*Np];
+			dist[nr14] = fq;
+
+			// q= 14
+			fq = mrt_V1*rho+mrt_V9*m1
+					+mrt_V10*m2+0.1*(jz-jx)+0.025*(m8-m4)
+					+mrt_V7*m9+mrt_V11*m10-mrt_V8*m11
+					-mrt_V12*m12-0.25*m15+0.125*(m16+m18);
+			//nread = neighborList[n+12*Np];
+			dist[nr13] = fq;
+
+
+			// q = 15
+			fq = mrt_V1*rho+mrt_V9*m1
+					+mrt_V10*m2+0.1*(jy+jz)+0.025*(m6+m8)
+					-mrt_V6*m9-mrt_V7*m10+0.25*m14+0.125*(m17-m18);
+			nread = neighborList[n+15*Np];
+			dist[nread] = fq;
+
+			// q = 16
+			fq =  mrt_V1*rho+mrt_V9*m1
+					+mrt_V10*m2-0.1*(jy+jz)-0.025*(m6+m8)
+					-mrt_V6*m9-mrt_V7*m10+0.25*m14+0.125*(m18-m17);
+			nread = neighborList[n+14*Np];
+			dist[nread] = fq;
+
+
+			// q = 17
+			fq = mrt_V1*rho+mrt_V9*m1
+					+mrt_V10*m2+0.1*(jy-jz)+0.025*(m6-m8)
+					-mrt_V6*m9-mrt_V7*m10-0.25*m14+0.125*(m17+m18);
+			nread = neighborList[n+17*Np];
+			dist[nread] = fq;
+
+			// q = 18
+			fq = mrt_V1*rho+mrt_V9*m1
+					+mrt_V10*m2+0.1*(jz-jy)+0.025*(m8-m6)
+					-mrt_V6*m9-mrt_V7*m10-0.25*m14-0.125*(m17+m18);
+			nread = neighborList[n+16*Np];
+			dist[nread] = fq;
+			//........................................................................
+
+			// Instantiate mass transport distributions
+			// Stationary value - distribution 0
+			nAB = 1.0/(nA+nB);
+			Aq[n] = 0.3333333333333333*nA;
+			Bq[n] = 0.3333333333333333*nB;
+
+			//...............................................
+			// q = 0,2,4
+			// Cq = {1,0,0}, {0,1,0}, {0,0,1}
+			delta = beta*nA*nB*nAB*0.1111111111111111*nx;
+			if (!(nA*nB*nAB>0)) delta=0;
+            if (RecoloringOff==true && porosity !=1.0) delta=0;
+			a1 = nA*(0.1111111111111111*(1+4.5*ux))+delta;
+			b1 = nB*(0.1111111111111111*(1+4.5*ux))-delta;
+			a2 = nA*(0.1111111111111111*(1-4.5*ux))-delta;
+			b2 = nB*(0.1111111111111111*(1-4.5*ux))+delta;
+
+			// q = 1
+			//nread = neighborList[n+Np];
+			Aq[nr2] = a1;
+			Bq[nr2] = b1;
+			// q=2
+			//nread = neighborList[n];
+			Aq[nr1] = a2;
+			Bq[nr1] = b2;
+
+			//...............................................
+			// Cq = {0,1,0}
+			delta = beta*nA*nB*nAB*0.1111111111111111*ny;
+			if (!(nA*nB*nAB>0)) delta=0;
+            if (RecoloringOff==true && porosity !=1.0) delta=0;
+			a1 = nA*(0.1111111111111111*(1+4.5*uy))+delta;
+			b1 = nB*(0.1111111111111111*(1+4.5*uy))-delta;
+			a2 = nA*(0.1111111111111111*(1-4.5*uy))-delta;
+			b2 = nB*(0.1111111111111111*(1-4.5*uy))+delta;
+
+			// q = 3
+			//nread = neighborList[n+3*Np];
+			Aq[nr4] = a1;
+			Bq[nr4] = b1;
+			// q = 4
+			//nread = neighborList[n+2*Np];
+			Aq[nr3] = a2;
+			Bq[nr3] = b2;
+
+			//...............................................
+			// q = 4
+			// Cq = {0,0,1}
+			delta = beta*nA*nB*nAB*0.1111111111111111*nz;
+			if (!(nA*nB*nAB>0)) delta=0;
+            if (RecoloringOff==true && porosity !=1.0) delta=0;
+			a1 = nA*(0.1111111111111111*(1+4.5*uz))+delta;
+			b1 = nB*(0.1111111111111111*(1+4.5*uz))-delta;
+			a2 = nA*(0.1111111111111111*(1-4.5*uz))-delta;
+			b2 = nB*(0.1111111111111111*(1-4.5*uz))+delta;
+
+			// q = 5
+			//nread = neighborList[n+5*Np];
+			Aq[nr6] = a1;
+			Bq[nr6] = b1;
+			// q = 6
+			//nread = neighborList[n+4*Np];
+			Aq[nr5] = a2;
+			Bq[nr5] = b2;
+			//...............................................
+		}
+	}
+}
+
+//CP: capillary penalty
+// also turn off recoloring for grey nodes
+__global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dist, double *Aq, double *Bq, double *Den, 
+        double *Phi, double *GreySolidGrad, double *Poros,double *Perm, double *Velocity, double *Pressure, 
+        double rhoA, double rhoB, double tauA, double tauB,double tauA_eff,double tauB_eff, double alpha, double beta,
+		double Gx, double Gy, double Gz, bool RecoloringOff, int strideY, int strideZ, int start, int finish, int Np){
+	int ijk,nn,n;
+	double fq;
+	// conserved momemnts
+	double rho,jx,jy,jz;
+	//double vx,vy,vz,v_mag;
+    //double ux,uy,uz,u_mag;
+    double ux,uy,uz;
+	// non-conserved moments
+	double m1,m2,m4,m6,m8,m9,m10,m11,m12,m13,m14,m15,m16,m17,m18;
+	double m3,m5,m7;
+	double nA,nB; // number density
+	double a1,b1,a2,b2,nAB,delta;
+	double C,nx,ny,nz; //color gradient magnitude and direction
+	double phi,tau,rho0,rlx_setA,rlx_setB;
+    double Fcpx,Fcpy,Fcpz;
+
+    //double GeoFun=0.0;//geometric function from Guo's PRE 66, 036304 (2002)
+    double porosity;
+    double perm;//voxel permeability
+    double c0, c1; //Guo's model parameters
+    double tau_eff;
+    double mu_eff;//kinematic viscosity
+    double nx_gs,ny_gs,nz_gs;//grey-solid color gradient
+    double nx_phase,ny_phase,nz_phase,C_phase;
+    double Fx,Fy,Fz;
+
+	const double mrt_V1=0.05263157894736842;
+	const double mrt_V2=0.012531328320802;
+	const double mrt_V3=0.04761904761904762;
+	const double mrt_V4=0.004594820384294068;
+	const double mrt_V5=0.01587301587301587;
+	const double mrt_V6=0.0555555555555555555555555;
+	const double mrt_V7=0.02777777777777778;
+	const double mrt_V8=0.08333333333333333;
+	const double mrt_V9=0.003341687552213868;
+	const double mrt_V10=0.003968253968253968;
+	const double mrt_V11=0.01388888888888889;
+	const double mrt_V12=0.04166666666666666;
+
+	int S = Np/NBLOCKS/NTHREADS + 1;
+	for (int s=0; s<S; s++){
+		//........Get 1-D index for this thread....................
+		n =  S*blockIdx.x*blockDim.x + s*blockDim.x + threadIdx.x + start;
+		if (n<finish) {
+
+			// read the component number densities
+			nA = Den[n];
+			nB = Den[Np + n];
+            porosity = Poros[n];
+            perm = Perm[n];
+            nx_gs = GreySolidGrad[n+0*Np];
+            ny_gs = GreySolidGrad[n+1*Np];
+            nz_gs = GreySolidGrad[n+2*Np];
+
+			// compute phase indicator field
+			phi=(nA-nB)/(nA+nB);
+
+			// local density
+			rho0=rhoA + 0.5*(1.0-phi)*(rhoB-rhoA);
+			// local relaxation time
+			tau=tauA + 0.5*(1.0-phi)*(tauB-tauA);
+			tau_eff=tauA_eff + 0.5*(1.0-phi)*(tauB_eff-tauA_eff);
+			rlx_setA = 1.f/tau;
+			rlx_setB = 8.f*(2.f-rlx_setA)/(8.f-rlx_setA);
+            mu_eff = (tau_eff-0.5)/3.0;//kinematic viscosity
+
+			// Get the 1D index based on regular data layout
+			ijk = Map[n];
+			//					COMPUTE THE COLOR GRADIENT
+			//........................................................................
+			//.................Read Phase Indicator Values............................
+			//........................................................................
+			nn = ijk-1;							// neighbor index (get convention)
+			m1 = Phi[nn];						// get neighbor for phi - 1
+			//........................................................................
+			nn = ijk+1;							// neighbor index (get convention)
+			m2 = Phi[nn];						// get neighbor for phi - 2
+			//........................................................................
+			nn = ijk-strideY;							// neighbor index (get convention)
+			m3 = Phi[nn];					// get neighbor for phi - 3
+			//........................................................................
+			nn = ijk+strideY;							// neighbor index (get convention)
+			m4 = Phi[nn];					// get neighbor for phi - 4
+			//........................................................................
+			nn = ijk-strideZ;						// neighbor index (get convention)
+			m5 = Phi[nn];					// get neighbor for phi - 5
+			//........................................................................
+			nn = ijk+strideZ;						// neighbor index (get convention)
+			m6 = Phi[nn];					// get neighbor for phi - 6
+			//........................................................................
+			nn = ijk-strideY-1;						// neighbor index (get convention)
+			m7 = Phi[nn];					// get neighbor for phi - 7
+			//........................................................................
+			nn = ijk+strideY+1;						// neighbor index (get convention)
+			m8 = Phi[nn];					// get neighbor for phi - 8
+			//........................................................................
+			nn = ijk+strideY-1;						// neighbor index (get convention)
+			m9 = Phi[nn];					// get neighbor for phi - 9
+			//........................................................................
+			nn = ijk-strideY+1;						// neighbor index (get convention)
+			m10 = Phi[nn];					// get neighbor for phi - 10
+			//........................................................................
+			nn = ijk-strideZ-1;						// neighbor index (get convention)
+			m11 = Phi[nn];					// get neighbor for phi - 11
+			//........................................................................
+			nn = ijk+strideZ+1;						// neighbor index (get convention)
+			m12 = Phi[nn];					// get neighbor for phi - 12
+			//........................................................................
+			nn = ijk+strideZ-1;						// neighbor index (get convention)
+			m13 = Phi[nn];					// get neighbor for phi - 13
+			//........................................................................
+			nn = ijk-strideZ+1;						// neighbor index (get convention)
+			m14 = Phi[nn];					// get neighbor for phi - 14
+			//........................................................................
+			nn = ijk-strideZ-strideY;					// neighbor index (get convention)
+			m15 = Phi[nn];					// get neighbor for phi - 15
+			//........................................................................
+			nn = ijk+strideZ+strideY;					// neighbor index (get convention)
+			m16 = Phi[nn];					// get neighbor for phi - 16
+			//........................................................................
+			nn = ijk+strideZ-strideY;					// neighbor index (get convention)
+			m17 = Phi[nn];					// get neighbor for phi - 17
+			//........................................................................
+			nn = ijk-strideZ+strideY;					// neighbor index (get convention)
+			m18 = Phi[nn];					// get neighbor for phi - 18
+			//............Compute the Color Gradient...................................
+			nx_phase = -(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
+			ny_phase = -(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
+			nz_phase = -(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
+			C_phase = sqrt(nx_phase*nx_phase+ny_phase*ny_phase+nz_phase*nz_phase);
+
+            //correct the normal color gradient by considering the effect of grey solid
+            nx = nx_phase + (1.0-porosity)*nx_gs; 
+            ny = ny_phase + (1.0-porosity)*ny_gs; 
+            nz = nz_phase + (1.0-porosity)*nz_gs; 
+            if (C_phase==0.0){
+                nx = nx_phase; 
+                ny = ny_phase;
+                nz = nz_phase;
+            }
+
+			//...........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;		
+
+			// q=0
+			fq = dist[n];
+			rho = fq;
+			m1  = -30.0*fq;
+			m2  = 12.0*fq;
+
+			// q=1
+			fq = dist[2*Np+n];
+			rho += fq;
+			m1 -= 11.0*fq;
+			m2 -= 4.0*fq;
+			jx = fq;
+			m4 = -4.0*fq;
+			m9 = 2.0*fq;
+			m10 = -4.0*fq;
+
+			// f2 = dist[10*Np+n];
+			fq = dist[1*Np+n];
+			rho += fq;
+			m1 -= 11.0*(fq);
+			m2 -= 4.0*(fq);
+			jx -= fq;
+			m4 += 4.0*(fq);
+			m9 += 2.0*(fq);
+			m10 -= 4.0*(fq);
+
+			// q=3
+			fq = dist[4*Np+n];
+			rho += fq;
+			m1 -= 11.0*fq;
+			m2 -= 4.0*fq;
+			jy = fq;
+			m6 = -4.0*fq;
+			m9 -= fq;
+			m10 += 2.0*fq;
+			m11 = fq;
+			m12 = -2.0*fq;
+
+			// q = 4
+			fq = dist[3*Np+n];
+			rho+= fq;
+			m1 -= 11.0*fq;
+			m2 -= 4.0*fq;
+			jy -= fq;
+			m6 += 4.0*fq;
+			m9 -= fq;
+			m10 += 2.0*fq;
+			m11 += fq;
+			m12 -= 2.0*fq;
+
+			// q=5
+			fq = dist[6*Np+n];
+			rho += fq;
+			m1 -= 11.0*fq;
+			m2 -= 4.0*fq;
+			jz = fq;
+			m8 = -4.0*fq;
+			m9 -= fq;
+			m10 += 2.0*fq;
+			m11 -= fq;
+			m12 += 2.0*fq;
+
+			// q = 6
+			fq = dist[5*Np+n];
+			rho+= fq;
+			m1 -= 11.0*fq;
+			m2 -= 4.0*fq;
+			jz -= fq;
+			m8 += 4.0*fq;
+			m9 -= fq;
+			m10 += 2.0*fq;
+			m11 -= fq;
+			m12 += 2.0*fq;
+
+			// q=7
+			fq = dist[8*Np+n];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jx += fq;
+			m4 += fq;
+			jy += fq;
+			m6 += fq;
+			m9  += fq;
+			m10 += fq;
+			m11 += fq;
+			m12 += fq;
+			m13 = fq;
+			m16 = fq;
+			m17 = -fq;
+
+			// q = 8
+			fq = dist[7*Np+n];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jx -= fq;
+			m4 -= fq;
+			jy -= fq;
+			m6 -= fq;
+			m9 += fq;
+			m10 += fq;
+			m11 += fq;
+			m12 += fq;
+			m13 += fq;
+			m16 -= fq;
+			m17 += fq;
+
+			// q=9
+			fq = dist[10*Np+n];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jx += fq;
+			m4 += fq;
+			jy -= fq;
+			m6 -= fq;
+			m9 += fq;
+			m10 += fq;
+			m11 += fq;
+			m12 += fq;
+			m13 -= fq;
+			m16 += fq;
+			m17 += fq;
+
+			// q = 10
+			fq = dist[9*Np+n];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jx -= fq;
+			m4 -= fq;
+			jy += fq;
+			m6 += fq;
+			m9 += fq;
+			m10 += fq;
+			m11 += fq;
+			m12 += fq;
+			m13 -= fq;
+			m16 -= fq;
+			m17 -= fq;
+
+			// q=11
+			fq = dist[12*Np+n];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jx += fq;
+			m4 += fq;
+			jz += fq;
+			m8 += fq;
+			m9 += fq;
+			m10 += fq;
+			m11 -= fq;
+			m12 -= fq;
+			m15 = fq;
+			m16 -= fq;
+			m18 = fq;
+
+			// q=12
+			fq = dist[11*Np+n];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jx -= fq;
+			m4 -= fq;
+			jz -= fq;
+			m8 -= fq;
+			m9 += fq;
+			m10 += fq;
+			m11 -= fq;
+			m12 -= fq;
+			m15 += fq;
+			m16 += fq;
+			m18 -= fq;
+
+			// q=13
+			fq = dist[14*Np+n];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jx += fq;
+			m4 += fq;
+			jz -= fq;
+			m8 -= fq;
+			m9 += fq;
+			m10 += fq;
+			m11 -= fq;
+			m12 -= fq;
+			m15 -= fq;
+			m16 -= fq;
+			m18 -= fq;
+
+			// q=14
+			fq = dist[13*Np+n];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jx -= fq;
+			m4 -= fq;
+			jz += fq;
+			m8 += fq;
+			m9 += fq;
+			m10 += fq;
+			m11 -= fq;
+			m12 -= fq;
+			m15 -= fq;
+			m16 += fq;
+			m18 += fq;
+
+			// q=15
+			fq = dist[16*Np+n];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jy += fq;
+			m6 += fq;
+			jz += fq;
+			m8 += fq;
+			m9 -= 2.0*fq;
+			m10 -= 2.0*fq;
+			m14 = fq;
+			m17 += fq;
+			m18 -= fq;
+
+			// q=16
+			fq = dist[15*Np+n];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jy -= fq;
+			m6 -= fq;
+			jz -= fq;
+			m8 -= fq;
+			m9 -= 2.0*fq;
+			m10 -= 2.0*fq;
+			m14 += fq;
+			m17 -= fq;
+			m18 += fq;
+
+			// q=17
+			fq = dist[18*Np+n];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jy += fq;
+			m6 += fq;
+			jz -= fq;
+			m8 -= fq;
+			m9 -= 2.0*fq;
+			m10 -= 2.0*fq;
+			m14 -= fq;
+			m17 += fq;
+			m18 += fq;
+
+			// q=18
+			fq = dist[17*Np+n];
+			rho += fq;
+			m1 += 8.0*fq;
+			m2 += fq;
+			jy -= fq;
+			m6 -= fq;
+			jz += fq;
+			m8 += fq;
+			m9 -= 2.0*fq;
+			m10 -= 2.0*fq;
+			m14 -= fq;
+			m17 -= fq;
+			m18 -= fq;
+
+            // Compute greyscale related parameters
+            c0 = 0.5*(1.0+porosity*0.5*mu_eff/perm);
+            if (porosity==1.0) c0 = 0.5;//i.e. apparent pore nodes
+            //GeoFun = 1.75/sqrt(150.0*porosity*porosity*porosity);
+            c1 = porosity*0.5*GeoFun/sqrt(perm);
+            if (porosity==1.0) c1 = 0.0;//i.e. apparent pore nodes
+
+            vx = jx/rho0+0.5*(porosity*Gx);
+            vy = jy/rho0+0.5*(porosity*Gy);
+            vz = jz/rho0+0.5*(porosity*Gz);
+            v_mag=sqrt(vx*vx+vy*vy+vz*vz);
+            ux = vx/(c0+sqrt(c0*c0+c1*v_mag));
+            uy = vy/(c0+sqrt(c0*c0+c1*v_mag));
+            uz = vz/(c0+sqrt(c0*c0+c1*v_mag));
+            u_mag=sqrt(ux*ux+uy*uy+uz*uz);
+
+            //Update the total force to include linear (Darcy) and nonlinear (Forchheimer) drags due to the porous medium
+            Fx = rho0*(-porosity*mu_eff/perm*ux - porosity*GeoFun/sqrt(perm)*u_mag*ux + porosity*Gx);
+            Fy = rho0*(-porosity*mu_eff/perm*uy - porosity*GeoFun/sqrt(perm)*u_mag*uy + porosity*Gy);
+            Fz = rho0*(-porosity*mu_eff/perm*uz - porosity*GeoFun/sqrt(perm)*u_mag*uz + porosity*Gz);
+            if (porosity==1.0){
+                Fx=rho0*(Gx);
+                Fy=rho0*(Gy);
+                Fz=rho0*(Gz);
+            }
+
+			// write the velocity 
+			Velocity[n] = ux;
+			Velocity[Np+n] = uy;
+			Velocity[2*Np+n] = uz;
+            //Pressure[n] = rho/3.f/porosity;
+            Pressure[n] = rho/3.f;
+
+			//........................................................................
+			//..............carry out relaxation process..............................
+			//..........Toelke, Fruediger et. al. 2006................................
+            //---------------- NO higher-order force -------------------------------//
+			if (C == 0.0)	nx = ny = nz = 0.0;
+			m1 = m1 + rlx_setA*((19*(ux*ux+uy*uy+uz*uz)*rho0/porosity - 11*rho) -19*alpha*C - m1);
+			m2 = m2 + rlx_setA*((3*rho - 5.5*(ux*ux+uy*uy+uz*uz)*rho0/porosity)- m2);
+            jx = jx + Fx;
+			m4 = m4 + rlx_setB*((-0.6666666666666666*ux*rho0)- m4)
+                    + (1-0.5*rlx_setB)*(-0.6666666666666666*Fx);
+            jy = jy + Fy;
+			m6 = m6 + rlx_setB*((-0.6666666666666666*uy*rho0)- m6)
+                    + (1-0.5*rlx_setB)*(-0.6666666666666666*Fy);
+            jz = jz + Fz;
+			m8 = m8 + rlx_setB*((-0.6666666666666666*uz*rho0)- m8)
+                    + (1-0.5*rlx_setB)*(-0.6666666666666666*Fz);
+			m9 = m9 + rlx_setA*(((2*ux*ux-uy*uy-uz*uz)*rho0/porosity) + 0.5*alpha*C*(2*nx*nx-ny*ny-nz*nz) - m9);
+			m10 = m10 + rlx_setA*( - m10);
+            //m10 = m10 + rlx_setA*(-0.5*rho0*((2*ux*ux-uy*uy-uz*uz)/porosity)- m10);
+			m11 = m11 + rlx_setA*(((uy*uy-uz*uz)*rho0/porosity) + 0.5*alpha*C*(ny*ny-nz*nz)- m11);
+			m12 = m12 + rlx_setA*( - m12);
+            //m12 = m12 + rlx_setA*(-0.5*(rho0*(uy*uy-uz*uz)/porosity)- m12);
+			m13 = m13 + rlx_setA*( (ux*uy*rho0/porosity) + 0.5*alpha*C*nx*ny - m13);
+			m14 = m14 + rlx_setA*( (uy*uz*rho0/porosity) + 0.5*alpha*C*ny*nz - m14);
+			m15 = m15 + rlx_setA*( (ux*uz*rho0/porosity) + 0.5*alpha*C*nx*nz - m15);
+			m16 = m16 + rlx_setB*( - m16);
+			m17 = m17 + rlx_setB*( - m17);
+			m18 = m18 + rlx_setB*( - m18);
+            //----------------------------------------------------------------------//
+
+            //----------------With higher-order force ------------------------------//
+			//if (C == 0.0)	nx = ny = nz = 0.0;
+			//m1 = m1 + rlx_setA*((19*(ux*ux+uy*uy+uz*uz)*rho0/porosity - 11*rho) -19*alpha*C - m1)
+            //        + (1-0.5*rlx_setA)*38*(Fx*ux+Fy*uy+Fz*uz)/porosity;
+			//m2 = m2 + rlx_setA*((3*rho - 5.5*(ux*ux+uy*uy+uz*uz)*rho0/porosity)- m2)
+            //        + (1-0.5*rlx_setA)*11*(-Fx*ux-Fy*uy-Fz*uz)/porosity;
+            //jx = jx + Fx;
+			//m4 = m4 + rlx_setB*((-0.6666666666666666*ux*rho0)- m4)
+            //        + (1-0.5*rlx_setB)*(-0.6666666666666666*Fx);
+            //jy = jy + Fy;
+			//m6 = m6 + rlx_setB*((-0.6666666666666666*uy*rho0)- m6)
+            //        + (1-0.5*rlx_setB)*(-0.6666666666666666*Fy);
+            //jz = jz + Fz;
+			//m8 = m8 + rlx_setB*((-0.6666666666666666*uz*rho0)- m8)
+            //        + (1-0.5*rlx_setB)*(-0.6666666666666666*Fz);
+			//m9 = m9 + rlx_setA*(((2*ux*ux-uy*uy-uz*uz)*rho0/porosity) + 0.5*alpha*C*(2*nx*nx-ny*ny-nz*nz) - m9)
+            //        + (1-0.5*rlx_setA)*(4*Fx*ux-2*Fy*uy-2*Fz*uz)/porosity;
+			////m10 = m10 + rlx_setA*( - m10);
+            //m10 = m10 + rlx_setA*(-0.5*rho0*((2*ux*ux-uy*uy-uz*uz)/porosity)- m10)
+            //          + (1-0.5*rlx_setA)*(-2*Fx*ux+Fy*uy+Fz*uz)/porosity;
+			//m11 = m11 + rlx_setA*(((uy*uy-uz*uz)*rho0/porosity) + 0.5*alpha*C*(ny*ny-nz*nz)- m11)
+            //          + (1-0.5*rlx_setA)*(2*Fy*uy-2*Fz*uz)/porosity;
+			////m12 = m12 + rlx_setA*( - m12);
+            //m12 = m12 + rlx_setA*(-0.5*(rho0*(uy*uy-uz*uz)/porosity)- m12)
+            //          + (1-0.5*rlx_setA)*(-Fy*uy+Fz*uz)/porosity;
+			//m13 = m13 + rlx_setA*( (ux*uy*rho0/porosity) + 0.5*alpha*C*nx*ny - m13);
+            //          + (1-0.5*rlx_setA)*(Fy*ux+Fx*uy)/porosity;
+			//m14 = m14 + rlx_setA*( (uy*uz*rho0/porosity) + 0.5*alpha*C*ny*nz - m14);
+            //          + (1-0.5*rlx_setA)*(Fz*uy+Fy*uz)/porosity;
+			//m15 = m15 + rlx_setA*( (ux*uz*rho0/porosity) + 0.5*alpha*C*nx*nz - m15);
+            //          + (1-0.5*rlx_setA)*(Fz*ux+Fx*uz)/porosity;
+			//m16 = m16 + rlx_setB*( - m16);
+			//m17 = m17 + rlx_setB*( - m17);
+			//m18 = m18 + rlx_setB*( - m18);
+            //----------------------------------------------------------------------//
+
+			//.................inverse transformation......................................................
+			// q=0
+			fq = mrt_V1*rho-mrt_V2*m1+mrt_V3*m2;
+			dist[n] = fq;
+
+			// q = 1
+			fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(jx-m4)+mrt_V6*(m9-m10);
+			dist[1*Np+n] = fq;
+
+			// q=2
+			fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(m4-jx)+mrt_V6*(m9-m10);
+			dist[2*Np+n] = fq;
+
+			// q = 3
+			fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(jy-m6)+mrt_V7*(m10-m9)+mrt_V8*(m11-m12);
+			dist[3*Np+n] = fq;
+
+			// q = 4
+			fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(m6-jy)+mrt_V7*(m10-m9)+mrt_V8*(m11-m12);
+			dist[4*Np+n] = fq;
+
+			// q = 5
+			fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(jz-m8)+mrt_V7*(m10-m9)+mrt_V8*(m12-m11);
+			dist[5*Np+n] = fq;
+
+			// q = 6
+			fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(m8-jz)+mrt_V7*(m10-m9)+mrt_V8*(m12-m11);
+			dist[6*Np+n] = fq;
+
+			// q = 7
+			fq = mrt_V1*rho+mrt_V9*m1+mrt_V10*m2+0.1*(jx+jy)+0.025*(m4+m6)+
+					mrt_V7*m9+mrt_V11*m10+mrt_V8*m11+mrt_V12*m12+0.25*m13+0.125*(m16-m17);
+			dist[7*Np+n] = fq;
+
+
+			// q = 8
+			fq = mrt_V1*rho+mrt_V9*m1+mrt_V10*m2-0.1*(jx+jy)-0.025*(m4+m6) +mrt_V7*m9+mrt_V11*m10+mrt_V8*m11
+					+mrt_V12*m12+0.25*m13+0.125*(m17-m16);
+			dist[8*Np+n] = fq;
+
+			// q = 9
+			fq = mrt_V1*rho+mrt_V9*m1+mrt_V10*m2+0.1*(jx-jy)+0.025*(m4-m6)+
+					mrt_V7*m9+mrt_V11*m10+mrt_V8*m11+mrt_V12*m12-0.25*m13+0.125*(m16+m17);
+			dist[9*Np+n] = fq;
+
+			// q = 10
+			fq = mrt_V1*rho+mrt_V9*m1+mrt_V10*m2+0.1*(jy-jx)+0.025*(m6-m4)+
+					mrt_V7*m9+mrt_V11*m10+mrt_V8*m11+mrt_V12*m12-0.25*m13-0.125*(m16+m17);
+			dist[10*Np+n] = fq;
+
+
+			// q = 11
+			fq = mrt_V1*rho+mrt_V9*m1
+					+mrt_V10*m2+0.1*(jx+jz)+0.025*(m4+m8)
+					+mrt_V7*m9+mrt_V11*m10-mrt_V8*m11
+					-mrt_V12*m12+0.25*m15+0.125*(m18-m16);
+			dist[11*Np+n] = fq;
+
+			// q = 12
+			fq = mrt_V1*rho+mrt_V9*m1+mrt_V10*m2-0.1*(jx+jz)-0.025*(m4+m8)+
+					mrt_V7*m9+mrt_V11*m10-mrt_V8*m11-mrt_V12*m12+0.25*m15+0.125*(m16-m18);
+			dist[12*Np+n] = fq;
+
+			// q = 13
+			fq = mrt_V1*rho+mrt_V9*m1
+					+mrt_V10*m2+0.1*(jx-jz)+0.025*(m4-m8)
+					+mrt_V7*m9+mrt_V11*m10-mrt_V8*m11
+					-mrt_V12*m12-0.25*m15-0.125*(m16+m18);
+			dist[13*Np+n] = fq;
+
+			// q= 14
+			fq = mrt_V1*rho+mrt_V9*m1
+					+mrt_V10*m2+0.1*(jz-jx)+0.025*(m8-m4)
+					+mrt_V7*m9+mrt_V11*m10-mrt_V8*m11
+					-mrt_V12*m12-0.25*m15+0.125*(m16+m18);
+
+			dist[14*Np+n] = fq;
+
+			// q = 15
+			fq = mrt_V1*rho+mrt_V9*m1
+					+mrt_V10*m2+0.1*(jy+jz)+0.025*(m6+m8)
+					-mrt_V6*m9-mrt_V7*m10+0.25*m14+0.125*(m17-m18);
+			dist[15*Np+n] = fq;
+
+			// q = 16
+			fq =  mrt_V1*rho+mrt_V9*m1
+					+mrt_V10*m2-0.1*(jy+jz)-0.025*(m6+m8)
+					-mrt_V6*m9-mrt_V7*m10+0.25*m14+0.125*(m18-m17);
+			dist[16*Np+n] = fq;
+
+
+			// q = 17
+			fq = mrt_V1*rho+mrt_V9*m1
+					+mrt_V10*m2+0.1*(jy-jz)+0.025*(m6-m8)
+					-mrt_V6*m9-mrt_V7*m10-0.25*m14+0.125*(m17+m18);
+			dist[17*Np+n] = fq;
+
+			// q = 18
+			fq = mrt_V1*rho+mrt_V9*m1
+					+mrt_V10*m2+0.1*(jz-jy)+0.025*(m8-m6)
+					-mrt_V6*m9-mrt_V7*m10-0.25*m14-0.125*(m17+m18);
+			dist[18*Np+n] = fq;
+			//........................................................................
+
+			// Instantiate mass transport distributions
+			// Stationary value - distribution 0
+			nAB = 1.0/(nA+nB);
+			Aq[n] = 0.3333333333333333*nA;
+			Bq[n] = 0.3333333333333333*nB;
+
+			//...............................................
+			// q = 0,2,4
+			// Cq = {1,0,0}, {0,1,0}, {0,0,1}
+			delta = beta*nA*nB*nAB*0.1111111111111111*nx;
+			if (!(nA*nB*nAB>0)) delta=0;
+			a1 = nA*(0.1111111111111111*(1+4.5*ux))+delta;
+			b1 = nB*(0.1111111111111111*(1+4.5*ux))-delta;
+			a2 = nA*(0.1111111111111111*(1-4.5*ux))-delta;
+			b2 = nB*(0.1111111111111111*(1-4.5*ux))+delta;
+
+			Aq[1*Np+n] = a1;
+			Bq[1*Np+n] = b1;
+			Aq[2*Np+n] = a2;
+			Bq[2*Np+n] = b2;
+
+			//...............................................
+			// q = 2
+			// Cq = {0,1,0}
+			delta = beta*nA*nB*nAB*0.1111111111111111*ny;
+			if (!(nA*nB*nAB>0)) delta=0;
+			a1 = nA*(0.1111111111111111*(1+4.5*uy))+delta;
+			b1 = nB*(0.1111111111111111*(1+4.5*uy))-delta;
+			a2 = nA*(0.1111111111111111*(1-4.5*uy))-delta;
+			b2 = nB*(0.1111111111111111*(1-4.5*uy))+delta;
+
+			Aq[3*Np+n] = a1;
+			Bq[3*Np+n] = b1;
+			Aq[4*Np+n] = a2;
+			Bq[4*Np+n] = b2;
+			//...............................................
+			// q = 4
+			// Cq = {0,0,1}
+			delta = beta*nA*nB*nAB*0.1111111111111111*nz;
+			if (!(nA*nB*nAB>0)) delta=0;
+			a1 = nA*(0.1111111111111111*(1+4.5*uz))+delta;
+			b1 = nB*(0.1111111111111111*(1+4.5*uz))-delta;
+			a2 = nA*(0.1111111111111111*(1-4.5*uz))-delta;
+			b2 = nB*(0.1111111111111111*(1-4.5*uz))+delta;
+
+			Aq[5*Np+n] = a1;
+			Bq[5*Np+n] = b1;
+			Aq[6*Np+n] = a2;
+			Bq[6*Np+n] = b2;
+			//...............................................
+
+		}
+	}
+}
+
 __global__ void dvc_ScaLBL_PhaseField_InitFromRestart(double *Den, double *Aq, double *Bq, int start, int finish, int Np){
 	int idx;
 	double nA,nB;

From dfac3e134954a3e32a095e828d79c5dd53a89716 Mon Sep 17 00:00:00 2001
From: Rex Zhe Li <zhe.rex.li@gmail.com>
Date: Sun, 28 Feb 2021 18:27:16 -0500
Subject: [PATCH 02/10] continue to add capillary penalty to pressure; to be
 compiled and tested

---
 common/ScaLBL.h                |  10 +++
 cuda/GreyscaleColor.cu         | 115 +++++++++++++++++++++------------
 models/GreyscaleColorModel.cpp |  50 ++++++++++----
 models/GreyscaleColorModel.h   |   2 +
 4 files changed, 121 insertions(+), 56 deletions(-)

diff --git a/common/ScaLBL.h b/common/ScaLBL.h
index 42c51525..007bb35f 100644
--- a/common/ScaLBL.h
+++ b/common/ScaLBL.h
@@ -87,6 +87,16 @@ extern "C" void ScaLBL_D3Q19_AAodd_GreyscaleColor(int *d_neighborList, int *Map,
         double rhoA, double rhoB, double tauA, double tauB, double tauA_eff,double tauB_eff, double alpha, double beta,
 		double Fx, double Fy, double Fz, int strideY, int strideZ, int start, int finish, int Np);
 
+extern "C" void ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dist, double *Aq, double *Bq, double *Den, 
+        double *Phi,double *GreySolidGrad, double *Poros,double *Perm,double *Vel, double *Pressure,
+        double rhoA, double rhoB, double tauA, double tauB,double tauA_eff,double tauB_eff, double alpha, double beta,
+		double Fx, double Fy, double Fz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np);
+
+extern "C" void ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *d_neighborList, int *Map, double *dist, double *Aq, double *Bq, double *Den, 
+		double *Phi, double *GreySolidGrad, double *Poros,double *Perm,double *Vel,double *Pressure, 
+        double rhoA, double rhoB, double tauA, double tauB, double tauA_eff,double tauB_eff, double alpha, double beta,
+		double Fx, double Fy, double Fz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np);
+
 // ION TRANSPORT MODEL
 
 extern "C" void ScaLBL_D3Q7_AAodd_IonConcentration(int *neighborList, double *dist, double *Den, int start, int finish, int Np);
diff --git a/cuda/GreyscaleColor.cu b/cuda/GreyscaleColor.cu
index 0cc49bcd..906782c9 100644
--- a/cuda/GreyscaleColor.cu
+++ b/cuda/GreyscaleColor.cu
@@ -1452,7 +1452,7 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor(int *Map, double *dist,
 __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int *Map, double *dist, double *Aq, double *Bq, double *Den,
 		 double *Phi, double *GreySolidGrad, double *Poros,double *Perm, double *Velocity, double *Pressure,
          double rhoA, double rhoB, double tauA, double tauB,double tauA_eff,double tauB_eff,double alpha, double beta,
-		double Gx, double Gy, double Gz, bool RecoloringOff, int strideY, int strideZ, int start, int finish, int Np){
+		double Gx, double Gy, double Gz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np){
 
 	int n,nn,ijk,nread;
 	int nr1,nr2,nr3,nr4,nr5,nr6;
@@ -1472,12 +1472,12 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 	double a1,b1,a2,b2,nAB,delta;
 	double C,nx,ny,nz; //color gradient magnitude and direction
 	double phi,tau,rho0,rlx_setA,rlx_setB;
-    double Fcpx,Fcpy,Fcpz;
+    double cp;//capillary pressure penalty - pressure term
 
     //double GeoFun=0.0;//geometric function from Guo's PRE 66, 036304 (2002)
     double porosity;
     double perm;//voxel permeability
-    double c0, c1; //Guo's model parameters
+    //double c0, c1; //Guo's model parameters
     double tau_eff;
     double mu_eff;//kinematic viscosity
     double nx_gs,ny_gs,nz_gs;//grey-solid color gradient
@@ -1606,16 +1606,6 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 			nx = nx/ColorMag;
 			ny = ny/ColorMag;
 			nz = nz/ColorMag;		
-            //----------- Introduce capillary penalty force -------------------------
-            //NOTE: apply only to grey nodes
-            Fcpx = 0.5*alpha*C*W/sqrt(perm)*nx;
-            Fcpy = 0.5*alpha*C*W/sqrt(perm)*ny;
-            Fcpz = 0.5*alpha*C*W/sqrt(perm)*nz;
-            if (porosity==1.0){
-                Fcpx = 0.0;
-                Fcpy = 0.0;
-                Fcpz = 0.0;
-            }
 
 			// q=0
 			fq = dist[n];
@@ -1939,14 +1929,14 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 			m18 -= fq;
 			
             //// Compute greyscale related parameters
-            ux = (jx/rho0+0.5*porosity*Gx+0.5*Fcpx/rho0)/(1.0+0.5*porosity*mu_eff/perm);
-            uy = (jy/rho0+0.5*porosity*Gy+0.5*Fcpy/rho0)/(1.0+0.5*porosity*mu_eff/perm);
-            uz = (jz/rho0+0.5*porosity*Gz+0.5*Fcpz/rho0)/(1.0+0.5*porosity*mu_eff/perm);
+            ux = (jx/rho0+0.5*porosity*Gx)/(1.0+0.5*porosity*mu_eff/perm);
+            uy = (jy/rho0+0.5*porosity*Gy)/(1.0+0.5*porosity*mu_eff/perm);
+            uz = (jz/rho0+0.5*porosity*Gz)/(1.0+0.5*porosity*mu_eff/perm);
 
             //Update the total force to include linear (Darcy) and nonlinear (Forchheimer) drags due to the porous medium
-            Fx = rho0*(-porosity*mu_eff/perm*ux + porosity*Gx)+Fcpx;
-            Fy = rho0*(-porosity*mu_eff/perm*uy + porosity*Gy)+Fcpy;
-            Fz = rho0*(-porosity*mu_eff/perm*uz + porosity*Gz)+Fcpz;
+            Fx = rho0*(-porosity*mu_eff/perm*ux + porosity*Gx);
+            Fy = rho0*(-porosity*mu_eff/perm*uy + porosity*Gy);
+            Fz = rho0*(-porosity*mu_eff/perm*uz + porosity*Gz);
             if (porosity==1.0){
                 Fx=rho0*(porosity*Gx);
                 Fy=rho0*(porosity*Gy);
@@ -1960,6 +1950,14 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
             //Pressure[n] = rho/3.f/porosity;
             Pressure[n] = rho/3.f;
 
+            //----------- Introduce capillary penalty force -------------------------
+            //NOTE: apply only to grey nodes
+            cp = 0.1*tanh(W*alpha*phi/sqrt(perm));//the extra factor of 0.1 is to make sure cp is bounded within [-0.1,0.1]
+            if (porosity==1.0){
+                cp = 0.0;
+            }
+            rho += cp;//pressure perturbation
+
 			//........................................................................
 			//..............carry out relaxation process..............................
 			//..........Toelke, Fruediger et. al. 2006................................
@@ -2220,7 +2218,7 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dist, double *Aq, double *Bq, double *Den, 
         double *Phi, double *GreySolidGrad, double *Poros,double *Perm, double *Velocity, double *Pressure, 
         double rhoA, double rhoB, double tauA, double tauB,double tauA_eff,double tauB_eff, double alpha, double beta,
-		double Gx, double Gy, double Gz, bool RecoloringOff, int strideY, int strideZ, int start, int finish, int Np){
+		double Gx, double Gy, double Gz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np){
 	int ijk,nn,n;
 	double fq;
 	// conserved momemnts
@@ -2235,12 +2233,12 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
 	double a1,b1,a2,b2,nAB,delta;
 	double C,nx,ny,nz; //color gradient magnitude and direction
 	double phi,tau,rho0,rlx_setA,rlx_setB;
-    double Fcpx,Fcpy,Fcpz;
+    double cp;//capillary pressure penalty - pressure term
 
     //double GeoFun=0.0;//geometric function from Guo's PRE 66, 036304 (2002)
     double porosity;
     double perm;//voxel permeability
-    double c0, c1; //Guo's model parameters
+    //double c0, c1; //Guo's model parameters
     double tau_eff;
     double mu_eff;//kinematic viscosity
     double nx_gs,ny_gs,nz_gs;//grey-solid color gradient
@@ -2640,30 +2638,19 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
 			m17 -= fq;
 			m18 -= fq;
 
-            // Compute greyscale related parameters
-            c0 = 0.5*(1.0+porosity*0.5*mu_eff/perm);
-            if (porosity==1.0) c0 = 0.5;//i.e. apparent pore nodes
-            //GeoFun = 1.75/sqrt(150.0*porosity*porosity*porosity);
-            c1 = porosity*0.5*GeoFun/sqrt(perm);
-            if (porosity==1.0) c1 = 0.0;//i.e. apparent pore nodes
-
-            vx = jx/rho0+0.5*(porosity*Gx);
-            vy = jy/rho0+0.5*(porosity*Gy);
-            vz = jz/rho0+0.5*(porosity*Gz);
-            v_mag=sqrt(vx*vx+vy*vy+vz*vz);
-            ux = vx/(c0+sqrt(c0*c0+c1*v_mag));
-            uy = vy/(c0+sqrt(c0*c0+c1*v_mag));
-            uz = vz/(c0+sqrt(c0*c0+c1*v_mag));
-            u_mag=sqrt(ux*ux+uy*uy+uz*uz);
+            //// Compute greyscale related parameters
+            ux = (jx/rho0+0.5*porosity*Gx)/(1.0+0.5*porosity*mu_eff/perm);
+            uy = (jy/rho0+0.5*porosity*Gy)/(1.0+0.5*porosity*mu_eff/perm);
+            uz = (jz/rho0+0.5*porosity*Gz)/(1.0+0.5*porosity*mu_eff/perm);
 
             //Update the total force to include linear (Darcy) and nonlinear (Forchheimer) drags due to the porous medium
-            Fx = rho0*(-porosity*mu_eff/perm*ux - porosity*GeoFun/sqrt(perm)*u_mag*ux + porosity*Gx);
-            Fy = rho0*(-porosity*mu_eff/perm*uy - porosity*GeoFun/sqrt(perm)*u_mag*uy + porosity*Gy);
-            Fz = rho0*(-porosity*mu_eff/perm*uz - porosity*GeoFun/sqrt(perm)*u_mag*uz + porosity*Gz);
+            Fx = rho0*(-porosity*mu_eff/perm*ux + porosity*Gx);
+            Fy = rho0*(-porosity*mu_eff/perm*uy + porosity*Gy);
+            Fz = rho0*(-porosity*mu_eff/perm*uz + porosity*Gz);
             if (porosity==1.0){
-                Fx=rho0*(Gx);
-                Fy=rho0*(Gy);
-                Fz=rho0*(Gz);
+                Fx=rho0*(porosity*Gx);
+                Fy=rho0*(porosity*Gy);
+                Fz=rho0*(porosity*Gz);
             }
 
 			// write the velocity 
@@ -2673,6 +2660,14 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
             //Pressure[n] = rho/3.f/porosity;
             Pressure[n] = rho/3.f;
 
+            //----------- Introduce capillary penalty force -------------------------
+            //NOTE: apply only to grey nodes
+            cp = 0.1*tanh(W*alpha*phi/sqrt(perm));//the extra factor of 0.1 is to make sure cp is bounded within [-0.1,0.1]
+            if (porosity==1.0){
+                cp = 0.0;
+            }
+            rho += cp;//pressure perturbation
+
 			//........................................................................
 			//..............carry out relaxation process..............................
 			//..........Toelke, Fruediger et. al. 2006................................
@@ -2854,6 +2849,7 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
 			// Cq = {1,0,0}, {0,1,0}, {0,0,1}
 			delta = beta*nA*nB*nAB*0.1111111111111111*nx;
 			if (!(nA*nB*nAB>0)) delta=0;
+            if (RecoloringOff==true && porosity !=1.0) delta=0;
 			a1 = nA*(0.1111111111111111*(1+4.5*ux))+delta;
 			b1 = nB*(0.1111111111111111*(1+4.5*ux))-delta;
 			a2 = nA*(0.1111111111111111*(1-4.5*ux))-delta;
@@ -2869,6 +2865,7 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
 			// Cq = {0,1,0}
 			delta = beta*nA*nB*nAB*0.1111111111111111*ny;
 			if (!(nA*nB*nAB>0)) delta=0;
+            if (RecoloringOff==true && porosity !=1.0) delta=0;
 			a1 = nA*(0.1111111111111111*(1+4.5*uy))+delta;
 			b1 = nB*(0.1111111111111111*(1+4.5*uy))-delta;
 			a2 = nA*(0.1111111111111111*(1-4.5*uy))-delta;
@@ -2883,6 +2880,7 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
 			// Cq = {0,0,1}
 			delta = beta*nA*nB*nAB*0.1111111111111111*nz;
 			if (!(nA*nB*nAB>0)) delta=0;
+            if (RecoloringOff==true && porosity !=1.0) delta=0;
 			a1 = nA*(0.1111111111111111*(1+4.5*uz))+delta;
 			b1 = nB*(0.1111111111111111*(1+4.5*uz))-delta;
 			a2 = nA*(0.1111111111111111*(1-4.5*uz))-delta;
@@ -4448,6 +4446,37 @@ extern "C" void ScaLBL_PhaseField_InitFromRestart(double *Den, double *Aq, doubl
 		printf("CUDA error in ScaLBL_PhaseField_InitFromRestart: %s \n",cudaGetErrorString(err));
 	}
 }
+
+//Model-1 & 4 with capillary pressure penalty
+extern "C" void ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dist, double *Aq, double *Bq, double *Den, 
+        double *Phi,double *GreySolidGrad, double *Poros,double *Perm,double *Vel, double *Pressure,
+        double rhoA, double rhoB, double tauA, double tauB,double tauA_eff,double tauB_eff, double alpha, double beta,
+		double Fx, double Fy, double Fz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np){
+
+	dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP<<<NBLOCKS,NTHREADS >>>(Map, dist, Aq, Bq, Den, Phi, GreySolidGrad, Poros, Perm, Vel, Pressure,
+            rhoA, rhoB, tauA, tauB, tauA_eff, tauB_eff, alpha, beta, Fx, Fy, Fz, RecoloringOff, W, strideY, strideZ, start, finish, Np);
+	cudaError_t err = cudaGetLastError();
+	if (cudaSuccess != err){
+		printf("CUDA error in ScaLBL_D3Q19_AAeven_GreyscaleColor_CP: %s \n",cudaGetErrorString(err));
+	}
+
+}
+
+//Model-1 & 4 with capillary pressure penalty
+extern "C" void ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *d_neighborList, int *Map, double *dist, double *Aq, double *Bq, double *Den, 
+		double *Phi, double *GreySolidGrad, double *Poros,double *Perm,double *Vel,double *Pressure, 
+        double rhoA, double rhoB, double tauA, double tauB, double tauA_eff,double tauB_eff, double alpha, double beta,
+		double Fx, double Fy, double Fz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np){
+
+	dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP<<<NBLOCKS,NTHREADS >>>(d_neighborList, Map, dist, Aq, Bq, Den, Phi,  GreySolidGrad, Poros, Perm,Vel,Pressure,
+			rhoA, rhoB, tauA, tauB, tauA_eff, tauB_eff,alpha, beta, Fx, Fy, Fz, RecoloringOff, W, strideY, strideZ, start, finish, Np);
+
+	cudaError_t err = cudaGetLastError();
+	if (cudaSuccess != err){
+		printf("CUDA error in ScaLBL_D3Q19_AAodd_GreyscaleColor_CP: %s \n",cudaGetErrorString(err));
+	}
+}
+
 ////Model-2&3
 //extern "C" void ScaLBL_D3Q19_AAeven_GreyscaleColor(int *Map, double *dist, double *Aq, double *Bq, double *Den, 
 //        double *Phi,double *GreySolidGrad, double *Poros,double *Perm,double *Vel, 
diff --git a/models/GreyscaleColorModel.cpp b/models/GreyscaleColorModel.cpp
index 5d2b4d07..d5ecea2d 100644
--- a/models/GreyscaleColorModel.cpp
+++ b/models/GreyscaleColorModel.cpp
@@ -17,7 +17,7 @@ void DeleteArray( const TYPE *p )
 
 ScaLBL_GreyscaleColorModel::ScaLBL_GreyscaleColorModel(int RANK, int NP, const Utilities::MPI& COMM):
 rank(RANK), nprocs(NP), Restart(0),timestep(0),timestepMax(0),tauA(0),tauB(0),tauA_eff(0),tauB_eff(0),rhoA(0),rhoB(0),alpha(0),beta(0),
-Fx(0),Fy(0),Fz(0),flux(0),din(0),dout(0),inletA(0),inletB(0),outletA(0),outletB(0),GreyPorosity(0),
+Fx(0),Fy(0),Fz(0),flux(0),din(0),dout(0),inletA(0),inletB(0),outletA(0),outletB(0),GreyPorosity(0),RecoloringOff(0),W(0),
 Nx(0),Ny(0),Nz(0),N(0),Np(0),nprocx(0),nprocy(0),nprocz(0),BoundaryCondition(0),Lx(0),Ly(0),Lz(0),comm(COMM)
 {
 	REVERSE_FLOW_DIRECTION = false;
@@ -43,6 +43,8 @@ void ScaLBL_GreyscaleColorModel::ReadParams(string filename){
 	Restart=false;
 	din=dout=1.0;
 	flux=0.0;
+    RecoloringOff = false;
+    W=1.0;
 	
 	// Color Model parameters
 	if (greyscaleColor_db->keyExists( "timestepMax" )){
@@ -85,6 +87,12 @@ void ScaLBL_GreyscaleColorModel::ReadParams(string filename){
 	if (greyscaleColor_db->keyExists( "flux" )){
 		flux = greyscaleColor_db->getScalar<double>( "flux" );
 	}
+	if (greyscaleColor_db->keyExists( "RecoloringOff" )){
+		RecoloringOff = greyscaleColor_db->getScalar<bool>( "RecoloringOff" );
+	}
+	if (greyscaleColor_db->keyExists( "W" )){
+		W = greyscaleColor_db->getScalar<double>( "W" );
+	}
 	inletA=1.f;
 	inletB=0.f;
 	outletA=0.f;
@@ -943,10 +951,14 @@ void ScaLBL_GreyscaleColorModel::Run(){
 		}
 		// Halo exchange for phase field
 		ScaLBL_Comm_Regular->SendHalo(Phi);
-        //Model-1&4
-        ScaLBL_D3Q19_AAodd_GreyscaleColor(NeighborList, dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
+        //Model-1&4 with capillary pressure penalty for grey nodes
+        ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(NeighborList, dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
                 rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff, 
-                alpha, beta, Fx, Fy, Fz, Nx, Nx*Ny, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
+                alpha, beta, Fx, Fy, Fz, RecoloringOff, W, Nx, Nx*Ny, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
+        //Model-1&4
+        //ScaLBL_D3Q19_AAodd_GreyscaleColor(NeighborList, dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
+        //        rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff, 
+        //        alpha, beta, Fx, Fy, Fz, Nx, Nx*Ny, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
         ////Model-2&3
         //ScaLBL_D3Q19_AAodd_GreyscaleColor(NeighborList, dvcMap, fq, Aq, Bq, Den, Phi,GreySolidPhi,Porosity_dvc,Permeability_dvc,Velocity, 
         //        rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff, 
@@ -968,10 +980,14 @@ void ScaLBL_GreyscaleColorModel::Run(){
 			ScaLBL_Comm->D3Q19_Reflection_BC_Z(fq);
 		}
 
-        //Model-1&4
-        ScaLBL_D3Q19_AAodd_GreyscaleColor(NeighborList, dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
+        //Model-1&4 with capillary pressure penalty for grey nodes
+        ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(NeighborList, dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
                 rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
-                alpha, beta, Fx, Fy, Fz, Nx, Nx*Ny, 0, ScaLBL_Comm->LastExterior(), Np);
+                alpha, beta, Fx, Fy, Fz, RecoloringOff, W, Nx, Nx*Ny, 0, ScaLBL_Comm->LastExterior(), Np);
+        //Model-1&4
+        //ScaLBL_D3Q19_AAodd_GreyscaleColor(NeighborList, dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
+        //        rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
+        //        alpha, beta, Fx, Fy, Fz, Nx, Nx*Ny, 0, ScaLBL_Comm->LastExterior(), Np);
         ////Model-2&3
         //ScaLBL_D3Q19_AAodd_GreyscaleColor(NeighborList, dvcMap, fq, Aq, Bq, Den, Phi,GreySolidPhi,Porosity_dvc,Permeability_dvc,Velocity,
         //        rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
@@ -995,10 +1011,14 @@ void ScaLBL_GreyscaleColorModel::Run(){
 			ScaLBL_Comm->Color_BC_Z(dvcMap, Phi, Den, outletA, outletB);
 		}
 		ScaLBL_Comm_Regular->SendHalo(Phi);
-        //Model-1&4
-        ScaLBL_D3Q19_AAeven_GreyscaleColor(dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure, 
+        //Model-1&4 with capillary pressure penalty for grey nodes
+        ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure, 
                 rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
-                alpha, beta, Fx, Fy, Fz,  Nx, Nx*Ny, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
+                alpha, beta, Fx, Fy, Fz, RecoloringOff, W, Nx, Nx*Ny, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
+        //Model-1&4
+        //ScaLBL_D3Q19_AAeven_GreyscaleColor(dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure, 
+        //        rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
+        //        alpha, beta, Fx, Fy, Fz,  Nx, Nx*Ny, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
         ////Model-2&3
         //ScaLBL_D3Q19_AAeven_GreyscaleColor(dvcMap, fq, Aq, Bq, Den, Phi,GreySolidPhi,Porosity_dvc,Permeability_dvc,Velocity, 
         //        rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
@@ -1020,10 +1040,14 @@ void ScaLBL_GreyscaleColorModel::Run(){
 			ScaLBL_Comm->D3Q19_Reflection_BC_Z(fq);
 		}
 
-        //Model-1&4
-        ScaLBL_D3Q19_AAeven_GreyscaleColor(dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
+        //Model-1&4 with capillary pressure penalty for grey nodes
+        ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
                 rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
-                alpha, beta, Fx, Fy, Fz, Nx, Nx*Ny, 0, ScaLBL_Comm->LastExterior(), Np);
+                alpha, beta, Fx, Fy, Fz, RecoloringOff, W, Nx, Nx*Ny, 0, ScaLBL_Comm->LastExterior(), Np);
+        //Model-1&4
+        //ScaLBL_D3Q19_AAeven_GreyscaleColor(dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
+        //        rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
+        //        alpha, beta, Fx, Fy, Fz, Nx, Nx*Ny, 0, ScaLBL_Comm->LastExterior(), Np);
         ////Model-2&3
         //ScaLBL_D3Q19_AAeven_GreyscaleColor(dvcMap, fq, Aq, Bq, Den, Phi,GreySolidPhi,Porosity_dvc,Permeability_dvc,Velocity,
         //        rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
diff --git a/models/GreyscaleColorModel.h b/models/GreyscaleColorModel.h
index 667099e9..dd25e613 100644
--- a/models/GreyscaleColorModel.h
+++ b/models/GreyscaleColorModel.h
@@ -39,6 +39,8 @@ public:
 	double Fx,Fy,Fz,flux;
 	double din,dout,inletA,inletB,outletA,outletB;
     double GreyPorosity;
+    bool RecoloringOff;//recoloring can be turn off for grey nodes if this is true
+    double W;//wetting strength paramter for capillary pressure penalty for grey nodes
 	
 	int Nx,Ny,Nz,N,Np;
 	int rank,nprocx,nprocy,nprocz,nprocs;

From 6aeb038fce252b33ac47058dfd00fbd094b61845 Mon Sep 17 00:00:00 2001
From: Rex Zhe Li <zhe.rex.li@gmail.com>
Date: Wed, 3 Mar 2021 20:25:26 -0500
Subject: [PATCH 03/10] update the cap penalty formulation; to be built and
 verified

---
 common/ScaLBL.h                |   3 +
 cuda/GreyscaleColor.cu         | 155 +++++++++++++++++++++++++--------
 models/GreyscaleColorModel.cpp |  94 ++++++++++++++++++--
 models/GreyscaleColorModel.h   |   1 +
 4 files changed, 209 insertions(+), 44 deletions(-)

diff --git a/common/ScaLBL.h b/common/ScaLBL.h
index 007bb35f..762efb20 100644
--- a/common/ScaLBL.h
+++ b/common/ScaLBL.h
@@ -97,6 +97,9 @@ extern "C" void ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *d_neighborList, int *M
         double rhoA, double rhoB, double tauA, double tauB, double tauA_eff,double tauB_eff, double alpha, double beta,
 		double Fx, double Fy, double Fz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np);
 
+extern "C" void ScaLBL_Update_GreyscalePotential(int *Map, double *Phi, double *Psi, double *Poro, double *Perm, double alpha, double W, 
+		int start, int finish);
+
 // ION TRANSPORT MODEL
 
 extern "C" void ScaLBL_D3Q7_AAodd_IonConcentration(int *neighborList, double *dist, double *Den, int start, int finish, int Np);
diff --git a/cuda/GreyscaleColor.cu b/cuda/GreyscaleColor.cu
index 906782c9..ee989bba 100644
--- a/cuda/GreyscaleColor.cu
+++ b/cuda/GreyscaleColor.cu
@@ -1450,9 +1450,9 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor(int *Map, double *dist,
 //CP: capillary penalty
 // also turn off recoloring for grey nodes
 __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int *Map, double *dist, double *Aq, double *Bq, double *Den,
-		 double *Phi, double *GreySolidGrad, double *Poros,double *Perm, double *Velocity, double *Pressure,
+		 double *Phi, double *Psi, double *GreySolidGrad, double *Poros,double *Perm, double *Velocity, double *Pressure,
          double rhoA, double rhoB, double tauA, double tauB,double tauA_eff,double tauB_eff,double alpha, double beta,
-		double Gx, double Gy, double Gz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np){
+		double Gx, double Gy, double Gz, bool RecoloringOff, int strideY, int strideZ, int start, int finish, int Np){
 
 	int n,nn,ijk,nread;
 	int nr1,nr2,nr3,nr4,nr5,nr6;
@@ -1472,7 +1472,7 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 	double a1,b1,a2,b2,nAB,delta;
 	double C,nx,ny,nz; //color gradient magnitude and direction
 	double phi,tau,rho0,rlx_setA,rlx_setB;
-    double cp;//capillary pressure penalty - pressure term
+    double psi;//greyscale potential
 
     //double GeoFun=0.0;//geometric function from Guo's PRE 66, 036304 (2002)
     double porosity;
@@ -1483,6 +1483,9 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
     double nx_gs,ny_gs,nz_gs;//grey-solid color gradient
     double nx_phase,ny_phase,nz_phase,C_phase;
     double Fx,Fy,Fz;
+	double gp1,gp2,gp4,gp6,gp8,gp9,gp10,gp11,gp12,gp13,gp14,gp15,gp16,gp17,gp18;
+	double gp3,gp5,gp7;
+    double Fcpx,Fcpy,Fcpz;//capillary penalty force
 
 	const double mrt_V1=0.05263157894736842;
 	const double mrt_V2=0.012531328320802;
@@ -1532,62 +1535,85 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 			//........................................................................
 			nn = ijk-1;							// neighbor index (get convention)
 			m1 = Phi[nn];						// get neighbor for phi - 1
+            gp1 = Psi[nn];
 			//........................................................................
 			nn = ijk+1;							// neighbor index (get convention)
 			m2 = Phi[nn];						// get neighbor for phi - 2
+            gp2 = Psi[nn];
 			//........................................................................
 			nn = ijk-strideY;							// neighbor index (get convention)
 			m3 = Phi[nn];					// get neighbor for phi - 3
+            gp3 = Psi[nn];
 			//........................................................................
 			nn = ijk+strideY;							// neighbor index (get convention)
 			m4 = Phi[nn];					// get neighbor for phi - 4
+            gp4 = Psi[nn];
 			//........................................................................
 			nn = ijk-strideZ;						// neighbor index (get convention)
 			m5 = Phi[nn];					// get neighbor for phi - 5
+            gp5 = Psi[nn];
 			//........................................................................
 			nn = ijk+strideZ;						// neighbor index (get convention)
 			m6 = Phi[nn];					// get neighbor for phi - 6
+            gp6 = Psi[nn];
 			//........................................................................
 			nn = ijk-strideY-1;						// neighbor index (get convention)
 			m7 = Phi[nn];					// get neighbor for phi - 7
+            gp7 = Psi[nn];
 			//........................................................................
 			nn = ijk+strideY+1;						// neighbor index (get convention)
 			m8 = Phi[nn];					// get neighbor for phi - 8
+            gp8 = Psi[nn];
 			//........................................................................
 			nn = ijk+strideY-1;						// neighbor index (get convention)
 			m9 = Phi[nn];					// get neighbor for phi - 9
+            gp9 = Psi[nn];
 			//........................................................................
 			nn = ijk-strideY+1;						// neighbor index (get convention)
 			m10 = Phi[nn];					// get neighbor for phi - 10
+            gp10 = Psi[nn];
 			//........................................................................
 			nn = ijk-strideZ-1;						// neighbor index (get convention)
 			m11 = Phi[nn];					// get neighbor for phi - 11
+            gp11 = Psi[nn];
 			//........................................................................
 			nn = ijk+strideZ+1;						// neighbor index (get convention)
 			m12 = Phi[nn];					// get neighbor for phi - 12
+            gp12 = Psi[nn];
 			//........................................................................
 			nn = ijk+strideZ-1;						// neighbor index (get convention)
 			m13 = Phi[nn];					// get neighbor for phi - 13
+            gp13 = Psi[nn];
 			//........................................................................
 			nn = ijk-strideZ+1;						// neighbor index (get convention)
 			m14 = Phi[nn];					// get neighbor for phi - 14
+            gp14 = Psi[nn];
 			//........................................................................
 			nn = ijk-strideZ-strideY;					// neighbor index (get convention)
 			m15 = Phi[nn];					// get neighbor for phi - 15
+            gp15 = Psi[nn];
 			//........................................................................
 			nn = ijk+strideZ+strideY;					// neighbor index (get convention)
 			m16 = Phi[nn];					// get neighbor for phi - 16
+            gp16 = Psi[nn];
 			//........................................................................
 			nn = ijk+strideZ-strideY;					// neighbor index (get convention)
 			m17 = Phi[nn];					// get neighbor for phi - 17
+            gp17 = Psi[nn];
 			//........................................................................
 			nn = ijk-strideZ+strideY;					// neighbor index (get convention)
 			m18 = Phi[nn];					// get neighbor for phi - 18
+            gp18 = Psi[nn];
 			//............Compute the Color Gradient...................................
 			nx_phase = -(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
 			ny_phase = -(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
 			nz_phase = -(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
 			C_phase = sqrt(nx_phase*nx_phase+ny_phase*ny_phase+nz_phase*nz_phase);
+			//............Compute the Greyscale Potential Gradient.....................
+			Fcpx = -3.0/18.0*(gp1-gp2+0.5*(gp7-gp8+gp9-gp10+gp11-gp12+gp13-gp14));
+			Fcpy = -3.0/18.0*(gp3-gp4+0.5*(gp7-gp8-gp9+gp10+gp15-gp16+gp17-gp18));
+			Fcpz = -3.0/18.0*(gp5-gp6+0.5*(gp11-gp12-gp13+gp14+gp15-gp16-gp17+gp18));
+
 
             //correct the normal color gradient by considering the effect of grey solid
             nx = nx_phase + (1.0-porosity)*nx_gs; 
@@ -1928,15 +1954,20 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 			m17 -= fq;
 			m18 -= fq;
 			
-            //// Compute greyscale related parameters
-            ux = (jx/rho0+0.5*porosity*Gx)/(1.0+0.5*porosity*mu_eff/perm);
-            uy = (jy/rho0+0.5*porosity*Gy)/(1.0+0.5*porosity*mu_eff/perm);
-            uz = (jz/rho0+0.5*porosity*Gz)/(1.0+0.5*porosity*mu_eff/perm);
+            // Compute greyscale related parameters
+            ux = (jx/rho0+0.5*porosity*Gx+0.5*Fcpx/rho0)/(1.0+0.5*porosity*mu_eff/perm);
+            uy = (jy/rho0+0.5*porosity*Gy+0.5*Fcpy/rho0)/(1.0+0.5*porosity*mu_eff/perm);
+            uz = (jz/rho0+0.5*porosity*Gz+0.5*Fcpz/rho0)/(1.0+0.5*porosity*mu_eff/perm);
+            if (porosity==1.0){//i.e. open nodes
+                ux = (jx/rho0+0.5*porosity*Gx);
+                uy = (jy/rho0+0.5*porosity*Gy);
+                uz = (jz/rho0+0.5*porosity*Gz);
+            }
 
             //Update the total force to include linear (Darcy) and nonlinear (Forchheimer) drags due to the porous medium
-            Fx = rho0*(-porosity*mu_eff/perm*ux + porosity*Gx);
-            Fy = rho0*(-porosity*mu_eff/perm*uy + porosity*Gy);
-            Fz = rho0*(-porosity*mu_eff/perm*uz + porosity*Gz);
+            Fx = rho0*(-porosity*mu_eff/perm*ux + porosity*Gx)+Fcpx;
+            Fy = rho0*(-porosity*mu_eff/perm*uy + porosity*Gy)+Fcpy;
+            Fz = rho0*(-porosity*mu_eff/perm*uz + porosity*Gz)+Fcpz;
             if (porosity==1.0){
                 Fx=rho0*(porosity*Gx);
                 Fy=rho0*(porosity*Gy);
@@ -1950,14 +1981,6 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
             //Pressure[n] = rho/3.f/porosity;
             Pressure[n] = rho/3.f;
 
-            //----------- Introduce capillary penalty force -------------------------
-            //NOTE: apply only to grey nodes
-            cp = 0.1*tanh(W*alpha*phi/sqrt(perm));//the extra factor of 0.1 is to make sure cp is bounded within [-0.1,0.1]
-            if (porosity==1.0){
-                cp = 0.0;
-            }
-            rho += cp;//pressure perturbation
-
 			//........................................................................
 			//..............carry out relaxation process..............................
 			//..........Toelke, Fruediger et. al. 2006................................
@@ -2216,9 +2239,9 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 //CP: capillary penalty
 // also turn off recoloring for grey nodes
 __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dist, double *Aq, double *Bq, double *Den, 
-        double *Phi, double *GreySolidGrad, double *Poros,double *Perm, double *Velocity, double *Pressure, 
+        double *Phi, double *Psi, double *GreySolidGrad, double *Poros,double *Perm, double *Velocity, double *Pressure, 
         double rhoA, double rhoB, double tauA, double tauB,double tauA_eff,double tauB_eff, double alpha, double beta,
-		double Gx, double Gy, double Gz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np){
+		double Gx, double Gy, double Gz, bool RecoloringOff, int strideY, int strideZ, int start, int finish, int Np){
 	int ijk,nn,n;
 	double fq;
 	// conserved momemnts
@@ -2233,7 +2256,7 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
 	double a1,b1,a2,b2,nAB,delta;
 	double C,nx,ny,nz; //color gradient magnitude and direction
 	double phi,tau,rho0,rlx_setA,rlx_setB;
-    double cp;//capillary pressure penalty - pressure term
+    double psi;//greyscale potential
 
     //double GeoFun=0.0;//geometric function from Guo's PRE 66, 036304 (2002)
     double porosity;
@@ -2244,6 +2267,9 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
     double nx_gs,ny_gs,nz_gs;//grey-solid color gradient
     double nx_phase,ny_phase,nz_phase,C_phase;
     double Fx,Fy,Fz;
+	double gp1,gp2,gp4,gp6,gp8,gp9,gp10,gp11,gp12,gp13,gp14,gp15,gp16,gp17,gp18;
+	double gp3,gp5,gp7;
+    double Fcpx,Fcpy,Fcpz;//capillary penalty force
 
 	const double mrt_V1=0.05263157894736842;
 	const double mrt_V2=0.012531328320802;
@@ -2293,62 +2319,84 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
 			//........................................................................
 			nn = ijk-1;							// neighbor index (get convention)
 			m1 = Phi[nn];						// get neighbor for phi - 1
+            gp1 = Psi[nn];
 			//........................................................................
 			nn = ijk+1;							// neighbor index (get convention)
 			m2 = Phi[nn];						// get neighbor for phi - 2
+            gp2 = Psi[nn];
 			//........................................................................
 			nn = ijk-strideY;							// neighbor index (get convention)
 			m3 = Phi[nn];					// get neighbor for phi - 3
+            gp3 = Psi[nn];
 			//........................................................................
 			nn = ijk+strideY;							// neighbor index (get convention)
 			m4 = Phi[nn];					// get neighbor for phi - 4
+            gp4 = Psi[nn];
 			//........................................................................
 			nn = ijk-strideZ;						// neighbor index (get convention)
 			m5 = Phi[nn];					// get neighbor for phi - 5
+            gp5 = Psi[nn];
 			//........................................................................
 			nn = ijk+strideZ;						// neighbor index (get convention)
 			m6 = Phi[nn];					// get neighbor for phi - 6
+            gp6 = Psi[nn];
 			//........................................................................
 			nn = ijk-strideY-1;						// neighbor index (get convention)
 			m7 = Phi[nn];					// get neighbor for phi - 7
+            gp7 = Psi[nn];
 			//........................................................................
 			nn = ijk+strideY+1;						// neighbor index (get convention)
 			m8 = Phi[nn];					// get neighbor for phi - 8
+            gp8 = Psi[nn];
 			//........................................................................
 			nn = ijk+strideY-1;						// neighbor index (get convention)
 			m9 = Phi[nn];					// get neighbor for phi - 9
+            gp9 = Psi[nn];
 			//........................................................................
 			nn = ijk-strideY+1;						// neighbor index (get convention)
 			m10 = Phi[nn];					// get neighbor for phi - 10
+            gp10 = Psi[nn];
 			//........................................................................
 			nn = ijk-strideZ-1;						// neighbor index (get convention)
 			m11 = Phi[nn];					// get neighbor for phi - 11
+            gp11 = Psi[nn];
 			//........................................................................
 			nn = ijk+strideZ+1;						// neighbor index (get convention)
 			m12 = Phi[nn];					// get neighbor for phi - 12
+            gp12 = Psi[nn];
 			//........................................................................
 			nn = ijk+strideZ-1;						// neighbor index (get convention)
 			m13 = Phi[nn];					// get neighbor for phi - 13
+            gp13 = Psi[nn];
 			//........................................................................
 			nn = ijk-strideZ+1;						// neighbor index (get convention)
 			m14 = Phi[nn];					// get neighbor for phi - 14
+            gp14 = Psi[nn];
 			//........................................................................
 			nn = ijk-strideZ-strideY;					// neighbor index (get convention)
 			m15 = Phi[nn];					// get neighbor for phi - 15
+            gp15 = Psi[nn];
 			//........................................................................
 			nn = ijk+strideZ+strideY;					// neighbor index (get convention)
 			m16 = Phi[nn];					// get neighbor for phi - 16
+            gp16 = Psi[nn];
 			//........................................................................
 			nn = ijk+strideZ-strideY;					// neighbor index (get convention)
 			m17 = Phi[nn];					// get neighbor for phi - 17
+            gp17 = Psi[nn];
 			//........................................................................
 			nn = ijk-strideZ+strideY;					// neighbor index (get convention)
 			m18 = Phi[nn];					// get neighbor for phi - 18
+            gp18 = Psi[nn];
 			//............Compute the Color Gradient...................................
 			nx_phase = -(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
 			ny_phase = -(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
 			nz_phase = -(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
 			C_phase = sqrt(nx_phase*nx_phase+ny_phase*ny_phase+nz_phase*nz_phase);
+			//............Compute the Greyscale Potential Gradient.....................
+			Fcpx = -3.0/18.0*(gp1-gp2+0.5*(gp7-gp8+gp9-gp10+gp11-gp12+gp13-gp14));
+			Fcpy = -3.0/18.0*(gp3-gp4+0.5*(gp7-gp8-gp9+gp10+gp15-gp16+gp17-gp18));
+			Fcpz = -3.0/18.0*(gp5-gp6+0.5*(gp11-gp12-gp13+gp14+gp15-gp16-gp17+gp18));
 
             //correct the normal color gradient by considering the effect of grey solid
             nx = nx_phase + (1.0-porosity)*nx_gs; 
@@ -2638,15 +2686,20 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
 			m17 -= fq;
 			m18 -= fq;
 
-            //// Compute greyscale related parameters
-            ux = (jx/rho0+0.5*porosity*Gx)/(1.0+0.5*porosity*mu_eff/perm);
-            uy = (jy/rho0+0.5*porosity*Gy)/(1.0+0.5*porosity*mu_eff/perm);
-            uz = (jz/rho0+0.5*porosity*Gz)/(1.0+0.5*porosity*mu_eff/perm);
+            // Compute greyscale related parameters
+            ux = (jx/rho0+0.5*porosity*Gx+0.5*Fcpx/rho0)/(1.0+0.5*porosity*mu_eff/perm);
+            uy = (jy/rho0+0.5*porosity*Gy+0.5*Fcpy/rho0)/(1.0+0.5*porosity*mu_eff/perm);
+            uz = (jz/rho0+0.5*porosity*Gz+0.5*Fcpz/rho0)/(1.0+0.5*porosity*mu_eff/perm);
+            if (porosity==1.0){//i.e. open nodes
+                ux = (jx/rho0+0.5*porosity*Gx);
+                uy = (jy/rho0+0.5*porosity*Gy);
+                uz = (jz/rho0+0.5*porosity*Gz);
+            }
 
             //Update the total force to include linear (Darcy) and nonlinear (Forchheimer) drags due to the porous medium
-            Fx = rho0*(-porosity*mu_eff/perm*ux + porosity*Gx);
-            Fy = rho0*(-porosity*mu_eff/perm*uy + porosity*Gy);
-            Fz = rho0*(-porosity*mu_eff/perm*uz + porosity*Gz);
+            Fx = rho0*(-porosity*mu_eff/perm*ux + porosity*Gx)+Fcpx;
+            Fy = rho0*(-porosity*mu_eff/perm*uy + porosity*Gy)+Fcpy;
+            Fz = rho0*(-porosity*mu_eff/perm*uz + porosity*Gz)+Fcpz;
             if (porosity==1.0){
                 Fx=rho0*(porosity*Gx);
                 Fy=rho0*(porosity*Gy);
@@ -2660,14 +2713,6 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
             //Pressure[n] = rho/3.f/porosity;
             Pressure[n] = rho/3.f;
 
-            //----------- Introduce capillary penalty force -------------------------
-            //NOTE: apply only to grey nodes
-            cp = 0.1*tanh(W*alpha*phi/sqrt(perm));//the extra factor of 0.1 is to make sure cp is bounded within [-0.1,0.1]
-            if (porosity==1.0){
-                cp = 0.0;
-            }
-            rho += cp;//pressure perturbation
-
 			//........................................................................
 			//..............carry out relaxation process..............................
 			//..........Toelke, Fruediger et. al. 2006................................
@@ -2927,6 +2972,33 @@ __global__ void dvc_ScaLBL_PhaseField_InitFromRestart(double *Den, double *Aq, d
 		}
 	}
 }
+
+__global__  void dvc_ScaLBL_Update_GreyscalePotential(int *Map, double *Phi, double *Psi, double *Poro, double *Perm, double alpha, double W, 
+                                                      int start, int finish){
+	int idx,n;
+    double phi,psi;
+    double cap_penalty;
+    double porosity,perm;
+
+	int S = Np/NBLOCKS/NTHREADS + 1;
+	for (int s=0; s<S; s++){
+		//........Get 1-D index for this thread....................
+		n =  S*blockIdx.x*blockDim.x + s*blockDim.x + threadIdx.x + start;
+		if (n<finish) {
+			idx = Map[n];
+			phi = Phi[idx];
+            porosity = Poro[n];
+            perm = Perm[n];    
+            cap_penalty = 1.0;
+
+            if (porosity!=1.0) cap_penalty = alpha*W/sqrt(perm);
+            psi = cap_penalty*phi;
+            Psi[idx] = psi;
+		}
+	}
+}
+
+
 ////Model-2&3
 //__global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor(int *neighborList, int *Map, double *dist, double *Aq, double *Bq, double *Den,
 //		 double *Phi, double *GreySolidGrad, double *Poros,double *Perm, double *Velocity, 
@@ -4477,6 +4549,17 @@ extern "C" void ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *d_neighborList, int *M
 	}
 }
 
+extern "C" void ScaLBL_Update_GreyscalePotential(int *Map, double *Phi, double *Psi, double *Poro, double *Perm, double alpha, double W, 
+		int start, int finish){
+
+	dvc_ScaLBL_Update_GreyscalePotential<<<NBLOCKS,NTHREADS >>>(Map, Phi, Psi, Poro, Perm, alpha, W, start, finish);
+
+	cudaError_t err = cudaGetLastError();
+	if (cudaSuccess != err){
+		printf("CUDA error in ScaLBL_Update_GreyscalePotential: %s \n",cudaGetErrorString(err));
+	}
+}
+
 ////Model-2&3
 //extern "C" void ScaLBL_D3Q19_AAeven_GreyscaleColor(int *Map, double *dist, double *Aq, double *Bq, double *Den, 
 //        double *Phi,double *GreySolidGrad, double *Poros,double *Perm,double *Vel, 
diff --git a/models/GreyscaleColorModel.cpp b/models/GreyscaleColorModel.cpp
index d5ecea2d..2e2b2f4b 100644
--- a/models/GreyscaleColorModel.cpp
+++ b/models/GreyscaleColorModel.cpp
@@ -220,6 +220,7 @@ void ScaLBL_GreyscaleColorModel::ReadInput(){
 	
 }
 
+
 void ScaLBL_GreyscaleColorModel::AssignComponentLabels()
 {
     // Initialize impermeability solid nodes and grey nodes
@@ -264,6 +265,7 @@ void ScaLBL_GreyscaleColorModel::AssignComponentLabels()
 				      //printf("idx=%i, value=%i, %i, \n",idx, VALUE,LabelList[idx]);
 					if (VALUE == LabelList[idx]){
 						AFFINITY=AffinityList[idx];
+
 						label_count[idx] += 1.0;
 						idx = NLABELS;
 						//Mask->id[n] = 0; // set mask to zero since this is an immobile component
@@ -583,6 +585,71 @@ void ScaLBL_GreyscaleColorModel::AssignGreyPoroPermLabels()
     delete [] Permeability;
 }
 
+void ScaLBL_GreyscaleColorModel::AssignGreyscalePotential()
+{
+    double *psi;//greyscale potential
+	psi = new double[N];
+
+	size_t NLABELS=0;
+	signed char VALUE=0;
+	double AFFINITY=0.f;
+
+	auto LabelList = greyscaleColor_db->getVector<int>( "ComponentLabels" );
+	auto AffinityList = greyscaleColor_db->getVector<double>( "ComponentAffinity" );
+	NLABELS=LabelList.size();
+
+    //first, copy over normal phase field
+	for (int k=0;k<Nz;k++){
+		for (int j=0;j<Ny;j++){
+			for (int i=0;i<Nx;i++){
+				int n = k*Nx*Ny+j*Nx+i;
+				VALUE=id[n];
+				// Assign the affinity from the paired list
+				for (unsigned int idx=0; idx < NLABELS; idx++){
+				      //printf("idx=%i, value=%i, %i, \n",idx, VALUE,LabelList[idx]);
+					if (VALUE == LabelList[idx]){
+						AFFINITY=AffinityList[idx];
+						idx = NLABELS;
+					}
+				}
+				// fluid labels are reserved
+				if (VALUE == 1) AFFINITY=1.0;
+				else if (VALUE == 2) AFFINITY=-1.0;
+				psi[n] = AFFINITY;
+			}
+		}
+	}
+
+    //second, scale the phase field for grey nodes
+	double Cap_Penalty=1.f;
+	auto GreyLabelList = greyscaleColor_db->getVector<int>( "GreySolidLabels" );
+	auto PermeabilityList = greyscaleColor_db->getVector<double>( "PermeabilityList" );
+	NLABELS=GreyLabelList.size();
+
+	for (int k=0;k<Nz;k++){
+		for (int j=0;j<Ny;j++){
+			for (int i=0;i<Nx;i++){
+				int n = k*Nx*Ny+j*Nx+i;
+				VALUE=id[n];
+                Cap_Penalty=1.f;
+				// Assign the affinity from the paired list
+				for (unsigned int idx=0; idx < NLABELS; idx++){
+					if (VALUE == GreyLabelList[idx]){
+						Cap_Penalty=alpha*W/sqrt(PermeabilityList[idx]/Dm->voxel_length/Dm->voxel_length);
+						idx = NLABELS;
+					}
+				}
+                //update greyscale potential
+                psi[n] = psi[n]*Cap_Penalty;
+			}
+		}
+	}
+
+	ScaLBL_CopyToDevice(Psi, psi, N*sizeof(double));
+	ScaLBL_Comm->Barrier();
+    delete [] psi;
+}
+
 void ScaLBL_GreyscaleColorModel::Create(){
 	/*
 	 *  This function creates the variables needed to run a LBM 
@@ -603,6 +670,7 @@ void ScaLBL_GreyscaleColorModel::Create(){
 	// ScaLBL_Communicator ScaLBL_Comm(Mask); // original
 	ScaLBL_Comm  = std::shared_ptr<ScaLBL_Communicator>(new ScaLBL_Communicator(Mask));
 	ScaLBL_Comm_Regular  = std::shared_ptr<ScaLBL_Communicator>(new ScaLBL_Communicator(Mask));
+	ScaLBL_Comm_Regular_2  = std::shared_ptr<ScaLBL_Communicator>(new ScaLBL_Communicator(Mask));
 
 	int Npad=(Np/16 + 2)*16;
 	if (rank==0)    printf ("Set up memory efficient layout, %i | %i | %i \n", Np, Npad, N);
@@ -627,6 +695,7 @@ void ScaLBL_GreyscaleColorModel::Create(){
 	ScaLBL_AllocateDeviceMemory((void **) &Bq, 7*dist_mem_size);
 	ScaLBL_AllocateDeviceMemory((void **) &Den, 2*dist_mem_size);
 	ScaLBL_AllocateDeviceMemory((void **) &Phi, sizeof(double)*Nx*Ny*Nz);		
+	ScaLBL_AllocateDeviceMemory((void **) &Psi, sizeof(double)*Nx*Ny*Nz);//greyscale potential		
 	ScaLBL_AllocateDeviceMemory((void **) &Pressure, sizeof(double)*Np);
 	ScaLBL_AllocateDeviceMemory((void **) &Velocity, 3*sizeof(double)*Np);
 	//ScaLBL_AllocateDeviceMemory((void **) &ColorGrad, 3*sizeof(double)*Np);
@@ -675,6 +744,7 @@ void ScaLBL_GreyscaleColorModel::Create(){
 	AssignComponentLabels();//do open/black/grey nodes initialization
     AssignGreySolidLabels();
     AssignGreyPoroPermLabels(); 
+    AssignGreyscalePotential(); 
 	Averages->SetParams(rhoA,rhoB,tauA,tauB,Fx,Fy,Fz,alpha,beta,GreyPorosity);
 	ScaLBL_Comm->RegularLayout(Map,Porosity_dvc,Averages->Porosity);//porosity doesn't change over time
 }        
@@ -939,9 +1009,11 @@ void ScaLBL_GreyscaleColorModel::Run(){
 		// Read for Aq, Bq happens in this routine (requires communication)
 		ScaLBL_Comm->BiSendD3Q7AA(Aq,Bq); //READ FROM NORMAL
 		ScaLBL_D3Q7_AAodd_PhaseField(NeighborList, dvcMap, Aq, Bq, Den, Phi, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
+        ScaLBL_Update_GreyscalePotential(dvcMap,Phi,Psi,Porosity_dvc,Permeability_dvc,alpha,W,ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior());
 		ScaLBL_Comm->BiRecvD3Q7AA(Aq,Bq); //WRITE INTO OPPOSITE
 		ScaLBL_Comm->Barrier();
 		ScaLBL_D3Q7_AAodd_PhaseField(NeighborList, dvcMap, Aq, Bq, Den, Phi, 0, ScaLBL_Comm->LastExterior(), Np);
+        ScaLBL_Update_GreyscalePotential(dvcMap,Phi,Psi,Porosity_dvc,Permeability_dvc,alpha,W,0,ScaLBL_Comm->LastExterior());
 
 		// Perform the collision operation
 		ScaLBL_Comm->SendD3Q19AA(fq); //READ FROM NORMAL
@@ -951,10 +1023,11 @@ void ScaLBL_GreyscaleColorModel::Run(){
 		}
 		// Halo exchange for phase field
 		ScaLBL_Comm_Regular->SendHalo(Phi);
+		ScaLBL_Comm_Regular_2->SendHalo(Psi);
         //Model-1&4 with capillary pressure penalty for grey nodes
-        ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(NeighborList, dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
+        ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(NeighborList, dvcMap, fq, Aq, Bq, Den, Phi, Psi, GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
                 rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff, 
-                alpha, beta, Fx, Fy, Fz, RecoloringOff, W, Nx, Nx*Ny, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
+                alpha, beta, Fx, Fy, Fz, RecoloringOff, Nx, Nx*Ny, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
         //Model-1&4
         //ScaLBL_D3Q19_AAodd_GreyscaleColor(NeighborList, dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
         //        rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff, 
@@ -963,6 +1036,7 @@ void ScaLBL_GreyscaleColorModel::Run(){
         //ScaLBL_D3Q19_AAodd_GreyscaleColor(NeighborList, dvcMap, fq, Aq, Bq, Den, Phi,GreySolidPhi,Porosity_dvc,Permeability_dvc,Velocity, 
         //        rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff, 
         //        alpha, beta, Fx, Fy, Fz, Nx, Nx*Ny, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
+		ScaLBL_Comm_Regular_2->RecvHalo(Psi);
 		ScaLBL_Comm_Regular->RecvHalo(Phi);
 		ScaLBL_Comm->RecvD3Q19AA(fq); //WRITE INTO OPPOSITE
 		ScaLBL_Comm->Barrier();
@@ -981,9 +1055,9 @@ void ScaLBL_GreyscaleColorModel::Run(){
 		}
 
         //Model-1&4 with capillary pressure penalty for grey nodes
-        ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(NeighborList, dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
+        ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(NeighborList, dvcMap, fq, Aq, Bq, Den, Phi, Psi, GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
                 rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
-                alpha, beta, Fx, Fy, Fz, RecoloringOff, W, Nx, Nx*Ny, 0, ScaLBL_Comm->LastExterior(), Np);
+                alpha, beta, Fx, Fy, Fz, RecoloringOff, Nx, Nx*Ny, 0, ScaLBL_Comm->LastExterior(), Np);
         //Model-1&4
         //ScaLBL_D3Q19_AAodd_GreyscaleColor(NeighborList, dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
         //        rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
@@ -999,9 +1073,11 @@ void ScaLBL_GreyscaleColorModel::Run(){
 		// Compute the Phase indicator field
 		ScaLBL_Comm->BiSendD3Q7AA(Aq,Bq); //READ FROM NORMAL
 		ScaLBL_D3Q7_AAeven_PhaseField(dvcMap, Aq, Bq, Den, Phi, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
+        ScaLBL_Update_GreyscalePotential(dvcMap,Phi,Psi,Porosity_dvc,Permeability_dvc,alpha,W,ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior());
 		ScaLBL_Comm->BiRecvD3Q7AA(Aq,Bq); //WRITE INTO OPPOSITE
 		ScaLBL_Comm->Barrier();
 		ScaLBL_D3Q7_AAeven_PhaseField(dvcMap, Aq, Bq, Den, Phi, 0, ScaLBL_Comm->LastExterior(), Np);
+        ScaLBL_Update_GreyscalePotential(dvcMap,Phi,Psi,Porosity_dvc,Permeability_dvc,alpha,W,0,ScaLBL_Comm->LastExterior());
 
 		// Perform the collision operation
 		ScaLBL_Comm->SendD3Q19AA(fq); //READ FORM NORMAL
@@ -1011,10 +1087,11 @@ void ScaLBL_GreyscaleColorModel::Run(){
 			ScaLBL_Comm->Color_BC_Z(dvcMap, Phi, Den, outletA, outletB);
 		}
 		ScaLBL_Comm_Regular->SendHalo(Phi);
+		ScaLBL_Comm_Regular_2->SendHalo(Psi);
         //Model-1&4 with capillary pressure penalty for grey nodes
-        ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure, 
+        ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(dvcMap, fq, Aq, Bq, Den, Phi, Psi, GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure, 
                 rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
-                alpha, beta, Fx, Fy, Fz, RecoloringOff, W, Nx, Nx*Ny, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
+                alpha, beta, Fx, Fy, Fz, RecoloringOff, Nx, Nx*Ny, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
         //Model-1&4
         //ScaLBL_D3Q19_AAeven_GreyscaleColor(dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure, 
         //        rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
@@ -1023,6 +1100,7 @@ void ScaLBL_GreyscaleColorModel::Run(){
         //ScaLBL_D3Q19_AAeven_GreyscaleColor(dvcMap, fq, Aq, Bq, Den, Phi,GreySolidPhi,Porosity_dvc,Permeability_dvc,Velocity, 
         //        rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
         //        alpha, beta, Fx, Fy, Fz,  Nx, Nx*Ny, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
+		ScaLBL_Comm_Regular_2->RecvHalo(Psi);
 		ScaLBL_Comm_Regular->RecvHalo(Phi);
 		ScaLBL_Comm->RecvD3Q19AA(fq); //WRITE INTO OPPOSITE
 		ScaLBL_Comm->Barrier();
@@ -1041,9 +1119,9 @@ void ScaLBL_GreyscaleColorModel::Run(){
 		}
 
         //Model-1&4 with capillary pressure penalty for grey nodes
-        ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
+        ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(dvcMap, fq, Aq, Bq, Den, Phi, Psi, GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
                 rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
-                alpha, beta, Fx, Fy, Fz, RecoloringOff, W, Nx, Nx*Ny, 0, ScaLBL_Comm->LastExterior(), Np);
+                alpha, beta, Fx, Fy, Fz, RecoloringOff, Nx, Nx*Ny, 0, ScaLBL_Comm->LastExterior(), Np);
         //Model-1&4
         //ScaLBL_D3Q19_AAeven_GreyscaleColor(dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
         //        rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
diff --git a/models/GreyscaleColorModel.h b/models/GreyscaleColorModel.h
index dd25e613..8285ee56 100644
--- a/models/GreyscaleColorModel.h
+++ b/models/GreyscaleColorModel.h
@@ -88,6 +88,7 @@ private:
     void AssignComponentLabels();
     void AssignGreySolidLabels();
     void AssignGreyPoroPermLabels();
+    void AssignGreyscalePotential();
     void ImageInit(std::string filename);
     double MorphInit(const double beta, const double morph_delta);
     double SeedPhaseField(const double seed_water_in_oil);

From 1ff5951ce0dfb368a06f1bb88509f1d30326703e Mon Sep 17 00:00:00 2001
From: Rex Zhe Li <zhe.rex.li@gmail.com>
Date: Wed, 3 Mar 2021 20:53:36 -0500
Subject: [PATCH 04/10] build pass; model to be validated

---
 common/ScaLBL.h                | 10 +++++-----
 cuda/GreyscaleColor.cu         | 24 +++++++++++-------------
 models/GreyscaleColorModel.cpp |  8 ++++----
 models/GreyscaleColorModel.h   |  4 +++-
 4 files changed, 23 insertions(+), 23 deletions(-)

diff --git a/common/ScaLBL.h b/common/ScaLBL.h
index 762efb20..24a5e370 100644
--- a/common/ScaLBL.h
+++ b/common/ScaLBL.h
@@ -88,17 +88,17 @@ extern "C" void ScaLBL_D3Q19_AAodd_GreyscaleColor(int *d_neighborList, int *Map,
 		double Fx, double Fy, double Fz, int strideY, int strideZ, int start, int finish, int Np);
 
 extern "C" void ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dist, double *Aq, double *Bq, double *Den, 
-        double *Phi,double *GreySolidGrad, double *Poros,double *Perm,double *Vel, double *Pressure,
+        double *Phi,double *Psi, double *GreySolidGrad, double *Poros,double *Perm,double *Vel, double *Pressure,
         double rhoA, double rhoB, double tauA, double tauB,double tauA_eff,double tauB_eff, double alpha, double beta,
-		double Fx, double Fy, double Fz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np);
+		double Fx, double Fy, double Fz, bool RecoloringOff, int strideY, int strideZ, int start, int finish, int Np);
 
 extern "C" void ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *d_neighborList, int *Map, double *dist, double *Aq, double *Bq, double *Den, 
-		double *Phi, double *GreySolidGrad, double *Poros,double *Perm,double *Vel,double *Pressure, 
+		double *Phi,double *Psi, double *GreySolidGrad, double *Poros,double *Perm,double *Vel,double *Pressure, 
         double rhoA, double rhoB, double tauA, double tauB, double tauA_eff,double tauB_eff, double alpha, double beta,
-		double Fx, double Fy, double Fz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np);
+		double Fx, double Fy, double Fz, bool RecoloringOff, int strideY, int strideZ, int start, int finish, int Np);
 
 extern "C" void ScaLBL_Update_GreyscalePotential(int *Map, double *Phi, double *Psi, double *Poro, double *Perm, double alpha, double W, 
-		int start, int finish);
+		int start, int finish, int Np);
 
 // ION TRANSPORT MODEL
 
diff --git a/cuda/GreyscaleColor.cu b/cuda/GreyscaleColor.cu
index ee989bba..d9541cc4 100644
--- a/cuda/GreyscaleColor.cu
+++ b/cuda/GreyscaleColor.cu
@@ -1472,7 +1472,6 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 	double a1,b1,a2,b2,nAB,delta;
 	double C,nx,ny,nz; //color gradient magnitude and direction
 	double phi,tau,rho0,rlx_setA,rlx_setB;
-    double psi;//greyscale potential
 
     //double GeoFun=0.0;//geometric function from Guo's PRE 66, 036304 (2002)
     double porosity;
@@ -2256,7 +2255,6 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
 	double a1,b1,a2,b2,nAB,delta;
 	double C,nx,ny,nz; //color gradient magnitude and direction
 	double phi,tau,rho0,rlx_setA,rlx_setB;
-    double psi;//greyscale potential
 
     //double GeoFun=0.0;//geometric function from Guo's PRE 66, 036304 (2002)
     double porosity;
@@ -2974,7 +2972,7 @@ __global__ void dvc_ScaLBL_PhaseField_InitFromRestart(double *Den, double *Aq, d
 }
 
 __global__  void dvc_ScaLBL_Update_GreyscalePotential(int *Map, double *Phi, double *Psi, double *Poro, double *Perm, double alpha, double W, 
-                                                      int start, int finish){
+                                                      int start, int finish, int Np){
 	int idx,n;
     double phi,psi;
     double cap_penalty;
@@ -4521,12 +4519,12 @@ extern "C" void ScaLBL_PhaseField_InitFromRestart(double *Den, double *Aq, doubl
 
 //Model-1 & 4 with capillary pressure penalty
 extern "C" void ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dist, double *Aq, double *Bq, double *Den, 
-        double *Phi,double *GreySolidGrad, double *Poros,double *Perm,double *Vel, double *Pressure,
+        double *Phi,double *Psi, double *GreySolidGrad, double *Poros,double *Perm,double *Vel, double *Pressure,
         double rhoA, double rhoB, double tauA, double tauB,double tauA_eff,double tauB_eff, double alpha, double beta,
-		double Fx, double Fy, double Fz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np){
+		double Fx, double Fy, double Fz, bool RecoloringOff, int strideY, int strideZ, int start, int finish, int Np){
 
-	dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP<<<NBLOCKS,NTHREADS >>>(Map, dist, Aq, Bq, Den, Phi, GreySolidGrad, Poros, Perm, Vel, Pressure,
-            rhoA, rhoB, tauA, tauB, tauA_eff, tauB_eff, alpha, beta, Fx, Fy, Fz, RecoloringOff, W, strideY, strideZ, start, finish, Np);
+	dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP<<<NBLOCKS,NTHREADS >>>(Map, dist, Aq, Bq, Den, Phi, Psi, GreySolidGrad, Poros, Perm, Vel, Pressure,
+            rhoA, rhoB, tauA, tauB, tauA_eff, tauB_eff, alpha, beta, Fx, Fy, Fz, RecoloringOff, strideY, strideZ, start, finish, Np);
 	cudaError_t err = cudaGetLastError();
 	if (cudaSuccess != err){
 		printf("CUDA error in ScaLBL_D3Q19_AAeven_GreyscaleColor_CP: %s \n",cudaGetErrorString(err));
@@ -4536,12 +4534,12 @@ extern "C" void ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dist, do
 
 //Model-1 & 4 with capillary pressure penalty
 extern "C" void ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *d_neighborList, int *Map, double *dist, double *Aq, double *Bq, double *Den, 
-		double *Phi, double *GreySolidGrad, double *Poros,double *Perm,double *Vel,double *Pressure, 
+		double *Phi, double *Psi, double *GreySolidGrad, double *Poros,double *Perm,double *Vel,double *Pressure, 
         double rhoA, double rhoB, double tauA, double tauB, double tauA_eff,double tauB_eff, double alpha, double beta,
-		double Fx, double Fy, double Fz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np){
+		double Fx, double Fy, double Fz, bool RecoloringOff, int strideY, int strideZ, int start, int finish, int Np){
 
-	dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP<<<NBLOCKS,NTHREADS >>>(d_neighborList, Map, dist, Aq, Bq, Den, Phi,  GreySolidGrad, Poros, Perm,Vel,Pressure,
-			rhoA, rhoB, tauA, tauB, tauA_eff, tauB_eff,alpha, beta, Fx, Fy, Fz, RecoloringOff, W, strideY, strideZ, start, finish, Np);
+	dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP<<<NBLOCKS,NTHREADS >>>(d_neighborList, Map, dist, Aq, Bq, Den, Phi, Psi, GreySolidGrad, Poros, Perm,Vel,Pressure,
+			rhoA, rhoB, tauA, tauB, tauA_eff, tauB_eff,alpha, beta, Fx, Fy, Fz, RecoloringOff, strideY, strideZ, start, finish, Np);
 
 	cudaError_t err = cudaGetLastError();
 	if (cudaSuccess != err){
@@ -4550,9 +4548,9 @@ extern "C" void ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *d_neighborList, int *M
 }
 
 extern "C" void ScaLBL_Update_GreyscalePotential(int *Map, double *Phi, double *Psi, double *Poro, double *Perm, double alpha, double W, 
-		int start, int finish){
+		int start, int finish, int Np){
 
-	dvc_ScaLBL_Update_GreyscalePotential<<<NBLOCKS,NTHREADS >>>(Map, Phi, Psi, Poro, Perm, alpha, W, start, finish);
+	dvc_ScaLBL_Update_GreyscalePotential<<<NBLOCKS,NTHREADS >>>(Map, Phi, Psi, Poro, Perm, alpha, W, start, finish, Np);
 
 	cudaError_t err = cudaGetLastError();
 	if (cudaSuccess != err){
diff --git a/models/GreyscaleColorModel.cpp b/models/GreyscaleColorModel.cpp
index 2e2b2f4b..f4cc6f95 100644
--- a/models/GreyscaleColorModel.cpp
+++ b/models/GreyscaleColorModel.cpp
@@ -1009,11 +1009,11 @@ void ScaLBL_GreyscaleColorModel::Run(){
 		// Read for Aq, Bq happens in this routine (requires communication)
 		ScaLBL_Comm->BiSendD3Q7AA(Aq,Bq); //READ FROM NORMAL
 		ScaLBL_D3Q7_AAodd_PhaseField(NeighborList, dvcMap, Aq, Bq, Den, Phi, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
-        ScaLBL_Update_GreyscalePotential(dvcMap,Phi,Psi,Porosity_dvc,Permeability_dvc,alpha,W,ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior());
+        ScaLBL_Update_GreyscalePotential(dvcMap,Phi,Psi,Porosity_dvc,Permeability_dvc,alpha,W,ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
 		ScaLBL_Comm->BiRecvD3Q7AA(Aq,Bq); //WRITE INTO OPPOSITE
 		ScaLBL_Comm->Barrier();
 		ScaLBL_D3Q7_AAodd_PhaseField(NeighborList, dvcMap, Aq, Bq, Den, Phi, 0, ScaLBL_Comm->LastExterior(), Np);
-        ScaLBL_Update_GreyscalePotential(dvcMap,Phi,Psi,Porosity_dvc,Permeability_dvc,alpha,W,0,ScaLBL_Comm->LastExterior());
+        ScaLBL_Update_GreyscalePotential(dvcMap,Phi,Psi,Porosity_dvc,Permeability_dvc,alpha,W,0,ScaLBL_Comm->LastExterior(), Np);
 
 		// Perform the collision operation
 		ScaLBL_Comm->SendD3Q19AA(fq); //READ FROM NORMAL
@@ -1073,11 +1073,11 @@ void ScaLBL_GreyscaleColorModel::Run(){
 		// Compute the Phase indicator field
 		ScaLBL_Comm->BiSendD3Q7AA(Aq,Bq); //READ FROM NORMAL
 		ScaLBL_D3Q7_AAeven_PhaseField(dvcMap, Aq, Bq, Den, Phi, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
-        ScaLBL_Update_GreyscalePotential(dvcMap,Phi,Psi,Porosity_dvc,Permeability_dvc,alpha,W,ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior());
+        ScaLBL_Update_GreyscalePotential(dvcMap,Phi,Psi,Porosity_dvc,Permeability_dvc,alpha,W,ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
 		ScaLBL_Comm->BiRecvD3Q7AA(Aq,Bq); //WRITE INTO OPPOSITE
 		ScaLBL_Comm->Barrier();
 		ScaLBL_D3Q7_AAeven_PhaseField(dvcMap, Aq, Bq, Den, Phi, 0, ScaLBL_Comm->LastExterior(), Np);
-        ScaLBL_Update_GreyscalePotential(dvcMap,Phi,Psi,Porosity_dvc,Permeability_dvc,alpha,W,0,ScaLBL_Comm->LastExterior());
+        ScaLBL_Update_GreyscalePotential(dvcMap,Phi,Psi,Porosity_dvc,Permeability_dvc,alpha,W,0,ScaLBL_Comm->LastExterior(), Np);
 
 		// Perform the collision operation
 		ScaLBL_Comm->SendD3Q19AA(fq); //READ FORM NORMAL
diff --git a/models/GreyscaleColorModel.h b/models/GreyscaleColorModel.h
index 8285ee56..be17c910 100644
--- a/models/GreyscaleColorModel.h
+++ b/models/GreyscaleColorModel.h
@@ -50,7 +50,8 @@ public:
 	std::shared_ptr<Domain> Mask; // this domain is for lbm
 	std::shared_ptr<ScaLBL_Communicator> ScaLBL_Comm;
 	std::shared_ptr<ScaLBL_Communicator> ScaLBL_Comm_Regular;
-        std::shared_ptr<GreyPhaseAnalysis> Averages;
+	std::shared_ptr<ScaLBL_Communicator> ScaLBL_Comm_Regular_2;
+    std::shared_ptr<GreyPhaseAnalysis> Averages;
     
     // input database
     std::shared_ptr<Database> db;
@@ -72,6 +73,7 @@ public:
 	double *Pressure;
     double *Porosity_dvc;
     double *Permeability_dvc;
+    double *Psi;
 		
 private:
 	Utilities::MPI comm;

From 5645f6125fe4c1e5c017956d3311db26de725564 Mon Sep 17 00:00:00 2001
From: Rex Zhe Li <zhe.rex.li@gmail.com>
Date: Mon, 8 Mar 2021 20:39:37 -0500
Subject: [PATCH 05/10] add debug output for greyscale potential

---
 models/GreyscaleColorModel.cpp | 7 +++++++
 1 file changed, 7 insertions(+)

diff --git a/models/GreyscaleColorModel.cpp b/models/GreyscaleColorModel.cpp
index f4cc6f95..de51270a 100644
--- a/models/GreyscaleColorModel.cpp
+++ b/models/GreyscaleColorModel.cpp
@@ -1677,6 +1677,13 @@ void ScaLBL_GreyscaleColorModel::WriteDebug(){
 	fwrite(PhaseField.data(),8,N,OUTFILE);
 	fclose(OUTFILE);
 
+	ScaLBL_CopyToHost(PhaseField.data(), Psi, sizeof(double)*N);
+	FILE *PSIFILE;
+	sprintf(LocalRankFilename,"Psi.%05i.raw",rank);
+	PSIFILE = fopen(LocalRankFilename,"wb");
+	fwrite(PhaseField.data(),8,N,PSIFILE);
+	fclose(PSIFILE);
+
     ScaLBL_Comm->RegularLayout(Map,&Den[0],PhaseField);
 	FILE *AFILE;
 	sprintf(LocalRankFilename,"A.%05i.raw",rank);

From 3ed949efb1e680b9b77529c95348732f34919e7b Mon Sep 17 00:00:00 2001
From: Rex Zhe Li <zhe.rex.li@gmail.com>
Date: Mon, 8 Mar 2021 22:28:19 -0500
Subject: [PATCH 06/10] slightly change the implementation of capillary
 penalty, to be built and tested

---
 common/ScaLBL.h                |  4 +--
 cuda/GreyscaleColor.cu         | 46 +++++++++++++++++++++++++---------
 models/GreyscaleColorModel.cpp | 16 ++++++------
 3 files changed, 44 insertions(+), 22 deletions(-)

diff --git a/common/ScaLBL.h b/common/ScaLBL.h
index 24a5e370..5c4d4730 100644
--- a/common/ScaLBL.h
+++ b/common/ScaLBL.h
@@ -90,12 +90,12 @@ extern "C" void ScaLBL_D3Q19_AAodd_GreyscaleColor(int *d_neighborList, int *Map,
 extern "C" void ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dist, double *Aq, double *Bq, double *Den, 
         double *Phi,double *Psi, double *GreySolidGrad, double *Poros,double *Perm,double *Vel, double *Pressure,
         double rhoA, double rhoB, double tauA, double tauB,double tauA_eff,double tauB_eff, double alpha, double beta,
-		double Fx, double Fy, double Fz, bool RecoloringOff, int strideY, int strideZ, int start, int finish, int Np);
+		double Fx, double Fy, double Fz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np);
 
 extern "C" void ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *d_neighborList, int *Map, double *dist, double *Aq, double *Bq, double *Den, 
 		double *Phi,double *Psi, double *GreySolidGrad, double *Poros,double *Perm,double *Vel,double *Pressure, 
         double rhoA, double rhoB, double tauA, double tauB, double tauA_eff,double tauB_eff, double alpha, double beta,
-		double Fx, double Fy, double Fz, bool RecoloringOff, int strideY, int strideZ, int start, int finish, int Np);
+		double Fx, double Fy, double Fz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np);
 
 extern "C" void ScaLBL_Update_GreyscalePotential(int *Map, double *Phi, double *Psi, double *Poro, double *Perm, double alpha, double W, 
 		int start, int finish, int Np);
diff --git a/cuda/GreyscaleColor.cu b/cuda/GreyscaleColor.cu
index d9541cc4..e3a4c421 100644
--- a/cuda/GreyscaleColor.cu
+++ b/cuda/GreyscaleColor.cu
@@ -1452,7 +1452,7 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor(int *Map, double *dist,
 __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int *Map, double *dist, double *Aq, double *Bq, double *Den,
 		 double *Phi, double *Psi, double *GreySolidGrad, double *Poros,double *Perm, double *Velocity, double *Pressure,
          double rhoA, double rhoB, double tauA, double tauB,double tauA_eff,double tauB_eff,double alpha, double beta,
-		double Gx, double Gy, double Gz, bool RecoloringOff, int strideY, int strideZ, int start, int finish, int Np){
+		double Gx, double Gy, double Gz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np){
 
 	int n,nn,ijk,nread;
 	int nr1,nr2,nr3,nr4,nr5,nr6;
@@ -1609,9 +1609,20 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 			nz_phase = -(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
 			C_phase = sqrt(nx_phase*nx_phase+ny_phase*ny_phase+nz_phase*nz_phase);
 			//............Compute the Greyscale Potential Gradient.....................
-			Fcpx = -3.0/18.0*(gp1-gp2+0.5*(gp7-gp8+gp9-gp10+gp11-gp12+gp13-gp14));
-			Fcpy = -3.0/18.0*(gp3-gp4+0.5*(gp7-gp8-gp9+gp10+gp15-gp16+gp17-gp18));
-			Fcpz = -3.0/18.0*(gp5-gp6+0.5*(gp11-gp12-gp13+gp14+gp15-gp16-gp17+gp18));
+			Fcpx = 0.0;
+			Fcpy = 0.0;
+			Fcpz = 0.0;
+            if (porosity!=1.0){
+                //Fcpx = -3.0/18.0*(gp1-gp2+0.5*(gp7-gp8+gp9-gp10+gp11-gp12+gp13-gp14));
+                //Fcpy = -3.0/18.0*(gp3-gp4+0.5*(gp7-gp8-gp9+gp10+gp15-gp16+gp17-gp18));
+                //Fcpz = -3.0/18.0*(gp5-gp6+0.5*(gp11-gp12-gp13+gp14+gp15-gp16-gp17+gp18));
+                Fcpx = -3.0/18.0*(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
+                Fcpy = -3.0/18.0*(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
+                Fcpz = -3.0/18.0*(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
+                Fcpx *= alpha*W/sqrt(perm); 
+                Fcpy *= alpha*W/sqrt(perm); 
+                Fcpz *= alpha*W/sqrt(perm); 
+            }
 
 
             //correct the normal color gradient by considering the effect of grey solid
@@ -2240,7 +2251,7 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dist, double *Aq, double *Bq, double *Den, 
         double *Phi, double *Psi, double *GreySolidGrad, double *Poros,double *Perm, double *Velocity, double *Pressure, 
         double rhoA, double rhoB, double tauA, double tauB,double tauA_eff,double tauB_eff, double alpha, double beta,
-		double Gx, double Gy, double Gz, bool RecoloringOff, int strideY, int strideZ, int start, int finish, int Np){
+		double Gx, double Gy, double Gz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np){
 	int ijk,nn,n;
 	double fq;
 	// conserved momemnts
@@ -2392,9 +2403,20 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
 			nz_phase = -(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
 			C_phase = sqrt(nx_phase*nx_phase+ny_phase*ny_phase+nz_phase*nz_phase);
 			//............Compute the Greyscale Potential Gradient.....................
-			Fcpx = -3.0/18.0*(gp1-gp2+0.5*(gp7-gp8+gp9-gp10+gp11-gp12+gp13-gp14));
-			Fcpy = -3.0/18.0*(gp3-gp4+0.5*(gp7-gp8-gp9+gp10+gp15-gp16+gp17-gp18));
-			Fcpz = -3.0/18.0*(gp5-gp6+0.5*(gp11-gp12-gp13+gp14+gp15-gp16-gp17+gp18));
+			Fcpx = 0.0;
+			Fcpy = 0.0;
+			Fcpz = 0.0;
+            if (porosity!=1.0){
+                //Fcpx = -3.0/18.0*(gp1-gp2+0.5*(gp7-gp8+gp9-gp10+gp11-gp12+gp13-gp14));
+                //Fcpy = -3.0/18.0*(gp3-gp4+0.5*(gp7-gp8-gp9+gp10+gp15-gp16+gp17-gp18));
+                //Fcpz = -3.0/18.0*(gp5-gp6+0.5*(gp11-gp12-gp13+gp14+gp15-gp16-gp17+gp18));
+                Fcpx = -3.0/18.0*(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
+                Fcpy = -3.0/18.0*(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
+                Fcpz = -3.0/18.0*(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
+                Fcpx *= alpha*W/sqrt(perm); 
+                Fcpy *= alpha*W/sqrt(perm); 
+                Fcpz *= alpha*W/sqrt(perm); 
+            }
 
             //correct the normal color gradient by considering the effect of grey solid
             nx = nx_phase + (1.0-porosity)*nx_gs; 
@@ -4521,10 +4543,10 @@ extern "C" void ScaLBL_PhaseField_InitFromRestart(double *Den, double *Aq, doubl
 extern "C" void ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dist, double *Aq, double *Bq, double *Den, 
         double *Phi,double *Psi, double *GreySolidGrad, double *Poros,double *Perm,double *Vel, double *Pressure,
         double rhoA, double rhoB, double tauA, double tauB,double tauA_eff,double tauB_eff, double alpha, double beta,
-		double Fx, double Fy, double Fz, bool RecoloringOff, int strideY, int strideZ, int start, int finish, int Np){
+		double Fx, double Fy, double Fz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np){
 
 	dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP<<<NBLOCKS,NTHREADS >>>(Map, dist, Aq, Bq, Den, Phi, Psi, GreySolidGrad, Poros, Perm, Vel, Pressure,
-            rhoA, rhoB, tauA, tauB, tauA_eff, tauB_eff, alpha, beta, Fx, Fy, Fz, RecoloringOff, strideY, strideZ, start, finish, Np);
+            rhoA, rhoB, tauA, tauB, tauA_eff, tauB_eff, alpha, beta, Fx, Fy, Fz, RecoloringOff, W, strideY, strideZ, start, finish, Np);
 	cudaError_t err = cudaGetLastError();
 	if (cudaSuccess != err){
 		printf("CUDA error in ScaLBL_D3Q19_AAeven_GreyscaleColor_CP: %s \n",cudaGetErrorString(err));
@@ -4536,10 +4558,10 @@ extern "C" void ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dist, do
 extern "C" void ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *d_neighborList, int *Map, double *dist, double *Aq, double *Bq, double *Den, 
 		double *Phi, double *Psi, double *GreySolidGrad, double *Poros,double *Perm,double *Vel,double *Pressure, 
         double rhoA, double rhoB, double tauA, double tauB, double tauA_eff,double tauB_eff, double alpha, double beta,
-		double Fx, double Fy, double Fz, bool RecoloringOff, int strideY, int strideZ, int start, int finish, int Np){
+		double Fx, double Fy, double Fz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np){
 
 	dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP<<<NBLOCKS,NTHREADS >>>(d_neighborList, Map, dist, Aq, Bq, Den, Phi, Psi, GreySolidGrad, Poros, Perm,Vel,Pressure,
-			rhoA, rhoB, tauA, tauB, tauA_eff, tauB_eff,alpha, beta, Fx, Fy, Fz, RecoloringOff, strideY, strideZ, start, finish, Np);
+			rhoA, rhoB, tauA, tauB, tauA_eff, tauB_eff,alpha, beta, Fx, Fy, Fz, RecoloringOff, W, strideY, strideZ, start, finish, Np);
 
 	cudaError_t err = cudaGetLastError();
 	if (cudaSuccess != err){
diff --git a/models/GreyscaleColorModel.cpp b/models/GreyscaleColorModel.cpp
index de51270a..a4bb607e 100644
--- a/models/GreyscaleColorModel.cpp
+++ b/models/GreyscaleColorModel.cpp
@@ -1009,11 +1009,11 @@ void ScaLBL_GreyscaleColorModel::Run(){
 		// Read for Aq, Bq happens in this routine (requires communication)
 		ScaLBL_Comm->BiSendD3Q7AA(Aq,Bq); //READ FROM NORMAL
 		ScaLBL_D3Q7_AAodd_PhaseField(NeighborList, dvcMap, Aq, Bq, Den, Phi, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
-        ScaLBL_Update_GreyscalePotential(dvcMap,Phi,Psi,Porosity_dvc,Permeability_dvc,alpha,W,ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
+        //ScaLBL_Update_GreyscalePotential(dvcMap,Phi,Psi,Porosity_dvc,Permeability_dvc,alpha,W,ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
 		ScaLBL_Comm->BiRecvD3Q7AA(Aq,Bq); //WRITE INTO OPPOSITE
 		ScaLBL_Comm->Barrier();
 		ScaLBL_D3Q7_AAodd_PhaseField(NeighborList, dvcMap, Aq, Bq, Den, Phi, 0, ScaLBL_Comm->LastExterior(), Np);
-        ScaLBL_Update_GreyscalePotential(dvcMap,Phi,Psi,Porosity_dvc,Permeability_dvc,alpha,W,0,ScaLBL_Comm->LastExterior(), Np);
+        //ScaLBL_Update_GreyscalePotential(dvcMap,Phi,Psi,Porosity_dvc,Permeability_dvc,alpha,W,0,ScaLBL_Comm->LastExterior(), Np);
 
 		// Perform the collision operation
 		ScaLBL_Comm->SendD3Q19AA(fq); //READ FROM NORMAL
@@ -1027,7 +1027,7 @@ void ScaLBL_GreyscaleColorModel::Run(){
         //Model-1&4 with capillary pressure penalty for grey nodes
         ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(NeighborList, dvcMap, fq, Aq, Bq, Den, Phi, Psi, GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
                 rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff, 
-                alpha, beta, Fx, Fy, Fz, RecoloringOff, Nx, Nx*Ny, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
+                alpha, beta, Fx, Fy, Fz, RecoloringOff, W, Nx, Nx*Ny, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
         //Model-1&4
         //ScaLBL_D3Q19_AAodd_GreyscaleColor(NeighborList, dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
         //        rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff, 
@@ -1057,7 +1057,7 @@ void ScaLBL_GreyscaleColorModel::Run(){
         //Model-1&4 with capillary pressure penalty for grey nodes
         ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(NeighborList, dvcMap, fq, Aq, Bq, Den, Phi, Psi, GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
                 rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
-                alpha, beta, Fx, Fy, Fz, RecoloringOff, Nx, Nx*Ny, 0, ScaLBL_Comm->LastExterior(), Np);
+                alpha, beta, Fx, Fy, Fz, RecoloringOff, W, Nx, Nx*Ny, 0, ScaLBL_Comm->LastExterior(), Np);
         //Model-1&4
         //ScaLBL_D3Q19_AAodd_GreyscaleColor(NeighborList, dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
         //        rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
@@ -1073,11 +1073,11 @@ void ScaLBL_GreyscaleColorModel::Run(){
 		// Compute the Phase indicator field
 		ScaLBL_Comm->BiSendD3Q7AA(Aq,Bq); //READ FROM NORMAL
 		ScaLBL_D3Q7_AAeven_PhaseField(dvcMap, Aq, Bq, Den, Phi, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
-        ScaLBL_Update_GreyscalePotential(dvcMap,Phi,Psi,Porosity_dvc,Permeability_dvc,alpha,W,ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
+        //ScaLBL_Update_GreyscalePotential(dvcMap,Phi,Psi,Porosity_dvc,Permeability_dvc,alpha,W,ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
 		ScaLBL_Comm->BiRecvD3Q7AA(Aq,Bq); //WRITE INTO OPPOSITE
 		ScaLBL_Comm->Barrier();
 		ScaLBL_D3Q7_AAeven_PhaseField(dvcMap, Aq, Bq, Den, Phi, 0, ScaLBL_Comm->LastExterior(), Np);
-        ScaLBL_Update_GreyscalePotential(dvcMap,Phi,Psi,Porosity_dvc,Permeability_dvc,alpha,W,0,ScaLBL_Comm->LastExterior(), Np);
+        //ScaLBL_Update_GreyscalePotential(dvcMap,Phi,Psi,Porosity_dvc,Permeability_dvc,alpha,W,0,ScaLBL_Comm->LastExterior(), Np);
 
 		// Perform the collision operation
 		ScaLBL_Comm->SendD3Q19AA(fq); //READ FORM NORMAL
@@ -1091,7 +1091,7 @@ void ScaLBL_GreyscaleColorModel::Run(){
         //Model-1&4 with capillary pressure penalty for grey nodes
         ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(dvcMap, fq, Aq, Bq, Den, Phi, Psi, GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure, 
                 rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
-                alpha, beta, Fx, Fy, Fz, RecoloringOff, Nx, Nx*Ny, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
+                alpha, beta, Fx, Fy, Fz, RecoloringOff, W, Nx, Nx*Ny, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
         //Model-1&4
         //ScaLBL_D3Q19_AAeven_GreyscaleColor(dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure, 
         //        rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
@@ -1121,7 +1121,7 @@ void ScaLBL_GreyscaleColorModel::Run(){
         //Model-1&4 with capillary pressure penalty for grey nodes
         ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(dvcMap, fq, Aq, Bq, Den, Phi, Psi, GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
                 rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,
-                alpha, beta, Fx, Fy, Fz, RecoloringOff, Nx, Nx*Ny, 0, ScaLBL_Comm->LastExterior(), Np);
+                alpha, beta, Fx, Fy, Fz, RecoloringOff, W, Nx, Nx*Ny, 0, ScaLBL_Comm->LastExterior(), Np);
         //Model-1&4
         //ScaLBL_D3Q19_AAeven_GreyscaleColor(dvcMap, fq, Aq, Bq, Den, Phi,GreySolidGrad,Porosity_dvc,Permeability_dvc,Velocity,Pressure,
         //        rhoA, rhoB, tauA, tauB,tauA_eff, tauB_eff,

From 2cec6b260e8554ef826d2994ff8e0f6551e5441e Mon Sep 17 00:00:00 2001
From: Rex Zhe Li <zhe.rex.li@gmail.com>
Date: Wed, 10 Mar 2021 19:53:58 -0500
Subject: [PATCH 07/10] change color-gradient on greynodes to be the same as
 Fcp

---
 cuda/GreyscaleColor.cu | 88 ++++++++++++++++++++++++------------------
 1 file changed, 51 insertions(+), 37 deletions(-)

diff --git a/cuda/GreyscaleColor.cu b/cuda/GreyscaleColor.cu
index e3a4c421..01597f7f 100644
--- a/cuda/GreyscaleColor.cu
+++ b/cuda/GreyscaleColor.cu
@@ -1608,22 +1608,6 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 			ny_phase = -(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
 			nz_phase = -(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
 			C_phase = sqrt(nx_phase*nx_phase+ny_phase*ny_phase+nz_phase*nz_phase);
-			//............Compute the Greyscale Potential Gradient.....................
-			Fcpx = 0.0;
-			Fcpy = 0.0;
-			Fcpz = 0.0;
-            if (porosity!=1.0){
-                //Fcpx = -3.0/18.0*(gp1-gp2+0.5*(gp7-gp8+gp9-gp10+gp11-gp12+gp13-gp14));
-                //Fcpy = -3.0/18.0*(gp3-gp4+0.5*(gp7-gp8-gp9+gp10+gp15-gp16+gp17-gp18));
-                //Fcpz = -3.0/18.0*(gp5-gp6+0.5*(gp11-gp12-gp13+gp14+gp15-gp16-gp17+gp18));
-                Fcpx = -3.0/18.0*(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
-                Fcpy = -3.0/18.0*(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
-                Fcpz = -3.0/18.0*(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
-                Fcpx *= alpha*W/sqrt(perm); 
-                Fcpy *= alpha*W/sqrt(perm); 
-                Fcpz *= alpha*W/sqrt(perm); 
-            }
-
 
             //correct the normal color gradient by considering the effect of grey solid
             nx = nx_phase + (1.0-porosity)*nx_gs; 
@@ -1643,6 +1627,28 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 			ny = ny/ColorMag;
 			nz = nz/ColorMag;		
 
+			//............Compute the Greyscale Potential Gradient.....................
+			Fcpx = 0.0;
+			Fcpy = 0.0;
+			Fcpz = 0.0;
+            if (porosity!=1.0){
+                //Fcpx = -3.0/18.0*(gp1-gp2+0.5*(gp7-gp8+gp9-gp10+gp11-gp12+gp13-gp14));
+                //Fcpy = -3.0/18.0*(gp3-gp4+0.5*(gp7-gp8-gp9+gp10+gp15-gp16+gp17-gp18));
+                //Fcpz = -3.0/18.0*(gp5-gp6+0.5*(gp11-gp12-gp13+gp14+gp15-gp16-gp17+gp18));
+                Fcpx = -3.0/18.0*(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
+                Fcpy = -3.0/18.0*(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
+                Fcpz = -3.0/18.0*(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
+                Fcpx *= alpha*W/sqrt(perm); 
+                Fcpy *= alpha*W/sqrt(perm); 
+                Fcpz *= alpha*W/sqrt(perm); 
+                double Fcp_mag_temp = sqrt(Fcpx*Fcpx+Fcpy*Fcpy+Fcpz*Fcpz);
+                double Fcp_mag = Fcp_mag_temp;
+                if (Fcp_mag_temp==0.0) Fcp_mag=1.0;
+                nx = Fcpx/Fcp_mag;
+                ny = Fcpy/Fcp_mag;
+                nz = Fcpz/Fcp_mag;
+            }
+
 			// q=0
 			fq = dist[n];
 			rho = fq;
@@ -2188,7 +2194,7 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 			// Cq = {1,0,0}, {0,1,0}, {0,0,1}
 			delta = beta*nA*nB*nAB*0.1111111111111111*nx;
 			if (!(nA*nB*nAB>0)) delta=0;
-            if (RecoloringOff==true && porosity !=1.0) delta=0;
+            //if (RecoloringOff==true && porosity !=1.0) delta=0;
 			a1 = nA*(0.1111111111111111*(1+4.5*ux))+delta;
 			b1 = nB*(0.1111111111111111*(1+4.5*ux))-delta;
 			a2 = nA*(0.1111111111111111*(1-4.5*ux))-delta;
@@ -2207,7 +2213,7 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 			// Cq = {0,1,0}
 			delta = beta*nA*nB*nAB*0.1111111111111111*ny;
 			if (!(nA*nB*nAB>0)) delta=0;
-            if (RecoloringOff==true && porosity !=1.0) delta=0;
+            //if (RecoloringOff==true && porosity !=1.0) delta=0;
 			a1 = nA*(0.1111111111111111*(1+4.5*uy))+delta;
 			b1 = nB*(0.1111111111111111*(1+4.5*uy))-delta;
 			a2 = nA*(0.1111111111111111*(1-4.5*uy))-delta;
@@ -2227,7 +2233,7 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 			// Cq = {0,0,1}
 			delta = beta*nA*nB*nAB*0.1111111111111111*nz;
 			if (!(nA*nB*nAB>0)) delta=0;
-            if (RecoloringOff==true && porosity !=1.0) delta=0;
+            //if (RecoloringOff==true && porosity !=1.0) delta=0;
 			a1 = nA*(0.1111111111111111*(1+4.5*uz))+delta;
 			b1 = nB*(0.1111111111111111*(1+4.5*uz))-delta;
 			a2 = nA*(0.1111111111111111*(1-4.5*uz))-delta;
@@ -2402,21 +2408,6 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
 			ny_phase = -(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
 			nz_phase = -(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
 			C_phase = sqrt(nx_phase*nx_phase+ny_phase*ny_phase+nz_phase*nz_phase);
-			//............Compute the Greyscale Potential Gradient.....................
-			Fcpx = 0.0;
-			Fcpy = 0.0;
-			Fcpz = 0.0;
-            if (porosity!=1.0){
-                //Fcpx = -3.0/18.0*(gp1-gp2+0.5*(gp7-gp8+gp9-gp10+gp11-gp12+gp13-gp14));
-                //Fcpy = -3.0/18.0*(gp3-gp4+0.5*(gp7-gp8-gp9+gp10+gp15-gp16+gp17-gp18));
-                //Fcpz = -3.0/18.0*(gp5-gp6+0.5*(gp11-gp12-gp13+gp14+gp15-gp16-gp17+gp18));
-                Fcpx = -3.0/18.0*(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
-                Fcpy = -3.0/18.0*(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
-                Fcpz = -3.0/18.0*(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
-                Fcpx *= alpha*W/sqrt(perm); 
-                Fcpy *= alpha*W/sqrt(perm); 
-                Fcpz *= alpha*W/sqrt(perm); 
-            }
 
             //correct the normal color gradient by considering the effect of grey solid
             nx = nx_phase + (1.0-porosity)*nx_gs; 
@@ -2436,6 +2427,29 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
 			ny = ny/ColorMag;
 			nz = nz/ColorMag;		
 
+			//............Compute the Greyscale Potential Gradient.....................
+			Fcpx = 0.0;
+			Fcpy = 0.0;
+			Fcpz = 0.0;
+            if (porosity!=1.0){
+                //Fcpx = -3.0/18.0*(gp1-gp2+0.5*(gp7-gp8+gp9-gp10+gp11-gp12+gp13-gp14));
+                //Fcpy = -3.0/18.0*(gp3-gp4+0.5*(gp7-gp8-gp9+gp10+gp15-gp16+gp17-gp18));
+                //Fcpz = -3.0/18.0*(gp5-gp6+0.5*(gp11-gp12-gp13+gp14+gp15-gp16-gp17+gp18));
+                Fcpx = -3.0/18.0*(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
+                Fcpy = -3.0/18.0*(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
+                Fcpz = -3.0/18.0*(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
+                Fcpx *= alpha*W/sqrt(perm); 
+                Fcpy *= alpha*W/sqrt(perm); 
+                Fcpz *= alpha*W/sqrt(perm); 
+                double Fcp_mag_temp = sqrt(Fcpx*Fcpx+Fcpy*Fcpy+Fcpz*Fcpz);
+                double Fcp_mag = Fcp_mag_temp;
+                if (Fcp_mag_temp==0.0) Fcp_mag=1.0;
+                nx = Fcpx/Fcp_mag;
+                ny = Fcpy/Fcp_mag;
+                nz = Fcpz/Fcp_mag;
+            }
+
+
 			// q=0
 			fq = dist[n];
 			rho = fq;
@@ -2914,7 +2928,7 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
 			// Cq = {1,0,0}, {0,1,0}, {0,0,1}
 			delta = beta*nA*nB*nAB*0.1111111111111111*nx;
 			if (!(nA*nB*nAB>0)) delta=0;
-            if (RecoloringOff==true && porosity !=1.0) delta=0;
+            //if (RecoloringOff==true && porosity !=1.0) delta=0;
 			a1 = nA*(0.1111111111111111*(1+4.5*ux))+delta;
 			b1 = nB*(0.1111111111111111*(1+4.5*ux))-delta;
 			a2 = nA*(0.1111111111111111*(1-4.5*ux))-delta;
@@ -2930,7 +2944,7 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
 			// Cq = {0,1,0}
 			delta = beta*nA*nB*nAB*0.1111111111111111*ny;
 			if (!(nA*nB*nAB>0)) delta=0;
-            if (RecoloringOff==true && porosity !=1.0) delta=0;
+            //if (RecoloringOff==true && porosity !=1.0) delta=0;
 			a1 = nA*(0.1111111111111111*(1+4.5*uy))+delta;
 			b1 = nB*(0.1111111111111111*(1+4.5*uy))-delta;
 			a2 = nA*(0.1111111111111111*(1-4.5*uy))-delta;
@@ -2945,7 +2959,7 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
 			// Cq = {0,0,1}
 			delta = beta*nA*nB*nAB*0.1111111111111111*nz;
 			if (!(nA*nB*nAB>0)) delta=0;
-            if (RecoloringOff==true && porosity !=1.0) delta=0;
+            //if (RecoloringOff==true && porosity !=1.0) delta=0;
 			a1 = nA*(0.1111111111111111*(1+4.5*uz))+delta;
 			b1 = nB*(0.1111111111111111*(1+4.5*uz))-delta;
 			a2 = nA*(0.1111111111111111*(1-4.5*uz))-delta;

From 57158b539ec401341af426dd84648d88a21febdc Mon Sep 17 00:00:00 2001
From: Rex Zhe Li <zhe.rex.li@gmail.com>
Date: Thu, 11 Mar 2021 20:00:03 -0500
Subject: [PATCH 08/10] had accidentally comment out the recoloring check; put
 it back on

---
 cuda/GreyscaleColor.cu | 12 ++++++------
 1 file changed, 6 insertions(+), 6 deletions(-)

diff --git a/cuda/GreyscaleColor.cu b/cuda/GreyscaleColor.cu
index 01597f7f..5cc5d87e 100644
--- a/cuda/GreyscaleColor.cu
+++ b/cuda/GreyscaleColor.cu
@@ -2194,7 +2194,7 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 			// Cq = {1,0,0}, {0,1,0}, {0,0,1}
 			delta = beta*nA*nB*nAB*0.1111111111111111*nx;
 			if (!(nA*nB*nAB>0)) delta=0;
-            //if (RecoloringOff==true && porosity !=1.0) delta=0;
+            if (RecoloringOff==true && porosity !=1.0) delta=0;
 			a1 = nA*(0.1111111111111111*(1+4.5*ux))+delta;
 			b1 = nB*(0.1111111111111111*(1+4.5*ux))-delta;
 			a2 = nA*(0.1111111111111111*(1-4.5*ux))-delta;
@@ -2213,7 +2213,7 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 			// Cq = {0,1,0}
 			delta = beta*nA*nB*nAB*0.1111111111111111*ny;
 			if (!(nA*nB*nAB>0)) delta=0;
-            //if (RecoloringOff==true && porosity !=1.0) delta=0;
+            if (RecoloringOff==true && porosity !=1.0) delta=0;
 			a1 = nA*(0.1111111111111111*(1+4.5*uy))+delta;
 			b1 = nB*(0.1111111111111111*(1+4.5*uy))-delta;
 			a2 = nA*(0.1111111111111111*(1-4.5*uy))-delta;
@@ -2233,7 +2233,7 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 			// Cq = {0,0,1}
 			delta = beta*nA*nB*nAB*0.1111111111111111*nz;
 			if (!(nA*nB*nAB>0)) delta=0;
-            //if (RecoloringOff==true && porosity !=1.0) delta=0;
+            if (RecoloringOff==true && porosity !=1.0) delta=0;
 			a1 = nA*(0.1111111111111111*(1+4.5*uz))+delta;
 			b1 = nB*(0.1111111111111111*(1+4.5*uz))-delta;
 			a2 = nA*(0.1111111111111111*(1-4.5*uz))-delta;
@@ -2928,7 +2928,7 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
 			// Cq = {1,0,0}, {0,1,0}, {0,0,1}
 			delta = beta*nA*nB*nAB*0.1111111111111111*nx;
 			if (!(nA*nB*nAB>0)) delta=0;
-            //if (RecoloringOff==true && porosity !=1.0) delta=0;
+            if (RecoloringOff==true && porosity !=1.0) delta=0;
 			a1 = nA*(0.1111111111111111*(1+4.5*ux))+delta;
 			b1 = nB*(0.1111111111111111*(1+4.5*ux))-delta;
 			a2 = nA*(0.1111111111111111*(1-4.5*ux))-delta;
@@ -2944,7 +2944,7 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
 			// Cq = {0,1,0}
 			delta = beta*nA*nB*nAB*0.1111111111111111*ny;
 			if (!(nA*nB*nAB>0)) delta=0;
-            //if (RecoloringOff==true && porosity !=1.0) delta=0;
+            if (RecoloringOff==true && porosity !=1.0) delta=0;
 			a1 = nA*(0.1111111111111111*(1+4.5*uy))+delta;
 			b1 = nB*(0.1111111111111111*(1+4.5*uy))-delta;
 			a2 = nA*(0.1111111111111111*(1-4.5*uy))-delta;
@@ -2959,7 +2959,7 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
 			// Cq = {0,0,1}
 			delta = beta*nA*nB*nAB*0.1111111111111111*nz;
 			if (!(nA*nB*nAB>0)) delta=0;
-            //if (RecoloringOff==true && porosity !=1.0) delta=0;
+            if (RecoloringOff==true && porosity !=1.0) delta=0;
 			a1 = nA*(0.1111111111111111*(1+4.5*uz))+delta;
 			b1 = nB*(0.1111111111111111*(1+4.5*uz))-delta;
 			a2 = nA*(0.1111111111111111*(1-4.5*uz))-delta;

From fdf635bb57c11a7cb2b336393bf45a57a8a7d844 Mon Sep 17 00:00:00 2001
From: Rex Zhe Li <zhe.rex.li@gmail.com>
Date: Thu, 11 Mar 2021 21:11:57 -0500
Subject: [PATCH 09/10] revise the capillary penalty formulation

---
 cuda/GreyscaleColor.cu | 106 +++++++++++++++++++++++------------------
 1 file changed, 60 insertions(+), 46 deletions(-)

diff --git a/cuda/GreyscaleColor.cu b/cuda/GreyscaleColor.cu
index 5cc5d87e..089d8fdc 100644
--- a/cuda/GreyscaleColor.cu
+++ b/cuda/GreyscaleColor.cu
@@ -1604,9 +1604,9 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 			m18 = Phi[nn];					// get neighbor for phi - 18
             gp18 = Psi[nn];
 			//............Compute the Color Gradient...................................
-			nx_phase = -(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
-			ny_phase = -(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
-			nz_phase = -(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
+			nx_phase = -3.0/18.0*(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
+			ny_phase = -3.0/18.0*(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
+			nz_phase = -3.0/18.0*(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
 			C_phase = sqrt(nx_phase*nx_phase+ny_phase*ny_phase+nz_phase*nz_phase);
 
             //correct the normal color gradient by considering the effect of grey solid
@@ -1628,26 +1628,33 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int
 			nz = nz/ColorMag;		
 
 			//............Compute the Greyscale Potential Gradient.....................
-			Fcpx = 0.0;
-			Fcpy = 0.0;
-			Fcpz = 0.0;
-            if (porosity!=1.0){
-                //Fcpx = -3.0/18.0*(gp1-gp2+0.5*(gp7-gp8+gp9-gp10+gp11-gp12+gp13-gp14));
-                //Fcpy = -3.0/18.0*(gp3-gp4+0.5*(gp7-gp8-gp9+gp10+gp15-gp16+gp17-gp18));
-                //Fcpz = -3.0/18.0*(gp5-gp6+0.5*(gp11-gp12-gp13+gp14+gp15-gp16-gp17+gp18));
-                Fcpx = -3.0/18.0*(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
-                Fcpy = -3.0/18.0*(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
-                Fcpz = -3.0/18.0*(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
-                Fcpx *= alpha*W/sqrt(perm); 
-                Fcpy *= alpha*W/sqrt(perm); 
-                Fcpz *= alpha*W/sqrt(perm); 
-                double Fcp_mag_temp = sqrt(Fcpx*Fcpx+Fcpy*Fcpy+Fcpz*Fcpz);
-                double Fcp_mag = Fcp_mag_temp;
-                if (Fcp_mag_temp==0.0) Fcp_mag=1.0;
-                nx = Fcpx/Fcp_mag;
-                ny = Fcpy/Fcp_mag;
-                nz = Fcpz/Fcp_mag;
-            }
+//			Fcpx = 0.0;
+//			Fcpy = 0.0;
+//			Fcpz = 0.0;
+//            if (porosity!=1.0){
+//                //Fcpx = -3.0/18.0*(gp1-gp2+0.5*(gp7-gp8+gp9-gp10+gp11-gp12+gp13-gp14));
+//                //Fcpy = -3.0/18.0*(gp3-gp4+0.5*(gp7-gp8-gp9+gp10+gp15-gp16+gp17-gp18));
+//                //Fcpz = -3.0/18.0*(gp5-gp6+0.5*(gp11-gp12-gp13+gp14+gp15-gp16-gp17+gp18));
+//                Fcpx = -3.0/18.0*(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
+//                Fcpy = -3.0/18.0*(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
+//                Fcpz = -3.0/18.0*(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
+//                Fcpx *= alpha*W/sqrt(perm); 
+//                Fcpy *= alpha*W/sqrt(perm); 
+//                Fcpz *= alpha*W/sqrt(perm); 
+//                //double Fcp_mag_temp = sqrt(Fcpx*Fcpx+Fcpy*Fcpy+Fcpz*Fcpz);
+//                //double Fcp_mag = Fcp_mag_temp;
+//                //if (Fcp_mag_temp==0.0) Fcp_mag=1.0;
+//                //nx = Fcpx/Fcp_mag;
+//                //ny = Fcpy/Fcp_mag;
+//                //nz = Fcpz/Fcp_mag;
+//            }
+            Fcpx = -3.0/18.0*(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
+            Fcpy = -3.0/18.0*(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
+            Fcpz = -3.0/18.0*(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
+            //NOTE for open node (porosity=1.0),Fcp=0.0
+            Fcpx *= alpha*W*(1.0-porosity)/sqrt(perm); 
+            Fcpy *= alpha*W*(1.0-porosity)/sqrt(perm); 
+            Fcpz *= alpha*W*(1.0-porosity)/sqrt(perm); 
 
 			// q=0
 			fq = dist[n];
@@ -2404,9 +2411,9 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
 			m18 = Phi[nn];					// get neighbor for phi - 18
             gp18 = Psi[nn];
 			//............Compute the Color Gradient...................................
-			nx_phase = -(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
-			ny_phase = -(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
-			nz_phase = -(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
+			nx_phase = -3.0/18.0*(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
+			ny_phase = -3.0/18.0*(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
+			nz_phase = -3.0/18.0*(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
 			C_phase = sqrt(nx_phase*nx_phase+ny_phase*ny_phase+nz_phase*nz_phase);
 
             //correct the normal color gradient by considering the effect of grey solid
@@ -2428,26 +2435,33 @@ __global__  void dvc_ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dis
 			nz = nz/ColorMag;		
 
 			//............Compute the Greyscale Potential Gradient.....................
-			Fcpx = 0.0;
-			Fcpy = 0.0;
-			Fcpz = 0.0;
-            if (porosity!=1.0){
-                //Fcpx = -3.0/18.0*(gp1-gp2+0.5*(gp7-gp8+gp9-gp10+gp11-gp12+gp13-gp14));
-                //Fcpy = -3.0/18.0*(gp3-gp4+0.5*(gp7-gp8-gp9+gp10+gp15-gp16+gp17-gp18));
-                //Fcpz = -3.0/18.0*(gp5-gp6+0.5*(gp11-gp12-gp13+gp14+gp15-gp16-gp17+gp18));
-                Fcpx = -3.0/18.0*(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
-                Fcpy = -3.0/18.0*(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
-                Fcpz = -3.0/18.0*(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
-                Fcpx *= alpha*W/sqrt(perm); 
-                Fcpy *= alpha*W/sqrt(perm); 
-                Fcpz *= alpha*W/sqrt(perm); 
-                double Fcp_mag_temp = sqrt(Fcpx*Fcpx+Fcpy*Fcpy+Fcpz*Fcpz);
-                double Fcp_mag = Fcp_mag_temp;
-                if (Fcp_mag_temp==0.0) Fcp_mag=1.0;
-                nx = Fcpx/Fcp_mag;
-                ny = Fcpy/Fcp_mag;
-                nz = Fcpz/Fcp_mag;
-            }
+//			Fcpx = 0.0;
+//			Fcpy = 0.0;
+//			Fcpz = 0.0;
+//            if (porosity!=1.0){
+//                //Fcpx = -3.0/18.0*(gp1-gp2+0.5*(gp7-gp8+gp9-gp10+gp11-gp12+gp13-gp14));
+//                //Fcpy = -3.0/18.0*(gp3-gp4+0.5*(gp7-gp8-gp9+gp10+gp15-gp16+gp17-gp18));
+//                //Fcpz = -3.0/18.0*(gp5-gp6+0.5*(gp11-gp12-gp13+gp14+gp15-gp16-gp17+gp18));
+//                Fcpx = -3.0/18.0*(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
+//                Fcpy = -3.0/18.0*(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
+//                Fcpz = -3.0/18.0*(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
+//                Fcpx *= alpha*W/sqrt(perm); 
+//                Fcpy *= alpha*W/sqrt(perm); 
+//                Fcpz *= alpha*W/sqrt(perm); 
+//                double Fcp_mag_temp = sqrt(Fcpx*Fcpx+Fcpy*Fcpy+Fcpz*Fcpz);
+//                double Fcp_mag = Fcp_mag_temp;
+//                if (Fcp_mag_temp==0.0) Fcp_mag=1.0;
+//                nx = Fcpx/Fcp_mag;
+//                ny = Fcpy/Fcp_mag;
+//                nz = Fcpz/Fcp_mag;
+//            }
+            Fcpx = -3.0/18.0*(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
+            Fcpy = -3.0/18.0*(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
+            Fcpz = -3.0/18.0*(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
+            //NOTE for open node (porosity=1.0),Fcp=0.0
+            Fcpx *= alpha*W*(1.0-porosity)/sqrt(perm); 
+            Fcpy *= alpha*W*(1.0-porosity)/sqrt(perm); 
+            Fcpz *= alpha*W*(1.0-porosity)/sqrt(perm); 
 
 
 			// q=0

From 02d2e514ed2bfb6828e7736ecd8364a8b9ac74df Mon Sep 17 00:00:00 2001
From: Rex Zhe Li <zhe.rex.li@gmail.com>
Date: Thu, 18 Mar 2021 07:58:05 -0400
Subject: [PATCH 10/10] make CPU also ready for capillary penalty; to be built
 and verified

---
 common/ScaLBL.h        |    4 +-
 cpu/GreyscaleColor.cpp | 1456 ++++++++++++++++++++++++++++++++++++++++
 cuda/GreyscaleColor.cu |   71 +-
 3 files changed, 1495 insertions(+), 36 deletions(-)

diff --git a/common/ScaLBL.h b/common/ScaLBL.h
index 5c4d4730..22ef077c 100644
--- a/common/ScaLBL.h
+++ b/common/ScaLBL.h
@@ -97,8 +97,8 @@ extern "C" void ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *d_neighborList, int *M
         double rhoA, double rhoB, double tauA, double tauB, double tauA_eff,double tauB_eff, double alpha, double beta,
 		double Fx, double Fy, double Fz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np);
 
-extern "C" void ScaLBL_Update_GreyscalePotential(int *Map, double *Phi, double *Psi, double *Poro, double *Perm, double alpha, double W, 
-		int start, int finish, int Np);
+//extern "C" void ScaLBL_Update_GreyscalePotential(int *Map, double *Phi, double *Psi, double *Poro, double *Perm, double alpha, double W, 
+//		int start, int finish, int Np);
 
 // ION TRANSPORT MODEL
 
diff --git a/cpu/GreyscaleColor.cpp b/cpu/GreyscaleColor.cpp
index 057c8a7d..dfd755e3 100644
--- a/cpu/GreyscaleColor.cpp
+++ b/cpu/GreyscaleColor.cpp
@@ -1338,6 +1338,1462 @@ extern "C" void ScaLBL_D3Q19_AAeven_GreyscaleColor(int *Map, double *dist, doubl
 	}
 }
 
+//CP: capillary penalty
+// also turn off recoloring for grey nodes
+extern "C" void ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *neighborList, int *Map, double *dist, double *Aq, double *Bq, double *Den, 
+		double *Phi, double *Psi, double *GreySolidGrad, double *Poros,double *Perm,double *Velocity,double *Pressure,
+        double rhoA, double rhoB, double tauA, double tauB, double tauA_eff,double tauB_eff, double alpha, double beta,
+		double Gx, double Gy, double Gz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np){
+
+	int n,nn,ijk,nread;
+	int nr1,nr2,nr3,nr4,nr5,nr6;
+	int nr7,nr8,nr9,nr10;
+	int nr11,nr12,nr13,nr14;
+	//int nr15,nr16,nr17,nr18;
+	double fq;
+	// conserved momemnts
+	double rho,jx,jy,jz;
+	//double vx,vy,vz,v_mag;
+    //double ux,uy,uz,u_mag;
+    double ux,uy,uz;
+	// non-conserved moments
+	double m1,m2,m4,m6,m8,m9,m10,m11,m12,m13,m14,m15,m16,m17,m18;
+	double m3,m5,m7;
+	double nA,nB; // number density
+	double a1,b1,a2,b2,nAB,delta;
+	double C,nx,ny,nz; //color gradient magnitude and direction
+	double phi,tau,rho0,rlx_setA,rlx_setB;
+
+    //double GeoFun=0.0;//geometric function from Guo's PRE 66, 036304 (2002)
+    double porosity;
+    double perm;//voxel permeability
+    //double c0, c1; //Guo's model parameters
+    double tau_eff;
+    double mu_eff;//kinematic viscosity
+    double nx_gs,ny_gs,nz_gs;//grey-solid color gradient
+    double nx_phase,ny_phase,nz_phase,C_phase;
+    double Fx,Fy,Fz;
+	double gp1,gp2,gp4,gp6,gp8,gp9,gp10,gp11,gp12,gp13,gp14,gp15,gp16,gp17,gp18;
+	double gp3,gp5,gp7;
+    double Fcpx,Fcpy,Fcpz;//capillary penalty force
+
+	const double mrt_V1=0.05263157894736842;
+	const double mrt_V2=0.012531328320802;
+	const double mrt_V3=0.04761904761904762;
+	const double mrt_V4=0.004594820384294068;
+	const double mrt_V5=0.01587301587301587;
+	const double mrt_V6=0.0555555555555555555555555;
+	const double mrt_V7=0.02777777777777778;
+	const double mrt_V8=0.08333333333333333;
+	const double mrt_V9=0.003341687552213868;
+	const double mrt_V10=0.003968253968253968;
+	const double mrt_V11=0.01388888888888889;
+	const double mrt_V12=0.04166666666666666;
+
+	for (n=start; n<finish; n++){
+		// read the component number densities
+		nA = Den[n];
+		nB = Den[Np + n];
+        porosity = Poros[n];
+        perm = Perm[n];
+        nx_gs = GreySolidGrad[n+0*Np];
+        ny_gs = GreySolidGrad[n+1*Np];
+        nz_gs = GreySolidGrad[n+2*Np];
+
+		// compute phase indicator field
+		phi=(nA-nB)/(nA+nB);
+
+		// local density
+		rho0=rhoA + 0.5*(1.0-phi)*(rhoB-rhoA);
+		// local relaxation time
+		tau=tauA + 0.5*(1.0-phi)*(tauB-tauA);
+		tau_eff=tauA_eff + 0.5*(1.0-phi)*(tauB_eff-tauA_eff);
+		rlx_setA = 1.f/tau;
+		rlx_setB = 8.f*(2.f-rlx_setA)/(8.f-rlx_setA);
+        mu_eff = (tau_eff-0.5)/3.0;//kinematic viscosity
+		
+		// Get the 1D index based on regular data layout
+		ijk = Map[n];
+		//					COMPUTE THE COLOR GRADIENT
+		//........................................................................
+		//.................Read Phase Indicator Values............................
+		//........................................................................
+		nn = ijk-1;							// neighbor index (get convention)
+		m1 = Phi[nn];						// get neighbor for phi - 1
+        gp1 = Psi[nn];
+		//........................................................................
+		nn = ijk+1;							// neighbor index (get convention)
+		m2 = Phi[nn];						// get neighbor for phi - 2
+        gp2 = Psi[nn];
+		//........................................................................
+		nn = ijk-strideY;							// neighbor index (get convention)
+		m3 = Phi[nn];					// get neighbor for phi - 3
+        gp3 = Psi[nn];
+		//........................................................................
+		nn = ijk+strideY;							// neighbor index (get convention)
+		m4 = Phi[nn];					// get neighbor for phi - 4
+        gp4 = Psi[nn];
+		//........................................................................
+		nn = ijk-strideZ;						// neighbor index (get convention)
+		m5 = Phi[nn];					// get neighbor for phi - 5
+        gp5 = Psi[nn];
+		//........................................................................
+		nn = ijk+strideZ;						// neighbor index (get convention)
+		m6 = Phi[nn];					// get neighbor for phi - 6
+        gp6 = Psi[nn];
+		//........................................................................
+		nn = ijk-strideY-1;						// neighbor index (get convention)
+		m7 = Phi[nn];					// get neighbor for phi - 7
+        gp7 = Psi[nn];
+		//........................................................................
+		nn = ijk+strideY+1;						// neighbor index (get convention)
+		m8 = Phi[nn];					// get neighbor for phi - 8
+        gp8 = Psi[nn];
+		//........................................................................
+		nn = ijk+strideY-1;						// neighbor index (get convention)
+		m9 = Phi[nn];					// get neighbor for phi - 9
+        gp9 = Psi[nn];
+		//........................................................................
+		nn = ijk-strideY+1;						// neighbor index (get convention)
+		m10 = Phi[nn];					// get neighbor for phi - 10
+        gp10 = Psi[nn];
+		//........................................................................
+		nn = ijk-strideZ-1;						// neighbor index (get convention)
+		m11 = Phi[nn];					// get neighbor for phi - 11
+        gp11 = Psi[nn];
+		//........................................................................
+		nn = ijk+strideZ+1;						// neighbor index (get convention)
+		m12 = Phi[nn];					// get neighbor for phi - 12
+        gp12 = Psi[nn];
+		//........................................................................
+		nn = ijk+strideZ-1;						// neighbor index (get convention)
+		m13 = Phi[nn];					// get neighbor for phi - 13
+        gp13 = Psi[nn];
+		//........................................................................
+		nn = ijk-strideZ+1;						// neighbor index (get convention)
+		m14 = Phi[nn];					// get neighbor for phi - 14
+        gp14 = Psi[nn];
+		//........................................................................
+		nn = ijk-strideZ-strideY;					// neighbor index (get convention)
+		m15 = Phi[nn];					// get neighbor for phi - 15
+        gp15 = Psi[nn];
+		//........................................................................
+		nn = ijk+strideZ+strideY;					// neighbor index (get convention)
+		m16 = Phi[nn];					// get neighbor for phi - 16
+        gp16 = Psi[nn];
+		//........................................................................
+		nn = ijk+strideZ-strideY;					// neighbor index (get convention)
+		m17 = Phi[nn];					// get neighbor for phi - 17
+        gp17 = Psi[nn];
+		//........................................................................
+		nn = ijk-strideZ+strideY;					// neighbor index (get convention)
+		m18 = Phi[nn];					// get neighbor for phi - 18
+        gp18 = Psi[nn];
+		//............Compute the Color Gradient...................................
+		nx_phase = -3.0/18.0*(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
+		ny_phase = -3.0/18.0*(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
+		nz_phase = -3.0/18.0*(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
+		C_phase = sqrt(nx_phase*nx_phase+ny_phase*ny_phase+nz_phase*nz_phase);
+
+        //correct the normal color gradient by considering the effect of grey solid
+        nx = nx_phase + (1.0-porosity)*nx_gs; 
+        ny = ny_phase + (1.0-porosity)*ny_gs; 
+        nz = nz_phase + (1.0-porosity)*nz_gs; 
+        if (C_phase==0.0){
+            nx = nx_phase; 
+            ny = ny_phase;
+            nz = nz_phase;
+        }
+
+		//...........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 the Greyscale Potential Gradient.....................
+//			Fcpx = 0.0;
+//			Fcpy = 0.0;
+//			Fcpz = 0.0;
+//            if (porosity!=1.0){
+//                //Fcpx = -3.0/18.0*(gp1-gp2+0.5*(gp7-gp8+gp9-gp10+gp11-gp12+gp13-gp14));
+//                //Fcpy = -3.0/18.0*(gp3-gp4+0.5*(gp7-gp8-gp9+gp10+gp15-gp16+gp17-gp18));
+//                //Fcpz = -3.0/18.0*(gp5-gp6+0.5*(gp11-gp12-gp13+gp14+gp15-gp16-gp17+gp18));
+//                Fcpx = -3.0/18.0*(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
+//                Fcpy = -3.0/18.0*(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
+//                Fcpz = -3.0/18.0*(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
+//                Fcpx *= alpha*W/sqrt(perm); 
+//                Fcpy *= alpha*W/sqrt(perm); 
+//                Fcpz *= alpha*W/sqrt(perm); 
+//                //double Fcp_mag_temp = sqrt(Fcpx*Fcpx+Fcpy*Fcpy+Fcpz*Fcpz);
+//                //double Fcp_mag = Fcp_mag_temp;
+//                //if (Fcp_mag_temp==0.0) Fcp_mag=1.0;
+//                //nx = Fcpx/Fcp_mag;
+//                //ny = Fcpy/Fcp_mag;
+//                //nz = Fcpz/Fcp_mag;
+//            }
+        Fcpx = -3.0/18.0*(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
+        Fcpy = -3.0/18.0*(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
+        Fcpz = -3.0/18.0*(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
+        //NOTE for open node (porosity=1.0),Fcp=0.0
+        Fcpx *= alpha*W*(1.0-porosity)/sqrt(perm); 
+        Fcpy *= alpha*W*(1.0-porosity)/sqrt(perm); 
+        Fcpz *= alpha*W*(1.0-porosity)/sqrt(perm); 
+
+		// q=0
+		fq = dist[n];
+		rho = fq;
+		m1  = -30.0*fq;
+		m2  = 12.0*fq;
+
+		// q=1
+		//nread = neighborList[n]; // neighbor 2 
+		//fq = dist[nread]; // reading the f1 data into register fq		
+		nr1 = neighborList[n]; 
+		fq = dist[nr1]; // reading the f1 data into register fq
+		rho += fq;
+		m1 -= 11.0*fq;
+		m2 -= 4.0*fq;
+		jx = fq;
+		m4 = -4.0*fq;
+		m9 = 2.0*fq;
+		m10 = -4.0*fq;
+
+		// f2 = dist[10*Np+n];
+		//nread = neighborList[n+Np]; // neighbor 1 ( < 10Np => even part of dist)
+		//fq = dist[nread];  // reading the f2 data into register fq
+		nr2 = neighborList[n+Np]; // neighbor 1 ( < 10Np => even part of dist)
+		fq = dist[nr2];  // reading the f2 data into register fq
+		rho += fq;
+		m1 -= 11.0*(fq);
+		m2 -= 4.0*(fq);
+		jx -= fq;
+		m4 += 4.0*(fq);
+		m9 += 2.0*(fq);
+		m10 -= 4.0*(fq);
+
+		// q=3
+		//nread = neighborList[n+2*Np]; // neighbor 4
+		//fq = dist[nread];
+		nr3 = neighborList[n+2*Np]; // neighbor 4
+		fq = dist[nr3];
+		rho += fq;
+		m1 -= 11.0*fq;
+		m2 -= 4.0*fq;
+		jy = fq;
+		m6 = -4.0*fq;
+		m9 -= fq;
+		m10 += 2.0*fq;
+		m11 = fq;
+		m12 = -2.0*fq;
+
+		// q = 4
+		//nread = neighborList[n+3*Np]; // neighbor 3
+		//fq = dist[nread];
+		nr4 = neighborList[n+3*Np]; // neighbor 3
+		fq = dist[nr4];
+		rho+= fq;
+		m1 -= 11.0*fq;
+		m2 -= 4.0*fq;
+		jy -= fq;
+		m6 += 4.0*fq;
+		m9 -= fq;
+		m10 += 2.0*fq;
+		m11 += fq;
+		m12 -= 2.0*fq;
+
+		// q=5
+		//nread = neighborList[n+4*Np];
+		//fq = dist[nread];
+		nr5 = neighborList[n+4*Np];
+		fq = dist[nr5];
+		rho += fq;
+		m1 -= 11.0*fq;
+		m2 -= 4.0*fq;
+		jz = fq;
+		m8 = -4.0*fq;
+		m9 -= fq;
+		m10 += 2.0*fq;
+		m11 -= fq;
+		m12 += 2.0*fq;
+
+
+		// q = 6
+		//nread = neighborList[n+5*Np];
+		//fq = dist[nread];
+		nr6 = neighborList[n+5*Np];
+		fq = dist[nr6];
+		rho+= fq;
+		m1 -= 11.0*fq;
+		m2 -= 4.0*fq;
+		jz -= fq;
+		m8 += 4.0*fq;
+		m9 -= fq;
+		m10 += 2.0*fq;
+		m11 -= fq;
+		m12 += 2.0*fq;
+
+		// q=7
+		//nread = neighborList[n+6*Np];
+		//fq = dist[nread];
+		nr7 = neighborList[n+6*Np];
+		fq = dist[nr7];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jx += fq;
+		m4 += fq;
+		jy += fq;
+		m6 += fq;
+		m9  += fq;
+		m10 += fq;
+		m11 += fq;
+		m12 += fq;
+		m13 = fq;
+		m16 = fq;
+		m17 = -fq;
+
+		// q = 8
+		//nread = neighborList[n+7*Np];
+		//fq = dist[nread];
+		nr8 = neighborList[n+7*Np];
+		fq = dist[nr8];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jx -= fq;
+		m4 -= fq;
+		jy -= fq;
+		m6 -= fq;
+		m9 += fq;
+		m10 += fq;
+		m11 += fq;
+		m12 += fq;
+		m13 += fq;
+		m16 -= fq;
+		m17 += fq;
+
+		// q=9
+		//nread = neighborList[n+8*Np];
+		//fq = dist[nread];
+		nr9 = neighborList[n+8*Np];
+		fq = dist[nr9];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jx += fq;
+		m4 += fq;
+		jy -= fq;
+		m6 -= fq;
+		m9 += fq;
+		m10 += fq;
+		m11 += fq;
+		m12 += fq;
+		m13 -= fq;
+		m16 += fq;
+		m17 += fq;
+
+		// q = 10
+		//nread = neighborList[n+9*Np];
+		//fq = dist[nread];
+		nr10 = neighborList[n+9*Np];
+		fq = dist[nr10];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jx -= fq;
+		m4 -= fq;
+		jy += fq;
+		m6 += fq;
+		m9 += fq;
+		m10 += fq;
+		m11 += fq;
+		m12 += fq;
+		m13 -= fq;
+		m16 -= fq;
+		m17 -= fq;
+
+		// q=11
+		//nread = neighborList[n+10*Np];
+		//fq = dist[nread];
+		nr11 = neighborList[n+10*Np];
+		fq = dist[nr11];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jx += fq;
+		m4 += fq;
+		jz += fq;
+		m8 += fq;
+		m9 += fq;
+		m10 += fq;
+		m11 -= fq;
+		m12 -= fq;
+		m15 = fq;
+		m16 -= fq;
+		m18 = fq;
+
+		// q=12
+		//nread = neighborList[n+11*Np];
+		//fq = dist[nread];
+		nr12 = neighborList[n+11*Np];
+		fq = dist[nr12];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jx -= fq;
+		m4 -= fq;
+		jz -= fq;
+		m8 -= fq;
+		m9 += fq;
+		m10 += fq;
+		m11 -= fq;
+		m12 -= fq;
+		m15 += fq;
+		m16 += fq;
+		m18 -= fq;
+
+		// q=13
+		//nread = neighborList[n+12*Np];
+		//fq = dist[nread];
+		nr13 = neighborList[n+12*Np];
+		fq = dist[nr13];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jx += fq;
+		m4 += fq;
+		jz -= fq;
+		m8 -= fq;
+		m9 += fq;
+		m10 += fq;
+		m11 -= fq;
+		m12 -= fq;
+		m15 -= fq;
+		m16 -= fq;
+		m18 -= fq;
+
+		// q=14
+		//nread = neighborList[n+13*Np];
+		//fq = dist[nread];
+		nr14 = neighborList[n+13*Np];
+		fq = dist[nr14];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jx -= fq;
+		m4 -= fq;
+		jz += fq;
+		m8 += fq;
+		m9 += fq;
+		m10 += fq;
+		m11 -= fq;
+		m12 -= fq;
+		m15 -= fq;
+		m16 += fq;
+		m18 += fq;
+
+		// q=15
+		nread = neighborList[n+14*Np];
+		fq = dist[nread];
+		//fq = dist[17*Np+n];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jy += fq;
+		m6 += fq;
+		jz += fq;
+		m8 += fq;
+		m9 -= 2.0*fq;
+		m10 -= 2.0*fq;
+		m14 = fq;
+		m17 += fq;
+		m18 -= fq;
+
+		// q=16
+		nread = neighborList[n+15*Np];
+		fq = dist[nread];
+		//fq = dist[8*Np+n];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jy -= fq;
+		m6 -= fq;
+		jz -= fq;
+		m8 -= fq;
+		m9 -= 2.0*fq;
+		m10 -= 2.0*fq;
+		m14 += fq;
+		m17 -= fq;
+		m18 += fq;
+
+		// q=17
+		//fq = dist[18*Np+n];
+		nread = neighborList[n+16*Np];
+		fq = dist[nread];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jy += fq;
+		m6 += fq;
+		jz -= fq;
+		m8 -= fq;
+		m9 -= 2.0*fq;
+		m10 -= 2.0*fq;
+		m14 -= fq;
+		m17 += fq;
+		m18 += fq;
+
+		// q=18
+		nread = neighborList[n+17*Np];
+		fq = dist[nread];
+		//fq = dist[9*Np+n];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jy -= fq;
+		m6 -= fq;
+		jz += fq;
+		m8 += fq;
+		m9 -= 2.0*fq;
+		m10 -= 2.0*fq;
+		m14 -= fq;
+		m17 -= fq;
+		m18 -= fq;
+		
+        // Compute greyscale related parameters
+        ux = (jx/rho0+0.5*porosity*Gx+0.5*Fcpx/rho0)/(1.0+0.5*porosity*mu_eff/perm);
+        uy = (jy/rho0+0.5*porosity*Gy+0.5*Fcpy/rho0)/(1.0+0.5*porosity*mu_eff/perm);
+        uz = (jz/rho0+0.5*porosity*Gz+0.5*Fcpz/rho0)/(1.0+0.5*porosity*mu_eff/perm);
+        if (porosity==1.0){//i.e. open nodes
+            ux = (jx/rho0+0.5*porosity*Gx);
+            uy = (jy/rho0+0.5*porosity*Gy);
+            uz = (jz/rho0+0.5*porosity*Gz);
+        }
+
+        //Update the total force to include linear (Darcy) and nonlinear (Forchheimer) drags due to the porous medium
+        Fx = rho0*(-porosity*mu_eff/perm*ux + porosity*Gx)+Fcpx;
+        Fy = rho0*(-porosity*mu_eff/perm*uy + porosity*Gy)+Fcpy;
+        Fz = rho0*(-porosity*mu_eff/perm*uz + porosity*Gz)+Fcpz;
+        if (porosity==1.0){
+            Fx=rho0*(porosity*Gx);
+            Fy=rho0*(porosity*Gy);
+            Fz=rho0*(porosity*Gz);
+        }
+
+		// write the velocity 
+		Velocity[n] = ux;
+		Velocity[Np+n] = uy;
+		Velocity[2*Np+n] = uz;
+        //Pressure[n] = rho/3.f/porosity;
+        Pressure[n] = rho/3.f;
+
+		//........................................................................
+		//..............carry out relaxation process..............................
+		//..........Toelke, Fruediger et. al. 2006................................
+		if (C == 0.0)	nx = ny = nz = 0.0;
+		m1 = m1 + rlx_setA*((19*(ux*ux+uy*uy+uz*uz)*rho0/porosity - 11*rho) -19*alpha*C - m1);
+		m2 = m2 + rlx_setA*((3*rho - 5.5*(ux*ux+uy*uy+uz*uz)*rho0/porosity)- m2);
+        jx = jx + Fx;
+		m4 = m4 + rlx_setB*((-0.6666666666666666*ux*rho0)- m4)
+                + (1-0.5*rlx_setB)*(-0.6666666666666666*Fx);
+        jy = jy + Fy;
+		m6 = m6 + rlx_setB*((-0.6666666666666666*uy*rho0)- m6)
+                + (1-0.5*rlx_setB)*(-0.6666666666666666*Fy);
+        jz = jz + Fz;
+		m8 = m8 + rlx_setB*((-0.6666666666666666*uz*rho0)- m8)
+                + (1-0.5*rlx_setB)*(-0.6666666666666666*Fz);
+		m9 = m9 + rlx_setA*(((2*ux*ux-uy*uy-uz*uz)*rho0/porosity) + 0.5*alpha*C*(2*nx*nx-ny*ny-nz*nz) - m9);
+		m10 = m10 + rlx_setA*( - m10);
+        //m10 = m10 + rlx_setA*(-0.5*rho0*((2*ux*ux-uy*uy-uz*uz)/porosity)- m10);
+		m11 = m11 + rlx_setA*(((uy*uy-uz*uz)*rho0/porosity) + 0.5*alpha*C*(ny*ny-nz*nz)- m11);
+		m12 = m12 + rlx_setA*( - m12);
+        //m12 = m12 + rlx_setA*(-0.5*(rho0*(uy*uy-uz*uz)/porosity)- m12);
+		m13 = m13 + rlx_setA*( (ux*uy*rho0/porosity) + 0.5*alpha*C*nx*ny - m13);
+		m14 = m14 + rlx_setA*( (uy*uz*rho0/porosity) + 0.5*alpha*C*ny*nz - m14);
+		m15 = m15 + rlx_setA*( (ux*uz*rho0/porosity) + 0.5*alpha*C*nx*nz - m15);
+		m16 = m16 + rlx_setB*( - m16);
+		m17 = m17 + rlx_setB*( - m17);
+		m18 = m18 + rlx_setB*( - m18);
+
+		//.................inverse transformation......................................................
+		// q=0
+		fq = mrt_V1*rho-mrt_V2*m1+mrt_V3*m2;
+		dist[n] = fq;
+
+		// q = 1
+		fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(jx-m4)+mrt_V6*(m9-m10);
+		//nread = neighborList[n+Np];
+		dist[nr2] = fq;
+
+		// q=2
+		fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(m4-jx)+mrt_V6*(m9-m10);
+		//nread = neighborList[n];
+		dist[nr1] = fq;
+
+		// q = 3
+		fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(jy-m6)+mrt_V7*(m10-m9)+mrt_V8*(m11-m12);
+		//nread = neighborList[n+3*Np];
+		dist[nr4] = fq;
+
+		// q = 4
+		fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(m6-jy)+mrt_V7*(m10-m9)+mrt_V8*(m11-m12);
+		//nread = neighborList[n+2*Np];
+		dist[nr3] = fq;
+
+		// q = 5
+		fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(jz-m8)+mrt_V7*(m10-m9)+mrt_V8*(m12-m11);
+		//nread = neighborList[n+5*Np];
+		dist[nr6] = fq;
+
+		// q = 6
+		fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(m8-jz)+mrt_V7*(m10-m9)+mrt_V8*(m12-m11);
+		//nread = neighborList[n+4*Np];
+		dist[nr5] = fq;
+
+		// q = 7
+		fq = mrt_V1*rho+mrt_V9*m1+mrt_V10*m2+0.1*(jx+jy)+0.025*(m4+m6)+
+				mrt_V7*m9+mrt_V11*m10+mrt_V8*m11+mrt_V12*m12+0.25*m13+0.125*(m16-m17);
+		//nread = neighborList[n+7*Np];
+		dist[nr8] = fq;
+
+		// q = 8
+		fq = mrt_V1*rho+mrt_V9*m1+mrt_V10*m2-0.1*(jx+jy)-0.025*(m4+m6) +mrt_V7*m9+mrt_V11*m10+mrt_V8*m11
+				+mrt_V12*m12+0.25*m13+0.125*(m17-m16);
+		//nread = neighborList[n+6*Np];
+		dist[nr7] = fq;
+
+		// q = 9
+		fq = mrt_V1*rho+mrt_V9*m1+mrt_V10*m2+0.1*(jx-jy)+0.025*(m4-m6)+
+				mrt_V7*m9+mrt_V11*m10+mrt_V8*m11+mrt_V12*m12-0.25*m13+0.125*(m16+m17);
+		//nread = neighborList[n+9*Np];
+		dist[nr10] = fq;
+
+		// q = 10
+		fq = mrt_V1*rho+mrt_V9*m1+mrt_V10*m2+0.1*(jy-jx)+0.025*(m6-m4)+
+				mrt_V7*m9+mrt_V11*m10+mrt_V8*m11+mrt_V12*m12-0.25*m13-0.125*(m16+m17);
+		//nread = neighborList[n+8*Np];
+		dist[nr9] = fq;
+
+		// q = 11
+		fq = mrt_V1*rho+mrt_V9*m1
+				+mrt_V10*m2+0.1*(jx+jz)+0.025*(m4+m8)
+				+mrt_V7*m9+mrt_V11*m10-mrt_V8*m11
+				-mrt_V12*m12+0.25*m15+0.125*(m18-m16);
+		//nread = neighborList[n+11*Np];
+		dist[nr12] = fq;
+
+		// q = 12
+		fq = mrt_V1*rho+mrt_V9*m1+mrt_V10*m2-0.1*(jx+jz)-0.025*(m4+m8)+
+				mrt_V7*m9+mrt_V11*m10-mrt_V8*m11-mrt_V12*m12+0.25*m15+0.125*(m16-m18);
+		//nread = neighborList[n+10*Np];
+		dist[nr11]= fq;
+
+		// q = 13
+		fq = mrt_V1*rho+mrt_V9*m1
+				+mrt_V10*m2+0.1*(jx-jz)+0.025*(m4-m8)
+				+mrt_V7*m9+mrt_V11*m10-mrt_V8*m11
+				-mrt_V12*m12-0.25*m15-0.125*(m16+m18);
+		//nread = neighborList[n+13*Np];
+		dist[nr14] = fq;
+
+		// q= 14
+		fq = mrt_V1*rho+mrt_V9*m1
+				+mrt_V10*m2+0.1*(jz-jx)+0.025*(m8-m4)
+				+mrt_V7*m9+mrt_V11*m10-mrt_V8*m11
+				-mrt_V12*m12-0.25*m15+0.125*(m16+m18);
+		//nread = neighborList[n+12*Np];
+		dist[nr13] = fq;
+
+
+		// q = 15
+		fq = mrt_V1*rho+mrt_V9*m1
+				+mrt_V10*m2+0.1*(jy+jz)+0.025*(m6+m8)
+				-mrt_V6*m9-mrt_V7*m10+0.25*m14+0.125*(m17-m18);
+		nread = neighborList[n+15*Np];
+		dist[nread] = fq;
+
+		// q = 16
+		fq =  mrt_V1*rho+mrt_V9*m1
+				+mrt_V10*m2-0.1*(jy+jz)-0.025*(m6+m8)
+				-mrt_V6*m9-mrt_V7*m10+0.25*m14+0.125*(m18-m17);
+		nread = neighborList[n+14*Np];
+		dist[nread] = fq;
+
+
+		// q = 17
+		fq = mrt_V1*rho+mrt_V9*m1
+				+mrt_V10*m2+0.1*(jy-jz)+0.025*(m6-m8)
+				-mrt_V6*m9-mrt_V7*m10-0.25*m14+0.125*(m17+m18);
+		nread = neighborList[n+17*Np];
+		dist[nread] = fq;
+
+		// q = 18
+		fq = mrt_V1*rho+mrt_V9*m1
+				+mrt_V10*m2+0.1*(jz-jy)+0.025*(m8-m6)
+				-mrt_V6*m9-mrt_V7*m10-0.25*m14-0.125*(m17+m18);
+		nread = neighborList[n+16*Np];
+		dist[nread] = fq;
+		//........................................................................
+
+		// Instantiate mass transport distributions
+		// Stationary value - distribution 0
+		nAB = 1.0/(nA+nB);
+		Aq[n] = 0.3333333333333333*nA;
+		Bq[n] = 0.3333333333333333*nB;
+
+		//...............................................
+		// q = 0,2,4
+		// Cq = {1,0,0}, {0,1,0}, {0,0,1}
+		delta = beta*nA*nB*nAB*0.1111111111111111*nx;
+		if (!(nA*nB*nAB>0)) delta=0;
+        if (RecoloringOff==true && porosity !=1.0) delta=0;
+		a1 = nA*(0.1111111111111111*(1+4.5*ux))+delta;
+		b1 = nB*(0.1111111111111111*(1+4.5*ux))-delta;
+		a2 = nA*(0.1111111111111111*(1-4.5*ux))-delta;
+		b2 = nB*(0.1111111111111111*(1-4.5*ux))+delta;
+
+		// q = 1
+		//nread = neighborList[n+Np];
+		Aq[nr2] = a1;
+		Bq[nr2] = b1;
+		// q=2
+		//nread = neighborList[n];
+		Aq[nr1] = a2;
+		Bq[nr1] = b2;
+
+		//...............................................
+		// Cq = {0,1,0}
+		delta = beta*nA*nB*nAB*0.1111111111111111*ny;
+		if (!(nA*nB*nAB>0)) delta=0;
+        if (RecoloringOff==true && porosity !=1.0) delta=0;
+		a1 = nA*(0.1111111111111111*(1+4.5*uy))+delta;
+		b1 = nB*(0.1111111111111111*(1+4.5*uy))-delta;
+		a2 = nA*(0.1111111111111111*(1-4.5*uy))-delta;
+		b2 = nB*(0.1111111111111111*(1-4.5*uy))+delta;
+
+		// q = 3
+		//nread = neighborList[n+3*Np];
+		Aq[nr4] = a1;
+		Bq[nr4] = b1;
+		// q = 4
+		//nread = neighborList[n+2*Np];
+		Aq[nr3] = a2;
+		Bq[nr3] = b2;
+
+		//...............................................
+		// q = 4
+		// Cq = {0,0,1}
+		delta = beta*nA*nB*nAB*0.1111111111111111*nz;
+		if (!(nA*nB*nAB>0)) delta=0;
+        if (RecoloringOff==true && porosity !=1.0) delta=0;
+		a1 = nA*(0.1111111111111111*(1+4.5*uz))+delta;
+		b1 = nB*(0.1111111111111111*(1+4.5*uz))-delta;
+		a2 = nA*(0.1111111111111111*(1-4.5*uz))-delta;
+		b2 = nB*(0.1111111111111111*(1-4.5*uz))+delta;
+
+		// q = 5
+		//nread = neighborList[n+5*Np];
+		Aq[nr6] = a1;
+		Bq[nr6] = b1;
+		// q = 6
+		//nread = neighborList[n+4*Np];
+		Aq[nr5] = a2;
+		Bq[nr5] = b2;
+		//...............................................
+	}
+}
+
+//CP: capillary penalty
+// also turn off recoloring for grey nodes
+extern "C" void ScaLBL_D3Q19_AAeven_GreyscaleColor_CP(int *Map, double *dist, double *Aq, double *Bq, double *Den, 
+        double *Phi,double *Psi, double *GreySolidGrad, double *Poros,double *Perm,double *Velocity,double *Pressure,
+        double rhoA, double rhoB, double tauA, double tauB,double tauA_eff,double tauB_eff, double alpha, double beta,
+		double Gx, double Gy, double Gz, bool RecoloringOff, double W, int strideY, int strideZ, int start, int finish, int Np){
+
+	int ijk,nn,n;
+	double fq;
+	// conserved momemnts
+	double rho,jx,jy,jz;
+	//double vx,vy,vz,v_mag;
+    //double ux,uy,uz,u_mag;
+    double ux,uy,uz;
+	// non-conserved moments
+	double m1,m2,m4,m6,m8,m9,m10,m11,m12,m13,m14,m15,m16,m17,m18;
+	double m3,m5,m7;
+	double nA,nB; // number density
+	double a1,b1,a2,b2,nAB,delta;
+	double C,nx,ny,nz; //color gradient magnitude and direction
+	double phi,tau,rho0,rlx_setA,rlx_setB;
+
+    //double GeoFun=0.0;//geometric function from Guo's PRE 66, 036304 (2002)
+    double porosity;
+    double perm;//voxel permeability
+    //double c0, c1; //Guo's model parameters
+    double tau_eff;
+    double mu_eff;//kinematic viscosity
+    double nx_gs,ny_gs,nz_gs;//grey-solid color gradient
+    double nx_phase,ny_phase,nz_phase,C_phase;
+    double Fx,Fy,Fz;
+	double gp1,gp2,gp4,gp6,gp8,gp9,gp10,gp11,gp12,gp13,gp14,gp15,gp16,gp17,gp18;
+	double gp3,gp5,gp7;
+    double Fcpx,Fcpy,Fcpz;//capillary penalty force
+
+	const double mrt_V1=0.05263157894736842;
+	const double mrt_V2=0.012531328320802;
+	const double mrt_V3=0.04761904761904762;
+	const double mrt_V4=0.004594820384294068;
+	const double mrt_V5=0.01587301587301587;
+	const double mrt_V6=0.0555555555555555555555555;
+	const double mrt_V7=0.02777777777777778;
+	const double mrt_V8=0.08333333333333333;
+	const double mrt_V9=0.003341687552213868;
+	const double mrt_V10=0.003968253968253968;
+	const double mrt_V11=0.01388888888888889;
+	const double mrt_V12=0.04166666666666666;
+
+	for (n=start; n<finish; n++){
+
+		// read the component number densities
+		nA = Den[n];
+		nB = Den[Np + n];
+        porosity = Poros[n];
+        perm = Perm[n];
+        nx_gs = GreySolidGrad[n+0*Np];
+        ny_gs = GreySolidGrad[n+1*Np];
+        nz_gs = GreySolidGrad[n+2*Np];
+
+		// compute phase indicator field
+		phi=(nA-nB)/(nA+nB);
+
+		// local density
+		rho0=rhoA + 0.5*(1.0-phi)*(rhoB-rhoA);
+		// local relaxation time
+		tau=tauA + 0.5*(1.0-phi)*(tauB-tauA);
+		tau_eff=tauA_eff + 0.5*(1.0-phi)*(tauB_eff-tauA_eff);
+		rlx_setA = 1.f/tau;
+		rlx_setB = 8.f*(2.f-rlx_setA)/(8.f-rlx_setA);
+        mu_eff = (tau_eff-0.5)/3.0;//kinematic viscosity
+
+		// Get the 1D index based on regular data layout
+		ijk = Map[n];
+		//					COMPUTE THE COLOR GRADIENT
+		//........................................................................
+		//.................Read Phase Indicator Values............................
+		//........................................................................
+		nn = ijk-1;							// neighbor index (get convention)
+		m1 = Phi[nn];						// get neighbor for phi - 1
+        gp1 = Psi[nn];
+		//........................................................................
+		nn = ijk+1;							// neighbor index (get convention)
+		m2 = Phi[nn];						// get neighbor for phi - 2
+        gp2 = Psi[nn];
+		//........................................................................
+		nn = ijk-strideY;							// neighbor index (get convention)
+		m3 = Phi[nn];					// get neighbor for phi - 3
+        gp3 = Psi[nn];
+		//........................................................................
+		nn = ijk+strideY;							// neighbor index (get convention)
+		m4 = Phi[nn];					// get neighbor for phi - 4
+        gp4 = Psi[nn];
+		//........................................................................
+		nn = ijk-strideZ;						// neighbor index (get convention)
+		m5 = Phi[nn];					// get neighbor for phi - 5
+        gp5 = Psi[nn];
+		//........................................................................
+		nn = ijk+strideZ;						// neighbor index (get convention)
+		m6 = Phi[nn];					// get neighbor for phi - 6
+        gp6 = Psi[nn];
+		//........................................................................
+		nn = ijk-strideY-1;						// neighbor index (get convention)
+		m7 = Phi[nn];					// get neighbor for phi - 7
+        gp7 = Psi[nn];
+		//........................................................................
+		nn = ijk+strideY+1;						// neighbor index (get convention)
+		m8 = Phi[nn];					// get neighbor for phi - 8
+        gp8 = Psi[nn];
+		//........................................................................
+		nn = ijk+strideY-1;						// neighbor index (get convention)
+		m9 = Phi[nn];					// get neighbor for phi - 9
+        gp9 = Psi[nn];
+		//........................................................................
+		nn = ijk-strideY+1;						// neighbor index (get convention)
+		m10 = Phi[nn];					// get neighbor for phi - 10
+        gp10 = Psi[nn];
+		//........................................................................
+		nn = ijk-strideZ-1;						// neighbor index (get convention)
+		m11 = Phi[nn];					// get neighbor for phi - 11
+        gp11 = Psi[nn];
+		//........................................................................
+		nn = ijk+strideZ+1;						// neighbor index (get convention)
+		m12 = Phi[nn];					// get neighbor for phi - 12
+        gp12 = Psi[nn];
+		//........................................................................
+		nn = ijk+strideZ-1;						// neighbor index (get convention)
+		m13 = Phi[nn];					// get neighbor for phi - 13
+        gp13 = Psi[nn];
+		//........................................................................
+		nn = ijk-strideZ+1;						// neighbor index (get convention)
+		m14 = Phi[nn];					// get neighbor for phi - 14
+        gp14 = Psi[nn];
+		//........................................................................
+		nn = ijk-strideZ-strideY;					// neighbor index (get convention)
+		m15 = Phi[nn];					// get neighbor for phi - 15
+        gp15 = Psi[nn];
+		//........................................................................
+		nn = ijk+strideZ+strideY;					// neighbor index (get convention)
+		m16 = Phi[nn];					// get neighbor for phi - 16
+        gp16 = Psi[nn];
+		//........................................................................
+		nn = ijk+strideZ-strideY;					// neighbor index (get convention)
+		m17 = Phi[nn];					// get neighbor for phi - 17
+        gp17 = Psi[nn];
+		//........................................................................
+		nn = ijk-strideZ+strideY;					// neighbor index (get convention)
+		m18 = Phi[nn];					// get neighbor for phi - 18
+        gp18 = Psi[nn];
+		//............Compute the Color Gradient...................................
+		nx_phase = -3.0/18.0*(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
+		ny_phase = -3.0/18.0*(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
+		nz_phase = -3.0/18.0*(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
+		C_phase = sqrt(nx_phase*nx_phase+ny_phase*ny_phase+nz_phase*nz_phase);
+
+        //correct the normal color gradient by considering the effect of grey solid
+        nx = nx_phase + (1.0-porosity)*nx_gs; 
+        ny = ny_phase + (1.0-porosity)*ny_gs; 
+        nz = nz_phase + (1.0-porosity)*nz_gs; 
+        if (C_phase==0.0){
+            nx = nx_phase; 
+            ny = ny_phase;
+            nz = nz_phase;
+        }
+
+		//...........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 the Greyscale Potential Gradient.....................
+//			Fcpx = 0.0;
+//			Fcpy = 0.0;
+//			Fcpz = 0.0;
+//            if (porosity!=1.0){
+//                //Fcpx = -3.0/18.0*(gp1-gp2+0.5*(gp7-gp8+gp9-gp10+gp11-gp12+gp13-gp14));
+//                //Fcpy = -3.0/18.0*(gp3-gp4+0.5*(gp7-gp8-gp9+gp10+gp15-gp16+gp17-gp18));
+//                //Fcpz = -3.0/18.0*(gp5-gp6+0.5*(gp11-gp12-gp13+gp14+gp15-gp16-gp17+gp18));
+//                Fcpx = -3.0/18.0*(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
+//                Fcpy = -3.0/18.0*(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
+//                Fcpz = -3.0/18.0*(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
+//                Fcpx *= alpha*W/sqrt(perm); 
+//                Fcpy *= alpha*W/sqrt(perm); 
+//                Fcpz *= alpha*W/sqrt(perm); 
+//                double Fcp_mag_temp = sqrt(Fcpx*Fcpx+Fcpy*Fcpy+Fcpz*Fcpz);
+//                double Fcp_mag = Fcp_mag_temp;
+//                if (Fcp_mag_temp==0.0) Fcp_mag=1.0;
+//                nx = Fcpx/Fcp_mag;
+//                ny = Fcpy/Fcp_mag;
+//                nz = Fcpz/Fcp_mag;
+//            }
+        Fcpx = -3.0/18.0*(m1-m2+0.5*(m7-m8+m9-m10+m11-m12+m13-m14));
+        Fcpy = -3.0/18.0*(m3-m4+0.5*(m7-m8-m9+m10+m15-m16+m17-m18));
+        Fcpz = -3.0/18.0*(m5-m6+0.5*(m11-m12-m13+m14+m15-m16-m17+m18));
+        //NOTE for open node (porosity=1.0),Fcp=0.0
+        Fcpx *= alpha*W*(1.0-porosity)/sqrt(perm); 
+        Fcpy *= alpha*W*(1.0-porosity)/sqrt(perm); 
+        Fcpz *= alpha*W*(1.0-porosity)/sqrt(perm); 
+
+
+		// q=0
+		fq = dist[n];
+		rho = fq;
+		m1  = -30.0*fq;
+		m2  = 12.0*fq;
+
+		// q=1
+		fq = dist[2*Np+n];
+		rho += fq;
+		m1 -= 11.0*fq;
+		m2 -= 4.0*fq;
+		jx = fq;
+		m4 = -4.0*fq;
+		m9 = 2.0*fq;
+		m10 = -4.0*fq;
+
+		// f2 = dist[10*Np+n];
+		fq = dist[1*Np+n];
+		rho += fq;
+		m1 -= 11.0*(fq);
+		m2 -= 4.0*(fq);
+		jx -= fq;
+		m4 += 4.0*(fq);
+		m9 += 2.0*(fq);
+		m10 -= 4.0*(fq);
+
+		// q=3
+		fq = dist[4*Np+n];
+		rho += fq;
+		m1 -= 11.0*fq;
+		m2 -= 4.0*fq;
+		jy = fq;
+		m6 = -4.0*fq;
+		m9 -= fq;
+		m10 += 2.0*fq;
+		m11 = fq;
+		m12 = -2.0*fq;
+
+		// q = 4
+		fq = dist[3*Np+n];
+		rho+= fq;
+		m1 -= 11.0*fq;
+		m2 -= 4.0*fq;
+		jy -= fq;
+		m6 += 4.0*fq;
+		m9 -= fq;
+		m10 += 2.0*fq;
+		m11 += fq;
+		m12 -= 2.0*fq;
+
+		// q=5
+		fq = dist[6*Np+n];
+		rho += fq;
+		m1 -= 11.0*fq;
+		m2 -= 4.0*fq;
+		jz = fq;
+		m8 = -4.0*fq;
+		m9 -= fq;
+		m10 += 2.0*fq;
+		m11 -= fq;
+		m12 += 2.0*fq;
+
+		// q = 6
+		fq = dist[5*Np+n];
+		rho+= fq;
+		m1 -= 11.0*fq;
+		m2 -= 4.0*fq;
+		jz -= fq;
+		m8 += 4.0*fq;
+		m9 -= fq;
+		m10 += 2.0*fq;
+		m11 -= fq;
+		m12 += 2.0*fq;
+
+		// q=7
+		fq = dist[8*Np+n];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jx += fq;
+		m4 += fq;
+		jy += fq;
+		m6 += fq;
+		m9  += fq;
+		m10 += fq;
+		m11 += fq;
+		m12 += fq;
+		m13 = fq;
+		m16 = fq;
+		m17 = -fq;
+
+		// q = 8
+		fq = dist[7*Np+n];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jx -= fq;
+		m4 -= fq;
+		jy -= fq;
+		m6 -= fq;
+		m9 += fq;
+		m10 += fq;
+		m11 += fq;
+		m12 += fq;
+		m13 += fq;
+		m16 -= fq;
+		m17 += fq;
+
+		// q=9
+		fq = dist[10*Np+n];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jx += fq;
+		m4 += fq;
+		jy -= fq;
+		m6 -= fq;
+		m9 += fq;
+		m10 += fq;
+		m11 += fq;
+		m12 += fq;
+		m13 -= fq;
+		m16 += fq;
+		m17 += fq;
+
+		// q = 10
+		fq = dist[9*Np+n];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jx -= fq;
+		m4 -= fq;
+		jy += fq;
+		m6 += fq;
+		m9 += fq;
+		m10 += fq;
+		m11 += fq;
+		m12 += fq;
+		m13 -= fq;
+		m16 -= fq;
+		m17 -= fq;
+
+		// q=11
+		fq = dist[12*Np+n];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jx += fq;
+		m4 += fq;
+		jz += fq;
+		m8 += fq;
+		m9 += fq;
+		m10 += fq;
+		m11 -= fq;
+		m12 -= fq;
+		m15 = fq;
+		m16 -= fq;
+		m18 = fq;
+
+		// q=12
+		fq = dist[11*Np+n];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jx -= fq;
+		m4 -= fq;
+		jz -= fq;
+		m8 -= fq;
+		m9 += fq;
+		m10 += fq;
+		m11 -= fq;
+		m12 -= fq;
+		m15 += fq;
+		m16 += fq;
+		m18 -= fq;
+
+		// q=13
+		fq = dist[14*Np+n];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jx += fq;
+		m4 += fq;
+		jz -= fq;
+		m8 -= fq;
+		m9 += fq;
+		m10 += fq;
+		m11 -= fq;
+		m12 -= fq;
+		m15 -= fq;
+		m16 -= fq;
+		m18 -= fq;
+
+		// q=14
+		fq = dist[13*Np+n];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jx -= fq;
+		m4 -= fq;
+		jz += fq;
+		m8 += fq;
+		m9 += fq;
+		m10 += fq;
+		m11 -= fq;
+		m12 -= fq;
+		m15 -= fq;
+		m16 += fq;
+		m18 += fq;
+
+		// q=15
+		fq = dist[16*Np+n];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jy += fq;
+		m6 += fq;
+		jz += fq;
+		m8 += fq;
+		m9 -= 2.0*fq;
+		m10 -= 2.0*fq;
+		m14 = fq;
+		m17 += fq;
+		m18 -= fq;
+
+		// q=16
+		fq = dist[15*Np+n];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jy -= fq;
+		m6 -= fq;
+		jz -= fq;
+		m8 -= fq;
+		m9 -= 2.0*fq;
+		m10 -= 2.0*fq;
+		m14 += fq;
+		m17 -= fq;
+		m18 += fq;
+
+		// q=17
+		fq = dist[18*Np+n];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jy += fq;
+		m6 += fq;
+		jz -= fq;
+		m8 -= fq;
+		m9 -= 2.0*fq;
+		m10 -= 2.0*fq;
+		m14 -= fq;
+		m17 += fq;
+		m18 += fq;
+
+		// q=18
+		fq = dist[17*Np+n];
+		rho += fq;
+		m1 += 8.0*fq;
+		m2 += fq;
+		jy -= fq;
+		m6 -= fq;
+		jz += fq;
+		m8 += fq;
+		m9 -= 2.0*fq;
+		m10 -= 2.0*fq;
+		m14 -= fq;
+		m17 -= fq;
+		m18 -= fq;
+
+        // Compute greyscale related parameters
+        ux = (jx/rho0+0.5*porosity*Gx+0.5*Fcpx/rho0)/(1.0+0.5*porosity*mu_eff/perm);
+        uy = (jy/rho0+0.5*porosity*Gy+0.5*Fcpy/rho0)/(1.0+0.5*porosity*mu_eff/perm);
+        uz = (jz/rho0+0.5*porosity*Gz+0.5*Fcpz/rho0)/(1.0+0.5*porosity*mu_eff/perm);
+        if (porosity==1.0){//i.e. open nodes
+            ux = (jx/rho0+0.5*porosity*Gx);
+            uy = (jy/rho0+0.5*porosity*Gy);
+            uz = (jz/rho0+0.5*porosity*Gz);
+        }
+
+        //Update the total force to include linear (Darcy) and nonlinear (Forchheimer) drags due to the porous medium
+        Fx = rho0*(-porosity*mu_eff/perm*ux + porosity*Gx)+Fcpx;
+        Fy = rho0*(-porosity*mu_eff/perm*uy + porosity*Gy)+Fcpy;
+        Fz = rho0*(-porosity*mu_eff/perm*uz + porosity*Gz)+Fcpz;
+        if (porosity==1.0){
+            Fx=rho0*(porosity*Gx);
+            Fy=rho0*(porosity*Gy);
+            Fz=rho0*(porosity*Gz);
+        }
+
+		// write the velocity 
+		Velocity[n] = ux;
+		Velocity[Np+n] = uy;
+		Velocity[2*Np+n] = uz;
+        //Pressure[n] = rho/3.f/porosity;
+        Pressure[n] = rho/3.f;
+
+		//........................................................................
+		//..............carry out relaxation process..............................
+		//..........Toelke, Fruediger et. al. 2006................................
+		if (C == 0.0)	nx = ny = nz = 0.0;
+		m1 = m1 + rlx_setA*((19*(ux*ux+uy*uy+uz*uz)*rho0/porosity - 11*rho) -19*alpha*C - m1);
+		m2 = m2 + rlx_setA*((3*rho - 5.5*(ux*ux+uy*uy+uz*uz)*rho0/porosity)- m2);
+        jx = jx + Fx;
+		m4 = m4 + rlx_setB*((-0.6666666666666666*ux*rho0)- m4)
+                + (1-0.5*rlx_setB)*(-0.6666666666666666*Fx);
+        jy = jy + Fy;
+		m6 = m6 + rlx_setB*((-0.6666666666666666*uy*rho0)- m6)
+                + (1-0.5*rlx_setB)*(-0.6666666666666666*Fy);
+        jz = jz + Fz;
+		m8 = m8 + rlx_setB*((-0.6666666666666666*uz*rho0)- m8)
+                + (1-0.5*rlx_setB)*(-0.6666666666666666*Fz);
+		m9 = m9 + rlx_setA*(((2*ux*ux-uy*uy-uz*uz)*rho0/porosity) + 0.5*alpha*C*(2*nx*nx-ny*ny-nz*nz) - m9);
+		m10 = m10 + rlx_setA*( - m10);
+        //m10 = m10 + rlx_setA*(-0.5*rho0*((2*ux*ux-uy*uy-uz*uz)/porosity)- m10);
+		m11 = m11 + rlx_setA*(((uy*uy-uz*uz)*rho0/porosity) + 0.5*alpha*C*(ny*ny-nz*nz)- m11);
+		m12 = m12 + rlx_setA*( - m12);
+        //m12 = m12 + rlx_setA*(-0.5*(rho0*(uy*uy-uz*uz)/porosity)- m12);
+		m13 = m13 + rlx_setA*( (ux*uy*rho0/porosity) + 0.5*alpha*C*nx*ny - m13);
+		m14 = m14 + rlx_setA*( (uy*uz*rho0/porosity) + 0.5*alpha*C*ny*nz - m14);
+		m15 = m15 + rlx_setA*( (ux*uz*rho0/porosity) + 0.5*alpha*C*nx*nz - m15);
+		m16 = m16 + rlx_setB*( - m16);
+		m17 = m17 + rlx_setB*( - m17);
+		m18 = m18 + rlx_setB*( - m18);
+
+		//.................inverse transformation......................................................
+		// q=0
+		fq = mrt_V1*rho-mrt_V2*m1+mrt_V3*m2;
+		dist[n] = fq;
+
+		// q = 1
+		fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(jx-m4)+mrt_V6*(m9-m10);
+		dist[1*Np+n] = fq;
+
+		// q=2
+		fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(m4-jx)+mrt_V6*(m9-m10);
+		dist[2*Np+n] = fq;
+
+		// q = 3
+		fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(jy-m6)+mrt_V7*(m10-m9)+mrt_V8*(m11-m12);
+		dist[3*Np+n] = fq;
+
+		// q = 4
+		fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(m6-jy)+mrt_V7*(m10-m9)+mrt_V8*(m11-m12);
+		dist[4*Np+n] = fq;
+
+		// q = 5
+		fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(jz-m8)+mrt_V7*(m10-m9)+mrt_V8*(m12-m11);
+		dist[5*Np+n] = fq;
+
+		// q = 6
+		fq = mrt_V1*rho-mrt_V4*m1-mrt_V5*m2+0.1*(m8-jz)+mrt_V7*(m10-m9)+mrt_V8*(m12-m11);
+		dist[6*Np+n] = fq;
+
+		// q = 7
+		fq = mrt_V1*rho+mrt_V9*m1+mrt_V10*m2+0.1*(jx+jy)+0.025*(m4+m6)+
+				mrt_V7*m9+mrt_V11*m10+mrt_V8*m11+mrt_V12*m12+0.25*m13+0.125*(m16-m17);
+		dist[7*Np+n] = fq;
+
+
+		// q = 8
+		fq = mrt_V1*rho+mrt_V9*m1+mrt_V10*m2-0.1*(jx+jy)-0.025*(m4+m6) +mrt_V7*m9+mrt_V11*m10+mrt_V8*m11
+				+mrt_V12*m12+0.25*m13+0.125*(m17-m16);
+		dist[8*Np+n] = fq;
+
+		// q = 9
+		fq = mrt_V1*rho+mrt_V9*m1+mrt_V10*m2+0.1*(jx-jy)+0.025*(m4-m6)+
+				mrt_V7*m9+mrt_V11*m10+mrt_V8*m11+mrt_V12*m12-0.25*m13+0.125*(m16+m17);
+		dist[9*Np+n] = fq;
+
+		// q = 10
+		fq = mrt_V1*rho+mrt_V9*m1+mrt_V10*m2+0.1*(jy-jx)+0.025*(m6-m4)+
+				mrt_V7*m9+mrt_V11*m10+mrt_V8*m11+mrt_V12*m12-0.25*m13-0.125*(m16+m17);
+		dist[10*Np+n] = fq;
+
+
+		// q = 11
+		fq = mrt_V1*rho+mrt_V9*m1
+				+mrt_V10*m2+0.1*(jx+jz)+0.025*(m4+m8)
+				+mrt_V7*m9+mrt_V11*m10-mrt_V8*m11
+				-mrt_V12*m12+0.25*m15+0.125*(m18-m16);
+		dist[11*Np+n] = fq;
+
+		// q = 12
+		fq = mrt_V1*rho+mrt_V9*m1+mrt_V10*m2-0.1*(jx+jz)-0.025*(m4+m8)+
+				mrt_V7*m9+mrt_V11*m10-mrt_V8*m11-mrt_V12*m12+0.25*m15+0.125*(m16-m18);
+		dist[12*Np+n] = fq;
+
+		// q = 13
+		fq = mrt_V1*rho+mrt_V9*m1
+				+mrt_V10*m2+0.1*(jx-jz)+0.025*(m4-m8)
+				+mrt_V7*m9+mrt_V11*m10-mrt_V8*m11
+				-mrt_V12*m12-0.25*m15-0.125*(m16+m18);
+		dist[13*Np+n] = fq;
+
+		// q= 14
+		fq = mrt_V1*rho+mrt_V9*m1
+				+mrt_V10*m2+0.1*(jz-jx)+0.025*(m8-m4)
+				+mrt_V7*m9+mrt_V11*m10-mrt_V8*m11
+				-mrt_V12*m12-0.25*m15+0.125*(m16+m18);
+
+		dist[14*Np+n] = fq;
+
+		// q = 15
+		fq = mrt_V1*rho+mrt_V9*m1
+				+mrt_V10*m2+0.1*(jy+jz)+0.025*(m6+m8)
+				-mrt_V6*m9-mrt_V7*m10+0.25*m14+0.125*(m17-m18);
+		dist[15*Np+n] = fq;
+
+		// q = 16
+		fq =  mrt_V1*rho+mrt_V9*m1
+				+mrt_V10*m2-0.1*(jy+jz)-0.025*(m6+m8)
+				-mrt_V6*m9-mrt_V7*m10+0.25*m14+0.125*(m18-m17);
+		dist[16*Np+n] = fq;
+
+
+		// q = 17
+		fq = mrt_V1*rho+mrt_V9*m1
+				+mrt_V10*m2+0.1*(jy-jz)+0.025*(m6-m8)
+				-mrt_V6*m9-mrt_V7*m10-0.25*m14+0.125*(m17+m18);
+		dist[17*Np+n] = fq;
+
+		// q = 18
+		fq = mrt_V1*rho+mrt_V9*m1
+				+mrt_V10*m2+0.1*(jz-jy)+0.025*(m8-m6)
+				-mrt_V6*m9-mrt_V7*m10-0.25*m14-0.125*(m17+m18);
+		dist[18*Np+n] = fq;
+		//........................................................................
+
+		// Instantiate mass transport distributions
+		// Stationary value - distribution 0
+		nAB = 1.0/(nA+nB);
+		Aq[n] = 0.3333333333333333*nA;
+		Bq[n] = 0.3333333333333333*nB;
+
+		//...............................................
+		// q = 0,2,4
+		// Cq = {1,0,0}, {0,1,0}, {0,0,1}
+		delta = beta*nA*nB*nAB*0.1111111111111111*nx;
+		if (!(nA*nB*nAB>0)) delta=0;
+        if (RecoloringOff==true && porosity !=1.0) delta=0;
+		a1 = nA*(0.1111111111111111*(1+4.5*ux))+delta;
+		b1 = nB*(0.1111111111111111*(1+4.5*ux))-delta;
+		a2 = nA*(0.1111111111111111*(1-4.5*ux))-delta;
+		b2 = nB*(0.1111111111111111*(1-4.5*ux))+delta;
+
+		Aq[1*Np+n] = a1;
+		Bq[1*Np+n] = b1;
+		Aq[2*Np+n] = a2;
+		Bq[2*Np+n] = b2;
+
+		//...............................................
+		// q = 2
+		// Cq = {0,1,0}
+		delta = beta*nA*nB*nAB*0.1111111111111111*ny;
+		if (!(nA*nB*nAB>0)) delta=0;
+        if (RecoloringOff==true && porosity !=1.0) delta=0;
+		a1 = nA*(0.1111111111111111*(1+4.5*uy))+delta;
+		b1 = nB*(0.1111111111111111*(1+4.5*uy))-delta;
+		a2 = nA*(0.1111111111111111*(1-4.5*uy))-delta;
+		b2 = nB*(0.1111111111111111*(1-4.5*uy))+delta;
+
+		Aq[3*Np+n] = a1;
+		Bq[3*Np+n] = b1;
+		Aq[4*Np+n] = a2;
+		Bq[4*Np+n] = b2;
+		//...............................................
+		// q = 4
+		// Cq = {0,0,1}
+		delta = beta*nA*nB*nAB*0.1111111111111111*nz;
+		if (!(nA*nB*nAB>0)) delta=0;
+        if (RecoloringOff==true && porosity !=1.0) delta=0;
+		a1 = nA*(0.1111111111111111*(1+4.5*uz))+delta;
+		b1 = nB*(0.1111111111111111*(1+4.5*uz))-delta;
+		a2 = nA*(0.1111111111111111*(1-4.5*uz))-delta;
+		b2 = nB*(0.1111111111111111*(1-4.5*uz))+delta;
+
+		Aq[5*Np+n] = a1;
+		Bq[5*Np+n] = b1;
+		Aq[6*Np+n] = a2;
+		Bq[6*Np+n] = b2;
+		//...............................................
+	}
+}
+
+
 extern "C" void ScaLBL_PhaseField_InitFromRestart(double *Den, double *Aq, double *Bq, int start, int finish, int Np){
 	int idx;
 	double nA,nB;
diff --git a/cuda/GreyscaleColor.cu b/cuda/GreyscaleColor.cu
index 089d8fdc..de4b9e0d 100644
--- a/cuda/GreyscaleColor.cu
+++ b/cuda/GreyscaleColor.cu
@@ -3021,30 +3021,33 @@ __global__ void dvc_ScaLBL_PhaseField_InitFromRestart(double *Den, double *Aq, d
 	}
 }
 
-__global__  void dvc_ScaLBL_Update_GreyscalePotential(int *Map, double *Phi, double *Psi, double *Poro, double *Perm, double alpha, double W, 
-                                                      int start, int finish, int Np){
-	int idx,n;
-    double phi,psi;
-    double cap_penalty;
-    double porosity,perm;
-
-	int S = Np/NBLOCKS/NTHREADS + 1;
-	for (int s=0; s<S; s++){
-		//........Get 1-D index for this thread....................
-		n =  S*blockIdx.x*blockDim.x + s*blockDim.x + threadIdx.x + start;
-		if (n<finish) {
-			idx = Map[n];
-			phi = Phi[idx];
-            porosity = Poro[n];
-            perm = Perm[n];    
-            cap_penalty = 1.0;
-
-            if (porosity!=1.0) cap_penalty = alpha*W/sqrt(perm);
-            psi = cap_penalty*phi;
-            Psi[idx] = psi;
-		}
-	}
-}
+//NOTE: so far it seems that we don't need this greyscale potental update;
+//      if we compute a grey-potential first, and take its gradient to work out the capillary penalty force, it is highly unstable;
+//      this is because the grey-potential is simply a scaling of the normal phase field, but such scaling create some artificial gradient at the open-grey interface
+//__global__  void dvc_ScaLBL_Update_GreyscalePotential(int *Map, double *Phi, double *Psi, double *Poro, double *Perm, double alpha, double W, 
+//                                                      int start, int finish, int Np){
+//	int idx,n;
+//    double phi,psi;
+//    double cap_penalty;
+//    double porosity,perm;
+//
+//	int S = Np/NBLOCKS/NTHREADS + 1;
+//	for (int s=0; s<S; s++){
+//		//........Get 1-D index for this thread....................
+//		n =  S*blockIdx.x*blockDim.x + s*blockDim.x + threadIdx.x + start;
+//		if (n<finish) {
+//			idx = Map[n];
+//			phi = Phi[idx];
+//            porosity = Poro[n];
+//            perm = Perm[n];    
+//            cap_penalty = 1.0;
+//
+//            if (porosity!=1.0) cap_penalty = alpha*W/sqrt(perm);
+//            psi = cap_penalty*phi;
+//            Psi[idx] = psi;
+//		}
+//	}
+//}
 
 
 ////Model-2&3
@@ -4597,16 +4600,16 @@ extern "C" void ScaLBL_D3Q19_AAodd_GreyscaleColor_CP(int *d_neighborList, int *M
 	}
 }
 
-extern "C" void ScaLBL_Update_GreyscalePotential(int *Map, double *Phi, double *Psi, double *Poro, double *Perm, double alpha, double W, 
-		int start, int finish, int Np){
-
-	dvc_ScaLBL_Update_GreyscalePotential<<<NBLOCKS,NTHREADS >>>(Map, Phi, Psi, Poro, Perm, alpha, W, start, finish, Np);
-
-	cudaError_t err = cudaGetLastError();
-	if (cudaSuccess != err){
-		printf("CUDA error in ScaLBL_Update_GreyscalePotential: %s \n",cudaGetErrorString(err));
-	}
-}
+//extern "C" void ScaLBL_Update_GreyscalePotential(int *Map, double *Phi, double *Psi, double *Poro, double *Perm, double alpha, double W, 
+//		int start, int finish, int Np){
+//
+//	dvc_ScaLBL_Update_GreyscalePotential<<<NBLOCKS,NTHREADS >>>(Map, Phi, Psi, Poro, Perm, alpha, W, start, finish, Np);
+//
+//	cudaError_t err = cudaGetLastError();
+//	if (cudaSuccess != err){
+//		printf("CUDA error in ScaLBL_Update_GreyscalePotential: %s \n",cudaGetErrorString(err));
+//	}
+//}
 
 ////Model-2&3
 //extern "C" void ScaLBL_D3Q19_AAeven_GreyscaleColor(int *Map, double *dist, double *Aq, double *Bq, double *Den,