From c0a6a874f623c5c7f4d9eaabbe19eafbeccaec16 Mon Sep 17 00:00:00 2001
From: Anna Wellmann <a.wellmann@tu-bs.de>
Date: Mon, 4 Jul 2022 12:23:28 +0000
Subject: [PATCH] Add pressure bc to BoundaryConditionFactory
---
apps/gpu/LBM/Sphere/Sphere.cpp | 2 +
.../BoundaryConditionFactory.cpp | 29 ++
.../BoundaryConditionFactory.h | 19 +-
src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h | 79 +---
src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu | 345 ++++++------------
.../KernelManager/LBKernelManager.cpp | 92 +----
.../KernelManager/LBKernelManager.h | 1 +
7 files changed, 166 insertions(+), 401 deletions(-)
diff --git a/apps/gpu/LBM/Sphere/Sphere.cpp b/apps/gpu/LBM/Sphere/Sphere.cpp
index 355859760..3f4ff459c 100644
--- a/apps/gpu/LBM/Sphere/Sphere.cpp
+++ b/apps/gpu/LBM/Sphere/Sphere.cpp
@@ -172,10 +172,12 @@ void multipleLevel(const std::string &configPath)
gridBuilder->setSlipBoundaryCondition(SideType::PZ, 0.0, 0.0, 0.0);
gridBuilder->setSlipBoundaryCondition(SideType::MZ, 0.0, 0.0, 0.0);
gridBuilder->setVelocityBoundaryCondition(SideType::GEOMETRY, 0.0, 0.0, 0.0);
+ gridBuilder->setPressureBoundaryCondition(SideType::PX, 0.0); // set pressure boundary condition last
bcFactory.setNoSlipBoundaryCondition(BoundaryConditionFactory::NoSlipBC::NoSlipCompressible);
bcFactory.setVelocityBoundaryCondition(BoundaryConditionFactory::VelocityBC::VelocityCompressible);
bcFactory.setSlipBoundaryCondition(BoundaryConditionFactory::SlipBC::SlipCompressible);
+ bcFactory.setPressureBoundaryCondition(BoundaryConditionFactory::PressureBC::PressureNonEquilibriumCompressible);
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.cpp b/src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.cpp
index 2d97f1998..d9ce7853d 100644
--- a/src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.cpp
+++ b/src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.cpp
@@ -18,6 +18,11 @@ void BoundaryConditionFactory::setSlipBoundaryCondition(const SlipBC boundaryCon
this->slipBoundaryCondition = boundaryConditionType;
}
+void BoundaryConditionFactory::setPressureBoundaryCondition(const PressureBC boundaryConditionType)
+{
+ this->pressureBoundaryCondition = boundaryConditionType;
+}
+
boundaryCondition BoundaryConditionFactory::getVelocityBoundaryConditionPost() const
{
// for descriptions of the boundary conditions refer to the header
@@ -75,6 +80,30 @@ boundaryCondition BoundaryConditionFactory::getSlipBoundaryConditionPost() const
}
}
+boundaryCondition BoundaryConditionFactory::getPressureBoundaryConditionPre() const
+{
+ // for descriptions of the boundary conditions refer to the header
+ switch (this->pressureBoundaryCondition) {
+ case PressureBC::PressureEquilibrium:
+ return QPressDev27;
+ break;
+ case PressureBC::PressureEquilibrium2:
+ return QPressDevEQZ27;
+ break;
+ case PressureBC::PressureNonEquilibriumIncompressible:
+ return QPressDevIncompNEQ27;
+ break;
+ case PressureBC::PressureNonEquilibriumCompressible:
+ return QPressDevNEQ27;
+ break;
+ case PressureBC::OutflowNonReflective:
+ return QPressNoRhoDev27;
+ break;
+ default:
+ return nullptr;
+ }
+}
+
// boundaryCondition BoundaryConditionFactory::getGeometryBoundaryConditionPost() const{
// this->getNoSlipBoundaryConditionPost();
// }
\ No newline at end of file
diff --git a/src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.h b/src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.h
index 150bfcf0b..a09b8e322 100644
--- a/src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.h
+++ b/src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.h
@@ -44,16 +44,31 @@ public:
//! - SlipCompressible = interpolated slip boundary condition, based on subgrid distances
SlipCompressible,
//! - SlipCompressible = interpolated slip boundary condition, based on subgrid distances.
- //! With turbulent viscosity
+ //! With turbulent viscosity -> para->setUseTurbulentViscosity(true) has to be set to true
SlipCompressibleTurbulentViscosity
};
+ //! \brief An enumeration for selecting a pressure boundary condition
+ enum class PressureBC {
+ //! - PressureEquilibrium = pressure boundary condition based on equilibrium
+ PressureEquilibrium, // incorrect pressure :(
+ //! - PressureEquilibrium2 = pressure boundary condition based on equilibrium (potentially better?! than PressureEquilibrium)
+ PressureEquilibrium2, // is broken --> nan :(
+ //! - PressureNonEquilibriumIncompressible = pressure boundary condition based on non-equilibrium
+ PressureNonEquilibriumIncompressible,
+ //! - PressureNonEquilibriumCompressible = pressure boundary condition based on non-equilibrium
+ PressureNonEquilibriumCompressible,
+ //! - OutflowNonReflective = outflow boundary condition
+ OutflowNonReflective
+ };
+
// enum class OutflowBoundaryCondition {}; // TODO:
// https://git.rz.tu-bs.de/m.schoenherr/VirtualFluids_dev/-/issues/16
void setVelocityBoundaryCondition(const VelocityBC boundaryConditionType);
void setNoSlipBoundaryCondition(const NoSlipBC boundaryConditionType);
void setSlipBoundaryCondition(const SlipBC boundaryConditionType);
+ void setPressureBoundaryCondition(const PressureBC boundaryConditionType);
// void setGeometryBoundaryCondition(const std::variant<VelocityBC, NoSlipBC, SlipBC> boundaryConditionType);
// void setOutflowBoundaryCondition(...); // TODO:
@@ -62,11 +77,13 @@ public:
boundaryCondition getVelocityBoundaryConditionPost() const;
boundaryCondition getNoSlipBoundaryConditionPost() const;
boundaryCondition getSlipBoundaryConditionPost() const;
+ boundaryCondition getPressureBoundaryConditionPre() const;
private:
VelocityBC velocityBoundaryCondition;
NoSlipBC noSlipBoundaryCondition;
SlipBC slipBoundaryCondition;
+ PressureBC pressureBoundaryCondition;
// OutflowBoundaryConditon outflowBC // TODO: https://git.rz.tu-bs.de/m.schoenherr/VirtualFluids_dev/-/issues/16
};
diff --git a/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h b/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h
index 8e260bdbf..d007b6e59 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h
+++ b/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h
@@ -938,18 +938,7 @@ extern "C" void BBStressDev27( unsigned int numberOfThreads,
unsigned int size_Mat,
bool isEvenTimestep);
-extern "C" void QPressDev27(unsigned int numberOfThreads,
- real* rhoBC,
- 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 QPressDev27(LBMSimulationParameter* parameterDevice, QforBoundaryConditions* boundaryCondition);
extern "C" void QPressDevFixBackflow27(unsigned int numberOfThreads,
real* rhoBC,
@@ -975,31 +964,9 @@ extern "C" void QPressDevDirDepBot27(unsigned int numberOfThreads,
unsigned int size_Mat,
bool isEvenTimestep);
-extern "C" void QPressNoRhoDev27( unsigned int numberOfThreads,
- real* rhoBC,
- real* DD,
- int* k_Q,
- int* k_N,
- unsigned int numberOfBCnodes,
- real om1,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int size_Mat,
- bool isEvenTimestep);
+extern "C" void QPressNoRhoDev27(LBMSimulationParameter* parameterDevice, QforBoundaryConditions* boundaryCondition);
-extern "C" void QInflowScaleByPressDev27(unsigned int numberOfThreads,
- real* rhoBC,
- real* DD,
- int* k_Q,
- int* k_N,
- unsigned int numberOfBCnodes,
- real om1,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int size_Mat,
- bool isEvenTimestep);
+extern "C" void QInflowScaleByPressDev27(LBMSimulationParameter* parameterDevice, QforBoundaryConditions* boundaryCondition);
extern "C" void QPressDevOld27(unsigned int numberOfThreads,
real* rhoBC,
@@ -1014,45 +981,11 @@ extern "C" void QPressDevOld27(unsigned int numberOfThreads,
unsigned int size_Mat,
bool isEvenTimestep);
-extern "C" void QPressDevIncompNEQ27(unsigned int numberOfThreads,
- real* rhoBC,
- real* DD,
- int* k_Q,
- int* k_N,
- unsigned int numberOfBCnodes,
- real om1,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int size_Mat,
- bool isEvenTimestep);
+extern "C" void QPressDevIncompNEQ27(LBMSimulationParameter* parameterDevice, QforBoundaryConditions* boundaryCondition);
-extern "C" void QPressDevNEQ27(unsigned int numberOfThreads,
- real* rhoBC,
- real* distribution,
- int* bcNodeIndices,
- int* bcNeighborIndices,
- unsigned int numberOfBCnodes,
- real omega1,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int size_Mat,
- bool isEvenTimestep);
+extern "C" void QPressDevNEQ27(LBMSimulationParameter* parameterDevice, QforBoundaryConditions* boundaryCondition);
-extern "C" void QPressDevEQZ27(unsigned int numberOfThreads,
- real* rhoBC,
- real* DD,
- int* k_Q,
- int* k_N,
- real* kTestRE,
- unsigned int numberOfBCnodes,
- real om1,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int size_Mat,
- bool isEvenTimestep);
+extern "C" void QPressDevEQZ27(LBMSimulationParameter* parameterDevice, QforBoundaryConditions* boundaryCondition);
extern "C" void QPressDevZero27(unsigned int numberOfThreads,
real* DD,
diff --git a/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu b/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu
index db41645a0..c4b439f80 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu
+++ b/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu
@@ -3670,46 +3670,24 @@ extern "C" void BBStressDev27(unsigned int numberOfThreads,
getLastCudaError("BBStressDevice27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
-extern "C" void QPressDev27(unsigned int numberOfThreads,
- real* rhoBC,
- 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 QPressDev27(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 );
- QPressDevice27<<< gridQ, threads >>> ( rhoBC,
- DD,
- k_Q,
- QQ,
- numberOfBCnodes,
- om1,
- neighborX,
- neighborY,
- neighborZ,
- size_Mat,
- isEvenTimestep);
- getLastCudaError("QPressDevice27 execution failed");
+ QPressDevice27<<< grid, threads >>> (
+ boundaryCondition->RhoBC,
+ parameterDevice->distributions.f[0],
+ boundaryCondition->k,
+ boundaryCondition->q27[0],
+ boundaryCondition->numberOfBCnodes,
+ parameterDevice->omega,
+ parameterDevice->neighborX,
+ parameterDevice->neighborY,
+ parameterDevice->neighborZ,
+ parameterDevice->numberOfNodes,
+ parameterDevice->isEvenTimestep);
+ getLastCudaError("QPressDevice27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
extern "C" void QPressDevAntiBB27( unsigned int numberOfThreads,
@@ -3840,87 +3818,43 @@ extern "C" void QPressDevDirDepBot27( unsigned int numberOfThreads,
getLastCudaError("QPressDeviceDirDepBot27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
-extern "C" void QPressNoRhoDev27(unsigned int numberOfThreads,
- real* rhoBC,
- real* DD,
- int* k_Q,
- int* k_N,
- unsigned int numberOfBCnodes,
- real om1,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int size_Mat,
- bool isEvenTimestep)
+extern "C" void QPressNoRhoDev27(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 );
- QPressNoRhoDevice27<<< gridQ, threads >>> ( rhoBC,
- DD,
- k_Q,
- k_N,
- numberOfBCnodes,
- om1,
- neighborX,
- neighborY,
- neighborZ,
- size_Mat,
- isEvenTimestep);
- getLastCudaError("QPressNoRhoDevice27 execution failed");
+ QPressNoRhoDevice27<<< grid, threads >>> (
+ boundaryCondition->RhoBC,
+ parameterDevice->distributions.f[0],
+ boundaryCondition->k,
+ boundaryCondition->kN,
+ boundaryCondition->numberOfBCnodes,
+ parameterDevice->omega,
+ parameterDevice->neighborX,
+ parameterDevice->neighborY,
+ parameterDevice->neighborZ,
+ parameterDevice->numberOfNodes,
+ parameterDevice->isEvenTimestep);
+ getLastCudaError("QPressNoRhoDevice27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
-extern "C" void QInflowScaleByPressDev27(unsigned int numberOfThreads,
- real* rhoBC,
- real* DD,
- int* k_Q,
- int* k_N,
- unsigned int numberOfBCnodes,
- real om1,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int size_Mat,
- bool isEvenTimestep)
+extern "C" void QInflowScaleByPressDev27(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 );
- QInflowScaleByPressDevice27<<< gridQ, threads >>> ( rhoBC,
- DD,
- k_Q,
- k_N,
- numberOfBCnodes,
- om1,
- neighborX,
- neighborY,
- neighborZ,
- size_Mat,
- isEvenTimestep);
+ QInflowScaleByPressDevice27<<< grid, threads >>> (
+ boundaryCondition->RhoBC,
+ parameterDevice->distributions.f[0],
+ boundaryCondition->k,
+ boundaryCondition->kN,
+ boundaryCondition->numberOfBCnodes,
+ parameterDevice->omega,
+ parameterDevice->neighborX,
+ parameterDevice->neighborY,
+ parameterDevice->neighborZ,
+ parameterDevice->numberOfNodes,
+ parameterDevice->isEvenTimestep);
getLastCudaError("QInflowScaleByPressDevice27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
@@ -3966,131 +3900,64 @@ extern "C" void QPressDevOld27( unsigned int numberOfThreads,
getLastCudaError("QPressDeviceOld27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
-extern "C" void QPressDevIncompNEQ27(unsigned int numberOfThreads,
- real* rhoBC,
- real* DD,
- int* k_Q,
- int* k_N,
- unsigned int numberOfBCnodes,
- real om1,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int size_Mat,
- bool isEvenTimestep)
+extern "C" void QPressDevIncompNEQ27(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 );
- QPressDeviceIncompNEQ27<<< gridQ, threads >>> ( rhoBC,
- DD,
- k_Q,
- k_N,
- numberOfBCnodes,
- om1,
- neighborX,
- neighborY,
- neighborZ,
- size_Mat,
- isEvenTimestep);
- getLastCudaError("QPressDeviceIncompNEQ27 execution failed");
+ QPressDeviceIncompNEQ27<<< grid, threads >>> (
+ boundaryCondition->RhoBC,
+ parameterDevice->distributions.f[0],
+ boundaryCondition->k,
+ boundaryCondition->kN,
+ boundaryCondition->numberOfBCnodes,
+ parameterDevice->omega,
+ parameterDevice->neighborX,
+ parameterDevice->neighborY,
+ parameterDevice->neighborZ,
+ parameterDevice->numberOfNodes,
+ parameterDevice->isEvenTimestep);
+ getLastCudaError("QPressDeviceIncompNEQ27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
-extern "C" void QPressDevNEQ27( unsigned int numberOfThreads,
- real* rhoBC,
- real* distribution,
- int* bcNodeIndices,
- int* bcNeighborIndices,
- unsigned int numberOfBCnodes,
- real omega1,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int size_Mat,
- bool isEvenTimestep)
+extern "C" void QPressDevNEQ27(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 );
- QPressDeviceNEQ27<<< gridQ, threads >>> ( rhoBC,
- distribution,
- bcNodeIndices,
- bcNeighborIndices,
- numberOfBCnodes,
- omega1,
- neighborX,
- neighborY,
- neighborZ,
- size_Mat,
- isEvenTimestep);
- getLastCudaError("QPressDeviceOld27 execution failed");
+ QPressDeviceNEQ27<<< grid, threads >>> (
+ boundaryCondition->RhoBC,
+ parameterDevice->distributions.f[0],
+ boundaryCondition->k,
+ boundaryCondition->kN,
+ boundaryCondition->numberOfBCnodes,
+ parameterDevice->omega,
+ parameterDevice->neighborX,
+ parameterDevice->neighborY,
+ parameterDevice->neighborZ,
+ parameterDevice->numberOfNodes,
+ parameterDevice->isEvenTimestep);
+ getLastCudaError("QPressDevNEQ27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
-extern "C" void QPressDevEQZ27( unsigned int numberOfThreads,
- real* rhoBC,
- real* DD,
- int* k_Q,
- int* k_N,
- real* kTestRE,
- unsigned int numberOfBCnodes,
- real om1,
- unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int size_Mat,
- bool isEvenTimestep)
+extern "C" void QPressDevEQZ27(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 );
- QPressDeviceEQZ27<<< gridQ, threads >>> ( rhoBC,
- DD,
- k_Q,
- k_N,
- kTestRE,
- numberOfBCnodes,
- om1,
- neighborX,
- neighborY,
- neighborZ,
- size_Mat,
- isEvenTimestep);
+ QPressDeviceEQZ27<<< grid, threads >>> (
+ boundaryCondition->RhoBC,
+ parameterDevice->distributions.f[0],
+ boundaryCondition->k,
+ boundaryCondition->kN,
+ parameterDevice->kDistTestRE.f[0],
+ boundaryCondition->numberOfBCnodes,
+ parameterDevice->omega,
+ parameterDevice->neighborX,
+ parameterDevice->neighborY,
+ parameterDevice->neighborZ,
+ parameterDevice->numberOfNodes,
+ parameterDevice->isEvenTimestep);
getLastCudaError("QPressDeviceEQZ27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
@@ -4177,17 +4044,17 @@ extern "C" void BBDev27(LBMSimulationParameter* parameterDevice, QforBoundaryCon
dim3 grid = vf::cuda::getCudaGrid( parameterDevice->numberofthreads, boundaryCondition->numberOfBCnodes);
dim3 threads(parameterDevice->numberofthreads, 1, 1 );
- BBDevice27<<< grid, threads >>> (
- parameterDevice->distributions.f[0],
- boundaryCondition->k,
- boundaryCondition->q27[0],
- boundaryCondition->numberOfBCnodes,
- parameterDevice->neighborX,
- parameterDevice->neighborY,
- parameterDevice->neighborZ,
- parameterDevice->numberOfNodes,
- parameterDevice->isEvenTimestep);
- getLastCudaError("BBDevice27 execution failed");
+ BBDevice27<<< grid, threads >>> (
+ parameterDevice->distributions.f[0],
+ boundaryCondition->k,
+ boundaryCondition->q27[0],
+ boundaryCondition->numberOfBCnodes,
+ parameterDevice->neighborX,
+ parameterDevice->neighborY,
+ parameterDevice->neighborZ,
+ parameterDevice->numberOfNodes,
+ parameterDevice->isEvenTimestep);
+ getLastCudaError("BBDevice27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
extern "C" void QPressDev27_IntBB( unsigned int numberOfThreads,
diff --git a/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.cpp b/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.cpp
index 0ca174fac..f3b6d072e 100644
--- a/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.cpp
+++ b/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.cpp
@@ -45,6 +45,7 @@ LBKernelManager::LBKernelManager(SPtr<Parameter> parameter, BoundaryConditionFac
this->velocityBoundaryConditionPost = bcFactory->getVelocityBoundaryConditionPost();
this->noSlipBoundaryConditionPost = bcFactory->getNoSlipBoundaryConditionPost();
this->slipBoundaryConditionPost = bcFactory->getSlipBoundaryConditionPost();
+ this->pressureBoundaryConditionPre = bcFactory->getPressureBoundaryConditionPre();
}
void LBKernelManager::runLBMKernel(const int level) const
@@ -449,19 +450,7 @@ void LBKernelManager::runGeoBCKernelPost(const int level) const
void LBKernelManager::runOutflowBCKernelPre(const int level) const{
if (para->getParD(level)->outflowBC.numberOfBCnodes > 0)
{
- QPressNoRhoDev27(
- para->getParD(level)->numberofthreads,
- para->getParD(level)->outflowBC.RhoBC,
- para->getParD(level)->distributions.f[0],
- para->getParD(level)->outflowBC.k,
- para->getParD(level)->outflowBC.kN,
- para->getParD(level)->outflowBC.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);
+ QPressNoRhoDev27(para->getParD(level).get(), &(para->getParD(level)->outflowBC));
// TODO: https://git.rz.tu-bs.de/irmb/VirtualFluids_dev/-/issues/29
// if ( myid == numprocs - 1)
@@ -488,80 +477,7 @@ void LBKernelManager::runOutflowBCKernelPre(const int level) const{
void LBKernelManager::runPressureBCKernelPre(const int level) const{
if (para->getParD(level)->pressureBC.numberOfBCnodes > 0)
{
- QPressNoRhoDev27(
- para->getParD(level)->numberofthreads,
- para->getParD(level)->pressureBC.RhoBC,
- para->getParD(level)->distributions.f[0],
- para->getParD(level)->pressureBC.k,
- para->getParD(level)->pressureBC.kN,
- para->getParD(level)->pressureBC.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);
-
- // QPressDevEQZ27(
- // para->getParD(level)->numberofthreads,
- // para->getParD(level)->pressureBC.RhoBC,
- // para->getParD(level)->distributions.f[0],
- // para->getParD(level)->pressureBC.k,
- // para->getParD(level)->pressureBC.kN,
- // para->getParD(level)->kDistTestRE.f[0],
- // para->getParD(level)->pressureBC.numberOfBCnodes,
- // para->getParD(level)->omega,
- // para->getParD(level)->neighborX,
- // para->getParD(level)->neighborY,
- // para->getParD(0)->neighborZ,
- // para->getParD(level)->numberOfNodes,
- // para->getParD(level)->isEvenTimestep);
-
- // QInflowScaleByPressDev27(
- // para->getParD(level)->numberofthreads,
- // para->getParD(level)->pressureBC.RhoBC,
- // para->getParD(level)->distributions.f[0],
- // para->getParD(level)->pressureBC.k,
- // para->getParD(level)->pressureBC.kN,
- // para->getParD(level)->pressureBC.numberOfBCnodes,
- // para->getParD(0)->omega,
- // para->getParD(level)->neighborX,
- // para->getParD(level)->neighborY,
- // para->getParD(0)->neighborZ,
- // para->getParD(level)->numberOfNodes,
- // para->getParD(level)->isEvenTimestep);
-
- // //////////////////////////////////////////////////////////////////////////////
- // // press NEQ incompressible
- // QPressDevIncompNEQ27(
- // para->getParD(level)->numberofthreads,
- // para->getParD(level)->pressureBC.RhoBC,
- // para->getParD(level)->distributions.f[0],
- // para->getParD(level)->pressureBC.k,
- // para->getParD(level)->pressureBC.kN,
- // para->getParD(level)->pressureBC.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);
-
- // ////////////////////////////////////////////////////////////////////////////////
- // // press NEQ compressible
- // QPressDevNEQ27(
- // para->getParD(level)->numberofthreads,
- // para->getParD(level)->pressureBC.RhoBC,
- // para->getParD(level)->distributions.f[0],
- // para->getParD(level)->pressureBC.k,
- // para->getParD(level)->pressureBC.kN,
- // para->getParD(level)->pressureBC.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);
+ this->pressureBoundaryConditionPre(para->getParD(level).get(), &(para->getParD(level)->pressureBC));
}
}
@@ -666,4 +582,4 @@ void LBKernelManager::runNoSlipBCKernelPost(const int level) const{
// para->getParD()->distributions.f[0],
// para->getParD()->isEvenTimestep);
// }
-// }#
+// }
diff --git a/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.h b/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.h
index cb32e8dc6..e8b8a2cc0 100644
--- a/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.h
+++ b/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.h
@@ -100,5 +100,6 @@ private:
boundaryCondition velocityBoundaryConditionPost;
boundaryCondition noSlipBoundaryConditionPost;
boundaryCondition slipBoundaryConditionPost;
+ boundaryCondition pressureBoundaryConditionPre;
};
#endif
--
GitLab