From fa149556eeb2f88bb7026e1de1233a593127f8c0 Mon Sep 17 00:00:00 2001
From: Anna Wellmann <a.wellmann@tu-braunschweig.de>
Date: Thu, 5 Aug 2021 15:21:14 +0200
Subject: [PATCH] Add field for special (non-bulk) fluid nodes
---
src/gpu/GridGenerator/grid/Field.cu | 4 +--
src/gpu/GridGenerator/grid/Field.h | 2 +-
src/gpu/GridGenerator/grid/Grid.h | 4 ++-
.../grid/GridBuilder/GridBuilder.h | 2 ++
.../grid/GridBuilder/LevelGridBuilder.cpp | 23 ++++++++++----
.../grid/GridBuilder/LevelGridBuilder.h | 2 ++
src/gpu/GridGenerator/grid/GridImp.cu | 30 ++++++++++++++++---
src/gpu/GridGenerator/grid/GridImp.h | 5 ++++
.../Calculation/UpdateGrid27.cpp | 5 +++-
.../DataStructureInitializer/GridProvider.cpp | 5 ++++
.../DataStructureInitializer/GridProvider.h | 2 ++
.../GridReaderFiles/GridReader.cpp | 6 ++++
.../GridReaderFiles/GridReader.h | 1 +
.../GridReaderGenerator/GridGenerator.cpp | 12 ++++++--
.../GridReaderGenerator/GridGenerator.h | 1 +
.../GPU/CudaMemoryManager.cpp | 26 ++++++++++++++--
.../VirtualFluids_GPU/GPU/CudaMemoryManager.h | 4 ++-
src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp | 1 +
.../VirtualFluids_GPU/Parameter/Parameter.h | 2 ++
19 files changed, 117 insertions(+), 20 deletions(-)
diff --git a/src/gpu/GridGenerator/grid/Field.cu b/src/gpu/GridGenerator/grid/Field.cu
index eff280f5a..7982a23be 100644
--- a/src/gpu/GridGenerator/grid/Field.cu
+++ b/src/gpu/GridGenerator/grid/Field.cu
@@ -67,10 +67,10 @@ HOSTDEVICE bool Field::isFluid(uint index) const
return type == FLUID || type == FLUID_CFC || type == FLUID_CFF || type == FLUID_FCC || type == FLUID_FCF || isBoundaryConditionNode(index);
}
-HOSTDEVICE bool Field::isFluidBulk(uint index) const
+HOSTDEVICE bool Field::isFluidNodeOfSpecialInterest(uint index) const
{
const char type = field[index];
- return type == FLUID || type == FLUID_CFC || type == FLUID_CFF || type == FLUID_FCC || type == FLUID_FCF;
+ return isBoundaryConditionNode(index);
}
HOSTDEVICE bool Field::isInvalidSolid(uint index) const
diff --git a/src/gpu/GridGenerator/grid/Field.h b/src/gpu/GridGenerator/grid/Field.h
index 92b420d67..9e7513108 100644
--- a/src/gpu/GridGenerator/grid/Field.h
+++ b/src/gpu/GridGenerator/grid/Field.h
@@ -22,7 +22,7 @@ public:
HOSTDEVICE bool isCoarseToFineNode(uint index) const;
HOSTDEVICE bool isFineToCoarseNode(uint index) const;
HOSTDEVICE bool isFluid(uint index) const;
- HOSTDEVICE bool isFluidBulk(uint index) const;
+ HOSTDEVICE bool isFluidNodeOfSpecialInterest(uint index) const;
HOSTDEVICE bool isInvalidSolid(uint index) const;
HOSTDEVICE bool isQ(uint index) const;
HOSTDEVICE bool isBoundaryConditionNode(uint index) const;
diff --git a/src/gpu/GridGenerator/grid/Grid.h b/src/gpu/GridGenerator/grid/Grid.h
index 855196dba..32f571791 100644
--- a/src/gpu/GridGenerator/grid/Grid.h
+++ b/src/gpu/GridGenerator/grid/Grid.h
@@ -141,11 +141,13 @@ public:
virtual void repairCommunicationInices(int direction) = 0;
- // needed for CUDA Streams MultiGPU
+ // needed for CUDA Streams
virtual void findFluidNodeIndices(bool onlyBulk) = 0;
virtual uint getNumberOfFluidNodes() const = 0;;
virtual void getFluidNodeIndices(uint *fluidNodeIndices) const = 0;
+ virtual uint getNumberOfFluidNodesBorder() const = 0;
+ virtual void getFluidNodeIndicesBorder(uint *fluidNodeIndicesBorder) const = 0;
};
#endif
diff --git a/src/gpu/GridGenerator/grid/GridBuilder/GridBuilder.h b/src/gpu/GridGenerator/grid/GridBuilder/GridBuilder.h
index 0d12590c1..9ff99fad4 100644
--- a/src/gpu/GridGenerator/grid/GridBuilder/GridBuilder.h
+++ b/src/gpu/GridGenerator/grid/GridBuilder/GridBuilder.h
@@ -86,6 +86,8 @@ public:
virtual uint getNumberOfFluidNodes(unsigned int level) const = 0;
virtual void getFluidNodeIndices(uint *fluidNodeIndices, const int level) const = 0;
+ virtual uint getNumberOfFluidNodesBorder(unsigned int level) const = 0;
+ virtual void getFluidNodeIndicesBorder(uint *fluidNodeIndices, const int level) const = 0;
};
diff --git a/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.cpp b/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.cpp
index 46059f2e5..ba2555654 100644
--- a/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.cpp
+++ b/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.cpp
@@ -263,11 +263,6 @@ uint LevelGridBuilder::getNumberOfNodes(unsigned int level) const
return grids[level]->getSparseSize();
}
-uint LevelGridBuilder::getNumberOfFluidNodes(unsigned int level) const
-{
- return grids[level]->getNumberOfFluidNodes();
-}
-
std::shared_ptr<Grid> LevelGridBuilder::getGrid(int level, int box)
{
return this->grids[level];
@@ -298,10 +293,26 @@ void LevelGridBuilder::getNodeValues(real *xCoords, real *yCoords, real *zCoords
//TODO: add getSlipSize...
-GRIDGENERATOR_EXPORT void LevelGridBuilder::getFluidNodeIndices(uint *fluidNodeIndices, const int level) const {
+GRIDGENERATOR_EXPORT void LevelGridBuilder::getFluidNodeIndices(uint *fluidNodeIndices, const int level) const
+{
grids[level]->getFluidNodeIndices(fluidNodeIndices);
}
+GRIDGENERATOR_EXPORT void LevelGridBuilder::getFluidNodeIndicesBorder(uint *fluidNodeIndices, const int level) const
+{
+ grids[level]->getFluidNodeIndicesBorder(fluidNodeIndices);
+}
+
+uint LevelGridBuilder::getNumberOfFluidNodes(unsigned int level) const
+{
+ return grids[level]->getNumberOfFluidNodes();
+}
+
+GRIDGENERATOR_EXPORT uint LevelGridBuilder::getNumberOfFluidNodesBorder(unsigned int level) const
+{
+ return grids[level]->getNumberOfFluidNodesBorder();
+}
+
uint LevelGridBuilder::getVelocitySize(int level) const
{
uint size = 0;
diff --git a/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.h b/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.h
index 001c54e71..4cbb927a2 100644
--- a/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.h
+++ b/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.h
@@ -58,6 +58,8 @@ public:
GRIDGENERATOR_EXPORT virtual uint getNumberOfFluidNodes(unsigned int level) const override;
GRIDGENERATOR_EXPORT virtual void getFluidNodeIndices(uint* fluidNodeIndices, const int level) const override;
+ GRIDGENERATOR_EXPORT virtual uint getNumberOfFluidNodesBorder(unsigned int level) const override;
+ GRIDGENERATOR_EXPORT virtual void getFluidNodeIndicesBorder(uint *fluidNodeIndices, const int level) const override;
GRIDGENERATOR_EXPORT virtual void getNodeValues(real *xCoords, real *yCoords, real *zCoords,
uint *neighborX, uint *neighborY, uint *neighborZ, uint *neighborNegative,
diff --git a/src/gpu/GridGenerator/grid/GridImp.cu b/src/gpu/GridGenerator/grid/GridImp.cu
index 8bc404c3a..991ebc722 100644
--- a/src/gpu/GridGenerator/grid/GridImp.cu
+++ b/src/gpu/GridGenerator/grid/GridImp.cu
@@ -865,12 +865,22 @@ CUDA_HOST void GridImp::findFluidNodeIndices(bool onlyBulk)
int sparseIndex = this->getSparseIndex(index);
if (sparseIndex == -1)
continue;
- if ((onlyBulk && this->field.isFluidBulk(index)) || (!onlyBulk && this->field.isFluid(index)))
- this->fluidNodeIndices.push_back((uint)sparseIndex + 1); // + 1 for numbering shift between GridGenerator and VF_GPU
+
+ // + 1 for numbering shift between GridGenerator and VF_GPU
+ // When onlyBulk: push indices of fluid nodes in bulk to "fluidNodeIndices" and push indices of special fluid nodes (not in bulk) to fluidNodeIndicesBorder
+ // When not onlyBulk: push indices of all fluid nodes to "fluidNodeIndices"
+ if (this->field.isFluid(index)) {
+ if (onlyBulk)
+ if (this->field.isFluidNodeOfSpecialInterest(index))
+ this->fluidNodeIndicesBorder.push_back((uint)sparseIndex + 1);
+ else
+ this->fluidNodeIndices.push_back((uint)sparseIndex + 1);
+ else
+ this->fluidNodeIndices.push_back((uint)sparseIndex + 1);
+ }
}
}
-
HOSTDEVICE void GridImp::setNeighborIndices(uint index)
{
real x, y, z;
@@ -1956,11 +1966,23 @@ CUDA_HOST void GridImp::getNodeValues(real *xCoords, real *yCoords, real *zCoord
}
}
-CUDA_HOST void GridImp::getFluidNodeIndices(uint *fluidNodeIndices) const {
+CUDA_HOST void GridImp::getFluidNodeIndices(uint *fluidNodeIndices) const
+{
for (uint nodeNumber = 0; nodeNumber < (uint)this->fluidNodeIndices.size(); nodeNumber++)
fluidNodeIndices[nodeNumber] = this->fluidNodeIndices[nodeNumber];
}
+uint GridImp::getNumberOfFluidNodesBorder() const
+{
+ return this->fluidNodeIndicesBorder.size();
+}
+
+void GridImp::getFluidNodeIndicesBorder(uint *fluidNodeIndicesBorder) const
+{
+ for (uint nodeNumber = 0; nodeNumber < (uint)this->fluidNodeIndicesBorder.size(); nodeNumber++)
+ fluidNodeIndicesBorder[nodeNumber] = this->fluidNodeIndicesBorder[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 33953ab2a..ffdb60665 100644
--- a/src/gpu/GridGenerator/grid/GridImp.h
+++ b/src/gpu/GridGenerator/grid/GridImp.h
@@ -87,6 +87,7 @@ private:
int *sparseIndices;
std::vector<uint> fluidNodeIndices;
+ std::vector<uint> fluidNodeIndicesBorder;
uint *qIndices; //maps from matrix index to qIndex
real *qValues;
@@ -316,9 +317,13 @@ public:
void repairCommunicationInices(int direction) override;
void findFluidNodeIndices(bool onlyBulk) override;
+
uint getNumberOfFluidNodes() const override;
CUDA_HOST void getFluidNodeIndices(uint *fluidNodeIndices) const override;
+ uint getNumberOfFluidNodesBorder() const override;
+ void getFluidNodeIndicesBorder(uint *fluidNodeIndicesBorder) const override;
+
public:
diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
index 9cb31dac8..445244ac7 100644
--- a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
+++ b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
@@ -26,9 +26,12 @@ void updateGrid27(Parameter* para,
//////////////////////////////////////////////////////////////////////////
- if (para->useStreams)
+ if (para->useStreams) {
collision(para, pm, level, t, kernels, para->getParD(level)->fluidNodeIndices,
para->getParD(level)->numberOfFluidNodes);
+ collision(para, pm, level, t, kernels, para->getParD(level)->fluidNodeIndicesBorder,
+ para->getParD(level)->numberOffluidNodesBorder);
+ }
else
collision(para, pm, level, t, kernels);
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp
index 22d3d3a38..c3c9bc3d8 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp
@@ -35,6 +35,11 @@ void GridProvider::setNumberOfFluidNodes(const int numberOfNodes, const int leve
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;
+}
+
void GridProvider::setInitalNodeValues(const int numberOfNodes, const int level) const
{
for (int j = 1; j <= numberOfNodes; j++)
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.h b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.h
index d8f442115..d12a25733 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.h
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.h
@@ -26,6 +26,7 @@ public:
virtual void allocArrays_BoundaryQs() = 0;
virtual void allocArrays_OffsetScale() = 0;
virtual void allocArrays_fluidNodeIndices() = 0;
+ virtual void allocArrays_fluidNodeIndicesBorder() = 0;
virtual void setDimensions() = 0;
virtual void setBoundingBox() = 0;
@@ -42,6 +43,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;
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 c2e19b659..d5f7b6b21 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.cpp
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.cpp
@@ -217,9 +217,15 @@ void GridReader::allocArrays_OffsetScale()
}
void GridReader::allocArrays_fluidNodeIndices() {
+ std::cout << "GridReader::allocArrays_fluidNodeIndices not imlemented" << std::endl;
// TODO
}
+void GridReader::allocArrays_fluidNodeIndicesBorder() {
+ std::cout << "GridReader::allocArrays_fluidNodeIndicesBorder not imlemented" << std::endl;
+ // TODO
+}
+
void GridReader::setPressureValues(int channelSide) const
{
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.h b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.h
index cbb0fac42..ae592f9bd 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.h
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.h
@@ -40,6 +40,7 @@ public:
void allocArrays_BoundaryValues() override;
void allocArrays_OffsetScale() override;
void allocArrays_fluidNodeIndices() override;
+ void allocArrays_fluidNodeIndicesBorder() override;
void initalValuesDomainDecompostion(int level);
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
index ec7056460..b96646142 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
@@ -88,8 +88,7 @@ void GridGenerator::allocArrays_CoordNeighborGeo()
}
void GridGenerator::allocArrays_fluidNodeIndices() {
- const uint numberOfLevels = builder->getNumberOfGridLevels();
- for (uint level = 0; level < numberOfLevels; level++) {
+ for (uint level = 0; level < builder->getNumberOfGridLevels(); level++) {
setNumberOfFluidNodes(builder->getNumberOfFluidNodes(level), level);
cudaMemoryManager->cudaAllocFluidNodeIndices(level);
builder->getFluidNodeIndices(para->getParH(level)->fluidNodeIndices, level);
@@ -97,6 +96,15 @@ void GridGenerator::allocArrays_fluidNodeIndices() {
}
}
+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);
+ }
+}
+
void GridGenerator::allocArrays_BoundaryValues()
{
std::cout << "------read BoundaryValues------" << std::endl;
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.h b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.h
index 096bb424d..5659cad85 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.h
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.h
@@ -30,6 +30,7 @@ public:
void allocArrays_BoundaryQs() override;
void allocArrays_OffsetScale() override;
void allocArrays_fluidNodeIndices() override;
+ void allocArrays_fluidNodeIndicesBorder() override;
virtual void setDimensions() override;
virtual void setBoundingBox() override;
diff --git a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
index c101b9295..e4fe89cb9 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
+++ b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
@@ -2652,8 +2652,7 @@ void CudaMemoryManager::cudaAllocFluidNodeIndices(int lev) {
// Device
checkCudaErrors(cudaMalloc((void **)&(parameter->getParD(lev)->fluidNodeIndices), mem_size_geo_fluid_nodes));
//////////////////////////////////////////////////////////////////////////
- double tmp = (double)mem_size_geo_fluid_nodes;
- setMemsizeGPU(tmp, false);
+ setMemsizeGPU((double)mem_size_geo_fluid_nodes, false);
}
void CudaMemoryManager::cudaCopyFluidNodeIndices(int lev) {
@@ -2667,6 +2666,29 @@ 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));
+}
+
diff --git a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.h b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.h
index 0588b2c32..492d9b9be 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.h
+++ b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.h
@@ -321,11 +321,13 @@ public:
void cudaCopyProcessNeighborADZFsDH(int lev, unsigned int processNeighbor);
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);
diff --git a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
index 08703cdff..0f91ad960 100644
--- a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
+++ b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
@@ -87,6 +87,7 @@ void Simulation::init(SPtr<Parameter> para, SPtr<GridProvider> gridProvider, std
gridProvider->allocAndCopyForcing();
gridProvider->allocAndCopyQuadricLimiters();
gridProvider->allocArrays_fluidNodeIndices();
+ gridProvider->allocArrays_fluidNodeIndicesBorder();
gridProvider->setDimensions();
gridProvider->setBoundingBox();
diff --git a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
index d44c3abd6..9f70266e4 100644
--- a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
+++ b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
@@ -312,6 +312,8 @@ struct LBMSimulationParameter
uint *fluidNodeIndices;
uint numberOfFluidNodes;
+ uint *fluidNodeIndicesBorder;
+ uint numberOffluidNodesBorder;
};
class VIRTUALFLUIDS_GPU_EXPORT Parameter
--
GitLab