diff --git a/tests/lbpm_random_pp.cpp b/tests/lbpm_random_pp.cpp
index e565f494..c44bece9 100644
--- a/tests/lbpm_random_pp.cpp
+++ b/tests/lbpm_random_pp.cpp
@@ -20,6 +20,34 @@ inline void FlipID(char *ID, int N)
 	}
 }
 
+//*************************************************************************
+// Implementation of Two-Phase Immiscible LBM using CUDA
+//*************************************************************************
+inline void PackID(int *list, int count, char *sendbuf, char *ID){
+	// Fill in the phase ID values from neighboring processors
+	// This packs up the values that need to be sent from one processor to another
+	int idx,n;
+
+	for (idx=0; idx<count; idx++){
+		n = list[idx];
+		sendbuf[idx] = ID[n];
+	}
+}
+//***************************************************************************************
+
+inline void UnpackID(int *list, int count, char *recvbuf, char *ID){
+	// Fill in the phase ID values from neighboring processors
+	// This unpacks the values once they have been recieved from neighbors
+	int idx,n;
+
+	for (idx=0; idx<count; idx++){
+		n = list[idx];
+		ID[n] = recvbuf[idx];
+	}
+}
+
+//***************************************************************************************
+
 
 int main(int argc, char **argv)
 {
@@ -36,38 +64,38 @@ int main(int argc, char **argv)
 	//sscanf(argv[2],"%d",&InitialWetting);
 	Saturation=strtod(argv[1],NULL);
 	InitialWetting=atoi(argv[2]);
-		if (rank==0){
-			printf("Initializing wetting phase saturation of %f \n",Saturation);
-			if (InitialWetting == 1)
-				printf("Begin from connected wetting phase \n");
-			else
-				printf("Begin from connected non-wetting phase \n");
-		}
-		//	}
+	if (rank==0){
+		printf("Initializing wetting phase saturation of %f \n",Saturation);
+		if (InitialWetting == 1)
+			printf("Begin from connected wetting phase \n");
+		else
+			printf("Begin from connected non-wetting phase \n");
+	}
+	//	}
 
 	if (InitialWetting == 1)	Saturation=1.0-Saturation;
-    //.......................................................................
-    // Reading the domain information file
-    //.......................................................................
-    int nprocx, nprocy, nprocz, nx, ny, nz, nspheres;
-    double Lx, Ly, Lz;
-    int i,j,k,n;
+	//.......................................................................
+	// Reading the domain information file
+	//.......................................................................
+	int nprocx, nprocy, nprocz, nx, ny, nz, nspheres;
+	double Lx, Ly, Lz;
+	int i,j,k,n;
 	int BC=0;
 
-    if (rank==0){
-    	ifstream domain("Domain.in");
-    	domain >> nprocx;
-    	domain >> nprocy;
-    	domain >> nprocz;
-    	domain >> nx;
-    	domain >> ny;
-    	domain >> nz;
-    	domain >> nspheres;
-    	domain >> Lx;
-    	domain >> Ly;
-    	domain >> Lz;
+	if (rank==0){
+		ifstream domain("Domain.in");
+		domain >> nprocx;
+		domain >> nprocy;
+		domain >> nprocz;
+		domain >> nx;
+		domain >> ny;
+		domain >> nz;
+		domain >> nspheres;
+		domain >> Lx;
+		domain >> Ly;
+		domain >> Lz;
 
-    }
+	}
 	MPI_Barrier(MPI_COMM_WORLD);
 	// Computational domain
 	MPI_Bcast(&nx,1,MPI_INT,0,MPI_COMM_WORLD);
@@ -83,60 +111,72 @@ int main(int argc, char **argv)
 	//.................................................
 	MPI_Barrier(MPI_COMM_WORLD);
 
-    // Check that the number of processors >= the number of ranks
-    if ( rank==0 ) {
-        printf("Number of MPI ranks required: %i \n", nprocx*nprocy*nprocz);
-        printf("Number of MPI ranks used: %i \n", nprocs);
-        printf("Full domain size: %i x %i x %i  \n",nx*nprocx,ny*nprocy,nz*nprocz);
-    }
-    if ( nprocs < nprocx*nprocy*nprocz ){
-        ERROR("Insufficient number of processors");
-    }
+	// Check that the number of processors >= the number of ranks
+	if ( rank==0 ) {
+		printf("Number of MPI ranks required: %i \n", nprocx*nprocy*nprocz);
+		printf("Number of MPI ranks used: %i \n", nprocs);
+		printf("Full domain size: %i x %i x %i  \n",nx*nprocx,ny*nprocy,nz*nprocz);
+	}
+	if ( nprocs < nprocx*nprocy*nprocz ){
+		ERROR("Insufficient number of processors");
+	}
 
 	char LocalRankFilename[40];
 
 	int BoundaryCondition=0;
 	Domain Dm(nx,ny,nz,rank,nprocx,nprocy,nprocz,Lx,Ly,Lz,BoundaryCondition);
+	for
 
 	nx+=2; ny+=2; nz+=2;
 	int N = nx*ny*nz;
-    char *id;
-    id = new char[N];
+	char *id;
+	id = new char[N];
 
-    DoubleArray SignDist(nx,ny,nz);
-    // Read the signed distance from file
-    sprintf(LocalRankFilename,"SignDist.%05i",rank);
-    FILE *DIST = fopen(LocalRankFilename,"rb");
-    fread(SignDist.get(),8,N,DIST);
-    fclose(DIST);
+	// Define communication sub-domain -- everywhere
+	for (int k=0; k<nz; k++){
+		for (int j=0; j<ny; j++){
+			for (int i=0; i<nx; i++){
+				n = k*nx*ny+j*nx+i;
+				Dm.id[n] = 1;
+			}
+		}
+	}
+	Dm.CommInit(MPI_COMM_WORLD);
+
+	DoubleArray SignDist(nx,ny,nz);
+	// Read the signed distance from file
+	sprintf(LocalRankFilename,"SignDist.%05i",rank);
+	FILE *DIST = fopen(LocalRankFilename,"rb");
+	fread(SignDist.get(),8,N,DIST);
+	fclose(DIST);
 
 	int count,countGlobal,totalGlobal;
-    count = 0;
-    for (int k=0; k<nz; k++){
-        for (int j=0; j<ny; j++){
-            for (int i=0; i<nx; i++){
-            	n = k*nx*ny+j*nx+i;
-            	if (SignDist(i,j,k) < 0.0)  id[n] = 0;
-            	else{
-            		id[n] = 2;
-            		count++;
-            	}
-            }
-        }
-    }
+	count = 0;
+	for (int k=0; k<nz; k++){
+		for (int j=0; j<ny; j++){
+			for (int i=0; i<nx; i++){
+				n = k*nx*ny+j*nx+i;
+				if (SignDist(i,j,k) < 0.0)  id[n] = 0;
+				else{
+					id[n] = 2;
+					count++;
+				}
+			}
+		}
+	}
 	// total Global is the number of nodes in the pore-space
 	MPI_Allreduce(&count,&totalGlobal,1,MPI_INT,MPI_SUM,MPI_COMM_WORLD);
 
-    Dm.CommInit(MPI_COMM_WORLD);
-    int iproc = Dm.iproc;
-    int jproc = Dm.jproc;
-    int kproc = Dm.kproc;
+	Dm.CommInit(MPI_COMM_WORLD);
+	int iproc = Dm.iproc;
+	int jproc = Dm.jproc;
+	int kproc = Dm.kproc;
 
 	int bin, binCount;
 	ifstream Dist("BlobSize.in");
 	Dist >> binCount;
 	int *SizeX, *SizeY, *SizeZ;
-//	printf("Number of blob sizes: %i \n",binCount);
+	//	printf("Number of blob sizes: %i \n",binCount);
 	SizeX = new int [binCount];
 	SizeY = new int [binCount];
 	SizeZ = new int [binCount];
@@ -144,7 +184,7 @@ int main(int argc, char **argv)
 		Dist >> SizeX[bin];
 		Dist >> SizeY[bin];
 		Dist >> SizeZ[bin];
-	//	printf("Blob %i dimension: %i x %i x %i \n",bin, SizeX[bin], SizeY[bin], SizeZ[bin]);
+		//	printf("Blob %i dimension: %i x %i x %i \n",bin, SizeX[bin], SizeY[bin], SizeZ[bin]);
 	}
 	Dist.close();
 
@@ -219,7 +259,7 @@ int main(int argc, char **argv)
 		for (int k=1; k<Nz-1; k++){
 			for (int j=1; j<Ny-1; j++){
 				for (int i=1; i<Nx-1; i++){
-				  n=k*Nx*Ny+j*Nx+i;
+					n=k*Nx*Ny+j*Nx+i;
 					if (id[n] == 1){
 						count++;
 					}
@@ -234,28 +274,151 @@ int main(int argc, char **argv)
 
 	if (InitialWetting == 1)	FlipID(id,nx*ny*nz);
 
+	// Fill in the phase ID from neighboring processors
+	char *sendID_x, *sendID_y, *sendID_z, *sendID_X, *sendID_Y, *sendID_Z;
+	char *sendID_xy, *sendID_yz, *sendID_xz, *sendID_Xy, *sendID_Yz, *sendID_xZ;
+	char *sendID_xY, *sendID_yZ, *sendID_Xz, *sendID_XY, *sendID_YZ, *sendID_XZ;
+	char *recvID_x, *recvID_y, *recvID_z, *recvID_X, *recvID_Y, *recvID_Z;
+	char *recvID_xy, *recvID_yz, *recvID_xz, *recvID_Xy, *recvID_Yz, *recvID_xZ;
+	char *recvID_xY, *recvID_yZ, *recvID_Xz, *recvID_XY, *recvID_YZ, *recvID_XZ;
+	// send buffers
+	sendID_x = new char [Dm.sendCount_x];
+	sendID_y = new char [Dm.sendCount_y];
+	sendID_z = new char [Dm.sendCount_z];
+	sendID_X = new char [Dm.sendCount_X];
+	sendID_Y = new char [Dm.sendCount_Y];
+	sendID_Z = new char [Dm.sendCount_Z];
+	sendID_xy = new char [Dm.sendCount_xy];
+	sendID_yz = new char [Dm.sendCount_yz];
+	sendID_xz = new char [Dm.sendCount_xz];
+	sendID_Xy = new char [Dm.sendCount_Xy];
+	sendID_Yz = new char [Dm.sendCount_Yz];
+	sendID_xZ = new char [Dm.sendCount_xZ];
+	sendID_xY = new char [Dm.sendCount_xY];
+	sendID_yZ = new char [Dm.sendCount_yZ];
+	sendID_Xz = new char [Dm.sendCount_Xz];
+	sendID_XY = new char [Dm.sendCount_XY];
+	sendID_YZ = new char [Dm.sendCount_YZ];
+	sendID_XZ = new char [Dm.sendCount_XZ];
+	//......................................................................................
+	// recv buffers
+	recvID_x = new char [recvCount_x];
+	recvID_y = new char [recvCount_y];
+	recvID_z = new char [recvCount_z];
+	recvID_X = new char [recvCount_X];
+	recvID_Y = new char [recvCount_Y];
+	recvID_Z = new char [recvCount_Z];
+	recvID_xy = new char [recvCount_xy];
+	recvID_yz = new char [recvCount_yz];
+	recvID_xz = new char [recvCount_xz];
+	recvID_Xy = new char [recvCount_Xy];
+	recvID_xZ = new char [recvCount_xZ];
+	recvID_xY = new char [recvCount_xY];
+	recvID_yZ = new char [recvCount_yZ];
+	recvID_Yz = new char [recvCount_Yz];
+	recvID_Xz = new char [recvCount_Xz];
+	recvID_XY = new char [recvCount_XY];
+	recvID_YZ = new char [recvCount_YZ];
+	recvID_XZ = new char [recvCount_XZ];
+	//......................................................................................
+	sendtag = recvtag = 7;
+	PackID(Dm.sendList_x, Dm.sendCount_x ,sendID_x, id);
+	PackID(Dm.sendList_X, Dm.sendCount_X ,sendID_X, id);
+	PackID(Dm.sendList_y, Dm.sendCount_y ,sendID_y, id);
+	PackID(Dm.sendList_Y, Dm.sendCount_Y ,sendID_Y, id);
+	PackID(Dm.sendList_z, Dm.sendCount_z ,sendID_z, id);
+	PackID(Dm.sendList_Z, Dm.sendCount_Z ,sendID_Z, id);
+	PackID(Dm.sendList_xy, Dm.sendCount_xy ,sendID_xy, id);
+	PackID(Dm.sendList_Xy, Dm.sendCount_Xy ,sendID_Xy, id);
+	PackID(Dm.sendList_xY, Dm.sendCount_xY ,sendID_xY, id);
+	PackID(Dm.sendList_XY, Dm.sendCount_XY ,sendID_XY, id);
+	PackID(Dm.sendList_xz, Dm.sendCount_xz ,sendID_xz, id);
+	PackID(Dm.sendList_Xz, Dm.sendCount_Xz ,sendID_Xz, id);
+	PackID(Dm.sendList_xZ, Dm.sendCount_xZ ,sendID_xZ, id);
+	PackID(Dm.sendList_XZ, Dm.sendCount_XZ ,sendID_XZ, id);
+	PackID(Dm.sendList_yz, Dm.sendCount_yz ,sendID_yz, id);
+	PackID(Dm.sendList_Yz, Dm.sendCount_Yz ,sendID_Yz, id);
+	PackID(Dm.sendList_yZ, Dm.sendCount_yZ ,sendID_yZ, id);
+	PackID(Dm.sendList_YZ, Dm.sendCount_YZ ,sendID_YZ, id);
+	//......................................................................................
+	MPI_Sendrecv(sendID_x,Dm.sendCount_x,MPI_CHAR,Dm.rank_x,sendtag,
+			recvID_X,recvCount_X,MPI_CHAR,Dm.rank_X,recvtag,MPI_COMM_WORLD,MPI_STATUS_IGNORE);
+	MPI_Sendrecv(sendID_X,Dm.sendCount_X,MPI_CHAR,Dm.rank_X,sendtag,
+			recvID_x,recvCount_x,MPI_CHAR,Dm.rank_x,recvtag,MPI_COMM_WORLD,MPI_STATUS_IGNORE);
+	MPI_Sendrecv(sendID_y,Dm.sendCount_y,MPI_CHAR,Dm.rank_y,sendtag,
+			recvID_Y,recvCount_Y,MPI_CHAR,Dm.rank_Y,recvtag,MPI_COMM_WORLD,MPI_STATUS_IGNORE);
+	MPI_Sendrecv(sendID_Y,Dm.sendCount_Y,MPI_CHAR,Dm.rank_Y,sendtag,
+			recvID_y,recvCount_y,MPI_CHAR,Dm.rank_y,recvtag,MPI_COMM_WORLD,MPI_STATUS_IGNORE);
+	MPI_Sendrecv(sendID_z,Dm.sendCount_z,MPI_CHAR,Dm.rank_z,sendtag,
+			recvID_Z,recvCount_Z,MPI_CHAR,Dm.rank_Z,recvtag,MPI_COMM_WORLD,MPI_STATUS_IGNORE);
+	MPI_Sendrecv(sendID_Z,Dm.sendCount_Z,MPI_CHAR,Dm.rank_Z,sendtag,
+			recvID_z,recvCount_z,MPI_CHAR,Dm.rank_z,recvtag,MPI_COMM_WORLD,MPI_STATUS_IGNORE);
+	MPI_Sendrecv(sendID_xy,Dm.sendCount_xy,MPI_CHAR,Dm.rank_xy,sendtag,
+			recvID_XY,recvCount_XY,MPI_CHAR,Dm.rank_XY,recvtag,MPI_COMM_WORLD,MPI_STATUS_IGNORE);
+	MPI_Sendrecv(sendID_XY,Dm.sendCount_XY,MPI_CHAR,Dm.rank_XY,sendtag,
+			recvID_xy,recvCount_xy,MPI_CHAR,Dm.rank_xy,recvtag,MPI_COMM_WORLD,MPI_STATUS_IGNORE);
+	MPI_Sendrecv(sendID_Xy,Dm.sendCount_Xy,MPI_CHAR,Dm.rank_Xy,sendtag,
+			recvID_xY,recvCount_xY,MPI_CHAR,Dm.rank_xY,recvtag,MPI_COMM_WORLD,MPI_STATUS_IGNORE);
+	MPI_Sendrecv(sendID_xY,Dm.sendCount_xY,MPI_CHAR,Dm.rank_xY,sendtag,
+			recvID_Xy,recvCount_Xy,MPI_CHAR,Dm.rank_Xy,recvtag,MPI_COMM_WORLD,MPI_STATUS_IGNORE);
+	MPI_Sendrecv(sendID_xz,Dm.sendCount_xz,MPI_CHAR,Dm.rank_xz,sendtag,
+			recvID_XZ,recvCount_XZ,MPI_CHAR,Dm.rank_XZ,recvtag,MPI_COMM_WORLD,MPI_STATUS_IGNORE);
+	MPI_Sendrecv(sendID_XZ,Dm.sendCount_XZ,MPI_CHAR,Dm.rank_XZ,sendtag,
+			recvID_xz,recvCount_xz,MPI_CHAR,Dm.rank_xz,recvtag,MPI_COMM_WORLD,MPI_STATUS_IGNORE);
+	MPI_Sendrecv(sendID_Xz,Dm.sendCount_Xz,MPI_CHAR,Dm.rank_Xz,sendtag,
+			recvID_xZ,recvCount_xZ,MPI_CHAR,Dm.rank_xZ,recvtag,MPI_COMM_WORLD,MPI_STATUS_IGNORE);
+	MPI_Sendrecv(sendID_xZ,Dm.sendCount_xZ,MPI_CHAR,Dm.rank_xZ,sendtag,
+			recvID_Xz,recvCount_Xz,MPI_CHAR,Dm.rank_Xz,recvtag,MPI_COMM_WORLD,MPI_STATUS_IGNORE);
+	MPI_Sendrecv(sendID_yz,Dm.sendCount_yz,MPI_CHAR,Dm.rank_yz,sendtag,
+			recvID_YZ,recvCount_YZ,MPI_CHAR,Dm.rank_YZ,recvtag,MPI_COMM_WORLD,MPI_STATUS_IGNORE);
+	MPI_Sendrecv(sendID_YZ,Dm.sendCount_YZ,MPI_CHAR,Dm.rank_YZ,sendtag,
+			recvID_yz,recvCount_yz,MPI_CHAR,Dm.rank_yz,recvtag,MPI_COMM_WORLD,MPI_STATUS_IGNORE);
+	MPI_Sendrecv(sendID_Yz,Dm.sendCount_Yz,MPI_CHAR,Dm.rank_Yz,sendtag,
+			recvID_yZ,recvCount_yZ,MPI_CHAR,Dm.rank_yZ,recvtag,MPI_COMM_WORLD,MPI_STATUS_IGNORE);
+	MPI_Sendrecv(sendID_yZ,Dm.sendCount_yZ,MPI_CHAR,Dm.rank_yZ,sendtag,
+			recvID_Yz,recvCount_Yz,MPI_CHAR,Dm.rank_Yz,recvtag,MPI_COMM_WORLD,MPI_STATUS_IGNORE);
+	//......................................................................................
+	UnpackID(Dm.recvList_x, Dm.recvCount_x ,recvID_x, id);
+	UnpackID(Dm.recvList_X, Dm.recvCount_X ,recvID_X, id);
+	UnpackID(Dm.recvList_y, Dm.recvCount_y ,recvID_y, id);
+	UnpackID(Dm.recvList_Y, Dm.recvCount_Y ,recvID_Y, id);
+	UnpackID(Dm.recvList_z, Dm.recvCount_z ,recvID_z, id);
+	UnpackID(Dm.recvList_Z, Dm.recvCount_Z ,recvID_Z, id);
+	UnpackID(Dm.recvList_xy, Dm.recvCount_xy ,recvID_xy, id);
+	UnpackID(Dm.recvList_Xy, Dm.recvCount_Xy ,recvID_Xy, id);
+	UnpackID(Dm.recvList_xY, Dm.recvCount_xY ,recvID_xY, id);
+	UnpackID(Dm.recvList_XY, Dm.recvCount_XY ,recvID_XY, id);
+	UnpackID(Dm.recvList_xz, Dm.recvCount_xz ,recvID_xz, id);
+	UnpackID(Dm.recvList_Xz, Dm.recvCount_Xz ,recvID_Xz, id);
+	UnpackID(Dm.recvList_xZ, Dm.recvCount_xZ ,recvID_xZ, id);
+	UnpackID(Dm.recvList_XZ, Dm.recvCount_XZ ,recvID_XZ, id);
+	UnpackID(Dm.recvList_yz, Dm.recvCount_yz ,recvID_yz, id);
+	UnpackID(Dm.recvList_Yz, Dm.recvCount_Yz ,recvID_Yz, id);
+	UnpackID(Dm.recvList_yZ, Dm.recvCount_yZ ,recvID_yZ, id);
+	UnpackID(Dm.recvList_YZ, Dm.recvCount_YZ ,recvID_YZ, id);
+	//......................................................................................
 
-		count = 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++){
-				  n=k*Nx*Ny+j*Nx+i;
-					if (id[n] == 1){
-						count++;
-					}
+	count = 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++){
+				n=k*Nx*Ny+j*Nx+i;
+				if (id[n] == 1){
+					count++;
 				}
 			}
 		}
-		MPI_Allreduce(&count,&countGlobal,1,MPI_INT,MPI_SUM,MPI_COMM_WORLD);
-		sat = float(countGlobal)/totalGlobal;
+	}
+	MPI_Allreduce(&count,&countGlobal,1,MPI_INT,MPI_SUM,MPI_COMM_WORLD);
+	sat = float(countGlobal)/totalGlobal;
 	if (rank==0) printf("Final saturation=%f\n",sat);
 
-    sprintf(LocalRankFilename,"ID.%05i",rank);
-    FILE *ID = fopen(LocalRankFilename,"wb");
-    fwrite(id,1,N,ID);
-    fclose(ID);
+	sprintf(LocalRankFilename,"ID.%05i",rank);
+	FILE *ID = fopen(LocalRankFilename,"wb");
+	fwrite(id,1,N,ID);
+	fclose(ID);
 
-    MPI_Barrier(MPI_COMM_WORLD);
-    MPI_Finalize();
-    return 0;
+	MPI_Barrier(MPI_COMM_WORLD);
+	MPI_Finalize();
+	return 0;
 }