From 598197c5480e5e4c8f1d2cda71efa94a596fd42f Mon Sep 17 00:00:00 2001
From: HenrikAsmuth <henrik.asmuth@geo.uu.se>
Date: Thu, 13 Oct 2022 10:40:32 +0200
Subject: [PATCH] Add allocs and copy functions for tagged fluid nodes
---
src/gpu/GridGenerator/grid/Grid.h | 7 +++
.../grid/GridBuilder/GridBuilder.h | 7 +++
.../grid/GridBuilder/LevelGridBuilder.cpp | 34 +++++++++-
.../grid/GridBuilder/LevelGridBuilder.h | 8 +++
src/gpu/GridGenerator/grid/GridImp.cpp | 33 ++++++++++
src/gpu/GridGenerator/grid/GridImp.h | 7 +++
.../Calculation/CollisisionStrategy.cpp | 12 ++--
.../DataStructureInitializer/GridProvider.cpp | 11 +---
.../DataStructureInitializer/GridProvider.h | 6 +-
.../GridReaderFiles/GridReader.cpp | 7 +--
.../GridReaderFiles/GridReader.h | 3 +-
.../GridReaderGenerator/GridGenerator.cpp | 62 ++++++++++++++-----
.../GridReaderGenerator/GridGenerator.h | 3 +-
.../GPU/CudaMemoryManager.cpp | 47 ++++----------
.../VirtualFluids_GPU/GPU/CudaMemoryManager.h | 9 +--
.../CumulantK17Almighty.cu | 5 +-
src/gpu/VirtualFluids_GPU/LBM/LB.h | 18 +++---
src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp | 12 ++--
.../VirtualFluids_GPU/Parameter/Parameter.h | 28 ++++++---
.../PreCollisionInteractor/ActuatorLine.cu | 2 +-
.../PreCollisionInteractor/ActuatorLine.h | 2 +-
.../PreCollisionInteractor.h | 2 +-
.../Probes/PlaneProbe.cu | 2 +-
.../Probes/PlaneProbe.h | 2 +-
.../Probes/PointProbe.cu | 2 +-
.../Probes/PointProbe.h | 2 +-
.../PreCollisionInteractor/Probes/Probe.cu | 2 +-
.../PreCollisionInteractor/Probes/Probe.h | 2 +-
28 files changed, 215 insertions(+), 122 deletions(-)
diff --git a/src/gpu/GridGenerator/grid/Grid.h b/src/gpu/GridGenerator/grid/Grid.h
index 2c4a8b63e..85b19bedd 100644
--- a/src/gpu/GridGenerator/grid/Grid.h
+++ b/src/gpu/GridGenerator/grid/Grid.h
@@ -186,6 +186,13 @@ public:
virtual void sortFluidNodeIndicesApplyBodyForce() = 0;
virtual void sortFluidNodeIndicesAllFeatures() = 0;
+ virtual uint getNumberOfFluidNodeIndicesMacroVars() const = 0;
+ virtual uint getNumberOfFluidNodeIndicesApplyBodyForce() const = 0;
+ virtual uint getNumberOfFluidNodeIndicesAllFeatures() const = 0;
+ virtual void getFluidNodeIndicesMacroVars(uint *fluidNodeIndicesMacroVars) const = 0;
+ virtual void getFluidNodeIndicesApplyBodyForce(uint *fluidNodeIndicesApplyBodyForce) const = 0;
+ virtual void getFluidNodeIndicesAllFeatures(uint *fluidNodeIndicesAllFeatures) const = 0;
+
};
#endif
diff --git a/src/gpu/GridGenerator/grid/GridBuilder/GridBuilder.h b/src/gpu/GridGenerator/grid/GridBuilder/GridBuilder.h
index c1aad5a33..78222687b 100644
--- a/src/gpu/GridGenerator/grid/GridBuilder/GridBuilder.h
+++ b/src/gpu/GridGenerator/grid/GridBuilder/GridBuilder.h
@@ -143,6 +143,13 @@ public:
virtual void sortFluidNodeIndicesMacroVars(uint level) = 0;
virtual void sortFluidNodeIndicesApplyBodyForce(uint level) = 0;
virtual void sortFluidNodeIndicesAllFeatures(uint level) = 0;
+ virtual uint getNumberOfFluidNodesMacroVars(unsigned int level) const = 0;
+ virtual void getFluidNodeIndicesMacroVars(uint *fluidNodeIndicesMacroVars, const int level) const = 0;
+ virtual uint getNumberOfFluidNodesApplyBodyForce(unsigned int level) const = 0;
+ virtual void getFluidNodeIndicesApplyBodyForce(uint *fluidNodeIndicesApplyBodyForce, const int level) const = 0;
+ virtual uint getNumberOfFluidNodesAllFeatures(unsigned int level) const = 0;
+ virtual void getFluidNodeIndicesAllFeatures(uint *fluidNodeIndicesAllFeatures, const int level) const = 0;
+
};
diff --git a/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.cpp b/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.cpp
index 9c5f439b5..5476de241 100644
--- a/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.cpp
+++ b/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.cpp
@@ -694,15 +694,45 @@ void LevelGridBuilder::addFluidNodeIndicesAllFeatures(std::vector<uint> fluidNod
void LevelGridBuilder::sortFluidNodeIndicesMacroVars(uint level)
{
-
+ grids[level]->sortFluidNodeIndicesMacroVars();
}
void LevelGridBuilder::sortFluidNodeIndicesApplyBodyForce(uint level)
{
-
+ grids[level]->sortFluidNodeIndicesApplyBodyForce();
}
void LevelGridBuilder::sortFluidNodeIndicesAllFeatures(uint level)
{
+ grids[level]->sortFluidNodeIndicesAllFeatures();
+}
+
+uint LevelGridBuilder::getNumberOfFluidNodesMacroVars(unsigned int level) const
+{
+ return grids[level]->getNumberOfFluidNodeIndicesMacroVars();
+}
+void LevelGridBuilder::getFluidNodeIndicesMacroVars(uint *fluidNodeIndicesMacroVars, const int level) const
+{
+ grids[level]->getFluidNodeIndicesMacroVars(fluidNodeIndicesMacroVars);
+}
+
+uint LevelGridBuilder::getNumberOfFluidNodesApplyBodyForce(unsigned int level) const
+{
+ return grids[level]->getNumberOfFluidNodeIndicesApplyBodyForce();
+}
+
+void LevelGridBuilder::getFluidNodeIndicesApplyBodyForce(uint *fluidNodeIndicesApplyBodyForce, const int level) const
+{
+ grids[level]->getFluidNodeIndicesApplyBodyForce(fluidNodeIndicesApplyBodyForce);
+}
+
+uint LevelGridBuilder::getNumberOfFluidNodesAllFeatures(unsigned int level) const
+{
+ return grids[level]->getNumberOfFluidNodeIndicesAllFeatures();
+}
+
+void LevelGridBuilder::getFluidNodeIndicesAllFeatures(uint *fluidNodeIndicesAllFeatures, const int level) const
+{
+ grids[level]->getFluidNodeIndicesAllFeatures(fluidNodeIndicesAllFeatures);
}
\ No newline at end of file
diff --git a/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.h b/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.h
index 214829524..d87b118b2 100644
--- a/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.h
+++ b/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.h
@@ -198,9 +198,17 @@ public:
void addFluidNodeIndicesMacroVars(std::vector<uint> fluidNodeIndicesMacroVars, uint level) override;
void addFluidNodeIndicesApplyBodyForce(std::vector<uint> fluidNodeIndicesApplyBodyForce, uint level) override;
void addFluidNodeIndicesAllFeatures(std::vector<uint> fluidNodeIndicesAllFeatures, uint level) override;
+
void sortFluidNodeIndicesMacroVars(uint level) override;
void sortFluidNodeIndicesApplyBodyForce(uint level) override;
void sortFluidNodeIndicesAllFeatures(uint level) override;
+
+ uint getNumberOfFluidNodesMacroVars(unsigned int level) const override;
+ void getFluidNodeIndicesMacroVars(uint *fluidNodeIndicesMacroVars, const int level) const override;
+ uint getNumberOfFluidNodesApplyBodyForce(unsigned int level) const override;
+ void getFluidNodeIndicesApplyBodyForce(uint *fluidNodeIndicesApplyBodyForce, const int level) const override;
+ uint getNumberOfFluidNodesAllFeatures(unsigned int level) const override;
+ void getFluidNodeIndicesAllFeatures(uint *fluidNodeIndicesAllFeatures, const int level) const override;
};
#endif
diff --git a/src/gpu/GridGenerator/grid/GridImp.cpp b/src/gpu/GridGenerator/grid/GridImp.cpp
index e886cac20..cf908d0cd 100644
--- a/src/gpu/GridGenerator/grid/GridImp.cpp
+++ b/src/gpu/GridGenerator/grid/GridImp.cpp
@@ -2179,6 +2179,39 @@ void GridImp::sortFluidNodeIndicesAllFeatures()
}
}
+uint GridImp::getNumberOfFluidNodeIndicesMacroVars() const {
+ return (uint)this->fluidNodeIndicesMacroVars.size();
+}
+
+uint GridImp::getNumberOfFluidNodeIndicesApplyBodyForce() const {
+ return (uint)this->fluidNodeIndicesApplyBodyForce.size();
+}
+
+uint GridImp::getNumberOfFluidNodeIndicesAllFeatures() const {
+ return (uint)this->fluidNodeIndicesAllFeatures.size();
+}
+
+void GridImp::getFluidNodeIndicesMacroVars(uint *fluidNodeIndicesMacroVars) const
+{
+ for (uint nodeNumber = 0; nodeNumber < (uint)this->fluidNodeIndicesMacroVars.size(); nodeNumber++)
+ fluidNodeIndicesMacroVars[nodeNumber] = this->fluidNodeIndicesMacroVars[nodeNumber];
+}
+void GridImp::getFluidNodeIndicesApplyBodyForce(uint *fluidNodeIndicesApplyBodyForce) const
+{
+ for (uint nodeNumber = 0; nodeNumber < (uint)this->fluidNodeIndicesApplyBodyForce.size(); nodeNumber++)
+ fluidNodeIndicesApplyBodyForce[nodeNumber] = this->fluidNodeIndicesApplyBodyForce[nodeNumber];
+}
+void GridImp::getFluidNodeIndicesAllFeatures(uint *fluidNodeIndicesAllFeatures) const
+{
+ for (uint nodeNumber = 0; nodeNumber < (uint)this->fluidNodeIndicesAllFeatures.size(); nodeNumber++)
+ fluidNodeIndicesAllFeatures[nodeNumber] = this->fluidNodeIndicesAllFeatures[nodeNumber];
+}
+
+
+
+
+
+
void GridImp::print() const
{
printf("min: (%2.4f, %2.4f, %2.4f), max: (%2.4f, %2.4f, %2.4f), size: %d, delta: %2.4f\n", startX, startY, startZ,
diff --git a/src/gpu/GridGenerator/grid/GridImp.h b/src/gpu/GridGenerator/grid/GridImp.h
index d368f68e4..7a25ac1c8 100644
--- a/src/gpu/GridGenerator/grid/GridImp.h
+++ b/src/gpu/GridGenerator/grid/GridImp.h
@@ -373,6 +373,13 @@ public:
void sortFluidNodeIndicesApplyBodyForce() override;
void sortFluidNodeIndicesAllFeatures() override;
+ uint getNumberOfFluidNodeIndicesMacroVars() const override;
+ uint getNumberOfFluidNodeIndicesApplyBodyForce() const override;
+ uint getNumberOfFluidNodeIndicesAllFeatures() const override;
+ void getFluidNodeIndicesMacroVars(uint *fluidNodeIndicesMacroVars) const override;
+ void getFluidNodeIndicesApplyBodyForce(uint *fluidNodeIndicesApplyBodyForce) const override;
+ void getFluidNodeIndicesAllFeatures(uint *fluidNodeIndicesAllFeatures) const override;
+
public:
struct CommunicationIndices {
std::vector<uint> sendIndices;
diff --git a/src/gpu/VirtualFluids_GPU/Calculation/CollisisionStrategy.cpp b/src/gpu/VirtualFluids_GPU/Calculation/CollisisionStrategy.cpp
index 776ff31c1..c80be76bf 100644
--- a/src/gpu/VirtualFluids_GPU/Calculation/CollisisionStrategy.cpp
+++ b/src/gpu/VirtualFluids_GPU/Calculation/CollisisionStrategy.cpp
@@ -40,8 +40,8 @@ void CollisionAndExchange_noStreams_indexKernel::operator()(UpdateGrid27 *update
//! 1. run collision
//!
updateGrid->collisionUsingIndices( level, t,
- para->getParD(level)->fluidNodeIndices,
- para->getParD(level)->numberOfFluidNodes,
+ para->getParD(level)->taggedFluidNodeIndices[CollisionTemplate::Default],
+ para->getParD(level)->numberOfTaggedFluidNodes[CollisionTemplate::Default],
CollisionTemplate::Default, -1);
//! 2. exchange information between GPUs
@@ -71,8 +71,8 @@ void CollisionAndExchange_streams::operator()(UpdateGrid27 *updateGrid, Paramete
//! 1. run collision for nodes which are at the border of the gpus/processes, running with WriteMacroVars in case probes sample on these nodes
//!
updateGrid->collisionUsingIndices( level, t,
- para->getParD(level)->fluidNodeIndicesBorder,
- para->getParD(level)->numberOfFluidNodesBorder,
+ para->getParD(level)->taggedFluidNodeIndices[CollisionTemplate::Border],
+ para->getParD(level)->numberOfTaggedFluidNodes[CollisionTemplate::Border],
CollisionTemplate::WriteMacroVars,
borderStreamIndex);
@@ -86,8 +86,8 @@ void CollisionAndExchange_streams::operator()(UpdateGrid27 *updateGrid, Paramete
//!
para->getStreamManager()->waitOnStartBulkKernelEvent(bulkStreamIndex);
updateGrid->collisionUsingIndices( level, t,
- para->getParD(level)->fluidNodeIndices,
- para->getParD(level)->numberOfFluidNodes,
+ para->getParD(level)->taggedFluidNodeIndices[CollisionTemplate::Default],
+ para->getParD(level)->numberOfTaggedFluidNodes[CollisionTemplate::Default],
CollisionTemplate::Default,
bulkStreamIndex);
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp
index e197fb5c2..c996525ee 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp
@@ -29,15 +29,10 @@ void GridProvider::setNumberOfNodes(const int numberOfNodes, const int level) co
para->getParD(level)->mem_size_int_SP = sizeof(uint) * para->getParD(level)->numberOfNodes;
}
-void GridProvider::setNumberOfFluidNodes(const int numberOfNodes, const int level) const
+void GridProvider::setNumberOfTaggedFluidNodes(const int numberOfNodes, CollisionTemplate tag, const int level) const
{
- para->getParH(level)->numberOfFluidNodes = numberOfNodes;
- para->getParD(level)->numberOfFluidNodes = numberOfNodes;
-}
-
-void GridProvider::setNumberOfFluidNodesBorder(const int numberOfNodes, const int level) const {
- para->getParH(level)->numberOfFluidNodesBorder = numberOfNodes;
- para->getParD(level)->numberOfFluidNodesBorder = numberOfNodes;
+ para->getParH(level)->numberOfTaggedFluidNodes[tag] = numberOfNodes;
+ para->getParD(level)->numberOfTaggedFluidNodes[tag] = numberOfNodes;
}
void GridProvider::setInitalNodeValues(const int numberOfNodes, const int level) const
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.h b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.h
index 0ba30d121..45bf6d908 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.h
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.h
@@ -28,8 +28,7 @@ public:
virtual void allocArrays_BoundaryValues() = 0;
virtual void allocArrays_BoundaryQs() = 0;
virtual void allocArrays_OffsetScale() = 0;
- virtual void allocArrays_fluidNodeIndices() = 0;
- virtual void allocArrays_fluidNodeIndicesBorder() = 0;
+ virtual void allocArrays_taggedFluidNodes() = 0;
virtual void tagFluidNodeIndices(std::vector<uint> taggedFluidNodeIndices, CollisionTemplate tag, uint level) = 0;
virtual void sortFluidNodeTags() = 0;
@@ -48,8 +47,7 @@ public:
protected:
void setNumberOfNodes(const int numberOfNodes, const int level) const;
- void setNumberOfFluidNodes(const int numberOfNodes, const int level) const;
- void setNumberOfFluidNodesBorder(const int numberOfNodes, const int level) const;
+ void setNumberOfTaggedFluidNodes(const int numberOfNodes, CollisionTemplate tag, const int level) const;
virtual void setInitalNodeValues(const int numberOfNodes, const int level) const;
void setPressSizePerLevel(int level, int sizePerLevel) const;
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.cpp
index e423cf871..793400869 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.cpp
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.cpp
@@ -218,16 +218,11 @@ void GridReader::allocArrays_OffsetScale()
std::cout << "-----Ende OffsetScale------" << std::endl;
}
-void GridReader::allocArrays_fluidNodeIndices() {
+void GridReader::allocArrays_taggedFluidNodes() {
std::cout << "GridReader::allocArrays_fluidNodeIndices not implemented" << std::endl;
// TODO
}
-void GridReader::allocArrays_fluidNodeIndicesBorder() {
- std::cout << "GridReader::allocArrays_fluidNodeIndicesBorder not implemented" << std::endl;
- // TODO
-}
-
void GridReader::tagFluidNodeIndices(std::vector<uint> taggedFluidNodeIndices, CollisionTemplate tag, uint level){
std::cout << "GridReader::tagFluidNodeIndices not implemented" << std::endl;
// TODO
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.h b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.h
index 70ddb1f47..2c17c2802 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.h
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.h
@@ -39,8 +39,7 @@ public:
void allocArrays_CoordNeighborGeo() override;
void allocArrays_BoundaryValues() override;
void allocArrays_OffsetScale() override;
- void allocArrays_fluidNodeIndices() override;
- void allocArrays_fluidNodeIndicesBorder() override;
+ void allocArrays_taggedFluidNodes() override;
void tagFluidNodeIndices(std::vector<uint> taggedFluidNodeIndices, CollisionTemplate tag, uint level) override;
void sortFluidNodeTags() override;
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
index c9aaa1d1b..51c14a69e 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
@@ -95,21 +95,48 @@ void GridGenerator::allocArrays_CoordNeighborGeo()
std::cout << "-----finish Coord, Neighbor, Geo------" << std::endl;
}
-void GridGenerator::allocArrays_fluidNodeIndices() {
- for (uint level = 0; level < builder->getNumberOfGridLevels(); level++) {
- setNumberOfFluidNodes(builder->getNumberOfFluidNodes(level), level);
- cudaMemoryManager->cudaAllocFluidNodeIndices(level);
- builder->getFluidNodeIndices(para->getParH(level)->fluidNodeIndices, level);
- cudaMemoryManager->cudaCopyFluidNodeIndices(level);
- }
-}
+void GridGenerator::allocArrays_taggedFluidNodes() {
+ for (uint level = 0; level < builder->getNumberOfGridLevels(); level++)
+ {
+ for ( int tag = 0; tag < int(CollisionTemplate::LAST); tag++ )
+ { //TODO: Need to add CollisionTemplate to GridBuilder to allow as argument and get rid of indivual get funtions for fluid node indices
+ switch(static_cast<CollisionTemplate>(tag))
+ {
+ case CollisionTemplate::Default:
+ this->setNumberOfTaggedFluidNodes(builder->getNumberOfFluidNodes(level), CollisionTemplate::Default, level);
+ cudaMemoryManager->cudaAllocTaggedFluidNodeIndices(CollisionTemplate::Default, level);
+ builder->getFluidNodeIndices(para->getParH(level)->taggedFluidNodeIndices[CollisionTemplate::Default], level);
+ cudaMemoryManager->cudaCopyTaggedFluidNodeIndices(CollisionTemplate::Default, level);
+ break;
+ case CollisionTemplate::Border:
+ this->setNumberOfTaggedFluidNodes(builder->getNumberOfFluidNodesBorder(level), CollisionTemplate::Border, level);
+ cudaMemoryManager->cudaAllocTaggedFluidNodeIndices(CollisionTemplate::Border, level);
+ builder->getFluidNodeIndicesBorder(para->getParH(level)->taggedFluidNodeIndices[CollisionTemplate::Border], level);
+ cudaMemoryManager->cudaCopyTaggedFluidNodeIndices(CollisionTemplate::Border, level);
+ break;
+ case CollisionTemplate::WriteMacroVars:
+ this->setNumberOfTaggedFluidNodes(builder->getNumberOfFluidNodesMacroVars(level), CollisionTemplate::WriteMacroVars, level);
+ cudaMemoryManager->cudaAllocTaggedFluidNodeIndices(CollisionTemplate::WriteMacroVars, level);
+ builder->getFluidNodeIndicesMacroVars(para->getParH(level)->taggedFluidNodeIndices[CollisionTemplate::WriteMacroVars], level);
+ cudaMemoryManager->cudaCopyTaggedFluidNodeIndices(CollisionTemplate::WriteMacroVars, level);
+ break;
+ case CollisionTemplate::ApplyBodyForce:
+ this->setNumberOfTaggedFluidNodes(builder->getNumberOfFluidNodesApplyBodyForce(level), CollisionTemplate::ApplyBodyForce, level);
+ cudaMemoryManager->cudaAllocTaggedFluidNodeIndices(CollisionTemplate::ApplyBodyForce, level);
+ builder->getFluidNodeIndicesApplyBodyForce(para->getParH(level)->taggedFluidNodeIndices[CollisionTemplate::ApplyBodyForce], level);
+ cudaMemoryManager->cudaCopyTaggedFluidNodeIndices(CollisionTemplate::ApplyBodyForce, level);
+ break;
+ case CollisionTemplate::AllFeatures:
+ this->setNumberOfTaggedFluidNodes(builder->getNumberOfFluidNodesApplyBodyForce(level), CollisionTemplate::AllFeatures, level);
+ cudaMemoryManager->cudaAllocTaggedFluidNodeIndices(CollisionTemplate::AllFeatures, level);
+ builder->getFluidNodeIndicesAllFeatures(para->getParH(level)->taggedFluidNodeIndices[CollisionTemplate::AllFeatures], level);
+ cudaMemoryManager->cudaCopyTaggedFluidNodeIndices(CollisionTemplate::AllFeatures, level);
+ break;
+ default:
+ break;
+ }
-void GridGenerator::allocArrays_fluidNodeIndicesBorder() {
- for (uint level = 0; level < builder->getNumberOfGridLevels(); level++) {
- setNumberOfFluidNodesBorder(builder->getNumberOfFluidNodesBorder(level), level);
- cudaMemoryManager->cudaAllocFluidNodeIndicesBorder(level);
- builder->getFluidNodeIndicesBorder(para->getParH(level)->fluidNodeIndicesBorder, level);
- cudaMemoryManager->cudaCopyFluidNodeIndicesBorder(level);
+ }
}
}
@@ -126,7 +153,8 @@ void GridGenerator::tagFluidNodeIndices(std::vector<uint> taggedFluidNodeIndices
builder->addFluidNodeIndicesAllFeatures( taggedFluidNodeIndices, level );
break;
case CollisionTemplate::Default:
- throw std::runtime_error("Cannot tag fluid nodes as CollisionTemplate::Default!");
+ case CollisionTemplate::Border:
+ throw std::runtime_error("Cannot tag fluid nodes as Default or Border!");
default:
throw std::runtime_error("Tagging fluid nodes with invald tag!");
break;
@@ -138,7 +166,9 @@ void GridGenerator::tagFluidNodeIndices(std::vector<uint> taggedFluidNodeIndices
void GridGenerator::sortFluidNodeTags() {
for (uint level = 0; level < builder->getNumberOfGridLevels(); level++)
{
-
+ builder->sortFluidNodeIndicesAllFeatures(level); //has to be called first!
+ builder->sortFluidNodeIndicesMacroVars(level);
+ builder->sortFluidNodeIndicesApplyBodyForce(level);
}
}
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.h b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.h
index fed5ac7f4..82f6cedca 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.h
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.h
@@ -71,8 +71,7 @@ public:
//! \brief allocates and initialized the sub-grid distances at the boundary conditions
void allocArrays_BoundaryQs() override;
void allocArrays_OffsetScale() override;
- void allocArrays_fluidNodeIndices() override;
- void allocArrays_fluidNodeIndicesBorder() override;
+ void allocArrays_taggedFluidNodes() override;
void tagFluidNodeIndices(std::vector<uint> taggedFluidNodeIndices, CollisionTemplate tag, uint level) override;
void sortFluidNodeTags() override;
diff --git a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
index 221922169..a0845b49b 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
+++ b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
@@ -2921,48 +2921,25 @@ void CudaMemoryManager::cudaFree2ndOrderDerivitivesIsoTest(int lev)
}
-void CudaMemoryManager::cudaAllocFluidNodeIndices(int lev) {
- uint mem_size_geo_fluid_nodes = sizeof(uint) * parameter->getParH(lev)->numberOfFluidNodes;
+void CudaMemoryManager::cudaAllocTaggedFluidNodeIndices(CollisionTemplate tag, int lev) {
+ uint mem_size_tagged_fluid_nodes = sizeof(uint) * parameter->getParH(lev)->numberOfTaggedFluidNodes[tag];
// Host
- checkCudaErrors(cudaMallocHost((void **)&(parameter->getParH(lev)->fluidNodeIndices), mem_size_geo_fluid_nodes));
+ checkCudaErrors(cudaMallocHost((void **)&(parameter->getParH(lev)->taggedFluidNodeIndices[tag]), mem_size_tagged_fluid_nodes));
// Device
- checkCudaErrors(cudaMalloc((void **)&(parameter->getParD(lev)->fluidNodeIndices), mem_size_geo_fluid_nodes));
+ checkCudaErrors(cudaMalloc((void **)&(parameter->getParD(lev)->taggedFluidNodeIndices[tag]), mem_size_tagged_fluid_nodes));
//////////////////////////////////////////////////////////////////////////
- setMemsizeGPU((double)mem_size_geo_fluid_nodes, false);
+ setMemsizeGPU((double)mem_size_tagged_fluid_nodes, false);
}
-void CudaMemoryManager::cudaCopyFluidNodeIndices(int lev) {
- uint mem_size_geo_fluid_nodes = sizeof(uint) * parameter->getParH(lev)->numberOfFluidNodes;
- checkCudaErrors(cudaMemcpy(parameter->getParD(lev)->fluidNodeIndices,
- parameter->getParH(lev)->fluidNodeIndices,
- mem_size_geo_fluid_nodes, cudaMemcpyHostToDevice));
+void CudaMemoryManager::cudaCopyTaggedFluidNodeIndices(CollisionTemplate tag, int lev) {
+ uint mem_size_tagged_fluid_nodes = sizeof(uint) * parameter->getParH(lev)->numberOfTaggedFluidNodes[tag];
+ checkCudaErrors(cudaMemcpy(parameter->getParD(lev)->taggedFluidNodeIndices[tag],
+ parameter->getParH(lev)->taggedFluidNodeIndices[tag],
+ mem_size_tagged_fluid_nodes, cudaMemcpyHostToDevice));
}
-void CudaMemoryManager::cudaFreeFluidNodeIndices(int lev) {
- checkCudaErrors(cudaFreeHost(parameter->getParH(lev)->fluidNodeIndices));
-}
-
-void CudaMemoryManager::cudaAllocFluidNodeIndicesBorder(int lev) {
- uint mem_size_fluid_nodes_border = sizeof(uint) * parameter->getParH(lev)->numberOfFluidNodesBorder;
- // Host
- checkCudaErrors(
- cudaMallocHost((void **)&(parameter->getParH(lev)->fluidNodeIndicesBorder), mem_size_fluid_nodes_border));
- // Device
- checkCudaErrors(
- cudaMalloc((void **)&(parameter->getParD(lev)->fluidNodeIndicesBorder), mem_size_fluid_nodes_border));
- //////////////////////////////////////////////////////////////////////////
- setMemsizeGPU((double)mem_size_fluid_nodes_border, false);
-}
-
-void CudaMemoryManager::cudaCopyFluidNodeIndicesBorder(int lev) {
- uint mem_size_fluid_nodes_border = sizeof(uint) * parameter->getParH(lev)->numberOfFluidNodesBorder;
- checkCudaErrors(cudaMemcpy(parameter->getParD(lev)->fluidNodeIndicesBorder,
- parameter->getParH(lev)->fluidNodeIndicesBorder,
- mem_size_fluid_nodes_border, cudaMemcpyHostToDevice));
-}
-
-void CudaMemoryManager::cudaFreeFluidNodeIndicesBorder(int lev) {
- checkCudaErrors(cudaFreeHost(parameter->getParH(lev)->fluidNodeIndicesBorder));
+void CudaMemoryManager::cudaFreeTaggedFluidNodeIndices(CollisionTemplate tag, int lev) {
+ checkCudaErrors(cudaFreeHost(parameter->getParH(lev)->taggedFluidNodeIndices[tag]));
}
////////////////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.h b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.h
index 8df1b412c..1007d7400 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.h
+++ b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.h
@@ -345,12 +345,9 @@ public:
void cudaCopyProcessNeighborADZIndex(int lev, unsigned int processNeighbor);
void cudaFreeProcessNeighborADZ(int lev, unsigned int processNeighbor);
- void cudaAllocFluidNodeIndices(int lev);
- void cudaCopyFluidNodeIndices(int lev);
- void cudaFreeFluidNodeIndices(int lev);
- void cudaAllocFluidNodeIndicesBorder(int lev);
- void cudaCopyFluidNodeIndicesBorder(int lev);
- void cudaFreeFluidNodeIndicesBorder(int lev);
+ void cudaAllocTaggedFluidNodeIndices(CollisionTemplate tag, int lev);
+ void cudaCopyTaggedFluidNodeIndices(CollisionTemplate tag, int lev);
+ void cudaFreeTaggedFluidNodeIndices(CollisionTemplate tag, int lev);
// Actuator Line
void cudaAllocBladeRadii(ActuatorLine* actuatorLine);
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Almighty/CumulantK17Almighty.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Almighty/CumulantK17Almighty.cu
index 64db26341..449876899 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Almighty/CumulantK17Almighty.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Almighty/CumulantK17Almighty.cu
@@ -30,8 +30,8 @@ void CumulantK17Almighty<turbulenceModel>::run()
para->getParD(level)->forceX_SP, para->getParD(level)->forceY_SP, para->getParD(level)->forceZ_SP,
para->getQuadricLimitersDev(),
para->getParD(level)->isEvenTimestep,
- para->getParD(level)->fluidNodeIndices,
- para->getParD(level)->numberOfFluidNodes);
+ para->getParD(level)->taggedFluidNodeIndices[CollisionTemplate::Default],
+ para->getParD(level)->numberOfTaggedFluidNodes[CollisionTemplate::Default]);
getLastCudaError("LB_Kernel_CumulantK17Almighty execution failed");
}
@@ -83,6 +83,7 @@ void CumulantK17Almighty<turbulenceModel>::runOnIndices( const unsigned int *ind
size_indices);
break;
+ case CollisionTemplate::Border:
case CollisionTemplate::AllFeatures:
LB_Kernel_CumulantK17Almighty < turbulenceModel, true, true > <<< cudaGrid.grid, cudaGrid.threads, 0, stream >>>( para->getParD(level)->omega,
para->getParD(level)->typeOfGridNode,
diff --git a/src/gpu/VirtualFluids_GPU/LBM/LB.h b/src/gpu/VirtualFluids_GPU/LBM/LB.h
index 0b93abecb..7a13dd45a 100644
--- a/src/gpu/VirtualFluids_GPU/LBM/LB.h
+++ b/src/gpu/VirtualFluids_GPU/LBM/LB.h
@@ -67,13 +67,17 @@ enum class TurbulenceModel {
//! \brief An enumeration for selecting a template of the collision kernel (CumulantK17Almighty)
enum class CollisionTemplate {
//! - Default: plain collision without additional read/write
- Default,
- //! - WriteMacroVars: write out macroscopic variables in the collision kernel
- WriteMacroVars,
- //! - ApplyBodyForce: read and apply body force in the collision kernel
- ApplyBodyForce,
- //! - AllFeatures: write out macroscopic variables AND read and apply body force
- AllFeatures
+ Default = 0,
+ //! - Border: collision on border nodes
+ Border = 1,
+ //! - WriteMacroVars: collision \w write out macroscopic variables
+ WriteMacroVars = 2,
+ //! - ApplyBodyForce: collision \w read and apply body force in the collision kernel
+ ApplyBodyForce = 3,
+ //! - AllFeatures: collision \w write out macroscopic variables AND read and apply body force
+ AllFeatures = 4,
+ //!
+ LAST = 5
};
struct InitCondition
diff --git a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
index 3a0f251d2..4015b7fce 100644
--- a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
+++ b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
@@ -132,20 +132,18 @@ void Simulation::init(GridProvider &gridProvider, BoundaryConditionFactory *bcFa
for (SPtr<PreCollisionInteractor> actuator : para->getActuators()) {
actuator->init(para.get(), &gridProvider, cudaMemoryManager.get());
- actuator->getInteractorFluidNodes( para.get(), &gridProvider );
+ actuator->getTaggedFluidNodes( para.get(), &gridProvider );
}
for (SPtr<PreCollisionInteractor> probe : para->getProbes()) {
probe->init(para.get(), &gridProvider, cudaMemoryManager.get());
- probe->getInteractorFluidNodes( para.get(), &gridProvider );
+ probe->getTaggedFluidNodes( para.get(), &gridProvider );
}
//////////////////////////////////////////////////////////////////////////
- if (para->getKernelNeedsFluidNodeIndicesToRun()) {
- gridProvider.allocArrays_fluidNodeIndices();
- gridProvider.allocArrays_fluidNodeIndicesBorder();
- }
-
+ if (para->getKernelNeedsFluidNodeIndicesToRun())
+ gridProvider.allocArrays_taggedFluidNodes();
+
//////////////////////////////////////////////////////////////////////////
// Kernel init
//////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
index 77351d4b8..a2d26d65d 100644
--- a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
+++ b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
@@ -368,16 +368,24 @@ struct LBMSimulationParameter {
std::vector<EdgeNodePositions> edgeNodesYtoZ;
///////////////////////////////////////////////////////
- uint *fluidNodeIndices;
- uint numberOfFluidNodes;
- uint *fluidNodeIndicesBorder;
- uint numberOfFluidNodesBorder;
- uint *fluidNodeIndicesWriteMacroscopicVariables;
- uint numberOfFluidNodeIndicesWriteMacroscopicVariables;
- uint *fluidNodeIndicesApplyBodyForce;
- uint numberOfFluidNodeIndicesApplyBodyForce;
- uint *indicesWriteMacroscopicVariablesAndApplyBodyForce;
- uint numberOfIndicesWriteMacroscopicVariablesAndApplyBodyForce;
+ std::map<CollisionTemplate, uint*> taggedFluidNodeIndices = {{CollisionTemplate::Default, nullptr},
+ {CollisionTemplate::Border, nullptr},
+ {CollisionTemplate::WriteMacroVars, nullptr},
+ {CollisionTemplate::ApplyBodyForce, nullptr},
+ {CollisionTemplate::AllFeatures, nullptr}};
+ std::map<CollisionTemplate, uint > numberOfTaggedFluidNodes = {{CollisionTemplate::Default, 0},
+ {CollisionTemplate::Border, 0},
+ {CollisionTemplate::WriteMacroVars, 0},
+ {CollisionTemplate::ApplyBodyForce, 0},
+ {CollisionTemplate::AllFeatures, 0}};
+ // uint *fluidNodeIndicesBorder;
+ // uint numberOfFluidNodesBorder;
+ // uint *fluidNodeIndicesWriteMacroscopicVariables;
+ // uint numberOfFluidNodeIndicesWriteMacroscopicVariables;
+ // uint *fluidNodeIndicesApplyBodyForce;
+ // uint numberOfFluidNodeIndicesApplyBodyForce;
+ // uint *indicesWriteMacroscopicVariablesAndApplyBodyForce;
+ // uint numberOfIndicesWriteMacroscopicVariablesAndApplyBodyForce;
};
//! \brief Class for LBM-parameter management
diff --git a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorLine.cu b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorLine.cu
index b489269d0..81f06fc3a 100644
--- a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorLine.cu
+++ b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorLine.cu
@@ -258,7 +258,7 @@ void ActuatorLine::free(Parameter* para, CudaMemoryManager* cudaMemoryManager)
cudaMemoryManager->cudaFreeSphereIndices(this);
}
-void ActuatorLine::getInteractorFluidNodes(Parameter *para, GridProvider* gridProvider)
+void ActuatorLine::getTaggedFluidNodes(Parameter *para, GridProvider* gridProvider)
{
std::vector<uint> indicesInSphere(this->boundingSphereIndicesH, this->boundingSphereIndicesH+this->nIndices);
gridProvider->tagFluidNodeIndices(indicesInSphere, CollisionTemplate::AllFeatures, this->level);
diff --git a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorLine.h b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorLine.h
index c107f4576..f115aba09 100644
--- a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorLine.h
+++ b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorLine.h
@@ -49,7 +49,7 @@ public:
void interact(Parameter* para, CudaMemoryManager* cudaMemoryManager, int level, uint t) override;
void free(Parameter* para, CudaMemoryManager* cudaMemoryManager) override;
void write(uint t);
- void getInteractorFluidNodes(Parameter *para, GridProvider* gridProvider) override;
+ void getTaggedFluidNodes(Parameter *para, GridProvider* gridProvider) override;
uint getNBladeNodes(){ return this->nBladeNodes; };
diff --git a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/PreCollisionInteractor.h b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/PreCollisionInteractor.h
index 368b85c2f..f9a87f613 100644
--- a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/PreCollisionInteractor.h
+++ b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/PreCollisionInteractor.h
@@ -33,7 +33,7 @@ public:
virtual void init(Parameter *para, GridProvider *gridProvider, CudaMemoryManager *cudaMemoryManager) = 0;
virtual void interact(Parameter *para, CudaMemoryManager *cudaMemoryManager, int level, uint t) = 0;
virtual void free(Parameter *para, CudaMemoryManager *cudaMemoryManager) = 0;
- virtual void getInteractorFluidNodes(Parameter *para, GridProvider* gridProvider) = 0;
+ virtual void getTaggedFluidNodes(Parameter *para, GridProvider* gridProvider) = 0;
protected:
uint updateInterval;
diff --git a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PlaneProbe.cu b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PlaneProbe.cu
index e0c574bba..8e0bf1faf 100644
--- a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PlaneProbe.cu
+++ b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PlaneProbe.cu
@@ -108,7 +108,7 @@ void PlaneProbe::calculateQuantities(SPtr<ProbeStruct> probeStruct, Parameter* p
probeStruct->quantitiesD, probeStruct->arrayOffsetsD, probeStruct->quantitiesArrayD);
}
-void PlaneProbe::getInteractorFluidNodes(Parameter *para, GridProvider* gridProvider)
+void PlaneProbe::getTaggedFluidNodes(Parameter *para, GridProvider* gridProvider)
{
for(int level=0; level<=para->getMaxLevel(); level++)
{
diff --git a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PlaneProbe.h b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PlaneProbe.h
index 559ed7497..c53cb5020 100644
--- a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PlaneProbe.h
+++ b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PlaneProbe.h
@@ -72,7 +72,7 @@ public:
this->deltaZ = _deltaZ;
}
- void getInteractorFluidNodes(Parameter *para, GridProvider* gridProvider) override;
+ void getTaggedFluidNodes(Parameter *para, GridProvider* gridProvider) override;
private:
bool isAvailableStatistic(Statistic _variable) override;
diff --git a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PointProbe.cu b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PointProbe.cu
index 25956f3db..37beeed05 100644
--- a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PointProbe.cu
+++ b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PointProbe.cu
@@ -142,7 +142,7 @@ void PointProbe::addProbePointsFromXNormalPlane(real pos_x, real pos0_y, real po
}
-void PointProbe::getInteractorFluidNodes(Parameter *para, GridProvider* gridProvider)
+void PointProbe::getTaggedFluidNodes(Parameter *para, GridProvider* gridProvider)
{
for(int level=0; level<=para->getMaxLevel(); level++)
{
diff --git a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PointProbe.h b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PointProbe.h
index 8cdf225eb..08c359705 100644
--- a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PointProbe.h
+++ b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PointProbe.h
@@ -64,7 +64,7 @@ public:
void addProbePointsFromList(std::vector<real>& _pointCoordsX, std::vector<real>& _pointCoordsY, std::vector<real>& _pointCoordsZ);
void addProbePointsFromXNormalPlane(real pos_x, real pos0_y, real pos0_z, real pos1_y, real pos1_z, uint n_y, uint n_z);
- void getInteractorFluidNodes(Parameter *para, GridProvider* gridProvider) override;
+ void getTaggedFluidNodes(Parameter *para, GridProvider* gridProvider) override;
private:
bool isAvailableStatistic(Statistic _variable) override;
diff --git a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/Probe.cu b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/Probe.cu
index a4fd7c0ec..594ce2773 100644
--- a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/Probe.cu
+++ b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/Probe.cu
@@ -315,7 +315,7 @@ void Probe::free(Parameter* para, CudaMemoryManager* cudaMemoryManager)
}
}
-void Probe::getInteractorFluidNodes(Parameter *para, GridProvider* gridProvider)
+void Probe::getTaggedFluidNodes(Parameter *para, GridProvider* gridProvider)
{
// Do nothing
};
diff --git a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/Probe.h b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/Probe.h
index 3012fdbc5..7c2956e8c 100644
--- a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/Probe.h
+++ b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/Probe.h
@@ -159,7 +159,7 @@ public:
void init(Parameter* para, GridProvider* gridProvider, CudaMemoryManager* cudaMemoryManager) override;
void interact(Parameter* para, CudaMemoryManager* cudaMemoryManager, int level, uint t) override;
void free(Parameter* para, CudaMemoryManager* cudaMemoryManager) override;
- virtual void getInteractorFluidNodes(Parameter *para, GridProvider* gridProvider);
+ virtual void getTaggedFluidNodes(Parameter *para, GridProvider* gridProvider);
SPtr<ProbeStruct> getProbeStruct(int level){ return this->probeParams[level]; }
--
GitLab