merged lbpm_color_macro_simulator

analysis/eikonal.hpp

@@ -167,8 +167,10 @@ inline double Eikonal(DoubleArray &Distance, char *ID, Domain &Dm, int timesteps
 		count++;
 
 
-		if (count%50 == 0 && Dm.rank==0 )
+		if (count%50 == 0 && Dm.rank==0 ){
 			printf("Time=%i, Max variation=%f, Global variation=%f \n",count,GlobalMax,GlobalVar);
+			fflush(stdout);
+		}
 
 		if (fabs(GlobalMax) < 1e-5){
 			if (Dm.rank==0) printf("Exiting with max tolerance of 1e-5 \n");

common/ScaLBL.cpp

@@ -1,7 +1,6 @@
 #include "common/Domain.h"
 #include "common/ScaLBL.h"
 
-
 ScaLBL_Communicator::ScaLBL_Communicator(Domain &Dm){
 	//......................................................................................
 	Lock=false; // unlock the communicator
@@ -82,139 +81,139 @@ ScaLBL_Communicator::ScaLBL_Communicator(Domain &Dm){
 	BoundaryCondition = Dm.BoundaryCondition;
 	//......................................................................................
 
-	ScaLBL_AllocateDeviceMemory((void **) &sendbuf_x, 5*sendCount_x*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &sendbuf_X, 5*sendCount_X*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &sendbuf_y, 5*sendCount_y*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &sendbuf_Y, 5*sendCount_Y*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &sendbuf_z, 5*sendCount_z*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &sendbuf_Z, 5*sendCount_Z*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &sendbuf_xy, sendCount_xy*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &sendbuf_xY, sendCount_xY*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &sendbuf_Xy, sendCount_Xy*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &sendbuf_XY, sendCount_XY*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &sendbuf_xz, sendCount_xz*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &sendbuf_xZ, sendCount_xZ*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &sendbuf_Xz, sendCount_Xz*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &sendbuf_XZ, sendCount_XZ*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &sendbuf_yz, sendCount_yz*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &sendbuf_yZ, sendCount_yZ*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &sendbuf_Yz, sendCount_Yz*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &sendbuf_YZ, sendCount_YZ*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &sendbuf_x, 5*sendCount_x*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &sendbuf_X, 5*sendCount_X*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &sendbuf_y, 5*sendCount_y*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &sendbuf_Y, 5*sendCount_Y*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &sendbuf_z, 5*sendCount_z*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &sendbuf_Z, 5*sendCount_Z*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &sendbuf_xy, sendCount_xy*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &sendbuf_xY, sendCount_xY*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &sendbuf_Xy, sendCount_Xy*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &sendbuf_XY, sendCount_XY*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &sendbuf_xz, sendCount_xz*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &sendbuf_xZ, sendCount_xZ*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &sendbuf_Xz, sendCount_Xz*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &sendbuf_XZ, sendCount_XZ*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &sendbuf_yz, sendCount_yz*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &sendbuf_yZ, sendCount_yZ*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &sendbuf_Yz, sendCount_Yz*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &sendbuf_YZ, sendCount_YZ*sizeof(double));	// Allocate device memory
 	//......................................................................................
-	ScaLBL_AllocateDeviceMemory((void **) &recvbuf_x, 5*recvCount_x*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &recvbuf_X, 5*recvCount_X*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &recvbuf_y, 5*recvCount_y*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &recvbuf_Y, 5*recvCount_Y*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &recvbuf_z, 5*recvCount_z*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &recvbuf_Z, 5*recvCount_Z*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &recvbuf_xy, recvCount_xy*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &recvbuf_xY, recvCount_xY*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &recvbuf_Xy, recvCount_Xy*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &recvbuf_XY, recvCount_XY*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &recvbuf_xz, recvCount_xz*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &recvbuf_xZ, recvCount_xZ*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &recvbuf_Xz, recvCount_Xz*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &recvbuf_XZ, recvCount_XZ*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &recvbuf_yz, recvCount_yz*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &recvbuf_yZ, recvCount_yZ*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &recvbuf_Yz, recvCount_Yz*sizeof(double));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &recvbuf_YZ, recvCount_YZ*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &recvbuf_x, 5*recvCount_x*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &recvbuf_X, 5*recvCount_X*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &recvbuf_y, 5*recvCount_y*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &recvbuf_Y, 5*recvCount_Y*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &recvbuf_z, 5*recvCount_z*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &recvbuf_Z, 5*recvCount_Z*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &recvbuf_xy, recvCount_xy*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &recvbuf_xY, recvCount_xY*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &recvbuf_Xy, recvCount_Xy*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &recvbuf_XY, recvCount_XY*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &recvbuf_xz, recvCount_xz*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &recvbuf_xZ, recvCount_xZ*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &recvbuf_Xz, recvCount_Xz*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &recvbuf_XZ, recvCount_XZ*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &recvbuf_yz, recvCount_yz*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &recvbuf_yZ, recvCount_yZ*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &recvbuf_Yz, recvCount_Yz*sizeof(double));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &recvbuf_YZ, recvCount_YZ*sizeof(double));	// Allocate device memory
 	//......................................................................................
-	ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_x, sendCount_x*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_X, sendCount_X*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_y, sendCount_y*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_Y, sendCount_Y*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_z, sendCount_z*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_Z, sendCount_Z*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_xy, sendCount_xy*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_xY, sendCount_xY*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_Xy, sendCount_Xy*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_XY, sendCount_XY*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_xz, sendCount_xz*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_xZ, sendCount_xZ*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_Xz, sendCount_Xz*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_XZ, sendCount_XZ*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_yz, sendCount_yz*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_yZ, sendCount_yZ*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_Yz, sendCount_Yz*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_YZ, sendCount_YZ*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcSendList_x, sendCount_x*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcSendList_X, sendCount_X*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcSendList_y, sendCount_y*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcSendList_Y, sendCount_Y*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcSendList_z, sendCount_z*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcSendList_Z, sendCount_Z*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcSendList_xy, sendCount_xy*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcSendList_xY, sendCount_xY*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcSendList_Xy, sendCount_Xy*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcSendList_XY, sendCount_XY*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcSendList_xz, sendCount_xz*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcSendList_xZ, sendCount_xZ*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcSendList_Xz, sendCount_Xz*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcSendList_XZ, sendCount_XZ*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcSendList_yz, sendCount_yz*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcSendList_yZ, sendCount_yZ*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcSendList_Yz, sendCount_Yz*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcSendList_YZ, sendCount_YZ*sizeof(int));	// Allocate device memory
 	//......................................................................................
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_x, recvCount_x*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_X, recvCount_X*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_y, recvCount_y*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_Y, recvCount_Y*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_z, recvCount_z*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_Z, recvCount_Z*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_xy, recvCount_xy*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_xY, recvCount_xY*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_Xy, recvCount_Xy*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_XY, recvCount_XY*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_xz, recvCount_xz*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_xZ, recvCount_xZ*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_Xz, recvCount_Xz*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_XZ, recvCount_XZ*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_yz, recvCount_yz*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_yZ, recvCount_yZ*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_Yz, recvCount_Yz*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_YZ, recvCount_YZ*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvList_x, recvCount_x*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvList_X, recvCount_X*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvList_y, recvCount_y*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvList_Y, recvCount_Y*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvList_z, recvCount_z*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvList_Z, recvCount_Z*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvList_xy, recvCount_xy*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvList_xY, recvCount_xY*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvList_Xy, recvCount_Xy*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvList_XY, recvCount_XY*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvList_xz, recvCount_xz*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvList_xZ, recvCount_xZ*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvList_Xz, recvCount_Xz*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvList_XZ, recvCount_XZ*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvList_yz, recvCount_yz*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvList_yZ, recvCount_yZ*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvList_Yz, recvCount_Yz*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvList_YZ, recvCount_YZ*sizeof(int));	// Allocate device memory
 	//......................................................................................
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvDist_x, 5*recvCount_x*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvDist_X, 5*recvCount_X*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvDist_y, 5*recvCount_y*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvDist_Y, 5*recvCount_Y*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvDist_z, 5*recvCount_z*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvDist_Z, 5*recvCount_Z*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvDist_xy, recvCount_xy*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvDist_xY, recvCount_xY*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvDist_Xy, recvCount_Xy*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvDist_XY, recvCount_XY*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvDist_xz, recvCount_xz*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvDist_xZ, recvCount_xZ*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvDist_Xz, recvCount_Xz*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvDist_XZ, recvCount_XZ*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvDist_yz, recvCount_yz*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvDist_yZ, recvCount_yZ*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvDist_Yz, recvCount_Yz*sizeof(int));	// Allocate device memory
-	ScaLBL_AllocateDeviceMemory((void **) &dvcRecvDist_YZ, recvCount_YZ*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvDist_x, 5*recvCount_x*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvDist_X, 5*recvCount_X*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvDist_y, 5*recvCount_y*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvDist_Y, 5*recvCount_Y*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvDist_z, 5*recvCount_z*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvDist_Z, 5*recvCount_Z*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvDist_xy, recvCount_xy*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvDist_xY, recvCount_xY*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvDist_Xy, recvCount_Xy*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvDist_XY, recvCount_XY*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvDist_xz, recvCount_xz*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvDist_xZ, recvCount_xZ*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvDist_Xz, recvCount_Xz*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvDist_XZ, recvCount_XZ*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvDist_yz, recvCount_yz*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvDist_yZ, recvCount_yZ*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvDist_Yz, recvCount_Yz*sizeof(int));	// Allocate device memory
+	ScaLBL_AllocateZeroCopy((void **) &dvcRecvDist_YZ, recvCount_YZ*sizeof(int));	// Allocate device memory
 	//......................................................................................
 
-	ScaLBL_CopyToDevice(dvcSendList_x,Dm.sendList_x,sendCount_x*sizeof(int));
-	ScaLBL_CopyToDevice(dvcSendList_X,Dm.sendList_X,sendCount_X*sizeof(int));
-	ScaLBL_CopyToDevice(dvcSendList_y,Dm.sendList_y,sendCount_y*sizeof(int));
-	ScaLBL_CopyToDevice(dvcSendList_Y,Dm.sendList_Y,sendCount_Y*sizeof(int));
-	ScaLBL_CopyToDevice(dvcSendList_z,Dm.sendList_z,sendCount_z*sizeof(int));
-	ScaLBL_CopyToDevice(dvcSendList_Z,Dm.sendList_Z,sendCount_Z*sizeof(int));
-	ScaLBL_CopyToDevice(dvcSendList_xy,Dm.sendList_xy,sendCount_xy*sizeof(int));
-	ScaLBL_CopyToDevice(dvcSendList_XY,Dm.sendList_XY,sendCount_XY*sizeof(int));
-	ScaLBL_CopyToDevice(dvcSendList_xY,Dm.sendList_xY,sendCount_xY*sizeof(int));
-	ScaLBL_CopyToDevice(dvcSendList_Xy,Dm.sendList_Xy,sendCount_Xy*sizeof(int));
-	ScaLBL_CopyToDevice(dvcSendList_xz,Dm.sendList_xz,sendCount_xz*sizeof(int));
-	ScaLBL_CopyToDevice(dvcSendList_XZ,Dm.sendList_XZ,sendCount_XZ*sizeof(int));
-	ScaLBL_CopyToDevice(dvcSendList_xZ,Dm.sendList_xZ,sendCount_xZ*sizeof(int));
-	ScaLBL_CopyToDevice(dvcSendList_Xz,Dm.sendList_Xz,sendCount_Xz*sizeof(int));
-	ScaLBL_CopyToDevice(dvcSendList_yz,Dm.sendList_yz,sendCount_yz*sizeof(int));
-	ScaLBL_CopyToDevice(dvcSendList_YZ,Dm.sendList_YZ,sendCount_YZ*sizeof(int));
-	ScaLBL_CopyToDevice(dvcSendList_yZ,Dm.sendList_yZ,sendCount_yZ*sizeof(int));
-	ScaLBL_CopyToDevice(dvcSendList_Yz,Dm.sendList_Yz,sendCount_Yz*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcSendList_x,Dm.sendList_x,sendCount_x*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcSendList_X,Dm.sendList_X,sendCount_X*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcSendList_y,Dm.sendList_y,sendCount_y*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcSendList_Y,Dm.sendList_Y,sendCount_Y*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcSendList_z,Dm.sendList_z,sendCount_z*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcSendList_Z,Dm.sendList_Z,sendCount_Z*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcSendList_xy,Dm.sendList_xy,sendCount_xy*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcSendList_XY,Dm.sendList_XY,sendCount_XY*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcSendList_xY,Dm.sendList_xY,sendCount_xY*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcSendList_Xy,Dm.sendList_Xy,sendCount_Xy*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcSendList_xz,Dm.sendList_xz,sendCount_xz*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcSendList_XZ,Dm.sendList_XZ,sendCount_XZ*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcSendList_xZ,Dm.sendList_xZ,sendCount_xZ*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcSendList_Xz,Dm.sendList_Xz,sendCount_Xz*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcSendList_yz,Dm.sendList_yz,sendCount_yz*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcSendList_YZ,Dm.sendList_YZ,sendCount_YZ*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcSendList_yZ,Dm.sendList_yZ,sendCount_yZ*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcSendList_Yz,Dm.sendList_Yz,sendCount_Yz*sizeof(int));
 	//......................................................................................
-	ScaLBL_CopyToDevice(dvcRecvList_x,Dm.recvList_x,recvCount_x*sizeof(int));
-	ScaLBL_CopyToDevice(dvcRecvList_X,Dm.recvList_X,recvCount_X*sizeof(int));
-	ScaLBL_CopyToDevice(dvcRecvList_y,Dm.recvList_y,recvCount_y*sizeof(int));
-	ScaLBL_CopyToDevice(dvcRecvList_Y,Dm.recvList_Y,recvCount_Y*sizeof(int));
-	ScaLBL_CopyToDevice(dvcRecvList_z,Dm.recvList_z,recvCount_z*sizeof(int));
-	ScaLBL_CopyToDevice(dvcRecvList_Z,Dm.recvList_Z,recvCount_Z*sizeof(int));
-	ScaLBL_CopyToDevice(dvcRecvList_xy,Dm.recvList_xy,recvCount_xy*sizeof(int));
-	ScaLBL_CopyToDevice(dvcRecvList_XY,Dm.recvList_XY,recvCount_XY*sizeof(int));
-	ScaLBL_CopyToDevice(dvcRecvList_xY,Dm.recvList_xY,recvCount_xY*sizeof(int));
-	ScaLBL_CopyToDevice(dvcRecvList_Xy,Dm.recvList_Xy,recvCount_Xy*sizeof(int));
-	ScaLBL_CopyToDevice(dvcRecvList_xz,Dm.recvList_xz,recvCount_xz*sizeof(int));
-	ScaLBL_CopyToDevice(dvcRecvList_XZ,Dm.recvList_XZ,recvCount_XZ*sizeof(int));
-	ScaLBL_CopyToDevice(dvcRecvList_xZ,Dm.recvList_xZ,recvCount_xZ*sizeof(int));
-	ScaLBL_CopyToDevice(dvcRecvList_Xz,Dm.recvList_Xz,recvCount_Xz*sizeof(int));
-	ScaLBL_CopyToDevice(dvcRecvList_yz,Dm.recvList_yz,recvCount_yz*sizeof(int));
-	ScaLBL_CopyToDevice(dvcRecvList_YZ,Dm.recvList_YZ,recvCount_YZ*sizeof(int));
-	ScaLBL_CopyToDevice(dvcRecvList_yZ,Dm.recvList_yZ,recvCount_yZ*sizeof(int));
-	ScaLBL_CopyToDevice(dvcRecvList_Yz,Dm.recvList_Yz,recvCount_Yz*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcRecvList_x,Dm.recvList_x,recvCount_x*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcRecvList_X,Dm.recvList_X,recvCount_X*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcRecvList_y,Dm.recvList_y,recvCount_y*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcRecvList_Y,Dm.recvList_Y,recvCount_Y*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcRecvList_z,Dm.recvList_z,recvCount_z*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcRecvList_Z,Dm.recvList_Z,recvCount_Z*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcRecvList_xy,Dm.recvList_xy,recvCount_xy*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcRecvList_XY,Dm.recvList_XY,recvCount_XY*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcRecvList_xY,Dm.recvList_xY,recvCount_xY*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcRecvList_Xy,Dm.recvList_Xy,recvCount_Xy*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcRecvList_xz,Dm.recvList_xz,recvCount_xz*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcRecvList_XZ,Dm.recvList_XZ,recvCount_XZ*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcRecvList_xZ,Dm.recvList_xZ,recvCount_xZ*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcRecvList_Xz,Dm.recvList_Xz,recvCount_Xz*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcRecvList_yz,Dm.recvList_yz,recvCount_yz*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcRecvList_YZ,Dm.recvList_YZ,recvCount_YZ*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcRecvList_yZ,Dm.recvList_yZ,recvCount_yZ*sizeof(int));
+	ScaLBL_CopyToZeroCopy(dvcRecvList_Yz,Dm.recvList_Yz,recvCount_Yz*sizeof(int));
 	//......................................................................................
 
 	MPI_Barrier(MPI_COMM_SCALBL);
@@ -2373,7 +2372,7 @@ void ScaLBL_Communicator::MemoryOptimizedLayoutFull(IntArray &Map, int *neighbor
 }
 
 
-void ScaLBL_Communicator::MemoryOptimizedLayoutAA(IntArray &Map, int *neighborList, char *id, int Np){
+int ScaLBL_Communicator::MemoryOptimizedLayoutAA(IntArray &Map, int *neighborList, char *id, int Np){
 	/*
 	 * Generate a memory optimized layout
 	 *   id[n] == 0 implies that site n should be ignored (treat as a mask)
@@ -2398,6 +2397,7 @@ void ScaLBL_Communicator::MemoryOptimizedLayoutAA(IntArray &Map, int *neighborLi
 		}
 	}
 
+	// ********* Exterior **********
 	// Step 1/2: Index the outer walls of the grid only
 	idx=0;	next=0;
 	for (k=1; k<Nz-1; k++){
@@ -2422,6 +2422,10 @@ void ScaLBL_Communicator::MemoryOptimizedLayoutAA(IntArray &Map, int *neighborLi
 	next=idx;
 
 	
+	// ********* Interior **********
+	// align the next read
+	first_interior=(next/16 + 1)*16;
+	idx = first_interior;
 	// Step 2/2: Next loop over the domain interior in block-cyclic fashion
 	for (k=2; k<Nz-2; k++){
 		for (j=2; j<Ny-2; j++){
@@ -2435,6 +2439,7 @@ void ScaLBL_Communicator::MemoryOptimizedLayoutAA(IntArray &Map, int *neighborLi
 			}
 		}
 	}
+	last_interior=idx;
 
 	/*
 	int MemBlockSize=32;
@@ -2469,9 +2474,9 @@ void ScaLBL_Communicator::MemoryOptimizedLayoutAA(IntArray &Map, int *neighborLi
 	}
 
 	*/
-	if (idx > Np ){
-		ERROR("ScaLBL_Communicator::MemoryOptimizedLayout: Failed to create memory efficient layout!\n");
-	}
+//	if (idx > Np ){
+//		ERROR("ScaLBL_Communicator::MemoryOptimizedLayout: Failed to create memory efficient layout!\n");
+//	}
 	/*
 		for (k=1;k<Nz-1;k++){
 			printf("....k=%i .....\n",k);
@@ -2487,6 +2492,8 @@ void ScaLBL_Communicator::MemoryOptimizedLayoutAA(IntArray &Map, int *neighborLi
 		printf("\n\n");
 	 */
 	
+	Np = (last_interior/16 + 1)*16;
+	
 	// Now use Map to determine the neighbors for each lattice direction
 	for (k=1;k<Nz-1;k++){
 		for (j=1;j<Ny-1;j++){
@@ -2963,6 +2970,7 @@ void ScaLBL_Communicator::MemoryOptimizedLayoutAA(IntArray &Map, int *neighborLi
 	
 	// Clean up
 	delete [] TempBuffer;
+	return(Np);
 }
 
 

common/ScaLBL.h

@@ -19,6 +19,10 @@ extern "C" void ScaLBL_CopyToDevice(void* dest, const void* source, size_t size)
 
 extern "C" void ScaLBL_CopyToHost(void* dest, const void* source, size_t size);
 
+extern "C" void ScaLBL_AllocateZeroCopy(void** address, size_t size);
+
+extern "C" void ScaLBL_CopyToZeroCopy(void* dest, const void* source, size_t size);
+
 extern "C" void ScaLBL_DeviceBarrier();
 
 extern "C" void ScaLBL_D3Q19_Pack(int q, int *list, int start, int count, double *sendbuf, double *dist, int N);
@@ -236,6 +240,7 @@ public:
 	unsigned long int CommunicationCount,SendCount,RecvCount;
 	int Nx,Ny,Nz,N;
 	int next;
+	int first_interior,last_interior;
 	int BoundaryCondition;
 	//......................................................................................
 	//  Set up for D319 distributions
@@ -251,7 +256,7 @@ public:
 	double *recvbuf_xY, *recvbuf_yZ, *recvbuf_Xz, *recvbuf_XY, *recvbuf_YZ, *recvbuf_XZ;
 	//......................................................................................
 
-	void MemoryOptimizedLayoutAA(IntArray &Map, int *neighborList, char *id, int Np);
+	int MemoryOptimizedLayoutAA(IntArray &Map, int *neighborList, char *id, int Np);
 	void MemoryOptimizedLayout(IntArray &Map, int *neighborList, char *id, int Np);
 	void MemoryOptimizedLayoutFull(IntArray &Map, int *neighborList, char *id, int Np);
 	void MemoryDenseLayout(IntArray &Map, int *neighborList, char *id, int Np);

gpu/Extras.cu

@@ -22,6 +22,20 @@ extern "C" void ScaLBL_CopyToDevice(void* dest, const void* source, size_t size)
 	}
 }
 
+extern "C" void ScaLBL_AllocateZeroCopy(void** address, size_t size){
+    cudaMallocHost(address,size);
+        cudaError_t err = cudaGetLastError();
+        if (cudaSuccess != err){
+           printf("Error in cudaMallocHost: %s \n",cudaGetErrorString(err));
+        }
+}
+
+extern "C" void ScaLBL_CopyToZeroCopy(void* dest, const void* source, size_t size){
+        //cudaMemcpy(dest,source,size,cudaMemcpyHostToDevice);
+        //cudaError_t err = cudaGetLastError();
+        memcpy(dest, source, size);
+
+}
 
 extern "C" void ScaLBL_CopyToHost(void* dest, const void* source, size_t size){
 	cudaMemcpy(dest,source,size,cudaMemcpyDeviceToHost);

tests/CMakeLists.txt

@@ -12,6 +12,7 @@ ADD_LBPM_EXECUTABLE( lbpm_morphdrain_pp )
 ADD_LBPM_EXECUTABLE( lbpm_morphopen_pp )
 ADD_LBPM_EXECUTABLE( lbpm_morph_pp )
 ADD_LBPM_EXECUTABLE( lbpm_segmented_pp )
+ADD_LBPM_EXECUTABLE( lbpm_block_pp )
 ADD_LBPM_EXECUTABLE( lbpm_segmented_decomp )
 ADD_LBPM_EXECUTABLE( lbpm_serial_decomp )
 ADD_LBPM_EXECUTABLE( lbpm_disc_pp )
@@ -29,6 +30,7 @@ ADD_LBPM_EXECUTABLE( BlobAnalysis )
 ADD_LBPM_EXECUTABLE( BlobIdentify )
 ADD_LBPM_EXECUTABLE( BlobIdentifyParallel )
 ADD_LBPM_EXECUTABLE( convertIO )
+ADD_LBPM_EXECUTABLE( DataAggregator )
 #ADD_LBPM_EXECUTABLE( BlobAnalyzeParallel )
 
 CONFIGURE_FILE( ${CMAKE_CURRENT_SOURCE_DIR}/cylindertest ${CMAKE_CURRENT_BINARY_DIR}/cylindertest COPYONLY )

tests/DataAggregator.cpp

@@ -49,6 +49,9 @@ int main(int argc, char **argv){
 	Bx=NX/Nx+1;
 	By=NY/Ny+1;
 	Bz=NZ/Nz+1;
+	if (Bx>8) Bx=8;
+	if (By>8) By=8;
+	if (Bz>8) Bz=8;
 
 	printf("System size (output) is: %i x %i x %i \n",NX,NY,NZ);
 	printf("Block size (read) is: %i x %i x %i \n",Nx,Ny,Nz);
@@ -78,9 +81,7 @@ int main(int argc, char **argv){
 			for (unsigned int bx=0; bx<Bx; bx++){
 				sprintf(sx,"%d",bx);
 				sprintf(sy,"%d",by);
-				sprintf(sz,"%d",bz);
-				//sprintf(LocalRankFilename,"%s%s%s%s%s%s%s","a1_x",sx,"_y",sy,"_z",sz,".gbd");
-				sprintf(LocalRankFilename,"%s%s%s%s%s%s%s","dis_",sx,"x_",sy,"y_",sz,"z.gbd");
+				sprintf(LocalRankFilename,"%s%s%s%s%s%s%s","a2_x",sx,"_y",sy,"_z",sz,".gbd");
 				printf("Reading file: %s \n", LocalRankFilename);
 				fflush(stdout);
 

tests/lbpm_BGK_simulator.cpp

@@ -272,21 +272,19 @@ int main(int argc, char **argv)
 		// Create a communicator for the device
 
 		ScaLBL_Communicator ScaLBL_Comm(Mask);
+		int *neighborList;
+		IntArray Map(Nx,Ny,Nz);
+		neighborList= new int[18*(Np+32)];
+		Np = ScaLBL_Comm.MemoryOptimizedLayoutAA(Map,neighborList,Mask.id,Np);
+		MPI_Barrier(comm);
 		
 		// LBM variables
 		if (rank==0)	printf ("Allocating distributions \n");
 
-		int neighborSize=18*Np*sizeof(int);
-		int *neighborList;
-		IntArray Map(Nx,Ny,Nz);
-
-		neighborList= new int[18*Np];
-		ScaLBL_Comm.MemoryOptimizedLayoutAA(Map,neighborList,Mask.id,Np);
-
-		MPI_Barrier(comm);
 
 		//......................device distributions.................................
 		int dist_mem_size = Np*sizeof(double);
+		int neighborSize=18*(Np*sizeof(int));
 
 		int *NeighborList;
 		//		double *f_even,*f_odd;
@@ -348,14 +346,14 @@ int main(int argc, char **argv)
 
 			timestep++;
 			ScaLBL_Comm.SendD3Q19AA(dist); //READ FROM NORMAL
-			ScaLBL_D3Q19_AAodd_BGK(NeighborList, dist, ScaLBL_Comm.next, Np, Np, rlx, Fx, Fy, Fz);
+			ScaLBL_D3Q19_AAodd_BGK(NeighborList, dist, ScaLBL_Comm.first_interior, ScaLBL_Comm.last_interior, Np, rlx, Fx, Fy, Fz);
 			ScaLBL_Comm.RecvD3Q19AA(dist); //WRITE INTO OPPOSITE
 			ScaLBL_D3Q19_AAodd_BGK(NeighborList, dist, 0, ScaLBL_Comm.next, Np, rlx, Fx, Fy, Fz);
 			ScaLBL_DeviceBarrier(); MPI_Barrier(comm);
 
 			timestep++;
 			ScaLBL_Comm.SendD3Q19AA(dist); //READ FORM NORMAL
-			ScaLBL_D3Q19_AAeven_BGK(dist, ScaLBL_Comm.next, Np, Np, rlx, Fx, Fy, Fz);
+			ScaLBL_D3Q19_AAeven_BGK(dist, ScaLBL_Comm.first_interior, ScaLBL_Comm.last_interior, Np, rlx, Fx, Fy, Fz);
 			ScaLBL_Comm.RecvD3Q19AA(dist); //WRITE INTO OPPOSITE
 			ScaLBL_D3Q19_AAeven_BGK(dist, 0, ScaLBL_Comm.next, Np, rlx, Fx, Fy, Fz);
 			ScaLBL_DeviceBarrier(); MPI_Barrier(comm);

tests/lbpm_color_macro_simulator.cpp

@@ -80,7 +80,6 @@ int main(int argc, char **argv)
 		int InitialCondition;
 		//	bool pBC,Restart;
 		int i,j,k,n;
-		int dist_mem_size;
 		double din, dout, flux;
 		double inletA,inletB,outletA,outletB;
 		inletA=1.f;
@@ -435,11 +434,10 @@ int main(int argc, char **argv)
 		ScaLBL_Communicator ScaLBL_Comm_Regular(Mask);
 		
 		if (rank==0)	printf ("Set up memory efficient layout \n");
-		int neighborSize=18*Np*sizeof(int);
 		int *neighborList;
 		IntArray Map(Nx,Ny,Nz);
-		neighborList= new int[18*Np];
-		ScaLBL_Comm.MemoryOptimizedLayoutAA(Map,neighborList,Mask.id,Np);
+		neighborList= new int[18*(Np+32)];
+		Np = ScaLBL_Comm.MemoryOptimizedLayoutAA(Map,neighborList,Mask.id,Np);
 		MPI_Barrier(comm);
 
 		//...........................................................................
@@ -448,8 +446,8 @@ int main(int argc, char **argv)
 		// LBM variables
 		if (rank==0)	printf ("Allocating distributions \n");
 		//......................device distributions.................................
-		dist_mem_size = Np*sizeof(double);
-		if (rank==0)	printf ("Allocating distributions \n");
+		int dist_mem_size = Np*sizeof(double);
+		int neighborSize=18*(Np*sizeof(int));
 
 		int *NeighborList;
 		int *dvcMap;
@@ -571,7 +569,7 @@ int main(int argc, char **argv)
 			// Compute the Phase indicator field
 			// Read for Aq, Bq happens in this routine (requires communication)
 			ScaLBL_Comm.BiSendD3Q7AA(Aq,Bq); //READ FROM NORMAL
-			ScaLBL_D3Q7_AAodd_PhaseField(NeighborList, dvcMap, Aq, Bq, Den, Phi, ScaLBL_Comm.next, Np, Np);
+			ScaLBL_D3Q7_AAodd_PhaseField(NeighborList, dvcMap, Aq, Bq, Den, Phi, ScaLBL_Comm.first_interior, ScaLBL_Comm.last_interior, Np);
 			ScaLBL_Comm.BiRecvD3Q7AA(Aq,Bq); //WRITE INTO OPPOSITE
 			ScaLBL_D3Q7_AAodd_PhaseField(NeighborList, dvcMap, Aq, Bq, Den, Phi, 0, ScaLBL_Comm.next, Np);
 			
@@ -581,8 +579,7 @@ int main(int argc, char **argv)
 			ScaLBL_Comm_Regular.SendHalo(Phi);
 
 			ScaLBL_D3Q19_AAodd_Color(NeighborList, dvcMap, fq, Aq, Bq, Den, Phi, Velocity, rhoA, rhoB, tauA, tauB,
-					alpha, beta, Fx, Fy, Fz, Nx, Nx*Ny, ScaLBL_Comm.next, Np, Np);
-			ScaLBL_Comm_Regular.RecvHalo(Phi);
+					alpha, beta, Fx, Fy, Fz, Nx, Nx*Ny, ScaLBL_Comm.first_interior, ScaLBL_Comm.last_interior, Np);
 			ScaLBL_Comm.RecvD3Q19AA(fq); //WRITE INTO OPPOSITE
 			// Set BCs
 			if (BoundaryCondition > 0){
@@ -605,7 +602,7 @@ int main(int argc, char **argv)
 			timestep++;
 			// Compute the Phase indicator field
 			ScaLBL_Comm.BiSendD3Q7AA(Aq,Bq); //READ FROM NORMAL
-			ScaLBL_D3Q7_AAeven_PhaseField(dvcMap, Aq, Bq, Den, Phi, ScaLBL_Comm.next, Np, Np);
+			ScaLBL_D3Q7_AAeven_PhaseField(dvcMap, Aq, Bq, Den, Phi, ScaLBL_Comm.first_interior, ScaLBL_Comm.last_interior, Np);
 			ScaLBL_Comm.BiRecvD3Q7AA(Aq,Bq); //WRITE INTO OPPOSITE
 			ScaLBL_D3Q7_AAeven_PhaseField(dvcMap, Aq, Bq, Den, Phi, 0, ScaLBL_Comm.next, Np);
 
@@ -616,7 +613,7 @@ int main(int argc, char **argv)
 			// Perform the collision operation
 			ScaLBL_Comm.SendD3Q19AA(fq); //READ FORM NORMAL
 			ScaLBL_D3Q19_AAeven_Color(dvcMap, fq, Aq, Bq, Den, Phi, Velocity, rhoA, rhoB, tauA, tauB,
-					alpha, beta, Fx, Fy, Fz,  Nx, Nx*Ny, ScaLBL_Comm.next, Np, Np);
+					alpha, beta, Fx, Fy, Fz,  Nx, Nx*Ny, ScaLBL_Comm.first_interior, ScaLBL_Comm.last_interior, Np);
 			ScaLBL_Comm.RecvD3Q19AA(fq); //WRITE INTO OPPOSITE
 			// Set boundary conditions
 			if (BoundaryCondition > 0){

tests/lbpm_permeability_simulator.cpp

@@ -190,7 +190,6 @@ int main(int argc, char **argv)
 		double porosity, pore_vol;
 		//...........................................................................
 		if (rank == 0) cout << "Reading in domain from signed distance function..." << endl;
-
 		//.......................................................................
 		// Read the signed distance
 		sprintf(LocalRankString,"%05d",rank);
@@ -278,17 +277,15 @@ int main(int argc, char **argv)
 		// LBM variables
 		if (rank==0)	printf ("Allocating distributions \n");
 
-		int neighborSize=18*Np*sizeof(int);
 		int *neighborList;
 		IntArray Map(Nx,Ny,Nz);
-
-		neighborList= new int[18*Np];
-		ScaLBL_Comm.MemoryOptimizedLayoutAA(Map,neighborList,Mask.id,Np);
-
+		neighborList= new int[18*(Np+32)];
+		Np = ScaLBL_Comm.MemoryOptimizedLayoutAA(Map,neighborList,Mask.id,Np);
 		MPI_Barrier(comm);
 
 		//......................device distributions.................................
 		int dist_mem_size = Np*sizeof(double);
+		int neighborSize=18*(Np*sizeof(int));
 
 		int *NeighborList;
 		//		double *f_even,*f_odd;
@@ -350,14 +347,14 @@ int main(int argc, char **argv)
 
 			timestep++;
 			ScaLBL_Comm.SendD3Q19AA(dist); //READ FROM NORMAL
-			ScaLBL_D3Q19_AAodd_MRT(NeighborList, dist, ScaLBL_Comm.next, Np, Np, rlx_setA, rlx_setB, Fx, Fy, Fz);
+			ScaLBL_D3Q19_AAodd_MRT(NeighborList, dist, ScaLBL_Comm.first_interior, ScaLBL_Comm.last_interior, Np, rlx_setA, rlx_setB, Fx, Fy, Fz);
 			ScaLBL_Comm.RecvD3Q19AA(dist); //WRITE INTO OPPOSITE
 			ScaLBL_D3Q19_AAodd_MRT(NeighborList, dist, 0, ScaLBL_Comm.next, Np, rlx_setA, rlx_setB, Fx, Fy, Fz);
 			ScaLBL_DeviceBarrier(); MPI_Barrier(comm);
 
 			timestep++;
 			ScaLBL_Comm.SendD3Q19AA(dist); //READ FORM NORMAL
-			ScaLBL_D3Q19_AAeven_MRT(dist, ScaLBL_Comm.next, Np, Np, rlx_setA, rlx_setB, Fx, Fy, Fz);
+			ScaLBL_D3Q19_AAeven_MRT(dist, ScaLBL_Comm.first_interior, ScaLBL_Comm.last_interior, Np, rlx_setA, rlx_setB, Fx, Fy, Fz);
 			ScaLBL_Comm.RecvD3Q19AA(dist); //WRITE INTO OPPOSITE
 			ScaLBL_D3Q19_AAeven_MRT(dist, 0, ScaLBL_Comm.next, Np, rlx_setA, rlx_setB, Fx, Fy, Fz);
 			ScaLBL_DeviceBarrier(); MPI_Barrier(comm);

tests/lbpm_segmented_pp.cpp

@@ -28,6 +28,92 @@ inline void MeanFilter(DoubleArray &Mesh){
 }
 
 
+double ReadFromBlock( char *ID, int iproc, int jproc, int kproc, int Nx, int Ny, int Nz){
+
+	int x,y,z,i,j,k;
+	// Get the list of blocks that contain the local rank
+	int b0x =iproc*(Nx-2)/1024;
+	int b0y =jproc*(Ny-2)/1024;
+	int b0z =kproc*(Nz-2)/1024;
+	int Bx =(iproc+1)*(Nx-2)/1024+1;
+	int By =(jproc+1)*(Ny-2)/1024+1;
+	int Bz =(kproc+1)*(Nz-2)/1024+1;
+
+	// arrays to hold the strings 
+    char LocalRankFilename[40];
+    char sx[2];
+    char sy[2];
+    char sz[2];
+    char tmpstr[10];
+    
+	// array to store ids read from block
+	char *id;
+	id = new char [1024*1024*1024];
+
+	for (int bz=b0z; bz<Bz; bz++){
+		for (int by=b0y; by<By; by++){
+			for (int bx=b0x; bx<Bx; bx++){
+				sprintf(sx,"%d",bx);
+				sprintf(sy,"%d",by);
+				sprintf(sz,"%d",bz);
+				sprintf(LocalRankFilename,"%s%s%s%s%s%s%s","a2_x",sx,"_y",sy,"_z",sz,".gbd");
+				//sprintf(LocalRankFilename,"%s%s%s%s%s%s%s","dis_",sx,"x_",sy,"y_",sz,"z.gbd");
+				//printf("Reading file: %s \n", LocalRankFilename);
+				//fflush(stdout);
+				
+				// Read the file
+				size_t readID;
+				FILE *IDFILE = fopen(LocalRankFilename,"rb");
+				readID=fread(id,1,1024*1024*1024,IDFILE);
+				fclose(IDFILE);
+				//printf("Loading data ... \n");
+				
+				// loop over global index
+				for ( k=0;k<Nz;k++){
+					for ( j=0;j<Ny;j++){
+						for ( i=0;i<Nx;i++){
+							//  id for this process relative to current block
+							x = iproc*(Nx-2) + i - 1 - 1024*bx;
+							y = jproc*(Ny-2) + j - 1 - 1024*by;
+							z = kproc*(Nz-2) + k - 1 - 1024*bz;
+							if (!(x<0) && !(y<0) && !(z<0) && x < 1024  && y < 1024  && z < 1024){
+							  //char value=id[z*1024*1024 + y*1024 + x];
+							  char value=0;
+								  // check porespace nodes
+								if (id[z*1024*1024 + y*1024 + x] < 215)	value = 1;
+								// set the ID
+								ID[k*Nx*Ny + j*Nx + i] = value;
+							}
+						}
+					}
+				}
+			}
+		}
+	}
+
+	// Compute porosity 
+	uint64_t count=0;
+	for (k=1; k<Nz-1; k++){
+		for (j=1; j<Ny-1; j++){
+			for (i=1; i<Nx-1; i++){
+				if (ID[k*Nx*Ny+j*Nx+i] == 1) count++; 
+			}
+		}
+	}
+	double porosity = double(count)/double((Nx-2)*(Ny-2)*(Nz-2));
+	return porosity;
+	/*printf("Porosity is %f \n",double(count)/double(NX*NY*NZ));
+
+	printf("Done getting data -- writing main file \n");
+	FILE *OUT = fopen("FullData.raw","wb");
+	fwrite(ID,1,N,OUT);
+	fclose(OUT);
+	printf("Completed! \n");
+	 */
+}
+
+
+
 int main(int argc, char **argv)
 {
 	// Initialize MPI
@@ -46,6 +132,14 @@ int main(int argc, char **argv)
 		int i,j,k,n;
 		int BC=0;
 		//  char fluidValue,solidValue;
+		int MAXTIME=1000;
+		int READ_FROM_BLOCK=0;
+		if (argc > 1){
+			MAXTIME=atoi(argv[1]);
+		}
+		if (argc > 2){
+			READ_FROM_BLOCK=atoi(argv[2]);			
+		}
 
 		std::vector<char> solidValues;
 		std::vector<char> nwpValues;
@@ -85,6 +179,8 @@ int main(int argc, char **argv)
 			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);
+			fflush(stdout);
+
 		}
 		if ( nprocs < nprocx*nprocy*nprocz ){
 			ERROR("Insufficient number of processors");
@@ -96,18 +192,36 @@ int main(int argc, char **argv)
 		Domain Dm(nx,ny,nz,rank,nprocx,nprocy,nprocz,Lx,Ly,Lz,BC);
 		for (n=0; n<N; n++) Dm.id[n]=1;
 		Dm.CommInit(comm);
-
-		// Read the phase ID
-		size_t readID;
-		sprintf(LocalRankFilename,"ID.%05i",rank);
-		FILE *ID = fopen(LocalRankFilename,"rb");
-		readID=fread(Dm.id,1,N,ID);
-		if (readID != size_t(N)) printf("lbpm_segmented_pp: Error reading ID \n");
-		fclose(ID);
-
-		// Initialize the domain and communication
 		nx+=2; ny+=2; nz+=2;
 
+		// Read the phase ID
+		double porosity;
+		sprintf(LocalRankFilename,"ID.%05i",rank);
+		if (READ_FROM_BLOCK == 0){
+			size_t readID;
+			FILE *ID = fopen(LocalRankFilename,"rb");
+			readID=fread(Dm.id,1,N,ID);
+			if (readID != size_t(N)) printf("lbpm_segmented_pp: Error reading ID \n");
+			fclose(ID);
+		}
+		else{
+			// Read from the large block and write the local ID file
+			if (rank==0) printf("Reading ID file from blocks \n");
+			fflush(stdout);
+			porosity = ReadFromBlock(Dm.id,Dm.iproc,Dm.jproc,Dm.kproc,nx,ny,nz);
+			
+			MPI_Barrier(MPI_COMM_WORLD);
+			if (rank==0) printf("Writing local ID files (poros=%f) \n",porosity);
+			fflush(stdout);
+			FILE *ID = fopen(LocalRankFilename,"wb");
+			fwrite(Dm.id,1,N,ID);
+			fclose(ID);
+			if (rank==0) printf("Succeeded! \n");
+			fflush(stdout);
+
+		}
+
+		// Initialize the domain and communication
 		char *id;
 		id = new char [N];
 		TwoPhase Averages(Dm);
@@ -141,6 +255,7 @@ int main(int argc, char **argv)
 		double LocalVar, TotalVar;
 		if (rank==0) printf("Initialized solid phase -- Converting to Signed Distance function \n");
 		int Maxtime=10*max(max(Dm.Nx*Dm.nprocx,Dm.Ny*Dm.nprocy),Dm.Nz*Dm.nprocz);
+		Maxtime=min(Maxtime,MAXTIME);
 		LocalVar = Eikonal(Averages.SDs,id,Dm,Maxtime);
 
 		MPI_Allreduce(&LocalVar,&TotalVar,1,MPI_DOUBLE,MPI_SUM,comm);