diff --git a/src/gpu/GridGenerator/grid/MultipleGridBuilderFacade.h b/src/gpu/GridGenerator/grid/MultipleGridBuilderFacade.h index a22714e653a38252052e4b7a32ced9574c18f5cb..263f0ac15646b468da5cae3ef0dee5397918d435 100644 --- a/src/gpu/GridGenerator/grid/MultipleGridBuilderFacade.h +++ b/src/gpu/GridGenerator/grid/MultipleGridBuilderFacade.h @@ -30,8 +30,6 @@ //! \ingroup gpu_GridGenerator GridGenerator //! \{ //! \author Anna Wellmann -//! \brief A class that makes the setup of simulations on multiple gpus easier -//! \details Using this class is optional. //======================================================================================= @@ -58,8 +56,9 @@ class FileCollection; //! \class MultipleGridBuilderFacade //! \brief Simplifies the creation of the grids for a multi-gpu simulation +//! \details Using this class is optional. //! -//! \details Steps to set up the grids: +//! Steps to set up the grids: //! //! - 1. initialize class with a MultipleGridBuilder and the dimensions of the entire domain //! diff --git a/src/gpu/core/DataStructureInitializer/GridReaderGenerator/IndexRearrangementForStreams.cpp b/src/gpu/core/DataStructureInitializer/GridReaderGenerator/IndexRearrangementForStreams.cpp index a68f5633e490cfcd04f86538131182769770c3f9..a664803d6026008d93932cf4f848cb56638f4413 100644 --- a/src/gpu/core/DataStructureInitializer/GridReaderGenerator/IndexRearrangementForStreams.cpp +++ b/src/gpu/core/DataStructureInitializer/GridReaderGenerator/IndexRearrangementForStreams.cpp @@ -30,7 +30,6 @@ //! \ingroup gpu_core core //! \{ //! \author Anna Wellmann -//! \details See [master thesis of Anna Wellmann] //======================================================================================= #include "IndexRearrangementForStreams.h" diff --git a/src/gpu/core/DataStructureInitializer/GridReaderGenerator/IndexRearrangementForStreams.h b/src/gpu/core/DataStructureInitializer/GridReaderGenerator/IndexRearrangementForStreams.h index 5a99159048eb9b51da691e9d3b1ab912ea86eca0..9a2465cd21cc7a5d483c986018ed5e57e7ce0a04 100644 --- a/src/gpu/core/DataStructureInitializer/GridReaderGenerator/IndexRearrangementForStreams.h +++ b/src/gpu/core/DataStructureInitializer/GridReaderGenerator/IndexRearrangementForStreams.h @@ -30,7 +30,6 @@ //! \ingroup gpu_core core //! \{ //! \author Anna Wellmann -//! \details See [master thesis of Anna Wellmann] //======================================================================================= #ifndef IndexRearrangementForStreams_H #define IndexRearrangementForStreams_H @@ -47,10 +46,14 @@ namespace vf::parallel class Communicator; } +//! \brief class that is used to rearrange the arrays of node indices for communication between gpus. The rearrangement is +//! needed for communication hiding with cuda streams +//! \details This class changes the order of the node indices that are needed for communication between gpus. The indices are +//! reordered so that they can be split into two groups: nodes that are part if the interpolation between grid levels, and +//! nodes that are not. These groups are needed for communication hiding. For details see [master thesis of Anna Wellmann] class IndexRearrangementForStreams { public: - //! \brief Construct IndexRearrangementForStreams object IndexRearrangementForStreams(std::shared_ptr<Parameter> para, std::shared_ptr<GridBuilder> builder, vf::parallel::Communicator& communicator); virtual ~IndexRearrangementForStreams() = default; diff --git a/src/gpu/core/DataStructureInitializer/GridReaderGenerator/IndexRearrangementForStreamsTest.cpp b/src/gpu/core/DataStructureInitializer/GridReaderGenerator/IndexRearrangementForStreamsTest.cpp index 8ce86fd9c6c99ca99d3f8e74026552ad7a3af264..015396dc9f24b1f881df441a11747e130a06080d 100644 --- a/src/gpu/core/DataStructureInitializer/GridReaderGenerator/IndexRearrangementForStreamsTest.cpp +++ b/src/gpu/core/DataStructureInitializer/GridReaderGenerator/IndexRearrangementForStreamsTest.cpp @@ -30,7 +30,6 @@ //! \ingroup gpu_core_tests core //! \{ //! \author Anna Wellmann -//! \details See [master thesis of Anna Wellmann] //======================================================================================= #include <gmock/gmock.h> diff --git a/src/gpu/core/DataStructureInitializer/GridReaderGenerator/InterpolationCellGrouper.cpp b/src/gpu/core/DataStructureInitializer/GridReaderGenerator/InterpolationCellGrouper.cpp index a6088cbf607ffbdc2210c0d7f71f49e616b567a0..ec28550cff0ba81fceafd1dd720778d7f97f685f 100644 --- a/src/gpu/core/DataStructureInitializer/GridReaderGenerator/InterpolationCellGrouper.cpp +++ b/src/gpu/core/DataStructureInitializer/GridReaderGenerator/InterpolationCellGrouper.cpp @@ -30,7 +30,6 @@ //! \ingroup gpu_core core //! \{ //! \author Anna Wellmann -//! \details See [master thesis of Anna Wellmann] //======================================================================================= #include "InterpolationCellGrouper.h" diff --git a/src/gpu/core/DataStructureInitializer/GridReaderGenerator/InterpolationCellGrouper.h b/src/gpu/core/DataStructureInitializer/GridReaderGenerator/InterpolationCellGrouper.h index 82d74b6d4f49db16f80a9dbf0dfdf4f493bf0f84..55ca4f5e7f07257143d12be9bcaf0dba65f7f1ac 100644 --- a/src/gpu/core/DataStructureInitializer/GridReaderGenerator/InterpolationCellGrouper.h +++ b/src/gpu/core/DataStructureInitializer/GridReaderGenerator/InterpolationCellGrouper.h @@ -30,7 +30,6 @@ //! \ingroup gpu_core core //! \{ //! \author Anna Wellmann -//! \details See [master thesis of Anna Wellmann] //======================================================================================= #ifndef InterpolationCellGrouper_H @@ -46,9 +45,11 @@ class GridBuilder; using LBMSimulationParameters = std::vector<std::shared_ptr<LBMSimulationParameter>>; +//! \brief Split the interpolation cells into two groups: cells which are at the border between gpus and therefore involved +//! in the communication between gpus, and cells which are not directly related to the communication between gpus. +//! \details See [master thesis of Anna Wellmann] class InterpolationCellGrouper { public: - //! \brief Construct InterpolationCellGrouper object InterpolationCellGrouper(const LBMSimulationParameters &parHs, const LBMSimulationParameters &parDs, SPtr<GridBuilder> builder); diff --git a/src/gpu/core/DataStructureInitializer/GridReaderGenerator/InterpolationCellGrouperTest.cpp b/src/gpu/core/DataStructureInitializer/GridReaderGenerator/InterpolationCellGrouperTest.cpp index f9c2a8b53eab1aa09ca90fd737750e60e7dc7478..6dff1ab7a11fb341707fc1b2dc3b9bb606f39b41 100644 --- a/src/gpu/core/DataStructureInitializer/GridReaderGenerator/InterpolationCellGrouperTest.cpp +++ b/src/gpu/core/DataStructureInitializer/GridReaderGenerator/InterpolationCellGrouperTest.cpp @@ -30,7 +30,6 @@ //! \ingroup gpu_core_tests core //! \{ //! \author Anna Wellmann -//! \details See [master thesis of Anna Wellmann] //======================================================================================= #include <gmock/gmock.h> #include "Utilities/testUtilitiesGPU.h" diff --git a/src/gpu/core/Output/DistributionDebugInspector.h b/src/gpu/core/Output/DistributionDebugInspector.h index a83ddeb789931ad79923531e13e77d39bba4d407..1d3d83cb356492fa11da9cf13593b3a9e03aa9a8 100644 --- a/src/gpu/core/Output/DistributionDebugInspector.h +++ b/src/gpu/core/Output/DistributionDebugInspector.h @@ -31,12 +31,6 @@ //! \{ //! \author Henrik Asmuth //! \date 13/012/2022 -//! \brief Basic debugging class to print out f's in a certain area of the domain -//! -//! Basic debugging class. Needs to be directly added in UpdateGrid (could potentially also be added as a proper Probe in the -//! future) How to use: Define a part of the domain via min/max x, y, and z. The DistributionDebugInspector will print out -//! all f's in that area. -//! //======================================================================================= #ifndef DISTRIBUTION_INSPECTOR_H @@ -46,6 +40,11 @@ class Parameter; +//! \brief Basic debugging class to print out f's in a certain area of the domain +//! +//! Basic debugging class. Needs to be directly added in UpdateGrid (could potentially also be added as a proper Probe in the +//! future) How to use: Define a part of the domain via min/max x, y, and z. The DistributionDebugInspector will print out +//! all f's in that area. class DistributionDebugInspector { public: diff --git a/src/gpu/core/Parameter/EdgeNodeFinder.cpp b/src/gpu/core/Parameter/EdgeNodeFinder.cpp index 428b24e635e7bf8f1f23f31da3ff07901381b02a..93f6d13c80597eacdb1400d0b03f5b8e5133829b 100644 --- a/src/gpu/core/Parameter/EdgeNodeFinder.cpp +++ b/src/gpu/core/Parameter/EdgeNodeFinder.cpp @@ -29,10 +29,6 @@ //! \addtogroup gpu_Parameter Parameter //! \ingroup gpu_core core //! \{ -//! \author Anna Wellmann -//! \brief Functions for finding edge nodes in the multi-gpu implementation -//! \details Edge nodes are nodes, which are part of the communication in multiple directions -//! \ref master thesis of Anna Wellmann (p. 54-57) //======================================================================================= #include <vector> #include <optional> diff --git a/src/gpu/core/Parameter/EdgeNodeFinder.h b/src/gpu/core/Parameter/EdgeNodeFinder.h index 72d7d23983340d214a507f50d46b597bb5aae1de..8eb6fbc841a51ddb00a755168b681f05fc8c51cc 100644 --- a/src/gpu/core/Parameter/EdgeNodeFinder.h +++ b/src/gpu/core/Parameter/EdgeNodeFinder.h @@ -30,9 +30,6 @@ //! \ingroup gpu_core core //! \{ //! \author Anna Wellmann -//! \brief Functions for finding edge nodes in the multi-gpu implementation -//! \details Edge nodes are nodes, which are part of the communication in multiple directions -//! \ref master thesis of Anna Wellmann (p. 54-57) //======================================================================================= #ifndef GPU_EDGENODES_H #define GPU_EDGENODES_H @@ -43,7 +40,9 @@ class Parameter; namespace vf::gpu { -//! \brief Find nodes which are part of communication in multiple coordinate directions +//! \brief Function for finding edge nodes in the multi-gpu implementation +//! \details Edge nodes are nodes, which are part of the communication in multiple directions +//! \ref master thesis of Anna Wellmann (p. 54-57) void findEdgeNodesCommMultiGPU(Parameter& parameter); } // namespace vf::gpu diff --git a/src/gpu/core/Parameter/EdgeNodeFinderTest.cpp b/src/gpu/core/Parameter/EdgeNodeFinderTest.cpp index 57c14cd9dbe0353ef53fc208492d21e0fa758f5e..11c8d0b65e4dc64c79ca4dbfe9258e2b67622a9b 100644 --- a/src/gpu/core/Parameter/EdgeNodeFinderTest.cpp +++ b/src/gpu/core/Parameter/EdgeNodeFinderTest.cpp @@ -29,10 +29,6 @@ //! \addtogroup gpu_Parameter_tests Parameter //! \ingroup gpu_core_tests core //! \{ -//! \author Anna Wellmann -//! \brief Functions for finding edge nodes in the multi-gpu implementation -//! \details Edge nodes are nodes, which are part of the communication in multiple directions -//! \ref master thesis of Anna Wellmann (p. 54-57) //======================================================================================= #include <gmock/gmock.h> diff --git a/src/gpu/core/PostProcessor/DragLift.cpp b/src/gpu/core/PostProcessor/DragLift.cpp index 0caa50d837c7f5f9a8f3ba0331fc8cee1336eb24..e1470cc5398bf387825ac50108519e2998468cf8 100644 --- a/src/gpu/core/PostProcessor/DragLift.cpp +++ b/src/gpu/core/PostProcessor/DragLift.cpp @@ -30,8 +30,6 @@ //! \ingroup gpu_core core //! \{ //! \author Martin Schoenherr -//! \brief note, that the drag/lift calculations are build for being used with -//! \brief geometry boundary nodes and the related area has to be defined here //======================================================================================= #include "DragLift.h" @@ -49,6 +47,8 @@ using namespace std; +//! \brief Calculate drag and lift for a geometry +//! \details note, that the drag/lift calculations are build for being used with geometry boundary nodes and the related area has to be defined here void calcDragLift(Parameter* para, CudaMemoryManager* cudaMemoryManager, int lev) { ////////////////////////////////////////////////////////////////////////// diff --git a/src/gpu/core/PreCollisionInteractor/PrecursorWriter.h b/src/gpu/core/PreCollisionInteractor/PrecursorWriter.h index bb7dd2e95dad054d7717c64d62ef09def05b2787..a3844690841983bf2e9dbc68dbbf8a49a0bff730 100644 --- a/src/gpu/core/PreCollisionInteractor/PrecursorWriter.h +++ b/src/gpu/core/PreCollisionInteractor/PrecursorWriter.h @@ -31,10 +31,6 @@ //! \{ //! \author Henry Korb, Henrik Asmuth //! \date 05/12/2022 -//! \brief Probe writing planes of data to be used as inflow data in successor simulation using PrecursorBC -//! -//! The probe writes out yz-planes at a specific x position ( \param xPos ) of either velocity or distributions -//! that can be read by PrecursorBC as inflow data. //======================================================================================= @@ -87,6 +83,11 @@ struct PrecursorStruct cudaStream_t stream; }; +//! \brief Probe writing planes of data to be used as inflow data in successor simulation using PrecursorBC +//! +//! The probe writes out yz-planes at a specific x position ( \param xPos ) of either velocity or distributions +//! that can be read by PrecursorBC as inflow data. +//! class PrecursorWriter : public PreCollisionInteractor { public: diff --git a/src/gpu/core/PreCollisionInteractor/Probes/PlanarAverageProbe.h b/src/gpu/core/PreCollisionInteractor/Probes/PlanarAverageProbe.h index 800dd039343e4b484e16ed3e24fe02b3e90953ae..09bdce7de87675c6a97b8b98df4351282647bdb8 100644 --- a/src/gpu/core/PreCollisionInteractor/Probes/PlanarAverageProbe.h +++ b/src/gpu/core/PreCollisionInteractor/Probes/PlanarAverageProbe.h @@ -31,11 +31,6 @@ //! \{ //! \author Henrik Asmuth //! \date 13/05/2022 -//! \brief Probe computing statistics across planes spanning the entire domain -//! -//! Computes spatial statistics across x, y or z-normal planes defined by planeNormal. -//! The planes include all points of the domain at each respective position along that normal direction. -//! The spatial statistics can additionally be averaged in time. //! //======================================================================================= @@ -50,6 +45,11 @@ __global__ void moveIndicesInPosNormalDir( uint* pointIndices, uint nPoints, uin /////////////////////////////////////////////////////////////////////////////////// +//! \brief Probe computing statistics across planes spanning the entire domain +//! +//! Computes spatial statistics across x, y or z-normal planes defined by planeNormal. +//! The planes include all points of the domain at each respective position along that normal direction. +//! The spatial statistics can additionally be averaged in time. class PlanarAverageProbe : public Probe { public: diff --git a/src/gpu/core/PreCollisionInteractor/Probes/PlaneProbe.h b/src/gpu/core/PreCollisionInteractor/Probes/PlaneProbe.h index f329c03bef7ec92023f84bb52b509dc745780019..9594e4e34ab9d1ff3bb9689049e812b19768e5c9 100644 --- a/src/gpu/core/PreCollisionInteractor/Probes/PlaneProbe.h +++ b/src/gpu/core/PreCollisionInteractor/Probes/PlaneProbe.h @@ -31,11 +31,6 @@ //! \{ //! \author Henry Korb, Henrik Asmuth //! \date 13/05/2022 -//! \brief Probe computing point-wise statistics for a set of points across a plane -//! -//! The set of points can be defined by providing a list or on an x-normal plane. -//! All statistics are temporal. -//! //======================================================================================= #ifndef PlaneProbe_H @@ -45,6 +40,11 @@ #include "Probe.h" +//! \brief Probe computing point-wise statistics for a set of points across a plane +//! +//! The set of points can be defined by providing a list or on an x-normal plane. +//! All statistics are temporal. +//! class PlaneProbe : public Probe { public: diff --git a/src/gpu/core/PreCollisionInteractor/Probes/PointProbe.h b/src/gpu/core/PreCollisionInteractor/Probes/PointProbe.h index 19d2a2e11b6c1c07f21c3cfaed13bc1ac87ca98c..bb7cc02a4ce7ccb9551d70347bc916c6ce2321e6 100644 --- a/src/gpu/core/PreCollisionInteractor/Probes/PointProbe.h +++ b/src/gpu/core/PreCollisionInteractor/Probes/PointProbe.h @@ -31,11 +31,6 @@ //! \{ //! \author Henry Korb, Henrik Asmuth //! \date 13/05/2022 -//! \brief Probe computing statistics for a set of points in space -//! -//! The set of points can be defined by providing a list or on an x-normal plane (the latter being somewhat redundant with PlaneProbe) -//! All statistics are temporal. -//! //======================================================================================= #ifndef PointProbe_H @@ -43,6 +38,11 @@ #include "Probe.h" +//! \brief Probe computing statistics for a set of points in space +//! +//! The set of points can be defined by providing a list or on an x-normal plane (the latter being somewhat redundant with PlaneProbe) +//! All statistics are temporal. +//! class PointProbe: public Probe { public: diff --git a/src/gpu/core/PreCollisionInteractor/Probes/Probe.h b/src/gpu/core/PreCollisionInteractor/Probes/Probe.h index 64d85bc4425e1a950a7749f53645cff4a9fc6aea..40874cb22b1bc943864c2df5fd24495e724d84a9 100644 --- a/src/gpu/core/PreCollisionInteractor/Probes/Probe.h +++ b/src/gpu/core/PreCollisionInteractor/Probes/Probe.h @@ -31,15 +31,6 @@ //! \{ //! \author Henry Korb, Henrik Asmuth //! \date 13/05/2022 -//! \brief Base class for probes called in UpdateGrid27 -//! -//! Any probe should be initiated in the app and added via para->addProbe( someProbe ) -//! Note, that all probes generally require that macroscopic variables have been updated in the -//! time step they are called in. Most collision kernels (atm, all except K17CompressibleNavierStokes) -//! don't do this and would require an explicit call of calcMacroscopicQuantities. It does seem quite -//! inexpensive though to simply save vx, vy, etc., directly in the collider. -//! -//! \todo might have to adapt conversionFactors when using grid refinement //======================================================================================= #ifndef Probe_H @@ -137,6 +128,15 @@ __global__ void interpAndCalcQuantitiesKernel( uint* pointIndices, uint calcOldTimestep(uint currentTimestep, uint lastTimestepInOldSeries); +//! \brief Base class for probes called in UpdateGrid27 +//! +//! Any probe should be initiated in the app and added via para->addProbe( someProbe ) +//! Note, that all probes generally require that macroscopic variables have been updated in the +//! time step they are called in. Most collision kernels (atm, all except K17CompressibleNavierStokes) +//! don't do this and would require an explicit call of calcMacroscopicQuantities. It does seem quite +//! inexpensive though to simply save vx, vy, etc., directly in the collider. +//! +//! \todo might have to adapt conversionFactors when using grid refinement class Probe : public PreCollisionInteractor { public: diff --git a/src/gpu/core/PreCollisionInteractor/Probes/WallModelProbe.h b/src/gpu/core/PreCollisionInteractor/Probes/WallModelProbe.h index 376c77d8096432b0b8ea6763fec6eb56fc89b8ae..6bcccc643bcc293adc49c2a3a90fd1706b348977 100644 --- a/src/gpu/core/PreCollisionInteractor/Probes/WallModelProbe.h +++ b/src/gpu/core/PreCollisionInteractor/Probes/WallModelProbe.h @@ -31,11 +31,6 @@ //! \{ //! \author Henrik Asmuth //! \date 13/05/2022 -//! \brief Probe computing statistics of all relevant wall model quantities used in the StressBC kernels -//! -//! Computes spatial statistics for all grid points of the StressBC -//! The spatial statistics can additionally be averaged in time. -//! //======================================================================================= #ifndef WallModelProbe_H @@ -47,6 +42,11 @@ /////////////////////////////////////////////////////////////////////////////////// +//! \brief Probe computing statistics of all relevant wall model quantities used in the StressBC kernels +//! +//! Computes spatial statistics for all grid points of the StressBC +//! The spatial statistics can additionally be averaged in time. +//! class WallModelProbe : public Probe { public: diff --git a/src/lbm/collision/K17CompressibleNavierStokes.h b/src/lbm/collision/K17CompressibleNavierStokes.h index 1ba5aaa1dfa6e04538cb3568ff41ef877b99a47a..41212c2846c947a54bfe04fafe3c79d9395d5afc 100644 --- a/src/lbm/collision/K17CompressibleNavierStokes.h +++ b/src/lbm/collision/K17CompressibleNavierStokes.h @@ -29,19 +29,8 @@ //! \addtogroup collision //! \ingroup lbm //! \{ -//! \author Anna Wellmann, Martin Schönherr, Henry Korb, Henrik Asmuth +//! \author Martin Schönherr, Anna Wellmann, Henry Korb, Henrik Asmuth //! \date 05/12/2022 -//! \brief Kernel for CumulantK17 including different turbulence models and options for local body forces and writing -//! macroscopic variables -//! -//! CumulantK17 kernel using chimera transformations and quartic limiters as present in Geier et al. (2017). Additional -//! options are three different eddy-viscosity turbulence models (Smagorinsky, AMD, QR) that can be set via the template -//! parameter turbulenceModel (with default TurbulenceModel::None). The kernel is executed separately for each subset of -//! fluid node indices with a different tag CollisionTemplate. For each subset, only the locally required options are -//! switched on ( \param writeMacroscopicVariables and/or \param applyBodyForce) in order to minimize memory accesses. The -//! default refers to the plain cumlant kernel (CollisionTemplate::Default). Nodes are added to subsets (taggedFluidNodes) in -//! Simulation::init using a corresponding tag with different values of CollisionTemplate. These subsets are provided by the -//! utilized PostCollisionInteractiors depending on they specific requirements (e.g. writeMacroscopicVariables for probes). //======================================================================================= #include <basics/constants/NumericConstants.h> @@ -77,7 +66,20 @@ namespace vf::lbm { ////////////////////////////////////////////////////////////////////////// -//! Cumulant K17 Kernel is based on \ref +//! \brief Kernel for CumulantK17 including different turbulence models and options for local body forces and writing +//! macroscopic variables +//! +//! \details +//! CumulantK17 kernel using chimera transformations and quartic limiters as present in Geier et al. (2017). Additional +//! options are three different eddy-viscosity turbulence models (Smagorinsky, AMD, QR) that can be set via the template +//! parameter turbulenceModel (with default TurbulenceModel::None). The kernel is executed separately for each subset of +//! fluid node indices with a different tag CollisionTemplate. For each subset, only the locally required options are +//! switched on ( \param writeMacroscopicVariables and/or \param applyBodyForce) in order to minimize memory accesses. The +//! default refers to the plain cumlant kernel (CollisionTemplate::Default). Nodes are added to subsets (taggedFluidNodes) in +//! Simulation::init using a corresponding tag with different values of CollisionTemplate. These subsets are provided by the +//! utilized PostCollisionInteractiors depending on they specific requirements (e.g. writeMacroscopicVariables for probes). +//! +//! The CumulantK17 Kernel is based on \ref //! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), DOI:10.1016/j.jcp.2017.05.040 //! ]</b></a> and \ref <a href="https://doi.org/10.1016/j.jcp.2017.07.004"><b>[ M. Geier et al. (2017), //! DOI:10.1016/j.jcp.2017.07.004 ]</b></a>