diff --git a/src/gpu/GridGenerator/grid/Field.cu b/src/gpu/GridGenerator/grid/Field.cu index eff280f5aea2a8fb7a1cc751d24bd13a790d163d..7982a23be0fa62c7afefdb339a3eb0f1513528d2 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 92b420d6756bdc71012305b483a9efb4641713fc..9e7513108fa039cc6b14ba519fce6acf667ed2f6 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 855196dba6070cd3eeeea282cf13003e9228b89b..32f57179182b069c3f2804e71e60824e831a4102 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 0d12590c18dc0b3de7b25a242f70a6439cb31310..9ff99fad4e2577536f50ce112acdc0e2e6bf40d5 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 46059f2e56c8572033f53f8641818c3ac437d759..ba2555654549fc708b6bcb91e69e499153252ab4 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 001c54e7169399907e6a4f108577f351732aa729..4cbb927a24373f06014109b1589aacd367a4d3a7 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 8bc404c3aa220e6aa55d9badd86ce138b1705ab8..991ebc722b950e3922a6fc91ae2cbb8ea5d5265d 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 33953ab2a82f3b47993961941780cfcb0f4e8443..ffdb606659a27d2f5491f978cf122019ceefea01 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 9cb31dac883590a5e9782c55f2e7df1948e7b645..445244ac761ed7dfa49f9fcaa385987e31955fa0 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 22d3d3a38d1364205f2386e377b3a0502993387a..c3c9bc3d8064dab29a21cb6e814425ddc2020529 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 d8f442115a965af58110dcc778768dc9f4210847..d12a25733539c319cdfd6ead1d6aa169fe6ae52d 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 c2e19b659d401ec94e14532a5b71fbb86e2d0fe6..d5f7b6b210de7bc0c4625df461664a29c05ffcd2 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 cbb0fac42f2e5e62dea82c63552a0d524c858d83..ae592f9bde145b4008581d62772e0a31cb7237f8 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 ec7056460e58da0eca90024975c74701ef56531d..b96646142a18835f51107cd07ca40170cf91beeb 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 096bb424d19262186d9afb4f491b059be8e160a2..5659cad85c42e949e26f5d9994e3755c796cd1c7 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 c101b9295e4b261e5572775b37368a332d0c9a1e..e4fe89cb9b68ea2464832f182cfb74ac3bb1e3b3 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 0588b2c3268b87664a5a4ce51d6de8bb16f580ed..492d9b9beca76f523e5a0b16732b0af76b1d3d56 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 08703cdff85435c9e346927983d74a102df15a15..0f91ad9605f25f7d34653ea30c116f7e2a52c6fd 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 d44c3abd6f061306315283b5fe573be08e5969d8..9f70266e4656863ca18fed53afe30196410c5d40 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