From d28d4c7377e037158c253319be9b1d811712b436 Mon Sep 17 00:00:00 2001 From: James E McClure Date: Mon, 15 Jun 2015 21:37:07 -0400 Subject: [PATCH] Adding permeability simulator --- cpu/Color.cpp | 91 ---------------------- cpu/D3Q19.cpp | 92 ++++++++++++++++++++++ cpu/ScaLBL.h | 11 +++ gpu/Color.cu | 106 ------------------------- gpu/D3Q19.cu | 108 ++++++++++++++++++++++++++ tests/lbpm_permeability_simulator.cpp | 4 +- 6 files changed, 212 insertions(+), 200 deletions(-) create mode 100644 cpu/ScaLBL.h diff --git a/cpu/Color.cpp b/cpu/Color.cpp index 75e7c7d8..5a61af63 100644 --- a/cpu/Color.cpp +++ b/cpu/Color.cpp @@ -112,97 +112,6 @@ extern "C" void InitDenColorDistance(char *ID, double *Den, double *Phi, double } -extern "C" void Compute_VELOCITY(char *ID, double *disteven, double *distodd, double *vel, int Nx, int Ny, int Nz, int S) -{ - int n,N; - // distributions - double f1,f2,f3,f4,f5,f6,f7,f8,f9; - double f10,f11,f12,f13,f14,f15,f16,f17,f18; - double vx,vy,vz; - - N = Nx*Ny*Nz; - - for (n=0; n 0){ - //........................................................................ - // Registers to store the distributions - //........................................................................ - f2 = disteven[N+n]; - f4 = disteven[2*N+n]; - f6 = disteven[3*N+n]; - f8 = disteven[4*N+n]; - f10 = disteven[5*N+n]; - f12 = disteven[6*N+n]; - f14 = disteven[7*N+n]; - f16 = disteven[8*N+n]; - f18 = disteven[9*N+n]; - //........................................................................ - f1 = distodd[n]; - f3 = distodd[1*N+n]; - f5 = distodd[2*N+n]; - f7 = distodd[3*N+n]; - f9 = distodd[4*N+n]; - f11 = distodd[5*N+n]; - f13 = distodd[6*N+n]; - f15 = distodd[7*N+n]; - f17 = distodd[8*N+n]; - //.................Compute the velocity................................... - vx = f1-f2+f7-f8+f9-f10+f11-f12+f13-f14; - vy = f3-f4+f7-f8-f9+f10+f15-f16+f17-f18; - vz = f5-f6+f11-f12-f13+f14+f15-f16-f17+f18; - //..................Write the velocity..................................... - vel[n] = vx; - vel[N+n] = vy; - vel[2*N+n] = vz; - //........................................................................ - - } - } -} - -extern "C" void ComputePressureD3Q19(char *ID, double *disteven, double *distodd, double *Pressure, - int Nx, int Ny, int Nz) -{ - int n,N; - // distributions - double f0,f1,f2,f3,f4,f5,f6,f7,f8,f9; - double f10,f11,f12,f13,f14,f15,f16,f17,f18; - - N = Nx*Ny*Nz; - - for (n=0; n 0){ - //........................................................................ - // Registers to store the distributions - //........................................................................ - f0 = disteven[n]; - f2 = disteven[N+n]; - f4 = disteven[2*N+n]; - f6 = disteven[3*N+n]; - f8 = disteven[4*N+n]; - f10 = disteven[5*N+n]; - f12 = disteven[6*N+n]; - f14 = disteven[7*N+n]; - f16 = disteven[8*N+n]; - f18 = disteven[9*N+n]; - //........................................................................ - f1 = distodd[n]; - f3 = distodd[1*N+n]; - f5 = distodd[2*N+n]; - f7 = distodd[3*N+n]; - f9 = distodd[4*N+n]; - f11 = distodd[5*N+n]; - f13 = distodd[6*N+n]; - f15 = distodd[7*N+n]; - f17 = distodd[8*N+n]; - //.................Compute the velocity................................... - Pressure[n] = 0.3333333333333333*(f0+f2+f1+f4+f3+f6+f5+f8+f7+f10+ - f9+f12+f11+f14+f13+f16+f15+f18+f17); - - } - } -} //************************************************************************* extern "C" void ColorBC_inlet(double *Phi, double *Den, double *A_even, double *A_odd, diff --git a/cpu/D3Q19.cpp b/cpu/D3Q19.cpp index 90d919e0..cd590a80 100644 --- a/cpu/D3Q19.cpp +++ b/cpu/D3Q19.cpp @@ -226,3 +226,95 @@ extern "C" void SwapD3Q19(char *ID, double *disteven, double *distodd, int Nx, i } } } + +extern "C" void ComputeVelocityD319(char *ID, double *disteven, double *distodd, double *vel, int Nx, int Ny, int Nz, int S) +{ + int n,N; + // distributions + double f1,f2,f3,f4,f5,f6,f7,f8,f9; + double f10,f11,f12,f13,f14,f15,f16,f17,f18; + double vx,vy,vz; + + N = Nx*Ny*Nz; + + for (n=0; n 0){ + //........................................................................ + // Registers to store the distributions + //........................................................................ + f2 = disteven[N+n]; + f4 = disteven[2*N+n]; + f6 = disteven[3*N+n]; + f8 = disteven[4*N+n]; + f10 = disteven[5*N+n]; + f12 = disteven[6*N+n]; + f14 = disteven[7*N+n]; + f16 = disteven[8*N+n]; + f18 = disteven[9*N+n]; + //........................................................................ + f1 = distodd[n]; + f3 = distodd[1*N+n]; + f5 = distodd[2*N+n]; + f7 = distodd[3*N+n]; + f9 = distodd[4*N+n]; + f11 = distodd[5*N+n]; + f13 = distodd[6*N+n]; + f15 = distodd[7*N+n]; + f17 = distodd[8*N+n]; + //.................Compute the velocity................................... + vx = f1-f2+f7-f8+f9-f10+f11-f12+f13-f14; + vy = f3-f4+f7-f8-f9+f10+f15-f16+f17-f18; + vz = f5-f6+f11-f12-f13+f14+f15-f16-f17+f18; + //..................Write the velocity..................................... + vel[n] = vx; + vel[N+n] = vy; + vel[2*N+n] = vz; + //........................................................................ + + } + } +} + +extern "C" void ComputePressureD3Q19(char *ID, double *disteven, double *distodd, double *Pressure, + int Nx, int Ny, int Nz) +{ + int n,N; + // distributions + double f0,f1,f2,f3,f4,f5,f6,f7,f8,f9; + double f10,f11,f12,f13,f14,f15,f16,f17,f18; + + N = Nx*Ny*Nz; + + for (n=0; n 0){ + //........................................................................ + // Registers to store the distributions + //........................................................................ + f0 = disteven[n]; + f2 = disteven[N+n]; + f4 = disteven[2*N+n]; + f6 = disteven[3*N+n]; + f8 = disteven[4*N+n]; + f10 = disteven[5*N+n]; + f12 = disteven[6*N+n]; + f14 = disteven[7*N+n]; + f16 = disteven[8*N+n]; + f18 = disteven[9*N+n]; + //........................................................................ + f1 = distodd[n]; + f3 = distodd[1*N+n]; + f5 = distodd[2*N+n]; + f7 = distodd[3*N+n]; + f9 = distodd[4*N+n]; + f11 = distodd[5*N+n]; + f13 = distodd[6*N+n]; + f15 = distodd[7*N+n]; + f17 = distodd[8*N+n]; + //.................Compute the velocity................................... + Pressure[n] = 0.3333333333333333*(f0+f2+f1+f4+f3+f6+f5+f8+f7+f10+ + f9+f12+f11+f14+f13+f16+f15+f18+f17); + } + } +} + diff --git a/cpu/ScaLBL.h b/cpu/ScaLBL.h new file mode 100644 index 00000000..fe43e1bc --- /dev/null +++ b/cpu/ScaLBL.h @@ -0,0 +1,11 @@ +extern "C" void PackDist(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 InitD3Q19(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 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 InitD3Q19(char *ID, double *f_even, double *f_odd, int Nx, int Ny, int Nz); diff --git a/gpu/Color.cu b/gpu/Color.cu index 0ed32fa2..39f80e37 100644 --- a/gpu/Color.cu +++ b/gpu/Color.cu @@ -120,104 +120,6 @@ __global__ void dvc_InitDenColorDistance(char *ID, double *Den, double *Phi, do } } } - -__global__ void dvc_Compute_VELOCITY(char *ID, double *disteven, double *distodd, double *vel, int Nx, int Ny, int Nz) -{ - int n,N; - // distributions - double f1,f2,f3,f4,f5,f6,f7,f8,f9; - double f10,f11,f12,f13,f14,f15,f16,f17,f18; - double vx,vy,vz; - - N = Nx*Ny*Nz; - - int S = N/NBLOCKS/NTHREADS + 1; - for (int s=0; s 0){ - //........................................................................ - // Registers to store the distributions - //........................................................................ - f2 = disteven[N+n]; - f4 = disteven[2*N+n]; - f6 = disteven[3*N+n]; - f8 = disteven[4*N+n]; - f10 = disteven[5*N+n]; - f12 = disteven[6*N+n]; - f14 = disteven[7*N+n]; - f16 = disteven[8*N+n]; - f18 = disteven[9*N+n]; - //........................................................................ - f1 = distodd[n]; - f3 = distodd[1*N+n]; - f5 = distodd[2*N+n]; - f7 = distodd[3*N+n]; - f9 = distodd[4*N+n]; - f11 = distodd[5*N+n]; - f13 = distodd[6*N+n]; - f15 = distodd[7*N+n]; - f17 = distodd[8*N+n]; - //.................Compute the velocity................................... - vx = f1-f2+f7-f8+f9-f10+f11-f12+f13-f14; - vy = f3-f4+f7-f8-f9+f10+f15-f16+f17-f18; - vz = f5-f6+f11-f12-f13+f14+f15-f16-f17+f18; - //..................Write the velocity..................................... - vel[n] = vx; - vel[N+n] = vy; - vel[2*N+n] = vz; - //........................................................................ - - } - } -} - -__global__ void dvc_ComputePressureD3Q19(char *ID, double *disteven, double *distodd, double *Pressure, - int Nx, int Ny, int Nz) -{ - int n,N; - // distributions - double f0,f1,f2,f3,f4,f5,f6,f7,f8,f9; - double f10,f11,f12,f13,f14,f15,f16,f17,f18; - - N = Nx*Ny*Nz; - - int S = N/NBLOCKS/NTHREADS + 1; - for (int s=0; s 0){ - //........................................................................ - // Registers to store the distributions - //........................................................................ - f0 = disteven[n]; - f2 = disteven[N+n]; - f4 = disteven[2*N+n]; - f6 = disteven[3*N+n]; - f8 = disteven[4*N+n]; - f10 = disteven[5*N+n]; - f12 = disteven[6*N+n]; - f14 = disteven[7*N+n]; - f16 = disteven[8*N+n]; - f18 = disteven[9*N+n]; - //........................................................................ - f1 = distodd[n]; - f3 = distodd[1*N+n]; - f5 = distodd[2*N+n]; - f7 = distodd[3*N+n]; - f9 = distodd[4*N+n]; - f11 = distodd[5*N+n]; - f13 = distodd[6*N+n]; - f15 = distodd[7*N+n]; - f17 = distodd[8*N+n]; - //.................Compute the velocity................................... - Pressure[n] = 0.3333333333333333*(f0+f2+f1+f4+f3+f6+f5+f8+f7+f10+ - f9+f12+f11+f14+f13+f16+f15+f18+f17); - - } - } -} - //************************************************************************* __global__ void dvc_ColorBC_inlet(double *Phi, double *Den, double *A_even, double *A_odd, double *B_even, double *B_odd, int Nx, int Ny, int Nz) @@ -1469,14 +1371,6 @@ extern "C" void InitDenColorDistance(char *ID, double *Den, double *Phi, double dvc_InitDenColorDistance<<>>(ID, Den, Phi, Distance, das, dbs, beta, xp, Nx, Ny, Nz); } -extern "C" void Compute_VELOCITY(char *ID, double *disteven, double *distodd, double *vel, int Nx, int Ny, int Nz){ - - dvc_Compute_VELOCITY<<>>(ID, disteven, distodd, vel, Nx, Ny, Nz); -} -extern "C" void ComputePressureD3Q19(char *ID, double *disteven, double *distodd, double *Pressure, - int Nx, int Ny, int Nz){ - dvc_ComputePressureD3Q19<<< NBLOCKS,NTHREADS >>>(ID, disteven, distodd, Pressure, Nx, Ny, Nz); -} extern "C" void ComputeColorGradient(char *ID, double *phi, double *ColorGrad, int Nx, int Ny, int Nz){ dvc_ComputeColorGradient<<>>(ID, phi, ColorGrad, Nx, Ny, Nz); } diff --git a/gpu/D3Q19.cu b/gpu/D3Q19.cu index 82090b02..a47cdcd9 100644 --- a/gpu/D3Q19.cu +++ b/gpu/D3Q19.cu @@ -234,6 +234,106 @@ __global__ void dvc_SwapD3Q19(char *ID, double *disteven, double *distodd, int } } + +__global__ void dvc_ComputeVelocityD3Q19(char *ID, double *disteven, double *distodd, double *vel, int Nx, int Ny, int Nz) +{ + int n,N; + // distributions + double f1,f2,f3,f4,f5,f6,f7,f8,f9; + double f10,f11,f12,f13,f14,f15,f16,f17,f18; + double vx,vy,vz; + + N = Nx*Ny*Nz; + + int S = N/NBLOCKS/NTHREADS + 1; + for (int s=0; s 0){ + //........................................................................ + // Registers to store the distributions + //........................................................................ + f2 = disteven[N+n]; + f4 = disteven[2*N+n]; + f6 = disteven[3*N+n]; + f8 = disteven[4*N+n]; + f10 = disteven[5*N+n]; + f12 = disteven[6*N+n]; + f14 = disteven[7*N+n]; + f16 = disteven[8*N+n]; + f18 = disteven[9*N+n]; + //........................................................................ + f1 = distodd[n]; + f3 = distodd[1*N+n]; + f5 = distodd[2*N+n]; + f7 = distodd[3*N+n]; + f9 = distodd[4*N+n]; + f11 = distodd[5*N+n]; + f13 = distodd[6*N+n]; + f15 = distodd[7*N+n]; + f17 = distodd[8*N+n]; + //.................Compute the velocity................................... + vx = f1-f2+f7-f8+f9-f10+f11-f12+f13-f14; + vy = f3-f4+f7-f8-f9+f10+f15-f16+f17-f18; + vz = f5-f6+f11-f12-f13+f14+f15-f16-f17+f18; + //..................Write the velocity..................................... + vel[n] = vx; + vel[N+n] = vy; + vel[2*N+n] = vz; + //........................................................................ + + } + } +} + +__global__ void dvc_ComputePressureD3Q19(char *ID, double *disteven, double *distodd, double *Pressure, + int Nx, int Ny, int Nz) +{ + int n,N; + // distributions + double f0,f1,f2,f3,f4,f5,f6,f7,f8,f9; + double f10,f11,f12,f13,f14,f15,f16,f17,f18; + + N = Nx*Ny*Nz; + + int S = N/NBLOCKS/NTHREADS + 1; + for (int s=0; s 0){ + //........................................................................ + // Registers to store the distributions + //........................................................................ + f0 = disteven[n]; + f2 = disteven[N+n]; + f4 = disteven[2*N+n]; + f6 = disteven[3*N+n]; + f8 = disteven[4*N+n]; + f10 = disteven[5*N+n]; + f12 = disteven[6*N+n]; + f14 = disteven[7*N+n]; + f16 = disteven[8*N+n]; + f18 = disteven[9*N+n]; + //........................................................................ + f1 = distodd[n]; + f3 = distodd[1*N+n]; + f5 = distodd[2*N+n]; + f7 = distodd[3*N+n]; + f9 = distodd[4*N+n]; + f11 = distodd[5*N+n]; + f13 = distodd[6*N+n]; + f15 = distodd[7*N+n]; + f17 = distodd[8*N+n]; + //.................Compute the velocity................................... + Pressure[n] = 0.3333333333333333*(f0+f2+f1+f4+f3+f6+f5+f8+f7+f10+ + f9+f12+f11+f14+f13+f16+f15+f18+f17); + + } + } +} + + + extern "C" void PackDist(int q, int *list, int start, int count, double *sendbuf, double *dist, int N){ int GRID = count / 512 + 1; dvc_PackDist <<>>(q, list, start, count, sendbuf, dist, N); @@ -251,4 +351,12 @@ extern "C" void SwapD3Q19(char *ID, double *disteven, double *distodd, int Nx, i dvc_SwapD3Q19<<>>(ID, disteven, distodd, Nx, Ny, Nz); } +extern "C" void ComputeVelocityD3Q19(char *ID, double *disteven, double *distodd, double *vel, int Nx, int \ +Ny, int Nz){ + dvc_ComputeVelocityD3Q19<<>>(ID, disteven, distodd, vel, Nx, Ny, Nz); +} +extern "C" void ComputePressureD3Q19(char *ID, double *disteven, double *distodd, double *Pressure, + int Nx, int Ny, int Nz){ + dvc_ComputePressureD3Q19<<< NBLOCKS,NTHREADS >>>(ID, disteven, distodd, Pressure, Nx, Ny, Nz); +} diff --git a/tests/lbpm_permeability_simulator.cpp b/tests/lbpm_permeability_simulator.cpp index 48b58346..7e27948b 100644 --- a/tests/lbpm_permeability_simulator.cpp +++ b/tests/lbpm_permeability_simulator.cpp @@ -955,8 +955,6 @@ int main(int argc, char **argv) PressureBC_outlet(f_even,f_odd,dout,Nx,Ny,Nz,Nx*Ny*(Nz-2)); } - - int timestep = 0; if (rank==0) printf("********************************************************\n"); if (rank==0) printf("No. of timesteps: %i \n", timestepMax); @@ -1186,8 +1184,8 @@ int main(int argc, char **argv) //........................................................................... DeviceBarrier(); ComputePressureD3Q19(ID,f_even,f_odd,Pressure,Nx,Ny,Nz); + ComputeVelocityD3Q19(ID,f_even,f_odd,Velocity,Nx,Ny,Nz); MPI_Barrier(MPI_COMM_WORLD); - } } //************************************************************************/