diff --git a/apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp b/apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp
index 9303cd37cce1607c82c7536ee4cd350f09ab2bd7..57349a82072dbf8e64c17a1f678224a9a6e177cb 100644
--- a/apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp
+++ b/apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp
@@ -44,6 +44,7 @@
#include "VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.h"
#include "VirtualFluids_GPU/Parameter/Parameter.h"
#include "VirtualFluids_GPU/Output/FileWriter.h"
+#include "VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.h"
#include "VirtualFluids_GPU/GPU/CudaMemoryManager.h"
@@ -156,6 +157,7 @@ void multipleLevel(const std::string& configPath)
config.load(configPath);
SPtr<Parameter> para = std::make_shared<Parameter>(config, communicator.getNummberOfProcess(), communicator.getPID());
+ BoundaryConditionFactory bcFactory = BoundaryConditionFactory();
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
@@ -208,6 +210,7 @@ void multipleLevel(const std::string& configPath)
gridBuilder->setNoSlipBoundaryCondition(SideType::MY);
gridBuilder->setVelocityBoundaryCondition(SideType::PZ, vx, vx, 0.0);
gridBuilder->setNoSlipBoundaryCondition(SideType::MZ);
+ bcFactory.setVelocityBoundaryCondition(BoundaryConditionFactory::VelocityBC::VelocitySimpleBounceBackCompressible);
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
@@ -219,7 +222,7 @@ void multipleLevel(const std::string& configPath)
auto gridGenerator = GridProvider::makeGridGenerator(gridBuilder, para, cudaMemoryManager, communicator);
- Simulation sim(para, cudaMemoryManager, communicator, *gridGenerator);
+ Simulation sim(para, cudaMemoryManager, communicator, *gridGenerator, &bcFactory);
sim.run();
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.cpp b/src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.cpp
new file mode 100644
index 0000000000000000000000000000000000000000..ea17e1f4fce6a9e24ffb60e1a379fbbab1ef9ed2
--- /dev/null
+++ b/src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.cpp
@@ -0,0 +1,33 @@
+#include "BoundaryConditionFactory.h"
+#include "GPU/GPU_Interface.h"
+#include "Parameter/Parameter.h"
+
+void BoundaryConditionFactory::setVelocityBoundaryCondition(VelocityBC boundaryConditionType)
+{
+ this->velocityBoundaryCondition = boundaryConditionType;
+}
+
+boundaryCondition BoundaryConditionFactory::getVelocityBoundaryConditionPost() const
+{
+ switch (this->velocityBoundaryCondition) {
+ case VelocityBC::VelocitySimpleBounceBackCompressible:
+ // VelocitySimpleBounceBackCompressible = plain bounce back velocity boundary condition
+ return QVelDevicePlainBB27;
+ break;
+ case VelocityBC::VelocityIncompressible:
+ // VelocityIncompressible = interpolated velocity boundary condition, based on subgrid distances
+ return QVelDev27;
+ break;
+ case VelocityBC::VelocityCompressible:
+ // VelocityCompressible = interpolated velocity boundary condition, based on subgrid distances
+ return QVelDevComp27;
+ break;
+ case VelocityBC::VelocityAndPressureCompressible:
+ // VelocityAndPressureCompressible = interpolated velocity boundary condition, based on subgrid distances.
+ // Also sets the pressure to the bulk pressure.
+ return QVelDevCompZeroPress27;
+ break;
+ default:
+ return nullptr;
+ }
+}
\ No newline at end of file
diff --git a/src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.h b/src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.h
new file mode 100644
index 0000000000000000000000000000000000000000..9db2857511df85aee03069d6f7fa87f353902bb9
--- /dev/null
+++ b/src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.h
@@ -0,0 +1,37 @@
+#ifndef BC_FACTORY
+#define BC_FACTORY
+
+#include <functional>
+#include <map>
+#include <string>
+
+#include "LBM/LB.h"
+
+class LBMSimulationParameter;
+
+using boundaryCondition = std::function<void(LBMSimulationParameter *, QforBoundaryConditions *)>;
+
+class BoundaryConditionFactory
+{
+public:
+ enum class VelocityBC {
+ VelocitySimpleBounceBackCompressible,
+ VelocityIncompressible,
+ VelocityCompressible,
+ VelocityAndPressureCompressible
+ };
+ // enum class OutflowBoundaryConditon {}; // TODO:
+ // https://git.rz.tu-bs.de/m.schoenherr/VirtualFluids_dev/-/issues/16
+
+ void setVelocityBoundaryCondition(const VelocityBC boundaryConditionType);
+ // void setOutflowBoundaryCondition(...); // TODO:
+ // https://git.rz.tu-bs.de/m.schoenherr/VirtualFluids_dev/-/issues/16
+
+ boundaryCondition getVelocityBoundaryConditionPost() const;
+
+private:
+ VelocityBC velocityBoundaryCondition;
+ // OutflowBoundaryConditon outflowBC // TODO: https://git.rz.tu-bs.de/m.schoenherr/VirtualFluids_dev/-/issues/16
+};
+
+#endif
diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
index ace81490e6880d688690f647f50172c1a94b5bfb..426bb14d6f3df1752fdda0a6ca9c8662457ffd36 100644
--- a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
+++ b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
@@ -493,17 +493,16 @@ void calcTurbulentViscosity(Parameter* para, int level)
UpdateGrid27::UpdateGrid27(SPtr<Parameter> para, vf::gpu::Communicator &comm, SPtr<CudaMemoryManager> cudaManager,
- std::vector<std::shared_ptr<PorousMedia>> &pm, std::vector<SPtr<Kernel>> &kernels)
+ std::vector<std::shared_ptr<PorousMedia>> &pm, std::vector<SPtr<Kernel>> &kernels , BoundaryConditionFactory* bcFactory)
: para(para), comm(comm), cudaMemoryManager(cudaManager), pm(pm), kernels(kernels)
{
chooseFunctionForCollisionAndExchange();
chooseFunctionForRefinementAndExchange();
- this->lbKernelManager = LBKernelManager::make(para);
+ this->lbKernelManager = std::make_shared<LBKernelManager>(para, bcFactory);
this->adKernelManager = std::make_shared<ADKernelManager>(para);
this->gridScalingKernelManager = std::make_shared<GridScalingKernelManager>(para);
}
-
void UpdateGrid27::chooseFunctionForCollisionAndExchange()
{
std::cout << "Function used for collisionAndExchange: ";
diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h
index 5ed5347ec72c3dc30d38cc98acfed229877b8d2c..db82166a78a88aab0c4031aaac08ad2cb9d95325 100644
--- a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h
+++ b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h
@@ -12,15 +12,15 @@ class LBKernelManager;
class ADKernelManager;
class GridScalingKernelManager;
class Kernel;
+class BoundaryConditionFactory;
class UpdateGrid27
{
public:
UpdateGrid27(SPtr<Parameter> para, vf::gpu::Communicator &comm, SPtr<CudaMemoryManager> cudaManager,
- std::vector<std::shared_ptr<PorousMedia>> &pm, std::vector<SPtr<Kernel>> &kernels);
+ std::vector<std::shared_ptr<PorousMedia>> &pm, std::vector<SPtr<Kernel>> &kernels, BoundaryConditionFactory* bcFactory);
void updateGrid(int level, unsigned int t);
-
private:
void postCollisionBC(int level);
void preCollisionBC(int level, unsigned int t);
diff --git a/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h b/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h
index f85640484c43af0e155dc56348d920835c6016e5..3d232dcb8d2c1259eca51d5279b32b346e21470c 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h
+++ b/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h
@@ -5,16 +5,17 @@
// |___/_/_/ \__/\__,_/\__,_/_/_/ /_/\__,_/_/\__,_/____/ \____/_/ \_____/
//
//////////////////////////////////////////////////////////////////////////
-//random numbers
-#include <curand.h>
+#ifndef GPU_INTERFACE_H
+#define GPU_INTERFACE_H
+
+#include <curand.h> //random numbers
#include <curand_kernel.h>
#include <cuda_runtime.h>
#include <DataTypes.h>
#include "LBM/LB.h"
-#ifndef GPU_INTERFACE_H
-#define GPU_INTERFACE_H
+class LBMSimulationParameter;
//////////////////////////////////////////////////////////////////////////
//Kernel
@@ -723,20 +724,7 @@ extern "C" void QDevCompHighNu27(unsigned int numberOfThreads,
unsigned int size_Mat,
bool isEvenTimestep);
-extern "C" void QVelDevicePlainBB27(unsigned int numberOfThreads,
- real* vx,
- real* vy,
- real* vz,
- real* DD,
- int* k_Q,
- real* QQ,
- unsigned int numberOfBCnodes,
- real om1,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int size_Mat,
- bool isEvenTimestep);
+extern "C" void QVelDevicePlainBB27(LBMSimulationParameter* parameterDevice, QforBoundaryConditions* boundaryCondition);
extern "C" void QVelDeviceCouette27(unsigned int numberOfThreads,
real* vx,
@@ -775,22 +763,7 @@ extern "C" void QVelDevice1h27( unsigned int numberOfThreads,
unsigned int size_Mat,
bool isEvenTimestep);
-extern "C" void QVelDev27(unsigned int numberOfThreads,
- int nx,
- int ny,
- real* vx,
- real* vy,
- real* vz,
- real* DD,
- int* k_Q,
- real* QQ,
- unsigned int numberOfBCnodes,
- real om1,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int size_Mat,
- bool isEvenTimestep);
+extern "C" void QVelDev27(LBMSimulationParameter* parameterDevice, QforBoundaryConditions* boundaryCondition);
extern "C" void QVelDevCompPlusSlip27(unsigned int numberOfThreads,
real* vx,
@@ -807,20 +780,7 @@ extern "C" void QVelDevCompPlusSlip27(unsigned int numberOfThreads,
unsigned int size_Mat,
bool isEvenTimestep);
-extern "C" void QVelDevComp27(unsigned int numberOfThreads,
- real* velocityX,
- real* velocityY,
- real* velocityZ,
- real* distribution,
- int* subgridDistanceIndices,
- real* subgridDistances,
- unsigned int numberOfBCnodes,
- real omega,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int numberOfLBnodes,
- bool isEvenTimestep);
+extern "C" void QVelDevComp27(LBMSimulationParameter* parameterDevice, QforBoundaryConditions* boundaryCondition);
extern "C" void QVelDevCompThinWalls27(unsigned int numberOfThreads,
real* vx,
@@ -839,21 +799,7 @@ extern "C" void QVelDevCompThinWalls27(unsigned int numberOfThreads,
unsigned int size_Mat,
bool isEvenTimestep);
-extern "C" void QVelDevCompZeroPress27(
- unsigned int numberOfThreads,
- real* velocityX,
- real* velocityY,
- real* velocityZ,
- real* distribution,
- int* subgridDistanceIndices,
- real* subgridDistances,
- unsigned int numberOfBCnodes,
- real omega,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int numberOfLBnodes,
- bool isEvenTimestep);
+extern "C" void QVelDevCompZeroPress27(LBMSimulationParameter* parameterDevice, QforBoundaryConditions* boundaryCondition);
extern "C" void QVelDevIncompHighNu27( unsigned int numberOfThreads,
real* vx,
diff --git a/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu b/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu
index 42e5bd12586a49905b0cad494ee7400136fb112a..74e091bbe2598e5a9fb8f70d4be159947d10fc5f 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu
+++ b/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu
@@ -15,6 +15,8 @@
// includes, kernels
#include "GPU/GPU_Kernels.cuh"
+
+#include "Parameter/Parameter.h"
//////////////////////////////////////////////////////////////////////////
extern "C" void KernelCas27( unsigned int grid_nx,
unsigned int grid_ny,
@@ -2915,50 +2917,25 @@ extern "C" void QDevCompHighNu27( unsigned int numberOfThreads,
getLastCudaError("QDevice27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
-extern "C" void QVelDevicePlainBB27(unsigned int numberOfThreads,
- real* vx,
- real* vy,
- real* vz,
- real* DD,
- int* k_Q,
- real* QQ,
- unsigned int numberOfBCnodes,
- real om1,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int size_Mat,
- bool isEvenTimestep)
+extern "C" void QVelDevicePlainBB27(LBMSimulationParameter* parameterDevice, QforBoundaryConditions* boundaryCondition)
{
- int Grid = (numberOfBCnodes / numberOfThreads)+1;
- int Grid1, Grid2;
- if (Grid>512)
- {
- Grid1 = 512;
- Grid2 = (Grid/Grid1)+1;
- }
- else
- {
- Grid1 = 1;
- Grid2 = Grid;
- }
- dim3 gridQ(Grid1, Grid2);
- dim3 threads(numberOfThreads, 1, 1 );
+ dim3 grid = vf::cuda::getCudaGrid( parameterDevice->numberofthreads, boundaryCondition->numberOfBCnodes);
+ dim3 threads(parameterDevice->numberofthreads, 1, 1 );
- QVelDevPlainBB27<<< gridQ, threads >>> ( vx,
- vy,
- vz,
- DD,
- k_Q,
- QQ,
- numberOfBCnodes,
- om1,
- neighborX,
- neighborY,
- neighborZ,
- size_Mat,
- isEvenTimestep);
- getLastCudaError("QVelDevicePlainBB27 execution failed");
+ QVelDevPlainBB27<<< grid, threads >>> (
+ boundaryCondition->Vx,
+ boundaryCondition->Vy,
+ boundaryCondition->Vz,
+ parameterDevice->distributions.f[0],
+ boundaryCondition->k,
+ boundaryCondition->q27[0],
+ boundaryCondition->numberOfBCnodes,
+ parameterDevice->neighborX,
+ parameterDevice->neighborY,
+ parameterDevice->neighborZ,
+ parameterDevice->numberOfNodes,
+ parameterDevice->isEvenTimestep);
+ getLastCudaError("QVelDevicePlainBB27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
extern "C" void QVelDeviceCouette27(unsigned int numberOfThreads,
@@ -3067,53 +3044,27 @@ extern "C" void QVelDevice1h27( unsigned int numberOfThreads,
getLastCudaError("QVelDevice27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
-extern "C" void QVelDev27(unsigned int numberOfThreads,
- int nx,
- int ny,
- real* vx,
- real* vy,
- real* vz,
- real* DD,
- int* k_Q,
- real* QQ,
- unsigned int numberOfBCnodes,
- real om1,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int size_Mat,
- bool isEvenTimestep)
-{
- int Grid = (numberOfBCnodes / numberOfThreads)+1;
- int Grid1, Grid2;
- if (Grid>512)
- {
- Grid1 = 512;
- Grid2 = (Grid/Grid1)+1;
- }
- else
- {
- Grid1 = 1;
- Grid2 = Grid;
- }
- dim3 gridQ(Grid1, Grid2);
- dim3 threads(numberOfThreads, 1, 1 );
-
- QVelDevice27<<< gridQ, threads >>> (nx,
- ny,
- vx,
- vy,
- vz,
- DD,
- k_Q,
- QQ,
- numberOfBCnodes,
- om1,
- neighborX,
- neighborY,
- neighborZ,
- size_Mat,
- isEvenTimestep);
+extern "C" void QVelDev27(LBMSimulationParameter* parameterDevice, QforBoundaryConditions* boundaryCondition)
+{
+ dim3 grid = vf::cuda::getCudaGrid( parameterDevice->numberofthreads, boundaryCondition->numberOfBCnodes);
+ dim3 threads(parameterDevice->numberofthreads, 1, 1 );
+
+ QVelDevice27<<< grid, threads >>> (
+ parameterDevice->nx,
+ parameterDevice->ny,
+ boundaryCondition->Vx,
+ boundaryCondition->Vy,
+ boundaryCondition->Vz,
+ parameterDevice->distributions.f[0],
+ boundaryCondition->k,
+ boundaryCondition->q27[0],
+ boundaryCondition->numberOfBCnodes,
+ parameterDevice->omega,
+ parameterDevice->neighborX,
+ parameterDevice->neighborY,
+ parameterDevice->neighborZ,
+ parameterDevice->numberOfNodes,
+ parameterDevice->isEvenTimestep);
getLastCudaError("QVelDevice27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
@@ -3164,39 +3115,26 @@ extern "C" void QVelDevCompPlusSlip27(unsigned int numberOfThreads,
getLastCudaError("QVelDeviceCompPlusSlip27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
-extern "C" void QVelDevComp27(unsigned int numberOfThreads,
- real* velocityX,
- real* velocityY,
- real* velocityZ,
- real* distribution,
- int* subgridDistanceIndices,
- real* subgridDistances,
- unsigned int numberOfBCnodes,
- real omega,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int numberOfLBnodes,
- bool isEvenTimestep)
+extern "C" void QVelDevComp27(LBMSimulationParameter* parameterDevice, QforBoundaryConditions* boundaryCondition)
{
- dim3 grid = vf::cuda::getCudaGrid(numberOfThreads, numberOfBCnodes);
- dim3 threads(numberOfThreads, 1, 1 );
+ dim3 grid = vf::cuda::getCudaGrid(parameterDevice->numberofthreads, boundaryCondition->numberOfBCnodes);
+ dim3 threads(parameterDevice->numberofthreads, 1, 1 );
- QVelDeviceComp27<<< grid, threads >>> (
- velocityX,
- velocityY,
- velocityZ,
- distribution,
- subgridDistanceIndices,
- subgridDistances,
- numberOfBCnodes,
- omega,
- neighborX,
- neighborY,
- neighborZ,
- numberOfLBnodes,
- isEvenTimestep);
- getLastCudaError("QVelDeviceComp27 execution failed");
+ QVelDeviceComp27<<< grid, threads >>> (
+ boundaryCondition->Vx,
+ boundaryCondition->Vy,
+ boundaryCondition->Vz,
+ parameterDevice->distributions.f[0],
+ boundaryCondition->k,
+ boundaryCondition->q27[0],
+ boundaryCondition->numberOfBCnodes,
+ parameterDevice->omega,
+ parameterDevice->neighborX,
+ parameterDevice->neighborY,
+ parameterDevice->neighborZ,
+ parameterDevice->numberOfNodes,
+ parameterDevice->isEvenTimestep);
+ getLastCudaError("QVelDeviceComp27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
extern "C" void QVelDevCompThinWalls27(unsigned int numberOfThreads,
@@ -3261,39 +3199,25 @@ extern "C" void QVelDevCompThinWalls27(unsigned int numberOfThreads,
getLastCudaError("QThinWallsPartTwo27 execution failed");
}
-extern "C" void QVelDevCompZeroPress27(
- unsigned int numberOfThreads,
- real* velocityX,
- real* velocityY,
- real* velocityZ,
- real* distribution,
- int* subgridDistanceIndices,
- real* subgridDistances,
- unsigned int numberOfBCnodes,
- real omega,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int numberOfLBnodes,
- bool isEvenTimestep)
+extern "C" void QVelDevCompZeroPress27(LBMSimulationParameter* parameterDevice, QforBoundaryConditions* boundaryCondition)
{
- dim3 grid = vf::cuda::getCudaGrid(numberOfThreads, numberOfBCnodes);
- dim3 threads(numberOfThreads, 1, 1 );
+ dim3 grid = vf::cuda::getCudaGrid( parameterDevice->numberofthreads, boundaryCondition->numberOfBCnodes);
+ dim3 threads(parameterDevice->numberofthreads, 1, 1 );
QVelDeviceCompZeroPress27<<< grid, threads >>> (
- velocityX,
- velocityY,
- velocityZ,
- distribution,
- subgridDistanceIndices,
- subgridDistances,
- numberOfBCnodes,
- omega,
- neighborX,
- neighborY,
- neighborZ,
- numberOfLBnodes,
- isEvenTimestep);
+ boundaryCondition->Vx,
+ boundaryCondition->Vy,
+ boundaryCondition->Vz,
+ parameterDevice->distributions.f[0],
+ boundaryCondition->k,
+ boundaryCondition->q27[0],
+ boundaryCondition->numberOfBCnodes,
+ parameterDevice->omega,
+ parameterDevice->neighborX,
+ parameterDevice->neighborY,
+ parameterDevice->neighborZ,
+ parameterDevice->numberOfNodes,
+ parameterDevice->isEvenTimestep);
getLastCudaError("QVelDeviceCompZeroPress27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.cpp b/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.cpp
index 0c7a88e3c351f11754ae9632ec4766b53190179a..375e19c7113e13a8b71c473b3753354508ec4725 100644
--- a/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.cpp
+++ b/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.cpp
@@ -34,13 +34,18 @@
#include <helper_cuda.h>
#include "LBKernelManager.h"
+#include "Parameter/Parameter.h"
#include "GPU/GPU_Interface.h"
-#include "Parameter/Parameter.h"
#include "Calculation/DragLift.h"
#include "Calculation/Cp.h"
+#include "BoundaryConditions/BoundaryConditionFactory.h"
+LBKernelManager::LBKernelManager(SPtr<Parameter> parameter, BoundaryConditionFactory* bcFactory): para(parameter)
+{
+ this->velocityBoundaryConditionPost = bcFactory->getVelocityBoundaryConditionPost();
+}
-void LBKernelManager::runLBMKernel(int level)
+void LBKernelManager::runLBMKernel(const int level) const
{
// if (para->getIsADcalculationOn()) {
// CumulantK17LBMDeviceKernelAD(
@@ -70,7 +75,7 @@ void LBKernelManager::runLBMKernel(int level)
// }
}
-void LBKernelManager::runVelocityBCKernelPre(int level)
+void LBKernelManager::runVelocityBCKernelPre(const int level) const
{
if (para->getParD(level)->velocityBC.numberOfBCnodes > 0)
{
@@ -124,75 +129,13 @@ void LBKernelManager::runVelocityBCKernelPre(int level)
}
}
-void LBKernelManager::runVelocityBCKernelPost(int level)
+void LBKernelManager::runVelocityBCKernelPost(const int level) const
{
if (para->getParD(level)->velocityBC.numberOfBCnodes > 0)
{
- // QVelDevicePlainBB27(
- // para->getParD(level)->numberofthreads,
- // para->getParD(level)->velocityBC.Vx,
- // para->getParD(level)->velocityBC.Vy,
- // para->getParD(level)->velocityBC.Vz,
- // para->getParD(level)->distributions.f[0],
- // para->getParD(level)->velocityBC.k,
- // para->getParD(level)->velocityBC.q27[0],
- // para->getParD(level)->velocityBC.numberOfBCnodes,
- // para->getParD(level)->velocityBC.kArray,
- // para->getParD(level)->neighborX,
- // para->getParD(level)->neighborY,
- // para->getParD(level)->neighborZ,
- // para->getParD(level)->numberOfNodes,
- // para->getParD(level)->isEvenTimestep);
+ velocityBoundaryConditionPost(para->getParD(level).get(), &(para->getParD(level)->velocityBC));
- // QVelDev27(
- // para->getParD(level)->numberofthreads,
- // para->getParD(level)->nx,
- // para->getParD(level)->ny,
- // para->getParD(level)->velocityBC.Vx,
- // para->getParD(level)->velocityBC.Vy,
- // para->getParD(level)->velocityBC.Vz,
- // para->getParD(level)->distributions.f[0],
- // para->getParD(level)->velocityBC.k,
- // para->getParD(level)->velocityBC.q27[0],
- // para->getParD(level)->velocityBC.numberOfBCnodes,
- // para->getParD(level)->omega,
- // para->getParD(level)->neighborX,
- // para->getParD(level)->neighborY,
- // para->getParD(level)->neighborZ,
- // para->getParD(level)->size_Mat,
- // para->getParD(level)->isEvenTimestep);
-
- // QVelDevComp27(
- // para->getParD(level)->numberofthreads,
- // para->getParD(level)->velocityBC.Vx,
- // para->getParD(level)->velocityBC.Vy,
- // para->getParD(level)->velocityBC.Vz,
- // para->getParD(level)->distributions.f[0],
- // para->getParD(level)->velocityBC.k,
- // para->getParD(level)->velocityBC.q27[0],
- // para->getParD(level)->velocityBC.numberOfBCnodes,
- // para->getParD(level)->omega,
- // para->getParD(level)->neighborX,
- // para->getParD(level)->neighborY,
- // para->getParD(level)->neighborZ,
- // para->getParD(level)->size_Mat,
- // para->getParD(level)->isEvenTimestep);
-
- QVelDevCompZeroPress27(
- para->getParD(level)->numberofthreads,
- para->getParD(level)->velocityBC.Vx,
- para->getParD(level)->velocityBC.Vy,
- para->getParD(level)->velocityBC.Vz,
- para->getParD(level)->distributions.f[0],
- para->getParD(level)->velocityBC.k,
- para->getParD(level)->velocityBC.q27[0],
- para->getParD(level)->velocityBC.numberOfBCnodes,
- para->getParD(level)->omega,
- para->getParD(level)->neighborX,
- para->getParD(level)->neighborY,
- para->getParD(level)->neighborZ,
- para->getParD(level)->numberOfNodes,
- para->getParD(level)->isEvenTimestep);
+ //QVelDevicePlainBB27(para->getParD(level).get(), &(para->getParD(level)->velocityBC));
//////////////////////////////////////////////////////////////////////////
// D E P R E C A T E D
@@ -205,12 +148,12 @@ void LBKernelManager::runVelocityBCKernelPost(int level)
// para->getPhi(), para->getAngularVelocity(),
// para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ,
// para->getParD(level)->coordinateX, para->getParD(level)->coordinateY, para->getParD(level)->coordinateZ,
- // para->getParD(level)->size_Mat, para->getParD(level)->isEvenTimestep);
+ // para->getParD(level)->numberOfNodes, para->getParD(level)->isEvenTimestep);
// getLastCudaError("QVelDev27 execution failed");
}
}
-void LBKernelManager::runGeoBCKernelPre(int level, unsigned int t, CudaMemoryManager* cudaMemoryManager){
+void LBKernelManager::runGeoBCKernelPre(const int level, unsigned int t, CudaMemoryManager* cudaMemoryManager) const{
if (para->getParD(level)->geometryBC.numberOfBCnodes > 0){
if (para->getCalcDragLift())
{
@@ -318,7 +261,7 @@ void LBKernelManager::runGeoBCKernelPre(int level, unsigned int t, CudaMemoryMan
}
}
-void LBKernelManager::runGeoBCKernelPost(int level)
+void LBKernelManager::runGeoBCKernelPost(const int level) const
{
if (para->getParD(level)->geometryBC.numberOfBCnodes > 0)
{
@@ -353,7 +296,7 @@ void LBKernelManager::runGeoBCKernelPost(int level)
// para->getParD(level)->neighborX,
// para->getParD(level)->neighborY,
// para->getParD(level)->neighborZ,
- // para->getParD(level)->size_Mat,
+ // para->getParD(level)->numberOfNodes,
// para->getParD(level)->isEvenTimestep);
// QDev27(
@@ -368,7 +311,7 @@ void LBKernelManager::runGeoBCKernelPost(int level)
// para->getParD(level)->neighborX,
// para->getParD(level)->neighborY,
// para->getParD(level)->neighborZ,
- // para->getParD(level)->size_Mat,
+ // para->getParD(level)->numberOfNodes,
// para->getParD(level)->isEvenTimestep);
// QVelDev27(
@@ -386,7 +329,7 @@ void LBKernelManager::runGeoBCKernelPost(int level)
// para->getParD(level)->neighborX,
// para->getParD(level)->neighborY,
// para->getParD(level)->neighborZ,
- // para->getParD(level)->size_Mat,
+ // para->getParD(level)->numberOfNodes,
// para->getParD(level)->isEvenTimestep);
// QDevComp27(
@@ -401,24 +344,24 @@ void LBKernelManager::runGeoBCKernelPost(int level)
// para->getParD(level)->neighborX,
// para->getParD(level)->neighborY,
// para->getParD(level)->neighborZ,
- // para->getParD(level)->size_Mat,
+ // para->getParD(level)->numberOfNodes,
// para->getParD(level)->isEvenTimestep);
- QVelDevComp27(
- para->getParD(level)->numberofthreads,
- para->getParD(level)->geometryBC.Vx,
- para->getParD(level)->geometryBC.Vy,
- para->getParD(level)->geometryBC.Vz,
- para->getParD(level)->distributions.f[0],
- para->getParD(level)->geometryBC.k,
- para->getParD(level)->geometryBC.q27[0],
- para->getParD(level)->geometryBC.numberOfBCnodes,
- para->getParD(level)->omega,
- para->getParD(level)->neighborX,
- para->getParD(level)->neighborY,
- para->getParD(level)->neighborZ,
- para->getParD(level)->numberOfNodes,
- para->getParD(level)->isEvenTimestep);
+ // QVelDevComp27(
+ // para->getParD(level)->numberofthreads,
+ // para->getParD(level)->geometryBC.Vx,
+ // para->getParD(level)->geometryBC.Vy,
+ // para->getParD(level)->geometryBC.Vz,
+ // para->getParD(level)->distributions.f[0],
+ // para->getParD(level)->geometryBC.k,
+ // para->getParD(level)->geometryBC.q27[0],
+ // para->getParD(level)->geometryBC.numberOfBCnodes,
+ // para->getParD(level)->omega,
+ // para->getParD(level)->neighborX,
+ // para->getParD(level)->neighborY,
+ // para->getParD(level)->neighborZ,
+ // para->getParD(level)->numberOfNodes,
+ // para->getParD(level)->isEvenTimestep);
// QVelDevCompZeroPress27(
// para->getParD(0)->numberofthreads,
@@ -433,7 +376,7 @@ void LBKernelManager::runGeoBCKernelPost(int level)
// para->getParD(0)->neighborX,
// para->getParD(0)->neighborY,
// para->getParD(0)->neighborZ,
- // para->getParD(0)->size_Mat,
+ // para->getParD(0)->numberOfNodes,
// para->getParD(0)->isEvenTimestep);
// QDev3rdMomentsComp27(
@@ -448,7 +391,7 @@ void LBKernelManager::runGeoBCKernelPost(int level)
// para->getParD(level)->neighborX,
// para->getParD(level)->neighborY,
// para->getParD(level)->neighborZ,
- // para->getParD(level)->size_Mat,
+ // para->getParD(level)->numberOfNodes,
// para->getParD(level)->isEvenTimestep);
// QSlipDev27(
@@ -461,7 +404,7 @@ void LBKernelManager::runGeoBCKernelPost(int level)
// para->getParD(level)->neighborX,
// para->getParD(level)->neighborY,
// para->getParD(level)->neighborZ,
- // para->getParD(level)->size_Mat,
+ // para->getParD(level)->numberOfNodes,
// para->getParD(level)->isEvenTimestep);
//////////////////////////////////////////////////////////////////////////
@@ -482,7 +425,7 @@ void LBKernelManager::runGeoBCKernelPost(int level)
// para->getParD(level)->neighborX,
// para->getParD(level)->neighborY,
// para->getParD(level)->neighborZ,
- // para->getParD(level)->size_Mat_SP,
+ // para->getParD(level)->numberOfNodes,
// para->getParD(level)->isEvenTimestep);
// QSlipNormDevComp27(
@@ -498,12 +441,12 @@ void LBKernelManager::runGeoBCKernelPost(int level)
// para->getParD(level)->neighborX,
// para->getParD(level)->neighborY,
// para->getParD(level)->neighborZ,
- // para->getParD(level)->size_Mat_SP,
+ // para->getParD(level)->numberOfNodes,
// para->getParD(level)->isEvenTimestep);
}
}
-void LBKernelManager::runOutflowBCKernelPre(int level){
+void LBKernelManager::runOutflowBCKernelPre(const int level) const{
if (para->getParD(level)->outflowBC.numberOfBCnodes > 0)
{
QPressNoRhoDev27(
@@ -542,7 +485,7 @@ void LBKernelManager::runOutflowBCKernelPre(int level){
}
}
-void LBKernelManager::runPressureBCKernelPre(int level){
+void LBKernelManager::runPressureBCKernelPre(const int level) const{
if (para->getParD(level)->pressureBC.numberOfBCnodes > 0)
{
QPressNoRhoDev27(
@@ -622,7 +565,7 @@ void LBKernelManager::runPressureBCKernelPre(int level){
}
}
-void LBKernelManager::runPressureBCKernelPost(int level){
+void LBKernelManager::runPressureBCKernelPost(const int level) const{
if (para->getParD(level)->pressureBC.numberOfBCnodes > 0)
{
// QPressDev27_IntBB(
@@ -641,7 +584,7 @@ void LBKernelManager::runPressureBCKernelPost(int level){
}
}
-void LBKernelManager::runStressWallModelKernel(int level){
+void LBKernelManager::runStressWallModelKernel(const int level) const{
if (para->getParD(level)->stressBC.numberOfBCnodes > 0)
{
// QStressDevComp27(para->getParD(level)->numberofthreads, para->getParD(level)->distributions.f[0],
@@ -656,7 +599,7 @@ void LBKernelManager::runStressWallModelKernel(int level){
// para->getHasWallModelMonitor(), para->getParD(level)->wallModel.u_star,
// para->getParD(level)->wallModel.Fx, para->getParD(level)->wallModel.Fy, para->getParD(level)->wallModel.Fz,
// para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ,
- // para->getParD(level)->size_Mat, para->getParD(level)->isEvenTimestep);
+ // para->getParD(level)->numberOfNodes, para->getParD(level)->isEvenTimestep);
BBStressDev27( para->getParD(level)->numberofthreads, para->getParD(level)->distributions.f[0],
para->getParD(level)->stressBC.k, para->getParD(level)->stressBC.kN,
@@ -674,7 +617,7 @@ void LBKernelManager::runStressWallModelKernel(int level){
}
-void LBKernelManager::runSlipBCKernel(int level){
+void LBKernelManager::runSlipBCKernel(const int level) const{
if (para->getParD(level)->slipBC.numberOfBCnodes > 0)
{
// QSlipDev27(
@@ -687,7 +630,7 @@ void LBKernelManager::runSlipBCKernel(int level){
// para->getParD(level)->neighborX,
// para->getParD(level)->neighborY,
// para->getParD(level)->neighborZ,
- // para->getParD(level)->size_Mat,
+ // para->getParD(level)->numberOfNodes,
// para->getParD(level)->isEvenTimestep);
QSlipDevComp27(
@@ -707,7 +650,7 @@ void LBKernelManager::runSlipBCKernel(int level){
}
}
-void LBKernelManager::runNoSlipBCKernel(int level){
+void LBKernelManager::runNoSlipBCKernel(const int level) const{
if (para->getParD(level)->noSlipBC.numberOfBCnodes > 0)
{
// QDev27(
@@ -722,7 +665,7 @@ void LBKernelManager::runNoSlipBCKernel(int level){
// para->getParD(level)->neighborX,
// para->getParD(level)->neighborY,
// para->getParD(level)->neighborZ,
- // para->getParD(level)->size_Mat,
+ // para->getParD(level)->numberOfNodes,
// para->getParD(level)->isEvenTimestep);
// BBDev27(
@@ -737,7 +680,7 @@ void LBKernelManager::runNoSlipBCKernel(int level){
// para->getParD(level)->neighborX,
// para->getParD(level)->neighborY,
// para->getParD(level)->neighborZ,
- // para->getParD(level)->size_Mat,
+ // para->getParD(level)->numberOfNodes,
// para->getParD(level)->isEvenTimestep);
QDevComp27(
@@ -755,7 +698,7 @@ void LBKernelManager::runNoSlipBCKernel(int level){
}
}
-// void LBKernelManager::calculateMacroscopicValues(int level)
+// void LBKernelManager::calculateMacroscopicValues(const int level) const
// {
// if (para->getIsADcalculationOn()) {
// CalcMacADCompSP27(
@@ -790,27 +733,4 @@ void LBKernelManager::runNoSlipBCKernel(int level){
// para->getParD()->distributions.f[0],
// para->getParD()->isEvenTimestep);
// }
-// }
-
-
-
-
-
-
-
-
-
-SPtr<LBKernelManager> LBKernelManager::make(SPtr<Parameter> parameter)
-{
- return SPtr<LBKernelManager>(new LBKernelManager(parameter));
-}
-
-LBKernelManager::LBKernelManager(SPtr<Parameter> parameter)
-{
- this->para = parameter;
-}
-
-LBKernelManager::LBKernelManager(const LBKernelManager&)
-{
-
-}
+// }#
diff --git a/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.h b/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.h
index 75c53f674e9b5b86e50aea9b12c4aa876d84e2f7..6bcdf6b25bf10d56c8d40f37bf68532d2fd91df6 100644
--- a/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.h
+++ b/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.h
@@ -33,70 +33,70 @@
#ifndef LBKernelManager_H
#define LBKernelManager_H
+#include <functional>
#include <memory>
+
#include "PointerDefinitions.h"
#include "VirtualFluids_GPU_export.h"
+#include "LBM/LB.h"
-//! \brief Class forwarding for Parameter
-class Parameter;
class CudaMemoryManager;
+class BoundaryConditionFactory;
+class Parameter;
+class LBMSimulationParameter;
+
+using boundaryCondition = std::function<void(LBMSimulationParameter *, QforBoundaryConditions *)>;
//! \class LBKernelManager
//! \brief manage the cuda kernel calls
class VIRTUALFLUIDS_GPU_EXPORT LBKernelManager
{
public:
- //! \brief makes an object of LBKernelManager
- //! \param para shared pointer to instance of class Parameter
- static SPtr<LBKernelManager> make(std::shared_ptr<Parameter> parameter);
-
- //! \brief calls the device function of the lattice Boltzmann kernel
- void runLBMKernel(int level);
+ //! Class constructor
+ //! \param parameter shared pointer to instance of class Parameter
+ LBKernelManager(SPtr<Parameter> parameter, BoundaryConditionFactory *bcFactory);
- //! \brief calls the device function of the velocity boundary condition (post-collision)
- void runVelocityBCKernelPost(int level);
+ void setBoundaryConditionKernels();
- //! \brief calls the device function of the velocity boundary condition (pre-collision)
- void runVelocityBCKernelPre(int level);
+ //! \brief calls the device function of the lattice Boltzmann kernel
+ void runLBMKernel(const int level) const;
- //! \brief calls the device function of the geometry boundary condition (post-collision)
- void runGeoBCKernelPost(int level);
+ //! \brief calls the device function of the velocity boundary condition (post-collision)
+ void runVelocityBCKernelPost(const int level) const;
- //! \brief calls the device function of the geometry boundary condition (pre-collision)
- void runGeoBCKernelPre(int level, unsigned int t, CudaMemoryManager* cudaMemoryManager);
+ //! \brief calls the device function of the velocity boundary condition (pre-collision)
+ void runVelocityBCKernelPre(const int level) const;
- //! \brief calls the device function of the slip boundary condition
- void runSlipBCKernel(int level);
+ //! \brief calls the device function of the geometry boundary condition (post-collision)
+ void runGeoBCKernelPost(const int level) const;
- //! \brief calls the device function of the no-slip boundary condition
- void runNoSlipBCKernel(int level);
+ //! \brief calls the device function of the geometry boundary condition (pre-collision)
+ void runGeoBCKernelPre(const int level, unsigned int t, CudaMemoryManager *cudaMemoryManager) const;
- //! \brief calls the device function of the pressure boundary condition (pre-collision)
- void runPressureBCKernelPre(int level);
+ //! \brief calls the device function of the slip boundary condition
+ void runSlipBCKernel(const int level) const;
- //! \brief calls the device function of the pressure boundary condition (post-collision)
- void runPressureBCKernelPost(int level);
+ //! \brief calls the device function of the no-slip boundary condition
+ void runNoSlipBCKernel(const int level) const;
- //! \brief calls the device function of the outflow boundary condition
- void runOutflowBCKernelPre(int level);
+ //! \brief calls the device function of the pressure boundary condition (pre-collision)
+ void runPressureBCKernelPre(const int level) const;
- //! \brief calls the device function of the stress wall model
- void runStressWallModelKernel(int level);
+ //! \brief calls the device function of the pressure boundary condition (post-collision)
+ void runPressureBCKernelPost(const int level) const;
- //! \brief calls the device function that calculates the macroscopic values
- void calculateMacroscopicValues(int level);
+ //! \brief calls the device function of the outflow boundary condition
+ void runOutflowBCKernelPre(const int level) const;
+ //! \brief calls the device function of the stress wall model
+ void runStressWallModelKernel(const int level) const;
-private:
- //! Class constructor
- //! \param parameter shared pointer to instance of class Parameter
- LBKernelManager(SPtr<Parameter> parameter);
- //! Class copy constructor
- //! \param LBKernelManager is a reference to LBKernelManager object
- LBKernelManager(const LBKernelManager&);
+ //! \brief calls the device function that calculates the macroscopic values
+ void calculateMacroscopicValues(const int level) const;
- //! \property para is a shared pointer to an object of Parameter
- SPtr<Parameter> para;
+private:
+ SPtr<Parameter> para;
+ boundaryCondition velocityBoundaryConditionPost;
};
#endif
diff --git a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
index 29725c5a30737161aeefd8d3642febaf0385bf53..4cdb9f1db6514d61ff90327a928822df924e7270 100644
--- a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
+++ b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
@@ -68,7 +68,7 @@ std::string getFileName(const std::string& fname, int step, int myID)
}
Simulation::Simulation(std::shared_ptr<Parameter> para, std::shared_ptr<CudaMemoryManager> memoryManager,
- vf::gpu::Communicator &communicator, GridProvider &gridProvider)
+ vf::gpu::Communicator &communicator, GridProvider &gridProvider, BoundaryConditionFactory* bcFactory)
: para(para), cudaMemoryManager(memoryManager), communicator(communicator), kernelFactory(std::make_unique<KernelFactoryImp>()),
preProcessorFactory(std::make_unique<PreProcessorFactoryImp>()), dataWriter(std::make_unique<FileWriter>())
{
@@ -377,7 +377,7 @@ Simulation::Simulation(std::shared_ptr<Parameter> para, std::shared_ptr<CudaMemo
//////////////////////////////////////////////////////////////////////////
// Init UpdateGrid
//////////////////////////////////////////////////////////////////////////
- this->updateGrid27 = std::make_unique<UpdateGrid27>(para, communicator, cudaMemoryManager, pm, kernels);
+ this->updateGrid27 = std::make_unique<UpdateGrid27>(para, communicator, cudaMemoryManager, pm, kernels, bcFactory);
//////////////////////////////////////////////////////////////////////////
// Print Init
diff --git a/src/gpu/VirtualFluids_GPU/LBM/Simulation.h b/src/gpu/VirtualFluids_GPU/LBM/Simulation.h
index 5c765bc992c247ebf9b3b51ac399711ff7db0651..e43e88e482ea56e3d2fecdcd1275b209575e3442 100644
--- a/src/gpu/VirtualFluids_GPU/LBM/Simulation.h
+++ b/src/gpu/VirtualFluids_GPU/LBM/Simulation.h
@@ -28,12 +28,13 @@ class TrafficMovementFactory;
class UpdateGrid27;
class KineticEnergyAnalyzer;
class EnstrophyAnalyzer;
+class BoundaryConditionFactory;
class Simulation
{
public:
Simulation(std::shared_ptr<Parameter> para, std::shared_ptr<CudaMemoryManager> memoryManager,
- vf::gpu::Communicator &communicator, GridProvider &gridProvider);
+ vf::gpu::Communicator &communicator, GridProvider &gridProvider, BoundaryConditionFactory* bcFactory);
~Simulation();
void run();