OilfieldToolsNet/LBPM
4
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

replacing ScaLBL functions in tests/*cpp

Browse Source
This commit is contained in:
James E McClure 2016-11-23 16:16:48 -05:00
parent 6ca3eb2079
commit b7a5ea342f
11 changed files with 1442 additions and 1329 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

586
common/ScaLBL.h
View File

@ -9,80 +9,74 @@
*/
#include "Domain.h"
extern "C" void AllocateDeviceMemory(void** address, size_t size);
extern "C" void ScaLBL_AllocateDeviceMemory(void** address, size_t size);
//extern "C" void FreeDeviceMemory(void** address);
extern "C" void CopyToDevice(void* dest, const void* source, size_t size);
extern "C" void ScaLBL_CopyToDevice(void* dest, const void* source, size_t size);
extern "C" void CopyToHost(void* dest, const void* source, size_t size);
extern "C" void ScaLBL_CopyToHost(void* dest, const void* source, size_t size);
extern "C" void DeviceBarrier();
extern "C" void ScaLBL_DeviceBarrier();
extern "C" void PackDist(int q, int *list, int start, int count, double *sendbuf, double *dist, int N);
extern "C" void ScaLBL_D3Q19_Pack(int q, int *list, int start, int count, double *sendbuf, double *dist, int N);
extern "C" void UnpackDist(int q, int Cqx, int Cqy, int Cqz, int *list, int start, int count,double *recvbuf, double *dist, int Nx, int Ny, int Nz);
extern "C" void ScaLBL_D3Q19_Unpack(int q, int Cqx, int Cqy, int Cqz, int *list, int start, int count,double *recvbuf, double *dist, int Nx, int Ny, int Nz);
extern "C" void PackValues(int *list, int count, double *sendbuf, double *Data, int N);
extern "C" void ScaLBL_Scalar_Pack(int *list, int count, double *sendbuf, double *Data, int N);
extern "C" void UnpackValues(int *list, int count, double *recvbuf, double *Data, int N);
extern "C" void ScaLBL_Scalar_Unpack(int *list, int count, double *recvbuf, double *Data, int N);
extern "C" void PackDenD3Q7(int *list, int count, double *sendbuf, int number, double *Data, int N);
extern "C" void ScaLBL_PackDenD3Q7(int *list, int count, double *sendbuf, int number, double *Data, int N);
extern "C" void UnpackDenD3Q7(int *list, int count, double *recvbuf, int number, double *Data, int N);
extern "C" void ScaLBL_UnpackDenD3Q7(int *list, int count, double *recvbuf, int number, double *Data, int N);
extern "C" void InitD3Q7(char *ID, double *f_even, double *f_odd, double *Den, int Nx, int Ny, int Nz);
extern "C" void ScaLBL_D3Q7_Init(char *ID, double *f_even, double *f_odd, double *Den, int Nx, int Ny, int Nz);
extern "C" void SwapD3Q7(char *ID, double *disteven, double *distodd, int Nx, int Ny, int Nz);
extern "C" void ScaLBL_D3Q7_Swap(char *ID, double *disteven, double *distodd, int Nx, int Ny, int Nz);
extern "C" void ComputeDensityD3Q7(char *ID, double *disteven, double *distodd, double *Den,
extern "C" void ScaLBL_D3Q7_Density(char *ID, double *disteven, double *distodd, double *Den,
int Nx, int Ny, int Nz);
extern "C" void InitD3Q19(char *ID, double *f_even, double *f_odd, int Nx, int Ny, int Nz);
extern "C" void ScaLBL_D3Q19_Init(char *ID, double *f_even, double *f_odd, int Nx, int Ny, int Nz);
extern "C" void SwapD3Q19(char *ID, double *disteven, double *distodd, int Nx, int Ny, int Nz);
extern "C" void ScaLBL_D3Q19_Swap(char *ID, double *disteven, double *distodd, int Nx, int Ny, int Nz);
extern "C" void MRT(char *ID, double *f_even, double *f_odd, double rlxA, double rlxB,
extern "C" void ScaLBL_D3Q19_MRT(char *ID, double *f_even, double *f_odd, double rlxA, double rlxB,
double Fx, double Fy, double Fz,int Nx, int Ny, int Nz);
extern "C" void ComputeVelocityD3Q19(char *ID, double *disteven, double *distodd, double *vel,
extern "C" void ScaLBL_D3Q19_Velocity(char *ID, double *disteven, double *distodd, double *vel,
int Nx, int Ny, int Nz);
extern "C" void ComputePressureD3Q19(const char *ID, const double *disteven, const double *distodd,
extern "C" void ScaLBL_D3Q19_Pressure(const char *ID, const double *disteven, const double *distodd,
double *Pressure, int Nx, int Ny, int Nz);
extern "C" void PressureBC_inlet(double *disteven, double *distodd, double din,
extern "C" void ScaLBL_D3Q19_Pressure_BC_z(double *disteven, double *distodd, double din,
int Nx, int Ny, int Nz);
extern "C" void PressureBC_outlet(double *disteven, double *distodd, double dout,
extern "C" void ScaLBL_D3Q19_Pressure_BC_Z(double *disteven, double *distodd, double dout,
int Nx, int Ny, int Nz, int outlet);
extern "C" void InitDenColor(char *ID, double *Den, double *Phi, double das, double dbs, int Nx, int Ny, int Nz);
extern "C" void ScaLBL_Color_Init(char *ID, double *Den, double *Phi, double das, double dbs, int Nx, int Ny, int Nz);
extern "C" void InitDenColorDistance(char *ID, double *Den, double *Phi, double *Distance,
extern "C" void ScaLBL_ColorDistance_Init(char *ID, double *Den, double *Phi, double *Distance,
double das, double dbs, double beta, double xp, int Nx, int Ny, int Nz);
extern "C" void ComputeColorGradient(char *ID, double *phi, double *ColorGrad, int Nx, int Ny, int Nz);
extern "C" void ScaLBL_D3Q19_ColorGradient(char *ID, double *phi, double *ColorGrad, int Nx, int Ny, int Nz);
extern "C" void ColorCollide( char *ID, double *disteven, double *distodd, double *ColorGrad,
double *Velocity, int Nx, int Ny, int Nz,double rlx_setA, double rlx_setB,
double alpha, double beta, double Fx, double Fy, double Fz, bool pBC);
extern "C" void ColorCollideOpt( char *ID, double *disteven, double *distodd, double *phi, double *ColorGrad,
extern "C" void ScaLBL_D3Q19_ColorCollide( char *ID, double *disteven, double *distodd, double *phi, double *ColorGrad,
double *Velocity, int Nx, int Ny, int Nz,double rlx_setA, double rlx_setB,
double alpha, double beta, double Fx, double Fy, double Fz);
extern "C" void DensityStreamD3Q7(char *ID, double *Den, double *Copy, double *Phi, double *ColorGrad, double *Velocity,
double beta, int Nx, int Ny, int Nz, bool pBC);
extern "C" void ComputePhi(char *ID, double *Phi, double *Den, int N);
extern "C" void MassColorCollideD3Q7(char *ID, double *A_even, double *A_odd, double *B_even, double *B_odd,
extern "C" void ScaLBL_D3Q7_ColorCollideMass(char *ID, double *A_even, double *A_odd, double *B_even, double *B_odd,
double *Den, double *Phi, double *ColorGrad, double *Velocity, double beta, int N, bool pBC);
extern "C" void ColorBC_inlet(double *Phi, double *Den, double *A_even, double *A_odd,
extern "C" void ScaLBL_ComputePhaseField(char *ID, double *Phi, double *Den, int N);
extern "C" void ScaLBL_Color_BC_z(double *Phi, double *Den, double *A_even, double *A_odd,
double *B_even, double *B_odd, int Nx, int Ny, int Nz);
extern "C" void ColorBC_outlet(double *Phi, double *Den, double *A_even, double *A_odd,
extern "C" void ScaLBL_Color_BC_Z(double *Phi, double *Den, double *A_even, double *A_odd,
double *B_even, double *B_odd, int Nx, int Ny, int Nz);
extern "C" void ScaLBL_D3Q19_Velocity_BC_z(double *disteven, double *distodd, double uz,
@ -91,7 +85,7 @@ extern "C" void ScaLBL_D3Q19_Velocity_BC_z(double *disteven, double *distodd, do
extern "C" void ScaLBL_D3Q19_Velocity_BC_Z(double *disteven, double *distodd, double uz,
int Nx, int Ny, int Nz, int outlet);
extern "C" void SetPhiSlice_z(double *Phi, double value, int Nx, int Ny, int Nz, int Slice);
extern "C" void ScaLBL_SetSlice_z(double *Phi, double value, int Nx, int Ny, int Nz, int Slice);
class ScaLBL_Communicator{
public:
@ -234,122 +228,122 @@ ScaLBL_Communicator::ScaLBL_Communicator(Domain &Dm){
recvCount_YZ=Dm.recvCount_YZ;
recvCount_XZ=Dm.recvCount_XZ;
//......................................................................................
AllocateDeviceMemory((void **) &sendbuf_x, 5*sendCount_x*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_X, 5*sendCount_X*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_y, 5*sendCount_y*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_Y, 5*sendCount_Y*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_z, 5*sendCount_z*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_Z, 5*sendCount_Z*sizeof(double)); // Allocatevoid * memory
AllocateDeviceMemory((void **) &sendbuf_xy, sendCount_xy*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_xY, sendCount_xY*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_Xy, sendCount_Xy*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_XY, sendCount_XY*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_xz, sendCount_xz*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_xZ, sendCount_xZ*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_Xz, sendCount_Xz*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_XZ, sendCount_XZ*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_yz, sendCount_yz*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_yZ, sendCount_yZ*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_Yz, sendCount_Yz*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_YZ, sendCount_YZ*sizeof(double)); // Allocate device memory
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)); // Allocatevoid * 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
//......................................................................................
AllocateDeviceMemory((void **) &recvbuf_x, 5*recvCount_x*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_X, 5*recvCount_X*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_y, 5*recvCount_y*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_Y, 5*recvCount_Y*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_z, 5*recvCount_z*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_Z, 5*recvCount_Z*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_xy, recvCount_xy*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_xY, recvCount_xY*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_Xy, recvCount_Xy*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_XY, recvCount_XY*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_xz, recvCount_xz*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_xZ, recvCount_xZ*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_Xz, recvCount_Xz*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_XZ, recvCount_XZ*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_yz, recvCount_yz*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_yZ, recvCount_yZ*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_Yz, recvCount_Yz*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_YZ, recvCount_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
//......................................................................................
AllocateDeviceMemory((void **) &dvcSendList_x, sendCount_x*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_X, sendCount_X*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_y, sendCount_y*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_Y, sendCount_Y*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_z, sendCount_z*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_Z, sendCount_Z*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_xy, sendCount_xy*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_xY, sendCount_xY*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_Xy, sendCount_Xy*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_XY, sendCount_XY*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_xz, sendCount_xz*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_xZ, sendCount_xZ*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_Xz, sendCount_Xz*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_XZ, sendCount_XZ*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_yz, sendCount_yz*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_yZ, sendCount_yZ*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_Yz, sendCount_Yz*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_YZ, sendCount_YZ*sizeof(int)); // 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
//......................................................................................
AllocateDeviceMemory((void **) &dvcRecvList_x, recvCount_x*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_X, recvCount_X*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_y, recvCount_y*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_Y, recvCount_Y*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_z, recvCount_z*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_Z, recvCount_Z*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_xy, recvCount_xy*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_xY, recvCount_xY*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_Xy, recvCount_Xy*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_XY, recvCount_XY*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_xz, recvCount_xz*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_xZ, recvCount_xZ*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_Xz, recvCount_Xz*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_XZ, recvCount_XZ*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_yz, recvCount_yz*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_yZ, recvCount_yZ*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_Yz, recvCount_Yz*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_YZ, recvCount_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
//......................................................................................
CopyToDevice(dvcSendList_x,Dm.sendList_x,sendCount_x*sizeof(int));
CopyToDevice(dvcSendList_X,Dm.sendList_X,sendCount_X*sizeof(int));
CopyToDevice(dvcSendList_y,Dm.sendList_y,sendCount_y*sizeof(int));
CopyToDevice(dvcSendList_Y,Dm.sendList_Y,sendCount_Y*sizeof(int));
CopyToDevice(dvcSendList_z,Dm.sendList_z,sendCount_z*sizeof(int));
CopyToDevice(dvcSendList_Z,Dm.sendList_Z,sendCount_Z*sizeof(int));
CopyToDevice(dvcSendList_xy,Dm.sendList_xy,sendCount_xy*sizeof(int));
CopyToDevice(dvcSendList_XY,Dm.sendList_XY,sendCount_XY*sizeof(int));
CopyToDevice(dvcSendList_xY,Dm.sendList_xY,sendCount_xY*sizeof(int));
CopyToDevice(dvcSendList_Xy,Dm.sendList_Xy,sendCount_Xy*sizeof(int));
CopyToDevice(dvcSendList_xz,Dm.sendList_xz,sendCount_xz*sizeof(int));
CopyToDevice(dvcSendList_XZ,Dm.sendList_XZ,sendCount_XZ*sizeof(int));
CopyToDevice(dvcSendList_xZ,Dm.sendList_xZ,sendCount_xZ*sizeof(int));
CopyToDevice(dvcSendList_Xz,Dm.sendList_Xz,sendCount_Xz*sizeof(int));
CopyToDevice(dvcSendList_yz,Dm.sendList_yz,sendCount_yz*sizeof(int));
CopyToDevice(dvcSendList_YZ,Dm.sendList_YZ,sendCount_YZ*sizeof(int));
CopyToDevice(dvcSendList_yZ,Dm.sendList_yZ,sendCount_yZ*sizeof(int));
CopyToDevice(dvcSendList_Yz,Dm.sendList_Yz,sendCount_Yz*sizeof(int));
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));
//......................................................................................
CopyToDevice(dvcRecvList_x,Dm.recvList_x,recvCount_x*sizeof(int));
CopyToDevice(dvcRecvList_X,Dm.recvList_X,recvCount_X*sizeof(int));
CopyToDevice(dvcRecvList_y,Dm.recvList_y,recvCount_y*sizeof(int));
CopyToDevice(dvcRecvList_Y,Dm.recvList_Y,recvCount_Y*sizeof(int));
CopyToDevice(dvcRecvList_z,Dm.recvList_z,recvCount_z*sizeof(int));
CopyToDevice(dvcRecvList_Z,Dm.recvList_Z,recvCount_Z*sizeof(int));
CopyToDevice(dvcRecvList_xy,Dm.recvList_xy,recvCount_xy*sizeof(int));
CopyToDevice(dvcRecvList_XY,Dm.recvList_XY,recvCount_XY*sizeof(int));
CopyToDevice(dvcRecvList_xY,Dm.recvList_xY,recvCount_xY*sizeof(int));
CopyToDevice(dvcRecvList_Xy,Dm.recvList_Xy,recvCount_Xy*sizeof(int));
CopyToDevice(dvcRecvList_xz,Dm.recvList_xz,recvCount_xz*sizeof(int));
CopyToDevice(dvcRecvList_XZ,Dm.recvList_XZ,recvCount_XZ*sizeof(int));
CopyToDevice(dvcRecvList_xZ,Dm.recvList_xZ,recvCount_xZ*sizeof(int));
CopyToDevice(dvcRecvList_Xz,Dm.recvList_Xz,recvCount_Xz*sizeof(int));
CopyToDevice(dvcRecvList_yz,Dm.recvList_yz,recvCount_yz*sizeof(int));
CopyToDevice(dvcRecvList_YZ,Dm.recvList_YZ,recvCount_YZ*sizeof(int));
CopyToDevice(dvcRecvList_yZ,Dm.recvList_yZ,recvCount_yZ*sizeof(int));
CopyToDevice(dvcRecvList_Yz,Dm.recvList_Yz,recvCount_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));
//......................................................................................
MPI_Barrier(MPI_COMM_SCALBL);
DeviceBarrier();
ScaLBL_DeviceBarrier();
//......................................................................................
SendCount = sendCount_x+sendCount_X+sendCount_y+sendCount_Y+sendCount_z+sendCount_Z+
sendCount_xy+sendCount_Xy+sendCount_xY+sendCount_XY+
@ -381,68 +375,68 @@ void ScaLBL_Communicator::SendD3Q19(double *f_even, double *f_odd){
}
// assign tag of 19 to D3Q19 communication
sendtag = recvtag = 19;
DeviceBarrier();
ScaLBL_DeviceBarrier();
// Pack the distributions
PackDist(1,dvcSendList_x,0,sendCount_x,sendbuf_x,f_even,N);
PackDist(4,dvcSendList_x,sendCount_x,sendCount_x,sendbuf_x,f_even,N);
PackDist(5,dvcSendList_x,2*sendCount_x,sendCount_x,sendbuf_x,f_even,N);
PackDist(6,dvcSendList_x,3*sendCount_x,sendCount_x,sendbuf_x,f_even,N);
PackDist(7,dvcSendList_x,4*sendCount_x,sendCount_x,sendbuf_x,f_even,N);
ScaLBL_D3Q19_Pack(1,dvcSendList_x,0,sendCount_x,sendbuf_x,f_even,N);
ScaLBL_D3Q19_Pack(4,dvcSendList_x,sendCount_x,sendCount_x,sendbuf_x,f_even,N);
ScaLBL_D3Q19_Pack(5,dvcSendList_x,2*sendCount_x,sendCount_x,sendbuf_x,f_even,N);
ScaLBL_D3Q19_Pack(6,dvcSendList_x,3*sendCount_x,sendCount_x,sendbuf_x,f_even,N);
ScaLBL_D3Q19_Pack(7,dvcSendList_x,4*sendCount_x,sendCount_x,sendbuf_x,f_even,N);
//...Packing for X face(1,7,9,11,13)................................
PackDist(0,dvcSendList_X,0,sendCount_X,sendbuf_X,f_odd,N);
PackDist(3,dvcSendList_X,sendCount_X,sendCount_X,sendbuf_X,f_odd,N);
PackDist(4,dvcSendList_X,2*sendCount_X,sendCount_X,sendbuf_X,f_odd,N);
PackDist(5,dvcSendList_X,3*sendCount_X,sendCount_X,sendbuf_X,f_odd,N);
PackDist(6,dvcSendList_X,4*sendCount_X,sendCount_X,sendbuf_X,f_odd,N);
ScaLBL_D3Q19_Pack(0,dvcSendList_X,0,sendCount_X,sendbuf_X,f_odd,N);
ScaLBL_D3Q19_Pack(3,dvcSendList_X,sendCount_X,sendCount_X,sendbuf_X,f_odd,N);
ScaLBL_D3Q19_Pack(4,dvcSendList_X,2*sendCount_X,sendCount_X,sendbuf_X,f_odd,N);
ScaLBL_D3Q19_Pack(5,dvcSendList_X,3*sendCount_X,sendCount_X,sendbuf_X,f_odd,N);
ScaLBL_D3Q19_Pack(6,dvcSendList_X,4*sendCount_X,sendCount_X,sendbuf_X,f_odd,N);
//...Packing for y face(4,8,9,16,18).................................
PackDist(2,dvcSendList_y,0,sendCount_y,sendbuf_y,f_even,N);
PackDist(4,dvcSendList_y,sendCount_y,sendCount_y,sendbuf_y,f_even,N);
PackDist(4,dvcSendList_y,2*sendCount_y,sendCount_y,sendbuf_y,f_odd,N);
PackDist(8,dvcSendList_y,3*sendCount_y,sendCount_y,sendbuf_y,f_even,N);
PackDist(9,dvcSendList_y,4*sendCount_y,sendCount_y,sendbuf_y,f_even,N);
ScaLBL_D3Q19_Pack(2,dvcSendList_y,0,sendCount_y,sendbuf_y,f_even,N);
ScaLBL_D3Q19_Pack(4,dvcSendList_y,sendCount_y,sendCount_y,sendbuf_y,f_even,N);
ScaLBL_D3Q19_Pack(4,dvcSendList_y,2*sendCount_y,sendCount_y,sendbuf_y,f_odd,N);
ScaLBL_D3Q19_Pack(8,dvcSendList_y,3*sendCount_y,sendCount_y,sendbuf_y,f_even,N);
ScaLBL_D3Q19_Pack(9,dvcSendList_y,4*sendCount_y,sendCount_y,sendbuf_y,f_even,N);
//...Packing for Y face(3,7,10,15,17).................................
PackDist(1,dvcSendList_Y,0,sendCount_Y,sendbuf_Y,f_odd,N);
PackDist(3,dvcSendList_Y,sendCount_Y,sendCount_Y,sendbuf_Y,f_odd,N);
PackDist(5,dvcSendList_Y,2*sendCount_Y,sendCount_Y,sendbuf_Y,f_even,N);
PackDist(7,dvcSendList_Y,3*sendCount_Y,sendCount_Y,sendbuf_Y,f_odd,N);
PackDist(8,dvcSendList_Y,4*sendCount_Y,sendCount_Y,sendbuf_Y,f_odd,N);
ScaLBL_D3Q19_Pack(1,dvcSendList_Y,0,sendCount_Y,sendbuf_Y,f_odd,N);
ScaLBL_D3Q19_Pack(3,dvcSendList_Y,sendCount_Y,sendCount_Y,sendbuf_Y,f_odd,N);
ScaLBL_D3Q19_Pack(5,dvcSendList_Y,2*sendCount_Y,sendCount_Y,sendbuf_Y,f_even,N);
ScaLBL_D3Q19_Pack(7,dvcSendList_Y,3*sendCount_Y,sendCount_Y,sendbuf_Y,f_odd,N);
ScaLBL_D3Q19_Pack(8,dvcSendList_Y,4*sendCount_Y,sendCount_Y,sendbuf_Y,f_odd,N);
//...Packing for z face(6,12,13,16,17)................................
PackDist(3,dvcSendList_z,0,sendCount_z,sendbuf_z,f_even,N);
PackDist(6,dvcSendList_z,sendCount_z,sendCount_z,sendbuf_z,f_even,N);
PackDist(6,dvcSendList_z,2*sendCount_z,sendCount_z,sendbuf_z,f_odd,N);
PackDist(8,dvcSendList_z,3*sendCount_z,sendCount_z,sendbuf_z,f_even,N);
PackDist(8,dvcSendList_z,4*sendCount_z,sendCount_z,sendbuf_z,f_odd,N);
ScaLBL_D3Q19_Pack(3,dvcSendList_z,0,sendCount_z,sendbuf_z,f_even,N);
ScaLBL_D3Q19_Pack(6,dvcSendList_z,sendCount_z,sendCount_z,sendbuf_z,f_even,N);
ScaLBL_D3Q19_Pack(6,dvcSendList_z,2*sendCount_z,sendCount_z,sendbuf_z,f_odd,N);
ScaLBL_D3Q19_Pack(8,dvcSendList_z,3*sendCount_z,sendCount_z,sendbuf_z,f_even,N);
ScaLBL_D3Q19_Pack(8,dvcSendList_z,4*sendCount_z,sendCount_z,sendbuf_z,f_odd,N);
//...Packing for Z face(5,11,14,15,18)................................
PackDist(2,dvcSendList_Z,0,sendCount_Z,sendbuf_Z,f_odd,N);
PackDist(5,dvcSendList_Z,sendCount_Z,sendCount_Z,sendbuf_Z,f_odd,N);
PackDist(7,dvcSendList_Z,2*sendCount_Z,sendCount_Z,sendbuf_Z,f_even,N);
PackDist(7,dvcSendList_Z,3*sendCount_Z,sendCount_Z,sendbuf_Z,f_odd,N);
PackDist(9,dvcSendList_Z,4*sendCount_Z,sendCount_Z,sendbuf_Z,f_even,N);
ScaLBL_D3Q19_Pack(2,dvcSendList_Z,0,sendCount_Z,sendbuf_Z,f_odd,N);
ScaLBL_D3Q19_Pack(5,dvcSendList_Z,sendCount_Z,sendCount_Z,sendbuf_Z,f_odd,N);
ScaLBL_D3Q19_Pack(7,dvcSendList_Z,2*sendCount_Z,sendCount_Z,sendbuf_Z,f_even,N);
ScaLBL_D3Q19_Pack(7,dvcSendList_Z,3*sendCount_Z,sendCount_Z,sendbuf_Z,f_odd,N);
ScaLBL_D3Q19_Pack(9,dvcSendList_Z,4*sendCount_Z,sendCount_Z,sendbuf_Z,f_even,N);
//...Pack the xy edge (8)................................
PackDist(4,dvcSendList_xy,0,sendCount_xy,sendbuf_xy,f_even,N);
ScaLBL_D3Q19_Pack(4,dvcSendList_xy,0,sendCount_xy,sendbuf_xy,f_even,N);
//...Pack the Xy edge (9)................................
PackDist(4,dvcSendList_Xy,0,sendCount_Xy,sendbuf_Xy,f_odd,N);
ScaLBL_D3Q19_Pack(4,dvcSendList_Xy,0,sendCount_Xy,sendbuf_Xy,f_odd,N);
//...Pack the xY edge (10)................................
PackDist(5,dvcSendList_xY,0,sendCount_xY,sendbuf_xY,f_even,N);
ScaLBL_D3Q19_Pack(5,dvcSendList_xY,0,sendCount_xY,sendbuf_xY,f_even,N);
//...Pack the XY edge (7)................................
PackDist(3,dvcSendList_XY,0,sendCount_XY,sendbuf_XY,f_odd,N);
ScaLBL_D3Q19_Pack(3,dvcSendList_XY,0,sendCount_XY,sendbuf_XY,f_odd,N);
//...Pack the xz edge (12)................................
PackDist(6,dvcSendList_xz,0,sendCount_xz,sendbuf_xz,f_even,N);
ScaLBL_D3Q19_Pack(6,dvcSendList_xz,0,sendCount_xz,sendbuf_xz,f_even,N);
//...Pack the xZ edge (14)................................
PackDist(7,dvcSendList_xZ,0,sendCount_xZ,sendbuf_xZ,f_even,N);
ScaLBL_D3Q19_Pack(7,dvcSendList_xZ,0,sendCount_xZ,sendbuf_xZ,f_even,N);
//...Pack the Xz edge (13)................................
PackDist(6,dvcSendList_Xz,0,sendCount_Xz,sendbuf_Xz,f_odd,N);
ScaLBL_D3Q19_Pack(6,dvcSendList_Xz,0,sendCount_Xz,sendbuf_Xz,f_odd,N);
//...Pack the XZ edge (11)................................
PackDist(5,dvcSendList_XZ,0,sendCount_XZ,sendbuf_XZ,f_odd,N);
ScaLBL_D3Q19_Pack(5,dvcSendList_XZ,0,sendCount_XZ,sendbuf_XZ,f_odd,N);
//...Pack the xz edge (12)................................
//...Pack the yz edge (16)................................
PackDist(8,dvcSendList_yz,0,sendCount_yz,sendbuf_yz,f_even,N);
ScaLBL_D3Q19_Pack(8,dvcSendList_yz,0,sendCount_yz,sendbuf_yz,f_even,N);
//...Pack the yZ edge (18)................................
PackDist(9,dvcSendList_yZ,0,sendCount_yZ,sendbuf_yZ,f_even,N);
ScaLBL_D3Q19_Pack(9,dvcSendList_yZ,0,sendCount_yZ,sendbuf_yZ,f_even,N);
//...Pack the Yz edge (17)................................
PackDist(8,dvcSendList_Yz,0,sendCount_Yz,sendbuf_Yz,f_odd,N);
ScaLBL_D3Q19_Pack(8,dvcSendList_Yz,0,sendCount_Yz,sendbuf_Yz,f_odd,N);
//...Pack the YZ edge (15)................................
PackDist(7,dvcSendList_YZ,0,sendCount_YZ,sendbuf_YZ,f_odd,N);
ScaLBL_D3Q19_Pack(7,dvcSendList_YZ,0,sendCount_YZ,sendbuf_YZ,f_odd,N);
//...................................................................................
//...................................................................................
@ -490,76 +484,76 @@ void ScaLBL_Communicator::RecvD3Q19(double *f_even, double *f_odd){
// Wait for completion of D3Q19 communication
MPI_Waitall(18,req1,stat1);
MPI_Waitall(18,req2,stat2);
DeviceBarrier();
ScaLBL_DeviceBarrier();
//...................................................................................
// Unpack the distributions on the device
//...................................................................................
//...Map recieve list for the X face: q=2,8,10,12,13 .................................
UnpackDist(0,-1,0,0,dvcRecvList_X,0,recvCount_X,recvbuf_X,f_odd,Nx,Ny,Nz);
UnpackDist(3,-1,-1,0,dvcRecvList_X,recvCount_X,recvCount_X,recvbuf_X,f_odd,Nx,Ny,Nz);
UnpackDist(4,-1,1,0,dvcRecvList_X,2*recvCount_X,recvCount_X,recvbuf_X,f_odd,Nx,Ny,Nz);
UnpackDist(5,-1,0,-1,dvcRecvList_X,3*recvCount_X,recvCount_X,recvbuf_X,f_odd,Nx,Ny,Nz);
UnpackDist(6,-1,0,1,dvcRecvList_X,4*recvCount_X,recvCount_X,recvbuf_X,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(0,-1,0,0,dvcRecvList_X,0,recvCount_X,recvbuf_X,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(3,-1,-1,0,dvcRecvList_X,recvCount_X,recvCount_X,recvbuf_X,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(4,-1,1,0,dvcRecvList_X,2*recvCount_X,recvCount_X,recvbuf_X,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(5,-1,0,-1,dvcRecvList_X,3*recvCount_X,recvCount_X,recvbuf_X,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(6,-1,0,1,dvcRecvList_X,4*recvCount_X,recvCount_X,recvbuf_X,f_odd,Nx,Ny,Nz);
//...................................................................................
//...Map recieve list for the x face: q=1,7,9,11,13..................................
UnpackDist(1,1,0,0,dvcRecvList_x,0,recvCount_x,recvbuf_x,f_even,Nx,Ny,Nz);
UnpackDist(4,1,1,0,dvcRecvList_x,recvCount_x,recvCount_x,recvbuf_x,f_even,Nx,Ny,Nz);
UnpackDist(5,1,-1,0,dvcRecvList_x,2*recvCount_x,recvCount_x,recvbuf_x,f_even,Nx,Ny,Nz);
UnpackDist(6,1,0,1,dvcRecvList_x,3*recvCount_x,recvCount_x,recvbuf_x,f_even,Nx,Ny,Nz);
UnpackDist(7,1,0,-1,dvcRecvList_x,4*recvCount_x,recvCount_x,recvbuf_x,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(1,1,0,0,dvcRecvList_x,0,recvCount_x,recvbuf_x,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(4,1,1,0,dvcRecvList_x,recvCount_x,recvCount_x,recvbuf_x,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(5,1,-1,0,dvcRecvList_x,2*recvCount_x,recvCount_x,recvbuf_x,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(6,1,0,1,dvcRecvList_x,3*recvCount_x,recvCount_x,recvbuf_x,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(7,1,0,-1,dvcRecvList_x,4*recvCount_x,recvCount_x,recvbuf_x,f_even,Nx,Ny,Nz);
//...................................................................................
//...Map recieve list for the y face: q=4,8,9,16,18 ...................................
UnpackDist(1,0,-1,0,dvcRecvList_Y,0,recvCount_Y,recvbuf_Y,f_odd,Nx,Ny,Nz);
UnpackDist(3,-1,-1,0,dvcRecvList_Y,recvCount_Y,recvCount_Y,recvbuf_Y,f_odd,Nx,Ny,Nz);
UnpackDist(5,1,-1,0,dvcRecvList_Y,2*recvCount_Y,recvCount_Y,recvbuf_Y,f_even,Nx,Ny,Nz);
UnpackDist(7,0,-1,-1,dvcRecvList_Y,3*recvCount_Y,recvCount_Y,recvbuf_Y,f_odd,Nx,Ny,Nz);
UnpackDist(8,0,-1,1,dvcRecvList_Y,4*recvCount_Y,recvCount_Y,recvbuf_Y,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(1,0,-1,0,dvcRecvList_Y,0,recvCount_Y,recvbuf_Y,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(3,-1,-1,0,dvcRecvList_Y,recvCount_Y,recvCount_Y,recvbuf_Y,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(5,1,-1,0,dvcRecvList_Y,2*recvCount_Y,recvCount_Y,recvbuf_Y,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(7,0,-1,-1,dvcRecvList_Y,3*recvCount_Y,recvCount_Y,recvbuf_Y,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(8,0,-1,1,dvcRecvList_Y,4*recvCount_Y,recvCount_Y,recvbuf_Y,f_odd,Nx,Ny,Nz);
//...................................................................................
//...Map recieve list for the Y face: q=3,7,10,15,17 ..................................
UnpackDist(2,0,1,0,dvcRecvList_y,0,recvCount_y,recvbuf_y,f_even,Nx,Ny,Nz);
UnpackDist(4,1,1,0,dvcRecvList_y,recvCount_y,recvCount_y,recvbuf_y,f_even,Nx,Ny,Nz);
UnpackDist(4,-1,1,0,dvcRecvList_y,2*recvCount_y,recvCount_y,recvbuf_y,f_odd,Nx,Ny,Nz);
UnpackDist(8,0,1,1,dvcRecvList_y,3*recvCount_y,recvCount_y,recvbuf_y,f_even,Nx,Ny,Nz);
UnpackDist(9,0,1,-1,dvcRecvList_y,4*recvCount_y,recvCount_y,recvbuf_y,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(2,0,1,0,dvcRecvList_y,0,recvCount_y,recvbuf_y,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(4,1,1,0,dvcRecvList_y,recvCount_y,recvCount_y,recvbuf_y,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(4,-1,1,0,dvcRecvList_y,2*recvCount_y,recvCount_y,recvbuf_y,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(8,0,1,1,dvcRecvList_y,3*recvCount_y,recvCount_y,recvbuf_y,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(9,0,1,-1,dvcRecvList_y,4*recvCount_y,recvCount_y,recvbuf_y,f_even,Nx,Ny,Nz);
//...................................................................................
//...Map recieve list for the z face<<<6,12,13,16,17)..............................................
UnpackDist(2,0,0,-1,dvcRecvList_Z,0,recvCount_Z,recvbuf_Z,f_odd,Nx,Ny,Nz);
UnpackDist(5,-1,0,-1,dvcRecvList_Z,recvCount_Z,recvCount_Z,recvbuf_Z,f_odd,Nx,Ny,Nz);
UnpackDist(7,1,0,-1,dvcRecvList_Z,2*recvCount_Z,recvCount_Z,recvbuf_Z,f_even,Nx,Ny,Nz);
UnpackDist(7,0,-1,-1,dvcRecvList_Z,3*recvCount_Z,recvCount_Z,recvbuf_Z,f_odd,Nx,Ny,Nz);
UnpackDist(9,0,1,-1,dvcRecvList_Z,4*recvCount_Z,recvCount_Z,recvbuf_Z,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(2,0,0,-1,dvcRecvList_Z,0,recvCount_Z,recvbuf_Z,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(5,-1,0,-1,dvcRecvList_Z,recvCount_Z,recvCount_Z,recvbuf_Z,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(7,1,0,-1,dvcRecvList_Z,2*recvCount_Z,recvCount_Z,recvbuf_Z,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(7,0,-1,-1,dvcRecvList_Z,3*recvCount_Z,recvCount_Z,recvbuf_Z,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(9,0,1,-1,dvcRecvList_Z,4*recvCount_Z,recvCount_Z,recvbuf_Z,f_even,Nx,Ny,Nz);
//...Map recieve list for the Z face<<<5,11,14,15,18)..............................................
UnpackDist(3,0,0,1,dvcRecvList_z,0,recvCount_z,recvbuf_z,f_even,Nx,Ny,Nz);
UnpackDist(6,1,0,1,dvcRecvList_z,recvCount_z,recvCount_z,recvbuf_z,f_even,Nx,Ny,Nz);
UnpackDist(6,-1,0,1,dvcRecvList_z,2*recvCount_z,recvCount_z,recvbuf_z,f_odd,Nx,Ny,Nz);
UnpackDist(8,0,1,1,dvcRecvList_z,3*recvCount_z,recvCount_z,recvbuf_z,f_even,Nx,Ny,Nz);
UnpackDist(8,0,-1,1,dvcRecvList_z,4*recvCount_z,recvCount_z,recvbuf_z,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(3,0,0,1,dvcRecvList_z,0,recvCount_z,recvbuf_z,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(6,1,0,1,dvcRecvList_z,recvCount_z,recvCount_z,recvbuf_z,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(6,-1,0,1,dvcRecvList_z,2*recvCount_z,recvCount_z,recvbuf_z,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(8,0,1,1,dvcRecvList_z,3*recvCount_z,recvCount_z,recvbuf_z,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(8,0,-1,1,dvcRecvList_z,4*recvCount_z,recvCount_z,recvbuf_z,f_odd,Nx,Ny,Nz);
//..................................................................................
//...Map recieve list for the xy edge <<<8)................................
UnpackDist(3,-1,-1,0,dvcRecvList_XY,0,recvCount_XY,recvbuf_XY,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(3,-1,-1,0,dvcRecvList_XY,0,recvCount_XY,recvbuf_XY,f_odd,Nx,Ny,Nz);
//...Map recieve list for the Xy edge <<<9)................................
UnpackDist(5,1,-1,0,dvcRecvList_xY,0,recvCount_xY,recvbuf_xY,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(5,1,-1,0,dvcRecvList_xY,0,recvCount_xY,recvbuf_xY,f_even,Nx,Ny,Nz);
//...Map recieve list for the xY edge <<<10)................................
UnpackDist(4,-1,1,0,dvcRecvList_Xy,0,recvCount_Xy,recvbuf_Xy,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(4,-1,1,0,dvcRecvList_Xy,0,recvCount_Xy,recvbuf_Xy,f_odd,Nx,Ny,Nz);
//...Map recieve list for the XY edge <<<7)................................
UnpackDist(4,1,1,0,dvcRecvList_xy,0,recvCount_xy,recvbuf_xy,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(4,1,1,0,dvcRecvList_xy,0,recvCount_xy,recvbuf_xy,f_even,Nx,Ny,Nz);
//...Map recieve list for the xz edge <<<12)................................
UnpackDist(5,-1,0,-1,dvcRecvList_XZ,0,recvCount_XZ,recvbuf_XZ,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(5,-1,0,-1,dvcRecvList_XZ,0,recvCount_XZ,recvbuf_XZ,f_odd,Nx,Ny,Nz);
//...Map recieve list for the xZ edge <<<14)................................
UnpackDist(6,-1,0,1,dvcRecvList_Xz,0,recvCount_Xz,recvbuf_Xz,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(6,-1,0,1,dvcRecvList_Xz,0,recvCount_Xz,recvbuf_Xz,f_odd,Nx,Ny,Nz);
//...Map recieve list for the Xz edge <<<13)................................
UnpackDist(7,1,0,-1,dvcRecvList_xZ,0,recvCount_xZ,recvbuf_xZ,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(7,1,0,-1,dvcRecvList_xZ,0,recvCount_xZ,recvbuf_xZ,f_even,Nx,Ny,Nz);
//...Map recieve list for the XZ edge <<<11)................................
UnpackDist(6,1,0,1,dvcRecvList_xz,0,recvCount_xz,recvbuf_xz,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(6,1,0,1,dvcRecvList_xz,0,recvCount_xz,recvbuf_xz,f_even,Nx,Ny,Nz);
//...Map recieve list for the yz edge <<<16)................................
UnpackDist(7,0,-1,-1,dvcRecvList_YZ,0,recvCount_YZ,recvbuf_YZ,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(7,0,-1,-1,dvcRecvList_YZ,0,recvCount_YZ,recvbuf_YZ,f_odd,Nx,Ny,Nz);
//...Map recieve list for the yZ edge <<<18)................................
UnpackDist(8,0,-1,1,dvcRecvList_Yz,0,recvCount_Yz,recvbuf_Yz,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(8,0,-1,1,dvcRecvList_Yz,0,recvCount_Yz,recvbuf_Yz,f_odd,Nx,Ny,Nz);
//...Map recieve list for the Yz edge <<<17)................................
UnpackDist(9,0,1,-1,dvcRecvList_yZ,0,recvCount_yZ,recvbuf_yZ,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(9,0,1,-1,dvcRecvList_yZ,0,recvCount_yZ,recvbuf_yZ,f_even,Nx,Ny,Nz);
//...Map recieve list for the YZ edge <<<15)................................
UnpackDist(8,0,1,1,dvcRecvList_yz,0,recvCount_yz,recvbuf_yz,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(8,0,1,1,dvcRecvList_yz,0,recvCount_yz,recvbuf_yz,f_even,Nx,Ny,Nz);
//...................................................................................
Lock=false; // unlock the communicator after communications complete
//...................................................................................
@ -573,27 +567,27 @@ void ScaLBL_Communicator::BiSendD3Q7(double *A_even, double *A_odd, double *B_ev
else{
Lock=true;
}
DeviceBarrier();
ScaLBL_DeviceBarrier();
//...................................................................................
sendtag = recvtag = 7;
//...................................................................................
PackDist(1,dvcSendList_x,0,sendCount_x,sendbuf_x,A_even,N);
PackDist(1,dvcSendList_x,sendCount_x,sendCount_x,sendbuf_x,B_even,N);
ScaLBL_D3Q19_Pack(1,dvcSendList_x,0,sendCount_x,sendbuf_x,A_even,N);
ScaLBL_D3Q19_Pack(1,dvcSendList_x,sendCount_x,sendCount_x,sendbuf_x,B_even,N);
//...Packing for X face(1,7,9,11,13)................................
PackDist(0,dvcSendList_X,0,sendCount_X,sendbuf_X,A_odd,N);
PackDist(0,dvcSendList_X,sendCount_X,sendCount_X,sendbuf_X,B_odd,N);
ScaLBL_D3Q19_Pack(0,dvcSendList_X,0,sendCount_X,sendbuf_X,A_odd,N);
ScaLBL_D3Q19_Pack(0,dvcSendList_X,sendCount_X,sendCount_X,sendbuf_X,B_odd,N);
//...Packing for y face(4,8,9,16,18).................................
PackDist(2,dvcSendList_y,0,sendCount_y,sendbuf_y,A_even,N);
PackDist(2,dvcSendList_y,sendCount_y,sendCount_y,sendbuf_y,B_even,N);
ScaLBL_D3Q19_Pack(2,dvcSendList_y,0,sendCount_y,sendbuf_y,A_even,N);
ScaLBL_D3Q19_Pack(2,dvcSendList_y,sendCount_y,sendCount_y,sendbuf_y,B_even,N);
//...Packing for Y face(3,7,10,15,17).................................
PackDist(1,dvcSendList_Y,0,sendCount_Y,sendbuf_Y,A_odd,N);
PackDist(1,dvcSendList_Y,sendCount_Y,sendCount_Y,sendbuf_Y,B_odd,N);
ScaLBL_D3Q19_Pack(1,dvcSendList_Y,0,sendCount_Y,sendbuf_Y,A_odd,N);
ScaLBL_D3Q19_Pack(1,dvcSendList_Y,sendCount_Y,sendCount_Y,sendbuf_Y,B_odd,N);
//...Packing for z face(6,12,13,16,17)................................
PackDist(3,dvcSendList_z,0,sendCount_z,sendbuf_z,A_even,N);
PackDist(3,dvcSendList_z,sendCount_z,sendCount_z,sendbuf_z,B_even,N);
ScaLBL_D3Q19_Pack(3,dvcSendList_z,0,sendCount_z,sendbuf_z,A_even,N);
ScaLBL_D3Q19_Pack(3,dvcSendList_z,sendCount_z,sendCount_z,sendbuf_z,B_even,N);
//...Packing for Z face(5,11,14,15,18)................................
PackDist(2,dvcSendList_Z,0,sendCount_Z,sendbuf_Z,A_odd,N);
PackDist(2,dvcSendList_Z,sendCount_Z,sendCount_Z,sendbuf_Z,B_odd,N);
ScaLBL_D3Q19_Pack(2,dvcSendList_Z,0,sendCount_Z,sendbuf_Z,A_odd,N);
ScaLBL_D3Q19_Pack(2,dvcSendList_Z,sendCount_Z,sendCount_Z,sendbuf_Z,B_odd,N);
//...................................................................................
//...................................................................................
@ -617,32 +611,32 @@ void ScaLBL_Communicator::BiRecvD3Q7(double *A_even, double *A_odd, double *B_ev
// Wait for completion of D3Q19 communication
MPI_Waitall(6,req1,stat1);
MPI_Waitall(6,req2,stat2);
DeviceBarrier();
ScaLBL_DeviceBarrier();
//...................................................................................
// Unpack the distributions on the device
//...................................................................................
//...Map recieve list for the X face: q=2,8,10,12,13 ................................
UnpackDist(0,-1,0,0,dvcRecvList_X,0,recvCount_X,recvbuf_X,A_odd,Nx,Ny,Nz);
UnpackDist(0,-1,0,0,dvcRecvList_X,recvCount_X,recvCount_X,recvbuf_X,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(0,-1,0,0,dvcRecvList_X,0,recvCount_X,recvbuf_X,A_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(0,-1,0,0,dvcRecvList_X,recvCount_X,recvCount_X,recvbuf_X,B_odd,Nx,Ny,Nz);
//...................................................................................
//...Map recieve list for the x face: q=1,7,9,11,13..................................
UnpackDist(1,1,0,0,dvcRecvList_x,0,recvCount_x,recvbuf_x,A_even,Nx,Ny,Nz);
UnpackDist(1,1,0,0,dvcRecvList_x,recvCount_x,recvCount_x,recvbuf_x,B_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(1,1,0,0,dvcRecvList_x,0,recvCount_x,recvbuf_x,A_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(1,1,0,0,dvcRecvList_x,recvCount_x,recvCount_x,recvbuf_x,B_even,Nx,Ny,Nz);
//...................................................................................
//...Map recieve list for the y face: q=4,8,9,16,18 .................................
UnpackDist(1,0,-1,0,dvcRecvList_Y,0,recvCount_Y,recvbuf_Y,A_odd,Nx,Ny,Nz);
UnpackDist(1,0,-1,0,dvcRecvList_Y,recvCount_Y,recvCount_Y,recvbuf_Y,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(1,0,-1,0,dvcRecvList_Y,0,recvCount_Y,recvbuf_Y,A_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(1,0,-1,0,dvcRecvList_Y,recvCount_Y,recvCount_Y,recvbuf_Y,B_odd,Nx,Ny,Nz);
//...................................................................................
//...Map recieve list for the Y face: q=3,7,10,15,17 ................................
UnpackDist(2,0,1,0,dvcRecvList_y,0,recvCount_y,recvbuf_y,A_even,Nx,Ny,Nz);
UnpackDist(2,0,1,0,dvcRecvList_y,recvCount_y,recvCount_y,recvbuf_y,B_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(2,0,1,0,dvcRecvList_y,0,recvCount_y,recvbuf_y,A_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(2,0,1,0,dvcRecvList_y,recvCount_y,recvCount_y,recvbuf_y,B_even,Nx,Ny,Nz);
//...................................................................................
//...Map recieve list for the z face<<<6,12,13,16,17)................................
UnpackDist(2,0,0,-1,dvcRecvList_Z,0,recvCount_Z,recvbuf_Z,A_odd,Nx,Ny,Nz);
UnpackDist(2,0,0,-1,dvcRecvList_Z,recvCount_Z,recvCount_Z,recvbuf_Z,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(2,0,0,-1,dvcRecvList_Z,0,recvCount_Z,recvbuf_Z,A_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(2,0,0,-1,dvcRecvList_Z,recvCount_Z,recvCount_Z,recvbuf_Z,B_odd,Nx,Ny,Nz);
//...Map recieve list for the Z face<<<5,11,14,15,18)................................
UnpackDist(3,0,0,1,dvcRecvList_z,0,recvCount_z,recvbuf_z,A_even,Nx,Ny,Nz);
UnpackDist(3,0,0,1,dvcRecvList_z,recvCount_z,recvCount_z,recvbuf_z,B_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(3,0,0,1,dvcRecvList_z,0,recvCount_z,recvbuf_z,A_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(3,0,0,1,dvcRecvList_z,recvCount_z,recvCount_z,recvbuf_z,B_even,Nx,Ny,Nz);
//..................................................................................
Lock=false; // unlock the communicator after communications complete
//...................................................................................
@ -655,28 +649,28 @@ void ScaLBL_Communicator::SendHalo(double *data){
else{
Lock=true;
}
DeviceBarrier();
ScaLBL_DeviceBarrier();
//...................................................................................
sendtag = recvtag = 1;
//...................................................................................
PackValues(dvcSendList_x, sendCount_x,sendbuf_x, data, N);
PackValues(dvcSendList_y, sendCount_y,sendbuf_y, data, N);
PackValues(dvcSendList_z, sendCount_z,sendbuf_z, data, N);
PackValues(dvcSendList_X, sendCount_X,sendbuf_X, data, N);
PackValues(dvcSendList_Y, sendCount_Y,sendbuf_Y, data, N);
PackValues(dvcSendList_Z, sendCount_Z,sendbuf_Z, data, N);
PackValues(dvcSendList_xy, sendCount_xy,sendbuf_xy, data, N);
PackValues(dvcSendList_xY, sendCount_xY,sendbuf_xY, data, N);
PackValues(dvcSendList_Xy, sendCount_Xy,sendbuf_Xy, data, N);
PackValues(dvcSendList_XY, sendCount_XY,sendbuf_XY, data, N);
PackValues(dvcSendList_xz, sendCount_xz,sendbuf_xz, data, N);
PackValues(dvcSendList_xZ, sendCount_xZ,sendbuf_xZ, data, N);
PackValues(dvcSendList_Xz, sendCount_Xz,sendbuf_Xz, data, N);
PackValues(dvcSendList_XZ, sendCount_XZ,sendbuf_XZ, data, N);
PackValues(dvcSendList_yz, sendCount_yz,sendbuf_yz, data, N);
PackValues(dvcSendList_yZ, sendCount_yZ,sendbuf_yZ, data, N);
PackValues(dvcSendList_Yz, sendCount_Yz,sendbuf_Yz, data, N);
PackValues(dvcSendList_YZ, sendCount_YZ,sendbuf_YZ, data, N);
ScaLBL_Scalar_Pack(dvcSendList_x, sendCount_x,sendbuf_x, data, N);
ScaLBL_Scalar_Pack(dvcSendList_y, sendCount_y,sendbuf_y, data, N);
ScaLBL_Scalar_Pack(dvcSendList_z, sendCount_z,sendbuf_z, data, N);
ScaLBL_Scalar_Pack(dvcSendList_X, sendCount_X,sendbuf_X, data, N);
ScaLBL_Scalar_Pack(dvcSendList_Y, sendCount_Y,sendbuf_Y, data, N);
ScaLBL_Scalar_Pack(dvcSendList_Z, sendCount_Z,sendbuf_Z, data, N);
ScaLBL_Scalar_Pack(dvcSendList_xy, sendCount_xy,sendbuf_xy, data, N);
ScaLBL_Scalar_Pack(dvcSendList_xY, sendCount_xY,sendbuf_xY, data, N);
ScaLBL_Scalar_Pack(dvcSendList_Xy, sendCount_Xy,sendbuf_Xy, data, N);
ScaLBL_Scalar_Pack(dvcSendList_XY, sendCount_XY,sendbuf_XY, data, N);
ScaLBL_Scalar_Pack(dvcSendList_xz, sendCount_xz,sendbuf_xz, data, N);
ScaLBL_Scalar_Pack(dvcSendList_xZ, sendCount_xZ,sendbuf_xZ, data, N);
ScaLBL_Scalar_Pack(dvcSendList_Xz, sendCount_Xz,sendbuf_Xz, data, N);
ScaLBL_Scalar_Pack(dvcSendList_XZ, sendCount_XZ,sendbuf_XZ, data, N);
ScaLBL_Scalar_Pack(dvcSendList_yz, sendCount_yz,sendbuf_yz, data, N);
ScaLBL_Scalar_Pack(dvcSendList_yZ, sendCount_yZ,sendbuf_yZ, data, N);
ScaLBL_Scalar_Pack(dvcSendList_Yz, sendCount_Yz,sendbuf_Yz, data, N);
ScaLBL_Scalar_Pack(dvcSendList_YZ, sendCount_YZ,sendbuf_YZ, data, N);
//...................................................................................
// Send / Recv all the phase indcator field values
//...................................................................................
@ -723,27 +717,27 @@ void ScaLBL_Communicator::RecvHalo(double *data){
//...................................................................................
MPI_Waitall(18,req1,stat1);
MPI_Waitall(18,req2,stat2);
DeviceBarrier();
ScaLBL_DeviceBarrier();
//...................................................................................
//...................................................................................
UnpackValues(dvcRecvList_x, recvCount_x,recvbuf_x, data, N);
UnpackValues(dvcRecvList_y, recvCount_y,recvbuf_y, data, N);
UnpackValues(dvcRecvList_z, recvCount_z,recvbuf_z, data, N);
UnpackValues(dvcRecvList_X, recvCount_X,recvbuf_X, data, N);
UnpackValues(dvcRecvList_Y, recvCount_Y,recvbuf_Y, data, N);
UnpackValues(dvcRecvList_Z, recvCount_Z,recvbuf_Z, data, N);
UnpackValues(dvcRecvList_xy, recvCount_xy,recvbuf_xy, data, N);
UnpackValues(dvcRecvList_xY, recvCount_xY,recvbuf_xY, data, N);
UnpackValues(dvcRecvList_Xy, recvCount_Xy,recvbuf_Xy, data, N);
UnpackValues(dvcRecvList_XY, recvCount_XY,recvbuf_XY, data, N);
UnpackValues(dvcRecvList_xz, recvCount_xz,recvbuf_xz, data, N);
UnpackValues(dvcRecvList_xZ, recvCount_xZ,recvbuf_xZ, data, N);
UnpackValues(dvcRecvList_Xz, recvCount_Xz,recvbuf_Xz, data, N);
UnpackValues(dvcRecvList_XZ, recvCount_XZ,recvbuf_XZ, data, N);
UnpackValues(dvcRecvList_yz, recvCount_yz,recvbuf_yz, data, N);
UnpackValues(dvcRecvList_yZ, recvCount_yZ,recvbuf_yZ, data, N);
UnpackValues(dvcRecvList_Yz, recvCount_Yz,recvbuf_Yz, data, N);
UnpackValues(dvcRecvList_YZ, recvCount_YZ,recvbuf_YZ, data, N);
ScaLBL_Scalar_Unpack(dvcRecvList_x, recvCount_x,recvbuf_x, data, N);
ScaLBL_Scalar_Unpack(dvcRecvList_y, recvCount_y,recvbuf_y, data, N);
ScaLBL_Scalar_Unpack(dvcRecvList_z, recvCount_z,recvbuf_z, data, N);
ScaLBL_Scalar_Unpack(dvcRecvList_X, recvCount_X,recvbuf_X, data, N);
ScaLBL_Scalar_Unpack(dvcRecvList_Y, recvCount_Y,recvbuf_Y, data, N);
ScaLBL_Scalar_Unpack(dvcRecvList_Z, recvCount_Z,recvbuf_Z, data, N);
ScaLBL_Scalar_Unpack(dvcRecvList_xy, recvCount_xy,recvbuf_xy, data, N);
ScaLBL_Scalar_Unpack(dvcRecvList_xY, recvCount_xY,recvbuf_xY, data, N);
ScaLBL_Scalar_Unpack(dvcRecvList_Xy, recvCount_Xy,recvbuf_Xy, data, N);
ScaLBL_Scalar_Unpack(dvcRecvList_XY, recvCount_XY,recvbuf_XY, data, N);
ScaLBL_Scalar_Unpack(dvcRecvList_xz, recvCount_xz,recvbuf_xz, data, N);
ScaLBL_Scalar_Unpack(dvcRecvList_xZ, recvCount_xZ,recvbuf_xZ, data, N);
ScaLBL_Scalar_Unpack(dvcRecvList_Xz, recvCount_Xz,recvbuf_Xz, data, N);
ScaLBL_Scalar_Unpack(dvcRecvList_XZ, recvCount_XZ,recvbuf_XZ, data, N);
ScaLBL_Scalar_Unpack(dvcRecvList_yz, recvCount_yz,recvbuf_yz, data, N);
ScaLBL_Scalar_Unpack(dvcRecvList_yZ, recvCount_yZ,recvbuf_yZ, data, N);
ScaLBL_Scalar_Unpack(dvcRecvList_Yz, recvCount_Yz,recvbuf_Yz, data, N);
ScaLBL_Scalar_Unpack(dvcRecvList_YZ, recvCount_YZ,recvbuf_YZ, data, N);
//...................................................................................
Lock=false; // unlock the communicator after communications complete
//...................................................................................

115
tests/Blob.cpp Normal file
View File

@ -0,0 +1,115 @@
#include <stdio.h>
#include <stdlib.h>
#include <iostream>
#include <fstream>
#define MAX_LOCAL_BLOB_COUNT 500
using namespace std;
struct BlobInfo{
BlobInfo(int size){
Values = new double [40*size];
Ranks = new int [19*size];
RankLabel = new int [19*size];
};
double * Values;
int * Ranks;
int * RankLabel;
// local rank and blob_id
int rank(int i){return Ranks[19*i];}
int rank_x(int i){return Ranks[19*i+1];}
int rank_y(int i){return Ranks[19*i+2];}
int rank_z(int i){return Ranks[19*i+3];}
int rank_X(int i){return Ranks[19*i+4];}
int rank_Y(int i){return Ranks[19*i+5];}
int rank_Z(int i){return Ranks[19*i+6];}
int rank_xy(int i){return Ranks[19*i+7];}
int rank_xY(int i){return Ranks[19*i+8];}
int rank_Xy(int i){return Ranks[19*i+9];}
int rank_XY(int i){return Ranks[19*i+10];}
int rank_xz(int i){return Ranks[19*i+11];}
int rank_xZ(int i){return Ranks[19*i+12];}
int rank_Xz(int i){return Ranks[19*i+13];}
int rank_XZ(int i){return Ranks[19*i+14];}
int rank_yz(int i){return Ranks[19*i+15];}
int rank_yZ(int i){return Ranks[19*i+16];}
int rank_Yz(int i){return Ranks[19*i+17];}
int rank_YZ(int i){return Ranks[19*i+18];}
int blobID(int i){return RankLabel[19*i];}
int blobID_x(int i){return RankLabel[19*i+1];}
int blobID_y(int i){return RankLabel[19*i+2];}
int blobID_z(int i){return RankLabel[19*i+3];}
int blobID_X(int i){return RankLabel[19*i+4];}
int blobID_Y(int i){return RankLabel[19*i+5];}
int blobID_Z(int i){return RankLabel[19*i+6];}
int blobID_xy(int i){return RankLabel[19*i+7];}
int blobID_xY(int i){return RankLabel[19*i+8];}
int blobID_Xy(int i){return RankLabel[19*i+9];}
int blobID_XY(int i){return RankLabel[19*i+10];}
int blobID_xz(int i){return RankLabel[19*i+11];}
int blobID_xZ(int i){return RankLabel[19*i+12];}
int blobID_Xz(int i){return RankLabel[19*i+13];}
int blobID_XZ(int i){return RankLabel[19*i+14];}
int blobID_yz(int i){return RankLabel[19*i+15];}
int blobID_yZ(int i){return RankLabel[19*i+16];}
int blobID_Yz(int i){return RankLabel[19*i+17];}
int blobID_YZ(int i){return RankLabel[19*i+18];}
double voln(int i){return Values[40*i];}
double pn(int i){return Values[40*i+1];}
double vnx(int i){return Values[40*i+2];}
double vny(int i){return Values[40*i+3];}
double vnz(int i){return Values[40*i+4];}
double vwnx(int i){return Values[40*i+5];}
double vwny(int i){return Values[40*i+6];}
double vwnz(int i){return Values[40*i+7];}
double Jwn(int i){return Values[40*i+8];}
double Kwn(int i){return Values[40*i+9];}
double awn(int i){return Values[40*i+10];}
double ans(int i){return Values[40*i+11];}
double Gwnxx(int i){return Values[40*i+12];}
double Gwnyy(int i){return Values[40*i+13];}
double Gwnzz(int i){return Values[40*i+14];}
double Gwnxy(int i){return Values[40*i+15];}
double Gwnxz(int i){return Values[40*i+16];}
double Gwnyz(int i){return Values[40*i+17];}
double Gnsxx(int i){return Values[40*i+18];}
double Gnsyy(int i){return Values[40*i+19];}
double Gnszz(int i){return Values[40*i+20];}
double Gnsxy(int i){return Values[40*i+21];}
double Gnsxz(int i){return Values[40*i+22];}
double Gnsyz(int i){return Values[40*i+23];}
double lwns(int i){return Values[40*i+24];}
double cospwns(int i){return Values[40*i+25];}
double vwnsx(int i){return Values[40*i+26];}
double vwnsy(int i){return Values[40*i+27];}
double vwnsz(int i){return Values[40*i+28];}
double cx(int i){return Values[40*i+29];}
double cy(int i){return Values[40*i+30];}
double cz(int i){return Values[40*i+31];}
};
int main(void){
int Nx,Ny,Nz,N;
Nx = Ny = Nz = 100;
N = Nx*Ny*Nz;
BlobInfo Blobs(MAX_LOCAL_BLOB_COUNT);
// LocalBlobs = new BlobInfo[MAX_LOCAL_BLOB_COUNT];
double * spheres;
spheres = new double [N];
int nblobs = 10;
for (int i=0; i<nblobs; i++){
Blobs.rank(i) = 0;
}
}

728
tests/TestBubble.cpp
View File

@ -654,43 +654,43 @@ int main(int argc, char **argv)
double *recvbuf_xy, *recvbuf_yz, *recvbuf_xz, *recvbuf_Xy, *recvbuf_Yz, *recvbuf_xZ;
double *recvbuf_xY, *recvbuf_yZ, *recvbuf_Xz, *recvbuf_XY, *recvbuf_YZ, *recvbuf_XZ;
//......................................................................................
AllocateDeviceMemory((void **) &sendbuf_x, 5*sendCount_x*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_X, 5*sendCount_X*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_y, 5*sendCount_y*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_Y, 5*sendCount_Y*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_z, 5*sendCount_z*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_Z, 5*sendCount_Z*sizeof(double)); // Allocatevoid * memory
AllocateDeviceMemory((void **) &sendbuf_xy, sendCount_xy*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_xY, sendCount_xY*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_Xy, sendCount_Xy*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_XY, sendCount_XY*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_xz, sendCount_xz*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_xZ, sendCount_xZ*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_Xz, sendCount_Xz*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_XZ, sendCount_XZ*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_yz, sendCount_yz*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_yZ, sendCount_yZ*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_Yz, sendCount_Yz*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &sendbuf_YZ, sendCount_YZ*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &sendbuf_x, 5*sendCount_x*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &sendbuf_X, 5*sendCount_X*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &sendbuf_y, 5*sendCount_y*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &sendbuf_Y, 5*sendCount_Y*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &sendbuf_z, 5*sendCount_z*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &sendbuf_Z, 5*sendCount_Z*sizeof(double)); // Allocatevoid * memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &sendbuf_xy, sendCount_xy*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &sendbuf_xY, sendCount_xY*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &sendbuf_Xy, sendCount_Xy*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &sendbuf_XY, sendCount_XY*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &sendbuf_xz, sendCount_xz*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &sendbuf_xZ, sendCount_xZ*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &sendbuf_Xz, sendCount_Xz*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &sendbuf_XZ, sendCount_XZ*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &sendbuf_yz, sendCount_yz*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &sendbuf_yZ, sendCount_yZ*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &sendbuf_Yz, sendCount_Yz*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &sendbuf_YZ, sendCount_YZ*sizeof(double)); // Allocate device memory
//......................................................................................
AllocateDeviceMemory((void **) &recvbuf_x, 5*recvCount_x*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_X, 5*recvCount_X*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_y, 5*recvCount_y*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_Y, 5*recvCount_Y*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_z, 5*recvCount_z*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_Z, 5*recvCount_Z*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_xy, recvCount_xy*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_xY, recvCount_xY*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_Xy, recvCount_Xy*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_XY, recvCount_XY*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_xz, recvCount_xz*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_xZ, recvCount_xZ*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_Xz, recvCount_Xz*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_XZ, recvCount_XZ*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_yz, recvCount_yz*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_yZ, recvCount_yZ*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_Yz, recvCount_Yz*sizeof(double)); // Allocate device memory
AllocateDeviceMemory((void **) &recvbuf_YZ, recvCount_YZ*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &recvbuf_x, 5*recvCount_x*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &recvbuf_X, 5*recvCount_X*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &recvbuf_y, 5*recvCount_y*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &recvbuf_Y, 5*recvCount_Y*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &recvbuf_z, 5*recvCount_z*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &recvbuf_Z, 5*recvCount_Z*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &recvbuf_xy, recvCount_xy*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &recvbuf_xY, recvCount_xY*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &recvbuf_Xy, recvCount_Xy*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &recvbuf_XY, recvCount_XY*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &recvbuf_xz, recvCount_xz*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &recvbuf_xZ, recvCount_xZ*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &recvbuf_Xz, recvCount_Xz*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &recvbuf_XZ, recvCount_XZ*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &recvbuf_yz, recvCount_yz*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &recvbuf_yZ, recvCount_yZ*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &recvbuf_Yz, recvCount_Yz*sizeof(double)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &recvbuf_YZ, recvCount_YZ*sizeof(double)); // Allocate device memory
//......................................................................................
int *dvcSendList_x, *dvcSendList_y, *dvcSendList_z, *dvcSendList_X, *dvcSendList_Y, *dvcSendList_Z;
int *dvcSendList_xy, *dvcSendList_yz, *dvcSendList_xz, *dvcSendList_Xy, *dvcSendList_Yz, *dvcSendList_xZ;
@ -700,83 +700,83 @@ int main(int argc, char **argv)
int *dvcRecvList_xy, *dvcRecvList_yz, *dvcRecvList_xz, *dvcRecvList_Xy, *dvcRecvList_Yz, *dvcRecvList_xZ;
int *dvcRecvList_xY, *dvcRecvList_yZ, *dvcRecvList_Xz, *dvcRecvList_XY, *dvcRecvList_YZ, *dvcRecvList_XZ;
//......................................................................................
AllocateDeviceMemory((void **) &dvcSendList_x, sendCount_x*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_X, sendCount_X*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_y, sendCount_y*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_Y, sendCount_Y*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_z, sendCount_z*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_Z, sendCount_Z*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_xy, sendCount_xy*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_xY, sendCount_xY*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_Xy, sendCount_Xy*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_XY, sendCount_XY*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_xz, sendCount_xz*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_xZ, sendCount_xZ*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_Xz, sendCount_Xz*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_XZ, sendCount_XZ*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_yz, sendCount_yz*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_yZ, sendCount_yZ*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_Yz, sendCount_Yz*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcSendList_YZ, sendCount_YZ*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_x, sendCount_x*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_X, sendCount_X*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_y, sendCount_y*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_Y, sendCount_Y*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_z, sendCount_z*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_Z, sendCount_Z*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_xy, sendCount_xy*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_xY, sendCount_xY*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_Xy, sendCount_Xy*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_XY, sendCount_XY*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_xz, sendCount_xz*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_xZ, sendCount_xZ*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_Xz, sendCount_Xz*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_XZ, sendCount_XZ*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_yz, sendCount_yz*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_yZ, sendCount_yZ*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_Yz, sendCount_Yz*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcSendList_YZ, sendCount_YZ*sizeof(int)); // Allocate device memory
//......................................................................................
AllocateDeviceMemory((void **) &dvcRecvList_x, recvCount_x*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_X, recvCount_X*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_y, recvCount_y*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_Y, recvCount_Y*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_z, recvCount_z*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_Z, recvCount_Z*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_xy, recvCount_xy*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_xY, recvCount_xY*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_Xy, recvCount_Xy*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_XY, recvCount_XY*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_xz, recvCount_xz*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_xZ, recvCount_xZ*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_Xz, recvCount_Xz*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_XZ, recvCount_XZ*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_yz, recvCount_yz*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_yZ, recvCount_yZ*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_Yz, recvCount_Yz*sizeof(int)); // Allocate device memory
AllocateDeviceMemory((void **) &dvcRecvList_YZ, recvCount_YZ*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_x, recvCount_x*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_X, recvCount_X*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_y, recvCount_y*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_Y, recvCount_Y*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_z, recvCount_z*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_Z, recvCount_Z*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_xy, recvCount_xy*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_xY, recvCount_xY*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_Xy, recvCount_Xy*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_XY, recvCount_XY*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_xz, recvCount_xz*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_xZ, recvCount_xZ*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_Xz, recvCount_Xz*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_XZ, recvCount_XZ*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_yz, recvCount_yz*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_yZ, recvCount_yZ*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_Yz, recvCount_Yz*sizeof(int)); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &dvcRecvList_YZ, recvCount_YZ*sizeof(int)); // Allocate device memory
//......................................................................................
MPI_Barrier(comm);
if (rank==0) printf ("Prepare to copy send/recv Lists to device \n");
CopyToDevice(dvcSendList_x,sendList_x,sendCount_x*sizeof(int));
CopyToDevice(dvcSendList_X,sendList_X,sendCount_X*sizeof(int));
CopyToDevice(dvcSendList_y,sendList_y,sendCount_y*sizeof(int));
CopyToDevice(dvcSendList_Y,sendList_Y,sendCount_Y*sizeof(int));
CopyToDevice(dvcSendList_z,sendList_z,sendCount_z*sizeof(int));
CopyToDevice(dvcSendList_Z,sendList_Z,sendCount_Z*sizeof(int));
CopyToDevice(dvcSendList_xy,sendList_xy,sendCount_xy*sizeof(int));
CopyToDevice(dvcSendList_XY,sendList_XY,sendCount_XY*sizeof(int));
CopyToDevice(dvcSendList_xY,sendList_xY,sendCount_xY*sizeof(int));
CopyToDevice(dvcSendList_Xy,sendList_Xy,sendCount_Xy*sizeof(int));
CopyToDevice(dvcSendList_xz,sendList_xz,sendCount_xz*sizeof(int));
CopyToDevice(dvcSendList_XZ,sendList_XZ,sendCount_XZ*sizeof(int));
CopyToDevice(dvcSendList_xZ,sendList_xZ,sendCount_xZ*sizeof(int));
CopyToDevice(dvcSendList_Xz,sendList_Xz,sendCount_Xz*sizeof(int));
CopyToDevice(dvcSendList_yz,sendList_yz,sendCount_yz*sizeof(int));
CopyToDevice(dvcSendList_YZ,sendList_YZ,sendCount_YZ*sizeof(int));
CopyToDevice(dvcSendList_yZ,sendList_yZ,sendCount_yZ*sizeof(int));
CopyToDevice(dvcSendList_Yz,sendList_Yz,sendCount_Yz*sizeof(int));
ScaLBL_CopyToDevice(dvcSendList_x,sendList_x,sendCount_x*sizeof(int));
ScaLBL_CopyToDevice(dvcSendList_X,sendList_X,sendCount_X*sizeof(int));
ScaLBL_CopyToDevice(dvcSendList_y,sendList_y,sendCount_y*sizeof(int));
ScaLBL_CopyToDevice(dvcSendList_Y,sendList_Y,sendCount_Y*sizeof(int));
ScaLBL_CopyToDevice(dvcSendList_z,sendList_z,sendCount_z*sizeof(int));
ScaLBL_CopyToDevice(dvcSendList_Z,sendList_Z,sendCount_Z*sizeof(int));
ScaLBL_CopyToDevice(dvcSendList_xy,sendList_xy,sendCount_xy*sizeof(int));
ScaLBL_CopyToDevice(dvcSendList_XY,sendList_XY,sendCount_XY*sizeof(int));
ScaLBL_CopyToDevice(dvcSendList_xY,sendList_xY,sendCount_xY*sizeof(int));
ScaLBL_CopyToDevice(dvcSendList_Xy,sendList_Xy,sendCount_Xy*sizeof(int));
ScaLBL_CopyToDevice(dvcSendList_xz,sendList_xz,sendCount_xz*sizeof(int));
ScaLBL_CopyToDevice(dvcSendList_XZ,sendList_XZ,sendCount_XZ*sizeof(int));
ScaLBL_CopyToDevice(dvcSendList_xZ,sendList_xZ,sendCount_xZ*sizeof(int));
ScaLBL_CopyToDevice(dvcSendList_Xz,sendList_Xz,sendCount_Xz*sizeof(int));
ScaLBL_CopyToDevice(dvcSendList_yz,sendList_yz,sendCount_yz*sizeof(int));
ScaLBL_CopyToDevice(dvcSendList_YZ,sendList_YZ,sendCount_YZ*sizeof(int));
ScaLBL_CopyToDevice(dvcSendList_yZ,sendList_yZ,sendCount_yZ*sizeof(int));
ScaLBL_CopyToDevice(dvcSendList_Yz,sendList_Yz,sendCount_Yz*sizeof(int));
//......................................................................................
CopyToDevice(dvcRecvList_x,recvList_x,recvCount_x*sizeof(int));
CopyToDevice(dvcRecvList_X,recvList_X,recvCount_X*sizeof(int));
CopyToDevice(dvcRecvList_y,recvList_y,recvCount_y*sizeof(int));
CopyToDevice(dvcRecvList_Y,recvList_Y,recvCount_Y*sizeof(int));
CopyToDevice(dvcRecvList_z,recvList_z,recvCount_z*sizeof(int));
CopyToDevice(dvcRecvList_Z,recvList_Z,recvCount_Z*sizeof(int));
CopyToDevice(dvcRecvList_xy,recvList_xy,recvCount_xy*sizeof(int));
CopyToDevice(dvcRecvList_XY,recvList_XY,recvCount_XY*sizeof(int));
CopyToDevice(dvcRecvList_xY,recvList_xY,recvCount_xY*sizeof(int));
CopyToDevice(dvcRecvList_Xy,recvList_Xy,recvCount_Xy*sizeof(int));
CopyToDevice(dvcRecvList_xz,recvList_xz,recvCount_xz*sizeof(int));
CopyToDevice(dvcRecvList_XZ,recvList_XZ,recvCount_XZ*sizeof(int));
CopyToDevice(dvcRecvList_xZ,recvList_xZ,recvCount_xZ*sizeof(int));
CopyToDevice(dvcRecvList_Xz,recvList_Xz,recvCount_Xz*sizeof(int));
CopyToDevice(dvcRecvList_yz,recvList_yz,recvCount_yz*sizeof(int));
CopyToDevice(dvcRecvList_YZ,recvList_YZ,recvCount_YZ*sizeof(int));
CopyToDevice(dvcRecvList_yZ,recvList_yZ,recvCount_yZ*sizeof(int));
CopyToDevice(dvcRecvList_Yz,recvList_Yz,recvCount_Yz*sizeof(int));
ScaLBL_CopyToDevice(dvcRecvList_x,recvList_x,recvCount_x*sizeof(int));
ScaLBL_CopyToDevice(dvcRecvList_X,recvList_X,recvCount_X*sizeof(int));
ScaLBL_CopyToDevice(dvcRecvList_y,recvList_y,recvCount_y*sizeof(int));
ScaLBL_CopyToDevice(dvcRecvList_Y,recvList_Y,recvCount_Y*sizeof(int));
ScaLBL_CopyToDevice(dvcRecvList_z,recvList_z,recvCount_z*sizeof(int));
ScaLBL_CopyToDevice(dvcRecvList_Z,recvList_Z,recvCount_Z*sizeof(int));
ScaLBL_CopyToDevice(dvcRecvList_xy,recvList_xy,recvCount_xy*sizeof(int));
ScaLBL_CopyToDevice(dvcRecvList_XY,recvList_XY,recvCount_XY*sizeof(int));
ScaLBL_CopyToDevice(dvcRecvList_xY,recvList_xY,recvCount_xY*sizeof(int));
ScaLBL_CopyToDevice(dvcRecvList_Xy,recvList_Xy,recvCount_Xy*sizeof(int));
ScaLBL_CopyToDevice(dvcRecvList_xz,recvList_xz,recvCount_xz*sizeof(int));
ScaLBL_CopyToDevice(dvcRecvList_XZ,recvList_XZ,recvCount_XZ*sizeof(int));
ScaLBL_CopyToDevice(dvcRecvList_xZ,recvList_xZ,recvCount_xZ*sizeof(int));
ScaLBL_CopyToDevice(dvcRecvList_Xz,recvList_Xz,recvCount_Xz*sizeof(int));
ScaLBL_CopyToDevice(dvcRecvList_yz,recvList_yz,recvCount_yz*sizeof(int));
ScaLBL_CopyToDevice(dvcRecvList_YZ,recvList_YZ,recvCount_YZ*sizeof(int));
ScaLBL_CopyToDevice(dvcRecvList_yZ,recvList_yZ,recvCount_yZ*sizeof(int));
ScaLBL_CopyToDevice(dvcRecvList_Yz,recvList_Yz,recvCount_Yz*sizeof(int));
//......................................................................................
// Fill in the phase ID from neighboring processors
char *sendID_x, *sendID_y, *sendID_z, *sendID_X, *sendID_Y, *sendID_Z;
@ -953,9 +953,9 @@ int main(int argc, char **argv)
//...........device phase ID.................................................
if (rank==0) printf ("Copying phase ID to device \n");
char *ID;
AllocateDeviceMemory((void **) &ID, N); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &ID, N); // Allocate device memory
// Copy to the device
CopyToDevice(ID, id, N);
ScaLBL_CopyToDevice(ID, id, N);
//...........................................................................
//...........................................................................
@ -969,23 +969,23 @@ int main(int argc, char **argv)
double *f_even,*f_odd;
double *A_even,*A_odd,*B_even,*B_odd;
//...........................................................................
AllocateDeviceMemory((void **) &f_even, 10*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &f_odd, 9*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &A_even, 4*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &A_odd, 3*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &B_even, 4*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &B_odd, 3*dist_mem_size); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &f_even, 10*dist_mem_size); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &f_odd, 9*dist_mem_size); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &A_even, 4*dist_mem_size); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &A_odd, 3*dist_mem_size); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &B_even, 4*dist_mem_size); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &B_odd, 3*dist_mem_size); // Allocate device memory
//...........................................................................
double *Phi,*Den;
// double *Copy;
double *ColorGrad, *Velocity, *Pressure;
//...........................................................................
AllocateDeviceMemory((void **) &Phi, dist_mem_size);
AllocateDeviceMemory((void **) &Pressure, dist_mem_size);
AllocateDeviceMemory((void **) &Den, 2*dist_mem_size);
// AllocateDeviceMemory((void **) &Copy, 2*dist_mem_size);
AllocateDeviceMemory((void **) &Velocity, 3*dist_mem_size);
AllocateDeviceMemory((void **) &ColorGrad, 3*dist_mem_size);
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &Phi, dist_mem_size);
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &Pressure, dist_mem_size);
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &Den, 2*dist_mem_size);
// ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &Copy, 2*dist_mem_size);
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &Velocity, 3*dist_mem_size);
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &ColorGrad, 3*dist_mem_size);
//...........................................................................
//...........................................................................
// Phase indicator (in array form as needed by PMMC algorithm)
@ -1192,16 +1192,16 @@ int main(int argc, char **argv)
}
}
// Copy the bubble to the device and initialize
CopyToDevice(ID, id, N);
ScaLBL_CopyToDevice(ID, id, N);
//...........................................................................
//...........................................................................
InitD3Q19(ID, f_even, f_odd, Nx, Ny, Nz);
ScaLBL_D3Q19_Init(ID, f_even, f_odd, Nx, Ny, Nz);
//......................................................................
// InitDenColorDistance(ID, Copy, Phi, SDs.data(), das, dbs, beta, xIntPos, Nx, Ny, Nz, S);
InitDenColorDistance(ID, Den, Phi, SDs.data(), das, dbs, beta, xIntPos, Nx, Ny, Nz);
InitD3Q7(ID, A_even, A_odd, &Den[0], Nx, Ny, Nz);
InitD3Q7(ID, B_even, B_odd, &Den[N], Nx, Ny, Nz);
// ScaLBL_Color_InitDistance(ID, Copy, Phi, SDs.data(), das, dbs, beta, xIntPos, Nx, Ny, Nz, S);
ScaLBL_Color_InitDistance(ID, Den, Phi, SDs.data(), das, dbs, beta, xIntPos, Nx, Ny, Nz);
ScaLBL_D3Q7_Init(ID, A_even, A_odd, &Den[0], Nx, Ny, Nz);
ScaLBL_D3Q7_Init(ID, B_even, B_odd, &Den[N], Nx, Ny, Nz);
//......................................................................
// Once phase has been initialized, map solid to account for 'smeared' interface
//......................................................................
@ -1219,40 +1219,40 @@ int main(int argc, char **argv)
// Read in the restart file to CPU buffers
ReadCheckpoint(LocalRestartFile, cDen, cDistEven, cDistOdd, N);
// Copy the restart data to the GPU
CopyToDevice(f_even,cDistEven,10*N*sizeof(double));
CopyToDevice(f_odd,cDistOdd,9*N*sizeof(double));
CopyToDevice(Den,cDen,2*N*sizeof(double));
DeviceBarrier();
ScaLBL_CopyToDevice(f_even,cDistEven,10*N*sizeof(double));
ScaLBL_CopyToDevice(f_odd,cDistOdd,9*N*sizeof(double));
ScaLBL_CopyToDevice(Den,cDen,2*N*sizeof(double));
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
}
//*************************************************************************
// Compute the phase indicator field and reset Copy, Den
//*************************************************************************
// ComputePhi(ID, Phi, Copy, Den, N);
ComputePhi(ID, Phi, Den, N);
// ScaLBL_ComputePhaseField(ID, Phi, Copy, Den, N);
ScaLBL_ComputePhaseField(ID, Phi, Den, N);
//*************************************************************************
//...................................................................................
PackValues(dvcSendList_x, sendCount_x,sendbuf_x, Phi, N);
PackValues(dvcSendList_y, sendCount_y,sendbuf_y, Phi, N);
PackValues(dvcSendList_z, sendCount_z,sendbuf_z, Phi, N);
PackValues(dvcSendList_X, sendCount_X,sendbuf_X, Phi, N);
PackValues(dvcSendList_Y, sendCount_Y,sendbuf_Y, Phi, N);
PackValues(dvcSendList_Z, sendCount_Z,sendbuf_Z, Phi, N);
PackValues(dvcSendList_xy, sendCount_xy,sendbuf_xy, Phi, N);
PackValues(dvcSendList_xY, sendCount_xY,sendbuf_xY, Phi, N);
PackValues(dvcSendList_Xy, sendCount_Xy,sendbuf_Xy, Phi, N);
PackValues(dvcSendList_XY, sendCount_XY,sendbuf_XY, Phi, N);
PackValues(dvcSendList_xz, sendCount_xz,sendbuf_xz, Phi, N);
PackValues(dvcSendList_xZ, sendCount_xZ,sendbuf_xZ, Phi, N);
PackValues(dvcSendList_Xz, sendCount_Xz,sendbuf_Xz, Phi, N);
PackValues(dvcSendList_XZ, sendCount_XZ,sendbuf_XZ, Phi, N);
PackValues(dvcSendList_yz, sendCount_yz,sendbuf_yz, Phi, N);
PackValues(dvcSendList_yZ, sendCount_yZ,sendbuf_yZ, Phi, N);
PackValues(dvcSendList_Yz, sendCount_Yz,sendbuf_Yz, Phi, N);
PackValues(dvcSendList_YZ, sendCount_YZ,sendbuf_YZ, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_x, sendCount_x,sendbuf_x, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_y, sendCount_y,sendbuf_y, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_z, sendCount_z,sendbuf_z, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_X, sendCount_X,sendbuf_X, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_Y, sendCount_Y,sendbuf_Y, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_Z, sendCount_Z,sendbuf_Z, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_xy, sendCount_xy,sendbuf_xy, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_xY, sendCount_xY,sendbuf_xY, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_Xy, sendCount_Xy,sendbuf_Xy, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_XY, sendCount_XY,sendbuf_XY, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_xz, sendCount_xz,sendbuf_xz, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_xZ, sendCount_xZ,sendbuf_xZ, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_Xz, sendCount_Xz,sendbuf_Xz, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_XZ, sendCount_XZ,sendbuf_XZ, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_yz, sendCount_yz,sendbuf_yz, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_yZ, sendCount_yZ,sendbuf_yZ, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_Yz, sendCount_Yz,sendbuf_Yz, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_YZ, sendCount_YZ,sendbuf_YZ, Phi, N);
DeviceBarrier();
ScaLBL_DeviceBarrier();
//...................................................................................
// Send / Recv all the phase indcator field values
//...................................................................................
@ -1298,34 +1298,34 @@ int main(int argc, char **argv)
//...................................................................................
MPI_Waitall(18,req1,stat1);
MPI_Waitall(18,req2,stat2);
DeviceBarrier();
ScaLBL_DeviceBarrier();
//...................................................................................
//...................................................................................
/* UnpackValues(faceGrid, packThreads, dvcSendList_x, sendCount_x,sendbuf_x, Phi, N);
UnpackValues(faceGrid, packThreads, dvcSendList_y, sendCount_y,sendbuf_y, Phi, N);
UnpackValues(faceGrid, packThreads, dvcSendList_z, sendCount_z,sendbuf_z, Phi, N);
UnpackValues(faceGrid, packThreads, dvcSendList_X, sendCount_X,sendbuf_X, Phi, N);
UnpackValues(faceGrid, packThreads, dvcSendList_Y, sendCount_Y,sendbuf_Y, Phi, N);
UnpackValues(faceGrid, packThreads, dvcSendList_Z, sendCount_Z,sendbuf_Z, Phi, N);
/* ScaLBL_Scalar_Unpack(faceGrid, packThreads, dvcSendList_x, sendCount_x,sendbuf_x, Phi, N);
ScaLBL_Scalar_Unpack(faceGrid, packThreads, dvcSendList_y, sendCount_y,sendbuf_y, Phi, N);
ScaLBL_Scalar_Unpack(faceGrid, packThreads, dvcSendList_z, sendCount_z,sendbuf_z, Phi, N);
ScaLBL_Scalar_Unpack(faceGrid, packThreads, dvcSendList_X, sendCount_X,sendbuf_X, Phi, N);
ScaLBL_Scalar_Unpack(faceGrid, packThreads, dvcSendList_Y, sendCount_Y,sendbuf_Y, Phi, N);
ScaLBL_Scalar_Unpack(faceGrid, packThreads, dvcSendList_Z, sendCount_Z,sendbuf_Z, Phi, N);
*/
UnpackValues(dvcRecvList_x, recvCount_x,recvbuf_x, Phi, N);
UnpackValues(dvcRecvList_y, recvCount_y,recvbuf_y, Phi, N);
UnpackValues(dvcRecvList_z, recvCount_z,recvbuf_z, Phi, N);
UnpackValues(dvcRecvList_X, recvCount_X,recvbuf_X, Phi, N);
UnpackValues(dvcRecvList_Y, recvCount_Y,recvbuf_Y, Phi, N);
UnpackValues(dvcRecvList_Z, recvCount_Z,recvbuf_Z, Phi, N);
UnpackValues(dvcRecvList_xy, recvCount_xy,recvbuf_xy, Phi, N);
UnpackValues(dvcRecvList_xY, recvCount_xY,recvbuf_xY, Phi, N);
UnpackValues(dvcRecvList_Xy, recvCount_Xy,recvbuf_Xy, Phi, N);
UnpackValues(dvcRecvList_XY, recvCount_XY,recvbuf_XY, Phi, N);
UnpackValues(dvcRecvList_xz, recvCount_xz,recvbuf_xz, Phi, N);
UnpackValues(dvcRecvList_xZ, recvCount_xZ,recvbuf_xZ, Phi, N);
UnpackValues(dvcRecvList_Xz, recvCount_Xz,recvbuf_Xz, Phi, N);
UnpackValues(dvcRecvList_XZ, recvCount_XZ,recvbuf_XZ, Phi, N);
UnpackValues(dvcRecvList_yz, recvCount_yz,recvbuf_yz, Phi, N);
UnpackValues(dvcRecvList_yZ, recvCount_yZ,recvbuf_yZ, Phi, N);
UnpackValues(dvcRecvList_Yz, recvCount_Yz,recvbuf_Yz, Phi, N);
UnpackValues(dvcRecvList_YZ, recvCount_YZ,recvbuf_YZ, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_x, recvCount_x,recvbuf_x, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_y, recvCount_y,recvbuf_y, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_z, recvCount_z,recvbuf_z, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_X, recvCount_X,recvbuf_X, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_Y, recvCount_Y,recvbuf_Y, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_Z, recvCount_Z,recvbuf_Z, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_xy, recvCount_xy,recvbuf_xy, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_xY, recvCount_xY,recvbuf_xY, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_Xy, recvCount_Xy,recvbuf_Xy, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_XY, recvCount_XY,recvbuf_XY, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_xz, recvCount_xz,recvbuf_xz, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_xZ, recvCount_xZ,recvbuf_xZ, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_Xz, recvCount_Xz,recvbuf_Xz, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_XZ, recvCount_XZ,recvbuf_XZ, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_yz, recvCount_yz,recvbuf_yz, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_yZ, recvCount_yZ,recvbuf_yZ, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_Yz, recvCount_Yz,recvbuf_Yz, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_YZ, recvCount_YZ,recvbuf_YZ, Phi, N);
//...................................................................................
if (rank==0 && pBC){
@ -1333,29 +1333,29 @@ int main(int argc, char **argv)
printf("Setting outlet pressure = %f \n", dout);
}
if (pBC && kproc == 0) {
PressureBC_inlet(f_even,f_odd,din,Nx,Ny,Nz);
ColorBC_inlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_z(f_even,f_odd,din,Nx,Ny,Nz);
ScaLBL_Color_BC_z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
}
if (pBC && kproc == nprocz-1){
PressureBC_outlet(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ColorBC_outlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_Z(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ScaLBL_Color_BC_Z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
}
//...........................................................................
// Copy the phase indicator field for the earlier timestep
DeviceBarrier();
CopyToHost(Phase_tplus.data(),Phi,N*sizeof(double));
ScaLBL_DeviceBarrier();
ScaLBL_CopyToHost(Phase_tplus.data(),Phi,N*sizeof(double));
//...........................................................................
//...........................................................................
// Copy the data for for the analysis timestep
//...........................................................................
// Copy the phase from the GPU -> CPU
//...........................................................................
DeviceBarrier();
ComputePressureD3Q19(ID,f_even,f_odd,Pressure,Nx,Ny,Nz);
CopyToHost(Phase.data(),Phi,N*sizeof(double));
CopyToHost(Press.data(),Pressure,N*sizeof(double));
ScaLBL_DeviceBarrier();
ScaLBL_D3Q19_Pressure(ID,f_even,f_odd,Pressure,Nx,Ny,Nz);
ScaLBL_CopyToHost(Phase.data(),Phi,N*sizeof(double));
ScaLBL_CopyToHost(Press.data(),Pressure,N*sizeof(double));
MPI_Barrier(comm);
//...........................................................................
@ -1370,71 +1370,71 @@ int main(int argc, char **argv)
//*************************************************************************
// Fused Color Gradient and Collision
//*************************************************************************
ColorCollideOpt( ID,f_even,f_odd,Phi,ColorGrad,
ScaLBL_D3Q19_ColorCollide( ID,f_even,f_odd,Phi,ColorGrad,
Velocity,Nx,Ny,Nz,rlxA,rlxB,alpha,beta,Fx,Fy,Fz);
//*************************************************************************
//...................................................................................
PackDist(1,dvcSendList_x,0,sendCount_x,sendbuf_x,f_even,N);
PackDist(4,dvcSendList_x,sendCount_x,sendCount_x,sendbuf_x,f_even,N);
PackDist(5,dvcSendList_x,2*sendCount_x,sendCount_x,sendbuf_x,f_even,N);
PackDist(6,dvcSendList_x,3*sendCount_x,sendCount_x,sendbuf_x,f_even,N);
PackDist(7,dvcSendList_x,4*sendCount_x,sendCount_x,sendbuf_x,f_even,N);
ScaLBL_D3Q19_Pack(1,dvcSendList_x,0,sendCount_x,sendbuf_x,f_even,N);
ScaLBL_D3Q19_Pack(4,dvcSendList_x,sendCount_x,sendCount_x,sendbuf_x,f_even,N);
ScaLBL_D3Q19_Pack(5,dvcSendList_x,2*sendCount_x,sendCount_x,sendbuf_x,f_even,N);
ScaLBL_D3Q19_Pack(6,dvcSendList_x,3*sendCount_x,sendCount_x,sendbuf_x,f_even,N);
ScaLBL_D3Q19_Pack(7,dvcSendList_x,4*sendCount_x,sendCount_x,sendbuf_x,f_even,N);
//...Packing for X face(1,7,9,11,13)................................
PackDist(0,dvcSendList_X,0,sendCount_X,sendbuf_X,f_odd,N);
PackDist(3,dvcSendList_X,sendCount_X,sendCount_X,sendbuf_X,f_odd,N);
PackDist(4,dvcSendList_X,2*sendCount_X,sendCount_X,sendbuf_X,f_odd,N);
PackDist(5,dvcSendList_X,3*sendCount_X,sendCount_X,sendbuf_X,f_odd,N);
PackDist(6,dvcSendList_X,4*sendCount_X,sendCount_X,sendbuf_X,f_odd,N);
ScaLBL_D3Q19_Pack(0,dvcSendList_X,0,sendCount_X,sendbuf_X,f_odd,N);
ScaLBL_D3Q19_Pack(3,dvcSendList_X,sendCount_X,sendCount_X,sendbuf_X,f_odd,N);
ScaLBL_D3Q19_Pack(4,dvcSendList_X,2*sendCount_X,sendCount_X,sendbuf_X,f_odd,N);
ScaLBL_D3Q19_Pack(5,dvcSendList_X,3*sendCount_X,sendCount_X,sendbuf_X,f_odd,N);
ScaLBL_D3Q19_Pack(6,dvcSendList_X,4*sendCount_X,sendCount_X,sendbuf_X,f_odd,N);
//...Packing for y face(4,8,9,16,18).................................
PackDist(2,dvcSendList_y,0,sendCount_y,sendbuf_y,f_even,N);
PackDist(4,dvcSendList_y,sendCount_y,sendCount_y,sendbuf_y,f_even,N);
PackDist(4,dvcSendList_y,2*sendCount_y,sendCount_y,sendbuf_y,f_odd,N);
PackDist(8,dvcSendList_y,3*sendCount_y,sendCount_y,sendbuf_y,f_even,N);
PackDist(9,dvcSendList_y,4*sendCount_y,sendCount_y,sendbuf_y,f_even,N);
ScaLBL_D3Q19_Pack(2,dvcSendList_y,0,sendCount_y,sendbuf_y,f_even,N);
ScaLBL_D3Q19_Pack(4,dvcSendList_y,sendCount_y,sendCount_y,sendbuf_y,f_even,N);
ScaLBL_D3Q19_Pack(4,dvcSendList_y,2*sendCount_y,sendCount_y,sendbuf_y,f_odd,N);
ScaLBL_D3Q19_Pack(8,dvcSendList_y,3*sendCount_y,sendCount_y,sendbuf_y,f_even,N);
ScaLBL_D3Q19_Pack(9,dvcSendList_y,4*sendCount_y,sendCount_y,sendbuf_y,f_even,N);
//...Packing for Y face(3,7,10,15,17).................................
PackDist(1,dvcSendList_Y,0,sendCount_Y,sendbuf_Y,f_odd,N);
PackDist(3,dvcSendList_Y,sendCount_Y,sendCount_Y,sendbuf_Y,f_odd,N);
PackDist(5,dvcSendList_Y,2*sendCount_Y,sendCount_Y,sendbuf_Y,f_even,N);
PackDist(7,dvcSendList_Y,3*sendCount_Y,sendCount_Y,sendbuf_Y,f_odd,N);
PackDist(8,dvcSendList_Y,4*sendCount_Y,sendCount_Y,sendbuf_Y,f_odd,N);
ScaLBL_D3Q19_Pack(1,dvcSendList_Y,0,sendCount_Y,sendbuf_Y,f_odd,N);
ScaLBL_D3Q19_Pack(3,dvcSendList_Y,sendCount_Y,sendCount_Y,sendbuf_Y,f_odd,N);
ScaLBL_D3Q19_Pack(5,dvcSendList_Y,2*sendCount_Y,sendCount_Y,sendbuf_Y,f_even,N);
ScaLBL_D3Q19_Pack(7,dvcSendList_Y,3*sendCount_Y,sendCount_Y,sendbuf_Y,f_odd,N);
ScaLBL_D3Q19_Pack(8,dvcSendList_Y,4*sendCount_Y,sendCount_Y,sendbuf_Y,f_odd,N);
//...Packing for z face(6,12,13,16,17)................................
PackDist(3,dvcSendList_z,0,sendCount_z,sendbuf_z,f_even,N);
PackDist(6,dvcSendList_z,sendCount_z,sendCount_z,sendbuf_z,f_even,N);
PackDist(6,dvcSendList_z,2*sendCount_z,sendCount_z,sendbuf_z,f_odd,N);
PackDist(8,dvcSendList_z,3*sendCount_z,sendCount_z,sendbuf_z,f_even,N);
PackDist(8,dvcSendList_z,4*sendCount_z,sendCount_z,sendbuf_z,f_odd,N);
ScaLBL_D3Q19_Pack(3,dvcSendList_z,0,sendCount_z,sendbuf_z,f_even,N);
ScaLBL_D3Q19_Pack(6,dvcSendList_z,sendCount_z,sendCount_z,sendbuf_z,f_even,N);
ScaLBL_D3Q19_Pack(6,dvcSendList_z,2*sendCount_z,sendCount_z,sendbuf_z,f_odd,N);
ScaLBL_D3Q19_Pack(8,dvcSendList_z,3*sendCount_z,sendCount_z,sendbuf_z,f_even,N);
ScaLBL_D3Q19_Pack(8,dvcSendList_z,4*sendCount_z,sendCount_z,sendbuf_z,f_odd,N);
//...Packing for Z face(5,11,14,15,18)................................
PackDist(2,dvcSendList_Z,0,sendCount_Z,sendbuf_Z,f_odd,N);
PackDist(5,dvcSendList_Z,sendCount_Z,sendCount_Z,sendbuf_Z,f_odd,N);
PackDist(7,dvcSendList_Z,2*sendCount_Z,sendCount_Z,sendbuf_Z,f_even,N);
PackDist(7,dvcSendList_Z,3*sendCount_Z,sendCount_Z,sendbuf_Z,f_odd,N);
PackDist(9,dvcSendList_Z,4*sendCount_Z,sendCount_Z,sendbuf_Z,f_even,N);
ScaLBL_D3Q19_Pack(2,dvcSendList_Z,0,sendCount_Z,sendbuf_Z,f_odd,N);
ScaLBL_D3Q19_Pack(5,dvcSendList_Z,sendCount_Z,sendCount_Z,sendbuf_Z,f_odd,N);
ScaLBL_D3Q19_Pack(7,dvcSendList_Z,2*sendCount_Z,sendCount_Z,sendbuf_Z,f_even,N);
ScaLBL_D3Q19_Pack(7,dvcSendList_Z,3*sendCount_Z,sendCount_Z,sendbuf_Z,f_odd,N);
ScaLBL_D3Q19_Pack(9,dvcSendList_Z,4*sendCount_Z,sendCount_Z,sendbuf_Z,f_even,N);
//...Pack the xy edge (8)................................
PackDist(4,dvcSendList_xy,0,sendCount_xy,sendbuf_xy,f_even,N);
ScaLBL_D3Q19_Pack(4,dvcSendList_xy,0,sendCount_xy,sendbuf_xy,f_even,N);
//...Pack the Xy edge (9)................................
PackDist(4,dvcSendList_Xy,0,sendCount_Xy,sendbuf_Xy,f_odd,N);
ScaLBL_D3Q19_Pack(4,dvcSendList_Xy,0,sendCount_Xy,sendbuf_Xy,f_odd,N);
//...Pack the xY edge (10)................................
PackDist(5,dvcSendList_xY,0,sendCount_xY,sendbuf_xY,f_even,N);
ScaLBL_D3Q19_Pack(5,dvcSendList_xY,0,sendCount_xY,sendbuf_xY,f_even,N);
//...Pack the XY edge (7)................................
PackDist(3,dvcSendList_XY,0,sendCount_XY,sendbuf_XY,f_odd,N);
ScaLBL_D3Q19_Pack(3,dvcSendList_XY,0,sendCount_XY,sendbuf_XY,f_odd,N);
//...Pack the xz edge (12)................................
PackDist(6,dvcSendList_xz,0,sendCount_xz,sendbuf_xz,f_even,N);
ScaLBL_D3Q19_Pack(6,dvcSendList_xz,0,sendCount_xz,sendbuf_xz,f_even,N);
//...Pack the xZ edge (14)................................
PackDist(7,dvcSendList_xZ,0,sendCount_xZ,sendbuf_xZ,f_even,N);
ScaLBL_D3Q19_Pack(7,dvcSendList_xZ,0,sendCount_xZ,sendbuf_xZ,f_even,N);
//...Pack the Xz edge (13)................................
PackDist(6,dvcSendList_Xz,0,sendCount_Xz,sendbuf_Xz,f_odd,N);
ScaLBL_D3Q19_Pack(6,dvcSendList_Xz,0,sendCount_Xz,sendbuf_Xz,f_odd,N);
//...Pack the XZ edge (11)................................
PackDist(5,dvcSendList_XZ,0,sendCount_XZ,sendbuf_XZ,f_odd,N);
ScaLBL_D3Q19_Pack(5,dvcSendList_XZ,0,sendCount_XZ,sendbuf_XZ,f_odd,N);
//...Pack the xz edge (12)................................
//...Pack the yz edge (16)................................
PackDist(8,dvcSendList_yz,0,sendCount_yz,sendbuf_yz,f_even,N);
ScaLBL_D3Q19_Pack(8,dvcSendList_yz,0,sendCount_yz,sendbuf_yz,f_even,N);
//...Pack the yZ edge (18)................................
PackDist(9,dvcSendList_yZ,0,sendCount_yZ,sendbuf_yZ,f_even,N);
ScaLBL_D3Q19_Pack(9,dvcSendList_yZ,0,sendCount_yZ,sendbuf_yZ,f_even,N);
//...Pack the Yz edge (17)................................
PackDist(8,dvcSendList_Yz,0,sendCount_Yz,sendbuf_Yz,f_odd,N);
ScaLBL_D3Q19_Pack(8,dvcSendList_Yz,0,sendCount_Yz,sendbuf_Yz,f_odd,N);
//...Pack the YZ edge (15)................................
PackDist(7,dvcSendList_YZ,0,sendCount_YZ,sendbuf_YZ,f_odd,N);
ScaLBL_D3Q19_Pack(7,dvcSendList_YZ,0,sendCount_YZ,sendbuf_YZ,f_odd,N);
//...................................................................................
//...................................................................................
@ -1482,14 +1482,14 @@ int main(int argc, char **argv)
//*************************************************************************
// DensityStreamD3Q7(ID, Den, Copy, Phi, ColorGrad, Velocity, beta, Nx, Ny, Nz, pBC);
//*************************************************************************
MassColorCollideD3Q7(ID, A_even, A_odd, B_even, B_odd, Den, Phi,
ScaLBL_D3Q7_ColorCollideMass(ID, A_even, A_odd, B_even, B_odd, Den, Phi,
ColorGrad, Velocity, beta, N, pBC);
//*************************************************************************
// Swap the distributions for momentum transport
//*************************************************************************
SwapD3Q19(ID, f_even, f_odd, Nx, Ny, Nz);
ScaLBL_D3Q19_Swap(ID, f_even, f_odd, Nx, Ny, Nz);
//*************************************************************************
//...................................................................................
@ -1501,90 +1501,90 @@ int main(int argc, char **argv)
// Unpack the distributions on the device
//...................................................................................
//...Map recieve list for the X face: q=2,8,10,12,13 .................................
UnpackDist(0,-1,0,0,dvcRecvList_X,0,recvCount_X,recvbuf_X,f_odd,Nx,Ny,Nz);
UnpackDist(3,-1,-1,0,dvcRecvList_X,recvCount_X,recvCount_X,recvbuf_X,f_odd,Nx,Ny,Nz);
UnpackDist(4,-1,1,0,dvcRecvList_X,2*recvCount_X,recvCount_X,recvbuf_X,f_odd,Nx,Ny,Nz);
UnpackDist(5,-1,0,-1,dvcRecvList_X,3*recvCount_X,recvCount_X,recvbuf_X,f_odd,Nx,Ny,Nz);
UnpackDist(6,-1,0,1,dvcRecvList_X,4*recvCount_X,recvCount_X,recvbuf_X,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(0,-1,0,0,dvcRecvList_X,0,recvCount_X,recvbuf_X,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(3,-1,-1,0,dvcRecvList_X,recvCount_X,recvCount_X,recvbuf_X,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(4,-1,1,0,dvcRecvList_X,2*recvCount_X,recvCount_X,recvbuf_X,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(5,-1,0,-1,dvcRecvList_X,3*recvCount_X,recvCount_X,recvbuf_X,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(6,-1,0,1,dvcRecvList_X,4*recvCount_X,recvCount_X,recvbuf_X,f_odd,Nx,Ny,Nz);
//...................................................................................
//...Map recieve list for the x face: q=1,7,9,11,13..................................
UnpackDist(1,1,0,0,dvcRecvList_x,0,recvCount_x,recvbuf_x,f_even,Nx,Ny,Nz);
UnpackDist(4,1,1,0,dvcRecvList_x,recvCount_x,recvCount_x,recvbuf_x,f_even,Nx,Ny,Nz);
UnpackDist(5,1,-1,0,dvcRecvList_x,2*recvCount_x,recvCount_x,recvbuf_x,f_even,Nx,Ny,Nz);
UnpackDist(6,1,0,1,dvcRecvList_x,3*recvCount_x,recvCount_x,recvbuf_x,f_even,Nx,Ny,Nz);
UnpackDist(7,1,0,-1,dvcRecvList_x,4*recvCount_x,recvCount_x,recvbuf_x,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(1,1,0,0,dvcRecvList_x,0,recvCount_x,recvbuf_x,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(4,1,1,0,dvcRecvList_x,recvCount_x,recvCount_x,recvbuf_x,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(5,1,-1,0,dvcRecvList_x,2*recvCount_x,recvCount_x,recvbuf_x,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(6,1,0,1,dvcRecvList_x,3*recvCount_x,recvCount_x,recvbuf_x,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(7,1,0,-1,dvcRecvList_x,4*recvCount_x,recvCount_x,recvbuf_x,f_even,Nx,Ny,Nz);
//...................................................................................
//...Map recieve list for the y face: q=4,8,9,16,18 ...................................
UnpackDist(1,0,-1,0,dvcRecvList_Y,0,recvCount_Y,recvbuf_Y,f_odd,Nx,Ny,Nz);
UnpackDist(3,-1,-1,0,dvcRecvList_Y,recvCount_Y,recvCount_Y,recvbuf_Y,f_odd,Nx,Ny,Nz);
UnpackDist(5,1,-1,0,dvcRecvList_Y,2*recvCount_Y,recvCount_Y,recvbuf_Y,f_even,Nx,Ny,Nz);
UnpackDist(7,0,-1,-1,dvcRecvList_Y,3*recvCount_Y,recvCount_Y,recvbuf_Y,f_odd,Nx,Ny,Nz);
UnpackDist(8,0,-1,1,dvcRecvList_Y,4*recvCount_Y,recvCount_Y,recvbuf_Y,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(1,0,-1,0,dvcRecvList_Y,0,recvCount_Y,recvbuf_Y,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(3,-1,-1,0,dvcRecvList_Y,recvCount_Y,recvCount_Y,recvbuf_Y,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(5,1,-1,0,dvcRecvList_Y,2*recvCount_Y,recvCount_Y,recvbuf_Y,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(7,0,-1,-1,dvcRecvList_Y,3*recvCount_Y,recvCount_Y,recvbuf_Y,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(8,0,-1,1,dvcRecvList_Y,4*recvCount_Y,recvCount_Y,recvbuf_Y,f_odd,Nx,Ny,Nz);
//...................................................................................
//...Map recieve list for the Y face: q=3,7,10,15,17 ..................................
UnpackDist(2,0,1,0,dvcRecvList_y,0,recvCount_y,recvbuf_y,f_even,Nx,Ny,Nz);
UnpackDist(4,1,1,0,dvcRecvList_y,recvCount_y,recvCount_y,recvbuf_y,f_even,Nx,Ny,Nz);
UnpackDist(4,-1,1,0,dvcRecvList_y,2*recvCount_y,recvCount_y,recvbuf_y,f_odd,Nx,Ny,Nz);
UnpackDist(8,0,1,1,dvcRecvList_y,3*recvCount_y,recvCount_y,recvbuf_y,f_even,Nx,Ny,Nz);
UnpackDist(9,0,1,-1,dvcRecvList_y,4*recvCount_y,recvCount_y,recvbuf_y,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(2,0,1,0,dvcRecvList_y,0,recvCount_y,recvbuf_y,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(4,1,1,0,dvcRecvList_y,recvCount_y,recvCount_y,recvbuf_y,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(4,-1,1,0,dvcRecvList_y,2*recvCount_y,recvCount_y,recvbuf_y,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(8,0,1,1,dvcRecvList_y,3*recvCount_y,recvCount_y,recvbuf_y,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(9,0,1,-1,dvcRecvList_y,4*recvCount_y,recvCount_y,recvbuf_y,f_even,Nx,Ny,Nz);
//...................................................................................
//...Map recieve list for the z face<<<6,12,13,16,17)..............................................
UnpackDist(2,0,0,-1,dvcRecvList_Z,0,recvCount_Z,recvbuf_Z,f_odd,Nx,Ny,Nz);
UnpackDist(5,-1,0,-1,dvcRecvList_Z,recvCount_Z,recvCount_Z,recvbuf_Z,f_odd,Nx,Ny,Nz);
UnpackDist(7,1,0,-1,dvcRecvList_Z,2*recvCount_Z,recvCount_Z,recvbuf_Z,f_even,Nx,Ny,Nz);
UnpackDist(7,0,-1,-1,dvcRecvList_Z,3*recvCount_Z,recvCount_Z,recvbuf_Z,f_odd,Nx,Ny,Nz);
UnpackDist(9,0,1,-1,dvcRecvList_Z,4*recvCount_Z,recvCount_Z,recvbuf_Z,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(2,0,0,-1,dvcRecvList_Z,0,recvCount_Z,recvbuf_Z,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(5,-1,0,-1,dvcRecvList_Z,recvCount_Z,recvCount_Z,recvbuf_Z,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(7,1,0,-1,dvcRecvList_Z,2*recvCount_Z,recvCount_Z,recvbuf_Z,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(7,0,-1,-1,dvcRecvList_Z,3*recvCount_Z,recvCount_Z,recvbuf_Z,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(9,0,1,-1,dvcRecvList_Z,4*recvCount_Z,recvCount_Z,recvbuf_Z,f_even,Nx,Ny,Nz);
//...Map recieve list for the Z face<<<5,11,14,15,18)..............................................
UnpackDist(3,0,0,1,dvcRecvList_z,0,recvCount_z,recvbuf_z,f_even,Nx,Ny,Nz);
UnpackDist(6,1,0,1,dvcRecvList_z,recvCount_z,recvCount_z,recvbuf_z,f_even,Nx,Ny,Nz);
UnpackDist(6,-1,0,1,dvcRecvList_z,2*recvCount_z,recvCount_z,recvbuf_z,f_odd,Nx,Ny,Nz);
UnpackDist(8,0,1,1,dvcRecvList_z,3*recvCount_z,recvCount_z,recvbuf_z,f_even,Nx,Ny,Nz);
UnpackDist(8,0,-1,1,dvcRecvList_z,4*recvCount_z,recvCount_z,recvbuf_z,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(3,0,0,1,dvcRecvList_z,0,recvCount_z,recvbuf_z,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(6,1,0,1,dvcRecvList_z,recvCount_z,recvCount_z,recvbuf_z,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(6,-1,0,1,dvcRecvList_z,2*recvCount_z,recvCount_z,recvbuf_z,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(8,0,1,1,dvcRecvList_z,3*recvCount_z,recvCount_z,recvbuf_z,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(8,0,-1,1,dvcRecvList_z,4*recvCount_z,recvCount_z,recvbuf_z,f_odd,Nx,Ny,Nz);
//..................................................................................
//...Map recieve list for the xy edge <<<8)................................
UnpackDist(3,-1,-1,0,dvcRecvList_XY,0,recvCount_XY,recvbuf_XY,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(3,-1,-1,0,dvcRecvList_XY,0,recvCount_XY,recvbuf_XY,f_odd,Nx,Ny,Nz);
//...Map recieve list for the Xy edge <<<9)................................
UnpackDist(5,1,-1,0,dvcRecvList_xY,0,recvCount_xY,recvbuf_xY,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(5,1,-1,0,dvcRecvList_xY,0,recvCount_xY,recvbuf_xY,f_even,Nx,Ny,Nz);
//...Map recieve list for the xY edge <<<10)................................
UnpackDist(4,-1,1,0,dvcRecvList_Xy,0,recvCount_Xy,recvbuf_Xy,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(4,-1,1,0,dvcRecvList_Xy,0,recvCount_Xy,recvbuf_Xy,f_odd,Nx,Ny,Nz);
//...Map recieve list for the XY edge <<<7)................................
UnpackDist(4,1,1,0,dvcRecvList_xy,0,recvCount_xy,recvbuf_xy,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(4,1,1,0,dvcRecvList_xy,0,recvCount_xy,recvbuf_xy,f_even,Nx,Ny,Nz);
//...Map recieve list for the xz edge <<<12)................................
UnpackDist(5,-1,0,-1,dvcRecvList_XZ,0,recvCount_XZ,recvbuf_XZ,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(5,-1,0,-1,dvcRecvList_XZ,0,recvCount_XZ,recvbuf_XZ,f_odd,Nx,Ny,Nz);
//...Map recieve list for the xZ edge <<<14)................................
UnpackDist(6,-1,0,1,dvcRecvList_Xz,0,recvCount_Xz,recvbuf_Xz,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(6,-1,0,1,dvcRecvList_Xz,0,recvCount_Xz,recvbuf_Xz,f_odd,Nx,Ny,Nz);
//...Map recieve list for the Xz edge <<<13)................................
UnpackDist(7,1,0,-1,dvcRecvList_xZ,0,recvCount_xZ,recvbuf_xZ,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(7,1,0,-1,dvcRecvList_xZ,0,recvCount_xZ,recvbuf_xZ,f_even,Nx,Ny,Nz);
//...Map recieve list for the XZ edge <<<11)................................
UnpackDist(6,1,0,1,dvcRecvList_xz,0,recvCount_xz,recvbuf_xz,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(6,1,0,1,dvcRecvList_xz,0,recvCount_xz,recvbuf_xz,f_even,Nx,Ny,Nz);
//...Map recieve list for the yz edge <<<16)................................
UnpackDist(7,0,-1,-1,dvcRecvList_YZ,0,recvCount_YZ,recvbuf_YZ,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(7,0,-1,-1,dvcRecvList_YZ,0,recvCount_YZ,recvbuf_YZ,f_odd,Nx,Ny,Nz);
//...Map recieve list for the yZ edge <<<18)................................
UnpackDist(8,0,-1,1,dvcRecvList_Yz,0,recvCount_Yz,recvbuf_Yz,f_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(8,0,-1,1,dvcRecvList_Yz,0,recvCount_Yz,recvbuf_Yz,f_odd,Nx,Ny,Nz);
//...Map recieve list for the Yz edge <<<17)................................
UnpackDist(9,0,1,-1,dvcRecvList_yZ,0,recvCount_yZ,recvbuf_yZ,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(9,0,1,-1,dvcRecvList_yZ,0,recvCount_yZ,recvbuf_yZ,f_even,Nx,Ny,Nz);
//...Map recieve list for the YZ edge <<<15)................................
UnpackDist(8,0,1,1,dvcRecvList_yz,0,recvCount_yz,recvbuf_yz,f_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(8,0,1,1,dvcRecvList_yz,0,recvCount_yz,recvbuf_yz,f_even,Nx,Ny,Nz);
//...................................................................................
//...................................................................................
PackDist(1,dvcSendList_x,0,sendCount_x,sendbuf_x,A_even,N);
PackDist(1,dvcSendList_x,sendCount_x,sendCount_x,sendbuf_x,B_even,N);
ScaLBL_D3Q19_Pack(1,dvcSendList_x,0,sendCount_x,sendbuf_x,A_even,N);
ScaLBL_D3Q19_Pack(1,dvcSendList_x,sendCount_x,sendCount_x,sendbuf_x,B_even,N);
//...Packing for X face(1,7,9,11,13)................................
PackDist(0,dvcSendList_X,0,sendCount_X,sendbuf_X,A_odd,N);
PackDist(0,dvcSendList_X,sendCount_X,sendCount_X,sendbuf_X,B_odd,N);
ScaLBL_D3Q19_Pack(0,dvcSendList_X,0,sendCount_X,sendbuf_X,A_odd,N);
ScaLBL_D3Q19_Pack(0,dvcSendList_X,sendCount_X,sendCount_X,sendbuf_X,B_odd,N);
//...Packing for y face(4,8,9,16,18).................................
PackDist(2,dvcSendList_y,0,sendCount_y,sendbuf_y,A_even,N);
PackDist(2,dvcSendList_y,sendCount_y,sendCount_y,sendbuf_y,B_even,N);
ScaLBL_D3Q19_Pack(2,dvcSendList_y,0,sendCount_y,sendbuf_y,A_even,N);
ScaLBL_D3Q19_Pack(2,dvcSendList_y,sendCount_y,sendCount_y,sendbuf_y,B_even,N);
//...Packing for Y face(3,7,10,15,17).................................
PackDist(1,dvcSendList_Y,0,sendCount_Y,sendbuf_Y,A_odd,N);
PackDist(1,dvcSendList_Y,sendCount_Y,sendCount_Y,sendbuf_Y,B_odd,N);
ScaLBL_D3Q19_Pack(1,dvcSendList_Y,0,sendCount_Y,sendbuf_Y,A_odd,N);
ScaLBL_D3Q19_Pack(1,dvcSendList_Y,sendCount_Y,sendCount_Y,sendbuf_Y,B_odd,N);
//...Packing for z face(6,12,13,16,17)................................
PackDist(3,dvcSendList_z,0,sendCount_z,sendbuf_z,A_even,N);
PackDist(3,dvcSendList_z,sendCount_z,sendCount_z,sendbuf_z,B_even,N);
ScaLBL_D3Q19_Pack(3,dvcSendList_z,0,sendCount_z,sendbuf_z,A_even,N);
ScaLBL_D3Q19_Pack(3,dvcSendList_z,sendCount_z,sendCount_z,sendbuf_z,B_even,N);
//...Packing for Z face(5,11,14,15,18)................................
PackDist(2,dvcSendList_Z,0,sendCount_Z,sendbuf_Z,A_odd,N);
PackDist(2,dvcSendList_Z,sendCount_Z,sendCount_Z,sendbuf_Z,B_odd,N);
ScaLBL_D3Q19_Pack(2,dvcSendList_Z,0,sendCount_Z,sendbuf_Z,A_odd,N);
ScaLBL_D3Q19_Pack(2,dvcSendList_Z,sendCount_Z,sendCount_Z,sendbuf_Z,B_odd,N);
//...................................................................................
//...................................................................................
@ -1603,8 +1603,8 @@ int main(int argc, char **argv)
MPI_Irecv(recvbuf_z, 2*recvCount_z,MPI_DOUBLE,rank_z,recvtag,comm,&req2[5]);
//...................................................................................
SwapD3Q7(ID, A_even, A_odd, Nx, Ny, Nz);
SwapD3Q7(ID, B_even, B_odd, Nx, Ny, Nz);
ScaLBL_D3Q7_Swap(ID, A_even, A_odd, Nx, Ny, Nz);
ScaLBL_D3Q7_Swap(ID, B_even, B_odd, Nx, Ny, Nz);
//...................................................................................
// Wait for completion of D3Q19 communication
@ -1614,59 +1614,59 @@ int main(int argc, char **argv)
// Unpack the distributions on the device
//...................................................................................
//...Map recieve list for the X face: q=2,8,10,12,13 .................................
UnpackDist(0,-1,0,0,dvcRecvList_X,0,recvCount_X,recvbuf_X,A_odd,Nx,Ny,Nz);
UnpackDist(0,-1,0,0,dvcRecvList_X,recvCount_X,recvCount_X,recvbuf_X,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(0,-1,0,0,dvcRecvList_X,0,recvCount_X,recvbuf_X,A_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(0,-1,0,0,dvcRecvList_X,recvCount_X,recvCount_X,recvbuf_X,B_odd,Nx,Ny,Nz);
//...................................................................................
//...Map recieve list for the x face: q=1,7,9,11,13..................................
UnpackDist(1,1,0,0,dvcRecvList_x,0,recvCount_x,recvbuf_x,A_even,Nx,Ny,Nz);
UnpackDist(1,1,0,0,dvcRecvList_x,recvCount_x,recvCount_x,recvbuf_x,B_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(1,1,0,0,dvcRecvList_x,0,recvCount_x,recvbuf_x,A_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(1,1,0,0,dvcRecvList_x,recvCount_x,recvCount_x,recvbuf_x,B_even,Nx,Ny,Nz);
//...................................................................................
//...Map recieve list for the y face: q=4,8,9,16,18 ...................................
UnpackDist(1,0,-1,0,dvcRecvList_Y,0,recvCount_Y,recvbuf_Y,A_odd,Nx,Ny,Nz);
UnpackDist(1,0,-1,0,dvcRecvList_Y,recvCount_Y,recvCount_Y,recvbuf_Y,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(1,0,-1,0,dvcRecvList_Y,0,recvCount_Y,recvbuf_Y,A_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(1,0,-1,0,dvcRecvList_Y,recvCount_Y,recvCount_Y,recvbuf_Y,B_odd,Nx,Ny,Nz);
//...................................................................................
//...Map recieve list for the Y face: q=3,7,10,15,17 ..................................
UnpackDist(2,0,1,0,dvcRecvList_y,0,recvCount_y,recvbuf_y,A_even,Nx,Ny,Nz);
UnpackDist(2,0,1,0,dvcRecvList_y,recvCount_y,recvCount_y,recvbuf_y,B_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(2,0,1,0,dvcRecvList_y,0,recvCount_y,recvbuf_y,A_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(2,0,1,0,dvcRecvList_y,recvCount_y,recvCount_y,recvbuf_y,B_even,Nx,Ny,Nz);
//...................................................................................
//...Map recieve list for the z face<<<6,12,13,16,17)..............................................
UnpackDist(2,0,0,-1,dvcRecvList_Z,0,recvCount_Z,recvbuf_Z,A_odd,Nx,Ny,Nz);
UnpackDist(2,0,0,-1,dvcRecvList_Z,recvCount_Z,recvCount_Z,recvbuf_Z,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(2,0,0,-1,dvcRecvList_Z,0,recvCount_Z,recvbuf_Z,A_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(2,0,0,-1,dvcRecvList_Z,recvCount_Z,recvCount_Z,recvbuf_Z,B_odd,Nx,Ny,Nz);
//...Map recieve list for the Z face<<<5,11,14,15,18)..............................................
UnpackDist(3,0,0,1,dvcRecvList_z,0,recvCount_z,recvbuf_z,A_even,Nx,Ny,Nz);
UnpackDist(3,0,0,1,dvcRecvList_z,recvCount_z,recvCount_z,recvbuf_z,B_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(3,0,0,1,dvcRecvList_z,0,recvCount_z,recvbuf_z,A_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(3,0,0,1,dvcRecvList_z,recvCount_z,recvCount_z,recvbuf_z,B_even,Nx,Ny,Nz);
//..................................................................................
//..................................................................................
ComputeDensityD3Q7(ID, A_even, A_odd, &Den[0], Nx, Ny, Nz);
ComputeDensityD3Q7(ID, B_even, B_odd, &Den[N], Nx, Ny, Nz);
ScaLBL_D3Q7_Density(ID, A_even, A_odd, &Den[0], Nx, Ny, Nz);
ScaLBL_D3Q7_Density(ID, B_even, B_odd, &Den[N], Nx, Ny, Nz);
//*************************************************************************
// Compute the phase indicator field
//*************************************************************************
// ComputePhi(ID, Phi, Copy, Den, N);
ComputePhi(ID, Phi, Den, N);
// ScaLBL_ComputePhaseField(ID, Phi, Copy, Den, N);
ScaLBL_ComputePhaseField(ID, Phi, Den, N);
//*************************************************************************
//...................................................................................
PackValues(dvcSendList_x, sendCount_x,sendbuf_x, Phi, N);
PackValues(dvcSendList_y, sendCount_y,sendbuf_y, Phi, N);
PackValues(dvcSendList_z, sendCount_z,sendbuf_z, Phi, N);
PackValues(dvcSendList_X, sendCount_X,sendbuf_X, Phi, N);
PackValues(dvcSendList_Y, sendCount_Y,sendbuf_Y, Phi, N);
PackValues(dvcSendList_Z, sendCount_Z,sendbuf_Z, Phi, N);
PackValues(dvcSendList_xy, sendCount_xy,sendbuf_xy, Phi, N);
PackValues(dvcSendList_xY, sendCount_xY,sendbuf_xY, Phi, N);
PackValues(dvcSendList_Xy, sendCount_Xy,sendbuf_Xy, Phi, N);
PackValues(dvcSendList_XY, sendCount_XY,sendbuf_XY, Phi, N);
PackValues(dvcSendList_xz, sendCount_xz,sendbuf_xz, Phi, N);
PackValues(dvcSendList_xZ, sendCount_xZ,sendbuf_xZ, Phi, N);
PackValues(dvcSendList_Xz, sendCount_Xz,sendbuf_Xz, Phi, N);
PackValues(dvcSendList_XZ, sendCount_XZ,sendbuf_XZ, Phi, N);
PackValues(dvcSendList_yz, sendCount_yz,sendbuf_yz, Phi, N);
PackValues(dvcSendList_yZ, sendCount_yZ,sendbuf_yZ, Phi, N);
PackValues(dvcSendList_Yz, sendCount_Yz,sendbuf_Yz, Phi, N);
PackValues(dvcSendList_YZ, sendCount_YZ,sendbuf_YZ, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_x, sendCount_x,sendbuf_x, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_y, sendCount_y,sendbuf_y, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_z, sendCount_z,sendbuf_z, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_X, sendCount_X,sendbuf_X, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_Y, sendCount_Y,sendbuf_Y, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_Z, sendCount_Z,sendbuf_Z, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_xy, sendCount_xy,sendbuf_xy, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_xY, sendCount_xY,sendbuf_xY, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_Xy, sendCount_Xy,sendbuf_Xy, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_XY, sendCount_XY,sendbuf_XY, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_xz, sendCount_xz,sendbuf_xz, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_xZ, sendCount_xZ,sendbuf_xZ, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_Xz, sendCount_Xz,sendbuf_Xz, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_XZ, sendCount_XZ,sendbuf_XZ, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_yz, sendCount_yz,sendbuf_yz, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_yZ, sendCount_yZ,sendbuf_yZ, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_Yz, sendCount_Yz,sendbuf_Yz, Phi, N);
ScaLBL_Scalar_Pack(dvcSendList_YZ, sendCount_YZ,sendbuf_YZ, Phi, N);
//...................................................................................
// Send / Recv all the phase indcator field values
//...................................................................................
@ -1712,44 +1712,44 @@ int main(int argc, char **argv)
//...................................................................................
MPI_Waitall(18,req1,stat1);
MPI_Waitall(18,req2,stat2);
DeviceBarrier();
ScaLBL_DeviceBarrier();
//...................................................................................
//...................................................................................
/* UnpackValues(faceGrid, packThreads, dvcSendList_x, sendCount_x,sendbuf_x, Phi, N);
UnpackValues(faceGrid, packThreads, dvcSendList_y, sendCount_y,sendbuf_y, Phi, N);
UnpackValues(faceGrid, packThreads, dvcSendList_z, sendCount_z,sendbuf_z, Phi, N);
UnpackValues(faceGrid, packThreads, dvcSendList_X, sendCount_X,sendbuf_X, Phi, N);
UnpackValues(faceGrid, packThreads, dvcSendList_Y, sendCount_Y,sendbuf_Y, Phi, N);
UnpackValues(faceGrid, packThreads, dvcSendList_Z, sendCount_Z,sendbuf_Z, Phi, N);
/* ScaLBL_Scalar_Unpack(faceGrid, packThreads, dvcSendList_x, sendCount_x,sendbuf_x, Phi, N);
ScaLBL_Scalar_Unpack(faceGrid, packThreads, dvcSendList_y, sendCount_y,sendbuf_y, Phi, N);
ScaLBL_Scalar_Unpack(faceGrid, packThreads, dvcSendList_z, sendCount_z,sendbuf_z, Phi, N);
ScaLBL_Scalar_Unpack(faceGrid, packThreads, dvcSendList_X, sendCount_X,sendbuf_X, Phi, N);
ScaLBL_Scalar_Unpack(faceGrid, packThreads, dvcSendList_Y, sendCount_Y,sendbuf_Y, Phi, N);
ScaLBL_Scalar_Unpack(faceGrid, packThreads, dvcSendList_Z, sendCount_Z,sendbuf_Z, Phi, N);
*/
UnpackValues(dvcRecvList_x, recvCount_x,recvbuf_x, Phi, N);
UnpackValues(dvcRecvList_y, recvCount_y,recvbuf_y, Phi, N);
UnpackValues(dvcRecvList_z, recvCount_z,recvbuf_z, Phi, N);
UnpackValues(dvcRecvList_X, recvCount_X,recvbuf_X, Phi, N);
UnpackValues(dvcRecvList_Y, recvCount_Y,recvbuf_Y, Phi, N);
UnpackValues(dvcRecvList_Z, recvCount_Z,recvbuf_Z, Phi, N);
UnpackValues(dvcRecvList_xy, recvCount_xy,recvbuf_xy, Phi, N);
UnpackValues(dvcRecvList_xY, recvCount_xY,recvbuf_xY, Phi, N);
UnpackValues(dvcRecvList_Xy, recvCount_Xy,recvbuf_Xy, Phi, N);
UnpackValues(dvcRecvList_XY, recvCount_XY,recvbuf_XY, Phi, N);
UnpackValues(dvcRecvList_xz, recvCount_xz,recvbuf_xz, Phi, N);
UnpackValues(dvcRecvList_xZ, recvCount_xZ,recvbuf_xZ, Phi, N);
UnpackValues(dvcRecvList_Xz, recvCount_Xz,recvbuf_Xz, Phi, N);
UnpackValues(dvcRecvList_XZ, recvCount_XZ,recvbuf_XZ, Phi, N);
UnpackValues(dvcRecvList_yz, recvCount_yz,recvbuf_yz, Phi, N);
UnpackValues(dvcRecvList_yZ, recvCount_yZ,recvbuf_yZ, Phi, N);
UnpackValues(dvcRecvList_Yz, recvCount_Yz,recvbuf_Yz, Phi, N);
UnpackValues(dvcRecvList_YZ, recvCount_YZ,recvbuf_YZ, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_x, recvCount_x,recvbuf_x, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_y, recvCount_y,recvbuf_y, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_z, recvCount_z,recvbuf_z, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_X, recvCount_X,recvbuf_X, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_Y, recvCount_Y,recvbuf_Y, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_Z, recvCount_Z,recvbuf_Z, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_xy, recvCount_xy,recvbuf_xy, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_xY, recvCount_xY,recvbuf_xY, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_Xy, recvCount_Xy,recvbuf_Xy, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_XY, recvCount_XY,recvbuf_XY, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_xz, recvCount_xz,recvbuf_xz, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_xZ, recvCount_xZ,recvbuf_xZ, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_Xz, recvCount_Xz,recvbuf_Xz, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_XZ, recvCount_XZ,recvbuf_XZ, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_yz, recvCount_yz,recvbuf_yz, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_yZ, recvCount_yZ,recvbuf_yZ, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_Yz, recvCount_Yz,recvbuf_Yz, Phi, N);
ScaLBL_Scalar_Unpack(dvcRecvList_YZ, recvCount_YZ,recvbuf_YZ, Phi, N);
//...................................................................................
if (pBC && kproc == 0) {
PressureBC_inlet(f_even,f_odd,din,Nx,Ny,Nz);
ColorBC_inlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_z(f_even,f_odd,din,Nx,Ny,Nz);
ScaLBL_Color_BC_z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
}
if (pBC && kproc == nprocz-1){
PressureBC_outlet(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ColorBC_outlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_Z(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ScaLBL_Color_BC_Z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
}
//...................................................................................
@ -1761,9 +1761,9 @@ int main(int argc, char **argv)
if (timestep%RESTART_INTERVAL == 0){
// Copy the data to the CPU
CopyToHost(cDistEven,f_even,10*N*sizeof(double));
CopyToHost(cDistOdd,f_odd,9*N*sizeof(double));
CopyToHost(cDen,Den,2*N*sizeof(double));
ScaLBL_CopyToHost(cDistEven,f_even,10*N*sizeof(double));
ScaLBL_CopyToHost(cDistOdd,f_odd,9*N*sizeof(double));
ScaLBL_CopyToHost(cDen,Den,2*N*sizeof(double));
// Read in the restart file to CPU buffers
WriteCheckpoint(LocalRestartFile, cDen, cDistEven, cDistOdd, N);
}
@ -1772,12 +1772,12 @@ int main(int argc, char **argv)
// End the bubble loop
//...........................................................................
// Copy the phase indicator field for the later timestep
DeviceBarrier();
ComputePressureD3Q19(ID,f_even,f_odd,Pressure,Nx,Ny,Nz);
CopyToHost(Phase_tminus.data(),Phi,N*sizeof(double));
CopyToHost(Phase_tplus.data(),Phi,N*sizeof(double));
CopyToHost(Phase.data(),Phi,N*sizeof(double));
CopyToHost(Press.data(),Pressure,N*sizeof(double));
ScaLBL_DeviceBarrier();
ScaLBL_D3Q19_Pressure(ID,f_even,f_odd,Pressure,Nx,Ny,Nz);
ScaLBL_CopyToHost(Phase_tminus.data(),Phi,N*sizeof(double));
ScaLBL_CopyToHost(Phase_tplus.data(),Phi,N*sizeof(double));
ScaLBL_CopyToHost(Phase.data(),Phi,N*sizeof(double));
ScaLBL_CopyToHost(Press.data(),Pressure,N*sizeof(double));
double temp=0.5/beta;
for (n=0; n<N; n++){
@ -2197,7 +2197,7 @@ int main(int argc, char **argv)
NULL_USE(sat_w);
//************************************************************************/
DeviceBarrier();
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
stoptime = MPI_Wtime();
if (rank==0) printf("-------------------------------------------------------------------\n");
@ -2218,7 +2218,7 @@ int main(int argc, char **argv)
//************************************************************************/
sprintf(LocalRankFilename,"%s%s","Phase.",LocalRankString);
// printf("Local File Name = %s \n",LocalRankFilename);
// CopyToHost(Phase.data(),Phi,N*sizeof(double));
// ScaLBL_CopyToHost(Phase.data(),Phi,N*sizeof(double));
FILE *PHASE;
PHASE = fopen(LocalRankFilename,"wb");
@ -2228,7 +2228,7 @@ int main(int argc, char **argv)
/* double *DensityValues;
DensityValues = new double [2*N];
CopyToHost(DensityValues,Copy,2*N*sizeof(double));
ScaLBL_CopyToHost(DensityValues,Copy,2*N*sizeof(double));
FILE *PHASE;
PHASE = fopen(LocalRankFilename,"wb");
fwrite(DensityValues,8,2*N,PHASE);

28
tests/TestCommD3Q19.cpp
View File

@ -11,7 +11,7 @@
using namespace std;
extern void GlobalFlipInitD3Q19(double *dist_even, double *dist_odd, int Nx, int Ny, int Nz,
extern void GlobalFlipScaLBL_D3Q19_Init(double *dist_even, double *dist_odd, int Nx, int Ny, int Nz,
int iproc, int jproc, int kproc, int nprocx, int nprocy, int nprocz)
{
// Set of Discrete velocities for the D3Q19 Model
@ -365,9 +365,9 @@ int main(int argc, char **argv)
//...........device phase ID.................................................
if (rank==0) printf ("Copying phase ID to device \n");
char *ID;
AllocateDeviceMemory((void **) &ID, N); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &ID, N); // Allocate device memory
// Copy to the device
CopyToDevice(ID, id, N);
ScaLBL_CopyToDevice(ID, id, N);
//...........................................................................
//...........................................................................
@ -378,8 +378,8 @@ int main(int argc, char **argv)
//......................device distributions.................................
double *f_even,*f_odd;
//...........................................................................
AllocateDeviceMemory((void **) &f_even, 10*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &f_odd, 9*dist_mem_size); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &f_even, 10*dist_mem_size); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &f_odd, 9*dist_mem_size); // Allocate device memory
//...........................................................................
double *f_even_host,*f_odd_host;
f_even_host = new double [10*N];
@ -416,10 +416,10 @@ int main(int argc, char **argv)
*/
if (rank==0) printf("Setting the distributions, size = : %i\n", N);
//...........................................................................
GlobalFlipInitD3Q19(f_even_host, f_odd_host, Nx-2, Ny-2, Nz-2,iproc,jproc,kproc,nprocx,nprocy,nprocz);
CopyToDevice(f_even, f_even_host, 10*dist_mem_size);
CopyToDevice(f_odd, f_odd_host, 9*dist_mem_size);
DeviceBarrier();
GlobalFlipScaLBL_D3Q19_Init(f_even_host, f_odd_host, Nx-2, Ny-2, Nz-2,iproc,jproc,kproc,nprocx,nprocy,nprocz);
ScaLBL_CopyToDevice(f_even, f_even_host, 10*dist_mem_size);
ScaLBL_CopyToDevice(f_odd, f_odd_host, 9*dist_mem_size);
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
//*************************************************************************
// Pack and send the D3Q19 distributions
@ -427,15 +427,15 @@ int main(int argc, char **argv)
//*************************************************************************
// Swap the distributions for momentum transport
//*************************************************************************
SwapD3Q19(ID, f_even, f_odd, Nx, Ny, Nz);
ScaLBL_D3Q19_Swap(ID, f_even, f_odd, Nx, Ny, Nz);
//*************************************************************************
// Wait for communications to complete and unpack the distributions
ScaLBL_Comm.RecvD3Q19(f_even, f_odd);
//*************************************************************************
//...........................................................................
CopyToHost(f_even_host,f_even,10*N*sizeof(double));
CopyToHost(f_odd_host,f_odd,9*N*sizeof(double));
ScaLBL_CopyToHost(f_even_host,f_even,10*N*sizeof(double));
ScaLBL_CopyToHost(f_odd_host,f_odd,9*N*sizeof(double));
check = GlobalCheckDebugDist(f_even_host, f_odd_host, Nx-2, Ny-2, Nz-2,iproc,jproc,kproc,nprocx,nprocy,nprocz);
//...........................................................................
@ -459,13 +459,13 @@ int main(int argc, char **argv)
//*************************************************************************
// Swap the distributions for momentum transport
//*************************************************************************
SwapD3Q19(ID, f_even, f_odd, Nx, Ny, Nz);
ScaLBL_D3Q19_Swap(ID, f_even, f_odd, Nx, Ny, Nz);
//*************************************************************************
// Wait for communications to complete and unpack the distributions
ScaLBL_Comm.RecvD3Q19(f_even, f_odd);
//*************************************************************************
DeviceBarrier();
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
// Iteration completed!
timestep++;

4
tests/TestDistanceFromSegmented.cpp Normal file
View File

@ -0,0 +1,4 @@
/* Generate a signed distance function from segmented 3D image data
*/

54
tests/TestMassConservationD3Q7.cpp
View File

@ -90,56 +90,56 @@ int main(int argc, char **argv)
//......................device distributions.................................
double *A_even,*A_odd,*B_even,*B_odd;
//...........................................................................
AllocateDeviceMemory((void **) &A_even, 4*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &A_odd, 3*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &B_even, 4*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &B_odd, 3*dist_mem_size); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &A_even, 4*dist_mem_size); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &A_odd, 3*dist_mem_size); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &B_even, 4*dist_mem_size); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &B_odd, 3*dist_mem_size); // Allocate device memory
//...........................................................................
double *Phi,*Den;
double *ColorGrad, *Velocity;
//...........................................................................
AllocateDeviceMemory((void **) &Phi, dist_mem_size);
AllocateDeviceMemory((void **) &Velocity, 3*dist_mem_size);
AllocateDeviceMemory((void **) &ColorGrad, 3*dist_mem_size);
AllocateDeviceMemory((void **) &Den, 2*dist_mem_size);
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &Phi, dist_mem_size);
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &Velocity, 3*dist_mem_size);
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &ColorGrad, 3*dist_mem_size);
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &Den, 2*dist_mem_size);
//...........device phase ID.................................................
if (rank==0) printf ("Copying phase ID to device \n");
char *ID;
AllocateDeviceMemory((void **) &ID, N); // Allocate device memory
ScaLBL_ScaLBL_AllocateDeviceMemory((void **) &ID, N); // Allocate device memory
// Copy to the device
CopyToDevice(ID, id, N);
CopyToDevice(Velocity, Vel, 3*N*sizeof(double));
ScaLBL_CopyToDevice(ID, id, N);
ScaLBL_CopyToDevice(Velocity, Vel, 3*N*sizeof(double));
//...........................................................................
InitDenColor(ID, Den, Phi, das, dbs, Nx, Ny, Nz);
CopyToHost(DenOriginal,Den,2*N*sizeof(double));
ScaLBL_Color_Init(ID, Den, Phi, das, dbs, Nx, Ny, Nz);
ScaLBL_CopyToHost(DenOriginal,Den,2*N*sizeof(double));
//......................................................................
InitD3Q7(ID, A_even, A_odd, &Den[0], Nx, Ny, Nz);
InitD3Q7(ID, B_even, B_odd, &Den[N], Nx, Ny, Nz);
ComputePhi(ID, Phi, Den, N);
ComputeColorGradient(ID,Phi,ColorGrad,Nx,Ny,Nz);
ScaLBL_D3Q7_Init(ID, A_even, A_odd, &Den[0], Nx, Ny, Nz);
ScaLBL_D3Q7_Init(ID, B_even, B_odd, &Den[N], Nx, Ny, Nz);
ScaLBL_ComputePhaseField(ID, Phi, Den, N);
ScaLBL_D3Q19_ColorGradient(ID,Phi,ColorGrad,Nx,Ny,Nz);
//..................................................................................
//*************************************************************************
// Carry out the density streaming step for mass transport
//*************************************************************************
MassColorCollideD3Q7(ID, A_even, A_odd, B_even, B_odd, Den, Phi,
ScaLBL_D3Q7_ColorCollideMass(ID, A_even, A_odd, B_even, B_odd, Den, Phi,
ColorGrad, Velocity, beta, N, pBC);
//*************************************************************************
//..................................................................................
ComputeDensityD3Q7(ID, A_even, A_odd, &Den[0], Nx, Ny, Nz);
ComputeDensityD3Q7(ID, B_even, B_odd, &Den[N], Nx, Ny, Nz);
ScaLBL_D3Q7_Density(ID, A_even, A_odd, &Den[0], Nx, Ny, Nz);
ScaLBL_D3Q7_Density(ID, B_even, B_odd, &Den[N], Nx, Ny, Nz);
//..................................................................................
ComputePhi(ID, Phi, Den, N);
ComputeColorGradient(ID,Phi,ColorGrad,Nx,Ny,Nz);
ScaLBL_ComputePhaseField(ID, Phi, Den, N);
ScaLBL_D3Q19_ColorGradient(ID,Phi,ColorGrad,Nx,Ny,Nz);
//..................................................................................
// Compare and make sure mass is conserved at every lattice site
bool CleanCheck = true;
double original,final;
CopyToHost(DenFinal,Den,2*N*sizeof(double));
ScaLBL_CopyToHost(DenFinal,Den,2*N*sizeof(double));
for (k=0;k<Nz;k++){
for (j=0;j<Ny;j++){
for (i=0;i<Nx;i++){
@ -172,10 +172,10 @@ int main(int argc, char **argv)
Aodd = new double[3*N];
Beven = new double[4*N];
Bodd = new double[3*N];
CopyToHost(Aeven,A_even,4*dist_mem_size);
CopyToHost(Aodd,A_odd,3*dist_mem_size);
CopyToHost(Beven,B_even,4*dist_mem_size);
CopyToHost(Bodd,B_odd,3*dist_mem_size);
ScaLBL_CopyToHost(Aeven,A_even,4*dist_mem_size);
ScaLBL_CopyToHost(Aodd,A_odd,3*dist_mem_size);
ScaLBL_CopyToHost(Beven,B_even,4*dist_mem_size);
ScaLBL_CopyToHost(Bodd,B_odd,3*dist_mem_size);
double rho,ux,uy,uz;
for (k=0;k<Nz;k++){
for (j=0;j<Ny;j++){

94
tests/lb2_CMT_wia.cpp
View File

@ -33,7 +33,7 @@ inline double NormProb(short int value, short int *mu, short int *sigma, int k)
return exp(-(value-m)*(value-m)/(2.0*s*s)) / sum;
}
extern "C" void CMT_MassColorCollideD3Q7(char *ID, double *A_even, double *A_odd, double *B_even, double *B_odd,
extern "C" void CMT_ScaLBL_D3Q7_ColorCollideMass(char *ID, double *A_even, double *A_odd, double *B_even, double *B_odd,
double *Den, double *Phi, double *ColorGrad, double beta, int N)
{
char id;
@ -171,16 +171,16 @@ int main(int argc, char **argv)
double *f_even,*f_odd;
double *A_even,*A_odd,*B_even,*B_odd;
//...........................................................................
AllocateDeviceMemory((void **) &f_even, 10*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &f_odd, 9*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &A_even, 4*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &A_odd, 3*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &B_even, 4*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &B_odd, 3*dist_mem_size); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &f_even, 10*dist_mem_size); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &f_odd, 9*dist_mem_size); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &A_even, 4*dist_mem_size); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &A_odd, 3*dist_mem_size); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &B_even, 4*dist_mem_size); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &B_odd, 3*dist_mem_size); // Allocate device memory
*/
printf("Set up ID \n");
char *ID;
AllocateDeviceMemory((void **) &ID, N);
ScaLBL_AllocateDeviceMemory((void **) &ID, N);
for (int k=0; k<Nz; k++){
for (int j=0; j<Ny; j++){
for (int i=0; i<Nx; i++){
@ -201,13 +201,13 @@ int main(int argc, char **argv)
printf("Allocate memory \n");
double *Phi,*Den, *ColorGrad;
AllocateDeviceMemory((void **) &Phi, dist_mem_size);
AllocateDeviceMemory((void **) &Den, NC*dist_mem_size);
AllocateDeviceMemory((void **) &ColorGrad, dist_mem_size);
ScaLBL_AllocateDeviceMemory((void **) &Phi, dist_mem_size);
ScaLBL_AllocateDeviceMemory((void **) &Den, NC*dist_mem_size);
ScaLBL_AllocateDeviceMemory((void **) &ColorGrad, dist_mem_size);
double *packed_even,*packed_odd;
AllocateDeviceMemory((void **) &packed_even, 4*dist_mem_size*NC); // Allocate device memory
AllocateDeviceMemory((void **) &packed_odd, 3*dist_mem_size*NC); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &packed_even, 4*dist_mem_size*NC); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &packed_odd, 3*dist_mem_size*NC); // Allocate device memory
//..............................................
// Read the input file
@ -255,9 +255,9 @@ int main(int argc, char **argv)
}
//..............................................
printf("Initialize distribution \n");
InitD3Q7(ID, &packed_even[0], &packed_odd[0], &Den[0], Nx, Ny, Nz);
InitD3Q7(ID, &packed_even[4*N], &packed_odd[3*N], &Den[N], Nx, Ny, Nz);
ComputePhi(ID, Phi, Den, N);
ScaLBL_D3Q7_Init(ID, &packed_even[0], &packed_odd[0], &Den[0], Nx, Ny, Nz);
ScaLBL_D3Q7_Init(ID, &packed_even[4*N], &packed_odd[3*N], &Den[N], Nx, Ny, Nz);
ScaLBL_ComputePhaseField(ID, Phi, Den, N);
int timestep=0;
int timestepMax=10;
@ -265,29 +265,29 @@ int main(int argc, char **argv)
while (timestep < timestepMax){
ComputeColorGradient(ID,Phi,ColorGrad,Nx,Ny,Nz);
ScaLBL_D3Q19_ColorGradient(ID,Phi,ColorGrad,Nx,Ny,Nz);
CMT_MassColorCollideD3Q7(ID, &packed_even[0], &packed_odd[0], &packed_even[4*N], &packed_odd[3*N], Den, Phi,
CMT_ScaLBL_D3Q7_ColorCollideMass(ID, &packed_even[0], &packed_odd[0], &packed_even[4*N], &packed_odd[3*N], Den, Phi,
ColorGrad, beta, N);
/* //...................................................................................
PackDist(1,dvcSendList_x,0,sendCount_x,sendbuf_x,A_even,N);
PackDist(1,dvcSendList_x,sendCount_x,sendCount_x,sendbuf_x,B_even,N);
ScaLBL_D3Q19_Pack(1,dvcSendList_x,0,sendCount_x,sendbuf_x,A_even,N);
ScaLBL_D3Q19_Pack(1,dvcSendList_x,sendCount_x,sendCount_x,sendbuf_x,B_even,N);
//...Packing for X face(1,7,9,11,13)................................
PackDist(0,dvcSendList_X,0,sendCount_X,sendbuf_X,A_odd,N);
PackDist(0,dvcSendList_X,sendCount_X,sendCount_X,sendbuf_X,B_odd,N);
ScaLBL_D3Q19_Pack(0,dvcSendList_X,0,sendCount_X,sendbuf_X,A_odd,N);
ScaLBL_D3Q19_Pack(0,dvcSendList_X,sendCount_X,sendCount_X,sendbuf_X,B_odd,N);
//...Packing for y face(4,8,9,16,18).................................
PackDist(2,dvcSendList_y,0,sendCount_y,sendbuf_y,A_even,N);
PackDist(2,dvcSendList_y,sendCount_y,sendCount_y,sendbuf_y,B_even,N);
ScaLBL_D3Q19_Pack(2,dvcSendList_y,0,sendCount_y,sendbuf_y,A_even,N);
ScaLBL_D3Q19_Pack(2,dvcSendList_y,sendCount_y,sendCount_y,sendbuf_y,B_even,N);
//...Packing for Y face(3,7,10,15,17).................................
PackDist(1,dvcSendList_Y,0,sendCount_Y,sendbuf_Y,A_odd,N);
PackDist(1,dvcSendList_Y,sendCount_Y,sendCount_Y,sendbuf_Y,B_odd,N);
ScaLBL_D3Q19_Pack(1,dvcSendList_Y,0,sendCount_Y,sendbuf_Y,A_odd,N);
ScaLBL_D3Q19_Pack(1,dvcSendList_Y,sendCount_Y,sendCount_Y,sendbuf_Y,B_odd,N);
//...Packing for z face(6,12,13,16,17)................................
PackDist(3,dvcSendList_z,0,sendCount_z,sendbuf_z,A_even,N);
PackDist(3,dvcSendList_z,sendCount_z,sendCount_z,sendbuf_z,B_even,N);
ScaLBL_D3Q19_Pack(3,dvcSendList_z,0,sendCount_z,sendbuf_z,A_even,N);
ScaLBL_D3Q19_Pack(3,dvcSendList_z,sendCount_z,sendCount_z,sendbuf_z,B_even,N);
//...Packing for Z face(5,11,14,15,18)................................
PackDist(2,dvcSendList_Z,0,sendCount_Z,sendbuf_Z,A_odd,N);
PackDist(2,dvcSendList_Z,sendCount_Z,sendCount_Z,sendbuf_Z,B_odd,N);
ScaLBL_D3Q19_Pack(2,dvcSendList_Z,0,sendCount_Z,sendbuf_Z,A_odd,N);
ScaLBL_D3Q19_Pack(2,dvcSendList_Z,sendCount_Z,sendCount_Z,sendbuf_Z,B_odd,N);
//...................................................................................
//...................................................................................
@ -306,8 +306,8 @@ int main(int argc, char **argv)
MPI_Irecv(recvbuf_z, 2*recvCount_z,MPI_DOUBLE,rank_z,recvtag,comm,&req2[5]);
*/ //...................................................................................
SwapD3Q7(ID, &packed_even[0], &packed_odd[0], Nx, Ny, Nz);
SwapD3Q7(ID, &packed_even[4*N], &packed_odd[3*N], Nx, Ny, Nz);
ScaLBL_D3Q7_Swap(ID, &packed_even[0], &packed_odd[0], Nx, Ny, Nz);
ScaLBL_D3Q7_Swap(ID, &packed_even[4*N], &packed_odd[3*N], Nx, Ny, Nz);
/* //...................................................................................
// Wait for completion of D3Q19 communication
@ -317,37 +317,37 @@ int main(int argc, char **argv)
// Unpack the distributions on the device
//...................................................................................
//...Map recieve list for the X face: q=2,8,10,12,13 .................................
UnpackDist(0,-1,0,0,dvcRecvList_X,0,recvCount_X,recvbuf_X,A_odd,Nx,Ny,Nz);
UnpackDist(0,-1,0,0,dvcRecvList_X,recvCount_X,recvCount_X,recvbuf_X,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(0,-1,0,0,dvcRecvList_X,0,recvCount_X,recvbuf_X,A_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(0,-1,0,0,dvcRecvList_X,recvCount_X,recvCount_X,recvbuf_X,B_odd,Nx,Ny,Nz);
//...................................................................................
//...Map recieve list for the x face: q=1,7,9,11,13..................................
UnpackDist(1,1,0,0,dvcRecvList_x,0,recvCount_x,recvbuf_x,A_even,Nx,Ny,Nz);
UnpackDist(1,1,0,0,dvcRecvList_x,recvCount_x,recvCount_x,recvbuf_x,B_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(1,1,0,0,dvcRecvList_x,0,recvCount_x,recvbuf_x,A_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(1,1,0,0,dvcRecvList_x,recvCount_x,recvCount_x,recvbuf_x,B_even,Nx,Ny,Nz);
//...................................................................................
//...Map recieve list for the y face: q=4,8,9,16,18 ...................................
UnpackDist(1,0,-1,0,dvcRecvList_Y,0,recvCount_Y,recvbuf_Y,A_odd,Nx,Ny,Nz);
UnpackDist(1,0,-1,0,dvcRecvList_Y,recvCount_Y,recvCount_Y,recvbuf_Y,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(1,0,-1,0,dvcRecvList_Y,0,recvCount_Y,recvbuf_Y,A_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(1,0,-1,0,dvcRecvList_Y,recvCount_Y,recvCount_Y,recvbuf_Y,B_odd,Nx,Ny,Nz);
//...................................................................................
//...Map recieve list for the Y face: q=3,7,10,15,17 ..................................
UnpackDist(2,0,1,0,dvcRecvList_y,0,recvCount_y,recvbuf_y,A_even,Nx,Ny,Nz);
UnpackDist(2,0,1,0,dvcRecvList_y,recvCount_y,recvCount_y,recvbuf_y,B_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(2,0,1,0,dvcRecvList_y,0,recvCount_y,recvbuf_y,A_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(2,0,1,0,dvcRecvList_y,recvCount_y,recvCount_y,recvbuf_y,B_even,Nx,Ny,Nz);
//...................................................................................
//...Map recieve list for the z face<<<6,12,13,16,17)..............................................
UnpackDist(2,0,0,-1,dvcRecvList_Z,0,recvCount_Z,recvbuf_Z,A_odd,Nx,Ny,Nz);
UnpackDist(2,0,0,-1,dvcRecvList_Z,recvCount_Z,recvCount_Z,recvbuf_Z,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(2,0,0,-1,dvcRecvList_Z,0,recvCount_Z,recvbuf_Z,A_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(2,0,0,-1,dvcRecvList_Z,recvCount_Z,recvCount_Z,recvbuf_Z,B_odd,Nx,Ny,Nz);
//...Map recieve list for the Z face<<<5,11,14,15,18)..............................................
UnpackDist(3,0,0,1,dvcRecvList_z,0,recvCount_z,recvbuf_z,A_even,Nx,Ny,Nz);
UnpackDist(3,0,0,1,dvcRecvList_z,recvCount_z,recvCount_z,recvbuf_z,B_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(3,0,0,1,dvcRecvList_z,0,recvCount_z,recvbuf_z,A_even,Nx,Ny,Nz);
ScaLBL_D3Q19_Unpack(3,0,0,1,dvcRecvList_z,recvCount_z,recvCount_z,recvbuf_z,B_even,Nx,Ny,Nz);
//..................................................................................
*/
//..................................................................................
ComputeDensityD3Q7(ID, &packed_even[0], &packed_odd[0], &Den[0], Nx, Ny, Nz);
ComputeDensityD3Q7(ID, &packed_even[4*N], &packed_odd[3*N], &Den[N], Nx, Ny, Nz);
ScaLBL_D3Q7_Density(ID, &packed_even[0], &packed_odd[0], &Den[0], Nx, Ny, Nz);
ScaLBL_D3Q7_Density(ID, &packed_even[4*N], &packed_odd[3*N], &Den[N], Nx, Ny, Nz);
//*************************************************************************
// Compute the phase indicator field
//*************************************************************************
ComputePhi(ID, Phi, Den, N);
ScaLBL_ComputePhaseField(ID, Phi, Den, N);
//*************************************************************************
timestep++;
}

748
tests/lb2_Color_blob_wia_mpi.cpp
View File

File diff suppressed because it is too large Load Diff

198
tests/lbpm_color_macro_simulator.cpp
View File

@ -494,7 +494,7 @@ int main(int argc, char **argv)
//...........device phase ID.................................................
if (rank==0) printf ("Copy phase ID to device \n");
char *ID;
AllocateDeviceMemory((void **) &ID, N); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &ID, N); // Allocate device memory
// Don't compute in the halo
for (k=0;k<Nz;k++){
for (j=0;j<Ny;j++){
@ -505,8 +505,8 @@ int main(int argc, char **argv)
}
}
// Copy to the device
CopyToDevice(ID, id, N);
DeviceBarrier();
ScaLBL_CopyToDevice(ID, id, N);
ScaLBL_DeviceBarrier();
//...........................................................................
//...........................................................................
@ -518,26 +518,26 @@ int main(int argc, char **argv)
double *f_even,*f_odd;
double *A_even,*A_odd,*B_even,*B_odd;
//...........................................................................
AllocateDeviceMemory((void **) &f_even, 10*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &f_odd, 9*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &A_even, 4*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &A_odd, 3*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &B_even, 4*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &B_odd, 3*dist_mem_size); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &f_even, 10*dist_mem_size); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &f_odd, 9*dist_mem_size); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &A_even, 4*dist_mem_size); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &A_odd, 3*dist_mem_size); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &B_even, 4*dist_mem_size); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &B_odd, 3*dist_mem_size); // Allocate device memory
//...........................................................................
double *Phi,*Den;
double *ColorGrad, *Velocity, *Pressure, *dvcSignDist;
//...........................................................................
AllocateDeviceMemory((void **) &Phi, dist_mem_size);
AllocateDeviceMemory((void **) &Pressure, dist_mem_size);
AllocateDeviceMemory((void **) &dvcSignDist, dist_mem_size);
AllocateDeviceMemory((void **) &Den, 2*dist_mem_size);
AllocateDeviceMemory((void **) &Velocity, 3*dist_mem_size);
AllocateDeviceMemory((void **) &ColorGrad, 3*dist_mem_size);
ScaLBL_AllocateDeviceMemory((void **) &Phi, dist_mem_size);
ScaLBL_AllocateDeviceMemory((void **) &Pressure, dist_mem_size);
ScaLBL_AllocateDeviceMemory((void **) &dvcSignDist, dist_mem_size);
ScaLBL_AllocateDeviceMemory((void **) &Den, 2*dist_mem_size);
ScaLBL_AllocateDeviceMemory((void **) &Velocity, 3*dist_mem_size);
ScaLBL_AllocateDeviceMemory((void **) &ColorGrad, 3*dist_mem_size);
//...........................................................................
// Copy signed distance for device initialization
CopyToDevice(dvcSignDist, Averages.SDs.data(), dist_mem_size);
ScaLBL_CopyToDevice(dvcSignDist, Averages.SDs.data(), dist_mem_size);
//...........................................................................
int logcount = 0; // number of surface write-outs
@ -549,10 +549,10 @@ int main(int argc, char **argv)
//...........................................................................
if (rank==0) printf("Setting the distributions, size = %i\n", N);
//...........................................................................
DeviceBarrier();
InitD3Q19(ID, f_even, f_odd, Nx, Ny, Nz);
InitDenColor(ID, Den, Phi, das, dbs, Nx, Ny, Nz);
DeviceBarrier();
ScaLBL_DeviceBarrier();
ScaLBL_D3Q19_Init(ID, f_even, f_odd, Nx, Ny, Nz);
ScaLBL_Color_Init(ID, Den, Phi, das, dbs, Nx, Ny, Nz);
ScaLBL_DeviceBarrier();
//......................................................................
if (Restart == true){
@ -578,10 +578,10 @@ int main(int argc, char **argv)
double *cDistOdd = new double[9*N];
ReadCheckpoint(LocalRestartFile, cDen, cDistEven, cDistOdd, N);
// Copy the restart data to the GPU
CopyToDevice(f_even,cDistEven,10*N*sizeof(double));
CopyToDevice(f_odd,cDistOdd,9*N*sizeof(double));
CopyToDevice(Den,cDen,2*N*sizeof(double));
DeviceBarrier();
ScaLBL_CopyToDevice(f_even,cDistEven,10*N*sizeof(double));
ScaLBL_CopyToDevice(f_odd,cDistOdd,9*N*sizeof(double));
ScaLBL_CopyToDevice(Den,cDen,2*N*sizeof(double));
ScaLBL_DeviceBarrier();
delete [] cDen;
delete [] cDistEven;
delete [] cDistOdd;
@ -589,9 +589,9 @@ int main(int argc, char **argv)
}
//......................................................................
InitD3Q7(ID, A_even, A_odd, &Den[0], Nx, Ny, Nz);
InitD3Q7(ID, B_even, B_odd, &Den[N], Nx, Ny, Nz);
DeviceBarrier();
ScaLBL_D3Q7_Init(ID, A_even, A_odd, &Den[0], Nx, Ny, Nz);
ScaLBL_D3Q7_Init(ID, B_even, B_odd, &Den[N], Nx, Ny, Nz);
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
//.......................................................................
// Once phase has been initialized, map solid to account for 'smeared' interface
@ -607,12 +607,12 @@ int main(int argc, char **argv)
//*************************************************************************
// Compute the phase indicator field and reset Copy, Den
//*************************************************************************
ComputePhi(ID, Phi, Den, N);
ScaLBL_ComputePhaseField(ID, Phi, Den, N);
//*************************************************************************
DeviceBarrier();
ScaLBL_DeviceBarrier();
ScaLBL_Comm.SendHalo(Phi);
ScaLBL_Comm.RecvHalo(Phi);
DeviceBarrier();
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
//*************************************************************************
@ -621,13 +621,13 @@ int main(int argc, char **argv)
printf("Setting outlet pressure = %f \n", dout);
}
if (BoundaryCondition==1 && Mask.kproc == 0) {
PressureBC_inlet(f_even,f_odd,din,Nx,Ny,Nz);
ColorBC_inlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_z(f_even,f_odd,din,Nx,Ny,Nz);
ScaLBL_Color_BC_z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
}
if (BoundaryCondition==1 && Mask.kproc == nprocz-1){
PressureBC_outlet(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ColorBC_outlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_Z(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ScaLBL_Color_BC_Z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
}
if (rank==0 && BoundaryCondition==2){
@ -636,13 +636,13 @@ int main(int argc, char **argv)
}
if (BoundaryCondition==2 && Mask.kproc == 0) {
ScaLBL_D3Q19_Velocity_BC_z(f_even,f_odd,din,Nx,Ny,Nz);
//ColorBC_inlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
//ScaLBL_Color_BC_z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
SetPhiSlice_z(Phi,1.0,Nx,Ny,Nz,0);
}
if (BoundaryCondition==2 && Mask.kproc == nprocz-1){
ScaLBL_D3Q19_Velocity_BC_Z(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
//ColorBC_outlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
//ScaLBL_Color_BC_Z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
SetPhiSlice_z(Phi,-1.0,Nx,Ny,Nz,Nz-1);
}
@ -658,35 +658,35 @@ int main(int argc, char **argv)
dout = 1.0-0.5*dp;
// set the initial boundary conditions
if (Mask.kproc == 0) {
PressureBC_inlet(f_even,f_odd,din,Nx,Ny,Nz);
ColorBC_inlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_z(f_even,f_odd,din,Nx,Ny,Nz);
ScaLBL_Color_BC_z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
}
if (Mask.kproc == nprocz-1){
PressureBC_outlet(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ColorBC_outlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_Z(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ScaLBL_Color_BC_Z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
}
}
ComputePressureD3Q19(ID,f_even,f_odd,Pressure,Nx,Ny,Nz);
ComputeVelocityD3Q19(ID,f_even,f_odd,Velocity,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure(ID,f_even,f_odd,Pressure,Nx,Ny,Nz);
ScaLBL_D3Q19_Velocity(ID,f_even,f_odd,Velocity,Nx,Ny,Nz);
//...........................................................................
// Copy the phase indicator field for the earlier timestep
DeviceBarrier();
CopyToHost(Averages.Phase_tplus.data(),Phi,N*sizeof(double));
ScaLBL_DeviceBarrier();
ScaLBL_CopyToHost(Averages.Phase_tplus.data(),Phi,N*sizeof(double));
//...........................................................................
//...........................................................................
// Copy the data for for the analysis timestep
//...........................................................................
// Copy the phase from the GPU -> CPU
//...........................................................................
DeviceBarrier();
ComputePressureD3Q19(ID,f_even,f_odd,Pressure,Nx,Ny,Nz);
CopyToHost(Averages.Phase.data(),Phi,N*sizeof(double));
CopyToHost(Averages.Press.data(),Pressure,N*sizeof(double));
CopyToHost(Averages.Vel_x.data(),&Velocity[0],N*sizeof(double));
CopyToHost(Averages.Vel_y.data(),&Velocity[N],N*sizeof(double));
CopyToHost(Averages.Vel_z.data(),&Velocity[2*N],N*sizeof(double));
ScaLBL_DeviceBarrier();
ScaLBL_D3Q19_Pressure(ID,f_even,f_odd,Pressure,Nx,Ny,Nz);
ScaLBL_CopyToHost(Averages.Phase.data(),Phi,N*sizeof(double));
ScaLBL_CopyToHost(Averages.Press.data(),Pressure,N*sizeof(double));
ScaLBL_CopyToHost(Averages.Vel_x.data(),&Velocity[0],N*sizeof(double));
ScaLBL_CopyToHost(Averages.Vel_y.data(),&Velocity[N],N*sizeof(double));
ScaLBL_CopyToHost(Averages.Vel_z.data(),&Velocity[2*N],N*sizeof(double));
//...........................................................................
if (rank==0) printf("********************************************************\n");
@ -694,7 +694,7 @@ int main(int argc, char **argv)
//.......create and start timer............
double starttime,stoptime,cputime;
DeviceBarrier();
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
starttime = MPI_Wtime();
//.........................................
@ -743,11 +743,11 @@ int main(int argc, char **argv)
//*************************************************************************
// Fused Color Gradient and Collision
//*************************************************************************
ColorCollideOpt( ID,f_even,f_odd,Phi,ColorGrad,
ScaLBL_D3Q19_ColorCollide( ID,f_even,f_odd,Phi,ColorGrad,
Velocity,Nx,Ny,Nz,rlxA,rlxB,alpha,beta,Fx,Fy,Fz);
//*************************************************************************
DeviceBarrier();
ScaLBL_DeviceBarrier();
//*************************************************************************
// Pack and send the D3Q19 distributions
ScaLBL_Comm.SendD3Q19(f_even, f_odd);
@ -756,36 +756,36 @@ int main(int argc, char **argv)
//*************************************************************************
// Carry out the density streaming step for mass transport
//*************************************************************************
MassColorCollideD3Q7(ID, A_even, A_odd, B_even, B_odd, Den, Phi,
ScaLBL_D3Q7_ColorCollideMass(ID, A_even, A_odd, B_even, B_odd, Den, Phi,
ColorGrad, Velocity, beta, N, pBC);
//*************************************************************************
DeviceBarrier();
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
//*************************************************************************
// Swap the distributions for momentum transport
//*************************************************************************
SwapD3Q19(ID, f_even, f_odd, Nx, Ny, Nz);
ScaLBL_D3Q19_Swap(ID, f_even, f_odd, Nx, Ny, Nz);
//*************************************************************************
DeviceBarrier();
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
//*************************************************************************
// Wait for communications to complete and unpack the distributions
ScaLBL_Comm.RecvD3Q19(f_even, f_odd);
//*************************************************************************
DeviceBarrier();
ScaLBL_DeviceBarrier();
//*************************************************************************
// Pack and send the D3Q7 distributions
ScaLBL_Comm.BiSendD3Q7(A_even, A_odd, B_even, B_odd);
//*************************************************************************
DeviceBarrier();
SwapD3Q7(ID, A_even, A_odd, Nx, Ny, Nz);
SwapD3Q7(ID, B_even, B_odd, Nx, Ny, Nz);
ScaLBL_DeviceBarrier();
ScaLBL_D3Q7_Swap(ID, A_even, A_odd, Nx, Ny, Nz);
ScaLBL_D3Q7_Swap(ID, B_even, B_odd, Nx, Ny, Nz);
DeviceBarrier();
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
//*************************************************************************
@ -793,44 +793,44 @@ int main(int argc, char **argv)
ScaLBL_Comm.BiRecvD3Q7(A_even, A_odd, B_even, B_odd);
//*************************************************************************
DeviceBarrier();
ScaLBL_DeviceBarrier();
//..................................................................................
ComputeDensityD3Q7(ID, A_even, A_odd, &Den[0], Nx, Ny, Nz);
ComputeDensityD3Q7(ID, B_even, B_odd, &Den[N], Nx, Ny, Nz);
ScaLBL_D3Q7_Density(ID, A_even, A_odd, &Den[0], Nx, Ny, Nz);
ScaLBL_D3Q7_Density(ID, B_even, B_odd, &Den[N], Nx, Ny, Nz);
//*************************************************************************
// Compute the phase indicator field
//*************************************************************************
DeviceBarrier();
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
ComputePhi(ID, Phi, Den, N);
ScaLBL_ComputePhaseField(ID, Phi, Den, N);
//*************************************************************************
ScaLBL_Comm.SendHalo(Phi);
DeviceBarrier();
ScaLBL_DeviceBarrier();
ScaLBL_Comm.RecvHalo(Phi);
//*************************************************************************
DeviceBarrier();
ScaLBL_DeviceBarrier();
// Pressure boundary conditions
if (BoundaryCondition==1 && Mask.kproc == 0) {
PressureBC_inlet(f_even,f_odd,din,Nx,Ny,Nz);
ColorBC_inlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_z(f_even,f_odd,din,Nx,Ny,Nz);
ScaLBL_Color_BC_z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
}
if (BoundaryCondition==1 && Mask.kproc == nprocz-1){
PressureBC_outlet(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ColorBC_outlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_Z(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ScaLBL_Color_BC_Z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
}
// Velocity boundary conditions
if (BoundaryCondition==2 && Mask.kproc == 0) {
ScaLBL_D3Q19_Velocity_BC_z(f_even,f_odd,din,Nx,Ny,Nz);
//ColorBC_inlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
//ScaLBL_Color_BC_z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
SetPhiSlice_z(Phi,1.0,Nx,Ny,Nz,0);
}
if (BoundaryCondition==2 && Mask.kproc == nprocz-1){
ScaLBL_D3Q19_Velocity_BC_Z(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
//ColorBC_outlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
//ScaLBL_Color_BC_Z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
SetPhiSlice_z(Phi,-1.0,Nx,Ny,Nz,Nz-1);
}
@ -841,12 +841,12 @@ int main(int argc, char **argv)
dout = 1.0-0.5*dp;
// set the initial boundary conditions
if (Mask.kproc == 0) {
PressureBC_inlet(f_even,f_odd,din,Nx,Ny,Nz);
ColorBC_inlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_z(f_even,f_odd,din,Nx,Ny,Nz);
ScaLBL_Color_BC_z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
}
if (Mask.kproc == nprocz-1){
PressureBC_outlet(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ColorBC_outlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_Z(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ScaLBL_Color_BC_Z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
}
}
@ -864,8 +864,8 @@ int main(int argc, char **argv)
if (timestep%1000 == 995){
//...........................................................................
// Copy the phase indicator field for the earlier timestep
DeviceBarrier();
CopyToHost(Averages.Phase_tplus.data(),Phi,N*sizeof(double));
ScaLBL_DeviceBarrier();
ScaLBL_CopyToHost(Averages.Phase_tplus.data(),Phi,N*sizeof(double));
// Averages.ColorToSignedDistance(beta,Averages.Phase,Averages.Phase_tplus);
//...........................................................................
}
@ -875,20 +875,20 @@ int main(int argc, char **argv)
//...........................................................................
// Copy the phase from the GPU -> CPU
//...........................................................................
DeviceBarrier();
ComputePressureD3Q19(ID,f_even,f_odd,Pressure,Nx,Ny,Nz);
CopyToHost(Averages.Phase.data(),Phi,N*sizeof(double));
CopyToHost(Averages.Press.data(),Pressure,N*sizeof(double));
CopyToHost(Averages.Vel_x.data(),&Velocity[0],N*sizeof(double));
CopyToHost(Averages.Vel_y.data(),&Velocity[N],N*sizeof(double));
CopyToHost(Averages.Vel_z.data(),&Velocity[2*N],N*sizeof(double));
ScaLBL_DeviceBarrier();
ScaLBL_D3Q19_Pressure(ID,f_even,f_odd,Pressure,Nx,Ny,Nz);
ScaLBL_CopyToHost(Averages.Phase.data(),Phi,N*sizeof(double));
ScaLBL_CopyToHost(Averages.Press.data(),Pressure,N*sizeof(double));
ScaLBL_CopyToHost(Averages.Vel_x.data(),&Velocity[0],N*sizeof(double));
ScaLBL_CopyToHost(Averages.Vel_y.data(),&Velocity[N],N*sizeof(double));
ScaLBL_CopyToHost(Averages.Vel_z.data(),&Velocity[2*N],N*sizeof(double));
MPI_Barrier(MPI_COMM_WORLD);
}
if (timestep%1000 == 5){
//...........................................................................
// Copy the phase indicator field for the later timestep
DeviceBarrier();
CopyToHost(Averages.Phase_tminus.data(),Phi,N*sizeof(double));
ScaLBL_DeviceBarrier();
ScaLBL_CopyToHost(Averages.Phase_tminus.data(),Phi,N*sizeof(double));
// Averages.ColorToSignedDistance(beta,Averages.Phase_tminus,Averages.Phase_tminus);
//....................................................................
Averages.Initialize();
@ -917,9 +917,9 @@ int main(int argc, char **argv)
double *cDistEven = new double[10*N];
double *cDistOdd = new double[9*N];
CopyToHost(cDistEven,f_even,10*N*sizeof(double));
CopyToHost(cDistOdd,f_odd,9*N*sizeof(double));
CopyToHost(cDen,Den,2*N*sizeof(double));
ScaLBL_CopyToHost(cDistEven,f_even,10*N*sizeof(double));
ScaLBL_CopyToHost(cDistOdd,f_odd,9*N*sizeof(double));
ScaLBL_CopyToHost(cDen,Den,2*N*sizeof(double));
// Read in the restart file to CPU buffers
WriteCheckpoint(LocalRestartFile, cDen, cDistEven, cDistOdd, N);
@ -936,7 +936,7 @@ int main(int argc, char **argv)
}
PROFILE_STOP("Loop");
//************************************************************************
DeviceBarrier();
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
stoptime = MPI_Wtime();
if (rank==0) printf("-------------------------------------------------------------------\n");
@ -962,8 +962,8 @@ int main(int argc, char **argv)
int NumberComponents_NWP = ComputeGlobalPhaseComponent(Mask.Nx-2,Mask.Ny-2,Mask.Nz-2,Mask.rank_info,Averages.PhaseID,1,Averages.Label_NWP);
printf("Number of non-wetting phase components: %i \n ",NumberComponents_NWP);
DeviceBarrier();
CopyToHost(Averages.Phase.data(),Phi,N*sizeof(double));
ScaLBL_DeviceBarrier();
ScaLBL_CopyToHost(Averages.Phase.data(),Phi,N*sizeof(double));
*/
/* Averages.WriteSurfaces(0);
@ -981,10 +981,10 @@ int main(int argc, char **argv)
fwrite(Averages.Press.data(),8,N,PRESS);
fclose(PRESS);
CopyToHost(Averages.Phase.data(),Phi,N*sizeof(double));
ScaLBL_CopyToHost(Averages.Phase.data(),Phi,N*sizeof(double));
double * Grad;
Grad = new double [3*N];
CopyToHost(Grad,ColorGrad,3*N*sizeof(double));
ScaLBL_CopyToHost(Grad,ColorGrad,3*N*sizeof(double));
sprintf(LocalRankFilename,"%s%s","ColorGrad.",LocalRankString);
FILE *GRAD;
GRAD = fopen(LocalRankFilename,"wb");

170
tests/lbpm_color_simulator.cpp
View File

@ -511,7 +511,7 @@ int main(int argc, char **argv)
//...........device phase ID.................................................
if (rank==0) printf ("Copy phase ID to device \n");
char *ID;
AllocateDeviceMemory((void **) &ID, N); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &ID, N); // Allocate device memory
// Don't compute in the halo
for (k=0;k<Nz;k++){
for (j=0;j<Ny;j++){
@ -522,8 +522,8 @@ int main(int argc, char **argv)
}
}
// Copy to the device
CopyToDevice(ID, id, N);
DeviceBarrier();
ScaLBL_CopyToDevice(ID, id, N);
ScaLBL_DeviceBarrier();
//...........................................................................
//...........................................................................
@ -535,26 +535,26 @@ int main(int argc, char **argv)
double *f_even,*f_odd;
double *A_even,*A_odd,*B_even,*B_odd;
//...........................................................................
AllocateDeviceMemory((void **) &f_even, 10*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &f_odd, 9*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &A_even, 4*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &A_odd, 3*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &B_even, 4*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &B_odd, 3*dist_mem_size); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &f_even, 10*dist_mem_size); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &f_odd, 9*dist_mem_size); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &A_even, 4*dist_mem_size); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &A_odd, 3*dist_mem_size); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &B_even, 4*dist_mem_size); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &B_odd, 3*dist_mem_size); // Allocate device memory
//...........................................................................
double *Phi,*Den;
double *ColorGrad, *Velocity, *Pressure, *dvcSignDist;
//...........................................................................
AllocateDeviceMemory((void **) &Phi, dist_mem_size);
AllocateDeviceMemory((void **) &Pressure, dist_mem_size);
AllocateDeviceMemory((void **) &dvcSignDist, dist_mem_size);
AllocateDeviceMemory((void **) &Den, 2*dist_mem_size);
AllocateDeviceMemory((void **) &Velocity, 3*dist_mem_size);
AllocateDeviceMemory((void **) &ColorGrad, 3*dist_mem_size);
ScaLBL_AllocateDeviceMemory((void **) &Phi, dist_mem_size);
ScaLBL_AllocateDeviceMemory((void **) &Pressure, dist_mem_size);
ScaLBL_AllocateDeviceMemory((void **) &dvcSignDist, dist_mem_size);
ScaLBL_AllocateDeviceMemory((void **) &Den, 2*dist_mem_size);
ScaLBL_AllocateDeviceMemory((void **) &Velocity, 3*dist_mem_size);
ScaLBL_AllocateDeviceMemory((void **) &ColorGrad, 3*dist_mem_size);
//...........................................................................
// Copy signed distance for device initialization
CopyToDevice(dvcSignDist, Averages->SDs.data(), dist_mem_size);
ScaLBL_CopyToDevice(dvcSignDist, Averages->SDs.data(), dist_mem_size);
//...........................................................................
int logcount = 0; // number of surface write-outs
@ -566,10 +566,10 @@ int main(int argc, char **argv)
//...........................................................................
if (rank==0) printf("Setting the distributions, size = %i\n", N);
//...........................................................................
DeviceBarrier();
InitD3Q19(ID, f_even, f_odd, Nx, Ny, Nz);
InitDenColor(ID, Den, Phi, das, dbs, Nx, Ny, Nz);
DeviceBarrier();
ScaLBL_DeviceBarrier();
ScaLBL_D3Q19_Init(ID, f_even, f_odd, Nx, Ny, Nz);
ScaLBL_Color_Init(ID, Den, Phi, das, dbs, Nx, Ny, Nz);
ScaLBL_DeviceBarrier();
//......................................................................
if (Restart == true){
@ -595,10 +595,10 @@ int main(int argc, char **argv)
double *cDistOdd = new double[9*N];
ReadCheckpoint(LocalRestartFile, cDen, cDistEven, cDistOdd, N);
// Copy the restart data to the GPU
CopyToDevice(f_even,cDistEven,10*N*sizeof(double));
CopyToDevice(f_odd,cDistOdd,9*N*sizeof(double));
CopyToDevice(Den,cDen,2*N*sizeof(double));
DeviceBarrier();
ScaLBL_CopyToDevice(f_even,cDistEven,10*N*sizeof(double));
ScaLBL_CopyToDevice(f_odd,cDistOdd,9*N*sizeof(double));
ScaLBL_CopyToDevice(Den,cDen,2*N*sizeof(double));
ScaLBL_DeviceBarrier();
delete [] cDen;
delete [] cDistEven;
delete [] cDistOdd;
@ -606,9 +606,9 @@ int main(int argc, char **argv)
}
//......................................................................
InitD3Q7(ID, A_even, A_odd, &Den[0], Nx, Ny, Nz);
InitD3Q7(ID, B_even, B_odd, &Den[N], Nx, Ny, Nz);
DeviceBarrier();
ScaLBL_D3Q7_Init(ID, A_even, A_odd, &Den[0], Nx, Ny, Nz);
ScaLBL_D3Q7_Init(ID, B_even, B_odd, &Den[N], Nx, Ny, Nz);
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
//.......................................................................
// Once phase has been initialized, map solid to account for 'smeared' interface
@ -624,12 +624,12 @@ int main(int argc, char **argv)
//*************************************************************************
// Compute the phase indicator field and reset Copy, Den
//*************************************************************************
ComputePhi(ID, Phi, Den, N);
ScaLBL_ComputePhaseField(ID, Phi, Den, N);
//*************************************************************************
DeviceBarrier();
ScaLBL_DeviceBarrier();
ScaLBL_Comm.SendHalo(Phi);
ScaLBL_Comm.RecvHalo(Phi);
DeviceBarrier();
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
//*************************************************************************
@ -638,13 +638,13 @@ int main(int argc, char **argv)
printf("Setting outlet pressure = %f \n", dout);
}
if (BoundaryCondition==1 && Mask.kproc == 0) {
PressureBC_inlet(f_even,f_odd,din,Nx,Ny,Nz);
ColorBC_inlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_z(f_even,f_odd,din,Nx,Ny,Nz);
ScaLBL_Color_BC_z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
}
if (BoundaryCondition==1 && Mask.kproc == nprocz-1){
PressureBC_outlet(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ColorBC_outlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_Z(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ScaLBL_Color_BC_Z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
}
if (rank==0 && BoundaryCondition==2){
@ -653,13 +653,13 @@ int main(int argc, char **argv)
}
if (BoundaryCondition==2 && Mask.kproc == 0) {
ScaLBL_D3Q19_Velocity_BC_z(f_even,f_odd,din,Nx,Ny,Nz);
//ColorBC_inlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
//ScaLBL_Color_BC_z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
SetPhiSlice_z(Phi,1.0,Nx,Ny,Nz,0);
}
if (BoundaryCondition==2 && Mask.kproc == nprocz-1){
ScaLBL_D3Q19_Velocity_BC_Z(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
//ColorBC_outlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
//ScaLBL_Color_BC_Z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
SetPhiSlice_z(Phi,-1.0,Nx,Ny,Nz,Nz-1);
}
@ -675,17 +675,17 @@ int main(int argc, char **argv)
dout = 1.0-0.5*dp;
// set the initial boundary conditions
if (Mask.kproc == 0) {
PressureBC_inlet(f_even,f_odd,din,Nx,Ny,Nz);
ColorBC_inlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_z(f_even,f_odd,din,Nx,Ny,Nz);
ScaLBL_Color_BC_z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
}
if (Mask.kproc == nprocz-1){
PressureBC_outlet(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ColorBC_outlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_Z(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ScaLBL_Color_BC_Z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
}
}
ComputePressureD3Q19(ID,f_even,f_odd,Pressure,Nx,Ny,Nz);
ComputeVelocityD3Q19(ID,f_even,f_odd,Velocity,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure(ID,f_even,f_odd,Pressure,Nx,Ny,Nz);
ScaLBL_D3Q19_Velocity(ID,f_even,f_odd,Velocity,Nx,Ny,Nz);
/* if (BoundaryCondition==1 && Mask.kproc == 0){
for (n=Nx*Ny; n<2*Nx*Ny; n++){
@ -701,21 +701,21 @@ int main(int argc, char **argv)
*/
//...........................................................................
// Copy the phase indicator field for the earlier timestep
DeviceBarrier();
CopyToHost(Averages->Phase_tplus.data(),Phi,N*sizeof(double));
ScaLBL_DeviceBarrier();
ScaLBL_CopyToHost(Averages->Phase_tplus.data(),Phi,N*sizeof(double));
//...........................................................................
//...........................................................................
// Copy the data for for the analysis timestep
//...........................................................................
// Copy the phase from the GPU -> CPU
//...........................................................................
DeviceBarrier();
ComputePressureD3Q19(ID,f_even,f_odd,Pressure,Nx,Ny,Nz);
CopyToHost(Averages->Phase.data(),Phi,N*sizeof(double));
CopyToHost(Averages->Press.data(),Pressure,N*sizeof(double));
CopyToHost(Averages->Vel_x.data(),&Velocity[0],N*sizeof(double));
CopyToHost(Averages->Vel_y.data(),&Velocity[N],N*sizeof(double));
CopyToHost(Averages->Vel_z.data(),&Velocity[2*N],N*sizeof(double));
ScaLBL_DeviceBarrier();
ScaLBL_D3Q19_Pressure(ID,f_even,f_odd,Pressure,Nx,Ny,Nz);
ScaLBL_CopyToHost(Averages->Phase.data(),Phi,N*sizeof(double));
ScaLBL_CopyToHost(Averages->Press.data(),Pressure,N*sizeof(double));
ScaLBL_CopyToHost(Averages->Vel_x.data(),&Velocity[0],N*sizeof(double));
ScaLBL_CopyToHost(Averages->Vel_y.data(),&Velocity[N],N*sizeof(double));
ScaLBL_CopyToHost(Averages->Vel_z.data(),&Velocity[2*N],N*sizeof(double));
//...........................................................................
if (rank==0) printf("********************************************************\n");
@ -723,7 +723,7 @@ int main(int argc, char **argv)
//.......create and start timer............
double starttime,stoptime,cputime;
DeviceBarrier();
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
starttime = MPI_Wtime();
//.........................................
@ -788,11 +788,11 @@ int main(int argc, char **argv)
//*************************************************************************
// Fused Color Gradient and Collision
//*************************************************************************
ColorCollideOpt( ID,f_even,f_odd,Phi,ColorGrad,
ScaLBL_D3Q19_ColorCollide( ID,f_even,f_odd,Phi,ColorGrad,
Velocity,Nx,Ny,Nz,rlxA,rlxB,alpha,beta,Fx,Fy,Fz);
//*************************************************************************
DeviceBarrier();
ScaLBL_DeviceBarrier();
//*************************************************************************
// Pack and send the D3Q19 distributions
ScaLBL_Comm.SendD3Q19(f_even, f_odd);
@ -801,36 +801,36 @@ int main(int argc, char **argv)
//*************************************************************************
// Carry out the density streaming step for mass transport
//*************************************************************************
MassColorCollideD3Q7(ID, A_even, A_odd, B_even, B_odd, Den, Phi,
ScaLBL_D3Q7_ColorCollideMass(ID, A_even, A_odd, B_even, B_odd, Den, Phi,
ColorGrad, Velocity, beta, N, pBC);
//*************************************************************************
DeviceBarrier();
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
//*************************************************************************
// Swap the distributions for momentum transport
//*************************************************************************
SwapD3Q19(ID, f_even, f_odd, Nx, Ny, Nz);
ScaLBL_D3Q19_Swap(ID, f_even, f_odd, Nx, Ny, Nz);
//*************************************************************************
DeviceBarrier();
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
//*************************************************************************
// Wait for communications to complete and unpack the distributions
ScaLBL_Comm.RecvD3Q19(f_even, f_odd);
//*************************************************************************
DeviceBarrier();
ScaLBL_DeviceBarrier();
//*************************************************************************
// Pack and send the D3Q7 distributions
ScaLBL_Comm.BiSendD3Q7(A_even, A_odd, B_even, B_odd);
//*************************************************************************
DeviceBarrier();
SwapD3Q7(ID, A_even, A_odd, Nx, Ny, Nz);
SwapD3Q7(ID, B_even, B_odd, Nx, Ny, Nz);
ScaLBL_DeviceBarrier();
ScaLBL_D3Q7_Swap(ID, A_even, A_odd, Nx, Ny, Nz);
ScaLBL_D3Q7_Swap(ID, B_even, B_odd, Nx, Ny, Nz);
DeviceBarrier();
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
//*************************************************************************
@ -838,44 +838,44 @@ int main(int argc, char **argv)
ScaLBL_Comm.BiRecvD3Q7(A_even, A_odd, B_even, B_odd);
//*************************************************************************
DeviceBarrier();
ScaLBL_DeviceBarrier();
//..................................................................................
ComputeDensityD3Q7(ID, A_even, A_odd, &Den[0], Nx, Ny, Nz);
ComputeDensityD3Q7(ID, B_even, B_odd, &Den[N], Nx, Ny, Nz);
ScaLBL_D3Q7_Density(ID, A_even, A_odd, &Den[0], Nx, Ny, Nz);
ScaLBL_D3Q7_Density(ID, B_even, B_odd, &Den[N], Nx, Ny, Nz);
//*************************************************************************
// Compute the phase indicator field
//*************************************************************************
DeviceBarrier();
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
ComputePhi(ID, Phi, Den, N);
ScaLBL_ComputePhaseField(ID, Phi, Den, N);
//*************************************************************************
ScaLBL_Comm.SendHalo(Phi);
DeviceBarrier();
ScaLBL_DeviceBarrier();
ScaLBL_Comm.RecvHalo(Phi);
//*************************************************************************
DeviceBarrier();
ScaLBL_DeviceBarrier();
// Pressure boundary conditions
if (BoundaryCondition==1 && Mask.kproc == 0) {
PressureBC_inlet(f_even,f_odd,din,Nx,Ny,Nz);
ColorBC_inlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_z(f_even,f_odd,din,Nx,Ny,Nz);
ScaLBL_Color_BC_z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
}
if (BoundaryCondition==1 && Mask.kproc == nprocz-1){
PressureBC_outlet(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ColorBC_outlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_Z(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ScaLBL_Color_BC_Z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
}
// Velocity boundary conditions
if (BoundaryCondition==2 && Mask.kproc == 0) {
ScaLBL_D3Q19_Velocity_BC_z(f_even,f_odd,din,Nx,Ny,Nz);
//ColorBC_inlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
//ScaLBL_Color_BC_z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
SetPhiSlice_z(Phi,1.0,Nx,Ny,Nz,0);
}
if (BoundaryCondition==2 && Mask.kproc == nprocz-1){
ScaLBL_D3Q19_Velocity_BC_Z(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
//ColorBC_outlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
//ScaLBL_Color_BC_Z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
SetPhiSlice_z(Phi,-1.0,Nx,Ny,Nz,Nz-1);
}
@ -886,12 +886,12 @@ int main(int argc, char **argv)
dout = 1.0-0.5*dp;
// set the initial boundary conditions
if (Mask.kproc == 0) {
PressureBC_inlet(f_even,f_odd,din,Nx,Ny,Nz);
ColorBC_inlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_z(f_even,f_odd,din,Nx,Ny,Nz);
ScaLBL_Color_BC_z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
}
if (Mask.kproc == nprocz-1){
PressureBC_outlet(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ColorBC_outlet(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_Z(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ScaLBL_Color_BC_Z(Phi,Den,A_even,A_odd,B_even,B_odd,Nx,Ny,Nz);
}
}
@ -923,7 +923,7 @@ int main(int argc, char **argv)
tpool.wait_pool_finished();
PROFILE_STOP("Loop");
//************************************************************************
DeviceBarrier();
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
stoptime = MPI_Wtime();
if (rank==0) printf("-------------------------------------------------------------------\n");
@ -949,8 +949,8 @@ int main(int argc, char **argv)
int NumberComponents_NWP = ComputeGlobalPhaseComponent(Mask.Nx-2,Mask.Ny-2,Mask.Nz-2,Mask.rank_info,Averages->PhaseID,1,Averages->Label_NWP);
printf("Number of non-wetting phase components: %i \n ",NumberComponents_NWP);
DeviceBarrier();
CopyToHost(Averages->Phase.data(),Phi,N*sizeof(double));
ScaLBL_DeviceBarrier();
ScaLBL_CopyToHost(Averages->Phase.data(),Phi,N*sizeof(double));
*/
/* Averages->WriteSurfaces(0);
@ -968,10 +968,10 @@ int main(int argc, char **argv)
fwrite(Averages->Press.data(),8,N,PRESS);
fclose(PRESS);
CopyToHost(Averages->Phase.data(),Phi,N*sizeof(double));
ScaLBL_CopyToHost(Averages->Phase.data(),Phi,N*sizeof(double));
double * Grad;
Grad = new double [3*N];
CopyToHost(Grad,ColorGrad,3*N*sizeof(double));
ScaLBL_CopyToHost(Grad,ColorGrad,3*N*sizeof(double));
sprintf(LocalRankFilename,"%s%s","ColorGrad.",LocalRankString);
FILE *GRAD;
GRAD = fopen(LocalRankFilename,"wb");

48
tests/lbpm_permeability_simulator.cpp
View File

@ -379,9 +379,9 @@ int main(int argc, char **argv)
//...........device phase ID.................................................
if (rank==0) printf ("Copying phase ID to device \n");
char *ID;
AllocateDeviceMemory((void **) &ID, N); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &ID, N); // Allocate device memory
// Copy to the device
CopyToDevice(ID, id, N);
ScaLBL_CopyToDevice(ID, id, N);
//...........................................................................
//...........................................................................
@ -392,18 +392,18 @@ int main(int argc, char **argv)
//......................device distributions.................................
double *f_even,*f_odd;
//...........................................................................
AllocateDeviceMemory((void **) &f_even, 10*dist_mem_size); // Allocate device memory
AllocateDeviceMemory((void **) &f_odd, 9*dist_mem_size); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &f_even, 10*dist_mem_size); // Allocate device memory
ScaLBL_AllocateDeviceMemory((void **) &f_odd, 9*dist_mem_size); // Allocate device memory
//...........................................................................
double *Velocity, *Pressure, *dvcSignDist;
//...........................................................................
AllocateDeviceMemory((void **) &Pressure, dist_mem_size);
AllocateDeviceMemory((void **) &dvcSignDist, dist_mem_size);
AllocateDeviceMemory((void **) &Velocity, 3*dist_mem_size);
ScaLBL_AllocateDeviceMemory((void **) &Pressure, dist_mem_size);
ScaLBL_AllocateDeviceMemory((void **) &dvcSignDist, dist_mem_size);
ScaLBL_AllocateDeviceMemory((void **) &Velocity, 3*dist_mem_size);
//...........................................................................
// Copy signed distance for device initialization
CopyToDevice(dvcSignDist, Averages.SDs.data(), dist_mem_size);
ScaLBL_CopyToDevice(dvcSignDist, Averages.SDs.data(), dist_mem_size);
//...........................................................................
int logcount = 0; // number of surface write-outs
@ -414,7 +414,7 @@ int main(int argc, char **argv)
//...........................................................................
if (rank==0) printf("Setting the distributions, size = %i\n", N);
//...........................................................................
InitD3Q19(ID, f_even, f_odd, Nx, Ny, Nz);
ScaLBL_D3Q19_Init(ID, f_even, f_odd, Nx, Ny, Nz);
//......................................................................
//.......................................................................
@ -446,11 +446,11 @@ int main(int argc, char **argv)
printf("Setting outlet pressure = %f \n", dout);
}
if (pBC && kproc == 0) {
PressureBC_inlet(f_even,f_odd,din,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_z(f_even,f_odd,din,Nx,Ny,Nz);
}
if (pBC && kproc == nprocz-1){
PressureBC_outlet(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ScaLBL_D3Q19_Pressure_BC_Z(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
}
int timestep = 0;
@ -472,7 +472,7 @@ int main(int argc, char **argv)
//*************************************************************************
// Fused Color Gradient and Collision
//*************************************************************************
MRT( ID,f_even,f_odd,rlxA,rlxB,Fx,Fy,Fz,Nx,Ny,Nz);
ScaLBL_D3Q19_MRT( ID,f_even,f_odd,rlxA,rlxB,Fx,Fy,Fz,Nx,Ny,Nz);
//*************************************************************************
//*************************************************************************
@ -481,21 +481,21 @@ int main(int argc, char **argv)
//*************************************************************************
// Swap the distributions for momentum transport
//*************************************************************************
SwapD3Q19(ID, f_even, f_odd, Nx, Ny, Nz);
ScaLBL_D3Q19_Swap(ID, f_even, f_odd, Nx, Ny, Nz);
//*************************************************************************
// Wait for communications to complete and unpack the distributions
ScaLBL_Comm.RecvD3Q19(f_even, f_odd);
//*************************************************************************
if (pBC && kproc == 0) {
PressureBC_inlet(f_even,f_odd,din,Nx,Ny,Nz);
ScaLBL_D3Q19_Pressure_BC_z(f_even,f_odd,din,Nx,Ny,Nz);
}
if (pBC && kproc == nprocz-1){
PressureBC_outlet(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
ScaLBL_D3Q19_Pressure_BC_Z(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2));
}
//...................................................................................
DeviceBarrier();
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
// Timestep completed!
@ -507,13 +507,13 @@ int main(int argc, char **argv)
//...........................................................................
// Copy the phase from the GPU -> CPU
//...........................................................................
DeviceBarrier();
ComputePressureD3Q19(ID,f_even,f_odd,Pressure,Nx,Ny,Nz);
ComputeVelocityD3Q19(ID,f_even,f_odd,Velocity,Nx,Ny,Nz);
CopyToHost(Averages.Press.data(),Pressure,N*sizeof(double));
CopyToHost(Averages.Vel_x.data(),&Velocity[0],N*sizeof(double));
CopyToHost(Averages.Vel_y.data(),&Velocity[N],N*sizeof(double));
CopyToHost(Averages.Vel_z.data(),&Velocity[2*N],N*sizeof(double));
ScaLBL_DeviceBarrier();
ScaLBL_D3Q19_Pressure(ID,f_even,f_odd,Pressure,Nx,Ny,Nz);
ScaLBL_D3Q19_Velocity(ID,f_even,f_odd,Velocity,Nx,Ny,Nz);
ScaLBL_CopyToHost(Averages.Press.data(),Pressure,N*sizeof(double));
ScaLBL_CopyToHost(Averages.Vel_x.data(),&Velocity[0],N*sizeof(double));
ScaLBL_CopyToHost(Averages.Vel_y.data(),&Velocity[N],N*sizeof(double));
ScaLBL_CopyToHost(Averages.Vel_z.data(),&Velocity[2*N],N*sizeof(double));
// Way more work than necessary -- this is just to get the solid interfacial area!!
Averages.Initialize();
@ -554,7 +554,7 @@ int main(int argc, char **argv)
}
}
//************************************************************************/
DeviceBarrier();
ScaLBL_DeviceBarrier();
MPI_Barrier(comm);
stoptime = MPI_Wtime();
if (rank==0) printf("-------------------------------------------------------------------\n");