diff --git a/apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp b/apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp
index 57349a82072dbf8e64c17a1f678224a9a6e177cb..2cde1330ba9832dc9ee7a68542b25541d104a2cc 100644
--- a/apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp
+++ b/apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp
@@ -210,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.setNoSlipBoundaryCondition(BoundaryConditionFactory::NoSlipBC::NoSlipIncompressible);
bcFactory.setVelocityBoundaryCondition(BoundaryConditionFactory::VelocityBC::VelocitySimpleBounceBackCompressible);
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.cpp b/src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.cpp
index 6bc2326e1929f4a2deaa48a1c7400f4a0f6d1639..881824928e08847a76c3b563e4154d67dd5c8dc6 100644
--- a/src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.cpp
+++ b/src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.cpp
@@ -7,9 +7,14 @@ void BoundaryConditionFactory::setVelocityBoundaryCondition(VelocityBC boundaryC
this->velocityBoundaryCondition = boundaryConditionType;
}
+void BoundaryConditionFactory::setNoSlipBoundaryCondition(const NoSlipBC boundaryConditionType)
+{
+ this->noSlipBoundaryCondition = boundaryConditionType;
+}
+
boundaryCondition BoundaryConditionFactory::getVelocityBoundaryConditionPost() const
{
- // for descriptions of the boundary conditions refer to the header ( VelocityBC)
+ // for descriptions of the boundary conditions refer to the header
switch (this->velocityBoundaryCondition) {
case VelocityBC::VelocitySimpleBounceBackCompressible:
return QVelDevicePlainBB27;
@@ -17,7 +22,7 @@ boundaryCondition BoundaryConditionFactory::getVelocityBoundaryConditionPost() c
case VelocityBC::VelocityIncompressible:
return QVelDev27;
break;
- case VelocityBC::VelocityCompressible:
+ case VelocityBC::VelocityCompressible:
return QVelDevComp27;
break;
case VelocityBC::VelocityAndPressureCompressible:
@@ -26,4 +31,26 @@ boundaryCondition BoundaryConditionFactory::getVelocityBoundaryConditionPost() c
default:
return nullptr;
}
-}
\ No newline at end of file
+}
+
+boundaryCondition BoundaryConditionFactory::getNoSlipBoundaryConditionPost() const
+{
+ // for descriptions of the boundary conditions refer to the header
+ switch (this->noSlipBoundaryCondition) {
+ case NoSlipBC::NoSlipBounceBack:
+ return BBDev27;
+ break;
+ case NoSlipBC::NoSlipIncompressible:
+ return QDev27;
+ break;
+ case NoSlipBC::NoSlipCompressible:
+ return QDevComp27;
+ 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 099423fcf08bd028fb667e4f584adaef2b4c071a..10544d996f425fb572e41b491c9c31e7bbe40c22 100644
--- a/src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.h
+++ b/src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionFactory.h
@@ -27,18 +27,32 @@ public:
VelocityAndPressureCompressible
};
+ //! \brief An enumeration for selecting a no-slip boundary condition
+ enum class NoSlipBC {
+ //! - NoSlipBounceBack = bounce back no-slip boundary condition
+ NoSlipBounceBack,
+ //! - NoSlipIncompressible = interpolated no-slip boundary condition, based on subgrid distances
+ NoSlipIncompressible,
+ //! - NoSlipCompressible = interpolated no-slip boundary condition, based on subgrid distances
+ NoSlipCompressible
+ };
+
// 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 setGeometryBoundaryCondition(const std::variant<VelocityBC, NoSlipBC, SlipBC> boundaryConditionType);
// void setOutflowBoundaryCondition(...); // TODO:
// https://git.rz.tu-bs.de/m.schoenherr/VirtualFluids_dev/-/issues/16
boundaryCondition getVelocityBoundaryConditionPost() const;
+ boundaryCondition getNoSlipBoundaryConditionPost() const;
private:
VelocityBC velocityBoundaryCondition;
+ NoSlipBC noSlipBoundaryCondition;
// 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 3f517369c8d1304579500453396b131d91b12e00..997948913654c2e373bc692f25d8be68023e3e6c 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h
+++ b/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h
@@ -646,31 +646,9 @@ extern "C" void BcVel27(int nx,
real u0x,
real om);
-extern "C" void QDev27( unsigned int numberOfThreads,
- int nx,
- int ny,
- 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 QDev27(LBMSimulationParameter* parameterDevice, QforBoundaryConditions* boundaryCondition);
-extern "C" void QDevComp27(unsigned int numberOfThreads,
- 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 QDevComp27(LBMSimulationParameter* parameterDevice, QforBoundaryConditions* boundaryCondition);
extern "C" void QDevCompThinWalls27(unsigned int numberOfThreads,
real* DD,
@@ -1119,19 +1097,7 @@ extern "C" void QPressDevFake27( unsigned int numberOfThreads,
unsigned int size_Mat,
bool isEvenTimestep);
-extern "C" void BBDev27( unsigned int numberOfThreads,
- int nx,
- int ny,
- 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 BBDev27(LBMSimulationParameter* parameterDevice, QforBoundaryConditions* boundaryCondition);
extern "C" void QPressDev27_IntBB( unsigned int numberOfThreads,
real* rho,
diff --git a/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu b/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu
index 74e091bbe2598e5a9fb8f70d4be159947d10fc5f..b557078bb89366f6f6c3cc77d12f083cbd55360d 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu
+++ b/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu
@@ -2650,88 +2650,44 @@ extern "C" void QADPressIncompDev27( unsigned int numberOfThreads,
getLastCudaError("QADPressIncomp27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
-extern "C" void QDev27( unsigned int numberOfThreads,
- int nx,
- int ny,
- 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 QDev27(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 );
+
+ QDevice27<<< grid, threads >>> (
+ parameterDevice->nx,
+ parameterDevice->ny,
+ parameterDevice->distributions.f[0],
+ boundaryCondition->k,
+ boundaryCondition->q27[0],
+ boundaryCondition->numberOfBCnodes,
+ parameterDevice->omega,
+ parameterDevice->neighborX,
+ parameterDevice->neighborY,
+ parameterDevice->neighborZ,
+ parameterDevice->numberOfNodes,
+ parameterDevice->isEvenTimestep);
- QDevice27<<< gridQ, threads >>> (nx,
- ny,
- DD,
- k_Q,
- QQ,
- numberOfBCnodes,
- om1,
- neighborX,
- neighborY,
- neighborZ,
- size_Mat,
- isEvenTimestep);
getLastCudaError("QDevice27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
-extern "C" void QDevComp27( unsigned int numberOfThreads,
- 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 QDevComp27(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 );
- QDeviceComp27<<< gridQ, threads >>> (
- distribution,
- subgridDistanceIndices,
- subgridDistances,
- numberOfBCnodes,
- omega,
- neighborX,
- neighborY,
- neighborZ,
- numberOfLBnodes,
- isEvenTimestep);
+ QDeviceComp27<<< grid, threads >>> (
+ parameterDevice->distributions.f[0],
+ boundaryCondition->k,
+ boundaryCondition->q27[0],
+ boundaryCondition->numberOfBCnodes,
+ parameterDevice->omega,
+ parameterDevice->neighborX,
+ parameterDevice->neighborY,
+ parameterDevice->neighborZ,
+ parameterDevice->numberOfNodes,
+ parameterDevice->isEvenTimestep);
getLastCudaError("QDeviceComp27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
@@ -4260,47 +4216,22 @@ extern "C" void QPressDevFake27( unsigned int numberOfThreads,
getLastCudaError("QPressDeviceFake27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
-extern "C" void BBDev27( unsigned int numberOfThreads,
- int nx,
- int ny,
- 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 BBDev27(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 );
- BBDevice27<<< gridQ, threads >>> ( nx,
- ny,
- DD,
- k_Q,
- QQ,
- numberOfBCnodes,
- om1,
- neighborX,
- neighborY,
- neighborZ,
- size_Mat,
- isEvenTimestep);
+ BBDevice27<<< grid, threads >>> (
+ parameterDevice->distributions.f[0],
+ boundaryCondition->k,
+ boundaryCondition->q27[0],
+ boundaryCondition->numberOfBCnodes,
+ parameterDevice->omega,
+ parameterDevice->neighborX,
+ parameterDevice->neighborY,
+ parameterDevice->neighborZ,
+ parameterDevice->numberOfNodes,
+ parameterDevice->isEvenTimestep);
getLastCudaError("BBDevice27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.cpp b/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.cpp
index 375e19c7113e13a8b71c473b3753354508ec4725..d6a53ba59cd0222293eccbd64a94d75c89481a85 100644
--- a/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.cpp
+++ b/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.cpp
@@ -43,6 +43,7 @@
LBKernelManager::LBKernelManager(SPtr<Parameter> parameter, BoundaryConditionFactory* bcFactory): para(parameter)
{
this->velocityBoundaryConditionPost = bcFactory->getVelocityBoundaryConditionPost();
+ this->noSlipBoundaryConditionPost = bcFactory->getNoSlipBoundaryConditionPost();
}
void LBKernelManager::runLBMKernel(const int level) const
@@ -135,8 +136,6 @@ void LBKernelManager::runVelocityBCKernelPost(const int level) const
{
velocityBoundaryConditionPost(para->getParD(level).get(), &(para->getParD(level)->velocityBC));
- //QVelDevicePlainBB27(para->getParD(level).get(), &(para->getParD(level)->velocityBC));
-
//////////////////////////////////////////////////////////////////////////
// D E P R E C A T E D
//////////////////////////////////////////////////////////////////////////
@@ -653,48 +652,7 @@ void LBKernelManager::runSlipBCKernel(const int level) const{
void LBKernelManager::runNoSlipBCKernel(const int level) const{
if (para->getParD(level)->noSlipBC.numberOfBCnodes > 0)
{
- // QDev27(
- // para->getParD(level)->numberofthreads,
- // para->getParD(level)->nx,
- // para->getParD(level)->ny,
- // para->getParD(level)->distributions.f[0],
- // para->getParD(level)->noSlipBC.k,
- // para->getParD(level)->noSlipBC.q27[0],
- // para->getParD(level)->noSlipBC.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);
-
- // BBDev27(
- // para->getParD(level)->numberofthreads,
- // para->getParD(level)->nx,
- // para->getParD(level)->ny,
- // para->getParD(level)->distributions.f[0],
- // para->getParD(level)->noSlipBC.k,
- // para->getParD(level)->noSlipBC.q27[0],
- // para->getParD(level)->noSlipBC.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);
-
- QDevComp27(
- para->getParD(level)->numberofthreads,
- para->getParD(level)->distributions.f[0],
- para->getParD(level)->noSlipBC.k,
- para->getParD(level)->noSlipBC.q27[0],
- para->getParD(level)->noSlipBC.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);
+ noSlipBoundaryConditionPost(para->getParD(level).get(), &(para->getParD(level)->noSlipBC));
}
}
diff --git a/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.h b/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.h
index b98469907e52241bed21758cbddaa126d97a2ebc..08a657d2dc3a0713e18ff4942dbdcab7a11490fb 100644
--- a/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.h
+++ b/src/gpu/VirtualFluids_GPU/KernelManager/LBKernelManager.h
@@ -98,5 +98,6 @@ private:
SPtr<Parameter> para;
boundaryCondition velocityBoundaryConditionPost;
+ boundaryCondition noSlipBoundaryConditionPost;
};
#endif