diff --git a/common/Domain.cpp b/common/Domain.cpp
index 7da902c6..03d0c5ca 100644
--- a/common/Domain.cpp
+++ b/common/Domain.cpp
@@ -435,7 +435,7 @@ void Domain::Decomp( const std::string& Filename )
 		}
 
 		if (inlet_layers_z > 0){
-			printf("Checkerboard pattern at z inlet for %i layers \n",inlet_layers_z);
+			printf("Checkerboard pattern at z inlet for %i layers, saturated with phase label=%i \n",inlet_layers_z,inlet_layers_phase);
 			// use checkerboard pattern
 			for (int k = zStart; k < zStart+inlet_layers_z; k++){
 				for (int j = 0; j<global_Ny; j++){
@@ -494,7 +494,7 @@ void Domain::Decomp( const std::string& Filename )
 		}
 
 		if (outlet_layers_z > 0){
-			printf("Checkerboard pattern at z outlet for %i layers \n",outlet_layers_z);
+			printf("Checkerboard pattern at z outlet for %i layers, saturated with phase label=%i \n",outlet_layers_z,outlet_layers_phase);
 			// use checkerboard pattern
 			for (int k = zStart + nz*nprocz - outlet_layers_z; k < zStart + nz*nprocz; k++){
 				for (int j = 0; j<global_Ny; j++){
@@ -596,23 +596,33 @@ void Domain::Decomp( const std::string& Filename )
 	//.........................................................
 	// If external boundary conditions are applied remove solid
 	if (BoundaryCondition >  0  && kproc() == 0){
-    	if (inlet_layers_z < 4)	inlet_layers_z=4;
+    	if (inlet_layers_z < 4){
+            inlet_layers_z=4;
+            if(RANK==0){
+                printf("NOTE:Non-periodic BC is applied, but the number of Z-inlet layers is not specified (or is smaller than 3 voxels) \n     the number of Z-inlet layer is reset to %i voxels, saturated with phase label=%i \n",inlet_layers_z-1,inlet_layers_phase);
+            } 
+        }	
 		for (int k=0; k<inlet_layers_z; k++){
 			for (int j=0;j<Ny;j++){
 				for (int i=0;i<Nx;i++){
 					int n = k*Nx*Ny+j*Nx+i;
-					id[n] = 1;
+					id[n] = inlet_layers_phase;
 				}                    
 			}
  		}
  	}
     if (BoundaryCondition >  0  && kproc() == nprocz-1){
-    	if (outlet_layers_z < 4)	outlet_layers_z=4;
+    	if (outlet_layers_z < 4){
+            outlet_layers_z=4;
+            if(RANK==nprocs-1){
+                printf("NOTE:Non-periodic BC is applied, but the number of Z-outlet layers is not specified (or is smaller than 3 voxels) \n     the number of Z-outlet layer is reset to %i voxels, saturated with phase label=%i \n",outlet_layers_z-1,outlet_layers_phase);
+            } 
+        }	
  		for (int k=Nz-outlet_layers_z; 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;
- 					id[n] = 2;
+ 					id[n] = outlet_layers_phase;
  				}                    
  			}
  		}
@@ -632,7 +642,6 @@ void Domain::Decomp( const std::string& Filename )
     porosity = sum*iVol_global;
     if (rank()==0) printf("Media porosity = %f \n",porosity);
  	//.........................................................
-
 }
 
 void Domain::AggregateLabels( const std::string& filename ){
diff --git a/common/ScaLBL.h b/common/ScaLBL.h
index 0d2ee0cf..901e0e3b 100644
--- a/common/ScaLBL.h
+++ b/common/ScaLBL.h
@@ -46,6 +46,7 @@ extern "C" void ScaLBL_UnpackDenD3Q7(int *list, int count, double *recvbuf, int
 
 extern "C" void ScaLBL_D3Q19_Init(double *Dist, int Np);
 
+
 extern "C" void ScaLBL_D3Q19_Momentum(double *dist, double *vel, int Np);
 
 extern "C" void ScaLBL_D3Q19_Pressure(double *dist, double *press, int Np);
@@ -55,6 +56,23 @@ extern "C" void ScaLBL_D3Q19_AAeven_BGK(double *dist, int start, int finish, int
 
 extern "C" void ScaLBL_D3Q19_AAodd_BGK(int *neighborList, double *dist, int start, int finish, int Np, double rlx, double Fx, double Fy, double Fz);
 
+// GREYSCALE MODEL
+
+extern "C" void ScaLBL_D3Q19_GreyIMRT_Init(double *Dist, int Np, double Den);
+
+extern "C" void ScaLBL_D3Q19_AAeven_Greyscale(double *dist, int start, int finish, int Np, double rlx, double rlx_eff, double Fx, double Fy, double Fz,
+                                              double *Poros,double *Perm, double *Velocity,double *Pressure);
+
+extern "C" void ScaLBL_D3Q19_AAodd_Greyscale(int *neighborList, double *dist, int start, int finish, int Np, double rlx, double rlx_eff, double Fx, double Fy, double Fz, 
+                                             double *Poros,double *Perm, double *Velocity,double *Pressure);
+
+extern "C" void ScaLBL_D3Q19_AAeven_Greyscale_IMRT(double *dist, int start, int finish, int Np, double rlx, double rlx_eff, double Fx, double Fy, double Fz,
+                                              double *Poros,double *Perm, double *Velocity,double Den,double *Pressure);
+
+extern "C" void ScaLBL_D3Q19_AAodd_Greyscale_IMRT(int *neighborList, double *dist, int start, int finish, int Np, double rlx, double rlx_eff, double Fx, double Fy, double Fz, 
+                                             double *Poros,double *Perm, double *Velocity,double Den,double *Pressure);
+
+
 // MRT MODEL
 extern "C" void ScaLBL_D3Q19_AAeven_MRT(double *dist, int start, int finish, int Np, double rlx_setA, double rlx_setB, double Fx,
 		double Fy, double Fz);
diff --git a/cpu/D3Q19.cpp b/cpu/D3Q19.cpp
index 564eb96d..2c0e686d 100644
--- a/cpu/D3Q19.cpp
+++ b/cpu/D3Q19.cpp
@@ -84,6 +84,33 @@ extern "C" void ScaLBL_D3Q19_Init(double *dist, int Np)
 	}
 }
 
+
+extern "C" void ScaLBL_D3Q19_GreyIMRT_Init(double *dist, int Np, double Den)
+{
+	int n;
+	for (n=0; n<Np; n++){
+		dist[n] = Den - 0.6666666666666667;
+		dist[Np+n] = 0.055555555555555555;		//double(100*n)+1.f;
+		dist[2*Np+n] = 0.055555555555555555;	//double(100*n)+2.f;
+		dist[3*Np+n] = 0.055555555555555555;	//double(100*n)+3.f;
+		dist[4*Np+n] = 0.055555555555555555;	//double(100*n)+4.f;
+		dist[5*Np+n] = 0.055555555555555555;	//double(100*n)+5.f;
+		dist[6*Np+n] = 0.055555555555555555;	//double(100*n)+6.f;
+		dist[7*Np+n] = 0.0277777777777778;   //double(100*n)+7.f;
+		dist[8*Np+n] = 0.0277777777777778;   //double(100*n)+8.f;
+		dist[9*Np+n] = 0.0277777777777778;   //double(100*n)+9.f;
+		dist[10*Np+n] = 0.0277777777777778;  //double(100*n)+10.f;
+		dist[11*Np+n] = 0.0277777777777778;  //double(100*n)+11.f;
+		dist[12*Np+n] = 0.0277777777777778;  //double(100*n)+12.f;
+		dist[13*Np+n] = 0.0277777777777778;  //double(100*n)+13.f;
+		dist[14*Np+n] = 0.0277777777777778;  //double(100*n)+14.f;
+		dist[15*Np+n] = 0.0277777777777778;  //double(100*n)+15.f;
+		dist[16*Np+n] = 0.0277777777777778;  //double(100*n)+16.f;
+		dist[17*Np+n] = 0.0277777777777778;  //double(100*n)+17.f;
+		dist[18*Np+n] = 0.0277777777777778;  //double(100*n)+18.f;
+	}
+}
+
 //*************************************************************************
 extern "C" void ScaLBL_D3Q19_Swap(char *ID, double *disteven, double *distodd, int Nx, int Ny, int Nz)
 {
diff --git a/cpu/Greyscale.cpp b/cpu/Greyscale.cpp
new file mode 100644
index 00000000..b4b017c8
--- /dev/null
+++ b/cpu/Greyscale.cpp
@@ -0,0 +1,1567 @@
+#include <math.h>
+
+extern "C" void ScaLBL_D3Q19_AAeven_Greyscale(double *dist, int start, int finish, int Np, double rlx, double rlx_eff, double Gx, double Gy, double Gz,
+                                              double *Poros,double *Perm, double *Velocity, double *Pressure){
+	int n;
+	// conserved momemnts
+	double rho,vx,vy,vz,v_mag;
+    double ux,uy,uz,u_mag;
+    double pressure;
+    //double uu;
+	// non-conserved moments
+	double f0,f1,f2,f3,f4,f5,f6,f7,f8,f9,f10,f11,f12,f13,f14,f15,f16,f17,f18;
+    double GeoFun;//geometric function from Guo's PRE 66, 036304 (2002)
+    double porosity;
+    double perm;//voxel permeability
+    double c0, c1; //Guo's model parameters
+    double mu_eff = (1.0/rlx_eff-0.5)/3.0;//kinematic viscosity
+    double Fx, Fy, Fz;//The total body force including Brinkman force and user-specified (Gx,Gy,Gz)
+
+	for (int n=start; n<finish; n++){
+		// q=0
+		f0 = dist[n];
+		f1 = dist[2*Np+n];
+		f2 = dist[1*Np+n];
+		f3 = dist[4*Np+n];
+		f4 = dist[3*Np+n];
+		f5 = dist[6*Np+n];
+		f6 = dist[5*Np+n];
+		f7 = dist[8*Np+n];
+		f8 = dist[7*Np+n];
+		f9 = dist[10*Np+n];
+		f10 = dist[9*Np+n];
+		f11 = dist[12*Np+n];
+		f12 = dist[11*Np+n];
+		f13 = dist[14*Np+n];
+		f14 = dist[13*Np+n];
+		f15 = dist[16*Np+n];
+		f16 = dist[15*Np+n];
+		f17 = dist[18*Np+n];
+		f18 = dist[17*Np+n];
+        
+        porosity = Poros[n];
+        perm = Perm[n];
+
+        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
+
+		rho = f0+f2+f1+f4+f3+f6+f5+f8+f7+f10+f9+f12+f11+f14+f13+f16+f15+f18+f17;
+        pressure = rho/porosity/3.0;
+		vx = (f1-f2+f7-f8+f9-f10+f11-f12+f13-f14)/rho+0.5*porosity*Gx;
+		vy = (f3-f4+f7-f8-f9+f10+f15-f16+f17-f18)/rho+0.5*porosity*Gy;
+		vz = (f5-f6+f11-f12-f13+f14+f15-f16-f17+f18)/rho+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 = -porosity*mu_eff/perm*ux - porosity*GeoFun/sqrt(perm)*u_mag*ux + porosity*Gx;
+        Fy = -porosity*mu_eff/perm*uy - porosity*GeoFun/sqrt(perm)*u_mag*uy + porosity*Gy;
+        Fz = -porosity*mu_eff/perm*uz - porosity*GeoFun/sqrt(perm)*u_mag*uz + porosity*Gz;
+        if (porosity==1.0){
+            Fx=Gx;
+            Fy=Gy;
+            Fz=Gz;
+        }
+
+        //------------------------ BGK collison where body force has higher-order terms ----------------------------------------------------------//
+//		// q=0
+//		dist[n] = f0*(1.0-rlx)+ rlx*0.3333333333333333*rho*(1. - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//                  + 0.3333333333333333*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(0. - (3.*uy)/porosity) + Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 1
+//		dist[1*Np+n] = f1*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*ux + (4.5*ux*ux)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//            +0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(3. + (6.*ux)/porosity) + Fy*(0. - (3.*uy)/porosity) + Fz*(0. - (3.*uz)/porosity));
+//
+//		// q=2
+//		dist[2*Np+n] = f2*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*ux + (4.5*ux*ux)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//            +0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(-3. + (6.*ux)/porosity) + Fy*(0. - (3.*uy)/porosity) + Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 3
+//		dist[3*Np+n] = f3*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*uy + (4.5*uy*uy)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(3. + (6.*uy)/porosity) + Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 4
+//		dist[4*Np+n] = f4*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*uy + (4.5*uy*uy)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(-3. + (6.*uy)/porosity) + Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 5
+//		dist[5*Np+n] = f5*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*uz + (4.5*uz*uz)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(0. - (3.*uy)/porosity) + Fz*(3. + (6.*uz)/porosity));
+//
+//		// q = 6
+//		dist[6*Np+n] = f6*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*uz + (4.5*uz*uz)/porosity - (1.5*(ux*ux+ uy*uy + uz*uz))/porosity)
+//				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(0. - (3.*uy)/porosity) + Fz*(-3. + (6.*uz)/porosity));
+//
+//		// q = 7
+//		dist[7*Np+n] = f7*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux + uy) + (4.5*(ux + uy)*(ux + uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3. - (3.*ux)/porosity + (9.*(ux + uy))/porosity) + Fy*(3. - (3.*uy)/porosity + (9.*(ux + uy))/porosity) + 
+//  Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 8
+//		dist[8*Np+n] = f8*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux - uy) + (4.5*(-ux - uy)*(-ux - uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3. - (3.*ux)/porosity - (9.*(-ux - uy))/porosity) + Fy*(-3. - (9.*(-ux - uy))/porosity - (3.*uy)/porosity) + 
+//  Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 9
+//		dist[9*Np+n] = f9*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux - uy) + (4.5*(ux - uy)*(ux - uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3. - (3.*ux)/porosity + (9.*(ux - uy))/porosity) + Fy*(-3. - (9.*(ux - uy))/porosity - (3.*uy)/porosity) + 
+//  Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 10
+//		dist[10*Np+n] = f10*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux + uy) + (4.5*(-ux + uy)*(-ux + uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3. - (3.*ux)/porosity - (9.*(-ux + uy))/porosity) + Fy*(3. - (3.*uy)/porosity + (9.*(-ux + uy))/porosity) + 
+//  Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 11
+//		dist[11*Np+n] = f11*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux + uz) + (4.5*(ux + uz)*(ux + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(0. - (3.*uy)/porosity) + Fx*(3. - (3.*ux)/porosity + (9.*(ux + uz))/porosity) + 
+//  Fz*(3. - (3.*uz)/porosity + (9.*(ux + uz))/porosity));
+//
+//		// q = 12
+//		dist[12*Np+n] = f12*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux - uz) + (4.5*(-ux - uz)*(-ux - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(0. - (3.*uy)/porosity) + Fx*(-3. - (3.*ux)/porosity - (9.*(-ux - uz))/porosity) + 
+//  Fz*(-3. - (9.*(-ux - uz))/porosity - (3.*uz)/porosity));
+//
+//		// q = 13
+//		dist[13*Np+n] = f13*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux - uz) + (4.5*(ux - uz)*(ux - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(0. - (3.*uy)/porosity) + Fx*(3. - (3.*ux)/porosity + (9.*(ux - uz))/porosity) + 
+//  Fz*(-3. - (9.*(ux - uz))/porosity - (3.*uz)/porosity));
+//
+//		// q= 14
+//		dist[14*Np+n] = f14*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux + uz) + (4.5*(-ux + uz)*(-ux + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(0. - (3.*uy)/porosity) + Fx*(-3. - (3.*ux)/porosity - (9.*(-ux + uz))/porosity) + 
+//  Fz*(3. - (3.*uz)/porosity + (9.*(-ux + uz))/porosity));
+//
+//		// q = 15
+//		dist[15*Np+n] = f15*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(uy + uz) + (4.5*(uy + uz)*(uy + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(3. - (3.*uy)/porosity + (9.*(uy + uz))/porosity) + 
+//  Fz*(3. - (3.*uz)/porosity + (9.*(uy + uz))/porosity));
+//
+//		// q = 16
+//		dist[16*Np+n] = f16*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-uy - uz) + (4.5*(-uy - uz)*(-uy - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(-3. - (3.*uy)/porosity - (9.*(-uy - uz))/porosity) + 
+//  Fz*(-3. - (9.*(-uy - uz))/porosity - (3.*uz)/porosity));
+//
+//		// q = 17
+//		dist[17*Np+n] = f17*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(uy - uz) + (4.5*(uy - uz)*(uy - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(3. - (3.*uy)/porosity + (9.*(uy - uz))/porosity) + 
+//  Fz*(-3. - (9.*(uy - uz))/porosity - (3.*uz)/porosity));
+//
+//		// q = 18
+//		dist[18*Np+n] = f18*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-uy + uz) + (4.5*(-uy + uz)*(-uy + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(-3. - (3.*uy)/porosity - (9.*(-uy + uz))/porosity) + 
+//  Fz*(3. - (3.*uz)/porosity + (9.*(-uy + uz))/porosity));
+        //----------------------------------------------------------------------------------------------------------------------------------------//
+
+        //------------------------ BGK collison where body force has NO higher-order terms ----------------------------------------------------------//
+		// q=0
+		dist[n] = f0*(1.0-rlx)+ rlx*0.3333333333333333*rho*(1. - (1.5*(ux*ux + uy*uy + uz*uz))/porosity);
+
+		// q = 1
+		dist[1*Np+n] = f1*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*ux + (4.5*ux*ux)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+            +0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(3.));
+
+		// q=2
+		dist[2*Np+n] = f2*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*ux + (4.5*ux*ux)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+            +0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(-3.));
+
+		// q = 3
+		dist[3*Np+n] = f3*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*uy + (4.5*uy*uy)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fy*(3.));
+
+		// q = 4
+		dist[4*Np+n] = f4*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*uy + (4.5*uy*uy)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fy*(-3.));
+
+		// q = 5
+		dist[5*Np+n] = f5*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*uz + (4.5*uz*uz)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fz*(3.));
+
+		// q = 6
+		dist[6*Np+n] = f6*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*uz + (4.5*uz*uz)/porosity - (1.5*(ux*ux+ uy*uy + uz*uz))/porosity)
+				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fz*(-3.));
+
+		// q = 7
+		dist[7*Np+n] = f7*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux + uy) + (4.5*(ux + uy)*(ux + uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3.) + Fy*(3.));
+
+		// q = 8
+		dist[8*Np+n] = f8*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux - uy) + (4.5*(-ux - uy)*(-ux - uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3.) + Fy*(-3.));
+
+		// q = 9
+		dist[9*Np+n] = f9*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux - uy) + (4.5*(ux - uy)*(ux - uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3.) + Fy*(-3.));
+
+		// q = 10
+		dist[10*Np+n] = f10*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux + uy) + (4.5*(-ux + uy)*(-ux + uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3.) + Fy*(3.));
+
+		// q = 11
+		dist[11*Np+n] = f11*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux + uz) + (4.5*(ux + uz)*(ux + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3.) + Fz*(3.));
+
+		// q = 12
+		dist[12*Np+n] = f12*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux - uz) + (4.5*(-ux - uz)*(-ux - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3.) + Fz*(-3.));
+
+		// q = 13
+		dist[13*Np+n] = f13*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux - uz) + (4.5*(ux - uz)*(ux - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3.) + Fz*(-3.));
+
+		// q= 14
+		dist[14*Np+n] = f14*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux + uz) + (4.5*(-ux + uz)*(-ux + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3.) + Fz*(3.));
+
+		// q = 15
+		dist[15*Np+n] = f15*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(uy + uz) + (4.5*(uy + uz)*(uy + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(3.) + Fz*(3.));
+
+		// q = 16
+		dist[16*Np+n] = f16*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-uy - uz) + (4.5*(-uy - uz)*(-uy - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(-3.) + Fz*(-3.));
+
+		// q = 17
+		dist[17*Np+n] = f17*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(uy - uz) + (4.5*(uy - uz)*(uy - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(3.) + Fz*(-3.));
+
+		// q = 18
+		dist[18*Np+n] = f18*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-uy + uz) + (4.5*(-uy + uz)*(-uy + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(-3.) + Fz*(3.));
+        //-------------------------------------------------------------------------------------------------------------------------------------------//
+        
+        //Update velocity on device
+		Velocity[0*Np+n] = ux;
+		Velocity[1*Np+n] = uy;
+		Velocity[2*Np+n] = uz;
+        //Update pressure on device
+        Pressure[n] = pressure;
+	}
+}
+
+extern "C" void ScaLBL_D3Q19_AAodd_Greyscale(int *neighborList, double *dist, int start, int finish, int Np, double rlx,  double rlx_eff, double Gx, double Gy, double Gz, 
+                                             double *Poros,double *Perm, double *Velocity,double *Pressure){
+	int n;
+	// conserved momemnts
+	double rho,vx,vy,vz,v_mag;
+    double ux,uy,uz,u_mag;
+    double pressure;
+    //double uu;
+	// non-conserved moments
+	double f0,f1,f2,f3,f4,f5,f6,f7,f8,f9,f10,f11,f12,f13,f14,f15,f16,f17,f18;
+	int nr1,nr2,nr3,nr4,nr5,nr6,nr7,nr8,nr9,nr10,nr11,nr12,nr13,nr14,nr15,nr16,nr17,nr18;
+    double GeoFun;//geometric function from Guo's PRE 66, 036304 (2002)
+    double porosity;
+    double perm;//voxel permeability
+    double c0, c1; //Guo's model parameters
+    double mu_eff = (1.0/rlx_eff-0.5)/3.0;//kinematic viscosity
+    double Fx, Fy, Fz;//The total body force including Brinkman force and user-specified (Gx,Gy,Gz)
+
+	int nread;
+	for (int n=start; n<finish; n++){
+		
+		// q=0
+		f0 = dist[n];
+		// q=1
+		nr1 = neighborList[n]; // neighbor 2 ( > 10Np => odd part of dist)
+		f1 = dist[nr1]; // reading the f1 data into register fq
+
+		nr2 = neighborList[n+Np]; // neighbor 1 ( < 10Np => even part of dist)
+		f2 = dist[nr2];  // reading the f2 data into register fq
+
+		// q=3
+		nr3 = neighborList[n+2*Np]; // neighbor 4
+		f3 = dist[nr3];
+
+		// q = 4
+		nr4 = neighborList[n+3*Np]; // neighbor 3
+		f4 = dist[nr4];
+
+		// q=5
+		nr5 = neighborList[n+4*Np];
+		f5 = dist[nr5];
+
+		// q = 6
+		nr6 = neighborList[n+5*Np];
+		f6 = dist[nr6];
+		
+		// q=7
+		nr7 = neighborList[n+6*Np];
+		f7 = dist[nr7];
+
+		// q = 8
+		nr8 = neighborList[n+7*Np];
+		f8 = dist[nr8];
+
+		// q=9
+		nr9 = neighborList[n+8*Np];
+		f9 = dist[nr9];
+
+		// q = 10
+		nr10 = neighborList[n+9*Np];
+		f10 = dist[nr10];
+
+		// q=11
+		nr11 = neighborList[n+10*Np];
+		f11 = dist[nr11];
+
+		// q=12
+		nr12 = neighborList[n+11*Np];
+		f12 = dist[nr12];
+
+		// q=13
+		nr13 = neighborList[n+12*Np];
+		f13 = dist[nr13];
+
+		// q=14
+		nr14 = neighborList[n+13*Np];
+		f14 = dist[nr14];
+
+		// q=15
+		nr15 = neighborList[n+14*Np];
+		f15 = dist[nr15];
+
+		// q=16
+		nr16 = neighborList[n+15*Np];
+		f16 = dist[nr16];
+
+		// q=17
+		//fq = dist[18*Np+n];
+		nr17 = neighborList[n+16*Np];
+		f17 = dist[nr17];
+
+		// q=18
+		nr18 = neighborList[n+17*Np];
+		f18 = dist[nr18];
+
+        porosity = Poros[n];
+        perm = Perm[n];
+
+        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
+
+		rho = f0+f2+f1+f4+f3+f6+f5+f8+f7+f10+f9+f12+f11+f14+f13+f16+f15+f18+f17;
+        pressure = rho/porosity/3.0;
+		vx = (f1-f2+f7-f8+f9-f10+f11-f12+f13-f14)/rho+0.5*porosity*Gx;
+		vy = (f3-f4+f7-f8-f9+f10+f15-f16+f17-f18)/rho+0.5*porosity*Gy;
+		vz = (f5-f6+f11-f12-f13+f14+f15-f16-f17+f18)/rho+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 = -porosity*mu_eff/perm*ux - porosity*GeoFun/sqrt(perm)*u_mag*ux + porosity*Gx;
+        Fy = -porosity*mu_eff/perm*uy - porosity*GeoFun/sqrt(perm)*u_mag*uy + porosity*Gy;
+        Fz = -porosity*mu_eff/perm*uz - porosity*GeoFun/sqrt(perm)*u_mag*uz + porosity*Gz;
+        if (porosity==1.0){
+            Fx=Gx;
+            Fy=Gy;
+            Fz=Gz;
+        }
+
+        //------------------------ BGK collison where body force has higher-order terms ----------------------------------------------------------//
+//		// q=0
+//		dist[n] = f0*(1.0-rlx) + rlx*0.3333333333333333*rho*(1. - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//                  + 0.3333333333333333*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(0. - (3.*uy)/porosity) + Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 1
+//		dist[nr2] = f1*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*ux + (4.5*ux*ux)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//            +0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(3. + (6.*ux)/porosity) + Fy*(0. - (3.*uy)/porosity) + Fz*(0. - (3.*uz)/porosity));
+//
+//		// q=2
+//		dist[nr1] = f2*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*ux + (4.5*ux*ux)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//            +0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(-3. + (6.*ux)/porosity) + Fy*(0. - (3.*uy)/porosity) + Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 3
+//		dist[nr4] = f3*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*uy + (4.5*uy*uy)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(3. + (6.*uy)/porosity) + Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 4
+//		dist[nr3] = f4*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*uy + (4.5*uy*uy)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)  
+//				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(-3. + (6.*uy)/porosity) + Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 5
+//		dist[nr6] = f5*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*uz + (4.5*uz*uz)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(0. - (3.*uy)/porosity) + Fz*(3. + (6.*uz)/porosity));
+//
+//		// q = 6
+//		dist[nr5] = f6*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*uz + (4.5*uz*uz)/porosity - (1.5*(ux*ux+ uy*uy + uz*uz))/porosity) 
+//				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(0. - (3.*uy)/porosity) + Fz*(-3. + (6.*uz)/porosity));
+//
+//		// q = 7
+//		dist[nr8] = f7*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux + uy) + (4.5*(ux + uy)*(ux + uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3. - (3.*ux)/porosity + (9.*(ux + uy))/porosity) + Fy*(3. - (3.*uy)/porosity + (9.*(ux + uy))/porosity) + 
+//  Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 8
+//		dist[nr7] = f8*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux - uy) + (4.5*(-ux - uy)*(-ux - uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3. - (3.*ux)/porosity - (9.*(-ux - uy))/porosity) + Fy*(-3. - (9.*(-ux - uy))/porosity - (3.*uy)/porosity) + 
+//  Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 9
+//		dist[nr10] = f9*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux - uy) + (4.5*(ux - uy)*(ux - uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3. - (3.*ux)/porosity + (9.*(ux - uy))/porosity) + Fy*(-3. - (9.*(ux - uy))/porosity - (3.*uy)/porosity) + 
+//  Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 10
+//		dist[nr9] = f10*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux + uy) + (4.5*(-ux + uy)*(-ux + uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3. - (3.*ux)/porosity - (9.*(-ux + uy))/porosity) + Fy*(3. - (3.*uy)/porosity + (9.*(-ux + uy))/porosity) + 
+//  Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 11
+//		dist[nr12] = f11*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux + uz) + (4.5*(ux + uz)*(ux + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(0. - (3.*uy)/porosity) + Fx*(3. - (3.*ux)/porosity + (9.*(ux + uz))/porosity) + 
+//  Fz*(3. - (3.*uz)/porosity + (9.*(ux + uz))/porosity));
+//
+//		// q = 12
+//		dist[nr11] = f12*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux - uz) + (4.5*(-ux - uz)*(-ux - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(0. - (3.*uy)/porosity) + Fx*(-3. - (3.*ux)/porosity - (9.*(-ux - uz))/porosity) + 
+//  Fz*(-3. - (9.*(-ux - uz))/porosity - (3.*uz)/porosity));
+//
+//		// q = 13
+//		dist[nr14] = f13*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux - uz) + (4.5*(ux - uz)*(ux - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(0. - (3.*uy)/porosity) + Fx*(3. - (3.*ux)/porosity + (9.*(ux - uz))/porosity) + 
+//  Fz*(-3. - (9.*(ux - uz))/porosity - (3.*uz)/porosity));
+//
+//		// q= 14
+//		dist[nr13] = f14*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux + uz) + (4.5*(-ux + uz)*(-ux + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(0. - (3.*uy)/porosity) + Fx*(-3. - (3.*ux)/porosity - (9.*(-ux + uz))/porosity) + 
+//  Fz*(3. - (3.*uz)/porosity + (9.*(-ux + uz))/porosity));
+//
+//		// q = 15
+//		dist[nr16] = f15*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(uy + uz) + (4.5*(uy + uz)*(uy + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(3. - (3.*uy)/porosity + (9.*(uy + uz))/porosity) + 
+//  Fz*(3. - (3.*uz)/porosity + (9.*(uy + uz))/porosity));
+//
+//		// q = 16
+//		dist[nr15] = f16*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-uy - uz) + (4.5*(-uy - uz)*(-uy - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(-3. - (3.*uy)/porosity - (9.*(-uy - uz))/porosity) + 
+//  Fz*(-3. - (9.*(-uy - uz))/porosity - (3.*uz)/porosity));
+//
+//		// q = 17
+//		dist[nr18] = f17*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(uy - uz) + (4.5*(uy - uz)*(uy - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(3. - (3.*uy)/porosity + (9.*(uy - uz))/porosity) + 
+//  Fz*(-3. - (9.*(uy - uz))/porosity - (3.*uz)/porosity));
+//
+//		// q = 18
+//		dist[nr17] = f18*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-uy + uz) + (4.5*(-uy + uz)*(-uy + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(-3. - (3.*uy)/porosity - (9.*(-uy + uz))/porosity) + 
+//  Fz*(3. - (3.*uz)/porosity + (9.*(-uy + uz))/porosity));
+        //----------------------------------------------------------------------------------------------------------------------------------------//
+
+
+
+        //------------------------ BGK collison where body force has NO higher-order terms ----------------------------------------------------------//
+		// q=0
+		dist[n] = f0*(1.0-rlx) + rlx*0.3333333333333333*rho*(1. - (1.5*(ux*ux + uy*uy + uz*uz))/porosity);
+
+		// q = 1
+		dist[nr2] = f1*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*ux + (4.5*ux*ux)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+            +0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(3.));
+
+		// q=2
+		dist[nr1] = f2*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*ux + (4.5*ux*ux)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+            +0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(-3.));
+
+		// q = 3
+		dist[nr4] = f3*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*uy + (4.5*uy*uy)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fy*(3.));
+
+		// q = 4
+		dist[nr3] = f4*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*uy + (4.5*uy*uy)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)  
+				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fy*(-3.));
+
+		// q = 5
+		dist[nr6] = f5*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*uz + (4.5*uz*uz)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fz*(3.));
+
+		// q = 6
+		dist[nr5] = f6*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*uz + (4.5*uz*uz)/porosity - (1.5*(ux*ux+ uy*uy + uz*uz))/porosity) 
+				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fz*(-3.));
+
+		// q = 7
+		dist[nr8] = f7*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux + uy) + (4.5*(ux + uy)*(ux + uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3.) + Fy*(3.));
+
+		// q = 8
+		dist[nr7] = f8*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux - uy) + (4.5*(-ux - uy)*(-ux - uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3.) + Fy*(-3.));
+
+		// q = 9
+		dist[nr10] = f9*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux - uy) + (4.5*(ux - uy)*(ux - uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3.) + Fy*(-3.));
+
+		// q = 10
+		dist[nr9] = f10*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux + uy) + (4.5*(-ux + uy)*(-ux + uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3.) + Fy*(3.));
+
+		// q = 11
+		dist[nr12] = f11*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux + uz) + (4.5*(ux + uz)*(ux + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3.) + Fz*(3.));
+
+		// q = 12
+		dist[nr11] = f12*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux - uz) + (4.5*(-ux - uz)*(-ux - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3.) + Fz*(-3.));
+
+		// q = 13
+		dist[nr14] = f13*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux - uz) + (4.5*(ux - uz)*(ux - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3.) + Fz*(-3.));
+
+		// q= 14
+		dist[nr13] = f14*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux + uz) + (4.5*(-ux + uz)*(-ux + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3.) + Fz*(3.));
+
+		// q = 15
+		dist[nr16] = f15*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(uy + uz) + (4.5*(uy + uz)*(uy + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(3.) + Fz*(3.));
+
+		// q = 16
+		dist[nr15] = f16*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-uy - uz) + (4.5*(-uy - uz)*(-uy - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(-3.) + Fz*(-3.));
+
+		// q = 17
+		dist[nr18] = f17*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(uy - uz) + (4.5*(uy - uz)*(uy - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(3.) + Fz*(-3.));
+
+		// q = 18
+		dist[nr17] = f18*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-uy + uz) + (4.5*(-uy + uz)*(-uy + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(-3.) + Fz*(3.));
+        //-------------------------------------------------------------------------------------------------------------------------------------------//
+
+
+
+        //Update velocity on device
+		Velocity[0*Np+n] = ux;
+		Velocity[1*Np+n] = uy;
+		Velocity[2*Np+n] = uz;
+        //Update pressure on device
+        Pressure[n] = pressure;
+	}
+}
+
+
+extern "C" void ScaLBL_D3Q19_AAeven_Greyscale_IMRT(double *dist, int start, int finish, int Np, double rlx,  double rlx_eff, double Gx, double Gy, double Gz,
+                                              double *Poros,double *Perm, double *Velocity, double Den,double *Pressure){
+	int n;
+	double vx,vy,vz,v_mag;
+    double ux,uy,uz,u_mag;
+    double pressure;//defined for this incompressible model
+	// conserved momemnts
+	double jx,jy,jz;
+	// non-conserved moments
+	double m1,m2,m4,m6,m8,m9,m10,m11,m12,m13,m14,m15,m16,m17,m18;
+    double fq;
+	//double f0,f1,f2,f3,f4,f5,f6,f7,f8,f9,f10,f11,f12,f13,f14,f15,f16,f17,f18;
+    double GeoFun;//geometric function from Guo's PRE 66, 036304 (2002)
+    double porosity;
+    double perm;//voxel permeability
+    double c0, c1; //Guo's model parameters
+    double mu_eff = (1.0/rlx_eff-0.5)/3.0;//kinematic viscosity
+    double Fx, Fy, Fz;//The total body force including Brinkman force and user-specified (Gx,Gy,Gz)
+    double rlx_setA = rlx;
+    double rlx_setB = 8.f*(2.f-rlx_setA)/(8.f-rlx_setA);
+
+	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 (int n=start; n<finish; n++){
+
+		//........................................................................
+		//					READ THE DISTRIBUTIONS
+		//		(read from opposite array due to previous swap operation)
+		//........................................................................
+		// q=0
+		fq = dist[n];
+		m1  = -30.0*fq;
+		m2  = 12.0*fq;
+
+		// q=1
+		fq = dist[2*Np+n];
+        pressure = 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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;
+        //---------------------------------------------------------------------//
+
+        porosity = Poros[n];
+        perm = Perm[n];
+
+        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/Den+0.5*porosity*Gx;
+		vy = jy/Den+0.5*porosity*Gy;
+		vz = jz/Den+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 = Den*(-porosity*mu_eff/perm*ux - porosity*GeoFun/sqrt(perm)*u_mag*ux + porosity*Gx);
+        Fy = Den*(-porosity*mu_eff/perm*uy - porosity*GeoFun/sqrt(perm)*u_mag*uy + porosity*Gy);
+        Fz = Den*(-porosity*mu_eff/perm*uz - porosity*GeoFun/sqrt(perm)*u_mag*uz + porosity*Gz);
+        if (porosity==1.0){
+            Fx=Den*Gx;
+            Fy=Den*Gy;
+            Fz=Den*Gz;
+        }
+
+        //Calculate pressure for Incompressible-MRT model
+        pressure=0.5/porosity*(pressure-0.5*Den*u_mag*u_mag/porosity);
+
+        //-------------------- IMRT collison where body force has higher-order terms -------------//
+//		//..............carry out relaxation process...............................................
+//		m1 = m1 + rlx_setA*((-30*Den+19*(ux*ux+uy*uy+uz*uz)/porosity + 57*pressure*porosity) - m1) 
+//                + (1-0.5*rlx_setA)*38*(Fx*ux+Fy*uy+Fz*uz)/porosity;
+//		m2 = m2 + rlx_setA*((12*Den - 5.5*(ux*ux+uy*uy+uz*uz)/porosity-27*pressure*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*Den) - m4)
+//                + (1-0.5*rlx_setB)*(-0.6666666666666666*Fx);
+//        jy = jy + Fy;
+//		m6 = m6 + rlx_setB*((-0.6666666666666666*uy*Den) - m6)
+//                + (1-0.5*rlx_setB)*(-0.6666666666666666*Fy);
+//        jz = jz + Fz;
+//		m8 = m8 + rlx_setB*((-0.6666666666666666*uz*Den) - m8)
+//                + (1-0.5*rlx_setB)*(-0.6666666666666666*Fz);
+//		m9 = m9 + rlx_setA*((Den*(2*ux*ux-uy*uy-uz*uz)/porosity) - m9)
+//                + (1-0.5*rlx_setA)*(4*Fx*ux-2*Fy*uy-2*Fz*uz)/porosity;
+//		m10 = m10 + rlx_setA*(-0.5*Den*((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*((Den*(uy*uy-uz*uz)/porosity) - m11)
+//                  + (1-0.5*rlx_setA)*(2*Fy*uy-2*Fz*uz)/porosity;
+//		m12 = m12 + rlx_setA*(-0.5*(Den*(uy*uy-uz*uz)/porosity)- m12)
+//                  + (1-0.5*rlx_setA)*(-Fy*uy+Fz*uz)/porosity;
+//		m13 = m13 + rlx_setA*((Den*ux*uy/porosity) - m13)
+//                  + (1-0.5*rlx_setA)*(Fy*ux+Fx*uy)/porosity;
+//		m14 = m14 + rlx_setA*((Den*uy*uz/porosity) - m14)
+//                  + (1-0.5*rlx_setA)*(Fz*uy+Fy*uz)/porosity;
+//		m15 = m15 + rlx_setA*((Den*ux*uz/porosity) - 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);
+//		//.......................................................................................................
+
+
+        //-------------------- IMRT collison where body force has NO higher-order terms -------------//
+		//..............carry out relaxation process...............................................
+		m1 = m1 + rlx_setA*((-30*Den+19*(ux*ux+uy*uy+uz*uz)/porosity + 57*pressure*porosity) - m1);
+		m2 = m2 + rlx_setA*((12*Den - 5.5*(ux*ux+uy*uy+uz*uz)/porosity-27*pressure*porosity) - m2);
+        jx = jx + Fx;
+		m4 = m4 + rlx_setB*((-0.6666666666666666*ux*Den) - m4)
+                + (1-0.5*rlx_setB)*(-0.6666666666666666*Fx);
+        jy = jy + Fy;
+		m6 = m6 + rlx_setB*((-0.6666666666666666*uy*Den) - m6)
+                + (1-0.5*rlx_setB)*(-0.6666666666666666*Fy);
+        jz = jz + Fz;
+		m8 = m8 + rlx_setB*((-0.6666666666666666*uz*Den) - m8)
+                + (1-0.5*rlx_setB)*(-0.6666666666666666*Fz);
+		m9 = m9 + rlx_setA*((Den*(2*ux*ux-uy*uy-uz*uz)/porosity) - m9);
+		m10 = m10 + rlx_setA*(-0.5*Den*((2*ux*ux-uy*uy-uz*uz)/porosity)- m10);
+		m11 = m11 + rlx_setA*((Den*(uy*uy-uz*uz)/porosity) - m11);
+		m12 = m12 + rlx_setA*(-0.5*(Den*(uy*uy-uz*uz)/porosity)- m12);
+		m13 = m13 + rlx_setA*((Den*ux*uy/porosity) - m13);
+		m14 = m14 + rlx_setA*((Den*uy*uz/porosity) - m14);
+		m15 = m15 + rlx_setA*((Den*ux*uz/porosity) - m15);
+		m16 = m16 + rlx_setB*( - m16);
+		m17 = m17 + rlx_setB*( - m17);
+		m18 = m18 + rlx_setB*( - m18);
+		//.......................................................................................................
+
+
+		//.................inverse transformation......................................................
+		// q=0
+		fq = mrt_V1*Den-mrt_V2*m1+mrt_V3*m2;
+		dist[n] = fq;
+
+		// q = 1
+		fq = mrt_V1*Den-mrt_V4*m1-mrt_V5*m2+0.1*(jx-m4)+mrt_V6*(m9-m10);
+		dist[1*Np+n] = fq;
+
+		// q=2
+		fq = mrt_V1*Den-mrt_V4*m1-mrt_V5*m2+0.1*(m4-jx)+mrt_V6*(m9-m10);
+		dist[2*Np+n] = fq;
+
+		// q = 3
+		fq = mrt_V1*Den-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*Den-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*Den-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*Den-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*Den+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*Den+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*Den+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*Den+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*Den+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*Den+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*Den+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*Den+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*Den+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*Den+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*Den+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*Den+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;
+		//........................................................................
+
+        //Update velocity on device
+		Velocity[0*Np+n] = ux;
+		Velocity[1*Np+n] = uy;
+		Velocity[2*Np+n] = uz;
+        //Update pressure on device
+        Pressure[n] = pressure;
+	}
+}
+
+extern "C" void ScaLBL_D3Q19_AAodd_Greyscale_IMRT(int *neighborList, double *dist, int start, int finish, int Np, double rlx,  double rlx_eff, double Gx, double Gy, double Gz, 
+                                             double *Poros,double *Perm, double *Velocity, double Den,double *Pressure){
+	int n, nread;
+	double vx,vy,vz,v_mag;
+    double ux,uy,uz,u_mag;
+    double pressure;//defined for this incompressible model
+	// conserved momemnts
+	double jx,jy,jz;
+	// non-conserved moments
+	double m1,m2,m4,m6,m8,m9,m10,m11,m12,m13,m14,m15,m16,m17,m18;
+    double fq;
+	//double f0,f1,f2,f3,f4,f5,f6,f7,f8,f9,f10,f11,f12,f13,f14,f15,f16,f17,f18;
+    double GeoFun;//geometric function from Guo's PRE 66, 036304 (2002)
+    double porosity;
+    double perm;//voxel permeability
+    double c0, c1; //Guo's model parameters
+    double mu_eff = (1.0/rlx_eff-0.5)/3.0;//kinematic viscosity
+    double Fx, Fy, Fz;//The total body force including Brinkman force and user-specified (Gx,Gy,Gz)
+    double rlx_setA = rlx;
+    double rlx_setB = 8.f*(2.f-rlx_setA)/(8.f-rlx_setA);
+
+	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 (int n=start; n<finish; n++){
+
+		//........................................................................
+		//					READ THE DISTRIBUTIONS
+		//		(read from opposite array due to previous swap operation)
+		//........................................................................
+		// q=0
+		fq = dist[n];
+		m1  = -30.0*fq;
+		m2  = 12.0*fq;
+
+		// q=1
+		nread = neighborList[n]; // neighbor 2 ( > 10Np => odd part of dist)
+		fq = dist[nread]; // reading the f1 data into register fq
+        pressure = fq;
+		m1 -= 11.0*fq;
+		m2 -= 4.0*fq;
+		jx = fq;
+		m4 = -4.0*fq;
+		m9 = 2.0*fq;
+		m10 = -4.0*fq;
+
+		// q=2
+		nread = neighborList[n+Np]; // neighbor 1 ( < 10Np => even part of dist)
+		fq = dist[nread];  // reading the f2 data into register fq
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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];
+		pressure += 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
+		nread = neighborList[n+16*Np];
+		fq = dist[nread];
+		pressure += 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];
+		pressure += 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;
+        //---------------------------------------------------------------------//
+
+        porosity = Poros[n];
+        perm = Perm[n];
+
+        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/Den+0.5*porosity*Gx;
+		vy = jy/Den+0.5*porosity*Gy;
+		vz = jz/Den+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 = Den*(-porosity*mu_eff/perm*ux - porosity*GeoFun/sqrt(perm)*u_mag*ux + porosity*Gx);
+        Fy = Den*(-porosity*mu_eff/perm*uy - porosity*GeoFun/sqrt(perm)*u_mag*uy + porosity*Gy);
+        Fz = Den*(-porosity*mu_eff/perm*uz - porosity*GeoFun/sqrt(perm)*u_mag*uz + porosity*Gz);
+        if (porosity==1.0){
+            Fx=Den*Gx;
+            Fy=Den*Gy;
+            Fz=Den*Gz;
+        }
+
+        //Calculate pressure for Incompressible-MRT model
+        pressure=0.5/porosity*(pressure-0.5*Den*u_mag*u_mag/porosity);
+
+        //-------------------- IMRT collison where body force has higher-order terms -------------//
+//		//..............carry out relaxation process...............................................
+//		m1 = m1 + rlx_setA*((-30*Den+19*(ux*ux+uy*uy+uz*uz)/porosity + 57*pressure*porosity) - m1) 
+//                + (1-0.5*rlx_setA)*38*(Fx*ux+Fy*uy+Fz*uz)/porosity;
+//		m2 = m2 + rlx_setA*((12*Den - 5.5*(ux*ux+uy*uy+uz*uz)/porosity-27*pressure*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*Den) - m4)
+//                + (1-0.5*rlx_setB)*(-0.6666666666666666*Fx);
+//        jy = jy + Fy;
+//		m6 = m6 + rlx_setB*((-0.6666666666666666*uy*Den) - m6)
+//                + (1-0.5*rlx_setB)*(-0.6666666666666666*Fy);
+//        jz = jz + Fz;
+//		m8 = m8 + rlx_setB*((-0.6666666666666666*uz*Den) - m8)
+//                + (1-0.5*rlx_setB)*(-0.6666666666666666*Fz);
+//		m9 = m9 + rlx_setA*((Den*(2*ux*ux-uy*uy-uz*uz)/porosity) - m9)
+//                + (1-0.5*rlx_setA)*(4*Fx*ux-2*Fy*uy-2*Fz*uz)/porosity;
+//		m10 = m10 + rlx_setA*(-0.5*Den*((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*((Den*(uy*uy-uz*uz)/porosity) - m11)
+//                  + (1-0.5*rlx_setA)*(2*Fy*uy-2*Fz*uz)/porosity;
+//		m12 = m12 + rlx_setA*(-0.5*(Den*(uy*uy-uz*uz)/porosity)- m12)
+//                  + (1-0.5*rlx_setA)*(-Fy*uy+Fz*uz)/porosity;
+//		m13 = m13 + rlx_setA*((Den*ux*uy/porosity) - m13)
+//                  + (1-0.5*rlx_setA)*(Fy*ux+Fx*uy)/porosity;
+//		m14 = m14 + rlx_setA*((Den*uy*uz/porosity) - m14)
+//                  + (1-0.5*rlx_setA)*(Fz*uy+Fy*uz)/porosity;
+//		m15 = m15 + rlx_setA*((Den*ux*uz/porosity) - 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);
+//		//.......................................................................................................
+       
+
+        //-------------------- IMRT collison where body force has NO higher-order terms -------------//
+		//..............carry out relaxation process...............................................
+		m1 = m1 + rlx_setA*((-30*Den+19*(ux*ux+uy*uy+uz*uz)/porosity + 57*pressure*porosity) - m1);
+		m2 = m2 + rlx_setA*((12*Den - 5.5*(ux*ux+uy*uy+uz*uz)/porosity-27*pressure*porosity) - m2);
+        jx = jx + Fx;
+		m4 = m4 + rlx_setB*((-0.6666666666666666*ux*Den) - m4)
+                + (1-0.5*rlx_setB)*(-0.6666666666666666*Fx);
+        jy = jy + Fy;
+		m6 = m6 + rlx_setB*((-0.6666666666666666*uy*Den) - m6)
+                + (1-0.5*rlx_setB)*(-0.6666666666666666*Fy);
+        jz = jz + Fz;
+		m8 = m8 + rlx_setB*((-0.6666666666666666*uz*Den) - m8)
+                + (1-0.5*rlx_setB)*(-0.6666666666666666*Fz);
+		m9 = m9 + rlx_setA*((Den*(2*ux*ux-uy*uy-uz*uz)/porosity) - m9);
+		m10 = m10 + rlx_setA*(-0.5*Den*((2*ux*ux-uy*uy-uz*uz)/porosity)- m10);
+		m11 = m11 + rlx_setA*((Den*(uy*uy-uz*uz)/porosity) - m11);
+		m12 = m12 + rlx_setA*(-0.5*(Den*(uy*uy-uz*uz)/porosity)- m12);
+		m13 = m13 + rlx_setA*((Den*ux*uy/porosity) - m13);
+		m14 = m14 + rlx_setA*((Den*uy*uz/porosity) - m14);
+		m15 = m15 + rlx_setA*((Den*ux*uz/porosity) - m15);
+		m16 = m16 + rlx_setB*( - m16);
+		m17 = m17 + rlx_setB*( - m17);
+		m18 = m18 + rlx_setB*( - m18);
+		//.......................................................................................................
+
+
+		//.................inverse transformation......................................................
+		// q=0
+		fq = mrt_V1*Den-mrt_V2*m1+mrt_V3*m2;
+		dist[n] = fq;
+
+		// q = 1
+		fq = mrt_V1*Den-mrt_V4*m1-mrt_V5*m2+0.1*(jx-m4)+mrt_V6*(m9-m10);
+		nread = neighborList[n+Np];
+		dist[nread] = fq;
+
+		// q=2
+		fq = mrt_V1*Den-mrt_V4*m1-mrt_V5*m2+0.1*(m4-jx)+mrt_V6*(m9-m10);
+		nread = neighborList[n];
+		dist[nread] = fq;
+
+		// q = 3
+		fq = mrt_V1*Den-mrt_V4*m1-mrt_V5*m2+0.1*(jy-m6)+mrt_V7*(m10-m9)+mrt_V8*(m11-m12);
+		nread = neighborList[n+3*Np];
+		dist[nread] = fq;
+
+		// q = 4
+		fq = mrt_V1*Den-mrt_V4*m1-mrt_V5*m2+0.1*(m6-jy)+mrt_V7*(m10-m9)+mrt_V8*(m11-m12);
+		nread = neighborList[n+2*Np];
+		dist[nread] = fq;
+
+		// q = 5
+		fq = mrt_V1*Den-mrt_V4*m1-mrt_V5*m2+0.1*(jz-m8)+mrt_V7*(m10-m9)+mrt_V8*(m12-m11);
+		nread = neighborList[n+5*Np];
+		dist[nread] = fq;
+
+		// q = 6
+		fq = mrt_V1*Den-mrt_V4*m1-mrt_V5*m2+0.1*(m8-jz)+mrt_V7*(m10-m9)+mrt_V8*(m12-m11);
+		nread = neighborList[n+4*Np];
+		dist[nread] = fq;
+
+		// q = 7
+		fq = mrt_V1*Den+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[nread] = fq;
+
+		// q = 8
+		fq = mrt_V1*Den+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[nread] = fq;
+
+		// q = 9
+		fq = mrt_V1*Den+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[nread] = fq;
+
+		// q = 10
+		fq = mrt_V1*Den+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[nread] = fq;
+
+		// q = 11
+		fq = mrt_V1*Den+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[nread] = fq;
+
+		// q = 12
+		fq = mrt_V1*Den+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[nread]= fq;
+
+		// q = 13
+		fq = mrt_V1*Den+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[nread] = fq;
+
+		// q= 14
+		fq = mrt_V1*Den+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[nread] = fq;
+
+		// q = 15
+		fq = mrt_V1*Den+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*Den+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*Den+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*Den+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;
+		//........................................................................
+
+        //Update velocity on device
+		Velocity[0*Np+n] = ux;
+		Velocity[1*Np+n] = uy;
+		Velocity[2*Np+n] = uz;
+        //Update pressure on device
+        Pressure[n] = pressure;
+	}
+}
+
diff --git a/gpu/D3Q19.cu b/gpu/D3Q19.cu
index 133c8455..ccd125b2 100644
--- a/gpu/D3Q19.cu
+++ b/gpu/D3Q19.cu
@@ -267,6 +267,36 @@ __global__ void dvc_ScaLBL_D3Q19_Init(double *dist, int Np)
 	}
 }
 
+__global__ void dvc_ScaLBL_D3Q19_GreyIMRT_Init(double *dist, int Np, double Den)
+{
+	int n;
+	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;
+		if (n<Np ){
+			dist[n] = Den - 0.6666666666666667;
+			dist[Np+n] = 0.055555555555555555;		//double(100*n)+1.f;
+			dist[2*Np+n] = 0.055555555555555555;	//double(100*n)+2.f;
+			dist[3*Np+n] = 0.055555555555555555;	//double(100*n)+3.f;
+			dist[4*Np+n] = 0.055555555555555555;	//double(100*n)+4.f;
+			dist[5*Np+n] = 0.055555555555555555;	//double(100*n)+5.f;
+			dist[6*Np+n] = 0.055555555555555555;	//double(100*n)+6.f;
+			dist[7*Np+n] = 0.0277777777777778;   //double(100*n)+7.f;
+			dist[8*Np+n] = 0.0277777777777778;   //double(100*n)+8.f;
+			dist[9*Np+n] = 0.0277777777777778;   //double(100*n)+9.f;
+			dist[10*Np+n] = 0.0277777777777778;  //double(100*n)+10.f;
+			dist[11*Np+n] = 0.0277777777777778;  //double(100*n)+11.f;
+			dist[12*Np+n] = 0.0277777777777778;  //double(100*n)+12.f;
+			dist[13*Np+n] = 0.0277777777777778;  //double(100*n)+13.f;
+			dist[14*Np+n] = 0.0277777777777778;  //double(100*n)+14.f;
+			dist[15*Np+n] = 0.0277777777777778;  //double(100*n)+15.f;
+			dist[16*Np+n] = 0.0277777777777778;  //double(100*n)+16.f;
+			dist[17*Np+n] = 0.0277777777777778;  //double(100*n)+17.f;
+			dist[18*Np+n] = 0.0277777777777778;  //double(100*n)+18.f;
+		}
+	}
+}
 
 //*************************************************************************
 __global__  void dvc_ScaLBL_D3Q19_Swap_Compact(int *neighborList, double *disteven, double *distodd, int Np, int q){
@@ -2324,7 +2354,15 @@ extern "C" void ScaLBL_D3Q19_Init(double *dist, int Np){
 	dvc_ScaLBL_D3Q19_Init<<<NBLOCKS,NTHREADS >>>(dist, Np);
 	cudaError_t err = cudaGetLastError();
 	if (cudaSuccess != err){
-		printf("CUDA error in ScaLBL_D3Q19_AA_Init: %s \n",cudaGetErrorString(err));
+		printf("CUDA error in ScaLBL_D3Q19_Init: %s \n",cudaGetErrorString(err));
+	}
+}
+
+extern "C" void ScaLBL_D3Q19_GreyIMRT_Init(double *dist, int Np, double Den){
+	dvc_ScaLBL_D3Q19_GreyIMRT_Init<<<NBLOCKS,NTHREADS >>>(dist, Np, Den);
+	cudaError_t err = cudaGetLastError();
+	if (cudaSuccess != err){
+		printf("CUDA error in ScaLBL_D3Q19_GreyIMRT_Init: %s \n",cudaGetErrorString(err));
 	}
 }
 
diff --git a/gpu/Greyscale.cu b/gpu/Greyscale.cu
new file mode 100644
index 00000000..0a9a63e0
--- /dev/null
+++ b/gpu/Greyscale.cu
@@ -0,0 +1,1633 @@
+#include <stdio.h>
+
+#define NBLOCKS 1024
+#define NTHREADS 256
+
+__global__ void dvc_ScaLBL_D3Q19_AAeven_Greyscale(double *dist, int start, int finish, int Np, double rlx,  double rlx_eff, double Gx, double Gy, double Gz,
+                                                  double *Poros,double *Perm, double *Velocity, double *Pressure){
+	int n;
+	// conserved momemnts
+	double rho,vx,vy,vz,v_mag;
+    double ux,uy,uz,u_mag;
+    double pressure;
+    //double uu;
+	// non-conserved moments
+	double f0,f1,f2,f3,f4,f5,f6,f7,f8,f9,f10,f11,f12,f13,f14,f15,f16,f17,f18;
+    double GeoFun;//geometric function from Guo's PRE 66, 036304 (2002)
+    double porosity;
+    double perm;//voxel permeability
+    double c0, c1; //Guo's model parameters
+    double mu_eff = (1.0/rlx_eff-0.5)/3.0;//kinematic viscosity
+    double Fx, Fy, Fz;//The total body force including Brinkman force and user-specified (Gx,Gy,Gz)
+
+	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 ){
+		// q=0
+		f0 = dist[n];
+		f1 = dist[2*Np+n];
+		f2 = dist[1*Np+n];
+		f3 = dist[4*Np+n];
+		f4 = dist[3*Np+n];
+		f5 = dist[6*Np+n];
+		f6 = dist[5*Np+n];
+		f7 = dist[8*Np+n];
+		f8 = dist[7*Np+n];
+		f9 = dist[10*Np+n];
+		f10 = dist[9*Np+n];
+		f11 = dist[12*Np+n];
+		f12 = dist[11*Np+n];
+		f13 = dist[14*Np+n];
+		f14 = dist[13*Np+n];
+		f15 = dist[16*Np+n];
+		f16 = dist[15*Np+n];
+		f17 = dist[18*Np+n];
+		f18 = dist[17*Np+n];
+
+        porosity = Poros[n];
+        perm = Perm[n];
+
+        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
+
+		rho = f0+f2+f1+f4+f3+f6+f5+f8+f7+f10+f9+f12+f11+f14+f13+f16+f15+f18+f17;
+        pressure = rho/porosity/3.0;
+		vx = (f1-f2+f7-f8+f9-f10+f11-f12+f13-f14)/rho+0.5*porosity*Gx;
+		vy = (f3-f4+f7-f8-f9+f10+f15-f16+f17-f18)/rho+0.5*porosity*Gy;
+		vz = (f5-f6+f11-f12-f13+f14+f15-f16-f17+f18)/rho+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 = -porosity*mu_eff/perm*ux - porosity*GeoFun/sqrt(perm)*u_mag*ux + porosity*Gx;
+        Fy = -porosity*mu_eff/perm*uy - porosity*GeoFun/sqrt(perm)*u_mag*uy + porosity*Gy;
+        Fz = -porosity*mu_eff/perm*uz - porosity*GeoFun/sqrt(perm)*u_mag*uz + porosity*Gz;
+        if (porosity==1.0){
+            Fx=Gx;
+            Fy=Gy;
+            Fz=Gz;
+        }
+
+        //------------------------ BGK collison where body force has higher-order terms ----------------------------------------------------------//
+//		// q=0
+//		dist[n] = f0*(1.0-rlx)+ rlx*0.3333333333333333*rho*(1. - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//                  + 0.3333333333333333*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(0. - (3.*uy)/porosity) + Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 1
+//		dist[1*Np+n] = f1*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*ux + (4.5*ux*ux)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//            +0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(3. + (6.*ux)/porosity) + Fy*(0. - (3.*uy)/porosity) + Fz*(0. - (3.*uz)/porosity));
+//
+//		// q=2
+//		dist[2*Np+n] = f2*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*ux + (4.5*ux*ux)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//            +0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(-3. + (6.*ux)/porosity) + Fy*(0. - (3.*uy)/porosity) + Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 3
+//		dist[3*Np+n] = f3*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*uy + (4.5*uy*uy)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(3. + (6.*uy)/porosity) + Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 4
+//		dist[4*Np+n] = f4*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*uy + (4.5*uy*uy)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(-3. + (6.*uy)/porosity) + Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 5
+//		dist[5*Np+n] = f5*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*uz + (4.5*uz*uz)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(0. - (3.*uy)/porosity) + Fz*(3. + (6.*uz)/porosity));
+//
+//		// q = 6
+//		dist[6*Np+n] = f6*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*uz + (4.5*uz*uz)/porosity - (1.5*(ux*ux+ uy*uy + uz*uz))/porosity)
+//				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(0. - (3.*uy)/porosity) + Fz*(-3. + (6.*uz)/porosity));
+//
+//		// q = 7
+//		dist[7*Np+n] = f7*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux + uy) + (4.5*(ux + uy)*(ux + uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3. - (3.*ux)/porosity + (9.*(ux + uy))/porosity) + Fy*(3. - (3.*uy)/porosity + (9.*(ux + uy))/porosity) + 
+//  Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 8
+//		dist[8*Np+n] = f8*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux - uy) + (4.5*(-ux - uy)*(-ux - uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3. - (3.*ux)/porosity - (9.*(-ux - uy))/porosity) + Fy*(-3. - (9.*(-ux - uy))/porosity - (3.*uy)/porosity) + 
+//  Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 9
+//		dist[9*Np+n] = f9*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux - uy) + (4.5*(ux - uy)*(ux - uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3. - (3.*ux)/porosity + (9.*(ux - uy))/porosity) + Fy*(-3. - (9.*(ux - uy))/porosity - (3.*uy)/porosity) + 
+//  Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 10
+//		dist[10*Np+n] = f10*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux + uy) + (4.5*(-ux + uy)*(-ux + uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3. - (3.*ux)/porosity - (9.*(-ux + uy))/porosity) + Fy*(3. - (3.*uy)/porosity + (9.*(-ux + uy))/porosity) + 
+//  Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 11
+//		dist[11*Np+n] = f11*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux + uz) + (4.5*(ux + uz)*(ux + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(0. - (3.*uy)/porosity) + Fx*(3. - (3.*ux)/porosity + (9.*(ux + uz))/porosity) + 
+//  Fz*(3. - (3.*uz)/porosity + (9.*(ux + uz))/porosity));
+//
+//		// q = 12
+//		dist[12*Np+n] = f12*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux - uz) + (4.5*(-ux - uz)*(-ux - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(0. - (3.*uy)/porosity) + Fx*(-3. - (3.*ux)/porosity - (9.*(-ux - uz))/porosity) + 
+//  Fz*(-3. - (9.*(-ux - uz))/porosity - (3.*uz)/porosity));
+//
+//		// q = 13
+//		dist[13*Np+n] = f13*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux - uz) + (4.5*(ux - uz)*(ux - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(0. - (3.*uy)/porosity) + Fx*(3. - (3.*ux)/porosity + (9.*(ux - uz))/porosity) + 
+//  Fz*(-3. - (9.*(ux - uz))/porosity - (3.*uz)/porosity));
+//
+//		// q= 14
+//		dist[14*Np+n] = f14*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux + uz) + (4.5*(-ux + uz)*(-ux + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(0. - (3.*uy)/porosity) + Fx*(-3. - (3.*ux)/porosity - (9.*(-ux + uz))/porosity) + 
+//  Fz*(3. - (3.*uz)/porosity + (9.*(-ux + uz))/porosity));
+//
+//		// q = 15
+//		dist[15*Np+n] = f15*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(uy + uz) + (4.5*(uy + uz)*(uy + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(3. - (3.*uy)/porosity + (9.*(uy + uz))/porosity) + 
+//  Fz*(3. - (3.*uz)/porosity + (9.*(uy + uz))/porosity));
+//
+//		// q = 16
+//		dist[16*Np+n] = f16*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-uy - uz) + (4.5*(-uy - uz)*(-uy - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(-3. - (3.*uy)/porosity - (9.*(-uy - uz))/porosity) + 
+//  Fz*(-3. - (9.*(-uy - uz))/porosity - (3.*uz)/porosity));
+//
+//		// q = 17
+//		dist[17*Np+n] = f17*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(uy - uz) + (4.5*(uy - uz)*(uy - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(3. - (3.*uy)/porosity + (9.*(uy - uz))/porosity) + 
+//  Fz*(-3. - (9.*(uy - uz))/porosity - (3.*uz)/porosity));
+//
+//		// q = 18
+//		dist[18*Np+n] = f18*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-uy + uz) + (4.5*(-uy + uz)*(-uy + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(-3. - (3.*uy)/porosity - (9.*(-uy + uz))/porosity) + 
+//  Fz*(3. - (3.*uz)/porosity + (9.*(-uy + uz))/porosity));
+        //----------------------------------------------------------------------------------------------------------------------------------------//
+
+
+        //------------------------ BGK collison where body force has NO higher-order terms ----------------------------------------------------------//
+		// q=0
+		dist[n] = f0*(1.0-rlx)+ rlx*0.3333333333333333*rho*(1. - (1.5*(ux*ux + uy*uy + uz*uz))/porosity);
+
+		// q = 1
+		dist[1*Np+n] = f1*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*ux + (4.5*ux*ux)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+            +0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(3.));
+
+		// q=2
+		dist[2*Np+n] = f2*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*ux + (4.5*ux*ux)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+            +0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(-3.));
+
+		// q = 3
+		dist[3*Np+n] = f3*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*uy + (4.5*uy*uy)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fy*(3.));
+
+		// q = 4
+		dist[4*Np+n] = f4*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*uy + (4.5*uy*uy)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fy*(-3.));
+
+		// q = 5
+		dist[5*Np+n] = f5*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*uz + (4.5*uz*uz)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fz*(3.));
+
+		// q = 6
+		dist[6*Np+n] = f6*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*uz + (4.5*uz*uz)/porosity - (1.5*(ux*ux+ uy*uy + uz*uz))/porosity)
+				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fz*(-3.));
+
+		// q = 7
+		dist[7*Np+n] = f7*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux + uy) + (4.5*(ux + uy)*(ux + uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3.) + Fy*(3.));
+
+		// q = 8
+		dist[8*Np+n] = f8*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux - uy) + (4.5*(-ux - uy)*(-ux - uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3.) + Fy*(-3.));
+
+		// q = 9
+		dist[9*Np+n] = f9*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux - uy) + (4.5*(ux - uy)*(ux - uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3.) + Fy*(-3.));
+
+		// q = 10
+		dist[10*Np+n] = f10*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux + uy) + (4.5*(-ux + uy)*(-ux + uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3.) + Fy*(3.));
+
+		// q = 11
+		dist[11*Np+n] = f11*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux + uz) + (4.5*(ux + uz)*(ux + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3.) + Fz*(3.));
+
+		// q = 12
+		dist[12*Np+n] = f12*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux - uz) + (4.5*(-ux - uz)*(-ux - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3.) + Fz*(-3.));
+
+		// q = 13
+		dist[13*Np+n] = f13*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux - uz) + (4.5*(ux - uz)*(ux - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3.) + Fz*(-3.));
+
+		// q= 14
+		dist[14*Np+n] = f14*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux + uz) + (4.5*(-ux + uz)*(-ux + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3.) + Fz*(3.));
+
+		// q = 15
+		dist[15*Np+n] = f15*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(uy + uz) + (4.5*(uy + uz)*(uy + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(3.) + Fz*(3.));
+
+		// q = 16
+		dist[16*Np+n] = f16*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-uy - uz) + (4.5*(-uy - uz)*(-uy - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(-3.) + Fz*(-3.));
+
+		// q = 17
+		dist[17*Np+n] = f17*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(uy - uz) + (4.5*(uy - uz)*(uy - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(3.) + Fz*(-3.));
+
+		// q = 18
+		dist[18*Np+n] = f18*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-uy + uz) + (4.5*(-uy + uz)*(-uy + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(-3.) + Fz*(3.));
+        //-------------------------------------------------------------------------------------------------------------------------------------------//
+
+        //Update velocity on device
+		Velocity[0*Np+n] = ux;
+		Velocity[1*Np+n] = uy;
+		Velocity[2*Np+n] = uz;
+        //Update pressure on device
+        Pressure[n] = pressure;
+
+		}
+	}
+}
+
+__global__ void dvc_ScaLBL_D3Q19_AAodd_Greyscale(int *neighborList, double *dist, int start, int finish, int Np, double rlx,  double rlx_eff, double Gx, double Gy, double Gz,
+                                                 double *Poros,double *Perm, double *Velocity, double *Pressure){
+	int n;
+	// conserved momemnts
+	double rho,vx,vy,vz,v_mag;
+    double ux,uy,uz,u_mag;
+    double pressure;
+    //double uu;
+	// non-conserved moments
+	double f0,f1,f2,f3,f4,f5,f6,f7,f8,f9,f10,f11,f12,f13,f14,f15,f16,f17,f18;
+	int nr1,nr2,nr3,nr4,nr5,nr6,nr7,nr8,nr9,nr10,nr11,nr12,nr13,nr14,nr15,nr16,nr17,nr18;
+    double GeoFun;//geometric function from Guo's PRE 66, 036304 (2002)
+    double porosity;
+    double perm;//voxel permeability
+    double c0, c1; //Guo's model parameters
+    double mu_eff = (1.0/rlx_eff-0.5)/3.0;//kinematic viscosity
+    double Fx, Fy, Fz;//The total body force including Brinkman force and user-specified (Gx,Gy,Gz)
+
+	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 ){		
+		// q=0
+		f0 = dist[n];
+		// q=1
+		nr1 = neighborList[n]; // neighbor 2 ( > 10Np => odd part of dist)
+		f1 = dist[nr1]; // reading the f1 data into register fq
+
+		nr2 = neighborList[n+Np]; // neighbor 1 ( < 10Np => even part of dist)
+		f2 = dist[nr2];  // reading the f2 data into register fq
+
+		// q=3
+		nr3 = neighborList[n+2*Np]; // neighbor 4
+		f3 = dist[nr3];
+
+		// q = 4
+		nr4 = neighborList[n+3*Np]; // neighbor 3
+		f4 = dist[nr4];
+
+		// q=5
+		nr5 = neighborList[n+4*Np];
+		f5 = dist[nr5];
+
+		// q = 6
+		nr6 = neighborList[n+5*Np];
+		f6 = dist[nr6];
+		
+		// q=7
+		nr7 = neighborList[n+6*Np];
+		f7 = dist[nr7];
+
+		// q = 8
+		nr8 = neighborList[n+7*Np];
+		f8 = dist[nr8];
+
+		// q=9
+		nr9 = neighborList[n+8*Np];
+		f9 = dist[nr9];
+
+		// q = 10
+		nr10 = neighborList[n+9*Np];
+		f10 = dist[nr10];
+
+		// q=11
+		nr11 = neighborList[n+10*Np];
+		f11 = dist[nr11];
+
+		// q=12
+		nr12 = neighborList[n+11*Np];
+		f12 = dist[nr12];
+
+		// q=13
+		nr13 = neighborList[n+12*Np];
+		f13 = dist[nr13];
+
+		// q=14
+		nr14 = neighborList[n+13*Np];
+		f14 = dist[nr14];
+
+		// q=15
+		nr15 = neighborList[n+14*Np];
+		f15 = dist[nr15];
+
+		// q=16
+		nr16 = neighborList[n+15*Np];
+		f16 = dist[nr16];
+
+		// q=17
+		//fq = dist[18*Np+n];
+		nr17 = neighborList[n+16*Np];
+		f17 = dist[nr17];
+
+		// q=18
+		nr18 = neighborList[n+17*Np];
+		f18 = dist[nr18];
+
+        porosity = Poros[n];
+        perm = Perm[n];
+
+        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
+
+		rho = f0+f2+f1+f4+f3+f6+f5+f8+f7+f10+f9+f12+f11+f14+f13+f16+f15+f18+f17;
+        pressure = rho/porosity/3.0;
+		vx = (f1-f2+f7-f8+f9-f10+f11-f12+f13-f14)/rho+0.5*porosity*Gx;
+		vy = (f3-f4+f7-f8-f9+f10+f15-f16+f17-f18)/rho+0.5*porosity*Gy;
+		vz = (f5-f6+f11-f12-f13+f14+f15-f16-f17+f18)/rho+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 body force to include linear (Darcy) and nonlinear (Forchheimer) drags due to the porous medium
+        Fx = -porosity*mu_eff/perm*ux - porosity*GeoFun/sqrt(perm)*u_mag*ux + porosity*Gx;
+        Fy = -porosity*mu_eff/perm*uy - porosity*GeoFun/sqrt(perm)*u_mag*uy + porosity*Gy;
+        Fz = -porosity*mu_eff/perm*uz - porosity*GeoFun/sqrt(perm)*u_mag*uz + porosity*Gz;
+        if (porosity==1.0){
+            Fx=Gx;
+            Fy=Gy;
+            Fz=Gz;
+        }
+
+        //------------------------ BGK collison where body force has higher-order terms ----------------------------------------------------------//
+//		// q=0
+//		dist[n] = f0*(1.0-rlx) + rlx*0.3333333333333333*rho*(1. - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//                  + 0.3333333333333333*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(0. - (3.*uy)/porosity) + Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 1
+//		dist[nr2] = f1*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*ux + (4.5*ux*ux)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//            +0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(3. + (6.*ux)/porosity) + Fy*(0. - (3.*uy)/porosity) + Fz*(0. - (3.*uz)/porosity));
+//
+//		// q=2
+//		dist[nr1] = f2*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*ux + (4.5*ux*ux)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//            +0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(-3. + (6.*ux)/porosity) + Fy*(0. - (3.*uy)/porosity) + Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 3
+//		dist[nr4] = f3*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*uy + (4.5*uy*uy)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(3. + (6.*uy)/porosity) + Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 4
+//		dist[nr3] = f4*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*uy + (4.5*uy*uy)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)  
+//				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(-3. + (6.*uy)/porosity) + Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 5
+//		dist[nr6] = f5*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*uz + (4.5*uz*uz)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(0. - (3.*uy)/porosity) + Fz*(3. + (6.*uz)/porosity));
+//
+//		// q = 6
+//		dist[nr5] = f6*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*uz + (4.5*uz*uz)/porosity - (1.5*(ux*ux+ uy*uy + uz*uz))/porosity) 
+//				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(0. - (3.*uy)/porosity) + Fz*(-3. + (6.*uz)/porosity));
+//
+//		// q = 7
+//		dist[nr8] = f7*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux + uy) + (4.5*(ux + uy)*(ux + uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3. - (3.*ux)/porosity + (9.*(ux + uy))/porosity) + Fy*(3. - (3.*uy)/porosity + (9.*(ux + uy))/porosity) + 
+//  Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 8
+//		dist[nr7] = f8*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux - uy) + (4.5*(-ux - uy)*(-ux - uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3. - (3.*ux)/porosity - (9.*(-ux - uy))/porosity) + Fy*(-3. - (9.*(-ux - uy))/porosity - (3.*uy)/porosity) + 
+//  Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 9
+//		dist[nr10] = f9*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux - uy) + (4.5*(ux - uy)*(ux - uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3. - (3.*ux)/porosity + (9.*(ux - uy))/porosity) + Fy*(-3. - (9.*(ux - uy))/porosity - (3.*uy)/porosity) + 
+//  Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 10
+//		dist[nr9] = f10*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux + uy) + (4.5*(-ux + uy)*(-ux + uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3. - (3.*ux)/porosity - (9.*(-ux + uy))/porosity) + Fy*(3. - (3.*uy)/porosity + (9.*(-ux + uy))/porosity) + 
+//  Fz*(0. - (3.*uz)/porosity));
+//
+//		// q = 11
+//		dist[nr12] = f11*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux + uz) + (4.5*(ux + uz)*(ux + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(0. - (3.*uy)/porosity) + Fx*(3. - (3.*ux)/porosity + (9.*(ux + uz))/porosity) + 
+//  Fz*(3. - (3.*uz)/porosity + (9.*(ux + uz))/porosity));
+//
+//		// q = 12
+//		dist[nr11] = f12*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux - uz) + (4.5*(-ux - uz)*(-ux - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(0. - (3.*uy)/porosity) + Fx*(-3. - (3.*ux)/porosity - (9.*(-ux - uz))/porosity) + 
+//  Fz*(-3. - (9.*(-ux - uz))/porosity - (3.*uz)/porosity));
+//
+//		// q = 13
+//		dist[nr14] = f13*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux - uz) + (4.5*(ux - uz)*(ux - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(0. - (3.*uy)/porosity) + Fx*(3. - (3.*ux)/porosity + (9.*(ux - uz))/porosity) + 
+//  Fz*(-3. - (9.*(ux - uz))/porosity - (3.*uz)/porosity));
+//
+//		// q= 14
+//		dist[nr13] = f14*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux + uz) + (4.5*(-ux + uz)*(-ux + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(0. - (3.*uy)/porosity) + Fx*(-3. - (3.*ux)/porosity - (9.*(-ux + uz))/porosity) + 
+//  Fz*(3. - (3.*uz)/porosity + (9.*(-ux + uz))/porosity));
+//
+//		// q = 15
+//		dist[nr16] = f15*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(uy + uz) + (4.5*(uy + uz)*(uy + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(3. - (3.*uy)/porosity + (9.*(uy + uz))/porosity) + 
+//  Fz*(3. - (3.*uz)/porosity + (9.*(uy + uz))/porosity));
+//
+//		// q = 16
+//		dist[nr15] = f16*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-uy - uz) + (4.5*(-uy - uz)*(-uy - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(-3. - (3.*uy)/porosity - (9.*(-uy - uz))/porosity) + 
+//  Fz*(-3. - (9.*(-uy - uz))/porosity - (3.*uz)/porosity));
+//
+//		// q = 17
+//		dist[nr18] = f17*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(uy - uz) + (4.5*(uy - uz)*(uy - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(3. - (3.*uy)/porosity + (9.*(uy - uz))/porosity) + 
+//  Fz*(-3. - (9.*(uy - uz))/porosity - (3.*uz)/porosity));
+//
+//		// q = 18
+//		dist[nr17] = f18*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-uy + uz) + (4.5*(-uy + uz)*(-uy + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+//				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(0. - (3.*ux)/porosity) + Fy*(-3. - (3.*uy)/porosity - (9.*(-uy + uz))/porosity) + 
+//  Fz*(3. - (3.*uz)/porosity + (9.*(-uy + uz))/porosity));
+        //----------------------------------------------------------------------------------------------------------------------------------------//
+
+
+        //------------------------ BGK collison where body force has NO higher-order terms ----------------------------------------------------------//
+		// q=0
+		dist[n] = f0*(1.0-rlx) + rlx*0.3333333333333333*rho*(1. - (1.5*(ux*ux + uy*uy + uz*uz))/porosity);
+
+		// q = 1
+		dist[nr2] = f1*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*ux + (4.5*ux*ux)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+            +0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(3.));
+
+		// q=2
+		dist[nr1] = f2*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*ux + (4.5*ux*ux)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+            +0.05555555555555555*rho*(1. - 0.5*rlx)*(Fx*(-3.));
+
+		// q = 3
+		dist[nr4] = f3*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*uy + (4.5*uy*uy)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fy*(3.));
+
+		// q = 4
+		dist[nr3] = f4*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*uy + (4.5*uy*uy)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)  
+				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fy*(-3.));
+
+		// q = 5
+		dist[nr6] = f5*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 + 3.*uz + (4.5*uz*uz)/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fz*(3.));
+
+		// q = 6
+		dist[nr5] = f6*(1.0-rlx) + rlx*0.05555555555555555*rho*(1 - 3.*uz + (4.5*uz*uz)/porosity - (1.5*(ux*ux+ uy*uy + uz*uz))/porosity) 
+				+0.05555555555555555*rho*(1. - 0.5*rlx)*(Fz*(-3.));
+
+		// q = 7
+		dist[nr8] = f7*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux + uy) + (4.5*(ux + uy)*(ux + uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3.) + Fy*(3.));
+
+		// q = 8
+		dist[nr7] = f8*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux - uy) + (4.5*(-ux - uy)*(-ux - uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3.) + Fy*(-3.));
+
+		// q = 9
+		dist[nr10] = f9*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux - uy) + (4.5*(ux - uy)*(ux - uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3.) + Fy*(-3.));
+
+		// q = 10
+		dist[nr9] = f10*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux + uy) + (4.5*(-ux + uy)*(-ux + uy))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3.) + Fy*(3.));
+
+		// q = 11
+		dist[nr12] = f11*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux + uz) + (4.5*(ux + uz)*(ux + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3.) + Fz*(3.));
+
+		// q = 12
+		dist[nr11] = f12*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux - uz) + (4.5*(-ux - uz)*(-ux - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3.) + Fz*(-3.));
+
+		// q = 13
+		dist[nr14] = f13*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(ux - uz) + (4.5*(ux - uz)*(ux - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(3.) + Fz*(-3.));
+
+		// q= 14
+		dist[nr13] = f14*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-ux + uz) + (4.5*(-ux + uz)*(-ux + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fx*(-3.) + Fz*(3.));
+
+		// q = 15
+		dist[nr16] = f15*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(uy + uz) + (4.5*(uy + uz)*(uy + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(3.) + Fz*(3.));
+
+		// q = 16
+		dist[nr15] = f16*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-uy - uz) + (4.5*(-uy - uz)*(-uy - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(-3.) + Fz*(-3.));
+
+		// q = 17
+		dist[nr18] = f17*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(uy - uz) + (4.5*(uy - uz)*(uy - uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity) 
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(3.) + Fz*(-3.));
+
+		// q = 18
+		dist[nr17] = f18*(1.0-rlx) + rlx*0.027777777777777776*rho*(1 + 3.*(-uy + uz) + (4.5*(-uy + uz)*(-uy + uz))/porosity - (1.5*(ux*ux + uy*uy + uz*uz))/porosity)
+				+0.027777777777777776*rho*(1. - 0.5*rlx)*(Fy*(-3.) + Fz*(3.));
+        //-------------------------------------------------------------------------------------------------------------------------------------------//
+
+
+        //Update velocity on device
+		Velocity[0*Np+n] = ux;
+		Velocity[1*Np+n] = uy;
+		Velocity[2*Np+n] = uz;
+        //Update pressure on device
+        Pressure[n] = pressure;
+		}
+	}
+}
+
+__global__ void dvc_ScaLBL_D3Q19_AAeven_Greyscale_IMRT(double *dist, int start, int finish, int Np, double rlx,  double rlx_eff, double Gx, double Gy, double Gz,
+                                                  double *Poros,double *Perm, double *Velocity, double Den, double *Pressure){
+
+	int n;
+	double vx,vy,vz,v_mag;
+    double ux,uy,uz,u_mag;
+    double pressure;//defined for this incompressible model
+	// conserved momemnts
+	double jx,jy,jz;
+	// non-conserved moments
+	double m1,m2,m4,m6,m8,m9,m10,m11,m12,m13,m14,m15,m16,m17,m18;
+    double fq;
+	//double f0,f1,f2,f3,f4,f5,f6,f7,f8,f9,f10,f11,f12,f13,f14,f15,f16,f17,f18;
+    double GeoFun;//geometric function from Guo's PRE 66, 036304 (2002)
+    double porosity;
+    double perm;//voxel permeability
+    double c0, c1; //Guo's model parameters
+    double mu_eff = (1.0/rlx_eff-0.5)/3.0;//kinematic viscosity
+    double Fx, Fy, Fz;//The total body force including Brinkman force and user-specified (Gx,Gy,Gz)
+    double rlx_setA = rlx;
+    double rlx_setB = 8.f*(2.f-rlx_setA)/(8.f-rlx_setA);
+
+	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 DISTRIBUTIONS
+            //		(read from opposite array due to previous swap operation)
+            //........................................................................
+            // q=0
+            fq = dist[n];
+            m1  = -30.0*fq;
+            m2  = 12.0*fq;
+
+            // q=1
+            fq = dist[2*Np+n];
+            pressure = 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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;
+            //---------------------------------------------------------------------//
+
+            porosity = Poros[n];
+            perm = Perm[n];
+
+            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/Den+0.5*porosity*Gx;
+            vy = jy/Den+0.5*porosity*Gy;
+            vz = jz/Den+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 = Den*(-porosity*mu_eff/perm*ux - porosity*GeoFun/sqrt(perm)*u_mag*ux + porosity*Gx);
+            Fy = Den*(-porosity*mu_eff/perm*uy - porosity*GeoFun/sqrt(perm)*u_mag*uy + porosity*Gy);
+            Fz = Den*(-porosity*mu_eff/perm*uz - porosity*GeoFun/sqrt(perm)*u_mag*uz + porosity*Gz);
+            if (porosity==1.0){
+                Fx=Den*Gx;
+                Fy=Den*Gy;
+                Fz=Den*Gz;
+            }
+
+            //Calculate pressure for Incompressible-MRT model
+            pressure=0.5/porosity*(pressure-0.5*Den*u_mag*u_mag/porosity);
+
+//            //..............carry out relaxation process...............................................
+//            m1 = m1 + rlx_setA*((-30*Den+19*(ux*ux+uy*uy+uz*uz)/porosity + 57*pressure*porosity) - m1) 
+//                    + (1-0.5*rlx_setA)*38*(Fx*ux+Fy*uy+Fz*uz)/porosity;
+//            m2 = m2 + rlx_setA*((12*Den - 5.5*(ux*ux+uy*uy+uz*uz)/porosity-27*pressure*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*Den) - m4)
+//                    + (1-0.5*rlx_setB)*(-0.6666666666666666*Fx);
+//            jy = jy + Fy;
+//            m6 = m6 + rlx_setB*((-0.6666666666666666*uy*Den) - m6)
+//                    + (1-0.5*rlx_setB)*(-0.6666666666666666*Fy);
+//            jz = jz + Fz;
+//            m8 = m8 + rlx_setB*((-0.6666666666666666*uz*Den) - m8)
+//                    + (1-0.5*rlx_setB)*(-0.6666666666666666*Fz);
+//            m9 = m9 + rlx_setA*((Den*(2*ux*ux-uy*uy-uz*uz)/porosity) - m9)
+//                    + (1-0.5*rlx_setA)*(4*Fx*ux-2*Fy*uy-2*Fz*uz)/porosity;
+//            m10 = m10 + rlx_setA*(-0.5*Den*((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*((Den*(uy*uy-uz*uz)/porosity) - m11)
+//                      + (1-0.5*rlx_setA)*(2*Fy*uy-2*Fz*uz)/porosity;
+//            m12 = m12 + rlx_setA*(-0.5*(Den*(uy*uy-uz*uz)/porosity)- m12)
+//                      + (1-0.5*rlx_setA)*(-Fy*uy+Fz*uz)/porosity;
+//            m13 = m13 + rlx_setA*((Den*ux*uy/porosity) - m13)
+//                      + (1-0.5*rlx_setA)*(Fy*ux+Fx*uy)/porosity;
+//            m14 = m14 + rlx_setA*((Den*uy*uz/porosity) - m14)
+//                      + (1-0.5*rlx_setA)*(Fz*uy+Fy*uz)/porosity;
+//            m15 = m15 + rlx_setA*((Den*ux*uz/porosity) - 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);
+//            //.......................................................................................................
+
+            //-------------------- IMRT collison where body force has NO higher-order terms -------------//
+            //..............carry out relaxation process...............................................
+            m1 = m1 + rlx_setA*((-30*Den+19*(ux*ux+uy*uy+uz*uz)/porosity + 57*pressure*porosity) - m1);
+            m2 = m2 + rlx_setA*((12*Den - 5.5*(ux*ux+uy*uy+uz*uz)/porosity-27*pressure*porosity) - m2);
+            jx = jx + Fx;
+            m4 = m4 + rlx_setB*((-0.6666666666666666*ux*Den) - m4)
+                    + (1-0.5*rlx_setB)*(-0.6666666666666666*Fx);
+            jy = jy + Fy;
+            m6 = m6 + rlx_setB*((-0.6666666666666666*uy*Den) - m6)
+                    + (1-0.5*rlx_setB)*(-0.6666666666666666*Fy);
+            jz = jz + Fz;
+            m8 = m8 + rlx_setB*((-0.6666666666666666*uz*Den) - m8)
+                    + (1-0.5*rlx_setB)*(-0.6666666666666666*Fz);
+            m9 = m9 + rlx_setA*((Den*(2*ux*ux-uy*uy-uz*uz)/porosity) - m9);
+            m10 = m10 + rlx_setA*(-0.5*Den*((2*ux*ux-uy*uy-uz*uz)/porosity)- m10);
+            m11 = m11 + rlx_setA*((Den*(uy*uy-uz*uz)/porosity) - m11);
+            m12 = m12 + rlx_setA*(-0.5*(Den*(uy*uy-uz*uz)/porosity)- m12);
+            m13 = m13 + rlx_setA*((Den*ux*uy/porosity) - m13);
+            m14 = m14 + rlx_setA*((Den*uy*uz/porosity) - m14);
+            m15 = m15 + rlx_setA*((Den*ux*uz/porosity) - m15);
+            m16 = m16 + rlx_setB*( - m16);
+            m17 = m17 + rlx_setB*( - m17);
+            m18 = m18 + rlx_setB*( - m18);
+            //.......................................................................................................
+
+            //.................inverse transformation......................................................
+            // q=0
+            fq = mrt_V1*Den-mrt_V2*m1+mrt_V3*m2;
+            dist[n] = fq;
+
+            // q = 1
+            fq = mrt_V1*Den-mrt_V4*m1-mrt_V5*m2+0.1*(jx-m4)+mrt_V6*(m9-m10);
+            dist[1*Np+n] = fq;
+
+            // q=2
+            fq = mrt_V1*Den-mrt_V4*m1-mrt_V5*m2+0.1*(m4-jx)+mrt_V6*(m9-m10);
+            dist[2*Np+n] = fq;
+
+            // q = 3
+            fq = mrt_V1*Den-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*Den-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*Den-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*Den-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*Den+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*Den+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*Den+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*Den+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*Den+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*Den+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*Den+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*Den+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*Den+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*Den+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*Den+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*Den+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;
+            //........................................................................
+
+            //Update velocity on device
+            Velocity[0*Np+n] = ux;
+            Velocity[1*Np+n] = uy;
+            Velocity[2*Np+n] = uz;
+            //Update pressure on device
+            Pressure[n] = pressure;
+
+		}
+	}
+}
+
+
+__global__ void dvc_ScaLBL_D3Q19_AAodd_Greyscale_IMRT(int *neighborList, double *dist, int start, int finish, int Np, double rlx,  double rlx_eff, double Gx, double Gy, double Gz,
+                                                 double *Poros,double *Perm, double *Velocity,double Den, double *Pressure){
+
+	int n, nread;
+	double vx,vy,vz,v_mag;
+    double ux,uy,uz,u_mag;
+    double pressure;//defined for this incompressible model
+	// conserved momemnts
+	double jx,jy,jz;
+	// non-conserved moments
+	double m1,m2,m4,m6,m8,m9,m10,m11,m12,m13,m14,m15,m16,m17,m18;
+    double fq;
+	//double f0,f1,f2,f3,f4,f5,f6,f7,f8,f9,f10,f11,f12,f13,f14,f15,f16,f17,f18;
+    double GeoFun;//geometric function from Guo's PRE 66, 036304 (2002)
+    double porosity;
+    double perm;//voxel permeability
+    double c0, c1; //Guo's model parameters
+    double mu_eff = (1.0/rlx_eff-0.5)/3.0;//kinematic viscosity
+    double Fx, Fy, Fz;//The total body force including Brinkman force and user-specified (Gx,Gy,Gz)
+    double rlx_setA = rlx;
+    double rlx_setB = 8.f*(2.f-rlx_setA)/(8.f-rlx_setA);
+
+	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 DISTRIBUTIONS
+            //		(read from opposite array due to previous swap operation)
+            //........................................................................
+            // q=0
+            fq = dist[n];
+            m1  = -30.0*fq;
+            m2  = 12.0*fq;
+
+            // q=1
+            nread = neighborList[n]; // neighbor 2 ( > 10Np => odd part of dist)
+            fq = dist[nread]; // reading the f1 data into register fq
+            pressure = fq;
+            m1 -= 11.0*fq;
+            m2 -= 4.0*fq;
+            jx = fq;
+            m4 = -4.0*fq;
+            m9 = 2.0*fq;
+            m10 = -4.0*fq;
+
+            // q=2
+            nread = neighborList[n+Np]; // neighbor 1 ( < 10Np => even part of dist)
+            fq = dist[nread];  // reading the f2 data into register fq
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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];
+            pressure += 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
+            nread = neighborList[n+16*Np];
+            fq = dist[nread];
+            pressure += 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];
+            pressure += 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;
+            //---------------------------------------------------------------------//
+
+            porosity = Poros[n];
+            perm = Perm[n];
+
+            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/Den+0.5*porosity*Gx;
+            vy = jy/Den+0.5*porosity*Gy;
+            vz = jz/Den+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 = Den*(-porosity*mu_eff/perm*ux - porosity*GeoFun/sqrt(perm)*u_mag*ux + porosity*Gx);
+            Fy = Den*(-porosity*mu_eff/perm*uy - porosity*GeoFun/sqrt(perm)*u_mag*uy + porosity*Gy);
+            Fz = Den*(-porosity*mu_eff/perm*uz - porosity*GeoFun/sqrt(perm)*u_mag*uz + porosity*Gz);
+            if (porosity==1.0){
+                Fx=Den*Gx;
+                Fy=Den*Gy;
+                Fz=Den*Gz;
+            }
+
+            //Calculate pressure for Incompressible-MRT model
+            pressure=0.5/porosity*(pressure-0.5*Den*u_mag*u_mag/porosity);
+
+//            //..............carry out relaxation process...............................................
+//            m1 = m1 + rlx_setA*((-30*Den+19*(ux*ux+uy*uy+uz*uz)/porosity + 57*pressure*porosity) - m1) 
+//                    + (1-0.5*rlx_setA)*38*(Fx*ux+Fy*uy+Fz*uz)/porosity;
+//            m2 = m2 + rlx_setA*((12*Den - 5.5*(ux*ux+uy*uy+uz*uz)/porosity-27*pressure*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*Den) - m4)
+//                    + (1-0.5*rlx_setB)*(-0.6666666666666666*Fx);
+//            jy = jy + Fy;
+//            m6 = m6 + rlx_setB*((-0.6666666666666666*uy*Den) - m6)
+//                    + (1-0.5*rlx_setB)*(-0.6666666666666666*Fy);
+//            jz = jz + Fz;
+//            m8 = m8 + rlx_setB*((-0.6666666666666666*uz*Den) - m8)
+//                    + (1-0.5*rlx_setB)*(-0.6666666666666666*Fz);
+//            m9 = m9 + rlx_setA*((Den*(2*ux*ux-uy*uy-uz*uz)/porosity) - m9)
+//                    + (1-0.5*rlx_setA)*(4*Fx*ux-2*Fy*uy-2*Fz*uz)/porosity;
+//            m10 = m10 + rlx_setA*(-0.5*Den*((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*((Den*(uy*uy-uz*uz)/porosity) - m11)
+//                      + (1-0.5*rlx_setA)*(2*Fy*uy-2*Fz*uz)/porosity;
+//            m12 = m12 + rlx_setA*(-0.5*(Den*(uy*uy-uz*uz)/porosity)- m12)
+//                      + (1-0.5*rlx_setA)*(-Fy*uy+Fz*uz)/porosity;
+//            m13 = m13 + rlx_setA*((Den*ux*uy/porosity) - m13)
+//                      + (1-0.5*rlx_setA)*(Fy*ux+Fx*uy)/porosity;
+//            m14 = m14 + rlx_setA*((Den*uy*uz/porosity) - m14)
+//                      + (1-0.5*rlx_setA)*(Fz*uy+Fy*uz)/porosity;
+//            m15 = m15 + rlx_setA*((Den*ux*uz/porosity) - 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);
+//            //.......................................................................................................
+           
+            //-------------------- IMRT collison where body force has NO higher-order terms -------------//
+            //..............carry out relaxation process...............................................
+            m1 = m1 + rlx_setA*((-30*Den+19*(ux*ux+uy*uy+uz*uz)/porosity + 57*pressure*porosity) - m1);
+            m2 = m2 + rlx_setA*((12*Den - 5.5*(ux*ux+uy*uy+uz*uz)/porosity-27*pressure*porosity) - m2);
+            jx = jx + Fx;
+            m4 = m4 + rlx_setB*((-0.6666666666666666*ux*Den) - m4)
+                    + (1-0.5*rlx_setB)*(-0.6666666666666666*Fx);
+            jy = jy + Fy;
+            m6 = m6 + rlx_setB*((-0.6666666666666666*uy*Den) - m6)
+                    + (1-0.5*rlx_setB)*(-0.6666666666666666*Fy);
+            jz = jz + Fz;
+            m8 = m8 + rlx_setB*((-0.6666666666666666*uz*Den) - m8)
+                    + (1-0.5*rlx_setB)*(-0.6666666666666666*Fz);
+            m9 = m9 + rlx_setA*((Den*(2*ux*ux-uy*uy-uz*uz)/porosity) - m9);
+            m10 = m10 + rlx_setA*(-0.5*Den*((2*ux*ux-uy*uy-uz*uz)/porosity)- m10);
+            m11 = m11 + rlx_setA*((Den*(uy*uy-uz*uz)/porosity) - m11);
+            m12 = m12 + rlx_setA*(-0.5*(Den*(uy*uy-uz*uz)/porosity)- m12);
+            m13 = m13 + rlx_setA*((Den*ux*uy/porosity) - m13);
+            m14 = m14 + rlx_setA*((Den*uy*uz/porosity) - m14);
+            m15 = m15 + rlx_setA*((Den*ux*uz/porosity) - m15);
+            m16 = m16 + rlx_setB*( - m16);
+            m17 = m17 + rlx_setB*( - m17);
+            m18 = m18 + rlx_setB*( - m18);
+            //.......................................................................................................
+
+
+            //.................inverse transformation......................................................
+            // q=0
+            fq = mrt_V1*Den-mrt_V2*m1+mrt_V3*m2;
+            dist[n] = fq;
+
+            // q = 1
+            fq = mrt_V1*Den-mrt_V4*m1-mrt_V5*m2+0.1*(jx-m4)+mrt_V6*(m9-m10);
+            nread = neighborList[n+Np];
+            dist[nread] = fq;
+
+            // q=2
+            fq = mrt_V1*Den-mrt_V4*m1-mrt_V5*m2+0.1*(m4-jx)+mrt_V6*(m9-m10);
+            nread = neighborList[n];
+            dist[nread] = fq;
+
+            // q = 3
+            fq = mrt_V1*Den-mrt_V4*m1-mrt_V5*m2+0.1*(jy-m6)+mrt_V7*(m10-m9)+mrt_V8*(m11-m12);
+            nread = neighborList[n+3*Np];
+            dist[nread] = fq;
+
+            // q = 4
+            fq = mrt_V1*Den-mrt_V4*m1-mrt_V5*m2+0.1*(m6-jy)+mrt_V7*(m10-m9)+mrt_V8*(m11-m12);
+            nread = neighborList[n+2*Np];
+            dist[nread] = fq;
+
+            // q = 5
+            fq = mrt_V1*Den-mrt_V4*m1-mrt_V5*m2+0.1*(jz-m8)+mrt_V7*(m10-m9)+mrt_V8*(m12-m11);
+            nread = neighborList[n+5*Np];
+            dist[nread] = fq;
+
+            // q = 6
+            fq = mrt_V1*Den-mrt_V4*m1-mrt_V5*m2+0.1*(m8-jz)+mrt_V7*(m10-m9)+mrt_V8*(m12-m11);
+            nread = neighborList[n+4*Np];
+            dist[nread] = fq;
+
+            // q = 7
+            fq = mrt_V1*Den+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[nread] = fq;
+
+            // q = 8
+            fq = mrt_V1*Den+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[nread] = fq;
+
+            // q = 9
+            fq = mrt_V1*Den+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[nread] = fq;
+
+            // q = 10
+            fq = mrt_V1*Den+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[nread] = fq;
+
+            // q = 11
+            fq = mrt_V1*Den+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[nread] = fq;
+
+            // q = 12
+            fq = mrt_V1*Den+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[nread]= fq;
+
+            // q = 13
+            fq = mrt_V1*Den+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[nread] = fq;
+
+            // q= 14
+            fq = mrt_V1*Den+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[nread] = fq;
+
+            // q = 15
+            fq = mrt_V1*Den+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*Den+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*Den+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*Den+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;
+            //........................................................................
+
+            //Update velocity on device
+            Velocity[0*Np+n] = ux;
+            Velocity[1*Np+n] = uy;
+            Velocity[2*Np+n] = uz;
+            //Update pressure on device
+            Pressure[n] = pressure;
+
+		}
+	}
+}
+
+
+extern "C" void ScaLBL_D3Q19_AAeven_Greyscale(double *dist, int start, int finish, int Np, double rlx, double rlx_eff, double Fx, double Fy, double Fz,double *Poros,double *Perm, double *Velocity,double *Pressure){
+	
+    dvc_ScaLBL_D3Q19_AAeven_Greyscale<<<NBLOCKS,NTHREADS >>>(dist,start,finish,Np,rlx,rlx_eff,Fx,Fy,Fz,Poros,Perm,Velocity,Pressure);
+
+    cudaError_t err = cudaGetLastError();
+	if (cudaSuccess != err){
+		printf("CUDA error in ScaLBL_D3Q19_AAeven_Greyscale: %s \n",cudaGetErrorString(err));
+	}
+}
+
+extern "C" void ScaLBL_D3Q19_AAodd_Greyscale(int *neighborList, double *dist, int start, int finish, int Np, double rlx, double rlx_eff, double Fx, double Fy, double Fz,double *Poros,double *Perm, double *Velocity,double *Pressure){
+
+    dvc_ScaLBL_D3Q19_AAodd_Greyscale<<<NBLOCKS,NTHREADS >>>(neighborList,dist,start,finish,Np,rlx,rlx_eff,Fx,Fy,Fz,Poros,Perm,Velocity,Pressure);
+
+    cudaError_t err = cudaGetLastError();
+	if (cudaSuccess != err){
+		printf("CUDA error in ScaLBL_D3Q19_AAodd_Greyscale: %s \n",cudaGetErrorString(err));
+	}
+}
+
+extern "C" void ScaLBL_D3Q19_AAeven_Greyscale_IMRT(double *dist, int start, int finish, int Np, double rlx, double rlx_eff, double Fx, double Fy, double Fz,double *Poros,double *Perm, double *Velocity,double Den,double *Pressure){
+	
+    dvc_ScaLBL_D3Q19_AAeven_Greyscale_IMRT<<<NBLOCKS,NTHREADS >>>(dist,start,finish,Np,rlx,rlx_eff,Fx,Fy,Fz,Poros,Perm,Velocity,Den,Pressure);
+
+    cudaError_t err = cudaGetLastError();
+	if (cudaSuccess != err){
+		printf("CUDA error in ScaLBL_D3Q19_AAeven_Greyscale_IMRT: %s \n",cudaGetErrorString(err));
+	}
+}
+
+extern "C" void ScaLBL_D3Q19_AAodd_Greyscale_IMRT(int *neighborList, double *dist, int start, int finish, int Np, double rlx, double rlx_eff, double Fx, double Fy, double Fz,double *Poros,double *Perm, double *Velocity,double Den,double *Pressure){
+
+    dvc_ScaLBL_D3Q19_AAodd_Greyscale_IMRT<<<NBLOCKS,NTHREADS >>>(neighborList,dist,start,finish,Np,rlx,rlx_eff,Fx,Fy,Fz,Poros,Perm,Velocity,Den,Pressure);
+
+    cudaError_t err = cudaGetLastError();
+	if (cudaSuccess != err){
+		printf("CUDA error in ScaLBL_D3Q19_AAodd_Greyscale_IMRT: %s \n",cudaGetErrorString(err));
+	}
+}
+
diff --git a/models/GreyscaleModel.cpp b/models/GreyscaleModel.cpp
new file mode 100644
index 00000000..11d92c80
--- /dev/null
+++ b/models/GreyscaleModel.cpp
@@ -0,0 +1,898 @@
+/*
+Greyscale lattice boltzmann model
+ */
+#include "models/GreyscaleModel.h"
+#include "analysis/distance.h"
+#include "analysis/morphology.h"
+#include <stdlib.h>
+#include <time.h>
+
+template<class TYPE>
+void DeleteArray( const TYPE *p )
+{
+    delete [] p;
+}
+
+ScaLBL_GreyscaleModel::ScaLBL_GreyscaleModel(int RANK, int NP, MPI_Comm COMM):
+rank(RANK), nprocs(NP), Restart(0),timestep(0),timestepMax(0),tau(0),tau_eff(0),Den(0),Fx(0),Fy(0),Fz(0),flux(0),din(0),dout(0),GreyPorosity(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)
+{
+	SignDist.resize(Nx,Ny,Nz);           
+    SignDist.fill(0);
+
+}
+ScaLBL_GreyscaleModel::~ScaLBL_GreyscaleModel(){
+
+}
+
+void ScaLBL_GreyscaleModel::ReadParams(string filename){
+	// read the input database 
+	db = std::make_shared<Database>( filename );
+	domain_db = db->getDatabase( "Domain" );
+	greyscale_db =  db->getDatabase( "Greyscale" );
+	analysis_db = db->getDatabase( "Analysis" );
+	vis_db = db->getDatabase( "Visualization" );
+
+	// set defaults
+	timestepMax = 100000;
+	tau = 1.0;
+    tau_eff = tau;
+    Den = 1.0;//constant density
+	tolerance = 0.01;
+	Fx = Fy = Fz = 0.0;
+	Restart=false;
+	din=dout=1.0;
+	flux=0.0;
+    dp = 10.0; //unit of 'dp': voxel
+    CollisionType = 1; //1: IMRT; 2: BGK
+	
+	// ---------------------- Greyscale Model parameters -----------------------//
+	if (greyscale_db->keyExists( "timestepMax" )){
+		timestepMax = greyscale_db->getScalar<int>( "timestepMax" );
+	}
+	if (greyscale_db->keyExists( "tau" )){
+		tau = greyscale_db->getScalar<double>( "tau" );
+	}
+	tau_eff = greyscale_db->getWithDefault<double>( "tau_eff", tau );
+	if (greyscale_db->keyExists( "Den" )){
+		Den = greyscale_db->getScalar<double>( "Den" );
+	}
+	if (greyscale_db->keyExists( "dp" )){
+		dp = greyscale_db->getScalar<double>( "dp" );
+	}
+	if (greyscale_db->keyExists( "F" )){
+		Fx = greyscale_db->getVector<double>( "F" )[0];
+		Fy = greyscale_db->getVector<double>( "F" )[1];
+		Fz = greyscale_db->getVector<double>( "F" )[2];
+	}
+	if (greyscale_db->keyExists( "Restart" )){
+		Restart = greyscale_db->getScalar<bool>( "Restart" );
+	}
+	if (greyscale_db->keyExists( "din" )){
+		din = greyscale_db->getScalar<double>( "din" );
+	}
+	if (greyscale_db->keyExists( "dout" )){
+		dout = greyscale_db->getScalar<double>( "dout" );
+	}
+	if (greyscale_db->keyExists( "flux" )){
+		flux = greyscale_db->getScalar<double>( "flux" );
+	}
+	if (greyscale_db->keyExists( "tolerance" )){
+		tolerance = greyscale_db->getScalar<double>( "tolerance" );
+	}
+	auto collision = greyscale_db->getWithDefault<std::string>( "collision", "IMRT" );
+	if (collision == "BGK"){
+        CollisionType=2;
+	}
+	// ------------------------------------------------------------------------//
+    
+    //------------------------ Other Domain parameters ------------------------//
+	BoundaryCondition = 0;
+	if (domain_db->keyExists( "BC" )){
+		BoundaryCondition = domain_db->getScalar<int>( "BC" );
+	}
+	// ------------------------------------------------------------------------//
+}
+
+void ScaLBL_GreyscaleModel::SetDomain(){
+	Dm  = std::shared_ptr<Domain>(new Domain(domain_db,comm));      // full domain for analysis
+	Mask  = std::shared_ptr<Domain>(new Domain(domain_db,comm));    // mask domain removes immobile phases
+	// domain parameters
+	Nx = Dm->Nx;
+	Ny = Dm->Ny;
+	Nz = Dm->Nz;
+	Lx = Dm->Lx;
+	Ly = Dm->Ly;
+	Lz = Dm->Lz;
+	N = Nx*Ny*Nz;
+
+	SignDist.resize(Nx,Ny,Nz);
+	Velocity_x.resize(Nx,Ny,Nz);
+	Velocity_y.resize(Nx,Ny,Nz);
+	Velocity_z.resize(Nx,Ny,Nz);
+	PorosityMap.resize(Nx,Ny,Nz);
+	Pressure.resize(Nx,Ny,Nz);
+
+	id = new signed char [N];
+	for (int i=0; i<Nx*Ny*Nz; i++) Dm->id[i] = 1;               // initialize this way
+	MPI_Barrier(comm);
+	Dm->CommInit();
+	MPI_Barrier(comm);
+	// Read domain parameters
+	rank = Dm->rank();	
+	nprocx = Dm->nprocx();
+	nprocy = Dm->nprocy();
+	nprocz = Dm->nprocz();
+}
+
+void ScaLBL_GreyscaleModel::ReadInput(){
+	
+	sprintf(LocalRankString,"%05d",rank);
+	sprintf(LocalRankFilename,"%s%s","ID.",LocalRankString);
+	sprintf(LocalRestartFile,"%s%s","Restart.",LocalRankString);
+	
+	if (domain_db->keyExists( "Filename" )){
+		auto Filename = domain_db->getScalar<std::string>( "Filename" );
+		Mask->Decomp(Filename);
+	}
+	else{
+        if (rank==0) printf("Filename of input image is not found, reading ID.0* instead.");
+		Mask->ReadIDs();
+	}
+	for (int i=0; i<Nx*Ny*Nz; i++) id[i] = Mask->id[i];  // save what was read
+	
+	// Generate the signed distance map
+	// Initialize the domain and communication
+	Array<char> id_solid(Nx,Ny,Nz);
+	int count = 0;
+	// Solve for the position of the solid phase
+	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;
+				// Initialize the solid phase
+				signed char label = Mask->id[n];
+				if (label > 0)		id_solid(i,j,k) = 1;
+				else	     		id_solid(i,j,k) = 0;
+			}
+		}
+	}
+	// Initialize the signed distance function
+	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;
+				// Initialize distance to +/- 1
+				SignDist(i,j,k) = 2.0*double(id_solid(i,j,k))-1.0;
+			}
+		}
+	}
+//	MeanFilter(SignDist);
+	if (rank==0) printf("Initialized solid phase -- Converting to Signed Distance function \n");
+	CalcDist(SignDist,id_solid,*Mask);
+	
+	if (rank == 0) cout << "Domain set." << endl;
+}
+
+/********************************************************
+ * AssignComponentLabels                                 *
+ ********************************************************/
+void ScaLBL_GreyscaleModel::AssignComponentLabels(double *Porosity, double *Permeability)
+{
+	size_t NLABELS=0;
+	signed char VALUE=0;
+	double POROSITY=0.f;
+	double PERMEABILITY=0.f;
+
+	auto LabelList = greyscale_db->getVector<int>( "ComponentLabels" );
+	auto PorosityList = greyscale_db->getVector<double>( "PorosityList" );
+	auto PermeabilityList = greyscale_db->getVector<double>( "PermeabilityList" );
+
+	NLABELS=LabelList.size();
+	if (NLABELS != PorosityList.size()){
+		ERROR("Error: ComponentLabels and PorosityList must be the same length! \n");
+	}
+
+	double label_count[NLABELS];
+	double label_count_global[NLABELS];
+	// Assign the labels
+
+	for (int idx=0; idx<NLABELS; idx++) label_count[idx]=0;
+
+	for (int k=1;k<Nz-1;k++){
+		for (int j=1;j<Ny-1;j++){
+			for (int i=1;i<Nx-1;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]){
+						POROSITY=PorosityList[idx];
+						label_count[idx] += 1.0;
+						idx = NLABELS;
+						//Mask->id[n] = 0; // set mask to zero since this is an immobile component
+					}
+				}
+				int idx = Map(i,j,k);
+				if (!(idx < 0)){
+                    if (POROSITY<=0.0){
+                        ERROR("Error: Porosity for grey voxels must be 0.0 < Porosity <= 1.0 !\n");
+                    }
+                    else{
+					    Porosity[idx] = POROSITY;
+                    }
+                }
+			}
+		}
+	}
+
+	if (NLABELS != PermeabilityList.size()){
+		ERROR("Error: ComponentLabels and PermeabilityList must be the same length! \n");
+	}
+	for (int k=1;k<Nz-1;k++){
+		for (int j=1;j<Ny-1;j++){
+			for (int i=1;i<Nx-1;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]){
+						PERMEABILITY=PermeabilityList[idx];
+						idx = NLABELS;
+						//Mask->id[n] = 0; // set mask to zero since this is an immobile component
+					}
+				}
+				int idx = Map(i,j,k);
+				if (!(idx < 0)){
+                    if (PERMEABILITY<=0.0){
+                        ERROR("Error: Permeability for grey voxel must be > 0.0 ! \n");
+                    }
+                    else{
+					    Permeability[idx] = PERMEABILITY/Dm->voxel_length/Dm->voxel_length;
+                    }
+                }
+			}
+		}
+	}
+
+
+	// Set Dm to match Mask
+	for (int i=0; i<Nx*Ny*Nz; i++) Dm->id[i] = Mask->id[i]; 
+	
+	for (int idx=0; idx<NLABELS; idx++)		label_count_global[idx]=Dm->Comm.sumReduce(label_count[idx]);
+
+    //Initialize a weighted porosity after considering grey voxels
+    GreyPorosity=0.0;
+	for (unsigned int idx=0; idx<NLABELS; idx++){
+		double volume_fraction  = double(label_count_global[idx])/double((Nx-2)*(Ny-2)*(Nz-2)*nprocs);
+        GreyPorosity+=volume_fraction*PorosityList[idx];
+    }
+
+	if (rank==0){
+        printf("Image resolution: %.5g [um/voxel]\n",Dm->voxel_length);
+		printf("Component labels: %lu \n",NLABELS);
+		for (unsigned int idx=0; idx<NLABELS; idx++){
+			VALUE=LabelList[idx];
+			POROSITY=PorosityList[idx];
+			PERMEABILITY=PermeabilityList[idx];
+			double volume_fraction  = double(label_count_global[idx])/double((Nx-2)*(Ny-2)*(Nz-2)*nprocs);
+			printf("   label=%d: porosity=%.3g, permeability=%.3g [um^2] (=%.3g [voxel^2]), volume fraction=%.3g\n",
+                    VALUE,POROSITY,PERMEABILITY,PERMEABILITY/Dm->voxel_length/Dm->voxel_length,volume_fraction); 
+            printf("             effective porosity=%.3g\n",volume_fraction*POROSITY);
+		}
+        printf("The weighted porosity, considering both open and grey voxels, is %.3g\n",GreyPorosity);
+	}
+}
+
+
+void ScaLBL_GreyscaleModel::Create(){
+	/*
+	 *  This function creates the variables needed to run a LBM 
+	 */
+	//.........................................................
+	// don't perform computations at the eight corners
+	//id[0] = id[Nx-1] = id[(Ny-1)*Nx] = id[(Ny-1)*Nx + Nx-1] = 0;
+	//id[(Nz-1)*Nx*Ny] = id[(Nz-1)*Nx*Ny+Nx-1] = id[(Nz-1)*Nx*Ny+(Ny-1)*Nx] = id[(Nz-1)*Nx*Ny+(Ny-1)*Nx + Nx-1] = 0;
+
+	//.........................................................
+	// Initialize communication structures in averaging domain
+	for (int i=0; i<Nx*Ny*Nz; i++) Dm->id[i] = Mask->id[i];
+	Mask->CommInit();
+	Np=Mask->PoreCount();
+	//...........................................................................
+	if (rank==0)    printf ("Create ScaLBL_Communicator \n");
+	// Create a communicator for the device (will use optimized layout)
+	// ScaLBL_Communicator ScaLBL_Comm(Mask); // original
+	ScaLBL_Comm  = 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);
+	Map.resize(Nx,Ny,Nz);       Map.fill(-2);
+	auto neighborList= new int[18*Npad];
+	Np = ScaLBL_Comm->MemoryOptimizedLayoutAA(Map,neighborList,Mask->id,Np);
+	MPI_Barrier(comm);
+
+	//...........................................................................
+	//                MAIN  VARIABLES ALLOCATED HERE
+	//...........................................................................
+	// LBM variables
+	if (rank==0)    printf ("Allocating distributions \n");
+	//......................device distributions.................................
+	dist_mem_size = Np*sizeof(double);
+	neighborSize=18*(Np*sizeof(int));
+	//...........................................................................
+	ScaLBL_AllocateDeviceMemory((void **) &NeighborList, neighborSize);
+	ScaLBL_AllocateDeviceMemory((void **) &dvcMap, sizeof(int)*Np);
+	ScaLBL_AllocateDeviceMemory((void **) &fq, 19*dist_mem_size);
+	ScaLBL_AllocateDeviceMemory((void **) &Permeability, sizeof(double)*Np);		
+	ScaLBL_AllocateDeviceMemory((void **) &Porosity, sizeof(double)*Np);		
+	ScaLBL_AllocateDeviceMemory((void **) &Pressure_dvc, sizeof(double)*Np);
+	ScaLBL_AllocateDeviceMemory((void **) &Velocity, 3*sizeof(double)*Np);
+	//...........................................................................
+	// Update GPU data structures
+	if (rank==0)	printf ("Setting up device map and neighbor list \n");
+	fflush(stdout);
+	int *TmpMap;
+	TmpMap=new int[Np];
+	for (int k=1; k<Nz-1; k++){
+		for (int j=1; j<Ny-1; j++){
+			for (int i=1; i<Nx-1; i++){
+				int idx=Map(i,j,k);
+				if (!(idx < 0))
+					TmpMap[idx] = k*Nx*Ny+j*Nx+i;
+			}
+		}
+	}
+	// check that TmpMap is valid
+	for (int idx=0; idx<ScaLBL_Comm->LastExterior(); idx++){
+		int n = TmpMap[idx];
+		if (n > Nx*Ny*Nz){
+			printf("Bad value! idx=%i \n");
+			TmpMap[idx] = Nx*Ny*Nz-1;
+		}
+	}
+	for (int idx=ScaLBL_Comm->FirstInterior(); idx<ScaLBL_Comm->LastInterior(); idx++){
+		int n = TmpMap[idx];
+		if (n > Nx*Ny*Nz){
+			printf("Bad value! idx=%i \n");
+			TmpMap[idx] = Nx*Ny*Nz-1;
+		}
+	}
+	ScaLBL_CopyToDevice(dvcMap, TmpMap, sizeof(int)*Np);
+	ScaLBL_DeviceBarrier();
+	delete [] TmpMap;
+	
+	// copy the neighbor list 
+	ScaLBL_CopyToDevice(NeighborList, neighborList, neighborSize);
+	// initialize phi based on PhaseLabel (include solid component labels)
+	double *Poros, *Perm;
+	Poros = new double[Np];
+	Perm = new double[Np];
+	AssignComponentLabels(Poros,Perm);
+	ScaLBL_CopyToDevice(Porosity, Poros, Np*sizeof(double));
+	ScaLBL_CopyToDevice(Permeability, Perm, Np*sizeof(double));
+}        
+
+
+void ScaLBL_GreyscaleModel::Initialize(){
+	if (rank==0)	printf ("Initializing distributions \n");
+    //TODO: for BGK, you need to consider voxel porosity
+    //      for IMRT, the whole set of feq is different
+    //      if in the future you have different collison mode, need to write two set of initialization functions
+    if (CollisionType==1){
+	    ScaLBL_D3Q19_GreyIMRT_Init(fq, Np, Den);
+        if (rank==0) printf("Collision model: Incompressible MRT.\n");
+    }
+    else if (CollisionType==2){
+	    ScaLBL_D3Q19_Init(fq, Np);
+        if (rank==0) printf("Collision model: BGK.\n");
+    }
+    else{
+        if (rank==0) printf("Unknown collison type! IMRT collision is used.\n"); 
+	    ScaLBL_D3Q19_GreyIMRT_Init(fq, Np, Den);
+        CollisionType=1;
+		greyscale_db->putScalar<std::string>( "collision", "IMRT" );
+    }
+
+	if (Restart == true){
+		if (rank==0){
+			printf("Initializing distributions from Restart! \n");
+		}
+		// Read in the restart file to CPU buffers
+        std::shared_ptr<double> cfq;
+        cfq = std::shared_ptr<double>(new double[19*Np],DeleteArray<double>);
+        FILE *File;
+        File=fopen(LocalRestartFile,"rb");
+        fread(cfq.get(),sizeof(double),19*Np,File);
+        fclose(File);
+
+		// Copy the restart data to the GPU
+		ScaLBL_CopyToDevice(fq,cfq.get(),19*Np*sizeof(double));
+		ScaLBL_DeviceBarrier();
+
+		MPI_Barrier(comm);
+	}
+}
+
+void ScaLBL_GreyscaleModel::Run(){
+	int nprocs=nprocx*nprocy*nprocz;
+	const RankInfoStruct rank_info(rank,nprocx,nprocy,nprocz);
+	
+	int analysis_interval = 1000; 	// number of timesteps in between in situ analysis 
+	int visualization_interval = 1000; 	 
+	int restart_interval = 10000; 	// number of timesteps in between in saving distributions for restart 
+	if (analysis_db->keyExists( "analysis_interval" )){
+		analysis_interval = analysis_db->getScalar<int>( "analysis_interval" );
+	}
+	if (analysis_db->keyExists( "visualization_interval" )){
+		visualization_interval = analysis_db->getScalar<int>( "visualization_interval" );
+	}
+	if (analysis_db->keyExists( "restart_interval" )){
+		restart_interval = analysis_db->getScalar<int>( "restart_interval" );
+	}
+	if (greyscale_db->keyExists( "timestep" )){
+		timestep = greyscale_db->getScalar<int>( "timestep" );
+	}
+
+	if (rank==0){
+		printf("********************************************************\n");
+		printf("No. of timesteps: %i \n", timestepMax);
+		fflush(stdout);
+	}
+
+	//.......create and start timer............
+	double starttime,stoptime,cputime;
+	ScaLBL_DeviceBarrier();
+	MPI_Barrier(comm);
+	starttime = MPI_Wtime();
+	//.........................................
+	
+	Minkowski Morphology(Mask);
+
+	//************ MAIN ITERATION LOOP ***************************************/
+	PROFILE_START("Loop");
+	auto current_db = db->cloneDatabase();
+	double rlx = 1.0/tau;
+    double rlx_eff = 1.0/tau_eff;
+	double error = 1.0;
+	double flow_rate_previous = 0.0;
+	while (timestep < timestepMax && error > tolerance) {
+		//************************************************************************/
+		// *************ODD TIMESTEP*************//
+		timestep++;
+		ScaLBL_Comm->SendD3Q19AA(fq); //READ FROM NORMAL
+        switch (CollisionType){
+            case 1: 
+                    ScaLBL_D3Q19_AAodd_Greyscale_IMRT(NeighborList, fq,  ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np, rlx, rlx_eff, Fx, Fy, Fz,Porosity,Permeability,Velocity,Den,Pressure_dvc);
+                    break;
+            case 2: 
+                    ScaLBL_D3Q19_AAodd_Greyscale(NeighborList, fq,  ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np, rlx, rlx_eff, Fx, Fy, Fz,Porosity,Permeability,Velocity,Pressure_dvc);
+                    break;
+            default: 
+                    ScaLBL_D3Q19_AAodd_Greyscale_IMRT(NeighborList, fq,  ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np, rlx, rlx_eff, Fx, Fy, Fz,Porosity,Permeability,Velocity,Den,Pressure_dvc);
+                    break;
+        }
+		ScaLBL_Comm->RecvD3Q19AA(fq); //WRITE INTO OPPOSITE
+		ScaLBL_DeviceBarrier();
+		// Set BCs
+		if (BoundaryCondition == 3){
+			ScaLBL_Comm->D3Q19_Pressure_BC_z(NeighborList, fq, din, timestep);
+			ScaLBL_Comm->D3Q19_Pressure_BC_Z(NeighborList, fq, dout, timestep);
+		}
+        switch (CollisionType){
+            case 1: 
+		            ScaLBL_D3Q19_AAodd_Greyscale_IMRT(NeighborList, fq, 0, ScaLBL_Comm->LastExterior(), Np, rlx, rlx_eff, Fx, Fy, Fz,Porosity,Permeability,Velocity,Den,Pressure_dvc);
+                    break;
+            case 2: 
+		            ScaLBL_D3Q19_AAodd_Greyscale(NeighborList, fq, 0, ScaLBL_Comm->LastExterior(), Np, rlx, rlx_eff, Fx, Fy, Fz,Porosity,Permeability,Velocity,Pressure_dvc);
+                    break;
+            default: 
+		            ScaLBL_D3Q19_AAodd_Greyscale_IMRT(NeighborList, fq, 0, ScaLBL_Comm->LastExterior(), Np, rlx, rlx_eff, Fx, Fy, Fz,Porosity,Permeability,Velocity,Den,Pressure_dvc);
+                    break;
+        }
+		ScaLBL_DeviceBarrier(); MPI_Barrier(comm);
+
+		// *************EVEN TIMESTEP*************//
+		timestep++;
+		ScaLBL_Comm->SendD3Q19AA(fq); //READ FORM NORMAL
+        switch (CollisionType){
+            case 1: 
+		            ScaLBL_D3Q19_AAeven_Greyscale_IMRT(fq, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np, rlx, rlx_eff, Fx, Fy, Fz,Porosity,Permeability,Velocity,Den,Pressure_dvc);
+                    break;
+            case 2: 
+		            ScaLBL_D3Q19_AAeven_Greyscale(fq, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np, rlx, rlx_eff, Fx, Fy, Fz,Porosity,Permeability,Velocity,Pressure_dvc);
+                    break;
+            default: 
+		            ScaLBL_D3Q19_AAeven_Greyscale_IMRT(fq, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np, rlx, rlx_eff, Fx, Fy, Fz,Porosity,Permeability,Velocity,Den,Pressure_dvc);
+                    break;
+        }
+        ScaLBL_Comm->RecvD3Q19AA(fq); //WRITE INTO OPPOSITE
+		ScaLBL_DeviceBarrier();
+		// Set BCs
+		if (BoundaryCondition == 3){
+			ScaLBL_Comm->D3Q19_Pressure_BC_z(NeighborList, fq, din, timestep);
+			ScaLBL_Comm->D3Q19_Pressure_BC_Z(NeighborList, fq, dout, timestep);
+		}
+        switch (CollisionType){
+            case 1: 
+		            ScaLBL_D3Q19_AAeven_Greyscale_IMRT(fq, 0, ScaLBL_Comm->LastExterior(), Np, rlx, rlx_eff, Fx, Fy, Fz,Porosity,Permeability,Velocity,Den,Pressure_dvc);
+                    break;
+            case 2: 
+		            ScaLBL_D3Q19_AAeven_Greyscale(fq, 0, ScaLBL_Comm->LastExterior(), Np, rlx, rlx_eff, Fx, Fy, Fz,Porosity,Permeability,Velocity,Pressure_dvc);
+                    break;
+            default: 
+		            ScaLBL_D3Q19_AAeven_Greyscale_IMRT(fq, 0, ScaLBL_Comm->LastExterior(), Np, rlx, rlx_eff, Fx, Fy, Fz,Porosity,Permeability,Velocity,Den,Pressure_dvc);
+                    break;
+        }
+        ScaLBL_DeviceBarrier(); MPI_Barrier(comm);
+		//************************************************************************/
+		
+		if (timestep%analysis_interval==0){
+			ScaLBL_Comm->RegularLayout(Map,&Velocity[0],Velocity_x);
+			ScaLBL_Comm->RegularLayout(Map,&Velocity[Np],Velocity_y);
+			ScaLBL_Comm->RegularLayout(Map,&Velocity[2*Np],Velocity_z);
+			//ScaLBL_Comm->RegularLayout(Map,Porosity,PorosityMap);
+			//ScaLBL_Comm->RegularLayout(Map,Pressure_dvc,Pressure);
+			
+			double count_loc=0;
+			double count;
+			double vax,vay,vaz;
+			double vax_loc,vay_loc,vaz_loc;
+            //double px_loc,py_loc,pz_loc;
+            //double px,py,pz;
+            //double mass_loc,mass_glb;
+            
+            //parameters for domain average
+	        int64_t i,j,k,n,imin,jmin,kmin,kmax;
+            // If external boundary conditions are set, do not average over the inlet and outlet
+            kmin=1; kmax=Nz-1;
+            //In case user forgets to specify the inlet/outlet buffer layers for BC>0
+            if (BoundaryCondition > 0 && Dm->kproc() == 0) kmin=4;
+            if (BoundaryCondition > 0 && Dm->kproc() == Dm->nprocz()-1) kmax=Nz-4;
+
+            imin=jmin=1;
+            // If inlet/outlet layers exist use these as default
+            //if (Dm->inlet_layers_x > 0) imin = Dm->inlet_layers_x;
+            //if (Dm->inlet_layers_y > 0) jmin = Dm->inlet_layers_y;
+            if (BoundaryCondition > 0 && Dm->inlet_layers_z > 0 && Dm->kproc() == 0) kmin = 1 + Dm->inlet_layers_z;//"1" indicates the halo layer
+            if (BoundaryCondition > 0 && Dm->outlet_layers_z > 0 && Dm->kproc() == Dm->nprocz()-1) kmax = Nz-1 - Dm->outlet_layers_z; 
+
+//			px_loc = py_loc = pz_loc = 0.f;
+//            mass_loc = 0.f;
+//			for (int k=kmin; k<kmax; k++){
+//				for (int j=jmin; j<Ny-1; j++){
+//					for (int i=imin; i<Nx-1; i++){
+//						if (SignDist(i,j,k) > 0){
+//							px_loc   += Velocity_x(i,j,k)*Den*PorosityMap(i,j,k);
+//							py_loc   += Velocity_y(i,j,k)*Den*PorosityMap(i,j,k);
+//							pz_loc   += Velocity_z(i,j,k)*Den*PorosityMap(i,j,k);
+//							mass_loc += Den*PorosityMap(i,j,k);
+//						}
+//					}
+//				}
+//			}
+//			MPI_Allreduce(&px_loc,  &px,      1,MPI_DOUBLE,MPI_SUM,Mask->Comm);
+//			MPI_Allreduce(&py_loc,  &py,      1,MPI_DOUBLE,MPI_SUM,Mask->Comm);
+//			MPI_Allreduce(&pz_loc,  &pz,      1,MPI_DOUBLE,MPI_SUM,Mask->Comm);
+//			MPI_Allreduce(&mass_loc,&mass_glb,1,MPI_DOUBLE,MPI_SUM,Mask->Comm);
+//			
+//			vax = px/mass_glb;
+//			vay = py/mass_glb;
+//			vaz = pz/mass_glb;
+            
+			vax_loc = vay_loc = vaz_loc = 0.f;
+			for (int k=kmin; k<kmax; k++){
+				for (int j=jmin; j<Ny-1; j++){
+					for (int i=imin; i<Nx-1; i++){
+						if (SignDist(i,j,k) > 0){
+							vax_loc += Velocity_x(i,j,k);
+							vay_loc += Velocity_y(i,j,k);
+							vaz_loc += Velocity_z(i,j,k);
+							count_loc+=1.0;
+						}
+					}
+				}
+			}
+			//MPI_Allreduce(&vax_loc,&vax,1,MPI_DOUBLE,MPI_SUM,Mask->Comm);
+			//MPI_Allreduce(&vay_loc,&vay,1,MPI_DOUBLE,MPI_SUM,Mask->Comm);
+			//MPI_Allreduce(&vaz_loc,&vaz,1,MPI_DOUBLE,MPI_SUM,Mask->Comm);
+			//MPI_Allreduce(&count_loc,&count,1,MPI_DOUBLE,MPI_SUM,Mask->Comm);
+			
+            vax = Mask->Comm.sumReduce( vax_loc );
+            vay = Mask->Comm.sumReduce( vay_loc );
+            vaz = Mask->Comm.sumReduce( vaz_loc );
+            count = Mask->Comm.sumReduce( count_loc );
+
+			vax /= count;
+			vay /= count;
+			vaz /= count;
+
+			double force_mag = sqrt(Fx*Fx+Fy*Fy+Fz*Fz);
+			double dir_x = Fx/force_mag;
+			double dir_y = Fy/force_mag;
+			double dir_z = Fz/force_mag;
+			if (force_mag == 0.0){
+				// default to z direction
+				dir_x = 0.0;
+				dir_y = 0.0;
+				dir_z = 1.0;
+				force_mag = 1.0;
+			}
+			//double flow_rate = (px*dir_x + py*dir_y + pz*dir_z)/mass_glb;
+			double flow_rate = (vax*dir_x + vay*dir_y + vaz*dir_z);
+			
+			error = fabs(flow_rate - flow_rate_previous) / fabs(flow_rate);
+			flow_rate_previous = flow_rate;
+			
+			//if (rank==0) printf("Computing Minkowski functionals \n");
+			Morphology.ComputeScalar(SignDist,0.f);
+			//Morphology.PrintAll();
+			double mu = (tau-0.5)/3.f;
+			double Vs = Morphology.V();
+			double As = Morphology.A();
+			double Hs = Morphology.H();
+			double Xs = Morphology.X();
+			Vs = Dm->Comm.sumReduce( Vs);
+			As = Dm->Comm.sumReduce( As);
+			Hs = Dm->Comm.sumReduce( Hs);
+			Xs = Dm->Comm.sumReduce( Xs);
+
+			double h = Dm->voxel_length;
+			//double absperm = h*h*mu*Mask->Porosity()*flow_rate / force_mag;
+			double absperm = h*h*mu*GreyPorosity*flow_rate / force_mag;
+
+            if (rank==0){
+				printf("     AbsPerm = %.5g [micron^2]\n",absperm);
+                bool WriteHeader=false;
+                FILE * log_file = fopen("Permeability.csv","r");
+                if (log_file != NULL)
+                    fclose(log_file);
+                else
+                    WriteHeader=true;
+                log_file = fopen("Permeability.csv","a");
+                if (WriteHeader)
+                    fprintf(log_file,"timestep Fx Fy Fz mu Vs As Hs Xs vax vay vaz AbsPerm \n",
+                            timestep,Fx,Fy,Fz,mu,h*h*h*Vs,h*h*As,h*Hs,Xs,vax,vay,vaz,absperm);
+
+                fprintf(log_file,"%i %.8g %.8g %.8g %.8g %.8g %.8g %.8g %.8g %.8g %.8g %.8g %.8g\n",timestep, Fx, Fy, Fz, mu, 
+                        h*h*h*Vs,h*h*As,h*Hs,Xs,vax,vay,vaz, absperm);
+                fclose(log_file);
+            }
+		}
+
+		if (timestep%visualization_interval==0){
+            VelocityField();
+        }
+
+		if (timestep%restart_interval==0){
+            //Use rank=0 write out Restart.db
+            if (rank==0) {
+                greyscale_db->putScalar<int>("timestep",timestep);    		
+                greyscale_db->putScalar<bool>( "Restart", true );
+                current_db->putDatabase("Greyscale", greyscale_db);
+                std::ofstream OutStream("Restart.db");
+                current_db->print(OutStream, "");
+                OutStream.close();
+      
+            }
+            //Write out Restart data.
+            std::shared_ptr<double> cfq;
+            cfq = std::shared_ptr<double>(new double[19*Np],DeleteArray<double>);
+            ScaLBL_CopyToHost(cfq.get(),fq,19*Np*sizeof(double));// Copy restart data to the CPU
+
+            FILE *RESTARTFILE;
+            RESTARTFILE=fopen(LocalRestartFile,"wb");
+            fwrite(cfq.get(),sizeof(double),19*Np,RESTARTFILE);
+            fclose(RESTARTFILE);
+		    MPI_Barrier(comm);
+        }
+	}
+
+	PROFILE_STOP("Loop");
+	PROFILE_SAVE("lbpm_greyscale_simulator",1);
+	//************************************************************************
+	ScaLBL_DeviceBarrier();
+	MPI_Barrier(comm);
+	stoptime = MPI_Wtime();
+	if (rank==0) printf("-------------------------------------------------------------------\n");
+	// Compute the walltime per timestep
+	cputime = (stoptime - starttime)/timestep;
+	// Performance obtained from each node
+	double MLUPS = double(Np)/cputime/1000000;
+
+	if (rank==0) printf("********************************************************\n");
+	if (rank==0) printf("CPU time = %f \n", cputime);
+	if (rank==0) printf("Lattice update rate (per core)= %f MLUPS \n", MLUPS);
+	MLUPS *= nprocs;
+	if (rank==0) printf("Lattice update rate (total)= %f MLUPS \n", MLUPS);
+	if (rank==0) printf("********************************************************\n");
+
+	// ************************************************************************
+}
+
+void ScaLBL_GreyscaleModel::VelocityField(){
+
+/*	Minkowski Morphology(Mask);
+	int SIZE=Np*sizeof(double);
+	ScaLBL_D3Q19_Momentum(fq,Velocity, Np);
+	ScaLBL_DeviceBarrier(); MPI_Barrier(comm);
+	ScaLBL_CopyToHost(&VELOCITY[0],&Velocity[0],3*SIZE);
+
+	memcpy(Morphology.SDn.data(), Distance.data(), Nx*Ny*Nz*sizeof(double));
+	Morphology.Initialize();
+	Morphology.UpdateMeshValues();
+	Morphology.ComputeLocal();
+	Morphology.Reduce();
+	
+	double count_loc=0;
+	double count;
+	double vax,vay,vaz;
+	double vax_loc,vay_loc,vaz_loc;
+	vax_loc = vay_loc = vaz_loc = 0.f;
+	for (int n=0; n<ScaLBL_Comm->LastExterior(); n++){
+		vax_loc += VELOCITY[n];
+		vay_loc += VELOCITY[Np+n];
+		vaz_loc += VELOCITY[2*Np+n];
+		count_loc+=1.0;
+	}
+	
+	for (int n=ScaLBL_Comm->FirstInterior(); n<ScaLBL_Comm->LastInterior(); n++){
+		vax_loc += VELOCITY[n];
+		vay_loc += VELOCITY[Np+n];
+		vaz_loc += VELOCITY[2*Np+n];
+		count_loc+=1.0;
+	}
+	MPI_Allreduce(&vax_loc,&vax,1,MPI_DOUBLE,MPI_SUM,Mask->Comm);
+	MPI_Allreduce(&vay_loc,&vay,1,MPI_DOUBLE,MPI_SUM,Mask->Comm);
+	MPI_Allreduce(&vaz_loc,&vaz,1,MPI_DOUBLE,MPI_SUM,Mask->Comm);
+	MPI_Allreduce(&count_loc,&count,1,MPI_DOUBLE,MPI_SUM,Mask->Comm);
+	
+	vax /= count;
+	vay /= count;
+	vaz /= count;
+	
+	double mu = (tau-0.5)/3.f;
+	if (rank==0) printf("Fx Fy Fz mu Vs As Js Xs vx vy vz\n");
+	if (rank==0) printf("%.8g %.8g %.8g %.8g %.8g %.8g %.8g %.8g %.8g %.8g %.8g\n",Fx, Fy, Fz, mu, 
+						Morphology.V(),Morphology.A(),Morphology.J(),Morphology.X(),vax,vay,vaz);
+						*/
+	
+	std::vector<IO::MeshDataStruct> visData;
+	fillHalo<double> fillData(Dm->Comm,Dm->rank_info,{Dm->Nx-2,Dm->Ny-2,Dm->Nz-2},{1,1,1},0,1);
+
+	auto VxVar = std::make_shared<IO::Variable>();
+	auto VyVar = std::make_shared<IO::Variable>();
+	auto VzVar = std::make_shared<IO::Variable>();
+	auto SignDistVar = std::make_shared<IO::Variable>();
+	auto PressureVar = std::make_shared<IO::Variable>();
+
+	IO::initialize("","silo","false");
+	// Create the MeshDataStruct	
+	visData.resize(1);
+	visData[0].meshName = "domain";
+	visData[0].mesh = std::make_shared<IO::DomainMesh>( Dm->rank_info,Dm->Nx-2,Dm->Ny-2,Dm->Nz-2,Dm->Lx,Dm->Ly,Dm->Lz );
+	SignDistVar->name = "SignDist";
+	SignDistVar->type = IO::VariableType::VolumeVariable;
+	SignDistVar->dim = 1;
+	SignDistVar->data.resize(Dm->Nx-2,Dm->Ny-2,Dm->Nz-2);
+	visData[0].vars.push_back(SignDistVar);
+	
+	VxVar->name = "Velocity_x";
+	VxVar->type = IO::VariableType::VolumeVariable;
+	VxVar->dim = 1;
+	VxVar->data.resize(Dm->Nx-2,Dm->Ny-2,Dm->Nz-2);
+	visData[0].vars.push_back(VxVar);
+	VyVar->name = "Velocity_y";
+	VyVar->type = IO::VariableType::VolumeVariable;
+	VyVar->dim = 1;
+	VyVar->data.resize(Dm->Nx-2,Dm->Ny-2,Dm->Nz-2);
+	visData[0].vars.push_back(VyVar);
+	VzVar->name = "Velocity_z";
+	VzVar->type = IO::VariableType::VolumeVariable;
+	VzVar->dim = 1;
+	VzVar->data.resize(Dm->Nx-2,Dm->Ny-2,Dm->Nz-2);
+	visData[0].vars.push_back(VzVar);
+	
+	PressureVar->name = "Pressure";
+	PressureVar->type = IO::VariableType::VolumeVariable;
+	PressureVar->dim = 1;
+	PressureVar->data.resize(Dm->Nx-2,Dm->Ny-2,Dm->Nz-2);
+	visData[0].vars.push_back(PressureVar);
+
+	Array<double>& SignData  = visData[0].vars[0]->data;
+	Array<double>& VelxData = visData[0].vars[1]->data;
+	Array<double>& VelyData = visData[0].vars[2]->data;
+	Array<double>& VelzData = visData[0].vars[3]->data;
+	Array<double>& PressureData = visData[0].vars[4]->data;
+	
+    ASSERT(visData[0].vars[0]->name=="SignDist");
+    ASSERT(visData[0].vars[1]->name=="Velocity_x");
+    ASSERT(visData[0].vars[2]->name=="Velocity_y");
+    ASSERT(visData[0].vars[3]->name=="Velocity_z");
+    ASSERT(visData[0].vars[4]->name=="Pressure");
+	
+	ScaLBL_Comm->RegularLayout(Map,&Velocity[0],Velocity_x);
+	ScaLBL_Comm->RegularLayout(Map,&Velocity[Np],Velocity_y);
+	ScaLBL_Comm->RegularLayout(Map,&Velocity[2*Np],Velocity_z);
+	ScaLBL_Comm->RegularLayout(Map,Pressure_dvc,Pressure);
+
+    fillData.copy(SignDist,SignData);
+    fillData.copy(Velocity_x,VelxData);
+    fillData.copy(Velocity_y,VelyData);
+    fillData.copy(Velocity_z,VelzData);
+    fillData.copy(Pressure,PressureData);
+	
+    IO::writeData( timestep, visData, Dm->Comm );
+
+}
+
+void ScaLBL_GreyscaleModel::WriteDebug(){
+	// Copy back final phase indicator field and convert to regular layout
+	DoubleArray PhaseField(Nx,Ny,Nz);
+
+	//ScaLBL_CopyToHost(Porosity.data(), Poros, sizeof(double)*N);
+
+//	FILE *OUTFILE;
+//	sprintf(LocalRankFilename,"Phase.%05i.raw",rank);
+//	OUTFILE = fopen(LocalRankFilename,"wb");
+//	fwrite(PhaseField.data(),8,N,OUTFILE);
+//	fclose(OUTFILE);
+//
+//    ScaLBL_Comm->RegularLayout(Map,&Den[0],PhaseField);
+//	FILE *AFILE;
+//	sprintf(LocalRankFilename,"A.%05i.raw",rank);
+//	AFILE = fopen(LocalRankFilename,"wb");
+//	fwrite(PhaseField.data(),8,N,AFILE);
+//	fclose(AFILE);
+//
+//	ScaLBL_Comm->RegularLayout(Map,&Den[Np],PhaseField);
+//	FILE *BFILE;
+//	sprintf(LocalRankFilename,"B.%05i.raw",rank);
+//	BFILE = fopen(LocalRankFilename,"wb");
+//	fwrite(PhaseField.data(),8,N,BFILE);
+//	fclose(BFILE);
+//
+//	ScaLBL_Comm->RegularLayout(Map,Pressure,PhaseField);
+//	FILE *PFILE;
+//	sprintf(LocalRankFilename,"Pressure.%05i.raw",rank);
+//	PFILE = fopen(LocalRankFilename,"wb");
+//	fwrite(PhaseField.data(),8,N,PFILE);
+//	fclose(PFILE);
+
+	ScaLBL_Comm->RegularLayout(Map,&Velocity[0],PhaseField);
+	FILE *VELX_FILE;
+	sprintf(LocalRankFilename,"Velocity_X.%05i.raw",rank);
+	VELX_FILE = fopen(LocalRankFilename,"wb");
+	fwrite(PhaseField.data(),8,N,VELX_FILE);
+	fclose(VELX_FILE);
+
+	ScaLBL_Comm->RegularLayout(Map,&Velocity[Np],PhaseField);
+	FILE *VELY_FILE;
+	sprintf(LocalRankFilename,"Velocity_Y.%05i.raw",rank);
+	VELY_FILE = fopen(LocalRankFilename,"wb");
+	fwrite(PhaseField.data(),8,N,VELY_FILE);
+	fclose(VELY_FILE);
+
+	ScaLBL_Comm->RegularLayout(Map,&Velocity[2*Np],PhaseField);
+	FILE *VELZ_FILE;
+	sprintf(LocalRankFilename,"Velocity_Z.%05i.raw",rank);
+	VELZ_FILE = fopen(LocalRankFilename,"wb");
+	fwrite(PhaseField.data(),8,N,VELZ_FILE);
+	fclose(VELZ_FILE);
+
+	ScaLBL_Comm->RegularLayout(Map,&Porosity[0],PhaseField);
+	FILE *POROS_FILE;
+	sprintf(LocalRankFilename,"Porosity.%05i.raw",rank);
+	POROS_FILE = fopen(LocalRankFilename,"wb");
+	fwrite(PhaseField.data(),8,N,POROS_FILE);
+	fclose(POROS_FILE);
+
+	ScaLBL_Comm->RegularLayout(Map,&Permeability[0],PhaseField);
+	FILE *PERM_FILE;
+	sprintf(LocalRankFilename,"Permeability.%05i.raw",rank);
+	PERM_FILE = fopen(LocalRankFilename,"wb");
+	fwrite(PhaseField.data(),8,N,PERM_FILE);
+	fclose(PERM_FILE);
+}
diff --git a/models/GreyscaleModel.h b/models/GreyscaleModel.h
new file mode 100644
index 00000000..a99925b1
--- /dev/null
+++ b/models/GreyscaleModel.h
@@ -0,0 +1,92 @@
+/*
+Implementation of color lattice boltzmann model
+ */
+#include <stdio.h>
+#include <stdlib.h>
+#include <sys/stat.h>
+#include <iostream>
+#include <exception>
+#include <stdexcept>
+#include <fstream>
+
+#include "common/Communication.h"
+#include "common/MPI.h"
+#include "common/Database.h"
+#include "common/ScaLBL.h"
+#include "ProfilerApp.h"
+#include "threadpool/thread_pool.h"
+
+class ScaLBL_GreyscaleModel{
+public:
+	ScaLBL_GreyscaleModel(int RANK, int NP, MPI_Comm COMM);
+	~ScaLBL_GreyscaleModel();	
+	
+	// functions in they should be run
+	void ReadParams(string filename);
+	void ReadParams(std::shared_ptr<Database> db0);
+	void SetDomain();
+	void ReadInput();
+	void Create();
+	void Initialize();
+	void Run();
+	void WriteDebug();
+	void VelocityField();
+	
+	bool Restart,pBC;
+	int timestep,timestepMax;
+	int BoundaryCondition;
+    int CollisionType;
+	double tau;
+    double tau_eff;
+    double Den;//constant density
+	double tolerance;
+	double Fx,Fy,Fz,flux;
+	double din,dout;
+    double dp;//solid particle diameter, unit in voxel
+    double GreyPorosity;
+	
+	int Nx,Ny,Nz,N,Np;
+	int rank,nprocx,nprocy,nprocz,nprocs;
+	double Lx,Ly,Lz;
+
+	std::shared_ptr<Domain> Dm;   // this domain is for analysis
+	std::shared_ptr<Domain> Mask; // this domain is for lbm
+	std::shared_ptr<ScaLBL_Communicator> ScaLBL_Comm;
+    
+    // input database
+    std::shared_ptr<Database> db;
+    std::shared_ptr<Database> domain_db;
+    std::shared_ptr<Database> greyscale_db;
+    std::shared_ptr<Database> analysis_db;
+    std::shared_ptr<Database> vis_db;
+
+    signed char *id;    
+	int *NeighborList;
+	int *dvcMap;
+	double *fq;
+	double *Permeability;//grey voxel permeability
+	double *Porosity;
+	double *Velocity;
+	double *Pressure_dvc;
+    IntArray Map;
+    DoubleArray SignDist;
+    DoubleArray Velocity_x;
+    DoubleArray Velocity_y;
+    DoubleArray Velocity_z;
+    DoubleArray PorosityMap;
+    DoubleArray Pressure;
+		
+private:
+	MPI_Comm comm;
+    
+	int dist_mem_size;
+	int neighborSize;
+	// filenames
+    char LocalRankString[8];
+    char LocalRankFilename[40];
+    char LocalRestartFile[40];
+   
+    void AssignComponentLabels(double *Porosity, double *Permeablity);
+    
+};
+
diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt
index 8d600321..8b14a9dc 100755
--- a/tests/CMakeLists.txt
+++ b/tests/CMakeLists.txt
@@ -3,6 +3,7 @@
 #ADD_LBPM_EXECUTABLE( lbpm_nondarcy_simulator )
 ADD_LBPM_EXECUTABLE( lbpm_color_simulator )
 ADD_LBPM_EXECUTABLE( lbpm_permeability_simulator )
+ADD_LBPM_EXECUTABLE( lbpm_greyscale_simulator )
 #ADD_LBPM_EXECUTABLE( lbpm_BGK_simulator )
 #ADD_LBPM_EXECUTABLE( lbpm_color_macro_simulator )
 ADD_LBPM_EXECUTABLE( lbpm_dfh_simulator )
diff --git a/tests/lbpm_greyscale_simulator.cpp b/tests/lbpm_greyscale_simulator.cpp
new file mode 100644
index 00000000..a54b6fc4
--- /dev/null
+++ b/tests/lbpm_greyscale_simulator.cpp
@@ -0,0 +1,64 @@
+#include <stdio.h>
+#include <stdlib.h>
+#include <sys/stat.h>
+#include <iostream>
+#include <exception>
+#include <stdexcept>
+#include <fstream>
+
+#include "common/ScaLBL.h"
+#include "common/Communication.h"
+#include "common/MPI.h"
+#include "models/GreyscaleModel.h"
+//#define WRITE_SURFACES
+
+/*
+ * Simulator for two-phase flow in porous media
+ * James E. McClure 2013-2014
+ */
+
+using namespace std;
+
+
+int main(int argc, char **argv)
+{
+	//*****************************************
+	// ***** MPI STUFF ****************
+	//*****************************************
+	// Initialize MPI
+	int rank,nprocs;
+	MPI_Init(&argc,&argv);
+	MPI_Comm comm = MPI_COMM_WORLD;
+	MPI_Comm_rank(comm,&rank);
+	MPI_Comm_size(comm,&nprocs);
+	{
+		// parallel domain size (# of sub-domains)
+		int nprocx,nprocy,nprocz;
+		int iproc,jproc,kproc;
+
+		if (rank == 0){
+			printf("********************************************************\n");
+			printf("Running Greyscale Single Phase Permeability Calculation \n");
+			printf("********************************************************\n");
+		}
+		// Initialize compute device
+		int device=ScaLBL_SetDevice(rank);
+		ScaLBL_DeviceBarrier();
+		MPI_Barrier(comm);
+		
+		ScaLBL_GreyscaleModel Greyscale(rank,nprocs,comm);
+		auto filename = argv[1];
+		Greyscale.ReadParams(filename);
+		Greyscale.SetDomain();    // this reads in the domain 
+		Greyscale.ReadInput();
+		Greyscale.Create();       // creating the model will create data structure to match the pore structure and allocate variables
+		Greyscale.Initialize();   // initializing the model will set initial conditions for variables
+		Greyscale.Run();	 
+		//Greyscale.VelocityField();
+		//Greyscale.WriteDebug();
+	}
+	// ****************************************************
+	MPI_Barrier(comm);
+	MPI_Finalize();
+	// ****************************************************
+}