diff --git a/src/gpu/core/GPU/GPU_Interface.h b/src/gpu/core/GPU/GPU_Interface.h
index 00894db28d2aebd098862eb92e0f408b9fd25040..72e839bcf741b9966a4e38a7b78c1085b20132e2 100644
--- a/src/gpu/core/GPU/GPU_Interface.h
+++ b/src/gpu/core/GPU/GPU_Interface.h
@@ -27,53 +27,6 @@
struct LBMSimulationParameter;
class Parameter;
-//////////////////////////////////////////////////////////////////////////
-//Kernel
-//////////////////////////////////////////////////////////////////////////
-void Init27(int myid,
- int numprocs,
- real u0,
- unsigned int* geoD,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- real* vParab,
- unsigned long long numberOfLBnodes,
- unsigned int grid_nx,
- unsigned int grid_ny,
- unsigned int grid_nz,
- real* DD,
- int level,
- int maxlevel);
-
-void InitNonEqPartSP27(unsigned int numberOfThreads,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* neighborWSB,
- unsigned int* geoD,
- real* rho,
- real* ux,
- real* uy,
- real* uz,
- unsigned long long numberOfLBnodes,
- real* DD,
- real omega,
- bool EvenOrOdd);
-
-
-void InitADDev27( unsigned int numberOfThreads,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* geoD,
- real* Conc,
- real* ux,
- real* uy,
- real* uz,
- unsigned long long numberOfLBnodes,
- real* DD27,
- bool EvenOrOdd);
void CalcMac27( real* vxD,
real* vyD,
diff --git a/src/gpu/core/GPU/GPU_Kernels.cuh b/src/gpu/core/GPU/GPU_Kernels.cuh
index 2fd05b9e7d605958a3d7d222a1346b6913d7d92d..ab8ab0fa3564de523e8fc8e6985f6d98c09fb404 100644
--- a/src/gpu/core/GPU/GPU_Kernels.cuh
+++ b/src/gpu/core/GPU/GPU_Kernels.cuh
@@ -17,48 +17,6 @@
#include "LBM/LB.h"
-__global__ void LBInit27( int myid,
- int numprocs,
- real u0,
- unsigned int* geoD,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- real* vParabel,
- unsigned long long numberOfLBnodes,
- unsigned int grid_nx,
- unsigned int grid_ny,
- unsigned int grid_nz,
- real* DD,
- int lev,
- int maxlev);
-
-__global__ void LBInitNonEqPartSP27(unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* neighborWSB,
- unsigned int* geoD,
- real* rho,
- real* ux,
- real* uy,
- real* uz,
- unsigned long long numberOfLBnodes,
- real* DD,
- real omega,
- bool EvenOrOdd);
-
-__global__ void InitAD27(unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* geoD,
- real* Conc,
- real* ux,
- real* uy,
- real* uz,
- unsigned long long numberOfLBnodes,
- real* DD27,
- bool EvenOrOdd);
-
__global__ void LBCalcMac27( real* vxD,
real* vyD,
real* vzD,
diff --git a/src/gpu/core/GPU/Init27.cu b/src/gpu/core/GPU/Init27.cu
deleted file mode 100644
index 961045d3c91eff5de62eccb68ec3ce8b148ae906..0000000000000000000000000000000000000000
--- a/src/gpu/core/GPU/Init27.cu
+++ /dev/null
@@ -1,512 +0,0 @@
-/* Device code */
-#include "LBM/LB.h"
-#include "lbm/constants/D3Q27.h"
-#include <basics/constants/NumericConstants.h>
-
-using namespace vf::basics::constant;
-using namespace vf::lbm::dir;
-
-////////////////////////////////////////////////////////////////////////////////
-__global__ void LBInit27( int myid,
- int numprocs,
- real u0,
- unsigned int* geoD,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- real* vParabel,
- unsigned long long numberOfLBnodes,
- unsigned int grid_nx,
- unsigned int grid_ny,
- unsigned int grid_nz,
- real* DD,
- int lev,
- int maxlev)
-{
- Distributions27 D;
- D.f[dP00] = &DD[dP00 * numberOfLBnodes];
- D.f[dM00] = &DD[dM00 * numberOfLBnodes];
- D.f[d0P0] = &DD[d0P0 * numberOfLBnodes];
- D.f[d0M0] = &DD[d0M0 * numberOfLBnodes];
- D.f[d00P] = &DD[d00P * numberOfLBnodes];
- D.f[d00M] = &DD[d00M * numberOfLBnodes];
- D.f[dPP0] = &DD[dPP0 * numberOfLBnodes];
- D.f[dMM0] = &DD[dMM0 * numberOfLBnodes];
- D.f[dPM0] = &DD[dPM0 * numberOfLBnodes];
- D.f[dMP0] = &DD[dMP0 * numberOfLBnodes];
- D.f[dP0P] = &DD[dP0P * numberOfLBnodes];
- D.f[dM0M] = &DD[dM0M * numberOfLBnodes];
- D.f[dP0M] = &DD[dP0M * numberOfLBnodes];
- D.f[dM0P] = &DD[dM0P * numberOfLBnodes];
- D.f[d0PP] = &DD[d0PP * numberOfLBnodes];
- D.f[d0MM] = &DD[d0MM * numberOfLBnodes];
- D.f[d0PM] = &DD[d0PM * numberOfLBnodes];
- D.f[d0MP] = &DD[d0MP * numberOfLBnodes];
- D.f[d000] = &DD[d000 * numberOfLBnodes];
- D.f[dPPP] = &DD[dPPP * numberOfLBnodes];
- D.f[dMMP] = &DD[dMMP * numberOfLBnodes];
- D.f[dPMP] = &DD[dPMP * numberOfLBnodes];
- D.f[dMPP] = &DD[dMPP * numberOfLBnodes];
- D.f[dPPM] = &DD[dPPM * numberOfLBnodes];
- D.f[dMMM] = &DD[dMMM * numberOfLBnodes];
- D.f[dPMM] = &DD[dPMM * numberOfLBnodes];
- D.f[dMPM] = &DD[dMPM * numberOfLBnodes];
- ////////////////////////////////////////////////////////////////////////////////
- unsigned int k; // Zugriff auf arrays im device
- //
- unsigned int tx = threadIdx.x; // Thread index = lokaler i index
- unsigned int by = blockIdx.x; // Block index x
- unsigned int bz = blockIdx.y; // Block index y
- unsigned int x = tx + STARTOFFX; // Globaler x-Index
- unsigned int y = by + STARTOFFY; // Globaler y-Index
- unsigned int z = bz + STARTOFFZ; // Globaler z-Index
-
- const unsigned sizeX = blockDim.x;
- const unsigned sizeY = gridDim.x;
- const unsigned nx = sizeX + 2 * STARTOFFX;
- const unsigned ny = sizeY + 2 * STARTOFFY;
-
- k = nx*(ny*z + y) + x;
- //////////////////////////////////////////////////////////////////////////
- geoD[k] = GEO_FLUID;
- if (lev==0)
- {
- if( by == 0 || by == grid_ny-1 || tx == 0 || tx == grid_nx-1 )
- geoD[k] = GEO_SOLID;
- else if( bz == grid_nz-1 && myid == numprocs - 1 && geoD[k] != GEO_SOLID )
- geoD[k] = GEO_PRESS;
- else if( bz == 0 && myid == 0 && geoD[k] != GEO_SOLID)
- geoD[k] = GEO_SOLID;//GEO_VELO;
- }
- else if (lev==maxlev-1)
- {
- unsigned int centerX = grid_nx / 2;
- unsigned int centerY = grid_ny / 2;
- unsigned int centerZ = grid_nz / 2;
- real radius = grid_ny / 2.56;
-
- unsigned int distSq = (centerX-tx)*(centerX-tx)+(centerY-by)*(centerY-by)+(centerZ-bz)*(centerZ-bz);
- real radiSq = radius*radius;
-
- if( distSq < radiSq) geoD[k] = GEO_SOLID;
- }
- //////////////////////////////////////////////////////////////////////////
- real drho = c0o1;
- real vx1 = c0o1;
- real vx2 = c0o1;
- real vx3 = u0;
- vParabel[k] = vx3;
- ////////////////////////////////////////////////////////////////////////////////
- //index
- unsigned int nxny = nx*ny;
- ////////////////////////////////////////////////////////////////////////////////
- //neighborX[k] = k+1;
- //neighborY[k+1] = k+nx+1;
- //neighborZ[k+1] = k+nxny+1;
- //neighborY[k] = k+nx;
- //neighborX[k+nx] = k+nx+1;
- //neighborZ[k+nx] = k+nx+nxny;
- //neighborZ[k] = k+nxny;
- //neighborX[k+nxny] = k+nxny+1;
- //neighborY[k+nxny] = k+nxny+nx;
- ////////////////////////////////////////////////////////////////////////////////
- unsigned int kzero= k;
- unsigned int ke = k;
- unsigned int kw = k + 1;
- unsigned int kn = k;
- unsigned int ks = k + nx;
- unsigned int kt = k;
- unsigned int kb = k + nxny;
- unsigned int ksw = k + nx + 1;
- unsigned int kne = k;
- unsigned int kse = k + nx;
- unsigned int knw = k + 1;
- unsigned int kbw = k + nxny + 1;
- unsigned int kte = k;
- unsigned int kbe = k + nxny;
- unsigned int ktw = k + 1;
- unsigned int kbs = k + nxny + nx;
- unsigned int ktn = k;
- unsigned int kbn = k + nxny;
- unsigned int kts = k + nx;
- unsigned int ktse = k + nx;
- unsigned int kbnw = k + nxny + 1;
- unsigned int ktnw = k + 1;
- unsigned int kbse = k + nxny + nx;
- unsigned int ktsw = k + nx + 1;
- unsigned int kbne = k + nxny;
- unsigned int ktne = k;
- unsigned int kbsw = k + nxny + nx + 1;
- //////////////////////////////////////////////////////////////////////////
-
- real cu_sq=c3o2*(vx1*vx1+vx2*vx2+vx3*vx3);
-
- (D.f[d000])[kzero] = c8o27* (drho-cu_sq);
- (D.f[dP00])[ke ] = c2o27* (drho+c3o1*( vx1 )+c9o2*( vx1 )*( vx1 )-cu_sq);
- (D.f[dM00])[kw ] = c2o27* (drho+c3o1*(-vx1 )+c9o2*(-vx1 )*(-vx1 )-cu_sq);
- (D.f[d0P0])[kn ] = c2o27* (drho+c3o1*( vx2 )+c9o2*( vx2 )*( vx2 )-cu_sq);
- (D.f[d0M0])[ks ] = c2o27* (drho+c3o1*( -vx2 )+c9o2*( -vx2 )*( -vx2 )-cu_sq);
- (D.f[d00P])[kt ] = c2o27* (drho+c3o1*( vx3)+c9o2*( vx3)*( vx3)-cu_sq);
- (D.f[d00M])[kb ] = c2o27* (drho+c3o1*( -vx3)+c9o2*( -vx3)*( -vx3)-cu_sq);
- (D.f[dPP0])[kne ] = c1o54* (drho+c3o1*( vx1+vx2 )+c9o2*( vx1+vx2 )*( vx1+vx2 )-cu_sq);
- (D.f[dMM0])[ksw ] = c1o54* (drho+c3o1*(-vx1-vx2 )+c9o2*(-vx1-vx2 )*(-vx1-vx2 )-cu_sq);
- (D.f[dPM0])[kse ] = c1o54* (drho+c3o1*( vx1-vx2 )+c9o2*( vx1-vx2 )*( vx1-vx2 )-cu_sq);
- (D.f[dMP0])[knw ] = c1o54* (drho+c3o1*(-vx1+vx2 )+c9o2*(-vx1+vx2 )*(-vx1+vx2 )-cu_sq);
- (D.f[dP0P])[kte ] = c1o54* (drho+c3o1*( vx1 +vx3)+c9o2*( vx1 +vx3)*( vx1 +vx3)-cu_sq);
- (D.f[dM0M])[kbw ] = c1o54* (drho+c3o1*(-vx1 -vx3)+c9o2*(-vx1 -vx3)*(-vx1 -vx3)-cu_sq);
- (D.f[dP0M])[kbe ] = c1o54* (drho+c3o1*( vx1 -vx3)+c9o2*( vx1 -vx3)*( vx1 -vx3)-cu_sq);
- (D.f[dM0P])[ktw ] = c1o54* (drho+c3o1*(-vx1 +vx3)+c9o2*(-vx1 +vx3)*(-vx1 +vx3)-cu_sq);
- (D.f[d0PP])[ktn ] = c1o54* (drho+c3o1*( vx2+vx3)+c9o2*( vx2+vx3)*( vx2+vx3)-cu_sq);
- (D.f[d0MM])[kbs ] = c1o54* (drho+c3o1*( -vx2-vx3)+c9o2*( -vx2-vx3)*( -vx2-vx3)-cu_sq);
- (D.f[d0PM])[kbn ] = c1o54* (drho+c3o1*( vx2-vx3)+c9o2*( vx2-vx3)*( vx2-vx3)-cu_sq);
- (D.f[d0MP])[kts ] = c1o54* (drho+c3o1*( -vx2+vx3)+c9o2*( -vx2+vx3)*( -vx2+vx3)-cu_sq);
- (D.f[dPPP])[ktne ] = c1o216*(drho+c3o1*( vx1+vx2+vx3)+c9o2*( vx1+vx2+vx3)*( vx1+vx2+vx3)-cu_sq);
- (D.f[dMMM])[kbsw ] = c1o216*(drho+c3o1*(-vx1-vx2-vx3)+c9o2*(-vx1-vx2-vx3)*(-vx1-vx2-vx3)-cu_sq);
- (D.f[dPPM])[kbne ] = c1o216*(drho+c3o1*( vx1+vx2-vx3)+c9o2*( vx1+vx2-vx3)*( vx1+vx2-vx3)-cu_sq);
- (D.f[dMMP])[ktsw ] = c1o216*(drho+c3o1*(-vx1-vx2+vx3)+c9o2*(-vx1-vx2+vx3)*(-vx1-vx2+vx3)-cu_sq);
- (D.f[dPMP])[ktse ] = c1o216*(drho+c3o1*( vx1-vx2+vx3)+c9o2*( vx1-vx2+vx3)*( vx1-vx2+vx3)-cu_sq);
- (D.f[dMPM])[kbnw ] = c1o216*(drho+c3o1*(-vx1+vx2-vx3)+c9o2*(-vx1+vx2-vx3)*(-vx1+vx2-vx3)-cu_sq);
- (D.f[dPMM])[kbse ] = c1o216*(drho+c3o1*( vx1-vx2-vx3)+c9o2*( vx1-vx2-vx3)*( vx1-vx2-vx3)-cu_sq);
- (D.f[dMPP])[ktnw ] = c1o216*(drho+c3o1*(-vx1+vx2+vx3)+c9o2*(-vx1+vx2+vx3)*(-vx1+vx2+vx3)-cu_sq);
-
-}
-////////////////////////////////////////////////////////////////////////////////
-
-
-
-
-
-
-
-
-
-
-////////////////////////////////////////////////////////////////////////////////
-__global__ void LBInitNonEqPartSP27( unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* neighborWSB,
- unsigned int* geoD,
- real* rho,
- real* ux,
- real* uy,
- real* uz,
- unsigned long long numberOfLBnodes,
- real* DD,
- real omega,
- bool EvenOrOdd)
-{
- ////////////////////////////////////////////////////////////////////////////////
- const unsigned x = threadIdx.x; // Globaler x-Index
- const unsigned y = blockIdx.x; // Globaler y-Index
- const unsigned z = blockIdx.y; // Globaler z-Index
-
- const unsigned nx = blockDim.x;
- const unsigned ny = gridDim.x;
-
- const unsigned k = nx*(ny*z + y) + x;
- //////////////////////////////////////////////////////////////////////////
-
- if(k<numberOfLBnodes)
- {
- ////////////////////////////////////////////////////////////////////////////////
- unsigned int BC;
- BC = geoD[k];
-
- if( BC != GEO_SOLID && BC != GEO_VOID)
- {
- Distributions27 D;
- if (EvenOrOdd==true)
- {
- D.f[dP00] = &DD[dP00 * numberOfLBnodes];
- D.f[dM00] = &DD[dM00 * numberOfLBnodes];
- D.f[d0P0] = &DD[d0P0 * numberOfLBnodes];
- D.f[d0M0] = &DD[d0M0 * numberOfLBnodes];
- D.f[d00P] = &DD[d00P * numberOfLBnodes];
- D.f[d00M] = &DD[d00M * numberOfLBnodes];
- D.f[dPP0] = &DD[dPP0 * numberOfLBnodes];
- D.f[dMM0] = &DD[dMM0 * numberOfLBnodes];
- D.f[dPM0] = &DD[dPM0 * numberOfLBnodes];
- D.f[dMP0] = &DD[dMP0 * numberOfLBnodes];
- D.f[dP0P] = &DD[dP0P * numberOfLBnodes];
- D.f[dM0M] = &DD[dM0M * numberOfLBnodes];
- D.f[dP0M] = &DD[dP0M * numberOfLBnodes];
- D.f[dM0P] = &DD[dM0P * numberOfLBnodes];
- D.f[d0PP] = &DD[d0PP * numberOfLBnodes];
- D.f[d0MM] = &DD[d0MM * numberOfLBnodes];
- D.f[d0PM] = &DD[d0PM * numberOfLBnodes];
- D.f[d0MP] = &DD[d0MP * numberOfLBnodes];
- D.f[d000] = &DD[d000 * numberOfLBnodes];
- D.f[dPPP] = &DD[dPPP * numberOfLBnodes];
- D.f[dMMP] = &DD[dMMP * numberOfLBnodes];
- D.f[dPMP] = &DD[dPMP * numberOfLBnodes];
- D.f[dMPP] = &DD[dMPP * numberOfLBnodes];
- D.f[dPPM] = &DD[dPPM * numberOfLBnodes];
- D.f[dMMM] = &DD[dMMM * numberOfLBnodes];
- D.f[dPMM] = &DD[dPMM * numberOfLBnodes];
- D.f[dMPM] = &DD[dMPM * numberOfLBnodes];
- }
- else
- {
- D.f[dM00] = &DD[dP00 * numberOfLBnodes];
- D.f[dP00] = &DD[dM00 * numberOfLBnodes];
- D.f[d0M0] = &DD[d0P0 * numberOfLBnodes];
- D.f[d0P0] = &DD[d0M0 * numberOfLBnodes];
- D.f[d00M] = &DD[d00P * numberOfLBnodes];
- D.f[d00P] = &DD[d00M * numberOfLBnodes];
- D.f[dMM0] = &DD[dPP0 * numberOfLBnodes];
- D.f[dPP0] = &DD[dMM0 * numberOfLBnodes];
- D.f[dMP0] = &DD[dPM0 * numberOfLBnodes];
- D.f[dPM0] = &DD[dMP0 * numberOfLBnodes];
- D.f[dM0M] = &DD[dP0P * numberOfLBnodes];
- D.f[dP0P] = &DD[dM0M * numberOfLBnodes];
- D.f[dM0P] = &DD[dP0M * numberOfLBnodes];
- D.f[dP0M] = &DD[dM0P * numberOfLBnodes];
- D.f[d0MM] = &DD[d0PP * numberOfLBnodes];
- D.f[d0PP] = &DD[d0MM * numberOfLBnodes];
- D.f[d0MP] = &DD[d0PM * numberOfLBnodes];
- D.f[d0PM] = &DD[d0MP * numberOfLBnodes];
- D.f[d000] = &DD[d000 * numberOfLBnodes];
- D.f[dMMM] = &DD[dPPP * numberOfLBnodes];
- D.f[dPPM] = &DD[dMMP * numberOfLBnodes];
- D.f[dMPM] = &DD[dPMP * numberOfLBnodes];
- D.f[dPMM] = &DD[dMPP * numberOfLBnodes];
- D.f[dMMP] = &DD[dPPM * numberOfLBnodes];
- D.f[dPPP] = &DD[dMMM * numberOfLBnodes];
- D.f[dMPP] = &DD[dPMM * numberOfLBnodes];
- D.f[dPMP] = &DD[dMPM * numberOfLBnodes];
- }
- //////////////////////////////////////////////////////////////////////////
- real drho = rho[k];//0.0f;//
- real vx1 = ux[k]; //0.0f;//
- real vx2 = uy[k]; //0.0f;//
- real vx3 = uz[k]; //0.0f;//
- //////////////////////////////////////////////////////////////////////////
- //index
- //////////////////////////////////////////////////////////////////////////
- unsigned int kzero= k;
- unsigned int ke = k;
- unsigned int kw = neighborX[k];
- unsigned int kn = k;
- unsigned int ks = neighborY[k];
- unsigned int kt = k;
- unsigned int kb = neighborZ[k];
- unsigned int ksw = neighborY[kw];
- unsigned int kne = k;
- unsigned int kse = ks;
- unsigned int knw = kw;
- unsigned int kbw = neighborZ[kw];
- unsigned int kte = k;
- unsigned int kbe = kb;
- unsigned int ktw = kw;
- unsigned int kbs = neighborZ[ks];
- unsigned int ktn = k;
- unsigned int kbn = kb;
- unsigned int kts = ks;
- unsigned int ktse = ks;
- unsigned int kbnw = kbw;
- unsigned int ktnw = kw;
- unsigned int kbse = kbs;
- unsigned int ktsw = ksw;
- unsigned int kbne = kb;
- unsigned int ktne = k;
- unsigned int kbsw = neighborZ[ksw];
- //////////////////////////////////////////////////////////////////////////////
- //neighbor index
- uint kPx = neighborX[k];
- uint kPy = neighborY[k];
- uint kPz = neighborZ[k];
- uint kMxyz = neighborWSB[k];
- uint kMx = neighborZ[neighborY[kMxyz]];
- uint kMy = neighborZ[neighborX[kMxyz]];
- uint kMz = neighborY[neighborX[kMxyz]];
- //////////////////////////////////////////////////////////////////////////
- //getVeloX//
- real vx1NeighborPx = ux[kPx];
- real vx1NeighborMx = ux[kMx];
- real vx1NeighborPy = ux[kPy];
- real vx1NeighborMy = ux[kMy];
- real vx1NeighborPz = ux[kPz];
- real vx1NeighborMz = ux[kMz];
- //getVeloY//
- real vx2NeighborPx = uy[kPx];
- real vx2NeighborMx = uy[kMx];
- real vx2NeighborPy = uy[kPy];
- real vx2NeighborMy = uy[kMy];
- real vx2NeighborPz = uy[kPz];
- real vx2NeighborMz = uy[kMz];
- //getVeloZ//
- real vx3NeighborPx = uz[kPx];
- real vx3NeighborMx = uz[kMx];
- real vx3NeighborPy = uz[kPy];
- real vx3NeighborMy = uz[kMy];
- real vx3NeighborPz = uz[kPz];
- real vx3NeighborMz = uz[kMz];
- //////////////////////////////////////////////////////////////////////////
-
- real dvx1dx = (vx1NeighborPx - vx1NeighborMx) / c2o1;
- real dvx1dy = (vx1NeighborPy - vx1NeighborMy) / c2o1;
- real dvx1dz = (vx1NeighborPz - vx1NeighborMz) / c2o1;
-
- real dvx2dx = (vx2NeighborPx - vx2NeighborMx) / c2o1;
- real dvx2dy = (vx2NeighborPy - vx2NeighborMy) / c2o1;
- real dvx2dz = (vx2NeighborPz - vx2NeighborMz) / c2o1;
-
- real dvx3dx = (vx3NeighborPx - vx3NeighborMx) / c2o1;
- real dvx3dy = (vx3NeighborPy - vx3NeighborMy) / c2o1;
- real dvx3dz = (vx3NeighborPz - vx3NeighborMz) / c2o1;
-
- //////////////////////////////////////////////////////////////////////////
-
- // the following code is copy and pasted from VirtualFluidsVisitors/InitDistributionsBlockVisitor.cpp
- // i.e. Konstantins code
-
- real ax = dvx1dx;
- real ay = dvx1dy;
- real az = dvx1dz;
-
- real bx = dvx2dx;
- real by = dvx2dy;
- real bz = dvx2dz;
-
- real cx = dvx3dx;
- real cy = dvx3dy;
- real cz = dvx3dz;
-
- real eps_new = c1o1;
- real op = c1o1;
- real o = omega;
-
- real f_E = eps_new *((5.*ax*omega + 5.*by*o + 5.*cz*o - 8.*ax*op + 4.*by*op + 4.*cz*op)/(54.*o*op));
-
- real f_N = f_E + eps_new *((2.*(ax - by))/(9.*o));
- real f_T = f_E + eps_new *((2.*(ax - cz))/(9.*o));
- real f_NE = eps_new *(-(5.*cz*o + 3.*(ay + bx)*op - 2.*cz*op + ax*(5.*o + op) + by*(5.*o + op))/(54.*o*op));
- real f_SE = f_NE + eps_new *(( ay + bx )/(9.*o));
- real f_TE = eps_new *(-(5.*cz*o + by*(5.*o - 2.*op) + 3.*(az + cx)*op + cz*op + ax*(5.*o + op))/(54.*o*op));
- real f_BE = f_TE + eps_new *(( az + cx )/(9.*o));
- real f_TN = eps_new *(-(5.*ax*o + 5.*by*o + 5.*cz*o - 2.*ax*op + by*op + 3.*bz*op + 3.*cy*op + cz*op)/(54.*o*op));
- real f_BN = f_TN + eps_new *(( bz + cy )/(9.*o));
- real f_ZERO = eps_new *((5.*(ax + by + cz))/(9.*op));
- real f_TNE = eps_new *(-(ay + az + bx + bz + cx + cy)/(72.*o));
- real f_TSW = - f_TNE - eps_new *((ay + bx)/(36.*o));
- real f_TSE = - f_TNE - eps_new *((az + cx)/(36.*o));
- real f_TNW = - f_TNE - eps_new *((bz + cy)/(36.*o));
-
- //////////////////////////////////////////////////////////////////////////
- real cu_sq=c3o2*(vx1*vx1+vx2*vx2+vx3*vx3);
-
- (D.f[d000])[kzero] = c8o27* (drho-cu_sq);
- (D.f[dP00])[ke ] = c2o27* (drho+c3o1*( vx1 )+c9o2*( vx1 )*( vx1 )-cu_sq);
- (D.f[dM00])[kw ] = c2o27* (drho+c3o1*(-vx1 )+c9o2*(-vx1 )*(-vx1 )-cu_sq);
- (D.f[d0P0])[kn ] = c2o27* (drho+c3o1*( vx2 )+c9o2*( vx2 )*( vx2 )-cu_sq);
- (D.f[d0M0])[ks ] = c2o27* (drho+c3o1*( -vx2 )+c9o2*( -vx2 )*( -vx2 )-cu_sq);
- (D.f[d00P])[kt ] = c2o27* (drho+c3o1*( vx3)+c9o2*( vx3)*( vx3)-cu_sq);
- (D.f[d00M])[kb ] = c2o27* (drho+c3o1*( -vx3)+c9o2*( -vx3)*( -vx3)-cu_sq);
- (D.f[dPP0])[kne ] = c1o54* (drho+c3o1*( vx1+vx2 )+c9o2*( vx1+vx2 )*( vx1+vx2 )-cu_sq);
- (D.f[dMM0])[ksw ] = c1o54* (drho+c3o1*(-vx1-vx2 )+c9o2*(-vx1-vx2 )*(-vx1-vx2 )-cu_sq);
- (D.f[dPM0])[kse ] = c1o54* (drho+c3o1*( vx1-vx2 )+c9o2*( vx1-vx2 )*( vx1-vx2 )-cu_sq);
- (D.f[dMP0])[knw ] = c1o54* (drho+c3o1*(-vx1+vx2 )+c9o2*(-vx1+vx2 )*(-vx1+vx2 )-cu_sq);
- (D.f[dP0P])[kte ] = c1o54* (drho+c3o1*( vx1 +vx3)+c9o2*( vx1 +vx3)*( vx1 +vx3)-cu_sq);
- (D.f[dM0M])[kbw ] = c1o54* (drho+c3o1*(-vx1 -vx3)+c9o2*(-vx1 -vx3)*(-vx1 -vx3)-cu_sq);
- (D.f[dP0M])[kbe ] = c1o54* (drho+c3o1*( vx1 -vx3)+c9o2*( vx1 -vx3)*( vx1 -vx3)-cu_sq);
- (D.f[dM0P])[ktw ] = c1o54* (drho+c3o1*(-vx1 +vx3)+c9o2*(-vx1 +vx3)*(-vx1 +vx3)-cu_sq);
- (D.f[d0PP])[ktn ] = c1o54* (drho+c3o1*( vx2+vx3)+c9o2*( vx2+vx3)*( vx2+vx3)-cu_sq);
- (D.f[d0MM])[kbs ] = c1o54* (drho+c3o1*( -vx2-vx3)+c9o2*( -vx2-vx3)*( -vx2-vx3)-cu_sq);
- (D.f[d0PM])[kbn ] = c1o54* (drho+c3o1*( vx2-vx3)+c9o2*( vx2-vx3)*( vx2-vx3)-cu_sq);
- (D.f[d0MP])[kts ] = c1o54* (drho+c3o1*( -vx2+vx3)+c9o2*( -vx2+vx3)*( -vx2+vx3)-cu_sq);
- (D.f[dPPP])[ktne ] = c1o216*(drho+c3o1*( vx1+vx2+vx3)+c9o2*( vx1+vx2+vx3)*( vx1+vx2+vx3)-cu_sq);
- (D.f[dMMM])[kbsw ] = c1o216*(drho+c3o1*(-vx1-vx2-vx3)+c9o2*(-vx1-vx2-vx3)*(-vx1-vx2-vx3)-cu_sq);
- (D.f[dPPM])[kbne ] = c1o216*(drho+c3o1*( vx1+vx2-vx3)+c9o2*( vx1+vx2-vx3)*( vx1+vx2-vx3)-cu_sq);
- (D.f[dMMP])[ktsw ] = c1o216*(drho+c3o1*(-vx1-vx2+vx3)+c9o2*(-vx1-vx2+vx3)*(-vx1-vx2+vx3)-cu_sq);
- (D.f[dPMP])[ktse ] = c1o216*(drho+c3o1*( vx1-vx2+vx3)+c9o2*( vx1-vx2+vx3)*( vx1-vx2+vx3)-cu_sq);
- (D.f[dMPM])[kbnw ] = c1o216*(drho+c3o1*(-vx1+vx2-vx3)+c9o2*(-vx1+vx2-vx3)*(-vx1+vx2-vx3)-cu_sq);
- (D.f[dPMM])[kbse ] = c1o216*(drho+c3o1*( vx1-vx2-vx3)+c9o2*( vx1-vx2-vx3)*( vx1-vx2-vx3)-cu_sq);
- (D.f[dMPP])[ktnw ] = c1o216*(drho+c3o1*(-vx1+vx2+vx3)+c9o2*(-vx1+vx2+vx3)*(-vx1+vx2+vx3)-cu_sq);
-
- //////////////////////////////////////////////////////////////////////////
-
- (D.f[d000])[kzero] += f_ZERO;
- (D.f[dP00])[ke ] += f_E ;
- (D.f[dM00])[kw ] += f_E ;
- (D.f[d0P0])[kn ] += f_N ;
- (D.f[d0M0])[ks ] += f_N ;
- (D.f[d00P])[kt ] += f_T ;
- (D.f[d00M])[kb ] += f_T ;
- (D.f[dPP0])[kne ] += f_NE ;
- (D.f[dMM0])[ksw ] += f_NE ;
- (D.f[dPM0])[kse ] += f_SE ;
- (D.f[dMP0])[knw ] += f_SE ;
- (D.f[dP0P])[kte ] += f_TE ;
- (D.f[dM0M])[kbw ] += f_TE ;
- (D.f[dP0M])[kbe ] += f_BE ;
- (D.f[dM0P])[ktw ] += f_BE ;
- (D.f[d0PP])[ktn ] += f_TN ;
- (D.f[d0MM])[kbs ] += f_TN ;
- (D.f[d0PM])[kbn ] += f_BN ;
- (D.f[d0MP])[kts ] += f_BN ;
- (D.f[dPPP])[ktne ] += f_TNE ;
- (D.f[dMMM])[kbsw ] += f_TNE ;
- (D.f[dPPM])[kbne ] += f_TSW ;
- (D.f[dMMP])[ktsw ] += f_TSW ;
- (D.f[dPMP])[ktse ] += f_TSE ;
- (D.f[dMPM])[kbnw ] += f_TSE ;
- (D.f[dPMM])[kbse ] += f_TNW ;
- (D.f[dMPP])[ktnw ] += f_TNW ;
-
- //////////////////////////////////////////////////////////////////////////
- }
- else
- {
- //////////////////////////////////////////////////////////////////////////
- Distributions27 D;
- D.f[d000] = &DD[d000 * numberOfLBnodes];
- //////////////////////////////////////////////////////////////////////////
- (D.f[d000])[k] = c96o1;
- //////////////////////////////////////////////////////////////////////////
- }
- }
-}
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
diff --git a/src/gpu/core/GPU/InitAdvectionDiffusion27.cu b/src/gpu/core/GPU/InitAdvectionDiffusion27.cu
deleted file mode 100644
index 8e2118e25f0337cc0d3fb8898ea26dab74ad15c5..0000000000000000000000000000000000000000
--- a/src/gpu/core/GPU/InitAdvectionDiffusion27.cu
+++ /dev/null
@@ -1,529 +0,0 @@
-//=======================================================================================
-// ____ ____ __ ______ __________ __ __ __ __
-// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
-// \ \ | | | | | |_) | | | | | | | / \ | |
-// \ \ | | | | | _ / | | | | | | / /\ \ | |
-// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
-// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
-// \ \ | | ________________________________________________________________
-// \ \ | | | ______________________________________________________________|
-// \ \| | | | __ __ __ __ ______ _______
-// \ | | |_____ | | | | | | | | | _ \ / _____)
-// \ | | _____| | | | | | | | | | | \ \ \_______
-// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
-// \ _____| |__| |________| \_______/ |__| |______/ (_______/
-//
-// This file is part of VirtualFluids. VirtualFluids is free software: you can
-// redistribute it and/or modify it under the terms of the GNU General Public
-// License as published by the Free Software Foundation, either version 3 of
-// the License, or (at your option) any later version.
-//
-// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
-// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
-// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
-// for more details.
-//
-// You should have received a copy of the GNU General Public License along
-// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
-//
-//! \file InitAdvectionDiffusion.cu
-//! \ingroup GPU
-//! \author Martin Schoenherr
-//=======================================================================================
-/* Device code */
-#include "LBM/LB.h"
-#include "lbm/constants/D3Q27.h"
-#include <basics/constants/NumericConstants.h>
-
-using namespace vf::basics::constant;
-using namespace vf::lbm::dir;
-
-__global__ void InitAD27(
- uint* neighborX,
- uint* neighborY,
- uint* neighborZ,
- uint* typeOfGridNode,
- real* concentration,
- real* velocityX,
- real* velocityY,
- real* velocityZ,
- unsigned long long numberOfLBnodes,
- real* distributionsAD,
- bool isEvenTimestep)
-{
- //////////////////////////////////////////////////////////////////////////
- //! The initialization is executed in the following steps
- //!
- ////////////////////////////////////////////////////////////////////////////////
- //! - Get node index coordinates from threadIdx, blockIdx, blockDim and gridDim.
- //!
- const unsigned x = threadIdx.x; // Globaler x-Index
- const unsigned y = blockIdx.x; // Globaler y-Index
- const unsigned z = blockIdx.y; // Globaler z-Index
-
- const unsigned nx = blockDim.x;
- const unsigned ny = gridDim.x;
-
- const unsigned k = nx*(ny*z + y) + x;
-
- //////////////////////////////////////////////////////////////////////////
- // run for all indices in size_Mat and fluid nodes
- if ((k < numberOfLBnodes) && (typeOfGridNode[k] == GEO_FLUID))
- {
- //////////////////////////////////////////////////////////////////////////
- //! - Read distributions: style of reading and writing the distributions from/to stored arrays dependent on timestep is based on the esoteric twist algorithm \ref
- //! <a href="https://doi.org/10.3390/computation5020019"><b>[ M. Geier et al. (2017), DOI:10.3390/computation5020019 ]</b></a>
- //!
- Distributions27 distAD;
- if (isEvenTimestep)
- {
- distAD.f[dP00] = &distributionsAD[dP00 * numberOfLBnodes];
- distAD.f[dM00] = &distributionsAD[dM00 * numberOfLBnodes];
- distAD.f[d0P0] = &distributionsAD[d0P0 * numberOfLBnodes];
- distAD.f[d0M0] = &distributionsAD[d0M0 * numberOfLBnodes];
- distAD.f[d00P] = &distributionsAD[d00P * numberOfLBnodes];
- distAD.f[d00M] = &distributionsAD[d00M * numberOfLBnodes];
- distAD.f[dPP0] = &distributionsAD[dPP0 * numberOfLBnodes];
- distAD.f[dMM0] = &distributionsAD[dMM0 * numberOfLBnodes];
- distAD.f[dPM0] = &distributionsAD[dPM0 * numberOfLBnodes];
- distAD.f[dMP0] = &distributionsAD[dMP0 * numberOfLBnodes];
- distAD.f[dP0P] = &distributionsAD[dP0P * numberOfLBnodes];
- distAD.f[dM0M] = &distributionsAD[dM0M * numberOfLBnodes];
- distAD.f[dP0M] = &distributionsAD[dP0M * numberOfLBnodes];
- distAD.f[dM0P] = &distributionsAD[dM0P * numberOfLBnodes];
- distAD.f[d0PP] = &distributionsAD[d0PP * numberOfLBnodes];
- distAD.f[d0MM] = &distributionsAD[d0MM * numberOfLBnodes];
- distAD.f[d0PM] = &distributionsAD[d0PM * numberOfLBnodes];
- distAD.f[d0MP] = &distributionsAD[d0MP * numberOfLBnodes];
- distAD.f[d000] = &distributionsAD[d000 * numberOfLBnodes];
- distAD.f[dPPP] = &distributionsAD[dPPP * numberOfLBnodes];
- distAD.f[dMMP] = &distributionsAD[dMMP * numberOfLBnodes];
- distAD.f[dPMP] = &distributionsAD[dPMP * numberOfLBnodes];
- distAD.f[dMPP] = &distributionsAD[dMPP * numberOfLBnodes];
- distAD.f[dPPM] = &distributionsAD[dPPM * numberOfLBnodes];
- distAD.f[dMMM] = &distributionsAD[dMMM * numberOfLBnodes];
- distAD.f[dPMM] = &distributionsAD[dPMM * numberOfLBnodes];
- distAD.f[dMPM] = &distributionsAD[dMPM * numberOfLBnodes];
- }
- else
- {
- distAD.f[dM00] = &distributionsAD[dP00 * numberOfLBnodes];
- distAD.f[dP00] = &distributionsAD[dM00 * numberOfLBnodes];
- distAD.f[d0M0] = &distributionsAD[d0P0 * numberOfLBnodes];
- distAD.f[d0P0] = &distributionsAD[d0M0 * numberOfLBnodes];
- distAD.f[d00M] = &distributionsAD[d00P * numberOfLBnodes];
- distAD.f[d00P] = &distributionsAD[d00M * numberOfLBnodes];
- distAD.f[dMM0] = &distributionsAD[dPP0 * numberOfLBnodes];
- distAD.f[dPP0] = &distributionsAD[dMM0 * numberOfLBnodes];
- distAD.f[dMP0] = &distributionsAD[dPM0 * numberOfLBnodes];
- distAD.f[dPM0] = &distributionsAD[dMP0 * numberOfLBnodes];
- distAD.f[dM0M] = &distributionsAD[dP0P * numberOfLBnodes];
- distAD.f[dP0P] = &distributionsAD[dM0M * numberOfLBnodes];
- distAD.f[dM0P] = &distributionsAD[dP0M * numberOfLBnodes];
- distAD.f[dP0M] = &distributionsAD[dM0P * numberOfLBnodes];
- distAD.f[d0MM] = &distributionsAD[d0PP * numberOfLBnodes];
- distAD.f[d0PP] = &distributionsAD[d0MM * numberOfLBnodes];
- distAD.f[d0MP] = &distributionsAD[d0PM * numberOfLBnodes];
- distAD.f[d0PM] = &distributionsAD[d0MP * numberOfLBnodes];
- distAD.f[d000] = &distributionsAD[d000 * numberOfLBnodes];
- distAD.f[dMMM] = &distributionsAD[dPPP * numberOfLBnodes];
- distAD.f[dPPM] = &distributionsAD[dMMP * numberOfLBnodes];
- distAD.f[dMPM] = &distributionsAD[dPMP * numberOfLBnodes];
- distAD.f[dPMM] = &distributionsAD[dMPP * numberOfLBnodes];
- distAD.f[dMMP] = &distributionsAD[dPPM * numberOfLBnodes];
- distAD.f[dPPP] = &distributionsAD[dMMM * numberOfLBnodes];
- distAD.f[dMPP] = &distributionsAD[dPMM * numberOfLBnodes];
- distAD.f[dPMP] = &distributionsAD[dMPM * numberOfLBnodes];
- }
- //////////////////////////////////////////////////////////////////////////
- //! - Set local velocities and concetration
- //!
- real conc = concentration[k];
- real vx1 = velocityX[k];
- real vx2 = velocityY[k];
- real vx3 = velocityZ[k];
- //////////////////////////////////////////////////////////////////////////
- //! - Set neighbor indices (necessary for indirect addressing)
- //!
- uint kzero = k;
- uint ke = k;
- uint kw = neighborX[k];
- uint kn = k;
- uint ks = neighborY[k];
- uint kt = k;
- uint kb = neighborZ[k];
- uint ksw = neighborY[kw];
- uint kne = k;
- uint kse = ks;
- uint knw = kw;
- uint kbw = neighborZ[kw];
- uint kte = k;
- uint kbe = kb;
- uint ktw = kw;
- uint kbs = neighborZ[ks];
- uint ktn = k;
- uint kbn = kb;
- uint kts = ks;
- uint ktse = ks;
- uint kbnw = kbw;
- uint ktnw = kw;
- uint kbse = kbs;
- uint ktsw = ksw;
- uint kbne = kb;
- uint ktne = k;
- uint kbsw = neighborZ[ksw];
- //////////////////////////////////////////////////////////////////////////
- //! - Calculate the equilibrium and set the distributions
- //!
- real cu_sq = c3o2*(vx1*vx1 + vx2*vx2 + vx3*vx3);
-
- (distAD.f[d000])[kzero] = c8o27 * conc * (c1o1 - cu_sq);
- (distAD.f[dP00])[ke ] = c2o27 * conc * (c1o1 + c3o1 * ( vx1 ) + c9o2 * ( vx1 ) * ( vx1 ) - cu_sq);
- (distAD.f[dM00])[kw ] = c2o27 * conc * (c1o1 + c3o1 * (-vx1 ) + c9o2 * (-vx1 ) * (-vx1 ) - cu_sq);
- (distAD.f[d0P0])[kn ] = c2o27 * conc * (c1o1 + c3o1 * ( vx2 ) + c9o2 * ( vx2 ) * ( vx2 ) - cu_sq);
- (distAD.f[d0M0])[ks ] = c2o27 * conc * (c1o1 + c3o1 * ( - vx2 ) + c9o2 * ( - vx2 ) * ( - vx2 ) - cu_sq);
- (distAD.f[d00P])[kt ] = c2o27 * conc * (c1o1 + c3o1 * ( vx3) + c9o2 * ( vx3) * ( vx3) - cu_sq);
- (distAD.f[d00M])[kb ] = c2o27 * conc * (c1o1 + c3o1 * ( - vx3) + c9o2 * ( - vx3) * ( - vx3) - cu_sq);
- (distAD.f[dPP0])[kne ] = c1o54 * conc * (c1o1 + c3o1 * ( vx1 + vx2 ) + c9o2 * ( vx1 + vx2 ) * ( vx1 + vx2 ) - cu_sq);
- (distAD.f[dMM0])[ksw ] = c1o54 * conc * (c1o1 + c3o1 * (-vx1 - vx2 ) + c9o2 * (-vx1 - vx2 ) * (-vx1 - vx2 ) - cu_sq);
- (distAD.f[dPM0])[kse ] = c1o54 * conc * (c1o1 + c3o1 * ( vx1 - vx2 ) + c9o2 * ( vx1 - vx2 ) * ( vx1 - vx2 ) - cu_sq);
- (distAD.f[dMP0])[knw ] = c1o54 * conc * (c1o1 + c3o1 * (-vx1 + vx2 ) + c9o2 * (-vx1 + vx2 ) * (-vx1 + vx2 ) - cu_sq);
- (distAD.f[dP0P])[kte ] = c1o54 * conc * (c1o1 + c3o1 * ( vx1 + vx3) + c9o2 * ( vx1 + vx3) * ( vx1 + vx3) - cu_sq);
- (distAD.f[dM0M])[kbw ] = c1o54 * conc * (c1o1 + c3o1 * (-vx1 - vx3) + c9o2 * (-vx1 - vx3) * (-vx1 - vx3) - cu_sq);
- (distAD.f[dP0M])[kbe ] = c1o54 * conc * (c1o1 + c3o1 * ( vx1 - vx3) + c9o2 * ( vx1 - vx3) * ( vx1 - vx3) - cu_sq);
- (distAD.f[dM0P])[ktw ] = c1o54 * conc * (c1o1 + c3o1 * (-vx1 + vx3) + c9o2 * (-vx1 + vx3) * (-vx1 + vx3) - cu_sq);
- (distAD.f[d0PP])[ktn ] = c1o54 * conc * (c1o1 + c3o1 * ( vx2 + vx3) + c9o2 * ( vx2 + vx3) * ( vx2 + vx3) - cu_sq);
- (distAD.f[d0MM])[kbs ] = c1o54 * conc * (c1o1 + c3o1 * ( - vx2 - vx3) + c9o2 * ( - vx2 - vx3) * ( - vx2 - vx3) - cu_sq);
- (distAD.f[d0PM])[kbn ] = c1o54 * conc * (c1o1 + c3o1 * ( vx2 - vx3) + c9o2 * ( vx2 - vx3) * ( vx2 - vx3) - cu_sq);
- (distAD.f[d0MP])[kts ] = c1o54 * conc * (c1o1 + c3o1 * ( - vx2 + vx3) + c9o2 * ( - vx2 + vx3) * ( - vx2 + vx3) - cu_sq);
- (distAD.f[dPPP])[ktne ] = c1o216 * conc * (c1o1 + c3o1 * ( vx1 + vx2 + vx3) + c9o2 * ( vx1 + vx2 + vx3) * ( vx1 + vx2 + vx3) - cu_sq);
- (distAD.f[dMMM])[kbsw ] = c1o216 * conc * (c1o1 + c3o1 * (-vx1 - vx2 - vx3) + c9o2 * (-vx1 - vx2 - vx3) * (-vx1 - vx2 - vx3) - cu_sq);
- (distAD.f[dPPM])[kbne ] = c1o216 * conc * (c1o1 + c3o1 * ( vx1 + vx2 - vx3) + c9o2 * ( vx1 + vx2 - vx3) * ( vx1 + vx2 - vx3) - cu_sq);
- (distAD.f[dMMP])[ktsw ] = c1o216 * conc * (c1o1 + c3o1 * (-vx1 - vx2 + vx3) + c9o2 * (-vx1 - vx2 + vx3) * (-vx1 - vx2 + vx3) - cu_sq);
- (distAD.f[dPMP])[ktse ] = c1o216 * conc * (c1o1 + c3o1 * ( vx1 - vx2 + vx3) + c9o2 * ( vx1 - vx2 + vx3) * ( vx1 - vx2 + vx3) - cu_sq);
- (distAD.f[dMPM])[kbnw ] = c1o216 * conc * (c1o1 + c3o1 * (-vx1 + vx2 - vx3) + c9o2 * (-vx1 + vx2 - vx3) * (-vx1 + vx2 - vx3) - cu_sq);
- (distAD.f[dPMM])[kbse ] = c1o216 * conc * (c1o1 + c3o1 * ( vx1 - vx2 - vx3) + c9o2 * ( vx1 - vx2 - vx3) * ( vx1 - vx2 - vx3) - cu_sq);
- (distAD.f[dMPP])[ktnw ] = c1o216 * conc * (c1o1 + c3o1 * (-vx1 + vx2 + vx3) + c9o2 * (-vx1 + vx2 + vx3) * (-vx1 + vx2 + vx3) - cu_sq);
- }
-}
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-// DEPRECATED (2022)
-
-// ////////////////////////////////////////////////////////////////////////////////
-// __global__ void InitAD27(unsigned int* neighborX,
-// unsigned int* neighborY,
-// unsigned int* neighborZ,
-// unsigned int* geoD,
-// real* Conc,
-// real* ux,
-// real* uy,
-// real* uz,
-// unsigned int size_Mat,
-// real* DD27,
-// bool EvenOrOdd)
-// {
-// ////////////////////////////////////////////////////////////////////////////////
-// const unsigned x = threadIdx.x; // Globaler x-Index
-// const unsigned y = blockIdx.x; // Globaler y-Index
-// const unsigned z = blockIdx.y; // Globaler z-Index
-
-// const unsigned nx = blockDim.x;
-// const unsigned ny = gridDim.x;
-
-// const unsigned k = nx*(ny*z + y) + x;
-// //////////////////////////////////////////////////////////////////////////
-
-// if(k<size_Mat)
-// {
-// ////////////////////////////////////////////////////////////////////////////////
-// unsigned int BC;
-// BC = geoD[k];
-
-// if( BC != GEO_SOLID && BC != GEO_VOID)
-// {
-// Distributions27 D27;
-// if (EvenOrOdd==true)
-// {
-// D27.f[dP00] = &DD27[dP00 * size_Mat];
-// D27.f[dM00] = &DD27[dM00 * size_Mat];
-// D27.f[d0P0] = &DD27[d0P0 * size_Mat];
-// D27.f[d0M0] = &DD27[d0M0 * size_Mat];
-// D27.f[d00P] = &DD27[d00P * size_Mat];
-// D27.f[d00M] = &DD27[d00M * size_Mat];
-// D27.f[dPP0] = &DD27[dPP0 * size_Mat];
-// D27.f[dMM0] = &DD27[dMM0 * size_Mat];
-// D27.f[dPM0] = &DD27[dPM0 * size_Mat];
-// D27.f[dMP0] = &DD27[dMP0 * size_Mat];
-// D27.f[dP0P] = &DD27[dP0P * size_Mat];
-// D27.f[dM0M] = &DD27[dM0M * size_Mat];
-// D27.f[dP0M] = &DD27[dP0M * size_Mat];
-// D27.f[dM0P] = &DD27[dM0P * size_Mat];
-// D27.f[d0PP] = &DD27[d0PP * size_Mat];
-// D27.f[d0MM] = &DD27[d0MM * size_Mat];
-// D27.f[d0PM] = &DD27[d0PM * size_Mat];
-// D27.f[d0MP] = &DD27[d0MP * size_Mat];
-// D27.f[d000] = &DD27[d000 * size_Mat];
-// D27.f[dPPP] = &DD27[dPPP * size_Mat];
-// D27.f[dMMP] = &DD27[dMMP * size_Mat];
-// D27.f[dPMP] = &DD27[dPMP * size_Mat];
-// D27.f[dMPP] = &DD27[dMPP * size_Mat];
-// D27.f[dPPM] = &DD27[dPPM * size_Mat];
-// D27.f[dMMM] = &DD27[dMMM * size_Mat];
-// D27.f[dPMM] = &DD27[dPMM * size_Mat];
-// D27.f[dMPM] = &DD27[dMPM * size_Mat];
-// }
-// else
-// {
-// D27.f[dM00] = &DD27[dP00 * size_Mat];
-// D27.f[dP00] = &DD27[dM00 * size_Mat];
-// D27.f[d0M0] = &DD27[d0P0 * size_Mat];
-// D27.f[d0P0] = &DD27[d0M0 * size_Mat];
-// D27.f[d00M] = &DD27[d00P * size_Mat];
-// D27.f[d00P] = &DD27[d00M * size_Mat];
-// D27.f[dMM0] = &DD27[dPP0 * size_Mat];
-// D27.f[dPP0] = &DD27[dMM0 * size_Mat];
-// D27.f[dMP0] = &DD27[dPM0 * size_Mat];
-// D27.f[dPM0] = &DD27[dMP0 * size_Mat];
-// D27.f[dM0M] = &DD27[dP0P * size_Mat];
-// D27.f[dP0P] = &DD27[dM0M * size_Mat];
-// D27.f[dM0P] = &DD27[dP0M * size_Mat];
-// D27.f[dP0M] = &DD27[dM0P * size_Mat];
-// D27.f[d0MM] = &DD27[d0PP * size_Mat];
-// D27.f[d0PP] = &DD27[d0MM * size_Mat];
-// D27.f[d0MP] = &DD27[d0PM * size_Mat];
-// D27.f[d0PM] = &DD27[d0MP * size_Mat];
-// D27.f[d000] = &DD27[d000 * size_Mat];
-// D27.f[dMMM] = &DD27[dPPP * size_Mat];
-// D27.f[dPPM] = &DD27[dMMP * size_Mat];
-// D27.f[dMPM] = &DD27[dPMP * size_Mat];
-// D27.f[dPMM] = &DD27[dMPP * size_Mat];
-// D27.f[dMMP] = &DD27[dPPM * size_Mat];
-// D27.f[dPPP] = &DD27[dMMM * size_Mat];
-// D27.f[dMPP] = &DD27[dPMM * size_Mat];
-// D27.f[dPMP] = &DD27[dMPM * size_Mat];
-// }
-// //////////////////////////////////////////////////////////////////////////
-// real ConcD = Conc[k];
-// real vx1 = ux[k];
-// real vx2 = uy[k];
-// real vx3 = uz[k];
-// //real lambdaD = -three + sqrt(three);
-// //real Diffusivity = c1o20;
-// //real Lam = -(c1o2+one/lambdaD);
-// //real nue_d = Lam/three;
-// //real ae = Diffusivity/nue_d - one;
-// //real ux_sq = vx1 * vx1;
-// //real uy_sq = vx2 * vx2;
-// //real uz_sq = vx3 * vx3;
-// ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
-// //D3Q7
-// ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
-// //index
-// //unsigned int kzero= k;
-// //unsigned int ke = k;
-// //unsigned int kw = neighborX[k];
-// //unsigned int kn = k;
-// //unsigned int ks = neighborY[k];
-// //unsigned int kt = k;
-// //unsigned int kb = neighborZ[k];
-// //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
-// //(D7.f[0])[kzero] = ConcD*(c1o3*(ae*(-three))-(ux_sq+uy_sq+uz_sq));
-// //(D7.f[1])[ke ] = ConcD*(c1o6*(ae+one)+c1o2*(ux_sq)+vx1*c1o2);
-// //(D7.f[2])[kw ] = ConcD*(c1o6*(ae+one)+c1o2*(ux_sq)-vx1*c1o2);
-// //(D7.f[3])[kn ] = ConcD*(c1o6*(ae+one)+c1o2*(uy_sq)+vx2*c1o2);
-// //(D7.f[4])[ks ] = ConcD*(c1o6*(ae+one)+c1o2*(uy_sq)-vx2*c1o2);
-// //(D7.f[5])[kt ] = ConcD*(c1o6*(ae+one)+c1o2*(uz_sq)+vx3*c1o2);
-// //(D7.f[6])[kb ] = ConcD*(c1o6*(ae+one)+c1o2*(uz_sq)-vx3*c1o2);
-// ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
-
-
-// ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
-// //D3Q27
-// ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
-// //index
-// unsigned int kzero= k;
-// unsigned int ke = k;
-// unsigned int kw = neighborX[k];
-// unsigned int kn = k;
-// unsigned int ks = neighborY[k];
-// unsigned int kt = k;
-// unsigned int kb = neighborZ[k];
-// unsigned int ksw = neighborY[kw];
-// unsigned int kne = k;
-// unsigned int kse = ks;
-// unsigned int knw = kw;
-// unsigned int kbw = neighborZ[kw];
-// unsigned int kte = k;
-// unsigned int kbe = kb;
-// unsigned int ktw = kw;
-// unsigned int kbs = neighborZ[ks];
-// unsigned int ktn = k;
-// unsigned int kbn = kb;
-// unsigned int kts = ks;
-// unsigned int ktse = ks;
-// unsigned int kbnw = kbw;
-// unsigned int ktnw = kw;
-// unsigned int kbse = kbs;
-// unsigned int ktsw = ksw;
-// unsigned int kbne = kb;
-// unsigned int ktne = k;
-// unsigned int kbsw = neighborZ[ksw];
-// ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
-// real cu_sq=c3o2*(vx1*vx1+vx2*vx2+vx3*vx3);
-
-// (D27.f[d000])[kzero] = c8o27* ConcD*(c1o1-cu_sq);
-// (D27.f[dP00])[ke ] = c2o27* ConcD*(c1o1+c3o1*( vx1 )+c9o2*( vx1 )*( vx1 )-cu_sq);
-// (D27.f[dM00])[kw ] = c2o27* ConcD*(c1o1+c3o1*(-vx1 )+c9o2*(-vx1 )*(-vx1 )-cu_sq);
-// (D27.f[d0P0])[kn ] = c2o27* ConcD*(c1o1+c3o1*( vx2 )+c9o2*( vx2 )*( vx2 )-cu_sq);
-// (D27.f[d0M0])[ks ] = c2o27* ConcD*(c1o1+c3o1*( -vx2 )+c9o2*( -vx2 )*( -vx2 )-cu_sq);
-// (D27.f[d00P])[kt ] = c2o27* ConcD*(c1o1+c3o1*( vx3)+c9o2*( vx3)*( vx3)-cu_sq);
-// (D27.f[d00M])[kb ] = c2o27* ConcD*(c1o1+c3o1*( -vx3)+c9o2*( -vx3)*( -vx3)-cu_sq);
-// (D27.f[dPP0])[kne ] = c1o54* ConcD*(c1o1+c3o1*( vx1+vx2 )+c9o2*( vx1+vx2 )*( vx1+vx2 )-cu_sq);
-// (D27.f[dMM0])[ksw ] = c1o54* ConcD*(c1o1+c3o1*(-vx1-vx2 )+c9o2*(-vx1-vx2 )*(-vx1-vx2 )-cu_sq);
-// (D27.f[dPM0])[kse ] = c1o54* ConcD*(c1o1+c3o1*( vx1-vx2 )+c9o2*( vx1-vx2 )*( vx1-vx2 )-cu_sq);
-// (D27.f[dMP0])[knw ] = c1o54* ConcD*(c1o1+c3o1*(-vx1+vx2 )+c9o2*(-vx1+vx2 )*(-vx1+vx2 )-cu_sq);
-// (D27.f[dP0P])[kte ] = c1o54* ConcD*(c1o1+c3o1*( vx1 +vx3)+c9o2*( vx1 +vx3)*( vx1 +vx3)-cu_sq);
-// (D27.f[dM0M])[kbw ] = c1o54* ConcD*(c1o1+c3o1*(-vx1 -vx3)+c9o2*(-vx1 -vx3)*(-vx1 -vx3)-cu_sq);
-// (D27.f[dP0M])[kbe ] = c1o54* ConcD*(c1o1+c3o1*( vx1 -vx3)+c9o2*( vx1 -vx3)*( vx1 -vx3)-cu_sq);
-// (D27.f[dM0P])[ktw ] = c1o54* ConcD*(c1o1+c3o1*(-vx1 +vx3)+c9o2*(-vx1 +vx3)*(-vx1 +vx3)-cu_sq);
-// (D27.f[d0PP])[ktn ] = c1o54* ConcD*(c1o1+c3o1*( vx2+vx3)+c9o2*( vx2+vx3)*( vx2+vx3)-cu_sq);
-// (D27.f[d0MM])[kbs ] = c1o54* ConcD*(c1o1+c3o1*( -vx2-vx3)+c9o2*( -vx2-vx3)*( -vx2-vx3)-cu_sq);
-// (D27.f[d0PM])[kbn ] = c1o54* ConcD*(c1o1+c3o1*( vx2-vx3)+c9o2*( vx2-vx3)*( vx2-vx3)-cu_sq);
-// (D27.f[d0MP])[kts ] = c1o54* ConcD*(c1o1+c3o1*( -vx2+vx3)+c9o2*( -vx2+vx3)*( -vx2+vx3)-cu_sq);
-// (D27.f[dPPP])[ktne ] = c1o216*ConcD*(c1o1+c3o1*( vx1+vx2+vx3)+c9o2*( vx1+vx2+vx3)*( vx1+vx2+vx3)-cu_sq);
-// (D27.f[dMMM])[kbsw ] = c1o216*ConcD*(c1o1+c3o1*(-vx1-vx2-vx3)+c9o2*(-vx1-vx2-vx3)*(-vx1-vx2-vx3)-cu_sq);
-// (D27.f[dPPM])[kbne ] = c1o216*ConcD*(c1o1+c3o1*( vx1+vx2-vx3)+c9o2*( vx1+vx2-vx3)*( vx1+vx2-vx3)-cu_sq);
-// (D27.f[dMMP])[ktsw ] = c1o216*ConcD*(c1o1+c3o1*(-vx1-vx2+vx3)+c9o2*(-vx1-vx2+vx3)*(-vx1-vx2+vx3)-cu_sq);
-// (D27.f[dPMP])[ktse ] = c1o216*ConcD*(c1o1+c3o1*( vx1-vx2+vx3)+c9o2*( vx1-vx2+vx3)*( vx1-vx2+vx3)-cu_sq);
-// (D27.f[dMPM])[kbnw ] = c1o216*ConcD*(c1o1+c3o1*(-vx1+vx2-vx3)+c9o2*(-vx1+vx2-vx3)*(-vx1+vx2-vx3)-cu_sq);
-// (D27.f[dPMM])[kbse ] = c1o216*ConcD*(c1o1+c3o1*( vx1-vx2-vx3)+c9o2*( vx1-vx2-vx3)*( vx1-vx2-vx3)-cu_sq);
-// (D27.f[dMPP])[ktnw ] = c1o216*ConcD*(c1o1+c3o1*(-vx1+vx2+vx3)+c9o2*(-vx1+vx2+vx3)*(-vx1+vx2+vx3)-cu_sq);
-// ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
-// }
-// }
-// }
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-////////////////////////////////////////////////////////////////////////////////
-__global__ void InitAD7( unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* geoD,
- real* Conc,
- real* ux,
- real* uy,
- real* uz,
- unsigned long long numberOfLBnodes,
- real* DD7,
- bool EvenOrOdd)
-{
- ////////////////////////////////////////////////////////////////////////////////
- const unsigned x = threadIdx.x; // Globaler x-Index
- const unsigned y = blockIdx.x; // Globaler y-Index
- const unsigned z = blockIdx.y; // Globaler z-Index
-
- const unsigned nx = blockDim.x;
- const unsigned ny = gridDim.x;
-
- const unsigned k = nx*(ny*z + y) + x;
- //////////////////////////////////////////////////////////////////////////
-
- if(k<numberOfLBnodes)
- {
- ////////////////////////////////////////////////////////////////////////////////
- unsigned int BC;
- BC = geoD[k];
-
- if( BC != GEO_SOLID && BC != GEO_VOID)
- {
- Distributions7 D7;
- if (EvenOrOdd==true)
- {
- D7.f[0] = &DD7[0*numberOfLBnodes];
- D7.f[1] = &DD7[1*numberOfLBnodes];
- D7.f[2] = &DD7[2*numberOfLBnodes];
- D7.f[3] = &DD7[3*numberOfLBnodes];
- D7.f[4] = &DD7[4*numberOfLBnodes];
- D7.f[5] = &DD7[5*numberOfLBnodes];
- D7.f[6] = &DD7[6*numberOfLBnodes];
- }
- else
- {
- D7.f[0] = &DD7[0*numberOfLBnodes];
- D7.f[2] = &DD7[1*numberOfLBnodes];
- D7.f[1] = &DD7[2*numberOfLBnodes];
- D7.f[4] = &DD7[3*numberOfLBnodes];
- D7.f[3] = &DD7[4*numberOfLBnodes];
- D7.f[6] = &DD7[5*numberOfLBnodes];
- D7.f[5] = &DD7[6*numberOfLBnodes];
- }
- //////////////////////////////////////////////////////////////////////////
- real ConcD = Conc[k];
- real vx1 = ux[k];
- real vx2 = uy[k];
- real vx3 = uz[k];
- real lambdaD = -c3o1 + sqrt(c3o1);
- real Diffusivity = c1o20;
- real Lam = -(c1o2+c1o1/lambdaD);
- real nue_d = Lam/c3o1;
- real ae = Diffusivity/nue_d - c1o1;
- real ux_sq = vx1 * vx1;
- real uy_sq = vx2 * vx2;
- real uz_sq = vx3 * vx3;
- //////////////////////////////////////////////////////////////////////////
- //index
- //////////////////////////////////////////////////////////////////////////
- unsigned int kzero= k;
- unsigned int ke = k;
- unsigned int kw = neighborX[k];
- unsigned int kn = k;
- unsigned int ks = neighborY[k];
- unsigned int kt = k;
- unsigned int kb = neighborZ[k];
- //////////////////////////////////////////////////////////////////////////
-
- (D7.f[0])[kzero] = ConcD*(c1o3*(ae*(-c3o1))-(ux_sq+uy_sq+uz_sq));
- (D7.f[1])[ke ] = ConcD*(c1o6*(ae+c1o1)+c1o2*(ux_sq)+vx1*c1o2);
- (D7.f[2])[kw ] = ConcD*(c1o6*(ae+c1o1)+c1o2*(ux_sq)-vx1*c1o2);
- (D7.f[3])[kn ] = ConcD*(c1o6*(ae+c1o1)+c1o2*(uy_sq)+vx2*c1o2);
- (D7.f[4])[ks ] = ConcD*(c1o6*(ae+c1o1)+c1o2*(uy_sq)-vx2*c1o2);
- (D7.f[5])[kt ] = ConcD*(c1o6*(ae+c1o1)+c1o2*(uz_sq)+vx3*c1o2);
- (D7.f[6])[kb ] = ConcD*(c1o6*(ae+c1o1)+c1o2*(uz_sq)-vx3*c1o2);
- }
- }
-}
\ No newline at end of file
diff --git a/src/gpu/core/GPU/LBMKernel.cu b/src/gpu/core/GPU/LBMKernel.cu
index 67aa3aead41b0d14c8dacec833f3b1f94fd472ca..403ed748f3d525bb5eaefe35b44ef26391230fcd 100644
--- a/src/gpu/core/GPU/LBMKernel.cu
+++ b/src/gpu/core/GPU/LBMKernel.cu
@@ -18,111 +18,6 @@
#include "Parameter/Parameter.h"
//////////////////////////////////////////////////////////////////////////
-void Init27(
- int myid,
- int numprocs,
- real u0,
- unsigned int* geoD,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- real* vParab,
- unsigned long long numberOfLBnodes,
- unsigned int grid_nx,
- unsigned int grid_ny,
- unsigned int grid_nz,
- real* DD,
- int level,
- int maxlevel)
-{
- dim3 threads ( grid_nx, 1, 1 );
- dim3 grid ( grid_ny, grid_nz );
-
- LBInit27<<< grid, threads >>> (
- myid,
- numprocs,
- u0,
- geoD,
- neighborX,
- neighborY,
- neighborZ,
- vParab,
- numberOfLBnodes,
- grid_nx,
- grid_ny,
- grid_nz,
- DD,
- level,
- maxlevel);
- getLastCudaError("LBInit27 execution failed");
-}
-//////////////////////////////////////////////////////////////////////////
-void InitNonEqPartSP27(
- unsigned int numberOfThreads,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* neighborWSB,
- unsigned int* geoD,
- real* rho,
- real* ux,
- real* uy,
- real* uz,
- unsigned long long numberOfLBnodes,
- real* DD,
- real omega,
- bool EvenOrOdd)
-{
- vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, numberOfLBnodes);
-
- LBInitNonEqPartSP27<<< grid.grid, grid.threads >>>(
- neighborX,
- neighborY,
- neighborZ,
- neighborWSB,
- geoD,
- rho,
- ux,
- uy,
- uz,
- numberOfLBnodes,
- DD,
- omega,
- EvenOrOdd);
- getLastCudaError("LBInitNonEqPartSP27 execution failed");
-}
-//////////////////////////////////////////////////////////////////////////
-void InitADDev27(
- unsigned int numberOfThreads,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* geoD,
- real* Conc,
- real* ux,
- real* uy,
- real* uz,
- unsigned long long numberOfLBnodes,
- real* DD27,
- bool EvenOrOdd)
-{
- vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, numberOfLBnodes);
-
- InitAD27<<< grid.grid, grid.threads >>>(
- neighborX,
- neighborY,
- neighborZ,
- geoD,
- Conc,
- ux,
- uy,
- uz,
- numberOfLBnodes,
- DD27,
- EvenOrOdd);
- getLastCudaError("InitAD27 execution failed");
-}
-//////////////////////////////////////////////////////////////////////////
void CalcMac27(
real* vxD,
real* vyD,
diff --git a/src/gpu/core/Init/InitLattice.cpp b/src/gpu/core/Init/InitLattice.cpp
index ba4a5cae5259ff43a1d85e6ee97dd125e3da912f..f8e00006e3b6438080c23b3a24bb9d30aa17fe48 100644
--- a/src/gpu/core/Init/InitLattice.cpp
+++ b/src/gpu/core/Init/InitLattice.cpp
@@ -39,7 +39,7 @@
#include "Temperature/FindTemperature.h"
-void initLattice(SPtr<Parameter> para, SPtr<PreProcessor> preProcessor, SPtr<CudaMemoryManager> cudaMemoryManager)
+void initLattice(SPtr<Parameter> para, SPtr<PreProcessor> preProcessor, SPtr<PreProcessor> preProcessorAD, SPtr<CudaMemoryManager> cudaMemoryManager)
{
for (int lev = para->getFine(); lev >= para->getCoarse(); lev--) {
preProcessor->init(para, lev);
@@ -84,7 +84,8 @@ void initLattice(SPtr<Parameter> para, SPtr<PreProcessor> preProcessor, SPtr<Cud
for (size_t index = 0; index < para->getParH(lev)->numberOfNodes; index++) {
para->getParH(lev)->concentration[index] = para->getTemperatureInit();
}
- initTemperatur(para.get(), cudaMemoryManager.get(), lev);
+
+ preProcessorAD->init(para, lev);
}
}
}
diff --git a/src/gpu/core/Init/InitLattice.h b/src/gpu/core/Init/InitLattice.h
index 8de22f321899ba7f8a08f8809fbd904c2d9ec1ec..2e51f911ce93525cf3cbc418ea14fe02aa598cd0 100644
--- a/src/gpu/core/Init/InitLattice.h
+++ b/src/gpu/core/Init/InitLattice.h
@@ -39,6 +39,6 @@ class Parameter;
class PreProcessor;
class CudaMemoryManager;
-void initLattice(SPtr<Parameter> para, SPtr<PreProcessor> preProcessor, SPtr<CudaMemoryManager> cudaMemoryManager);
+void initLattice(SPtr<Parameter> para, SPtr<PreProcessor> preProcessor, SPtr<PreProcessor> preProcessorAD, SPtr<CudaMemoryManager> cudaMemoryManager);
#endif
diff --git a/src/gpu/core/Kernel/Compressible/AdvectionDiffusion/F16/F16CompressibleAdvectionDiffusion.cu b/src/gpu/core/Kernel/Compressible/AdvectionDiffusion/F16/F16CompressibleAdvectionDiffusion.cu
index e8ff6ea34e7ae03f191f3810dc31ba31c6819df2..d9b1027576beb8ceb016318e032b62eec60bebeb 100644
--- a/src/gpu/core/Kernel/Compressible/AdvectionDiffusion/F16/F16CompressibleAdvectionDiffusion.cu
+++ b/src/gpu/core/Kernel/Compressible/AdvectionDiffusion/F16/F16CompressibleAdvectionDiffusion.cu
@@ -32,7 +32,7 @@ F16CompressibleAdvectionDiffusion::F16CompressibleAdvectionDiffusion(std::shared
this->para = para;
this->level = level;
- myPreProcessorTypes.push_back(InitCompAD27);
+ myPreProcessorTypes.push_back(InitAdvectionDiffusionCompressible);
}
diff --git a/src/gpu/core/Kernel/Compressible/NavierStokes/B15/B15CompressibleNavierStokesBGKplus.cu b/src/gpu/core/Kernel/Compressible/NavierStokes/B15/B15CompressibleNavierStokesBGKplus.cu
index 6f6befbd2cca54038ebe7a3b27747fe9b47cec1b..6f51eee72ee8b7db62cf43bb827b6d117cf26ec9 100644
--- a/src/gpu/core/Kernel/Compressible/NavierStokes/B15/B15CompressibleNavierStokesBGKplus.cu
+++ b/src/gpu/core/Kernel/Compressible/NavierStokes/B15/B15CompressibleNavierStokesBGKplus.cu
@@ -30,7 +30,7 @@ B15CompressibleNavierStokesBGKplus::B15CompressibleNavierStokesBGKplus(std::shar
this->para = para;
this->level = level;
- myPreProcessorTypes.push_back(InitCompSP27);
+ myPreProcessorTypes.push_back(InitNavierStokesCompressible);
}
diff --git a/src/gpu/core/Kernel/Compressible/NavierStokes/B92/B92CompressibleNavierStokes.cu b/src/gpu/core/Kernel/Compressible/NavierStokes/B92/B92CompressibleNavierStokes.cu
index 5f2d069459a3cdfc61b79771316f2d81c3039f98..1dba9343c90a213fee5193d417b135e03e855e3c 100644
--- a/src/gpu/core/Kernel/Compressible/NavierStokes/B92/B92CompressibleNavierStokes.cu
+++ b/src/gpu/core/Kernel/Compressible/NavierStokes/B92/B92CompressibleNavierStokes.cu
@@ -30,7 +30,7 @@ B92CompressibleNavierStokes::B92CompressibleNavierStokes(std::shared_ptr<Paramet
this->para = para;
this->level = level;
- myPreProcessorTypes.push_back(InitCompSP27);
+ myPreProcessorTypes.push_back(InitNavierStokesCompressible);
}
diff --git a/src/gpu/core/Kernel/Compressible/NavierStokes/K15/K15CompressibleNavierStokes.cu b/src/gpu/core/Kernel/Compressible/NavierStokes/K15/K15CompressibleNavierStokes.cu
index 2ba1af7a34b219d8f06e44708c354c609efc819f..4cb59734754105e50a9d9ea5d707fa493404f1d5 100644
--- a/src/gpu/core/Kernel/Compressible/NavierStokes/K15/K15CompressibleNavierStokes.cu
+++ b/src/gpu/core/Kernel/Compressible/NavierStokes/K15/K15CompressibleNavierStokes.cu
@@ -48,7 +48,7 @@ K15CompressibleNavierStokes::K15CompressibleNavierStokes(std::shared_ptr<Paramet
this->para = para;
this->level = level;
- myPreProcessorTypes.push_back(InitCompSP27);
+ myPreProcessorTypes.push_back(InitNavierStokesCompressible);
}
\ No newline at end of file
diff --git a/src/gpu/core/Kernel/Compressible/NavierStokes/K17/K17CompressibleNavierStokes.cu b/src/gpu/core/Kernel/Compressible/NavierStokes/K17/K17CompressibleNavierStokes.cu
index 60d4a031451b73e1c4b8de1b317d6488d32e69e1..3fd0f5786e5ed9562f591bd4c264c06fb18e3f4d 100644
--- a/src/gpu/core/Kernel/Compressible/NavierStokes/K17/K17CompressibleNavierStokes.cu
+++ b/src/gpu/core/Kernel/Compressible/NavierStokes/K17/K17CompressibleNavierStokes.cu
@@ -116,7 +116,7 @@ void K17CompressibleNavierStokes<turbulenceModel>::runOnIndices(const unsigned i
template <vf::lbm::TurbulenceModel turbulenceModel>
K17CompressibleNavierStokes<turbulenceModel>::K17CompressibleNavierStokes(std::shared_ptr<Parameter> para, int level) : KernelImp(para, level)
{
- myPreProcessorTypes.push_back(InitCompSP27);
+ myPreProcessorTypes.push_back(InitNavierStokesCompressible);
this->cudaGrid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
this->kernelUsesFluidNodeIndices = true;
diff --git a/src/gpu/core/Kernel/Incompressible/AdvectionDiffusion/F16/F16IncompressibleAdvectionDiffusion.cu b/src/gpu/core/Kernel/Incompressible/AdvectionDiffusion/F16/F16IncompressibleAdvectionDiffusion.cu
index 2e28074650c3c6d8b4ecc668bf9605e3d0396370..9ea0fe2a84c5fe933680acb93ddde7eab6c661fa 100644
--- a/src/gpu/core/Kernel/Incompressible/AdvectionDiffusion/F16/F16IncompressibleAdvectionDiffusion.cu
+++ b/src/gpu/core/Kernel/Incompressible/AdvectionDiffusion/F16/F16IncompressibleAdvectionDiffusion.cu
@@ -32,7 +32,7 @@ F16IncompressibleAdvectionDiffusion::F16IncompressibleAdvectionDiffusion(std::sh
this->para = para;
this->level = level;
- myPreProcessorTypes.push_back(InitIncompAD27);
+ myPreProcessorTypes.push_back(InitAdvectionDiffusionIncompressible);
}
diff --git a/src/gpu/core/Kernel/Incompressible/AdvectionDiffusion/F16/F16IncompressibleAdvectionDiffusion_Device.cu b/src/gpu/core/Kernel/Incompressible/AdvectionDiffusion/F16/F16IncompressibleAdvectionDiffusion_Device.cu
index 3c1dd9524d9fcc02c298483d24dffd924297ef9c..ba2aec9b1d9482cf0faaf6f8412f7272a1c267ac 100644
--- a/src/gpu/core/Kernel/Incompressible/AdvectionDiffusion/F16/F16IncompressibleAdvectionDiffusion_Device.cu
+++ b/src/gpu/core/Kernel/Incompressible/AdvectionDiffusion/F16/F16IncompressibleAdvectionDiffusion_Device.cu
@@ -246,11 +246,6 @@ __global__ void F16IncompressibleAdvectionDiffusion_Device(
//!
////////////////////////////////////////////////////////////////////////////////////
// fluid component
- real drhoFluid =
- ((((fccc + faaa) + (faca + fcac)) + ((facc + fcaa) + (faac + fcca))) +
- (((fbac + fbca) + (fbaa + fbcc)) + ((fabc + fcba) + (faba + fcbc)) + ((facb + fcab) + (faab + fccb))) +
- ((fabb + fcbb) + (fbab + fbcb) + (fbba + fbbc))) + fbbb;
-
real vvx =
((((fccc - faaa) + (fcac - faca)) + ((fcaa - facc) + (fcca - faac))) +
(((fcba - fabc) + (fcbc - faba)) + ((fcab - facb) + (fccb - faab))) +
diff --git a/src/gpu/core/Kernel/Incompressible/NavierStokes/B15/B15IncompressibleNavierStokesBGKplus.cu b/src/gpu/core/Kernel/Incompressible/NavierStokes/B15/B15IncompressibleNavierStokesBGKplus.cu
index 251cf7544eae9710d3460cacfdd4b301c9c48d79..cc0b6d43e4b7bd5fb329a0cbbf1ba7e33f466a75 100644
--- a/src/gpu/core/Kernel/Incompressible/NavierStokes/B15/B15IncompressibleNavierStokesBGKplus.cu
+++ b/src/gpu/core/Kernel/Incompressible/NavierStokes/B15/B15IncompressibleNavierStokesBGKplus.cu
@@ -30,7 +30,7 @@ B15IncompressibleNavierStokesBGKplus::B15IncompressibleNavierStokesBGKplus(std::
this->para = para;
this->level = level;
- myPreProcessorTypes.push_back(InitSP27);
+ myPreProcessorTypes.push_back(InitNavierStokesIncompressible);
}
diff --git a/src/gpu/core/Kernel/Incompressible/NavierStokes/B92/B92IncompressibleNavierStokes.cu b/src/gpu/core/Kernel/Incompressible/NavierStokes/B92/B92IncompressibleNavierStokes.cu
index 4ff8d0d6ce700b0245171fe31790e320bcc9e725..49210b1fe31a0addf91401f3c86161846787b96c 100644
--- a/src/gpu/core/Kernel/Incompressible/NavierStokes/B92/B92IncompressibleNavierStokes.cu
+++ b/src/gpu/core/Kernel/Incompressible/NavierStokes/B92/B92IncompressibleNavierStokes.cu
@@ -30,7 +30,7 @@ B92IncompressibleNavierStokes::B92IncompressibleNavierStokes(std::shared_ptr<Par
this->para = para;
this->level = level;
- myPreProcessorTypes.push_back(InitSP27);
+ myPreProcessorTypes.push_back(InitNavierStokesIncompressible);
}
diff --git a/src/gpu/core/Kernel/Incompressible/NavierStokes/K15/K15IncompressibleNavierStokes.cu b/src/gpu/core/Kernel/Incompressible/NavierStokes/K15/K15IncompressibleNavierStokes.cu
index 40247d0f664ff842246a1a9066a69bcf5dff117f..5d64f3216ba65389a03941b55f61a7091170e0ea 100644
--- a/src/gpu/core/Kernel/Incompressible/NavierStokes/K15/K15IncompressibleNavierStokes.cu
+++ b/src/gpu/core/Kernel/Incompressible/NavierStokes/K15/K15IncompressibleNavierStokes.cu
@@ -30,7 +30,7 @@ K15IncompressibleNavierStokes::K15IncompressibleNavierStokes(std::shared_ptr<Par
this->para = para;
this->level = level;
- myPreProcessorTypes.push_back(InitSP27);
+ myPreProcessorTypes.push_back(InitNavierStokesIncompressible);
}
diff --git a/src/gpu/core/KernelManager/ADKernelManager.cpp b/src/gpu/core/KernelManager/ADKernelManager.cpp
index 0a531bd36eedbd0f90a1ef7f94aa29e5f1dd4ded..7b867dc84acc245633ac65d785f198b48e2979ee 100644
--- a/src/gpu/core/KernelManager/ADKernelManager.cpp
+++ b/src/gpu/core/KernelManager/ADKernelManager.cpp
@@ -38,56 +38,6 @@
ADKernelManager::ADKernelManager(SPtr<Parameter> parameter, std::vector<SPtr<AdvectionDiffusionKernel>>& adkernels): para(parameter), adkernels(adkernels){}
-void ADKernelManager::initAD(const int level) const
-{
- //////////////////////////////////////////////////////////////////////////
- // calculation of omega for diffusivity
- para->getParD(level)->omegaDiffusivity = (real)2.0 / ((real)6.0 * para->getParD(level)->diffusivity + (real)1.0);
- //////////////////////////////////////////////////////////////////////////
- para->getParD(level)->isEvenTimestep = true;
- //////////////////////////////////////////////////////////////////////////
- InitADDev27(
- para->getParD(level)->numberofthreads,
- para->getParD(level)->neighborX,
- para->getParD(level)->neighborY,
- para->getParD(level)->neighborZ,
- para->getParD(level)->typeOfGridNode,
- para->getParD(level)->concentration,
- para->getParD(level)->velocityX,
- para->getParD(level)->velocityY,
- para->getParD(level)->velocityZ,
- para->getParD(level)->numberOfNodes,
- para->getParD(level)->distributionsAD.f[0],
- para->getParD(level)->isEvenTimestep);
- //////////////////////////////////////////////////////////////////////////
- para->getParD(level)->isEvenTimestep = false;
- //////////////////////////////////////////////////////////////////////////
- InitADDev27(
- para->getParD(level)->numberofthreads,
- para->getParD(level)->neighborX,
- para->getParD(level)->neighborY,
- para->getParD(level)->neighborZ,
- para->getParD(level)->typeOfGridNode,
- para->getParD(level)->concentration,
- para->getParD(level)->velocityX,
- para->getParD(level)->velocityY,
- para->getParD(level)->velocityZ,
- para->getParD(level)->numberOfNodes,
- para->getParD(level)->distributionsAD.f[0],
- para->getParD(level)->isEvenTimestep);
- //////////////////////////////////////////////////////////////////////////
- CalcConcentration27(
- para->getParD(level)->numberofthreads,
- para->getParD(level)->concentration,
- para->getParD(level)->typeOfGridNode,
- para->getParD(level)->neighborX,
- para->getParD(level)->neighborY,
- para->getParD(level)->neighborZ,
- para->getParD(level)->numberOfNodes,
- para->getParD(level)->distributionsAD.f[0],
- para->getParD(level)->isEvenTimestep);
-}
-
////////////////////////////////////////////////////////////////////////////////
void ADKernelManager::setInitialNodeValuesAD(const int level, SPtr<CudaMemoryManager> cudaMemoryManager) const
{
diff --git a/src/gpu/core/KernelManager/ADKernelManager.h b/src/gpu/core/KernelManager/ADKernelManager.h
index 46d818b04e3a2f3fc2b53d5deef249900d324faa..86a59d58371377fa3b4468f1f9e8bac94e0184d4 100644
--- a/src/gpu/core/KernelManager/ADKernelManager.h
+++ b/src/gpu/core/KernelManager/ADKernelManager.h
@@ -54,9 +54,6 @@ public:
//! \param parameter shared pointer to instance of class Parameter
ADKernelManager(SPtr<Parameter> parameter, std::vector<SPtr<AdvectionDiffusionKernel>>& adkernels);
- //! \brief initialize the advection diffusion distributions
- void initAD(const int level) const;
-
//! \brief set initial concentration values at all nodes
//! \param cudaMemoryManager instance of class CudaMemoryManager
void setInitialNodeValuesAD(const int level, SPtr<CudaMemoryManager> cudaMemoryManager) const;
diff --git a/src/gpu/core/LBM/Simulation.cpp b/src/gpu/core/LBM/Simulation.cpp
index 2a7787155d46e9305b5d8df0dd2a809ca7b3b15c..3566c26f8849e586bcdadbc6c916f0815b8d54cf 100644
--- a/src/gpu/core/LBM/Simulation.cpp
+++ b/src/gpu/core/LBM/Simulation.cpp
@@ -156,6 +156,8 @@ void Simulation::init(GridProvider &gridProvider, BoundaryConditionFactory *bcFa
if (para->getDiffOn()) {
VF_LOG_INFO("make AD Kernels");
adKernels = kernelFactory->makeAdvDifKernels(para);
+ std::vector<PreProcessorType> preProADTypes = adKernels.at(0)->getPreProcessorTypes();
+ preProcessorAD = preProcessorFactory->makePreProcessor(preProADTypes, para);
}
//////////////////////////////////////////////////////////////////////////
@@ -270,7 +272,7 @@ void Simulation::init(GridProvider &gridProvider, BoundaryConditionFactory *bcFa
// VF_LOG_INFO("done.");
VF_LOG_INFO("init lattice...");
- initLattice(para, preProcessor, cudaMemoryManager);
+ initLattice(para, preProcessor, preProcessorAD, cudaMemoryManager);
VF_LOG_INFO("done");
// VF_LOG_INFO("set geo for Q...\n");
diff --git a/src/gpu/core/LBM/Simulation.h b/src/gpu/core/LBM/Simulation.h
index e3bd7cf7bc5f761727c581c0870e53af2f1750a5..4a9c89ba0aec09439bf8f440b9af2bc02054f167 100644
--- a/src/gpu/core/LBM/Simulation.h
+++ b/src/gpu/core/LBM/Simulation.h
@@ -82,6 +82,7 @@ private:
std::vector < SPtr< Kernel>> kernels;
std::vector < SPtr< AdvectionDiffusionKernel>> adKernels;
std::shared_ptr<PreProcessor> preProcessor;
+ std::shared_ptr<PreProcessor> preProcessorAD;
SPtr<TurbulenceModelFactory> tmFactory;
SPtr<RestartObject> restart_object;
diff --git a/src/gpu/core/PreProcessor/PreProcessor.h b/src/gpu/core/PreProcessor/PreProcessor.h
index 36560b55f3355ccf0d9156ee9509524cf903ab23..6024b5d09fd81a77dfe279066eabb99b78e660b6 100644
--- a/src/gpu/core/PreProcessor/PreProcessor.h
+++ b/src/gpu/core/PreProcessor/PreProcessor.h
@@ -1,3 +1,33 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
#ifndef PREPROCESSOR_H
#define PREPROCESSOR_H
diff --git a/src/gpu/core/PreProcessor/PreProcessorFactory/PreProcessorFactory.h b/src/gpu/core/PreProcessor/PreProcessorFactory/PreProcessorFactory.h
index 1ea65788267f6a5fffe70f64fd962de013638788..a8677c4af15d90cd24cfadb14218b2a1c03e978d 100644
--- a/src/gpu/core/PreProcessor/PreProcessorFactory/PreProcessorFactory.h
+++ b/src/gpu/core/PreProcessor/PreProcessorFactory/PreProcessorFactory.h
@@ -1,3 +1,33 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
#ifndef PREPROCESSOR_FACTORY_H
#define PREPROCESSOR_FACTORY_H
diff --git a/src/gpu/core/PreProcessor/PreProcessorFactory/PreProcessorFactoryImp.cpp b/src/gpu/core/PreProcessor/PreProcessorFactory/PreProcessorFactoryImp.cpp
index e92be775c79b19c7bf5ca0d1840622e10cccd0fd..18a1d4e7023e92cc0ac3be5dc2a57d1de64ac5cc 100644
--- a/src/gpu/core/PreProcessor/PreProcessorFactory/PreProcessorFactoryImp.cpp
+++ b/src/gpu/core/PreProcessor/PreProcessorFactory/PreProcessorFactoryImp.cpp
@@ -1,14 +1,44 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
#include "PreProcessorFactoryImp.h"
#include "PreProcessor/PreProcessorImp.h"
-#include "PreProcessor/PreProcessorStrategy/InitCompAD7/InitCompAD7.h"
-#include "PreProcessor/PreProcessorStrategy/InitCompAD27/InitCompAD27.h"
-#include "PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27.h"
-#include "PreProcessor/PreProcessorStrategy/InitF3/InitF3.h"
-#include "PreProcessor/PreProcessorStrategy/InitIncompAD27/InitIncompAD27.h"
-#include "PreProcessor/PreProcessorStrategy/InitIncompAD7/InitIncompAD7.h"
-#include "PreProcessor/PreProcessorStrategy/InitSP27/InitSP27.h"
+#include "PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7.h"
+#include "PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressible/InitAdvectionDiffusionCompressible.h"
+#include "PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible.h"
+#include "PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.h"
+#include "PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressible/InitAdvectionDiffusionIncompressible.h"
+#include "PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7.h"
+#include "PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible.h"
std::shared_ptr<PreProcessor> PreProcessorFactoryImp::makePreProcessor(std::vector<PreProcessorType> preProcessorTypes, std::shared_ptr<Parameter> para)
@@ -25,26 +55,26 @@ std::shared_ptr<PreProcessorStrategy> PreProcessorFactoryImp::makePreProcessorSt
{
switch (preProcessorType)
{
- case InitSP27:
- return InitSP27::getNewInstance(para);
+ case InitNavierStokesIncompressible:
+ return InitNavierStokesIncompressible::getNewInstance(para);
break;
- case InitCompSP27:
- return InitCompSP27::getNewInstance(para);
+ case InitNavierStokesCompressible:
+ return InitNavierStokesCompressible::getNewInstance(para);
break;
- case InitF3:
- return InitF3::getNewInstance(para);
+ case InitK18K20NavierStokesCompressible:
+ return InitK18K20NavierStokesCompressible::getNewInstance(para);
break;
- case InitIncompAD7:
- return InitIncompAD7::getNewInstance(para);
+ case InitAdvectionDiffusionIncompressibleD3Q7:
+ return InitAdvectionDiffusionIncompressibleD3Q7::getNewInstance(para);
break;
- case InitIncompAD27:
- return InitIncompAD27::getNewInstance(para);
+ case InitAdvectionDiffusionIncompressible:
+ return InitAdvectionDiffusionIncompressible::getNewInstance(para);
break;
- case InitCompAD7:
- return InitCompAD7::getNewInstance(para);
+ case InitAdvectionDiffusionCompressibleD3Q7:
+ return InitAdvectionDiffusionCompressibleD3Q7::getNewInstance(para);
break;
- case InitCompAD27:
- return InitCompAD27::getNewInstance(para);
+ case InitAdvectionDiffusionCompressible:
+ return InitAdvectionDiffusionCompressible::getNewInstance(para);
break;
default:
break;
diff --git a/src/gpu/core/PreProcessor/PreProcessorFactory/PreProcessorFactoryImp.h b/src/gpu/core/PreProcessor/PreProcessorFactory/PreProcessorFactoryImp.h
index 8b78537746461c773f2f170ebf2f97d166e1be4d..373217659b7c641e7e2e1d9c7dace3a7a28ce21f 100644
--- a/src/gpu/core/PreProcessor/PreProcessorFactory/PreProcessorFactoryImp.h
+++ b/src/gpu/core/PreProcessor/PreProcessorFactory/PreProcessorFactoryImp.h
@@ -1,3 +1,33 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
#ifndef PREPROCESSOR_FACTORY_IMP_H
#define PREPROCESSOR_FACTORY_IMP_H
diff --git a/src/gpu/core/PreProcessor/PreProcessorImp.cpp b/src/gpu/core/PreProcessor/PreProcessorImp.cpp
index 94f25bffc2d553375b25dc3eab80eaa00561ca40..0ee8492e52ec7227b5711da1a425cb22bb9abcea 100644
--- a/src/gpu/core/PreProcessor/PreProcessorImp.cpp
+++ b/src/gpu/core/PreProcessor/PreProcessorImp.cpp
@@ -1,3 +1,33 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
#include "PreProcessorImp.h"
#include "PreProcessorStrategy/PreProcessorStrategy.h"
diff --git a/src/gpu/core/PreProcessor/PreProcessorImp.h b/src/gpu/core/PreProcessor/PreProcessorImp.h
index 10233d5bb694d10262e59d1496eed0665e076da7..f939482dbf1abdd7b28c0c10fac7f4fce47d2a11 100644
--- a/src/gpu/core/PreProcessor/PreProcessorImp.h
+++ b/src/gpu/core/PreProcessor/PreProcessorImp.h
@@ -1,3 +1,33 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
#ifndef PREPROCESSOR_IMP_H
#define PREPROCESSOR_IMP_H
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressible/InitAdvectionDiffusionCompressible.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressible/InitAdvectionDiffusionCompressible.cu
new file mode 100644
index 0000000000000000000000000000000000000000..659edc342c566d41477d0c2ad34b1d137de383b1
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressible/InitAdvectionDiffusionCompressible.cu
@@ -0,0 +1,73 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#include "InitAdvectionDiffusionCompressible.h"
+
+#include "InitAdvectionDiffusionCompressible_Device.cuh"
+#include "Parameter/Parameter.h"
+#include <cuda_helper/CudaGrid.h>
+
+std::shared_ptr<PreProcessorStrategy> InitAdvectionDiffusionCompressible::getNewInstance(std::shared_ptr<Parameter> para)
+{
+ return std::shared_ptr<PreProcessorStrategy>(new InitAdvectionDiffusionCompressible(para));
+}
+
+void InitAdvectionDiffusionCompressible::init(int level)
+{
+ vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
+
+ InitAdvectionDiffusionCompressible_Device <<< grid.grid, grid.threads >>>(
+ para->getParD(level)->neighborX,
+ para->getParD(level)->neighborY,
+ para->getParD(level)->neighborZ,
+ para->getParD(level)->typeOfGridNode,
+ para->getParD(level)->concentration,
+ para->getParD(level)->velocityX,
+ para->getParD(level)->velocityY,
+ para->getParD(level)->velocityZ,
+ para->getParD(level)->numberOfNodes,
+ para->getParD(level)->distributionsAD.f[0],
+ para->getParD(level)->isEvenTimestep);
+ getLastCudaError("InitAdvectionDiffusionCompressible_Device execution failed");
+}
+
+bool InitAdvectionDiffusionCompressible::checkParameter()
+{
+ return false;
+}
+
+InitAdvectionDiffusionCompressible::InitAdvectionDiffusionCompressible(std::shared_ptr<Parameter> para)
+{
+ this->para = para;
+}
+
+InitAdvectionDiffusionCompressible::InitAdvectionDiffusionCompressible()
+{
+}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressible/InitAdvectionDiffusionCompressible.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressible/InitAdvectionDiffusionCompressible.h
new file mode 100644
index 0000000000000000000000000000000000000000..868ac3944e4599dd9fe6ade02f6e083daed26bd6
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressible/InitAdvectionDiffusionCompressible.h
@@ -0,0 +1,53 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#ifndef InitAdvectionDiffusionCompressible_H
+#define InitAdvectionDiffusionCompressible_H
+
+#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h"
+
+#include <memory>
+
+class Parameter;
+
+class InitAdvectionDiffusionCompressible : public PreProcessorStrategy
+{
+public:
+ static std::shared_ptr<PreProcessorStrategy> getNewInstance(std::shared_ptr< Parameter> para);
+ void init(int level);
+ bool checkParameter();
+
+private:
+ InitAdvectionDiffusionCompressible();
+ InitAdvectionDiffusionCompressible(std::shared_ptr< Parameter> para);
+ std::shared_ptr< Parameter> para;
+};
+
+#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressible/InitAdvectionDiffusionCompressible_Device.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressible/InitAdvectionDiffusionCompressible_Device.cu
new file mode 100644
index 0000000000000000000000000000000000000000..9d348a7358a347083fb44677b69831858656aeae
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressible/InitAdvectionDiffusionCompressible_Device.cu
@@ -0,0 +1,207 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#include "LBM/LB.h"
+#include "lbm/constants/D3Q27.h"
+#include <basics/constants/NumericConstants.h>
+
+using namespace vf::basics::constant;
+using namespace vf::lbm;
+using namespace vf::lbm::dir;
+
+
+__global__ void InitAdvectionDiffusionCompressible_Device(
+ uint* neighborX,
+ uint* neighborY,
+ uint* neighborZ,
+ uint* typeOfGridNode,
+ real* concentration,
+ real* velocityX,
+ real* velocityY,
+ real* velocityZ,
+ unsigned long long numberOfLBnodes,
+ real* distributionsAD,
+ bool isEvenTimestep)
+{
+ //////////////////////////////////////////////////////////////////////////
+ //! The initialization is executed in the following steps
+ //!
+ ////////////////////////////////////////////////////////////////////////////////
+ //! - Get node index coordinates from threadIdx, blockIdx, blockDim and gridDim.
+ //!
+ const unsigned x = threadIdx.x; // Globaler x-Index
+ const unsigned y = blockIdx.x; // Globaler y-Index
+ const unsigned z = blockIdx.y; // Globaler z-Index
+
+ const unsigned nx = blockDim.x;
+ const unsigned ny = gridDim.x;
+
+ const unsigned k = nx*(ny*z + y) + x;
+
+ //////////////////////////////////////////////////////////////////////////
+ // run for all indices in size_Mat and fluid nodes
+ if ((k < numberOfLBnodes) && (typeOfGridNode[k] == GEO_FLUID))
+ {
+ //////////////////////////////////////////////////////////////////////////
+ //! - Read distributions: style of reading and writing the distributions from/to stored arrays dependent on timestep is based on the esoteric twist algorithm \ref
+ //! <a href="https://doi.org/10.3390/computation5020019"><b>[ M. Geier et al. (2017), DOI:10.3390/computation5020019 ]</b></a>
+ //!
+ Distributions27 distAD;
+ if (isEvenTimestep)
+ {
+ distAD.f[dP00] = &distributionsAD[dP00 * numberOfLBnodes];
+ distAD.f[dM00] = &distributionsAD[dM00 * numberOfLBnodes];
+ distAD.f[d0P0] = &distributionsAD[d0P0 * numberOfLBnodes];
+ distAD.f[d0M0] = &distributionsAD[d0M0 * numberOfLBnodes];
+ distAD.f[d00P] = &distributionsAD[d00P * numberOfLBnodes];
+ distAD.f[d00M] = &distributionsAD[d00M * numberOfLBnodes];
+ distAD.f[dPP0] = &distributionsAD[dPP0 * numberOfLBnodes];
+ distAD.f[dMM0] = &distributionsAD[dMM0 * numberOfLBnodes];
+ distAD.f[dPM0] = &distributionsAD[dPM0 * numberOfLBnodes];
+ distAD.f[dMP0] = &distributionsAD[dMP0 * numberOfLBnodes];
+ distAD.f[dP0P] = &distributionsAD[dP0P * numberOfLBnodes];
+ distAD.f[dM0M] = &distributionsAD[dM0M * numberOfLBnodes];
+ distAD.f[dP0M] = &distributionsAD[dP0M * numberOfLBnodes];
+ distAD.f[dM0P] = &distributionsAD[dM0P * numberOfLBnodes];
+ distAD.f[d0PP] = &distributionsAD[d0PP * numberOfLBnodes];
+ distAD.f[d0MM] = &distributionsAD[d0MM * numberOfLBnodes];
+ distAD.f[d0PM] = &distributionsAD[d0PM * numberOfLBnodes];
+ distAD.f[d0MP] = &distributionsAD[d0MP * numberOfLBnodes];
+ distAD.f[d000] = &distributionsAD[d000 * numberOfLBnodes];
+ distAD.f[dPPP] = &distributionsAD[dPPP * numberOfLBnodes];
+ distAD.f[dMMP] = &distributionsAD[dMMP * numberOfLBnodes];
+ distAD.f[dPMP] = &distributionsAD[dPMP * numberOfLBnodes];
+ distAD.f[dMPP] = &distributionsAD[dMPP * numberOfLBnodes];
+ distAD.f[dPPM] = &distributionsAD[dPPM * numberOfLBnodes];
+ distAD.f[dMMM] = &distributionsAD[dMMM * numberOfLBnodes];
+ distAD.f[dPMM] = &distributionsAD[dPMM * numberOfLBnodes];
+ distAD.f[dMPM] = &distributionsAD[dMPM * numberOfLBnodes];
+ }
+ else
+ {
+ distAD.f[dM00] = &distributionsAD[dP00 * numberOfLBnodes];
+ distAD.f[dP00] = &distributionsAD[dM00 * numberOfLBnodes];
+ distAD.f[d0M0] = &distributionsAD[d0P0 * numberOfLBnodes];
+ distAD.f[d0P0] = &distributionsAD[d0M0 * numberOfLBnodes];
+ distAD.f[d00M] = &distributionsAD[d00P * numberOfLBnodes];
+ distAD.f[d00P] = &distributionsAD[d00M * numberOfLBnodes];
+ distAD.f[dMM0] = &distributionsAD[dPP0 * numberOfLBnodes];
+ distAD.f[dPP0] = &distributionsAD[dMM0 * numberOfLBnodes];
+ distAD.f[dMP0] = &distributionsAD[dPM0 * numberOfLBnodes];
+ distAD.f[dPM0] = &distributionsAD[dMP0 * numberOfLBnodes];
+ distAD.f[dM0M] = &distributionsAD[dP0P * numberOfLBnodes];
+ distAD.f[dP0P] = &distributionsAD[dM0M * numberOfLBnodes];
+ distAD.f[dM0P] = &distributionsAD[dP0M * numberOfLBnodes];
+ distAD.f[dP0M] = &distributionsAD[dM0P * numberOfLBnodes];
+ distAD.f[d0MM] = &distributionsAD[d0PP * numberOfLBnodes];
+ distAD.f[d0PP] = &distributionsAD[d0MM * numberOfLBnodes];
+ distAD.f[d0MP] = &distributionsAD[d0PM * numberOfLBnodes];
+ distAD.f[d0PM] = &distributionsAD[d0MP * numberOfLBnodes];
+ distAD.f[d000] = &distributionsAD[d000 * numberOfLBnodes];
+ distAD.f[dMMM] = &distributionsAD[dPPP * numberOfLBnodes];
+ distAD.f[dPPM] = &distributionsAD[dMMP * numberOfLBnodes];
+ distAD.f[dMPM] = &distributionsAD[dPMP * numberOfLBnodes];
+ distAD.f[dPMM] = &distributionsAD[dMPP * numberOfLBnodes];
+ distAD.f[dMMP] = &distributionsAD[dPPM * numberOfLBnodes];
+ distAD.f[dPPP] = &distributionsAD[dMMM * numberOfLBnodes];
+ distAD.f[dMPP] = &distributionsAD[dPMM * numberOfLBnodes];
+ distAD.f[dPMP] = &distributionsAD[dMPM * numberOfLBnodes];
+ }
+ //////////////////////////////////////////////////////////////////////////
+ //! - Set local velocities and concetration
+ //!
+ real conc = concentration[k];
+ real vx1 = velocityX[k];
+ real vx2 = velocityY[k];
+ real vx3 = velocityZ[k];
+ //////////////////////////////////////////////////////////////////////////
+ //! - Set neighbor indices (necessary for indirect addressing)
+ //!
+ uint kzero = k;
+ uint ke = k;
+ uint kw = neighborX[k];
+ uint kn = k;
+ uint ks = neighborY[k];
+ uint kt = k;
+ uint kb = neighborZ[k];
+ uint ksw = neighborY[kw];
+ uint kne = k;
+ uint kse = ks;
+ uint knw = kw;
+ uint kbw = neighborZ[kw];
+ uint kte = k;
+ uint kbe = kb;
+ uint ktw = kw;
+ uint kbs = neighborZ[ks];
+ uint ktn = k;
+ uint kbn = kb;
+ uint kts = ks;
+ uint ktse = ks;
+ uint kbnw = kbw;
+ uint ktnw = kw;
+ uint kbse = kbs;
+ uint ktsw = ksw;
+ uint kbne = kb;
+ uint ktne = k;
+ uint kbsw = neighborZ[ksw];
+ //////////////////////////////////////////////////////////////////////////
+ //! - Calculate the equilibrium and set the distributions
+ //!
+ real cu_sq = c3o2*(vx1*vx1 + vx2*vx2 + vx3*vx3);
+
+ (distAD.f[d000])[kzero] = c8o27 * conc * (c1o1 - cu_sq);
+ (distAD.f[dP00])[ke ] = c2o27 * conc * (c1o1 + c3o1 * ( vx1 ) + c9o2 * ( vx1 ) * ( vx1 ) - cu_sq);
+ (distAD.f[dM00])[kw ] = c2o27 * conc * (c1o1 + c3o1 * (-vx1 ) + c9o2 * (-vx1 ) * (-vx1 ) - cu_sq);
+ (distAD.f[d0P0])[kn ] = c2o27 * conc * (c1o1 + c3o1 * ( vx2 ) + c9o2 * ( vx2 ) * ( vx2 ) - cu_sq);
+ (distAD.f[d0M0])[ks ] = c2o27 * conc * (c1o1 + c3o1 * ( - vx2 ) + c9o2 * ( - vx2 ) * ( - vx2 ) - cu_sq);
+ (distAD.f[d00P])[kt ] = c2o27 * conc * (c1o1 + c3o1 * ( vx3) + c9o2 * ( vx3) * ( vx3) - cu_sq);
+ (distAD.f[d00M])[kb ] = c2o27 * conc * (c1o1 + c3o1 * ( - vx3) + c9o2 * ( - vx3) * ( - vx3) - cu_sq);
+ (distAD.f[dPP0])[kne ] = c1o54 * conc * (c1o1 + c3o1 * ( vx1 + vx2 ) + c9o2 * ( vx1 + vx2 ) * ( vx1 + vx2 ) - cu_sq);
+ (distAD.f[dMM0])[ksw ] = c1o54 * conc * (c1o1 + c3o1 * (-vx1 - vx2 ) + c9o2 * (-vx1 - vx2 ) * (-vx1 - vx2 ) - cu_sq);
+ (distAD.f[dPM0])[kse ] = c1o54 * conc * (c1o1 + c3o1 * ( vx1 - vx2 ) + c9o2 * ( vx1 - vx2 ) * ( vx1 - vx2 ) - cu_sq);
+ (distAD.f[dMP0])[knw ] = c1o54 * conc * (c1o1 + c3o1 * (-vx1 + vx2 ) + c9o2 * (-vx1 + vx2 ) * (-vx1 + vx2 ) - cu_sq);
+ (distAD.f[dP0P])[kte ] = c1o54 * conc * (c1o1 + c3o1 * ( vx1 + vx3) + c9o2 * ( vx1 + vx3) * ( vx1 + vx3) - cu_sq);
+ (distAD.f[dM0M])[kbw ] = c1o54 * conc * (c1o1 + c3o1 * (-vx1 - vx3) + c9o2 * (-vx1 - vx3) * (-vx1 - vx3) - cu_sq);
+ (distAD.f[dP0M])[kbe ] = c1o54 * conc * (c1o1 + c3o1 * ( vx1 - vx3) + c9o2 * ( vx1 - vx3) * ( vx1 - vx3) - cu_sq);
+ (distAD.f[dM0P])[ktw ] = c1o54 * conc * (c1o1 + c3o1 * (-vx1 + vx3) + c9o2 * (-vx1 + vx3) * (-vx1 + vx3) - cu_sq);
+ (distAD.f[d0PP])[ktn ] = c1o54 * conc * (c1o1 + c3o1 * ( vx2 + vx3) + c9o2 * ( vx2 + vx3) * ( vx2 + vx3) - cu_sq);
+ (distAD.f[d0MM])[kbs ] = c1o54 * conc * (c1o1 + c3o1 * ( - vx2 - vx3) + c9o2 * ( - vx2 - vx3) * ( - vx2 - vx3) - cu_sq);
+ (distAD.f[d0PM])[kbn ] = c1o54 * conc * (c1o1 + c3o1 * ( vx2 - vx3) + c9o2 * ( vx2 - vx3) * ( vx2 - vx3) - cu_sq);
+ (distAD.f[d0MP])[kts ] = c1o54 * conc * (c1o1 + c3o1 * ( - vx2 + vx3) + c9o2 * ( - vx2 + vx3) * ( - vx2 + vx3) - cu_sq);
+ (distAD.f[dPPP])[ktne ] = c1o216 * conc * (c1o1 + c3o1 * ( vx1 + vx2 + vx3) + c9o2 * ( vx1 + vx2 + vx3) * ( vx1 + vx2 + vx3) - cu_sq);
+ (distAD.f[dMMM])[kbsw ] = c1o216 * conc * (c1o1 + c3o1 * (-vx1 - vx2 - vx3) + c9o2 * (-vx1 - vx2 - vx3) * (-vx1 - vx2 - vx3) - cu_sq);
+ (distAD.f[dPPM])[kbne ] = c1o216 * conc * (c1o1 + c3o1 * ( vx1 + vx2 - vx3) + c9o2 * ( vx1 + vx2 - vx3) * ( vx1 + vx2 - vx3) - cu_sq);
+ (distAD.f[dMMP])[ktsw ] = c1o216 * conc * (c1o1 + c3o1 * (-vx1 - vx2 + vx3) + c9o2 * (-vx1 - vx2 + vx3) * (-vx1 - vx2 + vx3) - cu_sq);
+ (distAD.f[dPMP])[ktse ] = c1o216 * conc * (c1o1 + c3o1 * ( vx1 - vx2 + vx3) + c9o2 * ( vx1 - vx2 + vx3) * ( vx1 - vx2 + vx3) - cu_sq);
+ (distAD.f[dMPM])[kbnw ] = c1o216 * conc * (c1o1 + c3o1 * (-vx1 + vx2 - vx3) + c9o2 * (-vx1 + vx2 - vx3) * (-vx1 + vx2 - vx3) - cu_sq);
+ (distAD.f[dPMM])[kbse ] = c1o216 * conc * (c1o1 + c3o1 * ( vx1 - vx2 - vx3) + c9o2 * ( vx1 - vx2 - vx3) * ( vx1 - vx2 - vx3) - cu_sq);
+ (distAD.f[dMPP])[ktnw ] = c1o216 * conc * (c1o1 + c3o1 * (-vx1 + vx2 + vx3) + c9o2 * (-vx1 + vx2 + vx3) * (-vx1 + vx2 + vx3) - cu_sq);
+ }
+}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressible/InitAdvectionDiffusionCompressible_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressible/InitAdvectionDiffusionCompressible_Device.cuh
new file mode 100644
index 0000000000000000000000000000000000000000..eab37f6ffd49a8333e503190ad1899e6fb882987
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressible/InitAdvectionDiffusionCompressible_Device.cuh
@@ -0,0 +1,50 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#ifndef InitAdvectionDiffusionCompressible_Device_H
+#define InitAdvectionDiffusionCompressible_Device_H
+
+#include <DataTypes.h>
+#include <curand.h>
+
+__global__ void InitAdvectionDiffusionCompressible_Device(
+ uint* neighborX,
+ uint* neighborY,
+ uint* neighborZ,
+ uint* typeOfGridNode,
+ real* concentration,
+ real* velocityX,
+ real* velocityY,
+ real* velocityZ,
+ unsigned long long numberOfLBnodes,
+ real* distributionsAD,
+ bool isEvenTimestep);
+
+#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7.cu
new file mode 100644
index 0000000000000000000000000000000000000000..2db77de45d67a0f74d80c50b6d875aa39d7b2d6d
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7.cu
@@ -0,0 +1,73 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#include "InitAdvectionDiffusionCompressibleD3Q7.h"
+
+#include "InitAdvectionDiffusionCompressibleD3Q7_Device.cuh"
+#include "Parameter/Parameter.h"
+#include <cuda_helper/CudaGrid.h>
+
+std::shared_ptr<InitAdvectionDiffusionCompressibleD3Q7> InitAdvectionDiffusionCompressibleD3Q7::getNewInstance(std::shared_ptr<Parameter> para)
+{
+ return std::shared_ptr<InitAdvectionDiffusionCompressibleD3Q7>(new InitAdvectionDiffusionCompressibleD3Q7(para));
+}
+
+void InitAdvectionDiffusionCompressibleD3Q7::init(int level)
+{
+ vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
+
+ InitAdvectionDiffusionCompressibleD3Q7_Device <<< grid.grid, grid.threads >>>(
+ para->getParD(level)->neighborX,
+ para->getParD(level)->neighborY,
+ para->getParD(level)->neighborZ,
+ para->getParD(level)->typeOfGridNode,
+ para->getParD(level)->concentration,
+ para->getParD(level)->velocityX,
+ para->getParD(level)->velocityY,
+ para->getParD(level)->velocityZ,
+ para->getParD(level)->numberOfNodes,
+ para->getParD(level)->distributionsAD7.f[0],
+ para->getParD(level)->isEvenTimestep);
+ getLastCudaError("InitAdvectionDiffusionCompressibleD3Q7_Device execution failed");
+}
+
+bool InitAdvectionDiffusionCompressibleD3Q7::checkParameter()
+{
+ return false;
+}
+
+InitAdvectionDiffusionCompressibleD3Q7::InitAdvectionDiffusionCompressibleD3Q7(std::shared_ptr<Parameter> para)
+{
+ this->para = para;
+}
+
+InitAdvectionDiffusionCompressibleD3Q7::InitAdvectionDiffusionCompressibleD3Q7()
+{
+}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7.h
new file mode 100644
index 0000000000000000000000000000000000000000..2c68fca2fde9d70075ceb1dc9f0e0f238e39fb40
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7.h
@@ -0,0 +1,52 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#ifndef InitAdvectionDiffusionCompressibleD3Q7_H
+#define InitAdvectionDiffusionCompressibleD3Q7_H
+
+#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h"
+
+#include <memory>
+
+class Parameter;
+
+class InitAdvectionDiffusionCompressibleD3Q7 : public PreProcessorStrategy
+{
+public:
+ static std::shared_ptr<InitAdvectionDiffusionCompressibleD3Q7> getNewInstance(std::shared_ptr< Parameter> para);
+ void init(int level);
+ bool checkParameter();
+
+private:
+ InitAdvectionDiffusionCompressibleD3Q7();
+ InitAdvectionDiffusionCompressibleD3Q7(std::shared_ptr< Parameter> para);
+ std::shared_ptr< Parameter> para;
+};
+#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD7/InitCompAD7_Device.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7_Device.cu
similarity index 61%
rename from src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD7/InitCompAD7_Device.cu
rename to src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7_Device.cu
index 537cc4463d88cfcf919d793fc58918392fe54e46..b9265e83bb1202a1126a1f2bb342bae3f10e842a 100644
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD7/InitCompAD7_Device.cu
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7_Device.cu
@@ -1,3 +1,33 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
#include "LBM/LB.h"
#include "lbm/constants/D3Q27.h"
#include <basics/constants/NumericConstants.h>
@@ -6,7 +36,7 @@ using namespace vf::basics::constant;
using namespace vf::lbm::dir;
#include "math.h"
-__global__ void LB_Init_Comp_AD_7(unsigned int* neighborX,
+__global__ void InitAdvectionDiffusionCompressibleD3Q7_Device(unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned int* geoD,
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7_Device.cuh
new file mode 100644
index 0000000000000000000000000000000000000000..fca4430fe935cc082692d2da837d7e7ea3fe6bec
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7_Device.cuh
@@ -0,0 +1,50 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#ifndef InitAdvectionDiffusionCompressibleD3Q7_Device_H
+#define InitAdvectionDiffusionCompressibleD3Q7_Device_H
+
+#include <DataTypes.h>
+#include <curand.h>
+
+__global__ void InitAdvectionDiffusionCompressibleD3Q7_Device(
+ unsigned int* neighborX,
+ unsigned int* neighborY,
+ unsigned int* neighborZ,
+ unsigned int* geoD,
+ real* Conc,
+ real* ux,
+ real* uy,
+ real* uz,
+ unsigned int size_Mat,
+ real* DD7,
+ bool EvenOrOdd);
+
+#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressible/InitAdvectionDiffusionIncompressible.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressible/InitAdvectionDiffusionIncompressible.cu
new file mode 100644
index 0000000000000000000000000000000000000000..2120c1b4191075f654cee5d85adf43d92137e541
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressible/InitAdvectionDiffusionIncompressible.cu
@@ -0,0 +1,73 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#include "InitAdvectionDiffusionIncompressible.h"
+
+#include "InitAdvectionDiffusionIncompressible_Device.cuh"
+#include "Parameter/Parameter.h"
+#include <cuda_helper/CudaGrid.h>
+
+std::shared_ptr<PreProcessorStrategy> InitAdvectionDiffusionIncompressible::getNewInstance(std::shared_ptr<Parameter> para)
+{
+ return std::shared_ptr<PreProcessorStrategy>(new InitAdvectionDiffusionIncompressible(para));
+}
+
+void InitAdvectionDiffusionIncompressible::init(int level)
+{
+ vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
+
+ InitAdvectionDiffusionIncompressible_Device <<< grid.grid, grid.threads >>>(
+ para->getParD(level)->neighborX,
+ para->getParD(level)->neighborY,
+ para->getParD(level)->neighborZ,
+ para->getParD(level)->typeOfGridNode,
+ para->getParD(level)->concentration,
+ para->getParD(level)->velocityX,
+ para->getParD(level)->velocityY,
+ para->getParD(level)->velocityZ,
+ para->getParD(level)->numberOfNodes,
+ para->getParD(level)->distributionsAD.f[0],
+ para->getParD(level)->isEvenTimestep);
+ getLastCudaError("InitAdvectionDiffusionIncompressible_Device execution failed");
+}
+
+bool InitAdvectionDiffusionIncompressible::checkParameter()
+{
+ return false;
+}
+
+InitAdvectionDiffusionIncompressible::InitAdvectionDiffusionIncompressible(std::shared_ptr<Parameter> para)
+{
+ this->para = para;
+}
+
+InitAdvectionDiffusionIncompressible::InitAdvectionDiffusionIncompressible()
+{
+}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressible/InitAdvectionDiffusionIncompressible.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressible/InitAdvectionDiffusionIncompressible.h
new file mode 100644
index 0000000000000000000000000000000000000000..c60512e644514a68f3542b69c823618c6f148bf4
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressible/InitAdvectionDiffusionIncompressible.h
@@ -0,0 +1,53 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#ifndef InitAdvectionDiffusionIncompressible_H
+#define InitAdvectionDiffusionIncompressible_H
+
+#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h"
+
+#include <memory>
+
+class Parameter;
+
+class InitAdvectionDiffusionIncompressible : public PreProcessorStrategy
+{
+public:
+ static std::shared_ptr<PreProcessorStrategy> getNewInstance(std::shared_ptr< Parameter> para);
+ void init(int level);
+ bool checkParameter();
+
+private:
+ InitAdvectionDiffusionIncompressible();
+ InitAdvectionDiffusionIncompressible(std::shared_ptr< Parameter> para);
+ std::shared_ptr< Parameter> para;
+};
+
+#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressible/InitAdvectionDiffusionIncompressible_Device.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressible/InitAdvectionDiffusionIncompressible_Device.cu
new file mode 100644
index 0000000000000000000000000000000000000000..52586d9b6ce34e6c90d196529174d2190f0b1caa
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressible/InitAdvectionDiffusionIncompressible_Device.cu
@@ -0,0 +1,206 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#include "LBM/LB.h"
+#include "lbm/constants/D3Q27.h"
+#include <basics/constants/NumericConstants.h>
+
+using namespace vf::basics::constant;
+using namespace vf::lbm::dir;
+#include "math.h"
+
+__global__ void InitAdvectionDiffusionIncompressible_Device(
+ uint* neighborX,
+ uint* neighborY,
+ uint* neighborZ,
+ uint* typeOfGridNode,
+ real* concentration,
+ real* velocityX,
+ real* velocityY,
+ real* velocityZ,
+ unsigned long long numberOfLBnodes,
+ real* distributionsAD,
+ bool isEvenTimestep)
+{
+ //////////////////////////////////////////////////////////////////////////
+ //! The initialization is executed in the following steps
+ //!
+ ////////////////////////////////////////////////////////////////////////////////
+ //! - Get node index coordinates from threadIdx, blockIdx, blockDim and gridDim.
+ //!
+ const unsigned x = threadIdx.x; // Globaler x-Index
+ const unsigned y = blockIdx.x; // Globaler y-Index
+ const unsigned z = blockIdx.y; // Globaler z-Index
+
+ const unsigned nx = blockDim.x;
+ const unsigned ny = gridDim.x;
+
+ const unsigned k = nx*(ny*z + y) + x;
+
+ //////////////////////////////////////////////////////////////////////////
+ // run for all indices in size_Mat and fluid nodes
+ if ((k < numberOfLBnodes) && (typeOfGridNode[k] == GEO_FLUID))
+ {
+ //////////////////////////////////////////////////////////////////////////
+ //! - Read distributions: style of reading and writing the distributions from/to stored arrays dependent on timestep is based on the esoteric twist algorithm \ref
+ //! <a href="https://doi.org/10.3390/computation5020019"><b>[ M. Geier et al. (2017), DOI:10.3390/computation5020019 ]</b></a>
+ //!
+ Distributions27 distAD;
+ if (isEvenTimestep)
+ {
+ distAD.f[dP00] = &distributionsAD[dP00 * numberOfLBnodes];
+ distAD.f[dM00] = &distributionsAD[dM00 * numberOfLBnodes];
+ distAD.f[d0P0] = &distributionsAD[d0P0 * numberOfLBnodes];
+ distAD.f[d0M0] = &distributionsAD[d0M0 * numberOfLBnodes];
+ distAD.f[d00P] = &distributionsAD[d00P * numberOfLBnodes];
+ distAD.f[d00M] = &distributionsAD[d00M * numberOfLBnodes];
+ distAD.f[dPP0] = &distributionsAD[dPP0 * numberOfLBnodes];
+ distAD.f[dMM0] = &distributionsAD[dMM0 * numberOfLBnodes];
+ distAD.f[dPM0] = &distributionsAD[dPM0 * numberOfLBnodes];
+ distAD.f[dMP0] = &distributionsAD[dMP0 * numberOfLBnodes];
+ distAD.f[dP0P] = &distributionsAD[dP0P * numberOfLBnodes];
+ distAD.f[dM0M] = &distributionsAD[dM0M * numberOfLBnodes];
+ distAD.f[dP0M] = &distributionsAD[dP0M * numberOfLBnodes];
+ distAD.f[dM0P] = &distributionsAD[dM0P * numberOfLBnodes];
+ distAD.f[d0PP] = &distributionsAD[d0PP * numberOfLBnodes];
+ distAD.f[d0MM] = &distributionsAD[d0MM * numberOfLBnodes];
+ distAD.f[d0PM] = &distributionsAD[d0PM * numberOfLBnodes];
+ distAD.f[d0MP] = &distributionsAD[d0MP * numberOfLBnodes];
+ distAD.f[d000] = &distributionsAD[d000 * numberOfLBnodes];
+ distAD.f[dPPP] = &distributionsAD[dPPP * numberOfLBnodes];
+ distAD.f[dMMP] = &distributionsAD[dMMP * numberOfLBnodes];
+ distAD.f[dPMP] = &distributionsAD[dPMP * numberOfLBnodes];
+ distAD.f[dMPP] = &distributionsAD[dMPP * numberOfLBnodes];
+ distAD.f[dPPM] = &distributionsAD[dPPM * numberOfLBnodes];
+ distAD.f[dMMM] = &distributionsAD[dMMM * numberOfLBnodes];
+ distAD.f[dPMM] = &distributionsAD[dPMM * numberOfLBnodes];
+ distAD.f[dMPM] = &distributionsAD[dMPM * numberOfLBnodes];
+ }
+ else
+ {
+ distAD.f[dM00] = &distributionsAD[dP00 * numberOfLBnodes];
+ distAD.f[dP00] = &distributionsAD[dM00 * numberOfLBnodes];
+ distAD.f[d0M0] = &distributionsAD[d0P0 * numberOfLBnodes];
+ distAD.f[d0P0] = &distributionsAD[d0M0 * numberOfLBnodes];
+ distAD.f[d00M] = &distributionsAD[d00P * numberOfLBnodes];
+ distAD.f[d00P] = &distributionsAD[d00M * numberOfLBnodes];
+ distAD.f[dMM0] = &distributionsAD[dPP0 * numberOfLBnodes];
+ distAD.f[dPP0] = &distributionsAD[dMM0 * numberOfLBnodes];
+ distAD.f[dMP0] = &distributionsAD[dPM0 * numberOfLBnodes];
+ distAD.f[dPM0] = &distributionsAD[dMP0 * numberOfLBnodes];
+ distAD.f[dM0M] = &distributionsAD[dP0P * numberOfLBnodes];
+ distAD.f[dP0P] = &distributionsAD[dM0M * numberOfLBnodes];
+ distAD.f[dM0P] = &distributionsAD[dP0M * numberOfLBnodes];
+ distAD.f[dP0M] = &distributionsAD[dM0P * numberOfLBnodes];
+ distAD.f[d0MM] = &distributionsAD[d0PP * numberOfLBnodes];
+ distAD.f[d0PP] = &distributionsAD[d0MM * numberOfLBnodes];
+ distAD.f[d0MP] = &distributionsAD[d0PM * numberOfLBnodes];
+ distAD.f[d0PM] = &distributionsAD[d0MP * numberOfLBnodes];
+ distAD.f[d000] = &distributionsAD[d000 * numberOfLBnodes];
+ distAD.f[dMMM] = &distributionsAD[dPPP * numberOfLBnodes];
+ distAD.f[dPPM] = &distributionsAD[dMMP * numberOfLBnodes];
+ distAD.f[dMPM] = &distributionsAD[dPMP * numberOfLBnodes];
+ distAD.f[dPMM] = &distributionsAD[dMPP * numberOfLBnodes];
+ distAD.f[dMMP] = &distributionsAD[dPPM * numberOfLBnodes];
+ distAD.f[dPPP] = &distributionsAD[dMMM * numberOfLBnodes];
+ distAD.f[dMPP] = &distributionsAD[dPMM * numberOfLBnodes];
+ distAD.f[dPMP] = &distributionsAD[dMPM * numberOfLBnodes];
+ }
+ //////////////////////////////////////////////////////////////////////////
+ //! - Set local velocities and concetration
+ //!
+ real conc = concentration[k];
+ real vx1 = velocityX[k];
+ real vx2 = velocityY[k];
+ real vx3 = velocityZ[k];
+ //////////////////////////////////////////////////////////////////////////
+ //! - Set neighbor indices (necessary for indirect addressing)
+ //!
+ uint kzero = k;
+ uint ke = k;
+ uint kw = neighborX[k];
+ uint kn = k;
+ uint ks = neighborY[k];
+ uint kt = k;
+ uint kb = neighborZ[k];
+ uint ksw = neighborY[kw];
+ uint kne = k;
+ uint kse = ks;
+ uint knw = kw;
+ uint kbw = neighborZ[kw];
+ uint kte = k;
+ uint kbe = kb;
+ uint ktw = kw;
+ uint kbs = neighborZ[ks];
+ uint ktn = k;
+ uint kbn = kb;
+ uint kts = ks;
+ uint ktse = ks;
+ uint kbnw = kbw;
+ uint ktnw = kw;
+ uint kbse = kbs;
+ uint ktsw = ksw;
+ uint kbne = kb;
+ uint ktne = k;
+ uint kbsw = neighborZ[ksw];
+ //////////////////////////////////////////////////////////////////////////
+ //! - Calculate the equilibrium and set the distributions
+ //!
+ real cu_sq = c3o2*(vx1*vx1 + vx2*vx2 + vx3*vx3);
+
+ (distAD.f[d000])[kzero] = c8o27 * conc * (c1o1 - cu_sq);
+ (distAD.f[dP00])[ke ] = c2o27 * conc * (c1o1 + c3o1 * ( vx1 ) + c9o2 * ( vx1 ) * ( vx1 ) - cu_sq);
+ (distAD.f[dM00])[kw ] = c2o27 * conc * (c1o1 + c3o1 * (-vx1 ) + c9o2 * (-vx1 ) * (-vx1 ) - cu_sq);
+ (distAD.f[d0P0])[kn ] = c2o27 * conc * (c1o1 + c3o1 * ( vx2 ) + c9o2 * ( vx2 ) * ( vx2 ) - cu_sq);
+ (distAD.f[d0M0])[ks ] = c2o27 * conc * (c1o1 + c3o1 * ( - vx2 ) + c9o2 * ( - vx2 ) * ( - vx2 ) - cu_sq);
+ (distAD.f[d00P])[kt ] = c2o27 * conc * (c1o1 + c3o1 * ( vx3) + c9o2 * ( vx3) * ( vx3) - cu_sq);
+ (distAD.f[d00M])[kb ] = c2o27 * conc * (c1o1 + c3o1 * ( - vx3) + c9o2 * ( - vx3) * ( - vx3) - cu_sq);
+ (distAD.f[dPP0])[kne ] = c1o54 * conc * (c1o1 + c3o1 * ( vx1 + vx2 ) + c9o2 * ( vx1 + vx2 ) * ( vx1 + vx2 ) - cu_sq);
+ (distAD.f[dMM0])[ksw ] = c1o54 * conc * (c1o1 + c3o1 * (-vx1 - vx2 ) + c9o2 * (-vx1 - vx2 ) * (-vx1 - vx2 ) - cu_sq);
+ (distAD.f[dPM0])[kse ] = c1o54 * conc * (c1o1 + c3o1 * ( vx1 - vx2 ) + c9o2 * ( vx1 - vx2 ) * ( vx1 - vx2 ) - cu_sq);
+ (distAD.f[dMP0])[knw ] = c1o54 * conc * (c1o1 + c3o1 * (-vx1 + vx2 ) + c9o2 * (-vx1 + vx2 ) * (-vx1 + vx2 ) - cu_sq);
+ (distAD.f[dP0P])[kte ] = c1o54 * conc * (c1o1 + c3o1 * ( vx1 + vx3) + c9o2 * ( vx1 + vx3) * ( vx1 + vx3) - cu_sq);
+ (distAD.f[dM0M])[kbw ] = c1o54 * conc * (c1o1 + c3o1 * (-vx1 - vx3) + c9o2 * (-vx1 - vx3) * (-vx1 - vx3) - cu_sq);
+ (distAD.f[dP0M])[kbe ] = c1o54 * conc * (c1o1 + c3o1 * ( vx1 - vx3) + c9o2 * ( vx1 - vx3) * ( vx1 - vx3) - cu_sq);
+ (distAD.f[dM0P])[ktw ] = c1o54 * conc * (c1o1 + c3o1 * (-vx1 + vx3) + c9o2 * (-vx1 + vx3) * (-vx1 + vx3) - cu_sq);
+ (distAD.f[d0PP])[ktn ] = c1o54 * conc * (c1o1 + c3o1 * ( vx2 + vx3) + c9o2 * ( vx2 + vx3) * ( vx2 + vx3) - cu_sq);
+ (distAD.f[d0MM])[kbs ] = c1o54 * conc * (c1o1 + c3o1 * ( - vx2 - vx3) + c9o2 * ( - vx2 - vx3) * ( - vx2 - vx3) - cu_sq);
+ (distAD.f[d0PM])[kbn ] = c1o54 * conc * (c1o1 + c3o1 * ( vx2 - vx3) + c9o2 * ( vx2 - vx3) * ( vx2 - vx3) - cu_sq);
+ (distAD.f[d0MP])[kts ] = c1o54 * conc * (c1o1 + c3o1 * ( - vx2 + vx3) + c9o2 * ( - vx2 + vx3) * ( - vx2 + vx3) - cu_sq);
+ (distAD.f[dPPP])[ktne ] = c1o216 * conc * (c1o1 + c3o1 * ( vx1 + vx2 + vx3) + c9o2 * ( vx1 + vx2 + vx3) * ( vx1 + vx2 + vx3) - cu_sq);
+ (distAD.f[dMMM])[kbsw ] = c1o216 * conc * (c1o1 + c3o1 * (-vx1 - vx2 - vx3) + c9o2 * (-vx1 - vx2 - vx3) * (-vx1 - vx2 - vx3) - cu_sq);
+ (distAD.f[dPPM])[kbne ] = c1o216 * conc * (c1o1 + c3o1 * ( vx1 + vx2 - vx3) + c9o2 * ( vx1 + vx2 - vx3) * ( vx1 + vx2 - vx3) - cu_sq);
+ (distAD.f[dMMP])[ktsw ] = c1o216 * conc * (c1o1 + c3o1 * (-vx1 - vx2 + vx3) + c9o2 * (-vx1 - vx2 + vx3) * (-vx1 - vx2 + vx3) - cu_sq);
+ (distAD.f[dPMP])[ktse ] = c1o216 * conc * (c1o1 + c3o1 * ( vx1 - vx2 + vx3) + c9o2 * ( vx1 - vx2 + vx3) * ( vx1 - vx2 + vx3) - cu_sq);
+ (distAD.f[dMPM])[kbnw ] = c1o216 * conc * (c1o1 + c3o1 * (-vx1 + vx2 - vx3) + c9o2 * (-vx1 + vx2 - vx3) * (-vx1 + vx2 - vx3) - cu_sq);
+ (distAD.f[dPMM])[kbse ] = c1o216 * conc * (c1o1 + c3o1 * ( vx1 - vx2 - vx3) + c9o2 * ( vx1 - vx2 - vx3) * ( vx1 - vx2 - vx3) - cu_sq);
+ (distAD.f[dMPP])[ktnw ] = c1o216 * conc * (c1o1 + c3o1 * (-vx1 + vx2 + vx3) + c9o2 * (-vx1 + vx2 + vx3) * (-vx1 + vx2 + vx3) - cu_sq);
+ }
+}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressible/InitAdvectionDiffusionIncompressible_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressible/InitAdvectionDiffusionIncompressible_Device.cuh
new file mode 100644
index 0000000000000000000000000000000000000000..9623ca788fecb15509d79eb643d9debd11eb9229
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressible/InitAdvectionDiffusionIncompressible_Device.cuh
@@ -0,0 +1,50 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#ifndef InitAdvectionDiffusionIncompressible_Device_H
+#define InitAdvectionDiffusionIncompressible_Device_H
+
+#include <DataTypes.h>
+#include <curand.h>
+
+__global__ void InitAdvectionDiffusionIncompressible_Device(
+ uint* neighborX,
+ uint* neighborY,
+ uint* neighborZ,
+ uint* typeOfGridNode,
+ real* concentration,
+ real* velocityX,
+ real* velocityY,
+ real* velocityZ,
+ unsigned long long numberOfLBnodes,
+ real* distributionsAD,
+ bool isEvenTimestep);
+
+#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7.cu
new file mode 100644
index 0000000000000000000000000000000000000000..94d7da3bcdc6296bfaaa03c5ce81508c5c3c53e2
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7.cu
@@ -0,0 +1,73 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#include "InitAdvectionDiffusionIncompressibleD3Q7.h"
+
+#include "InitAdvectionDiffusionIncompressibleD3Q7_Device.cuh"
+#include "Parameter/Parameter.h"
+#include <cuda_helper/CudaGrid.h>
+
+std::shared_ptr<PreProcessorStrategy> InitAdvectionDiffusionIncompressibleD3Q7::getNewInstance(std::shared_ptr<Parameter> para)
+{
+ return std::shared_ptr<PreProcessorStrategy>(new InitAdvectionDiffusionIncompressibleD3Q7(para));
+}
+
+void InitAdvectionDiffusionIncompressibleD3Q7::init(int level)
+{
+ vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
+
+ InitAdvectionDiffusionIncompressibleD3Q7_Device <<< grid.grid, grid.threads >>>(
+ para->getParD(level)->neighborX,
+ para->getParD(level)->neighborY,
+ para->getParD(level)->neighborZ,
+ para->getParD(level)->typeOfGridNode,
+ para->getParD(level)->concentration,
+ para->getParD(level)->velocityX,
+ para->getParD(level)->velocityY,
+ para->getParD(level)->velocityZ,
+ para->getParD(level)->numberOfNodes,
+ para->getParD(level)->distributionsAD.f[0],
+ para->getParD(level)->isEvenTimestep);
+ getLastCudaError("InitAdvectionDiffusionIncompressibleD3Q7_Device execution failed");
+}
+
+bool InitAdvectionDiffusionIncompressibleD3Q7::checkParameter()
+{
+ return false;
+}
+
+InitAdvectionDiffusionIncompressibleD3Q7::InitAdvectionDiffusionIncompressibleD3Q7(std::shared_ptr<Parameter> para)
+{
+ this->para = para;
+}
+
+InitAdvectionDiffusionIncompressibleD3Q7::InitAdvectionDiffusionIncompressibleD3Q7()
+{
+}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7.h
new file mode 100644
index 0000000000000000000000000000000000000000..4ece3ee7dd8c889ec373bcc08a5e47c255b7be5b
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7.h
@@ -0,0 +1,53 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#ifndef InitAdvectionDiffusionIncompressibleD3Q7_H
+#define InitAdvectionDiffusionIncompressibleD3Q7_H
+
+#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h"
+
+#include <memory>
+
+class Parameter;
+
+class InitAdvectionDiffusionIncompressibleD3Q7 : public PreProcessorStrategy
+{
+public:
+ static std::shared_ptr<PreProcessorStrategy> getNewInstance(std::shared_ptr< Parameter> para);
+ void init(int level);
+ bool checkParameter();
+
+private:
+ InitAdvectionDiffusionIncompressibleD3Q7();
+ InitAdvectionDiffusionIncompressibleD3Q7(std::shared_ptr< Parameter> para);
+ std::shared_ptr< Parameter> para;
+};
+
+#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD7/InitIncompAD7_Device.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7_Device.cu
similarity index 60%
rename from src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD7/InitIncompAD7_Device.cu
rename to src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7_Device.cu
index a5898729aa1ee73991a77fbe1cdbd407813c0671..441e8c7adcc18d891615d0374fe5bfaa56d92652 100644
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD7/InitIncompAD7_Device.cu
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7_Device.cu
@@ -1,3 +1,33 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
#include "LBM/LB.h"
#include "lbm/constants/D3Q27.h"
#include <basics/constants/NumericConstants.h>
@@ -6,7 +36,8 @@ using namespace vf::basics::constant;
using namespace vf::lbm::dir;
#include "math.h"
-__global__ void LB_Init_Incomp_AD_7(unsigned int* neighborX,
+__global__ void InitAdvectionDiffusionIncompressibleD3Q7_Device(
+ unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned int* geoD,
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7_Device.cuh
new file mode 100644
index 0000000000000000000000000000000000000000..3d0efb3a80edc645020226ca7fe88256a6c5097e
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7_Device.cuh
@@ -0,0 +1,50 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#ifndef InitAdvectionDiffusionIncompressibleD3Q7_Device_H
+#define InitAdvectionDiffusionIncompressibleD3Q7_Device_H
+
+#include <DataTypes.h>
+#include <curand.h>
+
+__global__ void InitAdvectionDiffusionIncompressibleD3Q7_Device(
+ unsigned int* neighborX,
+ unsigned int* neighborY,
+ unsigned int* neighborZ,
+ unsigned int* geoD,
+ real* Conc,
+ real* ux,
+ real* uy,
+ real* uz,
+ unsigned int size_Mat,
+ real* DD7,
+ bool EvenOrOdd);
+
+#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD27/InitCompAD27.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD27/InitCompAD27.cu
deleted file mode 100644
index d343b8bfaa443148b6a1f41a15d3acb5fbbf1586..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD27/InitCompAD27.cu
+++ /dev/null
@@ -1,43 +0,0 @@
-#include "InitCompAD27.h"
-
-#include "InitCompAD27_Device.cuh"
-#include "Parameter/Parameter.h"
-#include <cuda_helper/CudaGrid.h>
-
-std::shared_ptr<PreProcessorStrategy> InitCompAD27::getNewInstance(std::shared_ptr<Parameter> para)
-{
- return std::shared_ptr<PreProcessorStrategy>(new InitCompAD27(para));
-}
-
-void InitCompAD27::init(int level)
-{
- vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
-
- LB_Init_Comp_AD_27 <<< grid.grid, grid.threads >>>(
- para->getParD(level)->neighborX,
- para->getParD(level)->neighborY,
- para->getParD(level)->neighborZ,
- para->getParD(level)->typeOfGridNode,
- para->getParD(level)->concentration,
- para->getParD(level)->velocityX,
- para->getParD(level)->velocityY,
- para->getParD(level)->velocityZ,
- para->getParD(level)->numberOfNodes,
- para->getParD(level)->distributionsAD.f[0],
- para->getParD(level)->isEvenTimestep);
- getLastCudaError("LB_Init_Comp_AD_27 execution failed");
-}
-
-bool InitCompAD27::checkParameter()
-{
- return false;
-}
-
-InitCompAD27::InitCompAD27(std::shared_ptr<Parameter> para)
-{
- this->para = para;
-}
-
-InitCompAD27::InitCompAD27()
-{
-}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD27/InitCompAD27.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD27/InitCompAD27.h
deleted file mode 100644
index 9990f12270426420d4f9af2bb332d2d76d75ebdc..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD27/InitCompAD27.h
+++ /dev/null
@@ -1,23 +0,0 @@
-#ifndef Init_COMP_AD_27_H
-#define Init_COMP_AD_27_H
-
-#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h"
-
-#include <memory>
-
-class Parameter;
-
-class InitCompAD27 : public PreProcessorStrategy
-{
-public:
- static std::shared_ptr<PreProcessorStrategy> getNewInstance(std::shared_ptr< Parameter> para);
- void init(int level);
- bool checkParameter();
-
-private:
- InitCompAD27();
- InitCompAD27(std::shared_ptr< Parameter> para);
- std::shared_ptr< Parameter> para;
-};
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD27/InitCompAD27_Device.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD27/InitCompAD27_Device.cu
deleted file mode 100644
index ccb72094eec63b5de539d59db924a74f7ecceb32..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD27/InitCompAD27_Device.cu
+++ /dev/null
@@ -1,200 +0,0 @@
-#include "LBM/LB.h"
-#include "lbm/constants/D3Q27.h"
-#include <basics/constants/NumericConstants.h>
-
-using namespace vf::basics::constant;
-using namespace vf::lbm::dir;
-
-
-__global__ void LB_Init_Comp_AD_27(unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* geoD,
- real* Conc,
- real* ux,
- real* uy,
- real* uz,
- unsigned int size_Mat,
- real* DD27,
- bool EvenOrOdd)
-{
- ////////////////////////////////////////////////////////////////////////////////
- const unsigned x = threadIdx.x; // Globaler x-Index
- const unsigned y = blockIdx.x; // Globaler y-Index
- const unsigned z = blockIdx.y; // Globaler z-Index
-
- const unsigned nx = blockDim.x;
- const unsigned ny = gridDim.x;
-
- const unsigned k = nx*(ny*z + y) + x;
- //////////////////////////////////////////////////////////////////////////
-
- if (k<size_Mat)
- {
- ////////////////////////////////////////////////////////////////////////////////
- unsigned int BC;
- BC = geoD[k];
-
- if (BC != GEO_SOLID && BC != GEO_VOID)
- {
- Distributions27 D27;
- if (EvenOrOdd == true)
- {
- D27.f[dP00] = &DD27[dP00 *size_Mat];
- D27.f[dM00] = &DD27[dM00 *size_Mat];
- D27.f[d0P0] = &DD27[d0P0 *size_Mat];
- D27.f[d0M0] = &DD27[d0M0 *size_Mat];
- D27.f[d00P] = &DD27[d00P *size_Mat];
- D27.f[d00M] = &DD27[d00M *size_Mat];
- D27.f[dPP0] = &DD27[dPP0 *size_Mat];
- D27.f[dMM0] = &DD27[dMM0 *size_Mat];
- D27.f[dPM0] = &DD27[dPM0 *size_Mat];
- D27.f[dMP0] = &DD27[dMP0 *size_Mat];
- D27.f[dP0P] = &DD27[dP0P *size_Mat];
- D27.f[dM0M] = &DD27[dM0M *size_Mat];
- D27.f[dP0M] = &DD27[dP0M *size_Mat];
- D27.f[dM0P] = &DD27[dM0P *size_Mat];
- D27.f[d0PP] = &DD27[d0PP *size_Mat];
- D27.f[d0MM] = &DD27[d0MM *size_Mat];
- D27.f[d0PM] = &DD27[d0PM *size_Mat];
- D27.f[d0MP] = &DD27[d0MP *size_Mat];
- D27.f[d000] = &DD27[d000*size_Mat];
- D27.f[dPPP] = &DD27[dPPP *size_Mat];
- D27.f[dMMP] = &DD27[dMMP *size_Mat];
- D27.f[dPMP] = &DD27[dPMP *size_Mat];
- D27.f[dMPP] = &DD27[dMPP *size_Mat];
- D27.f[dPPM] = &DD27[dPPM *size_Mat];
- D27.f[dMMM] = &DD27[dMMM *size_Mat];
- D27.f[dPMM]= &DD27[dPMM *size_Mat];
- D27.f[dMPM]= &DD27[dMPM *size_Mat];
- }
- else
- {
- D27.f[dM00] = &DD27[dP00 *size_Mat];
- D27.f[dP00] = &DD27[dM00 *size_Mat];
- D27.f[d0M0] = &DD27[d0P0 *size_Mat];
- D27.f[d0P0] = &DD27[d0M0 *size_Mat];
- D27.f[d00M] = &DD27[d00P *size_Mat];
- D27.f[d00P] = &DD27[d00M *size_Mat];
- D27.f[dMM0] = &DD27[dPP0 *size_Mat];
- D27.f[dPP0] = &DD27[dMM0 *size_Mat];
- D27.f[dMP0] = &DD27[dPM0 *size_Mat];
- D27.f[dPM0] = &DD27[dMP0 *size_Mat];
- D27.f[dM0M] = &DD27[dP0P *size_Mat];
- D27.f[dP0P] = &DD27[dM0M *size_Mat];
- D27.f[dM0P] = &DD27[dP0M *size_Mat];
- D27.f[dP0M] = &DD27[dM0P *size_Mat];
- D27.f[d0MM] = &DD27[d0PP *size_Mat];
- D27.f[d0PP] = &DD27[d0MM *size_Mat];
- D27.f[d0MP] = &DD27[d0PM *size_Mat];
- D27.f[d0PM] = &DD27[d0MP *size_Mat];
- D27.f[d000] = &DD27[d000*size_Mat];
- D27.f[dMMM] = &DD27[dPPP *size_Mat];
- D27.f[dPPM] = &DD27[dMMP *size_Mat];
- D27.f[dMPM]= &DD27[dPMP *size_Mat];
- D27.f[dPMM]= &DD27[dMPP *size_Mat];
- D27.f[dMMP] = &DD27[dPPM *size_Mat];
- D27.f[dPPP] = &DD27[dMMM *size_Mat];
- D27.f[dMPP] = &DD27[dPMM *size_Mat];
- D27.f[dPMP] = &DD27[dMPM *size_Mat];
- }
- //////////////////////////////////////////////////////////////////////////
- real ConcD = Conc[k];
- real vx1 = ux[k];
- real vx2 = uy[k];
- real vx3 = uz[k];
- //real lambdaD = -three + sqrt(three);
- //real Diffusivity = c1o20;
- //real Lam = -(c1o2+one/lambdaD);
- //real nue_d = Lam/three;
- //real ae = Diffusivity/nue_d - one;
- //real ux_sq = vx1 * vx1;
- //real uy_sq = vx2 * vx2;
- //real uz_sq = vx3 * vx3;
- ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
- //D3Q7
- ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
- //index
- //unsigned int kzero= k;
- //unsigned int ke = k;
- //unsigned int kw = neighborX[k];
- //unsigned int kn = k;
- //unsigned int ks = neighborY[k];
- //unsigned int kt = k;
- //unsigned int kb = neighborZ[k];
- //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
- //(D7.f[0])[kzero] = ConcD*(c1o3*(ae*(-three))-(ux_sq+uy_sq+uz_sq));
- //(D7.f[1])[ke ] = ConcD*(c1o6*(ae+one)+c1o2*(ux_sq)+vx1*c1o2);
- //(D7.f[2])[kw ] = ConcD*(c1o6*(ae+one)+c1o2*(ux_sq)-vx1*c1o2);
- //(D7.f[3])[kn ] = ConcD*(c1o6*(ae+one)+c1o2*(uy_sq)+vx2*c1o2);
- //(D7.f[4])[ks ] = ConcD*(c1o6*(ae+one)+c1o2*(uy_sq)-vx2*c1o2);
- //(D7.f[5])[kt ] = ConcD*(c1o6*(ae+one)+c1o2*(uz_sq)+vx3*c1o2);
- //(D7.f[6])[kb ] = ConcD*(c1o6*(ae+one)+c1o2*(uz_sq)-vx3*c1o2);
- ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
-
-
- ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
- //D3Q27
- ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
- //index
- unsigned int kzero = k;
- unsigned int ke = k;
- unsigned int kw = neighborX[k];
- unsigned int kn = k;
- unsigned int ks = neighborY[k];
- unsigned int kt = k;
- unsigned int kb = neighborZ[k];
- unsigned int ksw = neighborY[kw];
- unsigned int kne = k;
- unsigned int kse = ks;
- unsigned int knw = kw;
- unsigned int kbw = neighborZ[kw];
- unsigned int kte = k;
- unsigned int kbe = kb;
- unsigned int ktw = kw;
- unsigned int kbs = neighborZ[ks];
- unsigned int ktn = k;
- unsigned int kbn = kb;
- unsigned int kts = ks;
- unsigned int ktse = ks;
- unsigned int kbnw = kbw;
- unsigned int ktnw = kw;
- unsigned int kbse = kbs;
- unsigned int ktsw = ksw;
- unsigned int kbne = kb;
- unsigned int ktne = k;
- unsigned int kbsw = neighborZ[ksw];
- ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
- real cu_sq = c3o2*(vx1*vx1 + vx2*vx2 + vx3*vx3);
-
- (D27.f[d000])[kzero] = c8o27* ConcD*(c1o1 - cu_sq);
- (D27.f[dP00])[ke] = c2o27* ConcD*(c1o1 + c3o1*(vx1)+c9o2*(vx1)*(vx1)-cu_sq);
- (D27.f[dM00])[kw] = c2o27* ConcD*(c1o1 + c3o1*(-vx1) + c9o2*(-vx1)*(-vx1) - cu_sq);
- (D27.f[d0P0])[kn] = c2o27* ConcD*(c1o1 + c3o1*(vx2)+c9o2*(vx2)*(vx2)-cu_sq);
- (D27.f[d0M0])[ks] = c2o27* ConcD*(c1o1 + c3o1*(-vx2) + c9o2*(-vx2)*(-vx2) - cu_sq);
- (D27.f[d00P])[kt] = c2o27* ConcD*(c1o1 + c3o1*(vx3)+c9o2*(vx3)*(vx3)-cu_sq);
- (D27.f[d00M])[kb] = c2o27* ConcD*(c1o1 + c3o1*(-vx3) + c9o2*(-vx3)*(-vx3) - cu_sq);
- (D27.f[dPP0])[kne] = c1o54* ConcD*(c1o1 + c3o1*(vx1 + vx2) + c9o2*(vx1 + vx2)*(vx1 + vx2) - cu_sq);
- (D27.f[dMM0])[ksw] = c1o54* ConcD*(c1o1 + c3o1*(-vx1 - vx2) + c9o2*(-vx1 - vx2)*(-vx1 - vx2) - cu_sq);
- (D27.f[dPM0])[kse] = c1o54* ConcD*(c1o1 + c3o1*(vx1 - vx2) + c9o2*(vx1 - vx2)*(vx1 - vx2) - cu_sq);
- (D27.f[dMP0])[knw] = c1o54* ConcD*(c1o1 + c3o1*(-vx1 + vx2) + c9o2*(-vx1 + vx2)*(-vx1 + vx2) - cu_sq);
- (D27.f[dP0P])[kte] = c1o54* ConcD*(c1o1 + c3o1*(vx1 + vx3) + c9o2*(vx1 + vx3)*(vx1 + vx3) - cu_sq);
- (D27.f[dM0M])[kbw] = c1o54* ConcD*(c1o1 + c3o1*(-vx1 - vx3) + c9o2*(-vx1 - vx3)*(-vx1 - vx3) - cu_sq);
- (D27.f[dP0M])[kbe] = c1o54* ConcD*(c1o1 + c3o1*(vx1 - vx3) + c9o2*(vx1 - vx3)*(vx1 - vx3) - cu_sq);
- (D27.f[dM0P])[ktw] = c1o54* ConcD*(c1o1 + c3o1*(-vx1 + vx3) + c9o2*(-vx1 + vx3)*(-vx1 + vx3) - cu_sq);
- (D27.f[d0PP])[ktn] = c1o54* ConcD*(c1o1 + c3o1*(vx2 + vx3) + c9o2*(vx2 + vx3)*(vx2 + vx3) - cu_sq);
- (D27.f[d0MM])[kbs] = c1o54* ConcD*(c1o1 + c3o1*(-vx2 - vx3) + c9o2*(-vx2 - vx3)*(-vx2 - vx3) - cu_sq);
- (D27.f[d0PM])[kbn] = c1o54* ConcD*(c1o1 + c3o1*(vx2 - vx3) + c9o2*(vx2 - vx3)*(vx2 - vx3) - cu_sq);
- (D27.f[d0MP])[kts] = c1o54* ConcD*(c1o1 + c3o1*(-vx2 + vx3) + c9o2*(-vx2 + vx3)*(-vx2 + vx3) - cu_sq);
- (D27.f[dPPP])[ktne] = c1o216*ConcD*(c1o1 + c3o1*(vx1 + vx2 + vx3) + c9o2*(vx1 + vx2 + vx3)*(vx1 + vx2 + vx3) - cu_sq);
- (D27.f[dMMM])[kbsw] = c1o216*ConcD*(c1o1 + c3o1*(-vx1 - vx2 - vx3) + c9o2*(-vx1 - vx2 - vx3)*(-vx1 - vx2 - vx3) - cu_sq);
- (D27.f[dPPM])[kbne] = c1o216*ConcD*(c1o1 + c3o1*(vx1 + vx2 - vx3) + c9o2*(vx1 + vx2 - vx3)*(vx1 + vx2 - vx3) - cu_sq);
- (D27.f[dMMP])[ktsw] = c1o216*ConcD*(c1o1 + c3o1*(-vx1 - vx2 + vx3) + c9o2*(-vx1 - vx2 + vx3)*(-vx1 - vx2 + vx3) - cu_sq);
- (D27.f[dPMP])[ktse] = c1o216*ConcD*(c1o1 + c3o1*(vx1 - vx2 + vx3) + c9o2*(vx1 - vx2 + vx3)*(vx1 - vx2 + vx3) - cu_sq);
- (D27.f[dMPM])[kbnw] = c1o216*ConcD*(c1o1 + c3o1*(-vx1 + vx2 - vx3) + c9o2*(-vx1 + vx2 - vx3)*(-vx1 + vx2 - vx3) - cu_sq);
- (D27.f[dPMM])[kbse] = c1o216*ConcD*(c1o1 + c3o1*(vx1 - vx2 - vx3) + c9o2*(vx1 - vx2 - vx3)*(vx1 - vx2 - vx3) - cu_sq);
- (D27.f[dMPP])[ktnw] = c1o216*ConcD*(c1o1 + c3o1*(-vx1 + vx2 + vx3) + c9o2*(-vx1 + vx2 + vx3)*(-vx1 + vx2 + vx3) - cu_sq);
- ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
- }
- }
-}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD27/InitCompAD27_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD27/InitCompAD27_Device.cuh
deleted file mode 100644
index dff2bb8a8843dcf19c96bd01cd299b5672b6c8d5..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD27/InitCompAD27_Device.cuh
+++ /dev/null
@@ -1,19 +0,0 @@
-#ifndef LB_INIT_COMP_AD_27_H
-#define LB_INIT_COMP_AD_27_H
-
-#include <DataTypes.h>
-#include <curand.h>
-
-__global__ void LB_Init_Comp_AD_27(unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* geoD,
- real* Conc,
- real* ux,
- real* uy,
- real* uz,
- unsigned int size_Mat,
- real* DD27,
- bool EvenOrOdd);
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD7/InitCompAD7.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD7/InitCompAD7.cu
deleted file mode 100644
index b6a3ef4c499c808de6133be8948c82d602758f03..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD7/InitCompAD7.cu
+++ /dev/null
@@ -1,43 +0,0 @@
-#include "InitCompAD7.h"
-
-#include "InitCompAD7_Device.cuh"
-#include "Parameter/Parameter.h"
-#include <cuda_helper/CudaGrid.h>
-
-std::shared_ptr<InitCompAD7> InitCompAD7::getNewInstance(std::shared_ptr<Parameter> para)
-{
- return std::shared_ptr<InitCompAD7>(new InitCompAD7(para));
-}
-
-void InitCompAD7::init(int level)
-{
- vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
-
- LB_Init_Comp_AD_7 <<< grid.grid, grid.threads >>>(
- para->getParD(level)->neighborX,
- para->getParD(level)->neighborY,
- para->getParD(level)->neighborZ,
- para->getParD(level)->typeOfGridNode,
- para->getParD(level)->concentration,
- para->getParD(level)->velocityX,
- para->getParD(level)->velocityY,
- para->getParD(level)->velocityZ,
- para->getParD(level)->numberOfNodes,
- para->getParD(level)->distributionsAD7.f[0],
- para->getParD(level)->isEvenTimestep);
- getLastCudaError("LB_Init_Comp_AD_7 execution failed");
-}
-
-bool InitCompAD7::checkParameter()
-{
- return false;
-}
-
-InitCompAD7::InitCompAD7(std::shared_ptr<Parameter> para)
-{
- this->para = para;
-}
-
-InitCompAD7::InitCompAD7()
-{
-}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD7/InitCompAD7.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD7/InitCompAD7.h
deleted file mode 100644
index c1e76f8ec40a15394f1cd5017a7c5b3e95912bed..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD7/InitCompAD7.h
+++ /dev/null
@@ -1,22 +0,0 @@
-#ifndef INIT_COMP_AD_7_H
-#define INIT_COMP_AD_7_H
-
-#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h"
-
-#include <memory>
-
-class Parameter;
-
-class InitCompAD7 : public PreProcessorStrategy
-{
-public:
- static std::shared_ptr<InitCompAD7> getNewInstance(std::shared_ptr< Parameter> para);
- void init(int level);
- bool checkParameter();
-
-private:
- InitCompAD7();
- InitCompAD7(std::shared_ptr< Parameter> para);
- std::shared_ptr< Parameter> para;
-};
-#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD7/InitCompAD7_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD7/InitCompAD7_Device.cuh
deleted file mode 100644
index cd12cd1a9dfe8edaf35b3a6eff5391add4648a43..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompAD7/InitCompAD7_Device.cuh
+++ /dev/null
@@ -1,19 +0,0 @@
-#ifndef LB_INIT_COMP_AD_7_H
-#define LB_INIT_COMP_AD_7_H
-
-#include <DataTypes.h>
-#include <curand.h>
-
-__global__ void LB_Init_Comp_AD_7(unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* geoD,
- real* Conc,
- real* ux,
- real* uy,
- real* uz,
- unsigned int size_Mat,
- real* DD7,
- bool EvenOrOdd);
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27.cu
deleted file mode 100644
index e70d5a8fe61e7c3119d83e35d81bbe742f97fa4e..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27.cu
+++ /dev/null
@@ -1,68 +0,0 @@
-#include "InitCompSP27.h"
-
-#include "InitCompSP27_Device.cuh"
-#include "Parameter/Parameter.h"
-#include <cuda_helper/CudaGrid.h>
-
-std::shared_ptr<PreProcessorStrategy> InitCompSP27::getNewInstance(std::shared_ptr<Parameter> para)
-{
- return std::shared_ptr<PreProcessorStrategy>(new InitCompSP27(para));
-}
-
-void InitCompSP27::init(int level)
-{
- vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
-
- if( ! para->getUseInitNeq() )
- {
- LB_Init_Comp_SP_27 <<< grid.grid, grid.threads >>> (
- para->getParD(level)->neighborX,
- para->getParD(level)->neighborY,
- para->getParD(level)->neighborZ,
- para->getParD(level)->typeOfGridNode,
- para->getParD(level)->rho,
- para->getParD(level)->velocityX,
- para->getParD(level)->velocityY,
- para->getParD(level)->velocityZ,
- para->getParD(level)->numberOfNodes,
- para->getParD(level)->distributions.f[0],
- para->getParD(level)->isEvenTimestep);
- getLastCudaError("LB_Init_Comp_SP_27 execution failed");
- }
- else
- {
- LB_Init_Comp_Neq_SP_27 <<< grid.grid, grid.threads >>> (
- para->getParD(level)->neighborX,
- para->getParD(level)->neighborY,
- para->getParD(level)->neighborZ,
- para->getParD(level)->neighborInverse,
- para->getParD(level)->typeOfGridNode,
- para->getParD(level)->rho,
- para->getParD(level)->velocityX,
- para->getParD(level)->velocityY,
- para->getParD(level)->velocityZ,
- para->getParD(level)->numberOfNodes,
- para->getParD(level)->distributions.f[0],
- para->getParD(level)->omega,
- para->getParD(level)->isEvenTimestep);
- cudaDeviceSynchronize();
- getLastCudaError("LB_Init_Comp_Neq_SP_27 execution failed");
- }
-
-
-
-}
-
-bool InitCompSP27::checkParameter()
-{
- return false;
-}
-
-InitCompSP27::InitCompSP27(std::shared_ptr<Parameter> para)
-{
- this->para = para;
-}
-
-InitCompSP27::InitCompSP27()
-{
-}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27.h
deleted file mode 100644
index 601f1a446ade70766d6b37953db77a3a92f1f67f..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27.h
+++ /dev/null
@@ -1,23 +0,0 @@
-#ifndef INIT_COMP_SP27_H
-#define INIT_COMP_SP27_H
-
-#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h"
-
-#include <memory>
-
-class Parameter;
-
-class InitCompSP27 : public PreProcessorStrategy
-{
-public:
- static std::shared_ptr<PreProcessorStrategy> getNewInstance(std::shared_ptr< Parameter> para);
- void init(int level);
- bool checkParameter();
-
-private:
- InitCompSP27();
- InitCompSP27(std::shared_ptr< Parameter> para);
- std::shared_ptr< Parameter> para;
-};
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27_Device.cuh
deleted file mode 100644
index d7f40b88bf53652de893076d529f217647c28293..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27_Device.cuh
+++ /dev/null
@@ -1,33 +0,0 @@
-#ifndef LB_INIT_COMP_SP27_H
-#define LB_INIT_COMP_SP27_H
-
-#include <DataTypes.h>
-#include <curand.h>
-
-__global__ void LB_Init_Comp_SP_27(unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* geoD,
- real* rho,
- real* ux,
- real* uy,
- real* uz,
- unsigned int size_Mat,
- real* DD,
- bool EvenOrOdd);
-
-__global__ void LB_Init_Comp_Neq_SP_27(unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* neighborWSB,
- unsigned int* geoD,
- real* rho,
- real* ux,
- real* uy,
- real* uz,
- unsigned int size_Mat,
- real* DD,
- real omega,
- bool EvenOrOdd);
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3.cu
deleted file mode 100644
index 82084e354ddd87c685f1913b1236bb3a01ba863b..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3.cu
+++ /dev/null
@@ -1,43 +0,0 @@
-#include "InitF3.h"
-
-#include "InitF3_Device.cuh"
-#include "Parameter/Parameter.h"
-#include <cuda_helper/CudaGrid.h>
-
-std::shared_ptr<PreProcessorStrategy> InitF3::getNewInstance(std::shared_ptr<Parameter> para)
-{
- return std::shared_ptr<PreProcessorStrategy>(new InitF3(para));
-}
-
-void InitF3::init(int level)
-{
- vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
-
- LB_Init_F3 <<< grid.grid, grid.threads >>>(
- para->getParD(level)->neighborX,
- para->getParD(level)->neighborY,
- para->getParD(level)->neighborZ,
- para->getParD(level)->typeOfGridNode,
- para->getParD(level)->rho,
- para->getParD(level)->velocityX,
- para->getParD(level)->velocityY,
- para->getParD(level)->velocityZ,
- para->getParD(level)->numberOfNodes,
- para->getParD(level)->g6.g[0],
- para->getParD(level)->isEvenTimestep);
- getLastCudaError("LB_Init_F3 execution failed");
-}
-
-bool InitF3::checkParameter()
-{
- return false;
-}
-
-InitF3::InitF3(std::shared_ptr<Parameter> para)
-{
- this->para = para;
-}
-
-InitF3::InitF3()
-{
-}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3.h
deleted file mode 100644
index c254aa906a31f61d1d82f86fa40bf048b4443fb4..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3.h
+++ /dev/null
@@ -1,23 +0,0 @@
-#ifndef INITF3_H
-#define INITF3_H
-
-#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h"
-
-#include <memory>
-
-class Parameter;
-
-class InitF3 : public PreProcessorStrategy
-{
-public:
- static std::shared_ptr<PreProcessorStrategy> getNewInstance(std::shared_ptr< Parameter> para);
- void init(int level);
- bool checkParameter();
-
-private:
- InitF3();
- InitF3(std::shared_ptr< Parameter> para);
- std::shared_ptr< Parameter> para;
-};
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3_Device.cuh
deleted file mode 100644
index 991a5ec75c9a05eefdfb34e0faad921cbe35d6ea..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3_Device.cuh
+++ /dev/null
@@ -1,19 +0,0 @@
-#ifndef LB_INIT_F3_H
-#define LB_INIT_F3_H
-
-#include <DataTypes.h>
-#include <curand.h>
-
-__global__ void LB_Init_F3(unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* geoD,
- real* rho,
- real* ux,
- real* uy,
- real* uz,
- unsigned int size_Mat,
- real* G6,
- bool EvenOrOdd);
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD27/InitIncompAD27.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD27/InitIncompAD27.cu
deleted file mode 100644
index f09d9833692adad4cae2e24a1159714275569a22..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD27/InitIncompAD27.cu
+++ /dev/null
@@ -1,43 +0,0 @@
-#include "InitIncompAD27.h"
-
-#include "InitIncompAD27_Device.cuh"
-#include "Parameter/Parameter.h"
-#include <cuda_helper/CudaGrid.h>
-
-std::shared_ptr<PreProcessorStrategy> InitIncompAD27::getNewInstance(std::shared_ptr<Parameter> para)
-{
- return std::shared_ptr<PreProcessorStrategy>(new InitIncompAD27(para));
-}
-
-void InitIncompAD27::init(int level)
-{
- vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
-
- LB_Init_Incomp_AD_27 <<< grid.grid, grid.threads >>>(
- para->getParD(level)->neighborX,
- para->getParD(level)->neighborY,
- para->getParD(level)->neighborZ,
- para->getParD(level)->typeOfGridNode,
- para->getParD(level)->concentration,
- para->getParD(level)->velocityX,
- para->getParD(level)->velocityY,
- para->getParD(level)->velocityZ,
- para->getParD(level)->numberOfNodes,
- para->getParD(level)->distributionsAD.f[0],
- para->getParD(level)->isEvenTimestep);
- getLastCudaError("LB_Init_Incomp_AD_27 execution failed");
-}
-
-bool InitIncompAD27::checkParameter()
-{
- return false;
-}
-
-InitIncompAD27::InitIncompAD27(std::shared_ptr<Parameter> para)
-{
- this->para = para;
-}
-
-InitIncompAD27::InitIncompAD27()
-{
-}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD27/InitIncompAD27.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD27/InitIncompAD27.h
deleted file mode 100644
index 9654345cf2fdc1d272be4766b88f6c1f67f27397..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD27/InitIncompAD27.h
+++ /dev/null
@@ -1,23 +0,0 @@
-#ifndef INIT_INCOMP_AD27_H
-#define INIT_INCOMP_AD27_H
-
-#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h"
-
-#include <memory>
-
-class Parameter;
-
-class InitIncompAD27 : public PreProcessorStrategy
-{
-public:
- static std::shared_ptr<PreProcessorStrategy> getNewInstance(std::shared_ptr< Parameter> para);
- void init(int level);
- bool checkParameter();
-
-private:
- InitIncompAD27();
- InitIncompAD27(std::shared_ptr< Parameter> para);
- std::shared_ptr< Parameter> para;
-};
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD27/InitIncompAD27_Device.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD27/InitIncompAD27_Device.cu
deleted file mode 100644
index a2933ca6822d256e997459fc4f741d5d2c7fe14e..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD27/InitIncompAD27_Device.cu
+++ /dev/null
@@ -1,173 +0,0 @@
-#include "LBM/LB.h"
-#include "lbm/constants/D3Q27.h"
-#include <basics/constants/NumericConstants.h>
-
-using namespace vf::basics::constant;
-using namespace vf::lbm::dir;
-#include "math.h"
-
-__global__ void LB_Init_Incomp_AD_27(unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* geoD,
- real* Conc,
- real* ux,
- real* uy,
- real* uz,
- unsigned int size_Mat,
- real* DD27,
- bool EvenOrOdd)
-{
- ////////////////////////////////////////////////////////////////////////////////
- const unsigned x = threadIdx.x; // Globaler x-Index
- const unsigned y = blockIdx.x; // Globaler y-Index
- const unsigned z = blockIdx.y; // Globaler z-Index
-
- const unsigned nx = blockDim.x;
- const unsigned ny = gridDim.x;
-
- const unsigned k = nx*(ny*z + y) + x;
- //////////////////////////////////////////////////////////////////////////
-
- if(k<size_Mat)
- {
- ////////////////////////////////////////////////////////////////////////////////
- unsigned int BC;
- BC = geoD[k];
-
- if( BC != GEO_SOLID && BC != GEO_VOID)
- {
- Distributions27 D27;
- if (EvenOrOdd==true)
- {
- D27.f[dP00 ] = &DD27[dP00 *size_Mat];
- D27.f[dM00 ] = &DD27[dM00 *size_Mat];
- D27.f[d0P0 ] = &DD27[d0P0 *size_Mat];
- D27.f[d0M0 ] = &DD27[d0M0 *size_Mat];
- D27.f[d00P ] = &DD27[d00P *size_Mat];
- D27.f[d00M ] = &DD27[d00M *size_Mat];
- D27.f[dPP0 ] = &DD27[dPP0 *size_Mat];
- D27.f[dMM0 ] = &DD27[dMM0 *size_Mat];
- D27.f[dPM0 ] = &DD27[dPM0 *size_Mat];
- D27.f[dMP0 ] = &DD27[dMP0 *size_Mat];
- D27.f[dP0P ] = &DD27[dP0P *size_Mat];
- D27.f[dM0M ] = &DD27[dM0M *size_Mat];
- D27.f[dP0M ] = &DD27[dP0M *size_Mat];
- D27.f[dM0P ] = &DD27[dM0P *size_Mat];
- D27.f[d0PP ] = &DD27[d0PP *size_Mat];
- D27.f[d0MM ] = &DD27[d0MM *size_Mat];
- D27.f[d0PM ] = &DD27[d0PM *size_Mat];
- D27.f[d0MP ] = &DD27[d0MP *size_Mat];
- D27.f[d000] = &DD27[d000*size_Mat];
- D27.f[dPPP ] = &DD27[dPPP *size_Mat];
- D27.f[dMMP ] = &DD27[dMMP *size_Mat];
- D27.f[dPMP ] = &DD27[dPMP *size_Mat];
- D27.f[dMPP ] = &DD27[dMPP *size_Mat];
- D27.f[dPPM ] = &DD27[dPPM *size_Mat];
- D27.f[dMMM ] = &DD27[dMMM *size_Mat];
- D27.f[dPMM ] = &DD27[dPMM *size_Mat];
- D27.f[dMPM ] = &DD27[dMPM *size_Mat];
- }
- else
- {
- D27.f[dM00 ] = &DD27[dP00 *size_Mat];
- D27.f[dP00 ] = &DD27[dM00 *size_Mat];
- D27.f[d0M0 ] = &DD27[d0P0 *size_Mat];
- D27.f[d0P0 ] = &DD27[d0M0 *size_Mat];
- D27.f[d00M ] = &DD27[d00P *size_Mat];
- D27.f[d00P ] = &DD27[d00M *size_Mat];
- D27.f[dMM0 ] = &DD27[dPP0 *size_Mat];
- D27.f[dPP0 ] = &DD27[dMM0 *size_Mat];
- D27.f[dMP0 ] = &DD27[dPM0 *size_Mat];
- D27.f[dPM0 ] = &DD27[dMP0 *size_Mat];
- D27.f[dM0M ] = &DD27[dP0P *size_Mat];
- D27.f[dP0P ] = &DD27[dM0M *size_Mat];
- D27.f[dM0P ] = &DD27[dP0M *size_Mat];
- D27.f[dP0M ] = &DD27[dM0P *size_Mat];
- D27.f[d0MM ] = &DD27[d0PP *size_Mat];
- D27.f[d0PP ] = &DD27[d0MM *size_Mat];
- D27.f[d0MP ] = &DD27[d0PM *size_Mat];
- D27.f[d0PM ] = &DD27[d0MP *size_Mat];
- D27.f[d000] = &DD27[d000*size_Mat];
- D27.f[dMMM ] = &DD27[dPPP *size_Mat];
- D27.f[dPPM ] = &DD27[dMMP *size_Mat];
- D27.f[dMPM ] = &DD27[dPMP *size_Mat];
- D27.f[dPMM ] = &DD27[dMPP *size_Mat];
- D27.f[dMMP ] = &DD27[dPPM *size_Mat];
- D27.f[dPPP ] = &DD27[dMMM *size_Mat];
- D27.f[dMPP ] = &DD27[dPMM *size_Mat];
- D27.f[dPMP ] = &DD27[dMPM *size_Mat];
- }
- //////////////////////////////////////////////////////////////////////////
- real ConcD = Conc[k];
- real vx1 = ux[k];
- real vx2 = uy[k];
- real vx3 = uz[k];
- ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
-
-
- ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
- //D3Q27
- ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
- //index
- unsigned int kzero= k;
- unsigned int ke = k;
- unsigned int kw = neighborX[k];
- unsigned int kn = k;
- unsigned int ks = neighborY[k];
- unsigned int kt = k;
- unsigned int kb = neighborZ[k];
- unsigned int ksw = neighborY[kw];
- unsigned int kne = k;
- unsigned int kse = ks;
- unsigned int knw = kw;
- unsigned int kbw = neighborZ[kw];
- unsigned int kte = k;
- unsigned int kbe = kb;
- unsigned int ktw = kw;
- unsigned int kbs = neighborZ[ks];
- unsigned int ktn = k;
- unsigned int kbn = kb;
- unsigned int kts = ks;
- unsigned int ktse = ks;
- unsigned int kbnw = kbw;
- unsigned int ktnw = kw;
- unsigned int kbse = kbs;
- unsigned int ktsw = ksw;
- unsigned int kbne = kb;
- unsigned int ktne = k;
- unsigned int kbsw = neighborZ[ksw];
- ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
- real cu_sq=c3o2*(vx1*vx1+vx2*vx2+vx3*vx3);
-
- (D27.f[d000])[kzero] = c8o27* ConcD*(c1o1-cu_sq);
- (D27.f[dP00 ])[ke ] = c2o27* ConcD*(c1o1+c3o1*( vx1 )+c9o2*( vx1 )*( vx1 )-cu_sq);
- (D27.f[dM00 ])[kw ] = c2o27* ConcD*(c1o1+c3o1*(-vx1 )+c9o2*(-vx1 )*(-vx1 )-cu_sq);
- (D27.f[d0P0 ])[kn ] = c2o27* ConcD*(c1o1+c3o1*( vx2 )+c9o2*( vx2 )*( vx2 )-cu_sq);
- (D27.f[d0M0 ])[ks ] = c2o27* ConcD*(c1o1+c3o1*( -vx2 )+c9o2*( -vx2 )*( -vx2 )-cu_sq);
- (D27.f[d00P ])[kt ] = c2o27* ConcD*(c1o1+c3o1*( vx3)+c9o2*( vx3)*( vx3)-cu_sq);
- (D27.f[d00M ])[kb ] = c2o27* ConcD*(c1o1+c3o1*( -vx3)+c9o2*( -vx3)*( -vx3)-cu_sq);
- (D27.f[dPP0 ])[kne ] = c1o54* ConcD*(c1o1+c3o1*( vx1+vx2 )+c9o2*( vx1+vx2 )*( vx1+vx2 )-cu_sq);
- (D27.f[dMM0 ])[ksw ] = c1o54* ConcD*(c1o1+c3o1*(-vx1-vx2 )+c9o2*(-vx1-vx2 )*(-vx1-vx2 )-cu_sq);
- (D27.f[dPM0 ])[kse ] = c1o54* ConcD*(c1o1+c3o1*( vx1-vx2 )+c9o2*( vx1-vx2 )*( vx1-vx2 )-cu_sq);
- (D27.f[dMP0 ])[knw ] = c1o54* ConcD*(c1o1+c3o1*(-vx1+vx2 )+c9o2*(-vx1+vx2 )*(-vx1+vx2 )-cu_sq);
- (D27.f[dP0P ])[kte ] = c1o54* ConcD*(c1o1+c3o1*( vx1 +vx3)+c9o2*( vx1 +vx3)*( vx1 +vx3)-cu_sq);
- (D27.f[dM0M ])[kbw ] = c1o54* ConcD*(c1o1+c3o1*(-vx1 -vx3)+c9o2*(-vx1 -vx3)*(-vx1 -vx3)-cu_sq);
- (D27.f[dP0M ])[kbe ] = c1o54* ConcD*(c1o1+c3o1*( vx1 -vx3)+c9o2*( vx1 -vx3)*( vx1 -vx3)-cu_sq);
- (D27.f[dM0P ])[ktw ] = c1o54* ConcD*(c1o1+c3o1*(-vx1 +vx3)+c9o2*(-vx1 +vx3)*(-vx1 +vx3)-cu_sq);
- (D27.f[d0PP ])[ktn ] = c1o54* ConcD*(c1o1+c3o1*( vx2+vx3)+c9o2*( vx2+vx3)*( vx2+vx3)-cu_sq);
- (D27.f[d0MM ])[kbs ] = c1o54* ConcD*(c1o1+c3o1*( -vx2-vx3)+c9o2*( -vx2-vx3)*( -vx2-vx3)-cu_sq);
- (D27.f[d0PM ])[kbn ] = c1o54* ConcD*(c1o1+c3o1*( vx2-vx3)+c9o2*( vx2-vx3)*( vx2-vx3)-cu_sq);
- (D27.f[d0MP ])[kts ] = c1o54* ConcD*(c1o1+c3o1*( -vx2+vx3)+c9o2*( -vx2+vx3)*( -vx2+vx3)-cu_sq);
- (D27.f[dPPP ])[ktne ] = c1o216*ConcD*(c1o1+c3o1*( vx1+vx2+vx3)+c9o2*( vx1+vx2+vx3)*( vx1+vx2+vx3)-cu_sq);
- (D27.f[dMMM ])[kbsw ] = c1o216*ConcD*(c1o1+c3o1*(-vx1-vx2-vx3)+c9o2*(-vx1-vx2-vx3)*(-vx1-vx2-vx3)-cu_sq);
- (D27.f[dPPM ])[kbne ] = c1o216*ConcD*(c1o1+c3o1*( vx1+vx2-vx3)+c9o2*( vx1+vx2-vx3)*( vx1+vx2-vx3)-cu_sq);
- (D27.f[dMMP ])[ktsw ] = c1o216*ConcD*(c1o1+c3o1*(-vx1-vx2+vx3)+c9o2*(-vx1-vx2+vx3)*(-vx1-vx2+vx3)-cu_sq);
- (D27.f[dPMP ])[ktse ] = c1o216*ConcD*(c1o1+c3o1*( vx1-vx2+vx3)+c9o2*( vx1-vx2+vx3)*( vx1-vx2+vx3)-cu_sq);
- (D27.f[dMPM ])[kbnw ] = c1o216*ConcD*(c1o1+c3o1*(-vx1+vx2-vx3)+c9o2*(-vx1+vx2-vx3)*(-vx1+vx2-vx3)-cu_sq);
- (D27.f[dPMM ])[kbse ] = c1o216*ConcD*(c1o1+c3o1*( vx1-vx2-vx3)+c9o2*( vx1-vx2-vx3)*( vx1-vx2-vx3)-cu_sq);
- (D27.f[dMPP ])[ktnw ] = c1o216*ConcD*(c1o1+c3o1*(-vx1+vx2+vx3)+c9o2*(-vx1+vx2+vx3)*(-vx1+vx2+vx3)-cu_sq);
- ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
- }
- }
-}
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD27/InitIncompAD27_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD27/InitIncompAD27_Device.cuh
deleted file mode 100644
index b81cf31d3064aa3eaa5fbdc7dc3b9c3150bb61ca..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD27/InitIncompAD27_Device.cuh
+++ /dev/null
@@ -1,19 +0,0 @@
-#ifndef LB_INIT_INCOMP_AD27_H
-#define LB_INIT_INCOMP_AD27_H
-
-#include <DataTypes.h>
-#include <curand.h>
-
-__global__ void LB_Init_Incomp_AD_27(unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* geoD,
- real* Conc,
- real* ux,
- real* uy,
- real* uz,
- unsigned int size_Mat,
- real* DD27,
- bool EvenOrOdd);
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD7/InitIncompAD7.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD7/InitIncompAD7.cu
deleted file mode 100644
index d987cbeed114b85a950b1d0f911486c2e5491bde..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD7/InitIncompAD7.cu
+++ /dev/null
@@ -1,43 +0,0 @@
-#include "InitIncompAD7.h"
-
-#include "InitIncompAD7_Device.cuh"
-#include "Parameter/Parameter.h"
-#include <cuda_helper/CudaGrid.h>
-
-std::shared_ptr<PreProcessorStrategy> InitIncompAD7::getNewInstance(std::shared_ptr<Parameter> para)
-{
- return std::shared_ptr<PreProcessorStrategy>(new InitIncompAD7(para));
-}
-
-void InitIncompAD7::init(int level)
-{
- vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
-
- LB_Init_Incomp_AD_7 <<< grid.grid, grid.threads >>>(
- para->getParD(level)->neighborX,
- para->getParD(level)->neighborY,
- para->getParD(level)->neighborZ,
- para->getParD(level)->typeOfGridNode,
- para->getParD(level)->concentration,
- para->getParD(level)->velocityX,
- para->getParD(level)->velocityY,
- para->getParD(level)->velocityZ,
- para->getParD(level)->numberOfNodes,
- para->getParD(level)->distributionsAD.f[0],
- para->getParD(level)->isEvenTimestep);
- getLastCudaError("LB_Init_Incomp_AD_7 execution failed");
-}
-
-bool InitIncompAD7::checkParameter()
-{
- return false;
-}
-
-InitIncompAD7::InitIncompAD7(std::shared_ptr<Parameter> para)
-{
- this->para = para;
-}
-
-InitIncompAD7::InitIncompAD7()
-{
-}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD7/InitIncompAD7.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD7/InitIncompAD7.h
deleted file mode 100644
index 33bb82c7879590422ab111f4a8e518882be692c1..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD7/InitIncompAD7.h
+++ /dev/null
@@ -1,23 +0,0 @@
-#ifndef INIT_INCOMP_AD7_H
-#define INIT_INCOMP_AD7_H
-
-#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h"
-
-#include <memory>
-
-class Parameter;
-
-class InitIncompAD7 : public PreProcessorStrategy
-{
-public:
- static std::shared_ptr<PreProcessorStrategy> getNewInstance(std::shared_ptr< Parameter> para);
- void init(int level);
- bool checkParameter();
-
-private:
- InitIncompAD7();
- InitIncompAD7(std::shared_ptr< Parameter> para);
- std::shared_ptr< Parameter> para;
-};
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD7/InitIncompAD7_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD7/InitIncompAD7_Device.cuh
deleted file mode 100644
index f773ae81fa5c879790b2e116790c0ff4faa4941d..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitIncompAD7/InitIncompAD7_Device.cuh
+++ /dev/null
@@ -1,19 +0,0 @@
-#ifndef LB_INIT_INCOMP_AD7_H
-#define LB_INIT_INCOMP_AD7_H
-
-#include <DataTypes.h>
-#include <curand.h>
-
-__global__ void LB_Init_Incomp_AD_7(unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* geoD,
- real* Conc,
- real* ux,
- real* uy,
- real* uz,
- unsigned int size_Mat,
- real* DD7,
- bool EvenOrOdd);
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.cu
new file mode 100644
index 0000000000000000000000000000000000000000..b3f298988f64f4726743d8429f6df52da8eb8227
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.cu
@@ -0,0 +1,73 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#include "InitK18K20NavierStokesCompressible.h"
+
+#include "InitK18K20NavierStokesCompressible_Device.cuh"
+#include "Parameter/Parameter.h"
+#include <cuda_helper/CudaGrid.h>
+
+std::shared_ptr<PreProcessorStrategy> InitK18K20NavierStokesCompressible::getNewInstance(std::shared_ptr<Parameter> para)
+{
+ return std::shared_ptr<PreProcessorStrategy>(new InitK18K20NavierStokesCompressible(para));
+}
+
+void InitK18K20NavierStokesCompressible::init(int level)
+{
+ vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
+
+ InitK18K20NavierStokesCompressible_Device <<< grid.grid, grid.threads >>>(
+ para->getParD(level)->neighborX,
+ para->getParD(level)->neighborY,
+ para->getParD(level)->neighborZ,
+ para->getParD(level)->typeOfGridNode,
+ para->getParD(level)->rho,
+ para->getParD(level)->velocityX,
+ para->getParD(level)->velocityY,
+ para->getParD(level)->velocityZ,
+ para->getParD(level)->numberOfNodes,
+ para->getParD(level)->g6.g[0],
+ para->getParD(level)->isEvenTimestep);
+ getLastCudaError("InitK18K20NavierStokesCompressible_Device execution failed");
+}
+
+bool InitK18K20NavierStokesCompressible::checkParameter()
+{
+ return false;
+}
+
+InitK18K20NavierStokesCompressible::InitK18K20NavierStokesCompressible(std::shared_ptr<Parameter> para)
+{
+ this->para = para;
+}
+
+InitK18K20NavierStokesCompressible::InitK18K20NavierStokesCompressible()
+{
+}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.h
new file mode 100644
index 0000000000000000000000000000000000000000..403dc4f0cac889af75f40da0cf92f8b821692407
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.h
@@ -0,0 +1,53 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#ifndef InitK18K20NavierStokesCompressible_H
+#define InitK18K20NavierStokesCompressible_H
+
+#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h"
+
+#include <memory>
+
+class Parameter;
+
+class InitK18K20NavierStokesCompressible : public PreProcessorStrategy
+{
+public:
+ static std::shared_ptr<PreProcessorStrategy> getNewInstance(std::shared_ptr< Parameter> para);
+ void init(int level);
+ bool checkParameter();
+
+private:
+ InitK18K20NavierStokesCompressible();
+ InitK18K20NavierStokesCompressible(std::shared_ptr< Parameter> para);
+ std::shared_ptr< Parameter> para;
+};
+
+#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3_Device.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible_Device.cu
similarity index 54%
rename from src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3_Device.cu
rename to src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible_Device.cu
index 29050bd2807dc753051aaa78fc2d0a35fb4162ad..6beb4ad8574236857b16cebf035c2796312bdf82 100644
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3_Device.cu
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible_Device.cu
@@ -1,3 +1,33 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
#include "LBM/LB.h"
#include "lbm/constants/D3Q27.h"
#include <basics/constants/NumericConstants.h>
@@ -7,7 +37,7 @@ using namespace vf::lbm::dir;
#include "math.h"
-__global__ void LB_Init_F3(unsigned int* neighborX,
+__global__ void InitK18K20NavierStokesCompressible_Device(unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned int* geoD,
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible_Device.cuh
new file mode 100644
index 0000000000000000000000000000000000000000..90ecb070533a0bf40522bdf246f59921149caaa3
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible_Device.cuh
@@ -0,0 +1,50 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#ifndef InitK18K20NavierStokesCompressible_Device_H
+#define InitK18K20NavierStokesCompressible_Device_H
+
+#include <DataTypes.h>
+#include <curand.h>
+
+__global__ void InitK18K20NavierStokesCompressible_Device(
+ unsigned int* neighborX,
+ unsigned int* neighborY,
+ unsigned int* neighborZ,
+ unsigned int* geoD,
+ real* rho,
+ real* ux,
+ real* uy,
+ real* uz,
+ unsigned int size_Mat,
+ real* G6,
+ bool EvenOrOdd);
+
+#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible.cu
new file mode 100644
index 0000000000000000000000000000000000000000..995607fd3278bd79a9cf36d8d5d4666699be8ae5
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible.cu
@@ -0,0 +1,98 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#include "InitNavierStokesCompressible.h"
+
+#include "InitNavierStokesCompressible_Device.cuh"
+#include "Parameter/Parameter.h"
+#include <cuda_helper/CudaGrid.h>
+
+std::shared_ptr<PreProcessorStrategy> InitNavierStokesCompressible::getNewInstance(std::shared_ptr<Parameter> para)
+{
+ return std::shared_ptr<PreProcessorStrategy>(new InitNavierStokesCompressible(para));
+}
+
+void InitNavierStokesCompressible::init(int level)
+{
+ vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
+
+ if( ! para->getUseInitNeq() )
+ {
+ InitNavierStokesCompressible_Device <<< grid.grid, grid.threads >>> (
+ para->getParD(level)->neighborX,
+ para->getParD(level)->neighborY,
+ para->getParD(level)->neighborZ,
+ para->getParD(level)->typeOfGridNode,
+ para->getParD(level)->rho,
+ para->getParD(level)->velocityX,
+ para->getParD(level)->velocityY,
+ para->getParD(level)->velocityZ,
+ para->getParD(level)->numberOfNodes,
+ para->getParD(level)->distributions.f[0],
+ para->getParD(level)->isEvenTimestep);
+ getLastCudaError("InitNavierStokesCompressible_Device execution failed");
+ }
+ else
+ {
+ InitNavierStokesCompressibleNonEquilibrium_Device <<< grid.grid, grid.threads >>> (
+ para->getParD(level)->neighborX,
+ para->getParD(level)->neighborY,
+ para->getParD(level)->neighborZ,
+ para->getParD(level)->neighborInverse,
+ para->getParD(level)->typeOfGridNode,
+ para->getParD(level)->rho,
+ para->getParD(level)->velocityX,
+ para->getParD(level)->velocityY,
+ para->getParD(level)->velocityZ,
+ para->getParD(level)->numberOfNodes,
+ para->getParD(level)->distributions.f[0],
+ para->getParD(level)->omega,
+ para->getParD(level)->isEvenTimestep);
+ cudaDeviceSynchronize();
+ getLastCudaError("InitNavierStokesCompressibleNonEquilibrium_Device execution failed");
+ }
+
+
+
+}
+
+bool InitNavierStokesCompressible::checkParameter()
+{
+ return false;
+}
+
+InitNavierStokesCompressible::InitNavierStokesCompressible(std::shared_ptr<Parameter> para)
+{
+ this->para = para;
+}
+
+InitNavierStokesCompressible::InitNavierStokesCompressible()
+{
+}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible.h
new file mode 100644
index 0000000000000000000000000000000000000000..804b629a5d84291c57c4ed2a0b9ef9781cac9aae
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible.h
@@ -0,0 +1,53 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#ifndef InitNavierStokesCompressible_H
+#define InitNavierStokesCompressible_H
+
+#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h"
+
+#include <memory>
+
+class Parameter;
+
+class InitNavierStokesCompressible : public PreProcessorStrategy
+{
+public:
+ static std::shared_ptr<PreProcessorStrategy> getNewInstance(std::shared_ptr< Parameter> para);
+ void init(int level);
+ bool checkParameter();
+
+private:
+ InitNavierStokesCompressible();
+ InitNavierStokesCompressible(std::shared_ptr< Parameter> para);
+ std::shared_ptr< Parameter> para;
+};
+
+#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27_Device.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible_Device.cu
similarity index 91%
rename from src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27_Device.cu
rename to src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible_Device.cu
index f733282faf3791164bcab12354d61beea6659418..cb38eedebba41a15400c570f5520876433c90edc 100644
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27_Device.cu
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible_Device.cu
@@ -1,3 +1,33 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
#include "LBM/LB.h"
#include "lbm/constants/D3Q27.h"
#include <basics/constants/NumericConstants.h>
@@ -8,7 +38,8 @@ using namespace vf::lbm::dir;
#include <stdio.h>
-__global__ void LB_Init_Comp_SP_27(unsigned int* neighborX,
+__global__ void InitNavierStokesCompressible_Device(
+ unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned int* geoD,
@@ -179,7 +210,7 @@ __global__ void LB_Init_Comp_SP_27(unsigned int* neighborX,
////////////////////////////////////////////////////////////////////////////////
-__global__ void LB_Init_Comp_Neq_SP_27( unsigned int* neighborX,
+__global__ void InitNavierStokesCompressibleNonEquilibrium_Device( unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned int* neighborWSB,
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible_Device.cuh
new file mode 100644
index 0000000000000000000000000000000000000000..a454d7249be6948374bb614e090cc7d64b603bdd
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible_Device.cuh
@@ -0,0 +1,63 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#ifndef InitNavierStokesCompressible_Device_H
+#define InitNavierStokesCompressible_Device_H
+
+#include <DataTypes.h>
+#include <curand.h>
+
+__global__ void InitNavierStokesCompressible_Device(unsigned int* neighborX,
+ unsigned int* neighborY,
+ unsigned int* neighborZ,
+ unsigned int* geoD,
+ real* rho,
+ real* ux,
+ real* uy,
+ real* uz,
+ unsigned int size_Mat,
+ real* DD,
+ bool EvenOrOdd);
+
+__global__ void InitNavierStokesCompressibleNonEquilibrium_Device(unsigned int* neighborX,
+ unsigned int* neighborY,
+ unsigned int* neighborZ,
+ unsigned int* neighborWSB,
+ unsigned int* geoD,
+ real* rho,
+ real* ux,
+ real* uy,
+ real* uz,
+ unsigned int size_Mat,
+ real* DD,
+ real omega,
+ bool EvenOrOdd);
+
+#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible.cu
new file mode 100644
index 0000000000000000000000000000000000000000..b4c3f01cbb6ca3cca275aa28f55816e540490756
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible.cu
@@ -0,0 +1,73 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#include "InitNavierStokesIncompressible.h"
+
+#include "InitNavierStokesIncompressible_Device.cuh"
+#include "Parameter/Parameter.h"
+#include <cuda_helper/CudaGrid.h>
+
+std::shared_ptr<PreProcessorStrategy> InitNavierStokesIncompressible::getNewInstance(std::shared_ptr<Parameter> para)
+{
+ return std::shared_ptr<PreProcessorStrategy>(new InitNavierStokesIncompressible(para));
+}
+
+void InitNavierStokesIncompressible::init(int level)
+{
+ vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
+
+ InitNavierStokesIncompressible_Device<<<grid.grid, grid.threads>>>(
+ para->getParD(level)->neighborX,
+ para->getParD(level)->neighborY,
+ para->getParD(level)->neighborZ,
+ para->getParD(level)->typeOfGridNode,
+ para->getParD(level)->rho,
+ para->getParD(level)->velocityX,
+ para->getParD(level)->velocityY,
+ para->getParD(level)->velocityZ,
+ para->getParD(level)->numberOfNodes,
+ para->getParD(level)->distributions.f[0],
+ para->getParD(level)->isEvenTimestep);
+ getLastCudaError("LB_Init_SP_27 execution failed");
+}
+
+bool InitNavierStokesIncompressible::checkParameter()
+{
+ return false;
+}
+
+InitNavierStokesIncompressible::InitNavierStokesIncompressible(std::shared_ptr<Parameter> para)
+{
+ this->para = para;
+}
+
+InitNavierStokesIncompressible::InitNavierStokesIncompressible()
+{
+}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible.h
new file mode 100644
index 0000000000000000000000000000000000000000..1bf9f945e450945d660c2f996ada29b175c77886
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible.h
@@ -0,0 +1,53 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#ifndef INIT_SP27_H
+#define INIT_SP27_H
+
+#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h"
+
+#include <memory>
+
+class Parameter;
+
+class InitNavierStokesIncompressible : public PreProcessorStrategy
+{
+public:
+ static std::shared_ptr<PreProcessorStrategy> getNewInstance(std::shared_ptr< Parameter> para);
+ void init(int level);
+ bool checkParameter();
+
+private:
+ InitNavierStokesIncompressible();
+ InitNavierStokesIncompressible(std::shared_ptr< Parameter> para);
+ std::shared_ptr< Parameter> para;
+};
+
+#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27_Device.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible_Device.cu
similarity index 79%
rename from src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27_Device.cu
rename to src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible_Device.cu
index 8677e0f2eb3b1c5917d6550ef478c94c8400098b..3b49f44274717c41d10a9e724639f99395b6265b 100644
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27_Device.cu
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible_Device.cu
@@ -1,3 +1,33 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
#include "LBM/LB.h"
#include "lbm/constants/D3Q27.h"
#include <basics/constants/NumericConstants.h>
@@ -6,7 +36,7 @@ using namespace vf::basics::constant;
using namespace vf::lbm::dir;
#include "math.h"
-__global__ void LB_Init_SP_27(unsigned int* neighborX,
+__global__ void InitNavierStokesIncompressible_Device(unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned int* geoD,
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible_Device.cuh
new file mode 100644
index 0000000000000000000000000000000000000000..44015e0100e53ba70bee8b1d720869b93c2da813
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible_Device.cuh
@@ -0,0 +1,50 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#ifndef InitNavierStokesIncompressible_Device_H
+#define InitNavierStokesIncompressible_Device_H
+
+#include <DataTypes.h>
+#include <curand.h>
+
+__global__ void InitNavierStokesIncompressible_Device(
+ unsigned int* neighborX,
+ unsigned int* neighborY,
+ unsigned int* neighborZ,
+ unsigned int* geoD,
+ real* rho,
+ real* ux,
+ real* uy,
+ real* uz,
+ unsigned int size_Mat,
+ real* DD,
+ bool EvenOrOdd);
+
+#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27.cu
deleted file mode 100644
index 78994e997c34d6b47cf1b8b49dd0d70442f91588..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27.cu
+++ /dev/null
@@ -1,43 +0,0 @@
-#include "InitSP27.h"
-
-#include "InitSP27_Device.cuh"
-#include "Parameter/Parameter.h"
-#include <cuda_helper/CudaGrid.h>
-
-std::shared_ptr<PreProcessorStrategy> InitSP27::getNewInstance(std::shared_ptr<Parameter> para)
-{
- return std::shared_ptr<PreProcessorStrategy>(new InitSP27(para));
-}
-
-void InitSP27::init(int level)
-{
- vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
-
- LB_Init_SP_27 <<< grid.grid, grid.threads >>>(
- para->getParD(level)->neighborX,
- para->getParD(level)->neighborY,
- para->getParD(level)->neighborZ,
- para->getParD(level)->typeOfGridNode,
- para->getParD(level)->rho,
- para->getParD(level)->velocityX,
- para->getParD(level)->velocityY,
- para->getParD(level)->velocityZ,
- para->getParD(level)->numberOfNodes,
- para->getParD(level)->distributions.f[0],
- para->getParD(level)->isEvenTimestep);
- getLastCudaError("LB_Init_SP_27 execution failed");
-}
-
-bool InitSP27::checkParameter()
-{
- return false;
-}
-
-InitSP27::InitSP27(std::shared_ptr<Parameter> para)
-{
- this->para = para;
-}
-
-InitSP27::InitSP27()
-{
-}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27.h
deleted file mode 100644
index def7ccd175168ba6f463e6bc0b755db26e233212..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27.h
+++ /dev/null
@@ -1,23 +0,0 @@
-#ifndef INIT_SP27_H
-#define INIT_SP27_H
-
-#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h"
-
-#include <memory>
-
-class Parameter;
-
-class InitSP27 : public PreProcessorStrategy
-{
-public:
- static std::shared_ptr<PreProcessorStrategy> getNewInstance(std::shared_ptr< Parameter> para);
- void init(int level);
- bool checkParameter();
-
-private:
- InitSP27();
- InitSP27(std::shared_ptr< Parameter> para);
- std::shared_ptr< Parameter> para;
-};
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27_Device.cuh
deleted file mode 100644
index c00a191b4afd5fd9168485932cb162dfc5a05694..0000000000000000000000000000000000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27_Device.cuh
+++ /dev/null
@@ -1,19 +0,0 @@
-#ifndef LB_INIT_SP27_H
-#define LB_INIT_SP27_H
-
-#include <DataTypes.h>
-#include <curand.h>
-
-__global__ void LB_Init_SP_27(unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* geoD,
- real* rho,
- real* ux,
- real* uy,
- real* uz,
- unsigned int size_Mat,
- real* DD,
- bool EvenOrOdd);
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h
index 5a899eb040ea936d8acdbcd1e74a6c5607848c7a..27536a3470574f196378de7af158df42982464bb 100644
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h
@@ -1,3 +1,33 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
#ifndef PREPROCESSOR_STRATEGY_H
#define PREPROCESSOR_STRATEGY_H
diff --git a/src/gpu/core/PreProcessor/PreProcessorType.h b/src/gpu/core/PreProcessor/PreProcessorType.h
index 6bf54afc08fa99d0e77e3e57897cfa379392f585..20230c4b08817352f60cb5567dd99fdd19ce6438 100644
--- a/src/gpu/core/PreProcessor/PreProcessorType.h
+++ b/src/gpu/core/PreProcessor/PreProcessorType.h
@@ -1,14 +1,44 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
#ifndef PRE_PROCESSOR_TYPE_H
#define PRE_PROCESSOR_TYPE_H
enum PreProcessorType
{
- InitSP27,
- InitCompSP27,
- InitF3,
- InitIncompAD7,
- InitIncompAD27,
- InitCompAD7,
- InitCompAD27
+ InitNavierStokesIncompressible,
+ InitNavierStokesCompressible,
+ InitK18K20NavierStokesCompressible,
+ InitAdvectionDiffusionIncompressibleD3Q7,
+ InitAdvectionDiffusionIncompressible,
+ InitAdvectionDiffusionCompressibleD3Q7,
+ InitAdvectionDiffusionCompressible
};
#endif
\ No newline at end of file