diff --git a/.gitignore b/.gitignore
index 182a78aa1f85b2b69ac85452b2d87fbd87e68e56..d0a459bb12205d23608e5afd55251598e6a603ab 100644
--- a/.gitignore
+++ b/.gitignore
@@ -13,13 +13,20 @@ __pycache__/
# IDE
.vscode/
+.vscode-server/
.devcontainer/
.sync/
.idea/
+# Simulation input
+stl/
+
# Simulation results
output/
logs/
+# Scripts
+scripts/
+
# MacOS
.DS_Store
diff --git a/apps/gpu/FlowAroundSphere/CMakeLists.txt b/apps/gpu/FlowAroundSphere/CMakeLists.txt
new file mode 100644
index 0000000000000000000000000000000000000000..187c9cc57539ca1a6c74642612481c9cc7040237
--- /dev/null
+++ b/apps/gpu/FlowAroundSphere/CMakeLists.txt
@@ -0,0 +1,3 @@
+project(LidDrivenCavityGPU LANGUAGES CUDA CXX)
+
+vf_add_library(BUILDTYPE binary PRIVATE_LINK basics GridGenerator VirtualFluids_GPU GksMeshAdapter GksGpu FILES FlowAroundSphere.cpp)
diff --git a/apps/gpu/FlowAroundSphere/FlowAroundSphere.cpp b/apps/gpu/FlowAroundSphere/FlowAroundSphere.cpp
new file mode 100644
index 0000000000000000000000000000000000000000..10dfd23bafd08b4b31c153ebd44b527c12aaced6
--- /dev/null
+++ b/apps/gpu/FlowAroundSphere/FlowAroundSphere.cpp
@@ -0,0 +1,201 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \file LidDrivenCavity.cpp
+//! \ingroup Applications
+//! \author Martin Schoenherr, Stephan Lenz, Anna Wellmann
+//=======================================================================================
+#define _USE_MATH_DEFINES
+#include <exception>
+#include <fstream>
+#include <iostream>
+#include <math.h>
+#include <memory>
+#include <sstream>
+#include <stdexcept>
+#include <string>
+
+//////////////////////////////////////////////////////////////////////////
+
+#include "Core/DataTypes.h"
+#include "Core/LbmOrGks.h"
+#include "Core/Logger/Logger.h"
+#include "Core/VectorTypes.h"
+#include "PointerDefinitions.h"
+
+//////////////////////////////////////////////////////////////////////////
+
+#include "GridGenerator/grid/BoundaryConditions/Side.h"
+#include "GridGenerator/grid/GridBuilder/LevelGridBuilder.h"
+#include "GridGenerator/grid/GridBuilder/MultipleGridBuilder.h"
+#include "GridGenerator/grid/GridFactory.h"
+
+#include "GridGenerator/geometries/Sphere/Sphere.h"
+
+//////////////////////////////////////////////////////////////////////////
+
+#include "VirtualFluids_GPU/DataStructureInitializer/GridProvider.h"
+#include "VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.h"
+#include "VirtualFluids_GPU/GPU/CudaMemoryManager.h"
+#include "VirtualFluids_GPU/LBM/Simulation.h"
+#include "VirtualFluids_GPU/Output/FileWriter.h"
+#include "VirtualFluids_GPU/Parameter/Parameter.h"
+
+//////////////////////////////////////////////////////////////////////////
+
+int main(int argc, char *argv[])
+{
+ try {
+ //////////////////////////////////////////////////////////////////////////
+ // Simulation parameters
+ //////////////////////////////////////////////////////////////////////////
+ std::string outputPath("./output/Sphere");
+ std::string simulationName("Sphere");
+
+ const real L = 1.0;
+ const real Re = 1000.0;
+ const real velocity = 1.0;
+ const real dt = (real)0.5e-3;
+ const uint nx = 64;
+
+ const uint timeStepOut = 10;
+ const uint timeStepEnd = 100;
+
+ //////////////////////////////////////////////////////////////////////////
+ // setup logger
+ //////////////////////////////////////////////////////////////////////////
+
+ logging::Logger::addStream(&std::cout);
+ logging::Logger::setDebugLevel(logging::Logger::Level::INFO_LOW);
+ logging::Logger::timeStamp(logging::Logger::ENABLE);
+ logging::Logger::enablePrintedRankNumbers(logging::Logger::ENABLE);
+
+ //////////////////////////////////////////////////////////////////////////
+ // setup gridGenerator
+ //////////////////////////////////////////////////////////////////////////
+
+ auto gridBuilder = MultipleGridBuilder::makeShared();
+
+ //////////////////////////////////////////////////////////////////////////
+ // create grid
+ //////////////////////////////////////////////////////////////////////////
+
+ real dx = L / real(nx);
+
+ gridBuilder->addCoarseGrid(-1.0 * L, -1.0 * L, -1.0 * L, 1.0 * L, 1.0 * L, 1.0 * L, dx);
+
+ // use primitive
+ Object *sphere = new Sphere(0.0, 0.0, 0.0, 0.2);
+ // use stl
+ // Object* sphere = TriangularMesh::make("stl/sphere.stl");
+ gridBuilder->addGeometry(sphere);
+
+ gridBuilder->setPeriodicBoundaryCondition(false, false, false);
+
+ gridBuilder->buildGrids(LBM, false);
+
+ //////////////////////////////////////////////////////////////////////////
+ // setup simulation parameters
+ //////////////////////////////////////////////////////////////////////////
+
+ SPtr<Parameter> para = Parameter::make();
+
+ //////////////////////////////////////////////////////////////////////////
+ // compute parameters in lattice units
+ //////////////////////////////////////////////////////////////////////////
+
+ const real velocityLB = velocity * dt / dx; // LB units
+
+ const real vxLB = velocityLB / sqrt(2.0); // LB units
+
+ const real viscosityLB = nx * velocityLB / Re; // LB units
+
+ *logging::out << logging::Logger::INFO_HIGH << "velocity [dx/dt] = " << velocityLB << " \n";
+ *logging::out << logging::Logger::INFO_HIGH << "viscosity [dx^2/dt] = " << viscosityLB << "\n";
+
+ //////////////////////////////////////////////////////////////////////////
+ // set parameters
+ //////////////////////////////////////////////////////////////////////////
+
+ para->setOutputPath(outputPath);
+ para->setOutputPrefix(simulationName);
+
+ para->setPathAndFilename(para->getOutputPath() + "/" + para->getOutputPrefix());
+
+ para->setPrintFiles(true);
+
+ para->setVelocityLB(velocityLB);
+ para->setViscosityLB(viscosityLB);
+
+ para->setVelocityRatio(velocity / velocityLB);
+
+ para->setTimestepOut(timeStepOut);
+ para->setTimestepEnd(timeStepEnd);
+
+ //////////////////////////////////////////////////////////////////////////
+ // set boundary conditions
+ //////////////////////////////////////////////////////////////////////////
+
+ gridBuilder->setVelocityBoundaryCondition(SideType::MX, vxLB, 0.0, 0.0);
+ gridBuilder->setVelocityBoundaryCondition(SideType::MY, 0.0, 0.0, 0.0);
+ gridBuilder->setVelocityBoundaryCondition(SideType::PY, vxLB, 0.0, 0.0);
+ gridBuilder->setVelocityBoundaryCondition(SideType::MZ, vxLB, 0.0, 0.0);
+ gridBuilder->setVelocityBoundaryCondition(SideType::PZ, vxLB, 0.0, 0.0);
+ gridBuilder->setVelocityBoundaryCondition(SideType::GEOMETRY, 0.0, 0.0, 0.0);
+ gridBuilder->setPressureBoundaryCondition(SideType::PX, 0.0); // set pressure BC after velocity BCs
+
+ //////////////////////////////////////////////////////////////////////////
+ // set copy mesh to simulation
+ //////////////////////////////////////////////////////////////////////////
+
+ SPtr<CudaMemoryManager> cudaMemoryManager = CudaMemoryManager::make(para);
+
+ SPtr<GridProvider> gridGenerator = GridProvider::makeGridGenerator(gridBuilder, para, cudaMemoryManager);
+
+ //////////////////////////////////////////////////////////////////////////
+ // run simulation
+ //////////////////////////////////////////////////////////////////////////
+
+ Simulation sim;
+ SPtr<FileWriter> fileWriter = SPtr<FileWriter>(new FileWriter());
+ sim.init(para, gridGenerator, fileWriter, cudaMemoryManager);
+ sim.run();
+ sim.free();
+
+ } catch (const std::bad_alloc &e) {
+ *logging::out << logging::Logger::LOGGER_ERROR << "Bad Alloc:" << e.what() << "\n";
+ } catch (const std::exception &e) {
+ *logging::out << logging::Logger::LOGGER_ERROR << e.what() << "\n";
+ } catch (std::string &s) {
+ *logging::out << logging::Logger::LOGGER_ERROR << s << "\n";
+ } catch (...) {
+ *logging::out << logging::Logger::LOGGER_ERROR << "Unknown exception!\n";
+ }
+
+ return 0;
+}
diff --git a/gpu.cmake b/gpu.cmake
index ca3929cc1fdc5d15af4b82cd1158ed5b0ad93d10..440f0e3d5846a3af926291a9f463873140f269f4 100644
--- a/gpu.cmake
+++ b/gpu.cmake
@@ -16,6 +16,7 @@ add_subdirectory(src/gpu/GksMeshAdapter)
add_subdirectory(src/gpu/GksGpu)
add_subdirectory(apps/gpu/LidDrivenCavityGPU)
+add_subdirectory(apps/gpu/FlowAroundSphere)
add_subdirectory(apps/gpu/DiluteGravityCurrents_Case1)
add_subdirectory(apps/gpu/DiluteGravityCurrents_Case2)
add_subdirectory(apps/gpu/DiluteGravityCurrents_Case3)
diff --git a/src/gpu/GridGenerator/grid/BoundaryConditions/Side.cpp b/src/gpu/GridGenerator/grid/BoundaryConditions/Side.cpp
index 02d61a83456090a3b96b207de5761b70d9b30fca..4c5161acb778c05402301a10462220aeea5e73f4 100644
--- a/src/gpu/GridGenerator/grid/BoundaryConditions/Side.cpp
+++ b/src/gpu/GridGenerator/grid/BoundaryConditions/Side.cpp
@@ -101,12 +101,32 @@ void Side::setQs(SPtr<Grid> grid, SPtr<BoundaryCondition> boundaryCondition, uin
real x,y,z;
grid->transIndexToCoords( index, x, y, z );
- x += grid->getDirection()[dir * DIMENSION + 0] * grid->getDelta();
- y += grid->getDirection()[dir * DIMENSION + 1] * grid->getDelta();
- z += grid->getDirection()[dir * DIMENSION + 2] * grid->getDelta();
+ real coords[3] = {x,y,z};
- uint neighborIndex = grid->transCoordToIndex( x, y, z );
+ real neighborX = x + grid->getDirection()[dir * DIMENSION + 0] * grid->getDelta();
+ real neighborY = y + grid->getDirection()[dir * DIMENSION + 1] * grid->getDelta();
+ real neighborZ = z + grid->getDirection()[dir * DIMENSION + 2] * grid->getDelta();
+ // correct neighbor coordinates in case of periodic boundaries
+ if( grid->getPeriodicityX() && grid->getFieldEntry( grid->transCoordToIndex( neighborX, y, z ) ) == vf::gpu::STOPPER_OUT_OF_GRID_BOUNDARY )
+ {
+ if( neighborX > x ) neighborX = grid->getFirstFluidNode( coords, 0, grid->getStartX() );
+ else neighborX = grid->getLastFluidNode ( coords, 0, grid->getEndX() );
+ }
+
+ if( grid->getPeriodicityY() && grid->getFieldEntry( grid->transCoordToIndex( x, neighborY, z ) ) == vf::gpu::STOPPER_OUT_OF_GRID_BOUNDARY )
+ {
+ if( neighborY > y ) neighborY = grid->getFirstFluidNode( coords, 1, grid->getStartY() );
+ else neighborY = grid->getLastFluidNode ( coords, 1, grid->getEndY() );
+ }
+
+ if( grid->getPeriodicityZ() && grid->getFieldEntry( grid->transCoordToIndex( x, y, neighborZ ) ) == vf::gpu::STOPPER_OUT_OF_GRID_BOUNDARY )
+ {
+ if( neighborZ > z ) neighborZ = grid->getFirstFluidNode( coords, 2, grid->getStartZ() );
+ else neighborZ = grid->getLastFluidNode ( coords, 2, grid->getEndZ() );
+ }
+
+ uint neighborIndex = grid->transCoordToIndex( neighborX, neighborY, neighborZ );
if( grid->getFieldEntry(neighborIndex) == vf::gpu::STOPPER_OUT_OF_GRID_BOUNDARY ||
grid->getFieldEntry(neighborIndex) == vf::gpu::STOPPER_OUT_OF_GRID ||
grid->getFieldEntry(neighborIndex) == vf::gpu::STOPPER_SOLID )
@@ -129,6 +149,44 @@ uint Side::getIndex(SPtr<Grid> grid, std::string coord, real constant, real v1,
return INVALID_INDEX;
}
+void Geometry::addIndices(std::vector<SPtr<Grid> > grids, uint level, SPtr<BoundaryCondition> boundaryCondition)
+{
+ auto geometryBoundaryCondition = std::dynamic_pointer_cast<GeometryBoundaryCondition>(boundaryCondition);
+
+ std::vector<real> qNode(grids[level]->getEndDirection() + 1);
+
+ for (uint index = 0; index < grids[level]->getSize(); index++)
+ {
+ if (grids[level]->getFieldEntry(index) != vf::gpu::BC_SOLID)
+ continue;
+
+ for (int dir = 0; dir <= grids[level]->getEndDirection(); dir++)
+ {
+ const real q = grids[level]->getQValue(index, dir);
+
+ qNode[dir] = q;
+
+ // also the neighbor if any Qs are required
+ real x,y,z;
+ grids[level]->transIndexToCoords( index, x, y, z );
+
+ x += grids[level]->getDirection()[dir * DIMENSION + 0] * grids[level]->getDelta();
+ y += grids[level]->getDirection()[dir * DIMENSION + 1] * grids[level]->getDelta();
+ z += grids[level]->getDirection()[dir * DIMENSION + 2] * grids[level]->getDelta();
+
+ uint neighborIndex = grids[level]->transCoordToIndex( x, y, z );
+
+ if( qNode[dir] < -0.5 && ( grids[level]->getFieldEntry(neighborIndex) == vf::gpu::STOPPER_OUT_OF_GRID_BOUNDARY ||
+ grids[level]->getFieldEntry(neighborIndex) == vf::gpu::STOPPER_OUT_OF_GRID ||
+ grids[level]->getFieldEntry(neighborIndex) == vf::gpu::STOPPER_SOLID ) )
+ qNode[dir] = 0.5;
+ }
+
+ geometryBoundaryCondition->indices.push_back(index);
+ geometryBoundaryCondition->qs.push_back(qNode);
+ geometryBoundaryCondition->patches.push_back(grids[level]->getQPatch(index));
+ }
+}
void MX::addIndices(std::vector<SPtr<Grid> > grid, uint level, SPtr<BoundaryCondition> boundaryCondition)
{
diff --git a/src/gpu/GridGenerator/grid/BoundaryConditions/Side.h b/src/gpu/GridGenerator/grid/BoundaryConditions/Side.h
index 260bc18967296e63cc21fc0454c86cdcb65e8719..d4c9e3a4bcab73d368c863ee57d66f692126fa06 100644
--- a/src/gpu/GridGenerator/grid/BoundaryConditions/Side.h
+++ b/src/gpu/GridGenerator/grid/BoundaryConditions/Side.h
@@ -84,6 +84,27 @@ private:
static uint getIndex(SPtr<Grid> grid, std::string coord, real constant, real v1, real v2);
};
+class Geometry : public Side
+{
+public:
+ void addIndices(std::vector<SPtr<Grid> > grid, uint level, SPtr<gg::BoundaryCondition> boundaryCondition) override;
+
+ int getCoordinate() const override
+ {
+ return X_INDEX;
+ }
+
+ int getDirection() const override
+ {
+ return NEGATIVE_DIR;
+ }
+
+ SideType whoAmI() const override
+ {
+ return SideType::GEOMETRY;
+ }
+};
+
class MX : public Side
{
public:
@@ -233,7 +254,7 @@ public:
case SideType::PZ:
return SPtr<Side>(new PZ());
case SideType::GEOMETRY:
- throw std::runtime_error("SideFactory::make() - SideType::GEOMETRY not supported.");
+ return SPtr<Side>(new Geometry());
default:
throw std::runtime_error("SideFactory::make() - SideType not valid.");
}
diff --git a/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.cpp b/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.cpp
index 20c285c13aa2c7644fc1ac6ea6a87b5d8dc7f72c..aa593ad7ffc4daecb3a05ab0cc41865a09594933 100644
--- a/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.cpp
+++ b/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.cpp
@@ -28,7 +28,7 @@
//
//! \file LevelGridBuilder.cpp
//! \ingroup grid
-//! \author Soeren Peters, Stephan Lenz, Martin Schönherr
+//! \author Soeren Peters, Stephan Lenz, Martin Sch�nherr
//=======================================================================================
#include "LevelGridBuilder.h"
@@ -92,37 +92,42 @@ void LevelGridBuilder::setSlipBoundaryCondition(SideType sideType, real nomalX,
void LevelGridBuilder::setVelocityBoundaryCondition(SideType sideType, real vx, real vy, real vz)
{
- SPtr<VelocityBoundaryCondition> velocityBoundaryCondition = VelocityBoundaryCondition::make(vx, vy, vz);
+ if (sideType == SideType::GEOMETRY)
+ setVelocityGeometryBoundaryCondition(vx, vy, vz);
+ else
+ {
+ SPtr<VelocityBoundaryCondition> velocityBoundaryCondition = VelocityBoundaryCondition::make(vx, vy, vz);
- auto side = SideFactory::make(sideType);
+ auto side = SideFactory::make(sideType);
- velocityBoundaryCondition->side = side;
- velocityBoundaryCondition->side->addIndices(grids, 0, velocityBoundaryCondition);
+ velocityBoundaryCondition->side = side;
+ velocityBoundaryCondition->side->addIndices(grids, 0, velocityBoundaryCondition);
- velocityBoundaryCondition->fillVelocityLists();
+ velocityBoundaryCondition->fillVelocityLists();
- boundaryConditions[0]->velocityBoundaryConditions.push_back(velocityBoundaryCondition);
+ boundaryConditions[0]->velocityBoundaryConditions.push_back(velocityBoundaryCondition);
- *logging::out << logging::Logger::INFO_INTERMEDIATE << "Set Velocity BC on level " << 0 << " with " << (int)velocityBoundaryCondition->indices.size() <<"\n";
+ *logging::out << logging::Logger::INFO_INTERMEDIATE << "Set Velocity BC on level " << 0 << " with " << (int)velocityBoundaryCondition->indices.size() <<"\n";
+ }
}
void LevelGridBuilder::setVelocityGeometryBoundaryCondition(real vx, real vy, real vz)
{
geometryHasValues = true;
-
+
for (uint level = 0; level < getNumberOfGridLevels(); level++)
{
- if (boundaryConditions[level]->geometryBoundaryCondition != nullptr)
- {
- boundaryConditions[level]->geometryBoundaryCondition->vx = vx;
- boundaryConditions[level]->geometryBoundaryCondition->vy = vy;
- boundaryConditions[level]->geometryBoundaryCondition->vz = vz;
- boundaryConditions[level]->geometryBoundaryCondition->side->addIndices(grids, level, boundaryConditions[level]->geometryBoundaryCondition);
+ if (boundaryConditions[level]->geometryBoundaryCondition != nullptr)
+ {
+ boundaryConditions[level]->geometryBoundaryCondition->vx = vx;
+ boundaryConditions[level]->geometryBoundaryCondition->vy = vy;
+ boundaryConditions[level]->geometryBoundaryCondition->vz = vz;
+ boundaryConditions[level]->geometryBoundaryCondition->side->addIndices(grids, level, boundaryConditions[level]->geometryBoundaryCondition);
boundaryConditions[level]->geometryBoundaryCondition->fillVelocityLists();
*logging::out << logging::Logger::INFO_INTERMEDIATE << "Set Geometry Velocity BC on level " << level << " with " << (int)boundaryConditions[level]->geometryBoundaryCondition->indices.size() <<"\n";
- }
+ }
}
}
@@ -159,7 +164,10 @@ void LevelGridBuilder::setNoSlipBoundaryCondition(SideType sideType)
noSlipBoundaryCondition->side = side;
noSlipBoundaryCondition->side->addIndices(grids, level, noSlipBoundaryCondition);
- boundaryConditions[level]->noSlipBoundaryConditions.push_back(noSlipBoundaryCondition);
+ noSlipBoundaryCondition->fillVelocityLists();
+
+ // effectively just a wrapper for velocityBC with zero velocity. No distinction in GridGenerator.
+ boundaryConditions[level]->velocityBoundaryConditions.push_back(noSlipBoundaryCondition);
}
}
@@ -469,9 +477,9 @@ void LevelGridBuilder::getGeometryValues(real* vx, real* vy, real* vz, int level
{
for (uint i = 0; i < boundaryConditions[level]->geometryBoundaryCondition->indices.size(); i++)
{
- vx[i] = boundaryConditions[level]->geometryBoundaryCondition->getVx(i);
- vy[i] = boundaryConditions[level]->geometryBoundaryCondition->getVy(i);
- vz[i] = boundaryConditions[level]->geometryBoundaryCondition->getVz(i);
+ vx[i] = boundaryConditions[level]->geometryBoundaryCondition->getVx(i);
+ vy[i] = boundaryConditions[level]->geometryBoundaryCondition->getVy(i);
+ vz[i] = boundaryConditions[level]->geometryBoundaryCondition->getVz(i);
}
}
diff --git a/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.h b/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.h
index f2325435d99140f33eee9844c13908de87788558..b4a461ec764d0a97efd83ec7f4899814fe99a489 100644
--- a/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.h
+++ b/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.h
@@ -28,7 +28,7 @@
//
//! \file LevelGridBuilder.h
//! \ingroup grid
-//! \author Soeren Peters, Stephan Lenz, Martin Schönherr
+//! \author Soeren Peters, Stephan Lenz, Martin Sch�nherr
//=======================================================================================
#ifndef LEVEL_GRID_BUILDER_H
#define LEVEL_GRID_BUILDER_H
@@ -131,7 +131,7 @@ protected:
std::vector<SPtr<PressureBoundaryCondition>> pressureBoundaryConditions;
- std::vector<SPtr<VelocityBoundaryCondition> > noSlipBoundaryConditions;
+ std::vector<SPtr<VelocityBoundaryCondition>> noSlipBoundaryConditions;
SPtr<GeometryBoundaryCondition> geometryBoundaryCondition;
};
diff --git a/src/gpu/GridGenerator/grid/GridBuilder/MultipleGridBuilder.cpp b/src/gpu/GridGenerator/grid/GridBuilder/MultipleGridBuilder.cpp
index 565a02a2807ef15679bf08fa001ea9a03f78ca4e..c36a8cb70e6bb36dccb9179b9c014e5fb7464a30 100644
--- a/src/gpu/GridGenerator/grid/GridBuilder/MultipleGridBuilder.cpp
+++ b/src/gpu/GridGenerator/grid/GridBuilder/MultipleGridBuilder.cpp
@@ -28,7 +28,7 @@
//
//! \file MultipleGridBuilder.cpp
//! \ingroup grid
-//! \author Soeren Peters, Stephan Lenz, Martin Schönherr
+//! \author Soeren Peters, Stephan Lenz, Martin Sch�nherr
//=======================================================================================
#include "MultipleGridBuilder.h"
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp
index af1e123e513e56dd4d500fee1f98bbd05c5f6240..48fe29321f5f884a2e73118bab1ea352fcf5f6ec 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp
@@ -80,8 +80,8 @@ void GridProvider::setInitalNodeValues(const int numberOfNodes) const
void GridProvider::setVelocitySize(int size) const
{
- para->getParH()->numberOfInflowBCnodes = size;
- para->getParD()->numberOfInflowBCnodes = size;
+ para->getParH()->numberOfVeloBCnodes = size;
+ para->getParD()->numberOfVeloBCnodes = size;
}
void GridProvider::allocAndCopyForcing()
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
index 75458df1c59a3a032b634cd12cd24ff2fd456a80..a68e8543f8673329f51be9de97d108e6eba4233f 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
@@ -119,20 +119,20 @@ void GridGenerator::allocArrays_BoundaryValues()
std::cout << "size velocity: " << numberOfVelocityValues << std::endl;
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
blocks = (numberOfVelocityValues / para->getParH()->numberofthreads) + 1;
- para->getParH()->inflowBC.kArray = blocks * para->getParH()->numberofthreads;
- para->getParD()->inflowBC.kArray = para->getParH()->inflowBC.kArray;
+ para->getParH()->veloBC.kArray = blocks * para->getParH()->numberofthreads;
+ para->getParD()->veloBC.kArray = para->getParH()->veloBC.kArray;
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
- para->getParH()->numberOfInflowBCnodes = numberOfVelocityValues;
- para->getParD()->numberOfInflowBCnodes = numberOfVelocityValues;
+ para->getParH()->numberOfVeloBCnodes = numberOfVelocityValues;
+ para->getParD()->numberOfVeloBCnodes = numberOfVelocityValues;
if (numberOfVelocityValues > 1) {
cudaMemoryManager->cudaAllocVeloBC();
builder->getVelocityValues(
- para->getParH()->inflowBC.Vx,
- para->getParH()->inflowBC.Vy,
- para->getParH()->inflowBC.Vz,
- para->getParH()->inflowBC.k, 0);
+ para->getParH()->veloBC.Vx,
+ para->getParH()->veloBC.Vy,
+ para->getParH()->veloBC.Vz,
+ para->getParH()->veloBC.k, 0);
cudaMemoryManager->cudaCopyVeloBC();
}
@@ -192,8 +192,8 @@ void GridGenerator::allocArrays_BoundaryQs()
std::cout << "size velocity: " << numberOfVelocityNodes << std::endl;
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
// preprocessing
- real *QQ = para->getParH()->inflowBC.q27[0];
- unsigned int sizeQ = para->getParH()->numberOfInflowBCnodes;
+ real *QQ = para->getParH()->veloBC.q27[0];
+ unsigned int sizeQ = para->getParH()->numberOfVeloBCnodes;
QforBoundaryConditions Q;
Q.q27[dirE] = &QQ[dirE * sizeQ];
Q.q27[dirW] = &QQ[dirW * sizeQ];
diff --git a/src/gpu/VirtualFluids_GPU/GPU/CudaKernelManager.cpp b/src/gpu/VirtualFluids_GPU/GPU/CudaKernelManager.cpp
index 451174fa70031aca85c806f72f79e378d667220d..3202144810c83e913c2bf5881b967bf78b4675e2 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/CudaKernelManager.cpp
+++ b/src/gpu/VirtualFluids_GPU/GPU/CudaKernelManager.cpp
@@ -69,18 +69,18 @@ void CudaKernelManager::runLBMKernel(SPtr<Parameter> para)
void CudaKernelManager::runVelocityBCKernel(SPtr<Parameter> para)
{
- if (para->getParD()->numberOfInflowBCnodes > 0)
+ if (para->getParD()->numberOfVeloBCnodes > 0)
{
QVelDevicePlainBB27(
para->getParD()->numberofthreads,
- para->getParD()->inflowBC.Vx,
- para->getParD()->inflowBC.Vy,
- para->getParD()->inflowBC.Vz,
+ para->getParD()->veloBC.Vx,
+ para->getParD()->veloBC.Vy,
+ para->getParD()->veloBC.Vz,
para->getParD()->distributions.f[0],
- para->getParD()->inflowBC.k,
- para->getParD()->inflowBC.q27[0],
- para->getParD()->numberOfInflowBCnodes,
- para->getParD()->inflowBC.kArray,
+ para->getParD()->veloBC.k,
+ para->getParD()->veloBC.q27[0],
+ para->getParD()->numberOfVeloBCnodes,
+ para->getParD()->veloBC.kArray,
para->getParD()->neighborX,
para->getParD()->neighborY,
para->getParD()->neighborZ,
@@ -89,6 +89,14 @@ void CudaKernelManager::runVelocityBCKernel(SPtr<Parameter> para)
}
}
+void CudaKernelManager::runGeoBCKernel(SPtr<Parameter> para)
+{
+ if (para->getParD()->numberOfVeloBCnodes > 0)
+ {
+ // ...
+ }
+}
+
void CudaKernelManager::calculateMacroscopicValues(SPtr<Parameter> para)
{
if (para->getIsADcalculationOn()) {
diff --git a/src/gpu/VirtualFluids_GPU/GPU/CudaKernelManager.h b/src/gpu/VirtualFluids_GPU/GPU/CudaKernelManager.h
index c257875521318615e2b6e59537de8bbb45ad0b70..622bdf135b08ec74b3b45936312ab32e6ac603e6 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/CudaKernelManager.h
+++ b/src/gpu/VirtualFluids_GPU/GPU/CudaKernelManager.h
@@ -55,6 +55,9 @@ public:
//! \brief calls the device function of the velocity boundary condition
void runVelocityBCKernel(SPtr<Parameter> para);
+ //! \brief calls the device function of the geometry boundary condition
+ void runGeoBCKernel(SPtr<Parameter> para);
+
//! \brief calls the device function that calculates the macroscopic values
void calculateMacroscopicValues(SPtr<Parameter> para);
diff --git a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
index 73c6d1acc10119b1d3de46c540a5f8fa785fd4f8..364843deb94fd6730940055f5cba10179f057265 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
+++ b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
@@ -192,22 +192,22 @@ void CudaMemoryManager::cudaFreeSlipBC()
//velocity boundary condition
void CudaMemoryManager::cudaAllocVeloBC()
{
- unsigned int mem_size_inflow_BC_INT = sizeof(int)*parameter->getParH()->numberOfInflowBCnodes;
- unsigned int mem_size_inflow_BC_REAL = sizeof(real)*parameter->getParH()->numberOfInflowBCnodes;
+ unsigned int mem_size_inflow_BC_INT = sizeof(int)*parameter->getParH()->numberOfVeloBCnodes;
+ unsigned int mem_size_inflow_BC_REAL = sizeof(real)*parameter->getParH()->numberOfVeloBCnodes;
//Host
- checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH()->inflowBC.q27[0]), parameter->getD3Qxx() * mem_size_inflow_BC_REAL ));
- checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH()->inflowBC.k), mem_size_inflow_BC_INT ));
- checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH()->inflowBC.Vx), mem_size_inflow_BC_REAL ));
- checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH()->inflowBC.Vy), mem_size_inflow_BC_REAL ));
- checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH()->inflowBC.Vz), mem_size_inflow_BC_REAL ));
+ checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH()->veloBC.q27[0]), parameter->getD3Qxx() * mem_size_inflow_BC_REAL ));
+ checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH()->veloBC.k), mem_size_inflow_BC_INT ));
+ checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH()->veloBC.Vx), mem_size_inflow_BC_REAL ));
+ checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH()->veloBC.Vy), mem_size_inflow_BC_REAL ));
+ checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH()->veloBC.Vz), mem_size_inflow_BC_REAL ));
//Device
- checkCudaErrors( cudaMalloc((void**) &(parameter->getParD()->inflowBC.q27[0]), parameter->getD3Qxx() * mem_size_inflow_BC_REAL ));
- checkCudaErrors( cudaMalloc((void**) &(parameter->getParD()->inflowBC.k), mem_size_inflow_BC_INT ));
- checkCudaErrors( cudaMalloc((void**) &(parameter->getParD()->inflowBC.Vx), mem_size_inflow_BC_REAL ));
- checkCudaErrors( cudaMalloc((void**) &(parameter->getParD()->inflowBC.Vy), mem_size_inflow_BC_REAL ));
- checkCudaErrors( cudaMalloc((void**) &(parameter->getParD()->inflowBC.Vz), mem_size_inflow_BC_REAL ));
+ checkCudaErrors( cudaMalloc((void**) &(parameter->getParD()->veloBC.q27[0]), parameter->getD3Qxx() * mem_size_inflow_BC_REAL ));
+ checkCudaErrors( cudaMalloc((void**) &(parameter->getParD()->veloBC.k), mem_size_inflow_BC_INT ));
+ checkCudaErrors( cudaMalloc((void**) &(parameter->getParD()->veloBC.Vx), mem_size_inflow_BC_REAL ));
+ checkCudaErrors( cudaMalloc((void**) &(parameter->getParD()->veloBC.Vy), mem_size_inflow_BC_REAL ));
+ checkCudaErrors( cudaMalloc((void**) &(parameter->getParD()->veloBC.Vz), mem_size_inflow_BC_REAL ));
//////////////////////////////////////////////////////////////////////////
double tmp = (double)mem_size_inflow_BC_INT + 4. * (double)mem_size_inflow_BC_REAL + (double)parameter->getD3Qxx() * (double)mem_size_inflow_BC_REAL;
@@ -215,23 +215,23 @@ void CudaMemoryManager::cudaAllocVeloBC()
}
void CudaMemoryManager::cudaCopyVeloBC()
{
- unsigned int mem_size_inflow_BC_INT = sizeof(int)*parameter->getParH()->numberOfInflowBCnodes;
- unsigned int mem_size_inflow_BC_REAL = sizeof(real)*parameter->getParH()->numberOfInflowBCnodes;
+ unsigned int mem_size_inflow_BC_INT = sizeof(int)*parameter->getParH()->numberOfVeloBCnodes;
+ unsigned int mem_size_inflow_BC_REAL = sizeof(real)*parameter->getParH()->numberOfVeloBCnodes;
- checkCudaErrors( cudaMemcpy(parameter->getParD()->inflowBC.q27[0], parameter->getParH()->inflowBC.q27[0], parameter->getD3Qxx() * mem_size_inflow_BC_REAL, cudaMemcpyHostToDevice));
- checkCudaErrors( cudaMemcpy(parameter->getParD()->inflowBC.k, parameter->getParH()->inflowBC.k, mem_size_inflow_BC_INT , cudaMemcpyHostToDevice));
- checkCudaErrors( cudaMemcpy(parameter->getParD()->inflowBC.Vx, parameter->getParH()->inflowBC.Vx, mem_size_inflow_BC_REAL, cudaMemcpyHostToDevice));
- checkCudaErrors( cudaMemcpy(parameter->getParD()->inflowBC.Vy, parameter->getParH()->inflowBC.Vy, mem_size_inflow_BC_REAL, cudaMemcpyHostToDevice));
- checkCudaErrors( cudaMemcpy(parameter->getParD()->inflowBC.Vz, parameter->getParH()->inflowBC.Vz, mem_size_inflow_BC_REAL, cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getParD()->veloBC.q27[0], parameter->getParH()->veloBC.q27[0], parameter->getD3Qxx() * mem_size_inflow_BC_REAL, cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getParD()->veloBC.k, parameter->getParH()->veloBC.k, mem_size_inflow_BC_INT , cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getParD()->veloBC.Vx, parameter->getParH()->veloBC.Vx, mem_size_inflow_BC_REAL, cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getParD()->veloBC.Vy, parameter->getParH()->veloBC.Vy, mem_size_inflow_BC_REAL, cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getParD()->veloBC.Vz, parameter->getParH()->veloBC.Vz, mem_size_inflow_BC_REAL, cudaMemcpyHostToDevice));
}
void CudaMemoryManager::cudaFreeVeloBC()
{
- checkCudaErrors( cudaFreeHost(parameter->getParH()->inflowBC.q27[0] ));
- checkCudaErrors( cudaFreeHost(parameter->getParH()->inflowBC.k ));
- checkCudaErrors( cudaFreeHost(parameter->getParH()->inflowBC.Vx ));
- checkCudaErrors( cudaFreeHost(parameter->getParH()->inflowBC.Vy ));
- checkCudaErrors( cudaFreeHost(parameter->getParH()->inflowBC.Vz ));
+ checkCudaErrors( cudaFreeHost(parameter->getParH()->veloBC.q27[0] ));
+ checkCudaErrors( cudaFreeHost(parameter->getParH()->veloBC.k ));
+ checkCudaErrors( cudaFreeHost(parameter->getParH()->veloBC.Vx ));
+ checkCudaErrors( cudaFreeHost(parameter->getParH()->veloBC.Vy ));
+ checkCudaErrors( cudaFreeHost(parameter->getParH()->veloBC.Vz ));
}
diff --git a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
index 3e75d3c4cc29f8c963dfe4d2edc78fec884385d3..95d3f15dff5d3a31240c965a36d243096ba1a2ef 100644
--- a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
+++ b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
@@ -144,7 +144,7 @@ void Simulation::run()
////////////////////////////////////////////////////////////////////////////////
// velocity boundary condition
- if (para->getParD()->numberOfInflowBCnodes > 1)
+ if (para->getParD()->numberOfVeloBCnodes > 1)
cudaKernelManager->runVelocityBCKernel(para);
////////////////////////////////////////////////////////////////////////////////
@@ -194,7 +194,7 @@ void Simulation::free()
cudaMemoryManager->cudaFreeSP();
cudaMemoryManager->cudaFreeNeighborWSB();
- if (para->getParH()->numberOfInflowBCnodes > 1)
+ if (para->getParH()->numberOfVeloBCnodes > 1)
this->cudaMemoryManager->cudaFreeVeloBC();
if (para->getParH()->numberOfSlipBCnodes > 1)
diff --git a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp
index 14ed3fdbb63e0061b835d5f9235a442366dccc8f..b8598910a62198afe13251d34b19950dca53fe04 100644
--- a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp
+++ b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp
@@ -43,41 +43,41 @@ SPtr<Parameter> Parameter::make()
Parameter::Parameter()
{
- this->setOutputPath("C:/Output/");
+ this->setOutputPath("C:/Output/");
- this->setOutputPrefix("MyFile");
+ this->setOutputPrefix("MyFile");
- this->setPrintFiles(false);
+ this->setPrintFiles(false);
- this->setD3Qxx((int)27);
+ this->setD3Qxx((int)27);
- this->setTimestepEnd((uint)10);
+ this->setTimestepEnd((uint)10);
- this->setTimestepOut((uint)1);
+ this->setTimestepOut((uint)1);
- this->setTimestepStartOut((uint)0);
+ this->setTimestepStartOut((uint)0);
- this->setViscosityLB((real)0.001);
+ this->setViscosityLB((real)0.001);
- this->setVelocityLB((real)0.01);
+ this->setVelocityLB((real)0.01);
this->setSc((real)1.0);
this->setG_r((real)0.0);
- this->setViscosityRatio((real)1.0);
+ this->setViscosityRatio((real)1.0);
- this->setVelocityRatio((real)1.0);
+ this->setVelocityRatio((real)1.0);
- this->setDensityRatio((real)1.0);
+ this->setDensityRatio((real)1.0);
- this->setPressureRatio((real)1.0);
+ this->setPressureRatio((real)1.0);
- this->setPathAndFilename(this->getOutputPath() + "/" + this->getOutputPrefix());
+ this->setPathAndFilename(this->getOutputPath() + "/" + this->getOutputPrefix());
- this->setlimitOfNodesForVTK((uint)30000000);
+ this->setlimitOfNodesForVTK((uint)30000000);
- //////////////////////////////////////////////////////////////////////////
+ //////////////////////////////////////////////////////////////////////////
// Advection Diffusion
this->setIsADcalculationOn(false);
}
@@ -87,9 +87,9 @@ Parameter::~Parameter()
Parameter* Parameter::instance = 0;
Parameter* Parameter::getInstanz()
{
- if( instance == 0 )
- instance = new Parameter();
- return instance;
+ if( instance == 0 )
+ instance = new Parameter();
+ return instance;
}
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
@@ -97,19 +97,19 @@ Parameter* Parameter::getInstanz()
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
void Parameter::initParameter()
{
- //host
- parametersOnHost = new ParameterStruct;
- parametersOnHost->numberofthreads = 64;
- parametersOnHost->omega = (real)1.0/((real)3.0*this->viscosityLB+(real)0.5);
- parametersOnHost->isEvenTimestep = true;
+ //host
+ parametersOnHost = new ParameterStruct;
+ parametersOnHost->numberofthreads = 64;
+ parametersOnHost->omega = (real)1.0/((real)3.0*this->viscosityLB+(real)0.5);
+ parametersOnHost->isEvenTimestep = true;
// set diffusivity
parametersOnHost->diffusivity = (real)0.01;
- //device
- parametersOnDevice = new ParameterStruct;
- parametersOnDevice->numberofthreads = parametersOnHost->numberofthreads;
- parametersOnDevice->omega = parametersOnHost->omega;
- parametersOnDevice->isEvenTimestep = parametersOnHost->isEvenTimestep;
+ //device
+ parametersOnDevice = new ParameterStruct;
+ parametersOnDevice->numberofthreads = parametersOnHost->numberofthreads;
+ parametersOnDevice->omega = parametersOnHost->omega;
+ parametersOnDevice->isEvenTimestep = parametersOnHost->isEvenTimestep;
// set diffusivity
parametersOnDevice->diffusivity = parametersOnHost->diffusivity;
}
@@ -120,39 +120,39 @@ void Parameter::initParameter()
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
void Parameter::setlimitOfNodesForVTK(uint limitOfNodesForVTK)
{
- this->limitOfNodesForVTK = limitOfNodesForVTK;
+ this->limitOfNodesForVTK = limitOfNodesForVTK;
}
void Parameter::setD3Qxx(int d3qxx)
{
- this->D3Qxx = d3qxx;
+ this->D3Qxx = d3qxx;
}
void Parameter::setTimestepEnd(uint timestepEnd)
{
- this->timestepEnd = timestepEnd;
+ this->timestepEnd = timestepEnd;
}
void Parameter::setTimestepOut(uint timestepOut)
{
- this->timestepOut = timestepOut;
+ this->timestepOut = timestepOut;
}
void Parameter::setTimestepStartOut(uint timestepStartOut)
{
- this->timestepStartOut = timestepStartOut;
+ this->timestepStartOut = timestepStartOut;
}
void Parameter::setOutputPath(std::string outputPath)
{
- this->outputPath = outputPath;
+ this->outputPath = outputPath;
}
void Parameter::setOutputPrefix(std::string outputPrefix)
{
- this->outputPrefix = outputPrefix;
+ this->outputPrefix = outputPrefix;
}
void Parameter::setPathAndFilename(std::string pathAndFilename)
{
- this->pathAndFilename = pathAndFilename;
+ this->pathAndFilename = pathAndFilename;
}
void Parameter::setIsADcalculationOn(bool calcAD)
{
- this->calcAD = calcAD;
+ this->calcAD = calcAD;
}
void Parameter::setIsBodyForceOn(bool calcBF)
{
@@ -160,35 +160,35 @@ void Parameter::setIsBodyForceOn(bool calcBF)
}
void Parameter::setPrintFiles(bool printfiles)
{
- this->printFiles = printfiles;
+ this->printFiles = printfiles;
}
void Parameter::setViscosityLB(real viscosity)
{
- this->viscosityLB = viscosity;
+ this->viscosityLB = viscosity;
}
void Parameter::setVelocityLB(real velocity)
{
- this->velocityLB = velocity;
+ this->velocityLB = velocity;
}
void Parameter::setViscosityRatio(real viscosityRatio)
{
- this->viscosityRatio = viscosityRatio;
+ this->viscosityRatio = viscosityRatio;
}
void Parameter::setVelocityRatio(real velocityRatio)
{
- this->velocityRatio = velocityRatio;
+ this->velocityRatio = velocityRatio;
}
void Parameter::setDensityRatio(real densityRatio)
{
- this->densityRatio = densityRatio;
+ this->densityRatio = densityRatio;
}
void Parameter::setPressureRatio(real pressureRatio)
{
- this->pressRatio = pressureRatio;
+ this->pressRatio = pressureRatio;
}
void Parameter::setRe(real Re)
{
- this->Re = Re;
+ this->Re = Re;
}
void Parameter::setSc(real Sc)
{
@@ -204,14 +204,14 @@ void Parameter::setG_r(real g_r)
}
void Parameter::setMemsizeGPU(double addMemory, bool reset)
{
- if (reset == true)
- {
- this->memsizeGPU = 0.;
- }
- else
- {
- this->memsizeGPU += addMemory;
- }
+ if (reset == true)
+ {
+ this->memsizeGPU = 0.;
+ }
+ else
+ {
+ this->memsizeGPU += addMemory;
+ }
}
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
@@ -222,23 +222,23 @@ void Parameter::setMemsizeGPU(double addMemory, bool reset)
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
uint Parameter::getlimitOfNodesForVTK()
{
- return this->limitOfNodesForVTK;
+ return this->limitOfNodesForVTK;
}
ParameterStruct* Parameter::getParD()
{
- return this->parametersOnDevice;
+ return this->parametersOnDevice;
}
ParameterStruct* Parameter::getParH()
{
- return this->parametersOnHost;
+ return this->parametersOnHost;
}
bool Parameter::getIsTimestepEven()
{
- return this->parametersOnHost->isEvenTimestep;
+ return this->parametersOnHost->isEvenTimestep;
}
bool Parameter::getIsADcalculationOn()
{
- return this->calcAD;
+ return this->calcAD;
}
bool Parameter::getIsBodyForceOn()
{
@@ -246,67 +246,67 @@ bool Parameter::getIsBodyForceOn()
}
int Parameter::getD3Qxx()
{
- return this->D3Qxx;
+ return this->D3Qxx;
}
uint Parameter::getTimestepStart()
{
- return 1;
+ return 1;
}
uint Parameter::getTimestepEnd()
{
- return this->timestepEnd;
+ return this->timestepEnd;
}
uint Parameter::getTimestepOut()
{
- return this->timestepOut;
+ return this->timestepOut;
}
uint Parameter::getTimestepStartOut()
{
- return this->timestepStartOut;
+ return this->timestepStartOut;
}
std::string Parameter::getOutputPath()
{
- return this->outputPath;
+ return this->outputPath;
}
std::string Parameter::getOutputPrefix()
{
- return this->outputPrefix;
+ return this->outputPrefix;
}
std::string Parameter::getPathAndFilename()
{
- return this->pathAndFilename;
+ return this->pathAndFilename;
}
bool Parameter::getPrintFiles()
{
- return this->printFiles;
+ return this->printFiles;
}
real Parameter::getViscosityLB()
{
- return this->viscosityLB;
+ return this->viscosityLB;
}
real Parameter::getVelocityLB()
{
- return this->velocityLB;
+ return this->velocityLB;
}
real Parameter::getViscosityRatio()
{
- return this->viscosityRatio;
+ return this->viscosityRatio;
}
real Parameter::getVelocityRatio()
{
- return this->velocityRatio;
+ return this->velocityRatio;
}
real Parameter::getDensityRatio()
{
- return this->densityRatio;
+ return this->densityRatio;
}
real Parameter::getPressureRatio()
{
- return this->pressRatio;
+ return this->pressRatio;
}
real Parameter::getRe()
{
- return this->Re;
+ return this->Re;
}
real Parameter::getSc()
{
@@ -322,7 +322,7 @@ real Parameter::getG_r()
}
double Parameter::getMemsizeGPU()
{
- return this->memsizeGPU;
+ return this->memsizeGPU;
}
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
@@ -332,23 +332,23 @@ double Parameter::getMemsizeGPU()
// initial condition fluid
void Parameter::setInitialCondition(std::function<void(real, real, real, real&, real&, real&, real&)> initialCondition)
{
- this->initialCondition = initialCondition;
+ this->initialCondition = initialCondition;
}
std::function<void(real, real, real, real&, real&, real&, real&)>& Parameter::getInitialCondition()
{
- return this->initialCondition;
+ return this->initialCondition;
}
// initial condition concentration
void Parameter::setInitialConditionAD(std::function<void(real, real, real, real&)> initialConditionAD)
{
- this->initialConditionAD = initialConditionAD;
+ this->initialConditionAD = initialConditionAD;
}
std::function<void(real, real, real, real&)>& Parameter::getInitialConditionAD()
{
- return this->initialConditionAD;
+ return this->initialConditionAD;
}
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
index 24b472f5ac3230e2925389d66b7a840f813f5e31..3e212482de466c4c4451d9d02d2ce6433a02edce 100644
--- a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
+++ b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
@@ -42,50 +42,55 @@
//! \brief struct holds and manages the LB-parameter of the simulation that correspond to a certain level
//! \brief For this purpose it holds structures and pointer for host and device data, respectively.
struct ParameterStruct{
- //////////////////////////////////////////////////////////////////////////
- //! \brief decides if the simulation timestep is even or odd
- //! \brief this information is important for the esoteric twist
- bool isEvenTimestep;
- //////////////////////////////////////////////////////////////////////////
- //! \brief stores the number of threads per GPU block
- uint numberofthreads;
- //////////////////////////////////////////////////////////////////////////
- //! \brief store all distribution functions for the D3Q27
- Distributions27 distributions;
- //////////////////////////////////////////////////////////////////////////
- //! \brief stores the type for every lattice node (f.e. fluid node)
- uint *typeOfGridNode;
- //////////////////////////////////////////////////////////////////////////
- //! \brief store the neighbors in +X, +Y, +Z, and in diagonal negative direction
- //! \brief this information is important because we use an indirect addressing scheme
- uint *neighborX, *neighborY, *neighborZ, *neighborInverse;
- //////////////////////////////////////////////////////////////////////////
- //! \brief store the coordinates for every lattice node
- real *coordinateX, *coordinateY, *coordinateZ;
- //////////////////////////////////////////////////////////////////////////
- //! \brief store the macroscopic values (velocity, density, pressure)
- //! \brief for every lattice node
- real *velocityX, *velocityY, *velocityZ, *rho, *pressure;
- //! \brief stores the value for omega
- real omega;
- //////////////////////////////////////////////////////////////////////////
- //! \brief stores the number of nodes (based on indirect addressing scheme)
- uint numberOfNodes;
- //! \brief stores the size of the memory consumption for real/int values of the above arrays
- uint mem_size_real, mem_size_int;
+ //////////////////////////////////////////////////////////////////////////
+ //! \brief decides if the simulation timestep is even or odd
+ //! \brief this information is important for the esoteric twist
+ bool isEvenTimestep;
+ //////////////////////////////////////////////////////////////////////////
+ //! \brief stores the number of threads per GPU block
+ uint numberofthreads;
+ //////////////////////////////////////////////////////////////////////////
+ //! \brief store all distribution functions for the D3Q27
+ Distributions27 distributions;
+ //////////////////////////////////////////////////////////////////////////
+ //! \brief stores the type for every lattice node (f.e. fluid node)
+ uint *typeOfGridNode;
+ //////////////////////////////////////////////////////////////////////////
+ //! \brief store the neighbors in +X, +Y, +Z, and in diagonal negative direction
+ //! \brief this information is important because we use an indirect addressing scheme
+ uint *neighborX, *neighborY, *neighborZ, *neighborInverse;
+ //////////////////////////////////////////////////////////////////////////
+ //! \brief store the coordinates for every lattice node
+ real *coordinateX, *coordinateY, *coordinateZ;
+ //////////////////////////////////////////////////////////////////////////
+ //! \brief store the macroscopic values (velocity, density, pressure)
+ //! \brief for every lattice node
+ real *velocityX, *velocityY, *velocityZ, *rho, *pressure;
+ //! \brief stores the value for omega
+ real omega;
+ //////////////////////////////////////////////////////////////////////////
+ //! \brief stores the number of nodes (based on indirect addressing scheme)
+ uint numberOfNodes;
+ //! \brief stores the size of the memory consumption for real/int values of the above arrays
+ uint mem_size_real, mem_size_int;
//////////////////////////////////////////////////////////////////////////
//! \brief stores the slip boundary condition data
QforBoundaryConditions slipBC;
//! \brief number of lattice nodes for the slip boundary condition
uint numberOfSlipBCnodes;
//////////////////////////////////////////////////////////////////////////
- //! \brief stores the velocity boundary condition data
- QforBoundaryConditions inflowBC;
- //! \brief number of lattice nodes for the velocity boundary condition
- uint numberOfInflowBCnodes;
- //////////////////////////////////////////////////////////////////////////
- //! \brief sets the forcing uniform on every fluid node in all three space dimensions
- real *forcing;
+ //! \brief stores the velocity boundary condition data
+ QforBoundaryConditions veloBC;
+ //! \brief number of lattice nodes for the velocity boundary condition
+ uint numberOfVeloBCnodes;
+ //////////////////////////////////////////////////////////////////////////
+ //! \brief stores the geometry boundary condition data
+ QforBoundaryConditions geometryBC;
+ //! \brief number of lattice nodes for the velocity boundary condition
+ uint numberOfgeometryBCnodes;
+ //////////////////////////////////////////////////////////////////////////
+ //! \brief sets the forcing uniform on every fluid node in all three space dimensions
+ real *forcing;
//////////////////////////////////////////////////////////////////////////
// Advection Diffusion
@@ -105,47 +110,47 @@ struct ParameterStruct{
class VIRTUALFLUIDS_GPU_EXPORT Parameter
{
public:
- ////////////////////////////////////////////////////////////////////////////
- //! \brief generate a new instance of parameter object
- static SPtr<Parameter> make();
- //! \brief returns the instance generate a new instance of parameter object
- static Parameter* getInstanz();
- //! \brief Pointer to instance of ParameterStruct - stored on Host System
- ParameterStruct* getParH();
- //! \brief Pointer to instance of ParameterStruct - stored on Device (GPU)
- ParameterStruct* getParD();
+ ////////////////////////////////////////////////////////////////////////////
+ //! \brief generate a new instance of parameter object
+ static SPtr<Parameter> make();
+ //! \brief returns the instance generate a new instance of parameter object
+ static Parameter* getInstanz();
+ //! \brief Pointer to instance of ParameterStruct - stored on Host System
+ ParameterStruct* getParH();
+ //! \brief Pointer to instance of ParameterStruct - stored on Device (GPU)
+ ParameterStruct* getParD();
- //////////////////////////////////////////////////////////////////////////
- //! \brief initialization of necessary parameters at startup
- void initParameter();
+ //////////////////////////////////////////////////////////////////////////
+ //! \brief initialization of necessary parameters at startup
+ void initParameter();
- //////////////////////////////////////////////////////////////////////////
- //set methods
- //////////////////////////////////////////////////////////////////////////
- //! \brief sets the limit of nodes, that can be written to a binary unstructured grid VTK file
- //! \param limitOfNodesForVTK holds the maximum number of nodes
- void setlimitOfNodesForVTK(uint limitOfNodesForVTK);
- //! \brief sets the LBM stencil
- //! \param d3qxx holds the number of neighbors (f.e. 27)
- void setD3Qxx(int d3qxx);
- //! \brief sets timestep to stop the simulation
- //! \param timestepEnd holds the last timestep of the simulation
- void setTimestepEnd(uint timestepEnd);
- //! \brief sets time interval to write output files
- //! \param timestepOut holds the value for the output interval
- void setTimestepOut(uint timestepOut);
- //! \brief sets first timestep to write output files
- //! \param timestepStartOut holds the value for the first output timestep
- void setTimestepStartOut(uint timestepStartOut);
- //! \brief sets the path, where the vtk-files are stored
- //! \param string "oPath" represents the output path
- void setOutputPath(std::string outputPath);
- //! \brief sets the prefix of the vtk-files name
- //! \param string "oPrefix" represents the file-name-prefix
- void setOutputPrefix(std::string outputPrefix);
- //! \brief sets the complete string for the vtk-files with results
- //! \param string "fname" represents the combination of path and prefix
- void setPathAndFilename(std::string pathAndFilename);
+ //////////////////////////////////////////////////////////////////////////
+ //set methods
+ //////////////////////////////////////////////////////////////////////////
+ //! \brief sets the limit of nodes, that can be written to a binary unstructured grid VTK file
+ //! \param limitOfNodesForVTK holds the maximum number of nodes
+ void setlimitOfNodesForVTK(uint limitOfNodesForVTK);
+ //! \brief sets the LBM stencil
+ //! \param d3qxx holds the number of neighbors (f.e. 27)
+ void setD3Qxx(int d3qxx);
+ //! \brief sets timestep to stop the simulation
+ //! \param timestepEnd holds the last timestep of the simulation
+ void setTimestepEnd(uint timestepEnd);
+ //! \brief sets time interval to write output files
+ //! \param timestepOut holds the value for the output interval
+ void setTimestepOut(uint timestepOut);
+ //! \brief sets first timestep to write output files
+ //! \param timestepStartOut holds the value for the first output timestep
+ void setTimestepStartOut(uint timestepStartOut);
+ //! \brief sets the path, where the vtk-files are stored
+ //! \param string "oPath" represents the output path
+ void setOutputPath(std::string outputPath);
+ //! \brief sets the prefix of the vtk-files name
+ //! \param string "oPrefix" represents the file-name-prefix
+ void setOutputPrefix(std::string outputPrefix);
+ //! \brief sets the complete string for the vtk-files with results
+ //! \param string "fname" represents the combination of path and prefix
+ void setPathAndFilename(std::string pathAndFilename);
//! \brief sets the status, if Advection Diffusion should be calculated
//! \param if calcAD is true, the Advection Diffusion will be calculated
void setIsADcalculationOn(bool calcAD);
@@ -153,29 +158,29 @@ public:
//! \param if calcBF is true, the body force will be calculated
void setIsBodyForceOn(bool calcBF);
//! \brief sets the status, if the vtk files should be printed
- //! \param if printfiles is true, the vtk files will be printed
- void setPrintFiles(bool printfiles);
- //! \brief sets the viscosity in LB units
- //! \param viscosity in LB units
- void setViscosityLB(real viscosity);
- //! \brief sets the velocity in LB units
- //! \param velocity in LB units
- void setVelocityLB(real velocity);
- //! \brief sets the viscosity ratio between SI and LB units
- //! \param viscosityRatio SI/LB units
- void setViscosityRatio(real viscosityRatio);
- //! \brief sets the velocity ratio between SI and LB units
- //! \param velocityRatio SI/LB units
- void setVelocityRatio(real velocityRatio);
- //! \brief sets the density ratio between SI and LB units
- //! \param densityRatio SI/LB units
- void setDensityRatio(real densityRatio);
- //! \brief sets the pressure ratio between SI and LB units
- //! \param pressureRatio SI/LB units
- void setPressureRatio(real pressureRatio);
- //! \brief sets the Reynolds number (Re) for the simulation
- //! \param Reynolds number (Re)
- void setRe(real Re);
+ //! \param if printfiles is true, the vtk files will be printed
+ void setPrintFiles(bool printfiles);
+ //! \brief sets the viscosity in LB units
+ //! \param viscosity in LB units
+ void setViscosityLB(real viscosity);
+ //! \brief sets the velocity in LB units
+ //! \param velocity in LB units
+ void setVelocityLB(real velocity);
+ //! \brief sets the viscosity ratio between SI and LB units
+ //! \param viscosityRatio SI/LB units
+ void setViscosityRatio(real viscosityRatio);
+ //! \brief sets the velocity ratio between SI and LB units
+ //! \param velocityRatio SI/LB units
+ void setVelocityRatio(real velocityRatio);
+ //! \brief sets the density ratio between SI and LB units
+ //! \param densityRatio SI/LB units
+ void setDensityRatio(real densityRatio);
+ //! \brief sets the pressure ratio between SI and LB units
+ //! \param pressureRatio SI/LB units
+ void setPressureRatio(real pressureRatio);
+ //! \brief sets the Reynolds number (Re) for the simulation
+ //! \param Reynolds number (Re)
+ void setRe(real Re);
//! \brief sets the Schmidt number (Sc) for the simulation
//! \param sets the Schmidt number (Sc)
void setSc(real Sc);
@@ -186,53 +191,53 @@ public:
//! \param reduced acceleration due to gravity in buoyancy driven flow (LBM units)
void setG_r(real g_r);
//! \brief sets the necessary memory on the device(s)/GPU(s)
- //! \param addMemory is the amount of additional memory
- //! \param reset decides if the value for overall GPU memory should be set to zero
- void setMemsizeGPU(double addMemory, bool reset);
+ //! \param addMemory is the amount of additional memory
+ //! \param reset decides if the value for overall GPU memory should be set to zero
+ void setMemsizeGPU(double addMemory, bool reset);
- //////////////////////////////////////////////////////////////////////////
- //get methods
- //////////////////////////////////////////////////////////////////////////
- //! \brief return the limit of nodes, that can be written to a binary unstructured grid VTK file
- uint getlimitOfNodesForVTK();
- //! \brief return if the timestep is even or odd
- bool getIsTimestepEven();
+ //////////////////////////////////////////////////////////////////////////
+ //get methods
+ //////////////////////////////////////////////////////////////////////////
+ //! \brief return the limit of nodes, that can be written to a binary unstructured grid VTK file
+ uint getlimitOfNodesForVTK();
+ //! \brief return if the timestep is even or odd
+ bool getIsTimestepEven();
//! \brief return true if Advection Diffusion calculation is switched on
bool getIsADcalculationOn();
//! \brief return true if body force calculation is switched on
bool getIsBodyForceOn();
//! \brief return if the simulation should write VTK files
- bool getPrintFiles();
- //! \brief return the number of neighbors of a lattice node (stencil)
- int getD3Qxx();
- //! \brief return the output path
- std::string getOutputPath();
- //! \brief return the prefix of the output files
- std::string getOutputPrefix();
- //! \brief return the combination of output path and prefix
- std::string getPathAndFilename();
- //! \brief return the timestep to start the simulation (in this code version = 1)
- uint getTimestepStart();
- //! \brief return the timestep to end the simulation
- uint getTimestepEnd();
- //! \brief return the time interval to write output files
- uint getTimestepOut();
- //! \brief return the timestep to start writing output files
- uint getTimestepStartOut();
- //! \brief return the viscosity in LB units
- real getViscosityLB();
- //! \brief return the velocity in LB units
- real getVelocityLB();
- //! \brief return the viscosity ratio in SI/LB units
- real getViscosityRatio();
- //! \brief return the velocity ratio in SI/LB units
- real getVelocityRatio();
- //! \brief return the density ratio in SI/LB units
- real getDensityRatio();
- //! \brief return the pressure ratio in SI/LB units
- real getPressureRatio();
- //! \brief return the Reynolds number
- real getRe();
+ bool getPrintFiles();
+ //! \brief return the number of neighbors of a lattice node (stencil)
+ int getD3Qxx();
+ //! \brief return the output path
+ std::string getOutputPath();
+ //! \brief return the prefix of the output files
+ std::string getOutputPrefix();
+ //! \brief return the combination of output path and prefix
+ std::string getPathAndFilename();
+ //! \brief return the timestep to start the simulation (in this code version = 1)
+ uint getTimestepStart();
+ //! \brief return the timestep to end the simulation
+ uint getTimestepEnd();
+ //! \brief return the time interval to write output files
+ uint getTimestepOut();
+ //! \brief return the timestep to start writing output files
+ uint getTimestepStartOut();
+ //! \brief return the viscosity in LB units
+ real getViscosityLB();
+ //! \brief return the velocity in LB units
+ real getVelocityLB();
+ //! \brief return the viscosity ratio in SI/LB units
+ real getViscosityRatio();
+ //! \brief return the velocity ratio in SI/LB units
+ real getVelocityRatio();
+ //! \brief return the density ratio in SI/LB units
+ real getDensityRatio();
+ //! \brief return the pressure ratio in SI/LB units
+ real getPressureRatio();
+ //! \brief return the Reynolds number
+ real getRe();
//! \brief return the Schmidt number
real getSc();
//! \brief return number of nodes used to discretize characteristic length scale L
@@ -240,9 +245,9 @@ public:
//! \brief return reduced acceleration due to gravity in buoyancy driven flow (LBM units)
real getG_r();
//! \brief return the used device memory
- double getMemsizeGPU();
+ double getMemsizeGPU();
- ////////////////////////////////////////////////////////////////////////////
+ ////////////////////////////////////////////////////////////////////////////
// initial condition fluid
void setInitialCondition(std::function<void(real, real, real, real &, real &, real &, real &)> initialCondition);
std::function<void(real, real, real, real &, real &, real &, real &)> &getInitialCondition();
@@ -251,24 +256,24 @@ public:
std::function<void(real, real, real, real &)> &getInitialConditionAD();
//////////////////////////////////////////////////////////////////////////
- //! Class destructor
- ~Parameter();
+ //! Class destructor
+ ~Parameter();
protected:
private:
- //! \brief instance of parameter object
- static Parameter* instance;
- //! \brief stencil for the LB simulation, number of node neighbors
- int D3Qxx;
- //! \brief limit of nodes, that can be written to a binary unstructured grid VTK file
- uint limitOfNodesForVTK;
- //! \brief last timestep of the simulation
- uint timestepEnd;
- //! \brief time interval to write output files
- uint timestepOut;
- //! \brief timestep - start writing output files
- uint timestepStartOut;
- //! \brief Reynolds number
- real Re;
+ //! \brief instance of parameter object
+ static Parameter* instance;
+ //! \brief stencil for the LB simulation, number of node neighbors
+ int D3Qxx;
+ //! \brief limit of nodes, that can be written to a binary unstructured grid VTK file
+ uint limitOfNodesForVTK;
+ //! \brief last timestep of the simulation
+ uint timestepEnd;
+ //! \brief time interval to write output files
+ uint timestepOut;
+ //! \brief timestep - start writing output files
+ uint timestepStartOut;
+ //! \brief Reynolds number
+ real Re;
//! \brief Schmidt number
real Sc;
//! \brief number of nodes used to discretize characteristic length scale L
@@ -276,35 +281,35 @@ private:
//! \brief reduced acceleration due to gravity in buoyancy driven flow (LBM units)
real g_r;
//! \brief viscosity and velocity in LB units
- real viscosityLB, velocityLB;
- //! \brief ratio SI units / LB units for viscosity, velocity, density and pressure
- real viscosityRatio, velocityRatio;
- real densityRatio, pressRatio;
- //! \brief used device memory
- double memsizeGPU;
+ real viscosityLB, velocityLB;
+ //! \brief ratio SI units / LB units for viscosity, velocity, density and pressure
+ real viscosityRatio, velocityRatio;
+ real densityRatio, pressRatio;
+ //! \brief used device memory
+ double memsizeGPU;
//! \brief Advection Diffusion calculation on/off
bool calcAD;
//! \brief Include body forces on/off
bool calcBF;
//! \brief write output files on/off
- bool printFiles;
- //! \brief strings to store output path, prefix and combination of both
- std::string pathAndFilename, outputPath, outputPrefix;
+ bool printFiles;
+ //! \brief strings to store output path, prefix and combination of both
+ std::string pathAndFilename, outputPath, outputPrefix;
- ////////////////////////////////////////////////////////////////////////////
+ ////////////////////////////////////////////////////////////////////////////
//! \brief initial condition fluid
std::function<void(real, real, real, real &, real &, real &, real &)> initialCondition;
//! \brief initial condition concentration
std::function<void(real, real, real, real &)> initialConditionAD;
//! \brief pointer to LB-parameter struct on host system
- ParameterStruct* parametersOnHost;
- //! \brief pointer to LB-parameter struct on device/GPU
- ParameterStruct* parametersOnDevice;
+ ParameterStruct* parametersOnHost;
+ //! \brief pointer to LB-parameter struct on device/GPU
+ ParameterStruct* parametersOnDevice;
- //! Class default constructor
- Parameter();
- //Parameter(const Parameter&);
+ //! Class default constructor
+ Parameter();
+ //Parameter(const Parameter&);
};
#endif