From 8c55a511d1fe6273a61e289c38942c08cba2a2a2 Mon Sep 17 00:00:00 2001
From: Martin Schoenherr <m.schoenherr@tu-braunschweig.de>
Date: Tue, 14 Nov 2023 13:13:53 +0100
Subject: [PATCH] refactor all navier stokes inits
---
.../B15/B15CompressibleNavierStokesBGKplus.cu | 2 +-
.../B92/B92CompressibleNavierStokes.cu | 2 +-
.../K15/K15CompressibleNavierStokes.cu | 2 +-
.../K17/K17CompressibleNavierStokes.cu | 2 +-
.../B15IncompressibleNavierStokesBGKplus.cu | 2 +-
.../B92/B92IncompressibleNavierStokes.cu | 2 +-
.../K15/K15IncompressibleNavierStokes.cu | 2 +-
.../core/KernelManager/ADKernelManager.cpp | 4 +-
.../PreProcessorFactoryImp.cpp | 18 ++--
.../InitCompSP27/InitCompSP27.cu | 68 -------------
.../InitCompSP27/InitCompSP27.h | 23 -----
.../InitCompSP27/InitCompSP27_Device.cuh | 33 -------
.../PreProcessorStrategy/InitF3/InitF3.cu | 43 --------
.../PreProcessorStrategy/InitF3/InitF3.h | 23 -----
.../InitF3/InitF3_Device.cuh | 19 ----
.../InitK18K20NavierStokesCompressible.cu | 73 ++++++++++++++
.../InitK18K20NavierStokesCompressible.h | 53 ++++++++++
...tK18K20NavierStokesCompressible_Device.cu} | 32 +++++-
...tK18K20NavierStokesCompressible_Device.cuh | 50 ++++++++++
.../InitNavierStokesCompressible.cu | 98 +++++++++++++++++++
.../InitNavierStokesCompressible.h | 53 ++++++++++
.../InitNavierStokesCompressible_Device.cu} | 35 ++++++-
.../InitNavierStokesCompressible_Device.cuh | 63 ++++++++++++
.../InitNavierStokesIncompressible.cu | 73 ++++++++++++++
.../InitNavierStokesIncompressible.h | 53 ++++++++++
.../InitNavierStokesIncompressible_Device.cu} | 32 +++++-
.../InitNavierStokesIncompressible_Device.cuh | 50 ++++++++++
.../PreProcessorStrategy/InitSP27/InitSP27.cu | 43 --------
.../PreProcessorStrategy/InitSP27/InitSP27.h | 23 -----
.../InitSP27/InitSP27_Device.cuh | 19 ----
src/gpu/core/PreProcessor/PreProcessorType.h | 36 ++++++-
31 files changed, 712 insertions(+), 319 deletions(-)
delete mode 100644 src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27.cu
delete mode 100644 src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27.h
delete mode 100644 src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27_Device.cuh
delete mode 100644 src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3.cu
delete mode 100644 src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3.h
delete mode 100644 src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3_Device.cuh
create mode 100644 src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.cu
create mode 100644 src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.h
rename src/gpu/core/PreProcessor/PreProcessorStrategy/{InitF3/InitF3_Device.cu => InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible_Device.cu} (54%)
create mode 100644 src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible_Device.cuh
create mode 100644 src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible.cu
create mode 100644 src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible.h
rename src/gpu/core/PreProcessor/PreProcessorStrategy/{InitCompSP27/InitCompSP27_Device.cu => InitNavierStokesCompressible/InitNavierStokesCompressible_Device.cu} (91%)
create mode 100644 src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible_Device.cuh
create mode 100644 src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible.cu
create mode 100644 src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible.h
rename src/gpu/core/PreProcessor/PreProcessorStrategy/{InitSP27/InitSP27_Device.cu => InitNavierStokesIncompressible/InitNavierStokesIncompressible_Device.cu} (79%)
create mode 100644 src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible_Device.cuh
delete mode 100644 src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27.cu
delete mode 100644 src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27.h
delete mode 100644 src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27_Device.cuh
diff --git a/src/gpu/core/Kernel/Compressible/NavierStokes/B15/B15CompressibleNavierStokesBGKplus.cu b/src/gpu/core/Kernel/Compressible/NavierStokes/B15/B15CompressibleNavierStokesBGKplus.cu
index 6f6befbd2..6f51eee72 100644
--- a/src/gpu/core/Kernel/Compressible/NavierStokes/B15/B15CompressibleNavierStokesBGKplus.cu
+++ b/src/gpu/core/Kernel/Compressible/NavierStokes/B15/B15CompressibleNavierStokesBGKplus.cu
@@ -30,7 +30,7 @@ B15CompressibleNavierStokesBGKplus::B15CompressibleNavierStokesBGKplus(std::shar
this->para = para;
this->level = level;
- myPreProcessorTypes.push_back(InitCompSP27);
+ myPreProcessorTypes.push_back(InitNavierStokesCompressible);
}
diff --git a/src/gpu/core/Kernel/Compressible/NavierStokes/B92/B92CompressibleNavierStokes.cu b/src/gpu/core/Kernel/Compressible/NavierStokes/B92/B92CompressibleNavierStokes.cu
index 5f2d06945..1dba9343c 100644
--- a/src/gpu/core/Kernel/Compressible/NavierStokes/B92/B92CompressibleNavierStokes.cu
+++ b/src/gpu/core/Kernel/Compressible/NavierStokes/B92/B92CompressibleNavierStokes.cu
@@ -30,7 +30,7 @@ B92CompressibleNavierStokes::B92CompressibleNavierStokes(std::shared_ptr<Paramet
this->para = para;
this->level = level;
- myPreProcessorTypes.push_back(InitCompSP27);
+ myPreProcessorTypes.push_back(InitNavierStokesCompressible);
}
diff --git a/src/gpu/core/Kernel/Compressible/NavierStokes/K15/K15CompressibleNavierStokes.cu b/src/gpu/core/Kernel/Compressible/NavierStokes/K15/K15CompressibleNavierStokes.cu
index 2ba1af7a3..4cb597347 100644
--- a/src/gpu/core/Kernel/Compressible/NavierStokes/K15/K15CompressibleNavierStokes.cu
+++ b/src/gpu/core/Kernel/Compressible/NavierStokes/K15/K15CompressibleNavierStokes.cu
@@ -48,7 +48,7 @@ K15CompressibleNavierStokes::K15CompressibleNavierStokes(std::shared_ptr<Paramet
this->para = para;
this->level = level;
- myPreProcessorTypes.push_back(InitCompSP27);
+ myPreProcessorTypes.push_back(InitNavierStokesCompressible);
}
\ No newline at end of file
diff --git a/src/gpu/core/Kernel/Compressible/NavierStokes/K17/K17CompressibleNavierStokes.cu b/src/gpu/core/Kernel/Compressible/NavierStokes/K17/K17CompressibleNavierStokes.cu
index 60d4a0314..3fd0f5786 100644
--- a/src/gpu/core/Kernel/Compressible/NavierStokes/K17/K17CompressibleNavierStokes.cu
+++ b/src/gpu/core/Kernel/Compressible/NavierStokes/K17/K17CompressibleNavierStokes.cu
@@ -116,7 +116,7 @@ void K17CompressibleNavierStokes<turbulenceModel>::runOnIndices(const unsigned i
template <vf::lbm::TurbulenceModel turbulenceModel>
K17CompressibleNavierStokes<turbulenceModel>::K17CompressibleNavierStokes(std::shared_ptr<Parameter> para, int level) : KernelImp(para, level)
{
- myPreProcessorTypes.push_back(InitCompSP27);
+ myPreProcessorTypes.push_back(InitNavierStokesCompressible);
this->cudaGrid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
this->kernelUsesFluidNodeIndices = true;
diff --git a/src/gpu/core/Kernel/Incompressible/NavierStokes/B15/B15IncompressibleNavierStokesBGKplus.cu b/src/gpu/core/Kernel/Incompressible/NavierStokes/B15/B15IncompressibleNavierStokesBGKplus.cu
index 251cf7544..cc0b6d43e 100644
--- a/src/gpu/core/Kernel/Incompressible/NavierStokes/B15/B15IncompressibleNavierStokesBGKplus.cu
+++ b/src/gpu/core/Kernel/Incompressible/NavierStokes/B15/B15IncompressibleNavierStokesBGKplus.cu
@@ -30,7 +30,7 @@ B15IncompressibleNavierStokesBGKplus::B15IncompressibleNavierStokesBGKplus(std::
this->para = para;
this->level = level;
- myPreProcessorTypes.push_back(InitSP27);
+ myPreProcessorTypes.push_back(InitNavierStokesIncompressible);
}
diff --git a/src/gpu/core/Kernel/Incompressible/NavierStokes/B92/B92IncompressibleNavierStokes.cu b/src/gpu/core/Kernel/Incompressible/NavierStokes/B92/B92IncompressibleNavierStokes.cu
index 4ff8d0d6c..49210b1fe 100644
--- a/src/gpu/core/Kernel/Incompressible/NavierStokes/B92/B92IncompressibleNavierStokes.cu
+++ b/src/gpu/core/Kernel/Incompressible/NavierStokes/B92/B92IncompressibleNavierStokes.cu
@@ -30,7 +30,7 @@ B92IncompressibleNavierStokes::B92IncompressibleNavierStokes(std::shared_ptr<Par
this->para = para;
this->level = level;
- myPreProcessorTypes.push_back(InitSP27);
+ myPreProcessorTypes.push_back(InitNavierStokesIncompressible);
}
diff --git a/src/gpu/core/Kernel/Incompressible/NavierStokes/K15/K15IncompressibleNavierStokes.cu b/src/gpu/core/Kernel/Incompressible/NavierStokes/K15/K15IncompressibleNavierStokes.cu
index 40247d0f6..5d64f3216 100644
--- a/src/gpu/core/Kernel/Incompressible/NavierStokes/K15/K15IncompressibleNavierStokes.cu
+++ b/src/gpu/core/Kernel/Incompressible/NavierStokes/K15/K15IncompressibleNavierStokes.cu
@@ -30,7 +30,7 @@ K15IncompressibleNavierStokes::K15IncompressibleNavierStokes(std::shared_ptr<Par
this->para = para;
this->level = level;
- myPreProcessorTypes.push_back(InitSP27);
+ myPreProcessorTypes.push_back(InitNavierStokesIncompressible);
}
diff --git a/src/gpu/core/KernelManager/ADKernelManager.cpp b/src/gpu/core/KernelManager/ADKernelManager.cpp
index 0a531bd36..ac504a839 100644
--- a/src/gpu/core/KernelManager/ADKernelManager.cpp
+++ b/src/gpu/core/KernelManager/ADKernelManager.cpp
@@ -46,7 +46,7 @@ void ADKernelManager::initAD(const int level) const
//////////////////////////////////////////////////////////////////////////
para->getParD(level)->isEvenTimestep = true;
//////////////////////////////////////////////////////////////////////////
- InitADDev27(
+ initAdvectionDiffusion(
para->getParD(level)->numberofthreads,
para->getParD(level)->neighborX,
para->getParD(level)->neighborY,
@@ -62,7 +62,7 @@ void ADKernelManager::initAD(const int level) const
//////////////////////////////////////////////////////////////////////////
para->getParD(level)->isEvenTimestep = false;
//////////////////////////////////////////////////////////////////////////
- InitADDev27(
+ initAdvectionDiffusion(
para->getParD(level)->numberofthreads,
para->getParD(level)->neighborX,
para->getParD(level)->neighborY,
diff --git a/src/gpu/core/PreProcessor/PreProcessorFactory/PreProcessorFactoryImp.cpp b/src/gpu/core/PreProcessor/PreProcessorFactory/PreProcessorFactoryImp.cpp
index e92be775c..9f72e8515 100644
--- a/src/gpu/core/PreProcessor/PreProcessorFactory/PreProcessorFactoryImp.cpp
+++ b/src/gpu/core/PreProcessor/PreProcessorFactory/PreProcessorFactoryImp.cpp
@@ -4,11 +4,11 @@
#include "PreProcessor/PreProcessorStrategy/InitCompAD7/InitCompAD7.h"
#include "PreProcessor/PreProcessorStrategy/InitCompAD27/InitCompAD27.h"
-#include "PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27.h"
-#include "PreProcessor/PreProcessorStrategy/InitF3/InitF3.h"
+#include "PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible.h"
+#include "PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.h"
#include "PreProcessor/PreProcessorStrategy/InitIncompAD27/InitIncompAD27.h"
#include "PreProcessor/PreProcessorStrategy/InitIncompAD7/InitIncompAD7.h"
-#include "PreProcessor/PreProcessorStrategy/InitSP27/InitSP27.h"
+#include "PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible.h"
std::shared_ptr<PreProcessor> PreProcessorFactoryImp::makePreProcessor(std::vector<PreProcessorType> preProcessorTypes, std::shared_ptr<Parameter> para)
@@ -25,14 +25,14 @@ std::shared_ptr<PreProcessorStrategy> PreProcessorFactoryImp::makePreProcessorSt
{
switch (preProcessorType)
{
- case InitSP27:
- return InitSP27::getNewInstance(para);
+ case InitNavierStokesIncompressible:
+ return InitNavierStokesIncompressible::getNewInstance(para);
break;
- case InitCompSP27:
- return InitCompSP27::getNewInstance(para);
+ case InitNavierStokesCompressible:
+ return InitNavierStokesCompressible::getNewInstance(para);
break;
- case InitF3:
- return InitF3::getNewInstance(para);
+ case InitK18K20NavierStokesCompressible:
+ return InitK18K20NavierStokesCompressible::getNewInstance(para);
break;
case InitIncompAD7:
return InitIncompAD7::getNewInstance(para);
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27.cu
deleted file mode 100644
index e70d5a8fe..000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27.cu
+++ /dev/null
@@ -1,68 +0,0 @@
-#include "InitCompSP27.h"
-
-#include "InitCompSP27_Device.cuh"
-#include "Parameter/Parameter.h"
-#include <cuda_helper/CudaGrid.h>
-
-std::shared_ptr<PreProcessorStrategy> InitCompSP27::getNewInstance(std::shared_ptr<Parameter> para)
-{
- return std::shared_ptr<PreProcessorStrategy>(new InitCompSP27(para));
-}
-
-void InitCompSP27::init(int level)
-{
- vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
-
- if( ! para->getUseInitNeq() )
- {
- LB_Init_Comp_SP_27 <<< grid.grid, grid.threads >>> (
- para->getParD(level)->neighborX,
- para->getParD(level)->neighborY,
- para->getParD(level)->neighborZ,
- para->getParD(level)->typeOfGridNode,
- para->getParD(level)->rho,
- para->getParD(level)->velocityX,
- para->getParD(level)->velocityY,
- para->getParD(level)->velocityZ,
- para->getParD(level)->numberOfNodes,
- para->getParD(level)->distributions.f[0],
- para->getParD(level)->isEvenTimestep);
- getLastCudaError("LB_Init_Comp_SP_27 execution failed");
- }
- else
- {
- LB_Init_Comp_Neq_SP_27 <<< grid.grid, grid.threads >>> (
- para->getParD(level)->neighborX,
- para->getParD(level)->neighborY,
- para->getParD(level)->neighborZ,
- para->getParD(level)->neighborInverse,
- para->getParD(level)->typeOfGridNode,
- para->getParD(level)->rho,
- para->getParD(level)->velocityX,
- para->getParD(level)->velocityY,
- para->getParD(level)->velocityZ,
- para->getParD(level)->numberOfNodes,
- para->getParD(level)->distributions.f[0],
- para->getParD(level)->omega,
- para->getParD(level)->isEvenTimestep);
- cudaDeviceSynchronize();
- getLastCudaError("LB_Init_Comp_Neq_SP_27 execution failed");
- }
-
-
-
-}
-
-bool InitCompSP27::checkParameter()
-{
- return false;
-}
-
-InitCompSP27::InitCompSP27(std::shared_ptr<Parameter> para)
-{
- this->para = para;
-}
-
-InitCompSP27::InitCompSP27()
-{
-}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27.h
deleted file mode 100644
index 601f1a446..000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27.h
+++ /dev/null
@@ -1,23 +0,0 @@
-#ifndef INIT_COMP_SP27_H
-#define INIT_COMP_SP27_H
-
-#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h"
-
-#include <memory>
-
-class Parameter;
-
-class InitCompSP27 : public PreProcessorStrategy
-{
-public:
- static std::shared_ptr<PreProcessorStrategy> getNewInstance(std::shared_ptr< Parameter> para);
- void init(int level);
- bool checkParameter();
-
-private:
- InitCompSP27();
- InitCompSP27(std::shared_ptr< Parameter> para);
- std::shared_ptr< Parameter> para;
-};
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27_Device.cuh
deleted file mode 100644
index d7f40b88b..000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27_Device.cuh
+++ /dev/null
@@ -1,33 +0,0 @@
-#ifndef LB_INIT_COMP_SP27_H
-#define LB_INIT_COMP_SP27_H
-
-#include <DataTypes.h>
-#include <curand.h>
-
-__global__ void LB_Init_Comp_SP_27(unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* geoD,
- real* rho,
- real* ux,
- real* uy,
- real* uz,
- unsigned int size_Mat,
- real* DD,
- bool EvenOrOdd);
-
-__global__ void LB_Init_Comp_Neq_SP_27(unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* neighborWSB,
- unsigned int* geoD,
- real* rho,
- real* ux,
- real* uy,
- real* uz,
- unsigned int size_Mat,
- real* DD,
- real omega,
- bool EvenOrOdd);
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3.cu
deleted file mode 100644
index 82084e354..000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3.cu
+++ /dev/null
@@ -1,43 +0,0 @@
-#include "InitF3.h"
-
-#include "InitF3_Device.cuh"
-#include "Parameter/Parameter.h"
-#include <cuda_helper/CudaGrid.h>
-
-std::shared_ptr<PreProcessorStrategy> InitF3::getNewInstance(std::shared_ptr<Parameter> para)
-{
- return std::shared_ptr<PreProcessorStrategy>(new InitF3(para));
-}
-
-void InitF3::init(int level)
-{
- vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
-
- LB_Init_F3 <<< grid.grid, grid.threads >>>(
- para->getParD(level)->neighborX,
- para->getParD(level)->neighborY,
- para->getParD(level)->neighborZ,
- para->getParD(level)->typeOfGridNode,
- para->getParD(level)->rho,
- para->getParD(level)->velocityX,
- para->getParD(level)->velocityY,
- para->getParD(level)->velocityZ,
- para->getParD(level)->numberOfNodes,
- para->getParD(level)->g6.g[0],
- para->getParD(level)->isEvenTimestep);
- getLastCudaError("LB_Init_F3 execution failed");
-}
-
-bool InitF3::checkParameter()
-{
- return false;
-}
-
-InitF3::InitF3(std::shared_ptr<Parameter> para)
-{
- this->para = para;
-}
-
-InitF3::InitF3()
-{
-}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3.h
deleted file mode 100644
index c254aa906..000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3.h
+++ /dev/null
@@ -1,23 +0,0 @@
-#ifndef INITF3_H
-#define INITF3_H
-
-#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h"
-
-#include <memory>
-
-class Parameter;
-
-class InitF3 : public PreProcessorStrategy
-{
-public:
- static std::shared_ptr<PreProcessorStrategy> getNewInstance(std::shared_ptr< Parameter> para);
- void init(int level);
- bool checkParameter();
-
-private:
- InitF3();
- InitF3(std::shared_ptr< Parameter> para);
- std::shared_ptr< Parameter> para;
-};
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3_Device.cuh
deleted file mode 100644
index 991a5ec75..000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3_Device.cuh
+++ /dev/null
@@ -1,19 +0,0 @@
-#ifndef LB_INIT_F3_H
-#define LB_INIT_F3_H
-
-#include <DataTypes.h>
-#include <curand.h>
-
-__global__ void LB_Init_F3(unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* geoD,
- real* rho,
- real* ux,
- real* uy,
- real* uz,
- unsigned int size_Mat,
- real* G6,
- bool EvenOrOdd);
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.cu
new file mode 100644
index 000000000..b3f298988
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.cu
@@ -0,0 +1,73 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#include "InitK18K20NavierStokesCompressible.h"
+
+#include "InitK18K20NavierStokesCompressible_Device.cuh"
+#include "Parameter/Parameter.h"
+#include <cuda_helper/CudaGrid.h>
+
+std::shared_ptr<PreProcessorStrategy> InitK18K20NavierStokesCompressible::getNewInstance(std::shared_ptr<Parameter> para)
+{
+ return std::shared_ptr<PreProcessorStrategy>(new InitK18K20NavierStokesCompressible(para));
+}
+
+void InitK18K20NavierStokesCompressible::init(int level)
+{
+ vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
+
+ InitK18K20NavierStokesCompressible_Device <<< grid.grid, grid.threads >>>(
+ para->getParD(level)->neighborX,
+ para->getParD(level)->neighborY,
+ para->getParD(level)->neighborZ,
+ para->getParD(level)->typeOfGridNode,
+ para->getParD(level)->rho,
+ para->getParD(level)->velocityX,
+ para->getParD(level)->velocityY,
+ para->getParD(level)->velocityZ,
+ para->getParD(level)->numberOfNodes,
+ para->getParD(level)->g6.g[0],
+ para->getParD(level)->isEvenTimestep);
+ getLastCudaError("InitK18K20NavierStokesCompressible_Device execution failed");
+}
+
+bool InitK18K20NavierStokesCompressible::checkParameter()
+{
+ return false;
+}
+
+InitK18K20NavierStokesCompressible::InitK18K20NavierStokesCompressible(std::shared_ptr<Parameter> para)
+{
+ this->para = para;
+}
+
+InitK18K20NavierStokesCompressible::InitK18K20NavierStokesCompressible()
+{
+}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.h
new file mode 100644
index 000000000..403dc4f0c
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.h
@@ -0,0 +1,53 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#ifndef InitK18K20NavierStokesCompressible_H
+#define InitK18K20NavierStokesCompressible_H
+
+#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h"
+
+#include <memory>
+
+class Parameter;
+
+class InitK18K20NavierStokesCompressible : public PreProcessorStrategy
+{
+public:
+ static std::shared_ptr<PreProcessorStrategy> getNewInstance(std::shared_ptr< Parameter> para);
+ void init(int level);
+ bool checkParameter();
+
+private:
+ InitK18K20NavierStokesCompressible();
+ InitK18K20NavierStokesCompressible(std::shared_ptr< Parameter> para);
+ std::shared_ptr< Parameter> para;
+};
+
+#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3_Device.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible_Device.cu
similarity index 54%
rename from src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3_Device.cu
rename to src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible_Device.cu
index 29050bd28..6beb4ad85 100644
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitF3/InitF3_Device.cu
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible_Device.cu
@@ -1,3 +1,33 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
#include "LBM/LB.h"
#include "lbm/constants/D3Q27.h"
#include <basics/constants/NumericConstants.h>
@@ -7,7 +37,7 @@ using namespace vf::lbm::dir;
#include "math.h"
-__global__ void LB_Init_F3(unsigned int* neighborX,
+__global__ void InitK18K20NavierStokesCompressible_Device(unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned int* geoD,
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible_Device.cuh
new file mode 100644
index 000000000..90ecb0705
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible_Device.cuh
@@ -0,0 +1,50 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#ifndef InitK18K20NavierStokesCompressible_Device_H
+#define InitK18K20NavierStokesCompressible_Device_H
+
+#include <DataTypes.h>
+#include <curand.h>
+
+__global__ void InitK18K20NavierStokesCompressible_Device(
+ unsigned int* neighborX,
+ unsigned int* neighborY,
+ unsigned int* neighborZ,
+ unsigned int* geoD,
+ real* rho,
+ real* ux,
+ real* uy,
+ real* uz,
+ unsigned int size_Mat,
+ real* G6,
+ bool EvenOrOdd);
+
+#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible.cu
new file mode 100644
index 000000000..995607fd3
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible.cu
@@ -0,0 +1,98 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#include "InitNavierStokesCompressible.h"
+
+#include "InitNavierStokesCompressible_Device.cuh"
+#include "Parameter/Parameter.h"
+#include <cuda_helper/CudaGrid.h>
+
+std::shared_ptr<PreProcessorStrategy> InitNavierStokesCompressible::getNewInstance(std::shared_ptr<Parameter> para)
+{
+ return std::shared_ptr<PreProcessorStrategy>(new InitNavierStokesCompressible(para));
+}
+
+void InitNavierStokesCompressible::init(int level)
+{
+ vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
+
+ if( ! para->getUseInitNeq() )
+ {
+ InitNavierStokesCompressible_Device <<< grid.grid, grid.threads >>> (
+ para->getParD(level)->neighborX,
+ para->getParD(level)->neighborY,
+ para->getParD(level)->neighborZ,
+ para->getParD(level)->typeOfGridNode,
+ para->getParD(level)->rho,
+ para->getParD(level)->velocityX,
+ para->getParD(level)->velocityY,
+ para->getParD(level)->velocityZ,
+ para->getParD(level)->numberOfNodes,
+ para->getParD(level)->distributions.f[0],
+ para->getParD(level)->isEvenTimestep);
+ getLastCudaError("InitNavierStokesCompressible_Device execution failed");
+ }
+ else
+ {
+ InitNavierStokesCompressibleNonEquilibrium_Device <<< grid.grid, grid.threads >>> (
+ para->getParD(level)->neighborX,
+ para->getParD(level)->neighborY,
+ para->getParD(level)->neighborZ,
+ para->getParD(level)->neighborInverse,
+ para->getParD(level)->typeOfGridNode,
+ para->getParD(level)->rho,
+ para->getParD(level)->velocityX,
+ para->getParD(level)->velocityY,
+ para->getParD(level)->velocityZ,
+ para->getParD(level)->numberOfNodes,
+ para->getParD(level)->distributions.f[0],
+ para->getParD(level)->omega,
+ para->getParD(level)->isEvenTimestep);
+ cudaDeviceSynchronize();
+ getLastCudaError("InitNavierStokesCompressibleNonEquilibrium_Device execution failed");
+ }
+
+
+
+}
+
+bool InitNavierStokesCompressible::checkParameter()
+{
+ return false;
+}
+
+InitNavierStokesCompressible::InitNavierStokesCompressible(std::shared_ptr<Parameter> para)
+{
+ this->para = para;
+}
+
+InitNavierStokesCompressible::InitNavierStokesCompressible()
+{
+}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible.h
new file mode 100644
index 000000000..804b629a5
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible.h
@@ -0,0 +1,53 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#ifndef InitNavierStokesCompressible_H
+#define InitNavierStokesCompressible_H
+
+#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h"
+
+#include <memory>
+
+class Parameter;
+
+class InitNavierStokesCompressible : public PreProcessorStrategy
+{
+public:
+ static std::shared_ptr<PreProcessorStrategy> getNewInstance(std::shared_ptr< Parameter> para);
+ void init(int level);
+ bool checkParameter();
+
+private:
+ InitNavierStokesCompressible();
+ InitNavierStokesCompressible(std::shared_ptr< Parameter> para);
+ std::shared_ptr< Parameter> para;
+};
+
+#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27_Device.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible_Device.cu
similarity index 91%
rename from src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27_Device.cu
rename to src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible_Device.cu
index f733282fa..cb38eedeb 100644
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitCompSP27/InitCompSP27_Device.cu
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible_Device.cu
@@ -1,3 +1,33 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
#include "LBM/LB.h"
#include "lbm/constants/D3Q27.h"
#include <basics/constants/NumericConstants.h>
@@ -8,7 +38,8 @@ using namespace vf::lbm::dir;
#include <stdio.h>
-__global__ void LB_Init_Comp_SP_27(unsigned int* neighborX,
+__global__ void InitNavierStokesCompressible_Device(
+ unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned int* geoD,
@@ -179,7 +210,7 @@ __global__ void LB_Init_Comp_SP_27(unsigned int* neighborX,
////////////////////////////////////////////////////////////////////////////////
-__global__ void LB_Init_Comp_Neq_SP_27( unsigned int* neighborX,
+__global__ void InitNavierStokesCompressibleNonEquilibrium_Device( unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned int* neighborWSB,
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible_Device.cuh
new file mode 100644
index 000000000..a454d7249
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible_Device.cuh
@@ -0,0 +1,63 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#ifndef InitNavierStokesCompressible_Device_H
+#define InitNavierStokesCompressible_Device_H
+
+#include <DataTypes.h>
+#include <curand.h>
+
+__global__ void InitNavierStokesCompressible_Device(unsigned int* neighborX,
+ unsigned int* neighborY,
+ unsigned int* neighborZ,
+ unsigned int* geoD,
+ real* rho,
+ real* ux,
+ real* uy,
+ real* uz,
+ unsigned int size_Mat,
+ real* DD,
+ bool EvenOrOdd);
+
+__global__ void InitNavierStokesCompressibleNonEquilibrium_Device(unsigned int* neighborX,
+ unsigned int* neighborY,
+ unsigned int* neighborZ,
+ unsigned int* neighborWSB,
+ unsigned int* geoD,
+ real* rho,
+ real* ux,
+ real* uy,
+ real* uz,
+ unsigned int size_Mat,
+ real* DD,
+ real omega,
+ bool EvenOrOdd);
+
+#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible.cu
new file mode 100644
index 000000000..b4c3f01cb
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible.cu
@@ -0,0 +1,73 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#include "InitNavierStokesIncompressible.h"
+
+#include "InitNavierStokesIncompressible_Device.cuh"
+#include "Parameter/Parameter.h"
+#include <cuda_helper/CudaGrid.h>
+
+std::shared_ptr<PreProcessorStrategy> InitNavierStokesIncompressible::getNewInstance(std::shared_ptr<Parameter> para)
+{
+ return std::shared_ptr<PreProcessorStrategy>(new InitNavierStokesIncompressible(para));
+}
+
+void InitNavierStokesIncompressible::init(int level)
+{
+ vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
+
+ InitNavierStokesIncompressible_Device<<<grid.grid, grid.threads>>>(
+ para->getParD(level)->neighborX,
+ para->getParD(level)->neighborY,
+ para->getParD(level)->neighborZ,
+ para->getParD(level)->typeOfGridNode,
+ para->getParD(level)->rho,
+ para->getParD(level)->velocityX,
+ para->getParD(level)->velocityY,
+ para->getParD(level)->velocityZ,
+ para->getParD(level)->numberOfNodes,
+ para->getParD(level)->distributions.f[0],
+ para->getParD(level)->isEvenTimestep);
+ getLastCudaError("LB_Init_SP_27 execution failed");
+}
+
+bool InitNavierStokesIncompressible::checkParameter()
+{
+ return false;
+}
+
+InitNavierStokesIncompressible::InitNavierStokesIncompressible(std::shared_ptr<Parameter> para)
+{
+ this->para = para;
+}
+
+InitNavierStokesIncompressible::InitNavierStokesIncompressible()
+{
+}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible.h
new file mode 100644
index 000000000..1bf9f945e
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible.h
@@ -0,0 +1,53 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#ifndef INIT_SP27_H
+#define INIT_SP27_H
+
+#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h"
+
+#include <memory>
+
+class Parameter;
+
+class InitNavierStokesIncompressible : public PreProcessorStrategy
+{
+public:
+ static std::shared_ptr<PreProcessorStrategy> getNewInstance(std::shared_ptr< Parameter> para);
+ void init(int level);
+ bool checkParameter();
+
+private:
+ InitNavierStokesIncompressible();
+ InitNavierStokesIncompressible(std::shared_ptr< Parameter> para);
+ std::shared_ptr< Parameter> para;
+};
+
+#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27_Device.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible_Device.cu
similarity index 79%
rename from src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27_Device.cu
rename to src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible_Device.cu
index 8677e0f2e..3b49f4427 100644
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27_Device.cu
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible_Device.cu
@@ -1,3 +1,33 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
#include "LBM/LB.h"
#include "lbm/constants/D3Q27.h"
#include <basics/constants/NumericConstants.h>
@@ -6,7 +36,7 @@ using namespace vf::basics::constant;
using namespace vf::lbm::dir;
#include "math.h"
-__global__ void LB_Init_SP_27(unsigned int* neighborX,
+__global__ void InitNavierStokesIncompressible_Device(unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned int* geoD,
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible_Device.cuh
new file mode 100644
index 000000000..44015e010
--- /dev/null
+++ b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible_Device.cuh
@@ -0,0 +1,50 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
+#ifndef InitNavierStokesIncompressible_Device_H
+#define InitNavierStokesIncompressible_Device_H
+
+#include <DataTypes.h>
+#include <curand.h>
+
+__global__ void InitNavierStokesIncompressible_Device(
+ unsigned int* neighborX,
+ unsigned int* neighborY,
+ unsigned int* neighborZ,
+ unsigned int* geoD,
+ real* rho,
+ real* ux,
+ real* uy,
+ real* uz,
+ unsigned int size_Mat,
+ real* DD,
+ bool EvenOrOdd);
+
+#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27.cu
deleted file mode 100644
index 78994e997..000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27.cu
+++ /dev/null
@@ -1,43 +0,0 @@
-#include "InitSP27.h"
-
-#include "InitSP27_Device.cuh"
-#include "Parameter/Parameter.h"
-#include <cuda_helper/CudaGrid.h>
-
-std::shared_ptr<PreProcessorStrategy> InitSP27::getNewInstance(std::shared_ptr<Parameter> para)
-{
- return std::shared_ptr<PreProcessorStrategy>(new InitSP27(para));
-}
-
-void InitSP27::init(int level)
-{
- vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
-
- LB_Init_SP_27 <<< grid.grid, grid.threads >>>(
- para->getParD(level)->neighborX,
- para->getParD(level)->neighborY,
- para->getParD(level)->neighborZ,
- para->getParD(level)->typeOfGridNode,
- para->getParD(level)->rho,
- para->getParD(level)->velocityX,
- para->getParD(level)->velocityY,
- para->getParD(level)->velocityZ,
- para->getParD(level)->numberOfNodes,
- para->getParD(level)->distributions.f[0],
- para->getParD(level)->isEvenTimestep);
- getLastCudaError("LB_Init_SP_27 execution failed");
-}
-
-bool InitSP27::checkParameter()
-{
- return false;
-}
-
-InitSP27::InitSP27(std::shared_ptr<Parameter> para)
-{
- this->para = para;
-}
-
-InitSP27::InitSP27()
-{
-}
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27.h
deleted file mode 100644
index def7ccd17..000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27.h
+++ /dev/null
@@ -1,23 +0,0 @@
-#ifndef INIT_SP27_H
-#define INIT_SP27_H
-
-#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h"
-
-#include <memory>
-
-class Parameter;
-
-class InitSP27 : public PreProcessorStrategy
-{
-public:
- static std::shared_ptr<PreProcessorStrategy> getNewInstance(std::shared_ptr< Parameter> para);
- void init(int level);
- bool checkParameter();
-
-private:
- InitSP27();
- InitSP27(std::shared_ptr< Parameter> para);
- std::shared_ptr< Parameter> para;
-};
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27_Device.cuh
deleted file mode 100644
index c00a191b4..000000000
--- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitSP27/InitSP27_Device.cuh
+++ /dev/null
@@ -1,19 +0,0 @@
-#ifndef LB_INIT_SP27_H
-#define LB_INIT_SP27_H
-
-#include <DataTypes.h>
-#include <curand.h>
-
-__global__ void LB_Init_SP_27(unsigned int* neighborX,
- unsigned int* neighborY,
- unsigned int* neighborZ,
- unsigned int* geoD,
- real* rho,
- real* ux,
- real* uy,
- real* uz,
- unsigned int size_Mat,
- real* DD,
- bool EvenOrOdd);
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/core/PreProcessor/PreProcessorType.h b/src/gpu/core/PreProcessor/PreProcessorType.h
index 6bf54afc0..1fb9b78fc 100644
--- a/src/gpu/core/PreProcessor/PreProcessorType.h
+++ b/src/gpu/core/PreProcessor/PreProcessorType.h
@@ -1,11 +1,41 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \author Martin Schoenherr
+//=======================================================================================
#ifndef PRE_PROCESSOR_TYPE_H
#define PRE_PROCESSOR_TYPE_H
enum PreProcessorType
{
- InitSP27,
- InitCompSP27,
- InitF3,
+ InitNavierStokesIncompressible,
+ InitNavierStokesCompressible,
+ InitK18K20NavierStokesCompressible,
InitIncompAD7,
InitIncompAD27,
InitCompAD7,
--
GitLab