Add lbm libary.

CMakeLists.txt

@@ -87,6 +87,7 @@ find_package(MPI REQUIRED)
 
 
 add_subdirectory(src/basics)
+add_subdirectory(src/lbm)
 
 #################################################################################
 #  VIRTUAL FLUIDS CPU / GPU

src/cpu/VirtualFluidsCore/CMakeLists.txt

@@ -25,7 +25,7 @@ if(BUILD_USE_OPENMP)
    list(APPEND VF_LIBRARIES OpenMP::OpenMP_CXX)
 endif()
 
-vf_add_library(BUILDTYPE static PUBLIC_LINK basics muparser MPI::MPI_CXX ${VF_LIBRARIES})
+vf_add_library(BUILDTYPE static PUBLIC_LINK basics muparser MPI::MPI_CXX ${VF_LIBRARIES} lbm)
 
 
 vf_get_library_name(library_name)

src/cpu/VirtualFluidsCore/LBM/D3Q27System.h

@@ -42,6 +42,8 @@
 #include "UbException.h"
 #include "UbMath.h"
 
+#include "lbm/CalcMac.h"
+
 //! \brief namespace for global system-functions
 namespace D3Q27System
 {
@@ -152,30 +154,24 @@ static const int ET_BNE = 12;
 /*=====================================================================*/
 static LBMReal getDensity(const LBMReal *const &f /*[27]*/)
 {
-    return ((f[TNE] + f[BSW]) + (f[TSE] + f[BNW])) + ((f[BSE] + f[TNW]) + (f[TSW] + f[BNE])) +
-           (((f[NE] + f[SW]) + (f[SE] + f[NW])) + ((f[TE] + f[BW]) + (f[BE] + f[TW])) +
-            ((f[BN] + f[TS]) + (f[TN] + f[BS]))) +
-           ((f[E] + f[W]) + (f[N] + f[S]) + (f[T] + f[B])) + f[REST];
+    return LBM::getDensity(f);
 }
 /*=====================================================================*/
 static LBMReal getPressure(const LBMReal *const &f /*[27]*/) { return REAL_CAST(UbMath::c1o3) * getDensity(f); }
 /*=====================================================================*/
 static LBMReal getIncompVelocityX1(const LBMReal *const &f /*[27]*/)
 {
-    return ((((f[TNE] - f[BSW]) + (f[TSE] - f[BNW])) + ((f[BSE] - f[TNW]) + (f[BNE] - f[TSW]))) +
-            (((f[BE] - f[TW]) + (f[TE] - f[BW])) + ((f[SE] - f[NW]) + (f[NE] - f[SW]))) + (f[E] - f[W]));
+    return LBM::getIncompVelocityX1(f);
 }
 /*=====================================================================*/
 static LBMReal getIncompVelocityX2(const LBMReal *const &f /*[27]*/)
 {
-    return ((((f[TNE] - f[BSW]) + (f[BNW] - f[TSE])) + ((f[TNW] - f[BSE]) + (f[BNE] - f[TSW]))) +
-            (((f[BN] - f[TS]) + (f[TN] - f[BS])) + ((f[NW] - f[SE]) + (f[NE] - f[SW]))) + (f[N] - f[S]));
+    return LBM::getIncompVelocityX2(f);
 }
 /*=====================================================================*/
 static LBMReal getIncompVelocityX3(const LBMReal *const &f /*[27]*/)
 {
-    return ((((f[TNE] - f[BSW]) + (f[TSE] - f[BNW])) + ((f[TNW] - f[BSE]) + (f[TSW] - f[BNE]))) +
-            (((f[TS] - f[BN]) + (f[TN] - f[BS])) + ((f[TW] - f[BE]) + (f[TE] - f[BW]))) + (f[T] - f[B]));
+    return LBM::getIncompVelocityX3(f);
 }
 /*=====================================================================*/
 static void calcDensity(const LBMReal *const &f /*[27]*/, LBMReal &rho)

src/gpu/VirtualFluids_GPU/CMakeLists.txt

@@ -5,10 +5,13 @@ if(MSVC)
     set(additional_libraries ws2_32 Traffic) # ws_32 throws an error on Phoenix
 endif()
 
-vf_add_library(PRIVATE_LINK ${additional_libraries} GridGenerator basics MPI::MPI_CXX)
+vf_add_library(PRIVATE_LINK ${additional_libraries} GridGenerator basics lbm MPI::MPI_CXX)
 
 linkBoost(COMPONENTS "serialization")
 
 #SET(TPN_WIN32 "/EHsc")
 #https://stackoverflow.com/questions/6832666/lnk2019-when-including-asio-headers-solution-generated-with-cmake
 #https://stackoverflow.com/questions/27442885/syntax-error-with-stdnumeric-limitsmax
+
+
+set_target_properties(VirtualFluids_GPU PROPERTIES CUDA_SEPARABLE_COMPILATION ON)

src/gpu/VirtualFluids_GPU/GPU/CalcMac27.cu

@@ -10,63 +10,7 @@
 #include "LBM/D3Q27.h"
 #include "Core/RealConstants.h"
 
-static const int E    = 0;
-static const int W    = 1;
-static const int N    = 2;
-static const int S    = 3;
-static const int T    = 4;
-static const int B    = 5;
-static const int NE   = 6;
-static const int SW   = 7;
-static const int SE   = 8;
-static const int NW   = 9;
-static const int TE   = 10;
-static const int BW   = 11;
-static const int BE   = 12;
-static const int TW   = 13;
-static const int TN   = 14;
-static const int BS   = 15;
-static const int BN   = 16;
-static const int TS   = 17;
-static const int TNE  = 18;
-static const int TNW  = 19;
-static const int TSE  = 20;
-static const int TSW  = 21;
-static const int BNE  = 22;
-static const int BNW  = 23;
-static const int BSE  = 24;
-static const int BSW  = 25;
-static const int REST = 26;
-
-__device__ real getDensity(const real *const &f /*[27]*/)
-{
-    return ((f[TNE] + f[BSW]) + (f[TSE] + f[BNW])) + ((f[BSE] + f[TNW]) + (f[TSW] + f[BNE])) +
-           (((f[NE] + f[SW]) + (f[SE] + f[NW])) + ((f[TE] + f[BW]) + (f[BE] + f[TW])) +
-            ((f[BN] + f[TS]) + (f[TN] + f[BS]))) +
-           ((f[E] + f[W]) + (f[N] + f[S]) + (f[T] + f[B])) + f[REST];
-}
-
-
-__device__ real getIncompVelocityX1(const real *const &f /*[27]*/)
-{
-    return ((((f[TNE] - f[BSW]) + (f[TSE] - f[BNW])) + ((f[BSE] - f[TNW]) + (f[BNE] - f[TSW]))) +
-            (((f[BE] - f[TW]) + (f[TE] - f[BW])) + ((f[SE] - f[NW]) + (f[NE] - f[SW]))) + (f[E] - f[W]));
-}
-
-
-__device__ real getIncompVelocityX2(const real *const &f /*[27]*/)
-{
-    return ((((f[TNE] - f[BSW]) + (f[BNW] - f[TSE])) + ((f[TNW] - f[BSE]) + (f[BNE] - f[TSW]))) +
-            (((f[BN] - f[TS]) + (f[TN] - f[BS])) + ((f[NW] - f[SE]) + (f[NE] - f[SW]))) + (f[N] - f[S]));
-}
-
-
-__device__ real getIncompVelocityX3(const real *const &f /*[27]*/)
-{
-    return ((((f[TNE] - f[BSW]) + (f[TSE] - f[BNW])) + ((f[TNW] - f[BSE]) + (f[TSW] - f[BNE]))) +
-            (((f[TS] - f[BN]) + (f[TN] - f[BS])) + ((f[TW] - f[BE]) + (f[TE] - f[BW]))) + (f[T] - f[B]));
-}
-
+#include "lbm/CalcMac.h"
 
 
 __device__ Distributions27 getDistributions(real* DD, unsigned int size_Mat, bool evenOrOdd)
@@ -275,10 +219,10 @@ extern "C" __global__ void LBCalcMac27( real* vxD,
    {
        const auto distribution = getDistribution(DD, size_Mat, evenOrOdd, k, neighborX, neighborY, neighborZ);
 
-       rhoD[k] = getDensity(distribution.f);
-       vxD[k] = getIncompVelocityX1(distribution.f);
-       vyD[k] = getIncompVelocityX2(distribution.f);
-       vzD[k] = getIncompVelocityX3(distribution.f);
+       rhoD[k] = LBM::getDensity(distribution.f);
+       vxD[k] = LBM::getIncompVelocityX1(distribution.f);
+       vyD[k] = LBM::getIncompVelocityX2(distribution.f);
+       vzD[k] = LBM::getIncompVelocityX3(distribution.f);
    }
 }
 

src/lbm/CMakeLists.txt

+project(lbm LANGUAGES CXX)
+
+if(BUILD_VF_GPU)
+    enable_language(CUDA)
+endif()
+
+vf_add_library(PUBLIC_LINK basics)
+
+if(BUILD_VF_GPU)
+set_target_properties(lbm PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
+endif()

src/lbm/CalcMac.cu

+#include "CalcMac.h"
+
+
+static const int E    = 0;
+static const int W    = 1;
+static const int N    = 2;
+static const int S    = 3;
+static const int T    = 4;
+static const int B    = 5;
+static const int NE   = 6;
+static const int SW   = 7;
+static const int SE   = 8;
+static const int NW   = 9;
+static const int TE   = 10;
+static const int BW   = 11;
+static const int BE   = 12;
+static const int TW   = 13;
+static const int TN   = 14;
+static const int BS   = 15;
+static const int BN   = 16;
+static const int TS   = 17;
+static const int TNE  = 18;
+static const int TNW  = 19;
+static const int TSE  = 20;
+static const int TSW  = 21;
+static const int BNE  = 22;
+static const int BNW  = 23;
+static const int BSE  = 24;
+static const int BSW  = 25;
+static const int REST = 26;
+
+__host__ __device__ real LBM::getDensity(const real *const &f /*[27]*/)
+{
+    return ((f[TNE] + f[BSW]) + (f[TSE] + f[BNW])) + ((f[BSE] + f[TNW]) + (f[TSW] + f[BNE])) +
+           (((f[NE] + f[SW]) + (f[SE] + f[NW])) + ((f[TE] + f[BW]) + (f[BE] + f[TW])) +
+            ((f[BN] + f[TS]) + (f[TN] + f[BS]))) +
+           ((f[E] + f[W]) + (f[N] + f[S]) + (f[T] + f[B])) + f[REST];
+}
+
+
+__host__ __device__ real LBM::getIncompVelocityX1(const real *const &f /*[27]*/)
+{
+    return ((((f[TNE] - f[BSW]) + (f[TSE] - f[BNW])) + ((f[BSE] - f[TNW]) + (f[BNE] - f[TSW]))) +
+            (((f[BE] - f[TW]) + (f[TE] - f[BW])) + ((f[SE] - f[NW]) + (f[NE] - f[SW]))) + (f[E] - f[W]));
+}
+
+
+__host__ __device__ real LBM::getIncompVelocityX2(const real *const &f /*[27]*/)
+{
+    return ((((f[TNE] - f[BSW]) + (f[BNW] - f[TSE])) + ((f[TNW] - f[BSE]) + (f[BNE] - f[TSW]))) +
+            (((f[BN] - f[TS]) + (f[TN] - f[BS])) + ((f[NW] - f[SE]) + (f[NE] - f[SW]))) + (f[N] - f[S]));
+}
+
+
+__host__ __device__ real LBM::getIncompVelocityX3(const real *const &f /*[27]*/)
+{
+    return ((((f[TNE] - f[BSW]) + (f[TSE] - f[BNW])) + ((f[TNW] - f[BSE]) + (f[TSW] - f[BNE]))) +
+            (((f[TS] - f[BN]) + (f[TN] - f[BS])) + ((f[TW] - f[BE]) + (f[TE] - f[BW]))) + (f[T] - f[B]));
+}
+

src/lbm/CalcMac.h

+#ifndef LBM_CALCMAC_H
+#define LBM_CALCMAC_H
+
+#include "Core/DataTypes.h"
+
+#ifdef __CUDACC__
+#pragma message ( "C Preprocessor got here!" )
+#include <cuda_runtime.h>
+#else
+#ifndef __host__
+#define __host__
+#endif
+#ifndef __device__
+#define __device__
+#endif
+#endif 
+
+class LBM
+{
+public:
+
+__host__ __device__ static real getDensity(const real *const &f /*[27]*/);
+
+__host__ __device__ static real getIncompVelocityX1(const real *const &f /*[27]*/);
+
+__host__ __device__ static real getIncompVelocityX2(const real *const &f /*[27]*/);
+
+__host__ __device__ static real getIncompVelocityX3(const real *const &f /*[27]*/);
+
+};
+
+#endif