diff --git a/apps/gpu/LBM/ActuatorLine/ActuatorLine.cpp b/apps/gpu/LBM/ActuatorLine/ActuatorLine.cpp
index 254f7c8ce76999b6d731662638be8ea14714c005..aa6e52d1f2c20782ed34e3e4bba97ad873cc412e 100644
--- a/apps/gpu/LBM/ActuatorLine/ActuatorLine.cpp
+++ b/apps/gpu/LBM/ActuatorLine/ActuatorLine.cpp
@@ -83,6 +83,7 @@
#include "VirtualFluids_GPU/Factories/BoundaryConditionFactory.h"
#include "VirtualFluids_GPU/TurbulenceModels/TurbulenceModelFactory.h"
#include "VirtualFluids_GPU/Factories/GridScalingFactory.h"
+#include "VirtualFluids_GPU/Kernel/Utilities/KernelTypes.h"
#include "VirtualFluids_GPU/GPU/CudaMemoryManager.h"
@@ -192,7 +193,7 @@ void multipleLevel(const std::string& configPath)
para->setViscosityLB(viscosityLB);
para->setVelocityRatio( dx / dt );
para->setViscosityRatio( dx*dx/dt );
- para->setMainKernel("CumulantK17");
+ para->setMainKernel(vf::CollisionKernel::Compressible::CumulantK17);
para->setInitialCondition([&](real coordX, real coordY, real coordZ, real &rho, real &vx, real &vy, real &vz) {
rho = (real)0.0;
diff --git a/apps/gpu/LBM/BoundaryLayer/BoundaryLayer.cpp b/apps/gpu/LBM/BoundaryLayer/BoundaryLayer.cpp
index 4c91bf148d422838bdfe647b53e40c97166b9f4e..3921c85244ad27456e98c750fd64638453546ff5 100644
--- a/apps/gpu/LBM/BoundaryLayer/BoundaryLayer.cpp
+++ b/apps/gpu/LBM/BoundaryLayer/BoundaryLayer.cpp
@@ -89,6 +89,7 @@
#include "VirtualFluids_GPU/Factories/BoundaryConditionFactory.h"
#include "VirtualFluids_GPU/Factories/GridScalingFactory.h"
#include "VirtualFluids_GPU/TurbulenceModels/TurbulenceModelFactory.h"
+#include "VirtualFluids_GPU/Kernel/Utilities/KernelTypes.h"
#include "VirtualFluids_GPU/GPU/CudaMemoryManager.h"
@@ -230,7 +231,7 @@ void multipleLevel(const std::string& configPath)
bool useStreams = (nProcs > 1 ? true: false);
// useStreams=false;
para->setUseStreams(useStreams);
- para->setMainKernel("CumulantK17");
+ para->setMainKernel(vf::CollisionKernel::Compressible::CumulantK17);
para->setIsBodyForce( config.getValue<bool>("bodyForce") );
para->setTimestepStartOut(uint(tStartOut/dt) );
diff --git a/apps/gpu/LBM/ChannelFlow/ChannelFlow.cpp b/apps/gpu/LBM/ChannelFlow/ChannelFlow.cpp
index d61ee2755df6db086cb7c4d8384607dbe097a718..a05f3243040a2fbd0617daa65ac29322f45f7025 100644
--- a/apps/gpu/LBM/ChannelFlow/ChannelFlow.cpp
+++ b/apps/gpu/LBM/ChannelFlow/ChannelFlow.cpp
@@ -74,6 +74,7 @@
#include "VirtualFluids_GPU/LBM/Simulation.h"
#include "VirtualFluids_GPU/Output/FileWriter.h"
#include "VirtualFluids_GPU/Parameter/Parameter.h"
+#include "VirtualFluids_GPU/Kernel/Utilities/KernelTypes.h"
//////////////////////////////////////////////////////////////////////////
@@ -166,7 +167,7 @@ int main(int argc, char *argv[])
para->setTimestepEnd(timeStepEnd);
para->setOutputPrefix("ChannelFlow");
- para->setMainKernel("CumulantK17CompChimStream");
+ para->setMainKernel(vf::CollisionKernel::Compressible::CumulantK17);
const uint generatePart = vf::gpu::Communicator::getInstance().getPID();
real overlap = (real)8.0 * dx;
diff --git a/apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp b/apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp
index f543e04811468aa0e9b9e9b38b72096485bcda7c..a57191a4dd54c9b7ecb06048377acfe59d883277 100644
--- a/apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp
+++ b/apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp
@@ -69,6 +69,7 @@
#include "VirtualFluids_GPU/Output/FileWriter.h"
#include "VirtualFluids_GPU/Parameter/Parameter.h"
#include "VirtualFluids_GPU/Factories/GridScalingFactory.h"
+#include "VirtualFluids_GPU/Kernel/Utilities/KernelTypes.h"
//////////////////////////////////////////////////////////////////////////
@@ -153,7 +154,7 @@ int main()
para->setTimestepOut(timeStepOut);
para->setTimestepEnd(timeStepEnd);
- para->setMainKernel("CumulantK17");
+ para->setMainKernel(vf::CollisionKernel::Compressible::CumulantK17);
//////////////////////////////////////////////////////////////////////////
// set boundary conditions
diff --git a/apps/gpu/LBM/DrivenCavityMultiGPU/DrivenCavityMultiGPU.cpp b/apps/gpu/LBM/DrivenCavityMultiGPU/DrivenCavityMultiGPU.cpp
index 46e40044d7b3dca0cfc842ef775f680f0922e9b8..1bbb35310e3dcc0a1b56be7d486acfb7370a00f8 100755
--- a/apps/gpu/LBM/DrivenCavityMultiGPU/DrivenCavityMultiGPU.cpp
+++ b/apps/gpu/LBM/DrivenCavityMultiGPU/DrivenCavityMultiGPU.cpp
@@ -51,6 +51,7 @@
#include "VirtualFluids_GPU/PreProcessor/PreProcessorFactory/PreProcessorFactoryImp.h"
#include "VirtualFluids_GPU/Factories/BoundaryConditionFactory.h"
#include "VirtualFluids_GPU/Factories/GridScalingFactory.h"
+#include "VirtualFluids_GPU/Kernel/Utilities/KernelTypes.h"
#include "VirtualFluids_GPU/GPU/CudaMemoryManager.h"
@@ -125,7 +126,7 @@ void multipleLevel(std::filesystem::path& configPath)
para->setPrintFiles(true);
std::cout << "Write result files to " << para->getFName() << std::endl;
- para->setMainKernel("CumulantK17");
+ para->setMainKernel(vf::CollisionKernel::Compressible::CumulantK17);
scalingFactory.setScalingFactory(GridScalingFactory::GridScaling::ScaleCompressible);
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/apps/gpu/LBM/DrivenCavityUniform/DrivenCavity.cpp b/apps/gpu/LBM/DrivenCavityUniform/DrivenCavity.cpp
index d355a399f50f7a5ebb44690e57527f35b6aa81d8..83e5f41d2b7c9569744167ee1d1f674ab46c8439 100644
--- a/apps/gpu/LBM/DrivenCavityUniform/DrivenCavity.cpp
+++ b/apps/gpu/LBM/DrivenCavityUniform/DrivenCavity.cpp
@@ -69,6 +69,7 @@
#include "VirtualFluids_GPU/Output/FileWriter.h"
#include "VirtualFluids_GPU/Parameter/Parameter.h"
#include "VirtualFluids_GPU/Factories/GridScalingFactory.h"
+#include "VirtualFluids_GPU/Kernel/Utilities/KernelTypes.h"
//////////////////////////////////////////////////////////////////////////
@@ -154,7 +155,7 @@ int main()
para->setTimestepOut(timeStepOut);
para->setTimestepEnd(timeStepEnd);
- para->setMainKernel("CumulantK17");
+ para->setMainKernel(vf::CollisionKernel::Compressible::CumulantK17);
//////////////////////////////////////////////////////////////////////////
// set boundary conditions
diff --git a/apps/gpu/LBM/MusselOyster/MusselOyster.cpp b/apps/gpu/LBM/MusselOyster/MusselOyster.cpp
index 1004e2e4222af60a6c700799a01710e163d1a116..2e43e20f33c061c3d25da0ea2ff53e2351ad3cad 100644
--- a/apps/gpu/LBM/MusselOyster/MusselOyster.cpp
+++ b/apps/gpu/LBM/MusselOyster/MusselOyster.cpp
@@ -49,6 +49,7 @@
#include "VirtualFluids_GPU/Output/FileWriter.h"
#include "VirtualFluids_GPU/Parameter/Parameter.h"
#include "VirtualFluids_GPU/Factories/BoundaryConditionFactory.h"
+#include "VirtualFluids_GPU/Kernel/Utilities/KernelTypes.h"
//////////////////////////////////////////////////////////////////////////
@@ -160,8 +161,7 @@ void multipleLevel(std::filesystem::path &configPath)
std::cout << "Write result files to " << para->getFName() << std::endl;
para->setUseStreams(useStreams);
- // para->setMainKernel("CumulantK17CompChim");
- para->setMainKernel("CumulantK17CompChimStream");
+ para->setMainKernel(vf::CollisionKernel::Compressible::CumulantK17);
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/apps/gpu/LBM/SphereScaling/SphereScaling.cpp b/apps/gpu/LBM/SphereScaling/SphereScaling.cpp
index 43103a2f02522f7b3898027f9b807fdb84dd9d5c..c632e9649f7461e8af4a6e9e73d740406283edbf 100755
--- a/apps/gpu/LBM/SphereScaling/SphereScaling.cpp
+++ b/apps/gpu/LBM/SphereScaling/SphereScaling.cpp
@@ -55,6 +55,7 @@
#include "VirtualFluids_GPU/PreProcessor/PreProcessorFactory/PreProcessorFactoryImp.h"
#include "VirtualFluids_GPU/Factories/BoundaryConditionFactory.h"
#include "VirtualFluids_GPU/Factories/GridScalingFactory.h"
+#include "VirtualFluids_GPU/Kernel/Utilities/KernelTypes.h"
#include "VirtualFluids_GPU/GPU/CudaMemoryManager.h"
@@ -134,7 +135,7 @@ void multipleLevel(std::filesystem::path& configPath)
else
para->setMaxLevel(1);
- para->setMainKernel("CumulantK17");
+ para->setMainKernel(vf::CollisionKernel::Compressible::CumulantK17);
scalingFactory.setScalingFactory(GridScalingFactory::GridScaling::ScaleCompressible);
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernel.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernel.h
index 50b4460d774010ea7d7b98cfa6fa505cdfeb88c2..d83901a0f7d6a7df8120673a4b14371a6e935aef 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernel.h
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernel.h
@@ -5,7 +5,6 @@
#include "LBM/LB.h"
-#include "Kernel/Utilities/KernelGroup.h"
#include "PreProcessor/PreProcessorType.h"
#include "Parameter/CudaStreamManager.h"
@@ -20,6 +19,5 @@ public:
virtual bool checkParameter() = 0;
virtual std::vector<PreProcessorType> getPreProcessorTypes() = 0;
- virtual KernelGroup getKernelGroup() = 0;
};
#endif
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.cpp b/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.cpp
index 9bd3945aa81147d03be2b1eac3ddec7c24d71532..328cf8db260bc0092cb2081998961d1e9fb17233 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.cpp
+++ b/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.cpp
@@ -19,11 +19,6 @@ std::vector<PreProcessorType> KernelImp::getPreProcessorTypes()
return myPreProcessorTypes;
}
-KernelGroup KernelImp::getKernelGroup()
-{
- return myKernelGroup;
-}
-
void KernelImp::setCheckParameterStrategy(std::shared_ptr<CheckParameterStrategy> strategy)
{
this->checkStrategy = strategy;
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h b/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h
index a96c2c123472ca33f635273e06a5bf36a745654d..84e5f3f6ac08b92ccd92fbf142cceb3245de51d5 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h
+++ b/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h
@@ -20,7 +20,6 @@ public:
bool checkParameter();
std::vector<PreProcessorType> getPreProcessorTypes();
- KernelGroup getKernelGroup();
void setCheckParameterStrategy(std::shared_ptr<CheckParameterStrategy> strategy);
bool getKernelUsesFluidNodeIndices();
@@ -33,8 +32,6 @@ protected:
std::shared_ptr<CheckParameterStrategy> checkStrategy;
int level;
std::vector<PreProcessorType> myPreProcessorTypes;
- KernelGroup myKernelGroup;
-
vf::cuda::CudaGrid cudaGrid;
bool kernelUsesFluidNodeIndices = false;
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Compressible/Mod27/ADComp27/ADComp27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Compressible/Mod27/ADComp27/ADComp27.cu
index d4d6307f688da4c8fa37c54fb4958681d5ec4941..dd30516ac4229908a418d932177c1b63d8f5d685 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Compressible/Mod27/ADComp27/ADComp27.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Compressible/Mod27/ADComp27/ADComp27.cu
@@ -33,7 +33,6 @@ ADComp27::ADComp27(std::shared_ptr<Parameter> para, int level)
myPreProcessorTypes.push_back(InitCompAD27);
- myKernelGroup = ADKernel27;
}
ADComp27::ADComp27()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Compressible/Mod7/ADComp7/ADComp7.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Compressible/Mod7/ADComp7/ADComp7.cu
index 3ee06a1e9ea77c8443d94f44ea54d11ffe7304ac..d218489c754edc89f99277670f09536962ce62b2 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Compressible/Mod7/ADComp7/ADComp7.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Compressible/Mod7/ADComp7/ADComp7.cu
@@ -33,7 +33,6 @@ ADComp7::ADComp7(std::shared_ptr<Parameter> para, int level)
myPreProcessorTypes.push_back(InitCompAD7);
- myKernelGroup = ADKernel7;
}
ADComp7::ADComp7()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Incompressible/Mod27/ADIncomp27/ADIncomp27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Incompressible/Mod27/ADIncomp27/ADIncomp27.cu
index f2a9feaa998b628fb782844d1a7d946317e5af5f..150245a312509d50b77cf86fec18fdb063dbcc2c 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Incompressible/Mod27/ADIncomp27/ADIncomp27.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Incompressible/Mod27/ADIncomp27/ADIncomp27.cu
@@ -33,7 +33,6 @@ ADIncomp27::ADIncomp27(std::shared_ptr<Parameter> para, int level)
myPreProcessorTypes.push_back(InitIncompAD27);
- myKernelGroup = ADKernel27;
}
ADIncomp27::ADIncomp27()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Incompressible/Mod7/ADIncomp7/ADIncomp7.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Incompressible/Mod7/ADIncomp7/ADIncomp7.cu
index d0c6a6a24ab4d0ebebee9324bdafa1f9e3db51b9..71adc96eef733084e01fa963f6d0fad66a2e1062 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Incompressible/Mod7/ADIncomp7/ADIncomp7.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Incompressible/Mod7/ADIncomp7/ADIncomp7.cu
@@ -33,7 +33,6 @@ ADIncomp7::ADIncomp7(std::shared_ptr<Parameter> para, int level)
myPreProcessorTypes.push_back(InitIncompAD7);
- myKernelGroup = ADKernel7;
}
ADIncomp7::ADIncomp7()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGK/BGKCompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGK/BGKCompSP27.cu
index 8c99f3b030984aef6215d5479be4b321145ee54f..4aef26b7dd31435b2dadceb78ac1e0b7ebedf029 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGK/BGKCompSP27.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGK/BGKCompSP27.cu
@@ -31,7 +31,7 @@ BGKCompSP27::BGKCompSP27(std::shared_ptr<Parameter> para, int level)
this->level = level;
myPreProcessorTypes.push_back(InitCompSP27);
- myKernelGroup = BasicKernel;
+
}
BGKCompSP27::BGKCompSP27()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKPlus/BGKPlusCompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKPlus/BGKPlusCompSP27.cu
index a4b136d1c21b1e4c68432eef5e21ff8c968bdfec..00aaf3c27f16a5d53e7aee225214f05bd62a541a 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKPlus/BGKPlusCompSP27.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKPlus/BGKPlusCompSP27.cu
@@ -32,7 +32,7 @@ BGKPlusCompSP27::BGKPlusCompSP27(std::shared_ptr<Parameter> para, int level)
myPreProcessorTypes.push_back(InitCompSP27);
- myKernelGroup = BasicKernel;
+
}
BGKPlusCompSP27::BGKPlusCompSP27()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKUnified/BGKUnified.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKUnified/BGKUnified.cu
index 1107d343801f8ac3626b03a93ca92415217732ac..0a5ac6cf7a1b6564a61d0150b187b10b584222b8 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKUnified/BGKUnified.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKUnified/BGKUnified.cu
@@ -24,7 +24,7 @@ BGKUnified::BGKUnified(std::shared_ptr<Parameter> para, int level)
myPreProcessorTypes.push_back(InitCompSP27);
- myKernelGroup = BasicKernel;
+
this->cudaGrid = cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
}
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cascade/CascadeCompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cascade/CascadeCompSP27.cu
index dcfda06db462fd83120751a32a40365445d659ba..664b46fcebd277b0c93300d86b2171edf4f91b2a 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cascade/CascadeCompSP27.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cascade/CascadeCompSP27.cu
@@ -32,7 +32,7 @@ CascadeCompSP27::CascadeCompSP27(std::shared_ptr<Parameter> para, int level)
myPreProcessorTypes.push_back(InitCompSP27);
- myKernelGroup = BasicKernel;
+
}
CascadeCompSP27::CascadeCompSP27()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cumulant/CumulantCompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cumulant/CumulantCompSP27.cu
index 7817c398285dda131401bd14c3ccdd8c119c5680..218623b7c51099717f6aaa6f375a82516e0c0dae 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cumulant/CumulantCompSP27.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cumulant/CumulantCompSP27.cu
@@ -33,7 +33,7 @@ CumulantCompSP27::CumulantCompSP27(std::shared_ptr<Parameter> para, int level)
myPreProcessorTypes.push_back(InitCompSP27);
- myKernelGroup = BasicKernel;
+
}
CumulantCompSP27::CumulantCompSP27()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27.cu
index 1518dcc209de1edf8a88dae72c1f10c3d4666610..c8aad41b87ef39514f6cf5abc8b8bff42a869346 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27.cu
@@ -35,5 +35,5 @@ CumulantAll4CompSP27::CumulantAll4CompSP27(std::shared_ptr<Parameter> para, int
myPreProcessorTypes.push_back(InitCompSP27);
- myKernelGroup = BasicKernel;
+
}
\ No newline at end of file
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK15/CumulantK15Comp.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK15/CumulantK15Comp.cu
index 5a480e5d9c97126e491655b4bbe2aeefef3e7161..09a3aa1cdb1a3cf9c01002a9d335c5a907f94917 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK15/CumulantK15Comp.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK15/CumulantK15Comp.cu
@@ -50,5 +50,5 @@ CumulantK15Comp::CumulantK15Comp(std::shared_ptr<Parameter> para, int level)
myPreProcessorTypes.push_back(InitCompSP27);
- myKernelGroup = BasicKernel;
+
}
\ No newline at end of file
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK15Bulk/CumulantK15BulkComp.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK15Bulk/CumulantK15BulkComp.cu
index 51876f30b8c8e37d8cb3355edde5dcf2b04675d0..f0e29a9740438bc78d39574e1046d937cd7b86ce 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK15Bulk/CumulantK15BulkComp.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK15Bulk/CumulantK15BulkComp.cu
@@ -34,7 +34,7 @@ CumulantK15BulkComp::CumulantK15BulkComp(std::shared_ptr<Parameter> para, int le
myPreProcessorTypes.push_back(InitCompSP27);
- myKernelGroup = BasicKernel;
+
}
CumulantK15BulkComp::CumulantK15BulkComp()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK15Sponge/CumulantK15SpongeComp.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK15Sponge/CumulantK15SpongeComp.cu
index 613464125bafc572fe7951b8c372e3455ea5b21d..69f84b0671c11fad8ae15676230c491ee815153d 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK15Sponge/CumulantK15SpongeComp.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK15Sponge/CumulantK15SpongeComp.cu
@@ -35,7 +35,7 @@ CumulantK15SpongeComp::CumulantK15SpongeComp(std::shared_ptr<Parameter> para, in
myPreProcessorTypes.push_back(InitCompSP27);
- myKernelGroup = BasicKernel;
+
}
CumulantK15SpongeComp::CumulantK15SpongeComp()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK15Unified/CumulantK15Unified.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK15Unified/CumulantK15Unified.cu
index 24b0bbc6f43a63093da6b6dcb3ce401b8a614f75..c95289f15fe13decbbe173e17f5d4255b8ef80b5 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK15Unified/CumulantK15Unified.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK15Unified/CumulantK15Unified.cu
@@ -23,7 +23,7 @@ CumulantK15Unified::CumulantK15Unified(std::shared_ptr<Parameter> para, int leve
myPreProcessorTypes.push_back(InitCompSP27);
- myKernelGroup = BasicKernel;
+
this->cudaGrid = cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
}
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17/CumulantK17.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17/CumulantK17.cu
index ea3442fecca63fdcb45878d742a547ce492ab5c8..b31e4964b609bcee1c3015dcf950b540977f8333 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17/CumulantK17.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17/CumulantK17.cu
@@ -94,7 +94,9 @@ void CumulantK17<turbulenceModel>::runOnIndices( const unsigned int *indices, un
para->getParD(level)->isEvenTimestep,
indices,
size_indices);
- break; case CollisionTemplate::ApplyBodyForce:
+ break;
+
+ case CollisionTemplate::ApplyBodyForce:
LB_Kernel_CumulantK17 < turbulenceModel, false, true > <<< cudaGrid.grid, cudaGrid.threads, 0, stream >>>( para->getParD(level)->omega,
para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ,
para->getParD(level)->distributions.f[0],
@@ -110,7 +112,8 @@ void CumulantK17<turbulenceModel>::runOnIndices( const unsigned int *indices, un
para->getParD(level)->isEvenTimestep,
indices,
size_indices);
- break; default:
+ break;
+ default:
throw std::runtime_error("Invalid CollisionTemplate in CumulantK17::runOnIndices()");
break;
}
@@ -126,7 +129,7 @@ CumulantK17<turbulenceModel>::CumulantK17(std::shared_ptr<Parameter> para, int l
myPreProcessorTypes.push_back(InitCompSP27);
- myKernelGroup = BasicKernel;
+
this->cudaGrid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
this->kernelUsesFluidNodeIndices = true;
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Bulk/CumulantK17BulkComp.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Bulk/CumulantK17BulkComp.cu
index b9e25494490507bde5a6aa7d6dd588ac1a1f6c87..13b54723780fa16374b332c731fc35c5664d75b6 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Bulk/CumulantK17BulkComp.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Bulk/CumulantK17BulkComp.cu
@@ -39,7 +39,7 @@ CumulantK17BulkComp::CumulantK17BulkComp(std::shared_ptr<Parameter> para, int le
myPreProcessorTypes.push_back(InitCompSP27);
- myKernelGroup = BasicKernel;
+
}
CumulantK17BulkComp::CumulantK17BulkComp()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Unified/CumulantK17Unified.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Unified/CumulantK17Unified.cu
index 295804887f9c451120d463c7fcdd968bd2f24d12..59c405ae6e3bb46f608454ddb3a11bb0baac134f 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Unified/CumulantK17Unified.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Unified/CumulantK17Unified.cu
@@ -22,7 +22,7 @@ CumulantK17Unified::CumulantK17Unified(std::shared_ptr<Parameter> para, int leve
myPreProcessorTypes.push_back(InitCompSP27);
- myKernelGroup = BasicKernel;
+
this->cudaGrid = cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
}
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chim/CumulantK17CompChim.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chim/CumulantK17CompChim.cu
index bc058881e2a013effa417a149cf7a17bce646c6f..466b9f85999257196e860e84919ca6ccce6946b7 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chim/CumulantK17CompChim.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chim/CumulantK17CompChim.cu
@@ -33,6 +33,6 @@ void CumulantK17CompChim::run()
CumulantK17CompChim::CumulantK17CompChim(std::shared_ptr<Parameter> para, int level): KernelImp(para, level)
{
myPreProcessorTypes.push_back(InitCompSP27);
- myKernelGroup = BasicKernel;
+
this->cudaGrid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
}
\ No newline at end of file
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp.cu
index 2e0af0bdb85d3f008768f9f430e8b4e5d9719b0f..15d3509e735faa08b97d0876600c30876829c35f 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp.cu
@@ -37,5 +37,5 @@ CumulantK18Comp::CumulantK18Comp(std::shared_ptr<Parameter> para, int level)
myPreProcessorTypes.push_back(InitCompSP27);
myPreProcessorTypes.push_back(InitF3);
- myKernelGroup = F3Kernel;
+
}
\ No newline at end of file
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK20/CumulantK20Comp.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK20/CumulantK20Comp.cu
index d0d81eaac711d4d80284b66a1040e0e8404f5d4d..8181cdb690b5813c368eddadb9cda58a7d749302 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK20/CumulantK20Comp.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK20/CumulantK20Comp.cu
@@ -37,5 +37,5 @@ CumulantK20Comp::CumulantK20Comp(std::shared_ptr<Parameter> para, int level)
myPreProcessorTypes.push_back(InitCompSP27);
myPreProcessorTypes.push_back(InitF3);
- myKernelGroup = F3Kernel;
+
}
\ No newline at end of file
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/MRT/MRTCompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/MRT/MRTCompSP27.cu
index b576333f50304f5628e073d2eee16cf5b82c9d34..6e11bd97a2e76cca3983a83f785a2435d40f594b 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/MRT/MRTCompSP27.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/MRT/MRTCompSP27.cu
@@ -32,7 +32,7 @@ MRTCompSP27::MRTCompSP27(std::shared_ptr<Parameter> para, int level)
myPreProcessorTypes.push_back(InitCompSP27);
- myKernelGroup = BasicKernel;
+
}
MRTCompSP27::MRTCompSP27()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGK/BGKIncompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGK/BGKIncompSP27.cu
index 3fb9be28654f83a7a98bb7d6b3a8a46e9170e7a8..39bd1f3491d0d70e4734d04ef8a2d6e38cdc6448 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGK/BGKIncompSP27.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGK/BGKIncompSP27.cu
@@ -32,7 +32,7 @@ BGKIncompSP27::BGKIncompSP27(std::shared_ptr<Parameter> para, int level)
myPreProcessorTypes.push_back(InitSP27);
- myKernelGroup = BasicKernel;
+
}
BGKIncompSP27::BGKIncompSP27()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGKPlus/BGKPlusIncompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGKPlus/BGKPlusIncompSP27.cu
index f274f576a14fc193bcabd44d2c9078a2c98055bc..84a55b89d68f9a1e18c5114f8088a7dee24a4cd1 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGKPlus/BGKPlusIncompSP27.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGKPlus/BGKPlusIncompSP27.cu
@@ -32,7 +32,7 @@ BGKPlusIncompSP27::BGKPlusIncompSP27(std::shared_ptr<Parameter> para, int level)
myPreProcessorTypes.push_back(InitSP27);
- myKernelGroup = BasicKernel;
+
}
BGKPlusIncompSP27::BGKPlusIncompSP27()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cascade/CascadeIncompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cascade/CascadeIncompSP27.cu
index 3a6760b619d2ca1a7eb19771478eb9e5989ead0c..b060137f2d505886ee02a4b72e372ce8b4d48a78 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cascade/CascadeIncompSP27.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cascade/CascadeIncompSP27.cu
@@ -32,7 +32,7 @@ CascadeIncompSP27::CascadeIncompSP27(std::shared_ptr<Parameter> para, int level)
myPreProcessorTypes.push_back(InitSP27);
- myKernelGroup = BasicKernel;
+
}
CascadeIncompSP27::CascadeIncompSP27()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cumulant1hSP27/Cumulant1hIncompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cumulant1hSP27/Cumulant1hIncompSP27.cu
index 44beb8507d5664f01283130dd3087a788e4491ed..2cade430786b17567c47264f0638dba259b3192d 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cumulant1hSP27/Cumulant1hIncompSP27.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cumulant1hSP27/Cumulant1hIncompSP27.cu
@@ -37,7 +37,7 @@ Cumulant1hIncompSP27::Cumulant1hIncompSP27(std::shared_ptr<Parameter> para, int
myPreProcessorTypes.push_back(InitSP27);
- myKernelGroup = BasicKernel;
+
}
Cumulant1hIncompSP27::Cumulant1hIncompSP27()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantIsoSP27/CumulantIsoIncompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantIsoSP27/CumulantIsoIncompSP27.cu
index 3a740bef6d7fbaa2883b3d36930d49bf9bf0bb3e..840067da7f34a4415b1b14458ae0fc8d316e366d 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantIsoSP27/CumulantIsoIncompSP27.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantIsoSP27/CumulantIsoIncompSP27.cu
@@ -35,7 +35,7 @@ CumulantIsoIncompSP27::CumulantIsoIncompSP27(std::shared_ptr<Parameter> para, in
myPreProcessorTypes.push_back(InitSP27);
- myKernelGroup = BasicKernel;
+
}
CumulantIsoIncompSP27::CumulantIsoIncompSP27()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantK15/CumulantK15Incomp.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantK15/CumulantK15Incomp.cu
index 7ae17b97170b4d8474acd6777f7c27411a962681..c597924193d859a35dddaa7b37a56e21d265ceba 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantK15/CumulantK15Incomp.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantK15/CumulantK15Incomp.cu
@@ -32,7 +32,7 @@ CumulantK15Incomp::CumulantK15Incomp(std::shared_ptr<Parameter> para, int level)
myPreProcessorTypes.push_back(InitSP27);
- myKernelGroup = BasicKernel;
+
}
CumulantK15Incomp::CumulantK15Incomp()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/MRT/MRTIncompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/MRT/MRTIncompSP27.cu
index 7645703e0d40176b136762d6b48633f4a9c0d950..daa90091fe092a98741d0764e2327f3ce4c9d2bc 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/MRT/MRTIncompSP27.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/MRT/MRTIncompSP27.cu
@@ -32,7 +32,7 @@ MRTIncompSP27::MRTIncompSP27(std::shared_ptr<Parameter> para, int level)
myPreProcessorTypes.push_back(InitSP27);
- myKernelGroup = BasicKernel;
+
}
MRTIncompSP27::MRTIncompSP27()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/PorousMediaKernels/FluidFlow/Compressible/CumulantOne/PMCumulantOneCompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/PorousMediaKernels/FluidFlow/Compressible/CumulantOne/PMCumulantOneCompSP27.cu
index 43724f9165e2bb8dca1705ae0053612df92413ec..a8c1af64ebd4641a755bf9fed7e9fafa18e9cad7 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/PorousMediaKernels/FluidFlow/Compressible/CumulantOne/PMCumulantOneCompSP27.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/PorousMediaKernels/FluidFlow/Compressible/CumulantOne/PMCumulantOneCompSP27.cu
@@ -57,7 +57,7 @@ PMCumulantOneCompSP27::PMCumulantOneCompSP27(std::shared_ptr<Parameter> para, st
myPreProcessorTypes.push_back(InitSP27);
- myKernelGroup = BasicKernel;
+
}
PMCumulantOneCompSP27::PMCumulantOneCompSP27()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/WaleKernels/FluidFlow/Compressible/CumulantK15/WaleCumulantK15Comp.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/WaleKernels/FluidFlow/Compressible/CumulantK15/WaleCumulantK15Comp.cu
index 2b8a7d61e8966e2ed00022986311ae68ac0ca6d6..cfcc544aac2172cef2f4d58600931db8ccfa0189 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/WaleKernels/FluidFlow/Compressible/CumulantK15/WaleCumulantK15Comp.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/WaleKernels/FluidFlow/Compressible/CumulantK15/WaleCumulantK15Comp.cu
@@ -55,7 +55,7 @@ WaleCumulantK15Comp::WaleCumulantK15Comp(std::shared_ptr<Parameter> para, int le
myPreProcessorTypes.push_back(InitCompSP27);
- myKernelGroup = BasicWaleKernel;
+
}
WaleCumulantK15Comp::WaleCumulantK15Comp()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/WaleKernels/FluidFlow/Compressible/CumulantK15BySoniMalav/WaleBySoniMalavCumulantK15Comp.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/WaleKernels/FluidFlow/Compressible/CumulantK15BySoniMalav/WaleBySoniMalavCumulantK15Comp.cu
index 49ee20b44f37b01cd9bc837024a47c1428c00a18..05e257a52b38e2c31badcb1fb739de3ab0239f6e 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/WaleKernels/FluidFlow/Compressible/CumulantK15BySoniMalav/WaleBySoniMalavCumulantK15Comp.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/WaleKernels/FluidFlow/Compressible/CumulantK15BySoniMalav/WaleBySoniMalavCumulantK15Comp.cu
@@ -58,7 +58,7 @@ WaleBySoniMalavCumulantK15Comp::WaleBySoniMalavCumulantK15Comp(std::shared_ptr<P
myPreProcessorTypes.push_back(InitCompSP27);
- myKernelGroup = BasicWaleKernel;
+
}
WaleBySoniMalavCumulantK15Comp::WaleBySoniMalavCumulantK15Comp()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/WaleKernels/FluidFlow/Compressible/CumulantK17/WaleCumulantK17Comp.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/WaleKernels/FluidFlow/Compressible/CumulantK17/WaleCumulantK17Comp.cu
index c9c16e2d2d2259656248948f3f10977c8f18fd24..b7f4038c6b67cc4d1cf521bc7a904801650d1e8d 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/WaleKernels/FluidFlow/Compressible/CumulantK17/WaleCumulantK17Comp.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/WaleKernels/FluidFlow/Compressible/CumulantK17/WaleCumulantK17Comp.cu
@@ -60,7 +60,7 @@ WaleCumulantK17Comp::WaleCumulantK17Comp(std::shared_ptr<Parameter> para, int le
myPreProcessorTypes.push_back(InitCompSP27);
- myKernelGroup = BasicWaleKernel;
+
}
WaleCumulantK17Comp::WaleCumulantK17Comp()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/WaleKernels/FluidFlow/Compressible/CumulantK17Debug/WaleCumulantK17DebugComp.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/WaleKernels/FluidFlow/Compressible/CumulantK17Debug/WaleCumulantK17DebugComp.cu
index b3cdd494c02c6649d60818b6b264b8db8b79d426..5fe0284e675785691e51a58e7e0869ba4164ad5f 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/WaleKernels/FluidFlow/Compressible/CumulantK17Debug/WaleCumulantK17DebugComp.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/WaleKernels/FluidFlow/Compressible/CumulantK17Debug/WaleCumulantK17DebugComp.cu
@@ -70,7 +70,7 @@ WaleCumulantK17DebugComp::WaleCumulantK17DebugComp(std::shared_ptr<Parameter> pa
myPreProcessorTypes.push_back(InitCompSP27);
- myKernelGroup = BasicWaleKernel;
+
}
WaleCumulantK17DebugComp::WaleCumulantK17DebugComp()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Utilities/KernelFactory/KernelFactoryImp.cpp b/src/gpu/VirtualFluids_GPU/Kernel/Utilities/KernelFactory/KernelFactoryImp.cpp
index 5a2d8c9a426e5cb23ca75f91aaf6fbff75cba72b..27c061ce99f71c349ac1c479efb5f9b780cff3b2 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Utilities/KernelFactory/KernelFactoryImp.cpp
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Utilities/KernelFactory/KernelFactoryImp.cpp
@@ -2,6 +2,8 @@
#include "Parameter/Parameter.h"
+#include "Kernel/Utilities/KernelTypes.h"
+
//LBM kernel (compressible)
#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGK/BGKCompSP27.h"
#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKUnified/BGKUnified.h"
@@ -57,151 +59,153 @@
#include "Kernel/Kernels/PorousMediaKernels/FluidFlow/Compressible/PMFluidFlowCompStrategy.h"
#include "Kernel/Kernels/WaleKernels/FluidFlow/Compressible/WaleFluidFlowCompStrategy.h"
+using namespace vf;
+
std::vector<std::shared_ptr<Kernel>> KernelFactoryImp::makeKernels(std::shared_ptr<Parameter> para)
{
- std::vector< std::shared_ptr< Kernel>> kernels;
- for (int level = 0; level <= para->getMaxLevel(); level++)
- kernels.push_back(makeKernel(para, para->getMainKernel(), level));
-
- if (para->getMaxLevel() > 0)
- if (para->getMultiKernelOn())
- for (std::size_t i = 0; i < para->getMultiKernelLevel().size(); i++)
- setKernelAtLevel(kernels, para, para->getMultiKernel().at(i), para->getMultiKernelLevel().at(i));
- return kernels;
+ std::vector< std::shared_ptr< Kernel>> kernels;
+ for (int level = 0; level <= para->getMaxLevel(); level++)
+ kernels.push_back(makeKernel(para, para->getMainKernel(), level));
+
+ if (para->getMaxLevel() > 0)
+ if (para->getMultiKernelOn())
+ for (std::size_t i = 0; i < para->getMultiKernelLevel().size(); i++)
+ setKernelAtLevel(kernels, para, para->getMultiKernel().at(i), para->getMultiKernelLevel().at(i));
+ return kernels;
}
std::vector<std::shared_ptr<ADKernel>> KernelFactoryImp::makeAdvDifKernels(std::shared_ptr<Parameter> para)
{
- std::vector< std::shared_ptr< ADKernel>> aDKernels;
- for (int level = 0; level <= para->getMaxLevel(); level++)
- aDKernels.push_back(makeAdvDifKernel(para, para->getADKernel(), level));
- return aDKernels;
+ std::vector< std::shared_ptr< ADKernel>> aDKernels;
+ for (int level = 0; level <= para->getMaxLevel(); level++)
+ aDKernels.push_back(makeAdvDifKernel(para, para->getADKernel(), level));
+ return aDKernels;
}
void KernelFactoryImp::setPorousMedia(std::vector<std::shared_ptr<PorousMedia>> pm)
{
- this->pm = pm;
+ this->pm = pm;
}
void KernelFactoryImp::setKernelAtLevel(std::vector<std::shared_ptr<Kernel>> kernels, std::shared_ptr<Parameter> para, std::string kernel, int level)
{
- kernels.at(level) = makeKernel(para, kernel, level);
+ kernels.at(level) = makeKernel(para, kernel, level);
}
std::shared_ptr<Kernel> KernelFactoryImp::makeKernel(std::shared_ptr<Parameter> para, std::string kernel, int level)
{
printf("Instantiating Kernel: %s\n", kernel.c_str());
- std::shared_ptr<KernelImp> newKernel;
- std::shared_ptr<CheckParameterStrategy> checkStrategy;
+ std::shared_ptr<KernelImp> newKernel;
+ std::shared_ptr<CheckParameterStrategy> checkStrategy;
- if (kernel == "BGKCompSP27") {
- newKernel = BGKCompSP27::getNewInstance(para, level); // compressible
- checkStrategy = FluidFlowCompStrategy::getInstance(); // ||
- } else if (kernel == "BGKUnified") { // \/
+ if (kernel == CollisionKernel::Compressible::BGK) {
+ newKernel = BGKCompSP27::getNewInstance(para, level); // compressible
+ checkStrategy = FluidFlowCompStrategy::getInstance(); // ||
+ } else if (kernel == CollisionKernel::Compressible::BGKUnified) { // \/
newKernel = std::make_shared<vf::gpu::BGKUnified>(para, level);
checkStrategy = FluidFlowCompStrategy::getInstance();
- } else if (kernel == "BGKPlusCompSP27") {
+ } else if (kernel == CollisionKernel::Compressible::BGKPlus) {
newKernel = BGKPlusCompSP27::getNewInstance(para, level);
checkStrategy = FluidFlowCompStrategy::getInstance();
- } else if (kernel == "MRTCompSP27") {
+ } else if (kernel == CollisionKernel::Compressible::MRT) {
newKernel = MRTCompSP27::getNewInstance(para, level);
checkStrategy = FluidFlowCompStrategy::getInstance();
- } else if (kernel == "CascadeCompSP27") {
+ } else if (kernel == CollisionKernel::Compressible::Cascade) {
newKernel = CascadeCompSP27::getNewInstance(para, level);
checkStrategy = FluidFlowCompStrategy::getInstance();
- } else if (kernel == "CumulantCompSP27") {
+ } else if (kernel == CollisionKernel::Compressible::CumulantClassic) {
newKernel = CumulantCompSP27::getNewInstance(para, level);
checkStrategy = FluidFlowCompStrategy::getInstance();
- } else if (kernel == "CumulantK15Unified") {
+ } else if (kernel == CollisionKernel::Compressible::CumulantK15Unified) {
newKernel = std::make_shared<vf::gpu::CumulantK15Unified>(para, level);
checkStrategy = FluidFlowCompStrategy::getInstance();
- } else if (kernel == "CumulantK17Unified") {
+ } else if (kernel == CollisionKernel::Compressible::CumulantK17Unified) {
newKernel = std::make_shared<vf::gpu::CumulantK17Unified>(para, level);
checkStrategy = FluidFlowCompStrategy::getInstance();
- } else if (kernel == "CumulantK17BulkComp") {
+ } else if (kernel == CollisionKernel::Compressible::CumulantK17Bulk) {
newKernel = CumulantK17BulkComp::getNewInstance(para, level);
checkStrategy = FluidFlowCompStrategy::getInstance();
- } else if (kernel == "CumulantK17CompChim") {
+ } else if (kernel == CollisionKernel::Compressible::CumulantK17Chim) {
newKernel = CumulantK17CompChim::getNewInstance(para, level);
checkStrategy = FluidFlowCompStrategy::getInstance();
- } else if (kernel == "CumulantK17"){
- switch(para->getTurbulenceModel())
- {
+ } else if (kernel == CollisionKernel::Compressible::CumulantK17){
+ switch(para->getTurbulenceModel())
+ {
case TurbulenceModel::AMD:
- newKernel = CumulantK17<TurbulenceModel::AMD>::getNewInstance(para, level);
+ newKernel = CumulantK17<TurbulenceModel::AMD>::getNewInstance(para, level);
break;
case TurbulenceModel::Smagorinsky:
- newKernel = CumulantK17<TurbulenceModel::Smagorinsky>::getNewInstance(para, level);
+ newKernel = CumulantK17<TurbulenceModel::Smagorinsky>::getNewInstance(para, level);
break;
case TurbulenceModel::QR:
- newKernel = CumulantK17<TurbulenceModel::QR>::getNewInstance(para, level);
+ newKernel = CumulantK17<TurbulenceModel::QR>::getNewInstance(para, level);
break;
case TurbulenceModel::None:
- newKernel = CumulantK17<TurbulenceModel::None>::getNewInstance(para, level);
+ newKernel = CumulantK17<TurbulenceModel::None>::getNewInstance(para, level);
break;
default:
throw std::runtime_error("Unknown turbulence model!");
- break;
- }
- checkStrategy = FluidFlowCompStrategy::getInstance();
- } else if (kernel == "CumulantAll4CompSP27") {
+ break;
+ }
+ checkStrategy = FluidFlowCompStrategy::getInstance();
+ } else if (kernel == CollisionKernel::Compressible::CumulantAll4SP27) {
newKernel = CumulantAll4CompSP27::getNewInstance(para, level);
checkStrategy = FluidFlowCompStrategy::getInstance();
- } else if (kernel == "CumulantK18Comp") {
+ } else if (kernel == CollisionKernel::Compressible::CumulantK18) {
newKernel = CumulantK18Comp::getNewInstance(para, level);
checkStrategy = FluidFlowCompStrategy::getInstance();
- } else if (kernel == "CumulantK20Comp") {
+ } else if (kernel == CollisionKernel::Compressible::CumulantK20) {
newKernel = CumulantK20Comp::getNewInstance(para, level);
checkStrategy = FluidFlowCompStrategy::getInstance();
- } else if (kernel == "CumulantK15Comp") {
+ } else if (kernel == CollisionKernel::Compressible::CumulantK15) {
newKernel = CumulantK15Comp::getNewInstance(para, level);
checkStrategy = FluidFlowCompStrategy::getInstance();
- } else if (kernel == "CumulantK15BulkComp") {
+ } else if (kernel == CollisionKernel::Compressible::CumulantK15Bulk) {
newKernel = CumulantK15BulkComp::getNewInstance(para, level);
checkStrategy = FluidFlowCompStrategy::getInstance();
- } else if (kernel == "CumulantK15SpongeComp") { // /\ //
- newKernel = CumulantK15SpongeComp::getNewInstance(para, level); // ||
+ } else if (kernel == CollisionKernel::Compressible::CumulantK15Sponge) { // /\ //
+ newKernel = CumulantK15SpongeComp::getNewInstance(para, level); // ||
checkStrategy = FluidFlowCompStrategy::getInstance(); // compressible
- } //===============
- else if ( kernel == "BGKIncompSP27") { // incompressible
- newKernel = BGKIncompSP27::getNewInstance(para, level); // ||
+ } //===============
+ else if ( kernel == CollisionKernel::Incompressible::BGK) { // incompressible
+ newKernel = BGKIncompSP27::getNewInstance(para, level); // ||
checkStrategy = FluidFlowIncompStrategy::getInstance(); // \/
- } else if (kernel == "BGKPlusIncompSP27") {
+ } else if (kernel == CollisionKernel::Incompressible::BGKPlus) {
newKernel = BGKPlusIncompSP27::getNewInstance(para, level);
checkStrategy = FluidFlowIncompStrategy::getInstance();
- } else if (kernel == "MRTIncompSP27") {
+ } else if (kernel == CollisionKernel::Incompressible::MRT) {
newKernel = MRTIncompSP27::getNewInstance(para, level);
checkStrategy = FluidFlowIncompStrategy::getInstance();
- } else if (kernel == "CascadeIncompSP27") {
+ } else if (kernel == CollisionKernel::Incompressible::Cascade) {
newKernel = CascadeIncompSP27::getNewInstance(para, level);
checkStrategy = FluidFlowIncompStrategy::getInstance();
- } else if (kernel == "Cumulant1hIncompSP27") {
+ } else if (kernel == CollisionKernel::Incompressible::Cumulant1h) {
newKernel = Cumulant1hIncompSP27::getNewInstance(para, level);
checkStrategy = FluidFlowIncompStrategy::getInstance();
- } else if (kernel == "CumulantIsoIncompSP27") {
+ } else if (kernel == CollisionKernel::Incompressible::CumulantIsometric) {
newKernel = CumulantIsoIncompSP27::getNewInstance(para, level);
checkStrategy = FluidFlowIncompStrategy::getInstance();
- } else if (kernel == "CumulantK15Incomp") { // /\ //
- newKernel = CumulantK15Incomp::getNewInstance(para, level); // ||
- checkStrategy = FluidFlowIncompStrategy::getInstance(); // incompressible
- } //===============
- else if (kernel == "PMCumulantOneCompSP27") { // porous media
- newKernel = PMCumulantOneCompSP27::getNewInstance(para, pm, level); // ||
- checkStrategy = PMFluidFlowCompStrategy::getInstance(); // porous media
- } //===============
- else if (kernel == "WaleCumulantK17Comp") { // wale model
- newKernel = WaleCumulantK17Comp::getNewInstance(para, level); // ||
- checkStrategy = WaleFluidFlowCompStrategy::getInstance(); // \/
- } else if (kernel == "WaleCumulantK17DebugComp") {
+ } else if (kernel == CollisionKernel::Incompressible::CumulantK15) { // /\ //
+ newKernel = CumulantK15Incomp::getNewInstance(para, level); // ||
+ checkStrategy = FluidFlowIncompStrategy::getInstance(); // incompressible
+ } //===============
+ else if (kernel == CollisionKernel::PorousMedia::CumulantOne) { // porous media
+ newKernel = PMCumulantOneCompSP27::getNewInstance(para, pm, level); // ||
+ checkStrategy = PMFluidFlowCompStrategy::getInstance(); // porous media
+ } //===============
+ else if (kernel == CollisionKernel::Wale::CumulantK17) { // wale model
+ newKernel = WaleCumulantK17Comp::getNewInstance(para, level); // ||
+ checkStrategy = WaleFluidFlowCompStrategy::getInstance(); // \/
+ } else if (kernel == CollisionKernel::Wale::CumulantK17Debug) {
newKernel = WaleCumulantK17DebugComp::getNewInstance(para, level);
checkStrategy = WaleFluidFlowCompStrategy::getInstance();
- } else if (kernel == "WaleCumulantK15Comp") {
+ } else if (kernel == CollisionKernel::Wale::CumulantK15) {
newKernel = WaleCumulantK15Comp::getNewInstance(para, level);
checkStrategy = WaleFluidFlowCompStrategy::getInstance();
- } else if (kernel == "WaleBySoniMalavCumulantK15Comp") { // /\ //
- newKernel = WaleBySoniMalavCumulantK15Comp::getNewInstance(para, level);// ||
- checkStrategy = WaleFluidFlowCompStrategy::getInstance(); // wale model
- } //===============
+ } else if (kernel == CollisionKernel::Wale::CumulantK15SoniMalav) { // /\ //
+ newKernel = WaleBySoniMalavCumulantK15Comp::getNewInstance(para, level); // ||
+ checkStrategy = WaleFluidFlowCompStrategy::getInstance(); // wale model
+ } //===============
else {
throw std::runtime_error("KernelFactory does not know the KernelType.");
}
@@ -212,8 +216,8 @@ std::shared_ptr<Kernel> KernelFactoryImp::makeKernel(std::shared_ptr<Parameter>
std::shared_ptr<ADKernel> KernelFactoryImp::makeAdvDifKernel(std::shared_ptr<Parameter> para, std::string kernel, int level)
{
- std::shared_ptr<ADKernel> newKernel;
- std::shared_ptr<CheckParameterStrategy> checkStrategy;
+ std::shared_ptr<ADKernel> newKernel;
+ std::shared_ptr<CheckParameterStrategy> checkStrategy;
if (kernel == "ADComp27") {
newKernel = ADComp27::getNewInstance(para, level);
@@ -223,18 +227,18 @@ std::shared_ptr<ADKernel> KernelFactoryImp::makeAdvDifKernel(std::shared_ptr<Par
checkStrategy = ADMod7CompStrategy::getInstance();
} else if (kernel == "ADIncomp27") {
newKernel = ADIncomp27::getNewInstance(para, level);
- checkStrategy = ADMod7CompStrategy::getInstance();
+ checkStrategy = ADMod7IncompStrategy::getInstance();
} else if (kernel == "ADIncomp7") {
newKernel = ADIncomp7::getNewInstance(para, level);
- checkStrategy = ADMod7CompStrategy::getInstance();
+ checkStrategy = ADMod7IncompStrategy::getInstance();
} else {
throw std::runtime_error("KernelFactory does not know the KernelType.");
}
- if (newKernel) {
- newKernel->setCheckParameterStrategy(checkStrategy);
- return newKernel;
- }
- else
- throw std::runtime_error("KernelFactory does not know the KernelType.");
+ if (newKernel) {
+ newKernel->setCheckParameterStrategy(checkStrategy);
+ return newKernel;
+ }
+ else
+ throw std::runtime_error("KernelFactory does not know the KernelType.");
}
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Utilities/KernelGroup.h b/src/gpu/VirtualFluids_GPU/Kernel/Utilities/KernelGroup.h
deleted file mode 100644
index 0a6543ca0ac1d47bb6f8838d029769846c361868..0000000000000000000000000000000000000000
--- a/src/gpu/VirtualFluids_GPU/Kernel/Utilities/KernelGroup.h
+++ /dev/null
@@ -1,14 +0,0 @@
-#ifndef KERNEL_GROUP_H
-#define KERNEL_GROUP_H
-
-enum KernelGroup
-{
- BasicKernel,
- BasicWaleKernel,
- F3Kernel,
- F3WaleKernel,
- ADKernel7,
- ADKernel27
-};
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Utilities/KernelTypes.h b/src/gpu/VirtualFluids_GPU/Kernel/Utilities/KernelTypes.h
new file mode 100644
index 0000000000000000000000000000000000000000..f249c0bd595d21455b4338334763be4e08abeda9
--- /dev/null
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Utilities/KernelTypes.h
@@ -0,0 +1,51 @@
+#ifndef KERNEL_TYPES_H
+#define KERNEL_TYPES_H
+
+namespace vf::CollisionKernel::Compressible {
+ static const std::string BGK = "BGKCompSP27";
+ static const std::string BGKUnified = "BGKUnified";
+ static const std::string BGKPlus = "BGKPlusCompSP27";
+ static const std::string MRT = "MRTCompSP27";
+ static const std::string Cascade = "CascadeCompSP27";
+
+ static const std::string CumulantClassic = "CumulantCompSP27";
+
+ static const std::string CumulantK15Unified = "CumulantK15Unified";
+ static const std::string CumulantK17Unified = "CumulantK17Unified";
+
+ static const std::string CumulantK17Bulk = "CumulantK17BulkComp";
+ static const std::string CumulantK17Chim = "CumulantK17CompChim";
+ static const std::string CumulantK17 = "CumulantK17";
+
+ static const std::string CumulantAll4SP27 = "CumulantAll4CompSP27";
+ static const std::string CumulantK18 = "CumulantK18Comp";
+ static const std::string CumulantK20 = "CumulantK20Comp";
+
+ static const std::string CumulantK15 = "CumulantK15Comp";
+ static const std::string CumulantK15Bulk = "CumulantK15BulkComp";
+ static const std::string CumulantK15Sponge = "CumulantK15SpongeComp";
+}
+
+namespace vf::CollisionKernel::Incompressible {
+ static const std::string BGK = "BGKIncompSP27";
+ static const std::string BGKPlus = "BGKPlusIncompSP27";
+ static const std::string MRT = "MRTIncompSP27";
+ static const std::string Cascade = "CascadeIncompSP27";
+
+ static const std::string Cumulant1h = "Cumulant1hIncompSP27";
+ static const std::string CumulantIsometric = "CumulantIsoIncompSP27";
+ static const std::string CumulantK15 = "CumulantK15Incomp";
+}
+
+namespace vf::CollisionKernel::PorousMedia {
+ static const std::string CumulantOne = "CumulantOneCompSP27";
+}
+
+namespace vf::CollisionKernel::Wale {
+ static const std::string CumulantK17 = "WaleCumulantK17Comp";
+ static const std::string CumulantK17Debug = "WaleCumulantK17DebugComp";
+ static const std::string CumulantK15 = "WaleCumulantK15Comp";
+ static const std::string CumulantK15SoniMalav = "WaleBySoniMalavCumulantK15Comp";
+}
+
+#endif