diff --git a/src/cpu/VirtualFluidsCore/LBM/D3Q27System.cpp b/src/cpu/VirtualFluidsCore/LBM/D3Q27System.cpp index 616bc1ea318181dabe27e8641ef1075c279e54fa..fe635126152bfd94deeadf40c51c515220db342b 100644 --- a/src/cpu/VirtualFluidsCore/LBM/D3Q27System.cpp +++ b/src/cpu/VirtualFluidsCore/LBM/D3Q27System.cpp @@ -1,6 +1,6 @@ #include "D3Q27System.h" -#include "lbm/CalcMac.h" +#include "lbm/MacroscopicQuantities.h" namespace D3Q27System { @@ -25,22 +25,22 @@ const int INVDIR[] = { INV_E, INV_W, INV_N, INV_S, INV_T, INV_B, INV LBMReal getDensity(const LBMReal *const &f /*[27]*/) { - return LBM::getDensity(f); + return VF::LBM::getDensity(f); } LBMReal getIncompVelocityX1(const LBMReal *const &f /*[27]*/) { - return LBM::getIncompVelocityX1(f); + return VF::LBM::getIncompressibleVelocityX1(f); } LBMReal getIncompVelocityX2(const LBMReal *const &f /*[27]*/) { - return LBM::getIncompVelocityX2(f); + return VF::LBM::getIncompressibleVelocityX2(f); } LBMReal getIncompVelocityX3(const LBMReal *const &f /*[27]*/) { - return LBM::getIncompVelocityX3(f); + return VF::LBM::getIncompressibleVelocityX3(f); } diff --git a/src/gpu/VirtualFluids_GPU/GPU/CalcMac27.cu b/src/gpu/VirtualFluids_GPU/GPU/CalcMac27.cu index 8581dddad8769fbda787706bd5d06d74cca21a7e..2de824557371814dc7f3327c45af2838b8642275 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/CalcMac27.cu +++ b/src/gpu/VirtualFluids_GPU/GPU/CalcMac27.cu @@ -10,7 +10,7 @@ #include "LBM/D3Q27.h" #include "Core/RealConstants.h" -#include "lbm/CalcMac.h" +#include "lbm/MacroscopicQuantities.h" __device__ Distributions27 getDistributions(real* DD, unsigned int size_Mat, bool evenOrOdd) @@ -219,10 +219,10 @@ extern "C" __global__ void LBCalcMac27( real* vxD, { const auto distribution = getDistribution(DD, size_Mat, evenOrOdd, k, neighborX, neighborY, neighborZ); - rhoD[k] = LBM::getDensity(distribution.f); - vxD[k] = LBM::getIncompVelocityX1(distribution.f); - vyD[k] = LBM::getIncompVelocityX2(distribution.f); - vzD[k] = LBM::getIncompVelocityX3(distribution.f); + rhoD[k] = VF::LBM::getDensity(distribution.f); + vxD[k] = VF::LBM::getIncompressibleVelocityX1(distribution.f); + vyD[k] = VF::LBM::getIncompressibleVelocityX2(distribution.f); + vzD[k] = VF::LBM::getIncompressibleVelocityX3(distribution.f); } } diff --git a/src/lbm/CMakeLists.txt b/src/lbm/CMakeLists.txt index 991d1eeb0b93bb3b01fac8ba64e0cc125ae2f312..89458299ddc75e025a5545732aee7c23ea889411 100644 --- a/src/lbm/CMakeLists.txt +++ b/src/lbm/CMakeLists.txt @@ -6,4 +6,4 @@ if(BUILD_VF_GPU) add_subdirectory(cuda) endif() -vf_add_tests() \ No newline at end of file +vf_add_tests() diff --git a/src/lbm/CalcMac.cpp b/src/lbm/CalcMac.cpp deleted file mode 100644 index 3c120fd57cf41730e65b6d87c7f4c7e9ff3cb51e..0000000000000000000000000000000000000000 --- a/src/lbm/CalcMac.cpp +++ /dev/null @@ -1,60 +0,0 @@ -#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])); -} - diff --git a/src/lbm/CalcMac.h b/src/lbm/CalcMac.h deleted file mode 100644 index 941ab10c35a812312fa4912b77fd0b1ba6ade5bf..0000000000000000000000000000000000000000 --- a/src/lbm/CalcMac.h +++ /dev/null @@ -1,31 +0,0 @@ -#ifndef LBM_CALCMAC_H -#define LBM_CALCMAC_H - -#ifdef __CUDACC__ -#include <cuda_runtime.h> -#else -#ifndef __host__ -#define __host__ -#endif -#ifndef __device__ -#define __device__ -#endif -#endif - -#include "Core/DataTypes.h" - -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 diff --git a/src/lbm/CalcMacTests.cpp b/src/lbm/CalcMacTests.cpp deleted file mode 100644 index 0bead1fef1448106a7fe63e532fc7830c5fdbf24..0000000000000000000000000000000000000000 --- a/src/lbm/CalcMacTests.cpp +++ /dev/null @@ -1,12 +0,0 @@ -#include <gmock/gmock.h> - -#include "CalcMac.h" - -TEST(CalcMacTest, calcDensity) -{ - real f[27] = {1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1}; - - double density = LBM::getDensity(f); - - ASSERT_THAT(density, testing::DoubleEq(27)); -} \ No newline at end of file diff --git a/src/lbm/D3Q27.h b/src/lbm/D3Q27.h new file mode 100644 index 0000000000000000000000000000000000000000..b07c0e721ce845c2170709b283947aa2f6220e4c --- /dev/null +++ b/src/lbm/D3Q27.h @@ -0,0 +1,43 @@ +#ifndef LBM_D3Q27_H +#define LBM_D3Q27_H + +namespace VF +{ +namespace LBM +{ +namespace DIR +{ + +static constexpr int E = 0; +static constexpr int W = 1; +static constexpr int N = 2; +static constexpr int S = 3; +static constexpr int T = 4; +static constexpr int B = 5; +static constexpr int NE = 6; +static constexpr int SW = 7; +static constexpr int SE = 8; +static constexpr int NW = 9; +static constexpr int TE = 10; +static constexpr int BW = 11; +static constexpr int BE = 12; +static constexpr int TW = 13; +static constexpr int TN = 14; +static constexpr int BS = 15; +static constexpr int BN = 16; +static constexpr int TS = 17; +static constexpr int TNE = 18; +static constexpr int TNW = 19; +static constexpr int TSE = 20; +static constexpr int TSW = 21; +static constexpr int BNE = 22; +static constexpr int BNW = 23; +static constexpr int BSE = 24; +static constexpr int BSW = 25; +static constexpr int REST = 26; + +} +} +} + +#endif diff --git a/src/lbm/MacroscopicQuantities.cpp b/src/lbm/MacroscopicQuantities.cpp new file mode 100644 index 0000000000000000000000000000000000000000..84813be617e8c3d50207a4ad3e1cec2d2c0e6b12 --- /dev/null +++ b/src/lbm/MacroscopicQuantities.cpp @@ -0,0 +1,34 @@ +#include "MacroscopicQuantities.h" + +#include "D3Q27.h" + + +__host__ __device__ real VF::LBM::getDensity(const real *const &f /*[27]*/) +{ + return ((f[DIR::TNE] + f[DIR::BSW]) + (f[DIR::TSE] + f[DIR::BNW])) + ((f[DIR::BSE] + f[DIR::TNW]) + (f[DIR::TSW] + f[DIR::BNE])) + + (((f[DIR::NE] + f[DIR::SW]) + (f[DIR::SE] + f[DIR::NW])) + ((f[DIR::TE] + f[DIR::BW]) + (f[DIR::BE] + f[DIR::TW])) + + ((f[DIR::BN] + f[DIR::TS]) + (f[DIR::TN] + f[DIR::BS]))) + + ((f[DIR::E] + f[DIR::W]) + (f[DIR::N] + f[DIR::S]) + (f[DIR::T] + f[DIR::B])) + f[DIR::REST]; +} + + +real VF::LBM::getIncompressibleVelocityX1(const real *const &f /*[27]*/) +{ + return ((((f[DIR::TNE] - f[DIR::BSW]) + (f[DIR::TSE] - f[DIR::BNW])) + ((f[DIR::BSE] - f[DIR::TNW]) + (f[DIR::BNE] - f[DIR::TSW]))) + + (((f[DIR::BE] - f[DIR::TW]) + (f[DIR::TE] - f[DIR::BW])) + ((f[DIR::SE] - f[DIR::NW]) + (f[DIR::NE] - f[DIR::SW]))) + (f[DIR::E] - f[DIR::W])); +} + + +real VF::LBM::getIncompressibleVelocityX2(const real *const &f /*[27]*/) +{ + return ((((f[DIR::TNE] - f[DIR::BSW]) + (f[DIR::BNW] - f[DIR::TSE])) + ((f[DIR::TNW] - f[DIR::BSE]) + (f[DIR::BNE] - f[DIR::TSW]))) + + (((f[DIR::BN] - f[DIR::TS]) + (f[DIR::TN] - f[DIR::BS])) + ((f[DIR::NW] - f[DIR::SE]) + (f[DIR::NE] - f[DIR::SW]))) + (f[DIR::N] - f[DIR::S])); +} + + +real VF::LBM::getIncompressibleVelocityX3(const real *const &f /*[27]*/) +{ + return ((((f[DIR::TNE] - f[DIR::BSW]) + (f[DIR::TSE] - f[DIR::BNW])) + ((f[DIR::TNW] - f[DIR::BSE]) + (f[DIR::TSW] - f[DIR::BNE]))) + + (((f[DIR::TS] - f[DIR::BN]) + (f[DIR::TN] - f[DIR::BS])) + ((f[DIR::TW] - f[DIR::BE]) + (f[DIR::TE] - f[DIR::BW]))) + (f[DIR::T] - f[DIR::B])); +} + diff --git a/src/lbm/MacroscopicQuantities.h b/src/lbm/MacroscopicQuantities.h new file mode 100644 index 0000000000000000000000000000000000000000..df20a0537bf7eb6a6cc66b85cc13e203315d33a7 --- /dev/null +++ b/src/lbm/MacroscopicQuantities.h @@ -0,0 +1,33 @@ +#ifndef LBM_CALCMAC_H +#define LBM_CALCMAC_H + +#ifdef __CUDACC__ +#include <cuda_runtime.h> +#else +#ifndef __host__ +#define __host__ +#endif +#ifndef __device__ +#define __device__ +#endif +#endif + +#include "Core/DataTypes.h" + +namespace VF +{ +namespace LBM +{ + +__host__ __device__ real getDensity(const real *const &f /*[27]*/); + +__host__ __device__ real getIncompressibleVelocityX1(const real *const &f /*[27]*/); + +__host__ __device__ real getIncompressibleVelocityX2(const real *const &f /*[27]*/); + +__host__ __device__ real getIncompressibleVelocityX3(const real *const &f /*[27]*/); + +} +} + +#endif diff --git a/src/lbm/MacroscopicQuantitiesTests.cpp b/src/lbm/MacroscopicQuantitiesTests.cpp new file mode 100644 index 0000000000000000000000000000000000000000..fadafde8a579984fa78a8db4c01872deb929f34f --- /dev/null +++ b/src/lbm/MacroscopicQuantitiesTests.cpp @@ -0,0 +1,63 @@ +#include <gmock/gmock.h> + +#include "MacroscopicQuantities.h" +#include "D3Q27.h" + + +/* +* given distributions, which are all one. +*/ +real f[27] = {1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1}; + + + +TEST(MacroscopicQuantitiesTest, check_density) +{ + const double density = VF::LBM::getDensity(f); + + const double expected_density = 27.; + ASSERT_THAT(density, testing::DoubleEq(expected_density)); +} + +TEST(MacroscopicQuantitiesTest, whenFsAreEqual_velocityInEachDirectionShouldBeZero) +{ + const double velocityX1 = VF::LBM::getIncompressibleVelocityX1(f); + const double velocityX2 = VF::LBM::getIncompressibleVelocityX2(f); + const double velocityX3 = VF::LBM::getIncompressibleVelocityX3(f); + + const double expected_velocity = 0.; + EXPECT_THAT(velocityX1, testing::DoubleEq(expected_velocity)); + EXPECT_THAT(velocityX2, testing::DoubleEq(expected_velocity)); + EXPECT_THAT(velocityX3, testing::DoubleEq(expected_velocity)); +} + +TEST(MacroscopicQuantitiesTest, givenAllFsAreOne_when_Eis2_velocityInXShouldBeOne) +{ + f[VF::LBM::DIR::E] = 2.; + + const double velocity = VF::LBM::getIncompressibleVelocityX1(f); + + const double expected_velocity = 1.; + ASSERT_THAT(velocity, testing::DoubleEq(expected_velocity)); +} + +TEST(MacroscopicQuantitiesTest, givenAllFsAreOne_when_Nis2_velocityInX2ShouldBeOne) +{ + f[VF::LBM::DIR::N] = 2.; + + const double velocity = VF::LBM::getIncompressibleVelocityX2(f); + + const double expected_velocity = 1.; + ASSERT_THAT(velocity, testing::DoubleEq(expected_velocity)); +} + + +TEST(MacroscopicQuantitiesTest, givenAllFsAreOne_when_Tis2_velocityInX3ShouldBeOne) +{ + f[VF::LBM::DIR::T] = 2.; + + const double velocity = VF::LBM::getIncompressibleVelocityX3(f); + + const double expected_velocity = 1.; + ASSERT_THAT(velocity, testing::DoubleEq(expected_velocity)); +} diff --git a/src/lbm/cuda/CMakeLists.txt b/src/lbm/cuda/CMakeLists.txt index 4e78bce5857caafa0a1b7819e698ac8599652092..f024461cf3d17d15f511b82b5518c979675acf28 100644 --- a/src/lbm/cuda/CMakeLists.txt +++ b/src/lbm/cuda/CMakeLists.txt @@ -3,6 +3,8 @@ project(lbmCuda LANGUAGES CUDA CXX) vf_add_library(NAME lbmCuda BUILDTYPE static PUBLIC_LINK basics FOLDER ../../lbm) + set_target_properties(lbmCuda PROPERTIES CUDA_SEPARABLE_COMPILATION ON) -set_source_files_properties(../CalcMac.cpp PROPERTIES LANGUAGE CUDA) + +set_source_files_properties(../MacroscopicQuantities.cpp PROPERTIES LANGUAGE CUDA) diff --git a/src/lbmApp/main.cpp b/src/lbmApp/main.cpp index 2b3405b394e4087e1044a7f9fc0592391d17ae60..ce0d6f716aa7beaf6e25f30ded74eb502301ce9d 100644 --- a/src/lbmApp/main.cpp +++ b/src/lbmApp/main.cpp @@ -1,7 +1,7 @@ #include <iostream> #include <stdio.h> -#include "lbm/CalcMac.h" +#include "lbm/MacroscopicQuantities.h" //#include <cuda_runtime.h> /* __global__ */ void test() @@ -9,7 +9,7 @@ printf("Hello World from GPU!\n"); real f[27] = {1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1}; - double density = LBM::getDensity(f); + double density = VF::LBM::getDensity(f); printf("Hello density: %f \n", density); }