Skip to content
Snippets Groups Projects
Commit 848f64da authored by Anna Wellmann's avatar Anna Wellmann
Browse files

Merge branch 'develop' into feature/rotatingGrids

parents d768f81f 27bdb73c
No related branches found
No related tags found
1 merge request!307[GPU] Add a cylinder geometry
Hide whitespace changes
Inline Side-by-side
Showing
with 62 additions and 76 deletions
apps/gpu/ActuatorLine/configActuatorLine.txt
+ 6
− 6
View file @ 848f64da
...@@ -24,11 +24,11 @@ tStartOutProbe=100 ...@@ -24,11 +24,11 @@ tStartOutProbe=100
tOutProbe=100 tOutProbe=100
################################################## ##################################################
#TurbulenceModel = QR TurbulenceModel = QR
#SGSconstant = 0.3333333 SGSconstant = 0.3333333
#
#QuadricLimiterP = 100000.0 QuadricLimiterP = 10000.0
#QuadricLimiterM = 100000.0 QuadricLimiterM = 10000.0
#QuadricLimiterD = 100000.0 QuadricLimiterD = 10000.0
################################################## ##################################################
src/cpu/VirtualFluidsCore/LBM/K17CompressibleNavierStokes.cpp
+ 1
− 3
View file @ 848f64da
...@@ -153,11 +153,9 @@ void K17CompressibleNavierStokes::calculate(int step) ...@@ -153,11 +153,9 @@ void K17CompressibleNavierStokes::calculate(int step)
real quadricLimiter[3] = { 0.01, 0.01, 0.01 }; // TODO: Where do we configure the quadricLimiter? real quadricLimiter[3] = { 0.01, 0.01, 0.01 }; // TODO: Where do we configure the quadricLimiter?
parameter.quadricLimiter = quadricLimiter; parameter.quadricLimiter = quadricLimiter;
const bool writeMacroscopicVariables = false;
vf::lbm::MacroscopicValues mv; // not used vf::lbm::MacroscopicValues mv; // not used
vf::lbm::TurbulentViscosity tv; // not used vf::lbm::TurbulentViscosity tv; // not used
vf::lbm::runK17CompressibleNavierStokes<vf::lbm::TurbulenceModel::None, writeMacroscopicVariables>(parameter, vf::lbm::runK17CompressibleNavierStokes<vf::lbm::TurbulenceModel::None>(parameter, mv, tv);
mv, tv);
dataSet->getFdistributions()->setDistribution(parameter.distribution, x1, x2, x3); dataSet->getFdistributions()->setDistribution(parameter.distribution, x1, x2, x3);
} }
} }
......
src/gpu/VirtualFluids_GPU/GPU/GridScaling/scaleCF_compressible.cu
+ 17
− 16
View file @ 848f64da
...@@ -38,6 +38,7 @@ ...@@ -38,6 +38,7 @@
#include <lbm/refinement/InterpolationCF.h> #include <lbm/refinement/InterpolationCF.h>
#include <lbm/interpolation/InterpolationCoefficients.h> #include <lbm/interpolation/InterpolationCoefficients.h>
#include <lbm/collision/TurbulentViscosity.h>
...@@ -89,7 +90,7 @@ template <bool hasTurbulentViscosity> __device__ void interpolate( ...@@ -89,7 +90,7 @@ template <bool hasTurbulentViscosity> __device__ void interpolate(
//! - Set moments (zeroth to sixth order) on destination node //! - Set moments (zeroth to sixth order) on destination node
//! //!
real omegaF = omegaFine; real omegaF = omegaFine;
omegaF = hasTurbulentViscosity ? vf::gpu::calculateOmega(omegaFine, turbulentViscosityFine[indices.k_000]) : omegaFine; omegaF = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omegaFine, turbulentViscosityFine[indices.k_000]) : omegaFine;
const real epsilon_new = c1o2; // ratio of grid resolutions const real epsilon_new = c1o2; // ratio of grid resolutions
real f[27]; real f[27];
...@@ -120,7 +121,7 @@ template <bool hasTurbulentViscosity> __device__ void interpolate( ...@@ -120,7 +121,7 @@ template <bool hasTurbulentViscosity> __device__ void interpolate(
//////////////////////////////////////////////////////////////////////////////// ////////////////////////////////////////////////////////////////////////////////
// Set moments (zeroth to sixth orders) on destination node // Set moments (zeroth to sixth orders) on destination node
omegaF = hasTurbulentViscosity ? vf::gpu::calculateOmega(omegaFine, turbulentViscosityFine[indices.k_000]) : omegaFine; omegaF = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omegaFine, turbulentViscosityFine[indices.k_000]) : omegaFine;
vf::lbm::interpolateCF(f, omegaF, epsilon_new, coefficients, x, y, z); vf::lbm::interpolateCF(f, omegaF, epsilon_new, coefficients, x, y, z);
...@@ -148,7 +149,7 @@ template <bool hasTurbulentViscosity> __device__ void interpolate( ...@@ -148,7 +149,7 @@ template <bool hasTurbulentViscosity> __device__ void interpolate(
//////////////////////////////////////////////////////////////////////////////// ////////////////////////////////////////////////////////////////////////////////
// Set moments (zeroth to sixth orders) on destination node // Set moments (zeroth to sixth orders) on destination node
omegaF = hasTurbulentViscosity ? vf::gpu::calculateOmega(omegaFine, turbulentViscosityFine[indices.k_000]) : omegaFine; omegaF = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omegaFine, turbulentViscosityFine[indices.k_000]) : omegaFine;
vf::lbm::interpolateCF(f, omegaF, epsilon_new, coefficients, x, y, z); vf::lbm::interpolateCF(f, omegaF, epsilon_new, coefficients, x, y, z);
...@@ -176,7 +177,7 @@ template <bool hasTurbulentViscosity> __device__ void interpolate( ...@@ -176,7 +177,7 @@ template <bool hasTurbulentViscosity> __device__ void interpolate(
//////////////////////////////////////////////////////////////////////////////// ////////////////////////////////////////////////////////////////////////////////
// Set moments (zeroth to sixth orders) on destination node // Set moments (zeroth to sixth orders) on destination node
omegaF = hasTurbulentViscosity ? vf::gpu::calculateOmega(omegaFine, turbulentViscosityFine[indices.k_000]) : omegaFine; omegaF = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omegaFine, turbulentViscosityFine[indices.k_000]) : omegaFine;
vf::lbm::interpolateCF(f, omegaF, epsilon_new, coefficients, x, y, z); vf::lbm::interpolateCF(f, omegaF, epsilon_new, coefficients, x, y, z);
...@@ -215,7 +216,7 @@ template <bool hasTurbulentViscosity> __device__ void interpolate( ...@@ -215,7 +216,7 @@ template <bool hasTurbulentViscosity> __device__ void interpolate(
//////////////////////////////////////////////////////////////////////////////// ////////////////////////////////////////////////////////////////////////////////
// Set moments (zeroth to sixth orders) on destination node // Set moments (zeroth to sixth orders) on destination node
omegaF = hasTurbulentViscosity ? vf::gpu::calculateOmega(omegaFine, turbulentViscosityFine[indices.k_000]) : omegaFine; omegaF = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omegaFine, turbulentViscosityFine[indices.k_000]) : omegaFine;
vf::lbm::interpolateCF(f, omegaF, epsilon_new, coefficients, x, y, z); vf::lbm::interpolateCF(f, omegaF, epsilon_new, coefficients, x, y, z);
...@@ -239,7 +240,7 @@ template <bool hasTurbulentViscosity> __device__ void interpolate( ...@@ -239,7 +240,7 @@ template <bool hasTurbulentViscosity> __device__ void interpolate(
indices.k_0MM = neighborZfine[indices.k_0MM]; indices.k_0MM = neighborZfine[indices.k_0MM];
indices.k_MMM = neighborZfine[indices.k_MMM]; indices.k_MMM = neighborZfine[indices.k_MMM];
omegaF = hasTurbulentViscosity ? vf::gpu::calculateOmega(omegaFine, turbulentViscosityFine[indices.k_000]) : omegaFine; omegaF = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omegaFine, turbulentViscosityFine[indices.k_000]) : omegaFine;
vf::lbm::interpolateCF(f, omegaF, epsilon_new, coefficients, x, y, z); vf::lbm::interpolateCF(f, omegaF, epsilon_new, coefficients, x, y, z);
...@@ -263,7 +264,7 @@ template <bool hasTurbulentViscosity> __device__ void interpolate( ...@@ -263,7 +264,7 @@ template <bool hasTurbulentViscosity> __device__ void interpolate(
indices.k_0MM = indices.k_MMM; indices.k_0MM = indices.k_MMM;
indices.k_MMM = neighborXfine[indices.k_MMM]; indices.k_MMM = neighborXfine[indices.k_MMM];
omegaF = hasTurbulentViscosity ? vf::gpu::calculateOmega(omegaFine, turbulentViscosityFine[indices.k_000]) : omegaFine; omegaF = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omegaFine, turbulentViscosityFine[indices.k_000]) : omegaFine;
vf::lbm::interpolateCF(f, omegaF, epsilon_new, coefficients, x, y, z); vf::lbm::interpolateCF(f, omegaF, epsilon_new, coefficients, x, y, z);
...@@ -287,7 +288,7 @@ template <bool hasTurbulentViscosity> __device__ void interpolate( ...@@ -287,7 +288,7 @@ template <bool hasTurbulentViscosity> __device__ void interpolate(
indices.k_0M0 = k_base_MM0; indices.k_0M0 = k_base_MM0;
indices.k_MM0 = neighborXfine[k_base_MM0]; indices.k_MM0 = neighborXfine[k_base_MM0];
omegaF = hasTurbulentViscosity ? vf::gpu::calculateOmega(omegaFine, turbulentViscosityFine[indices.k_000]) : omegaFine; omegaF = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omegaFine, turbulentViscosityFine[indices.k_000]) : omegaFine;
vf::lbm::interpolateCF(f, omegaF, epsilon_new, coefficients, x, y, z); vf::lbm::interpolateCF(f, omegaF, epsilon_new, coefficients, x, y, z);
...@@ -365,7 +366,7 @@ template<bool hasTurbulentViscosity> __global__ void scaleCF_compressible( ...@@ -365,7 +366,7 @@ template<bool hasTurbulentViscosity> __global__ void scaleCF_compressible(
indices.k_0MM = k_base_0MM; indices.k_0MM = k_base_0MM;
indices.k_MMM = k_base_MMM; indices.k_MMM = k_base_MMM;
omegaC = hasTurbulentViscosity ? vf::gpu::calculateOmega(omegaCoarse, turbulentViscosityCoarse[indices.k_000]) : omegaCoarse; omegaC = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omegaCoarse, turbulentViscosityCoarse[indices.k_000]) : omegaCoarse;
real f_coarse[27]; real f_coarse[27];
...@@ -385,7 +386,7 @@ template<bool hasTurbulentViscosity> __global__ void scaleCF_compressible( ...@@ -385,7 +386,7 @@ template<bool hasTurbulentViscosity> __global__ void scaleCF_compressible(
indices.k_0MM = neighborZcoarse[indices.k_0MM]; indices.k_0MM = neighborZcoarse[indices.k_0MM];
indices.k_MMM = neighborZcoarse[indices.k_MMM]; indices.k_MMM = neighborZcoarse[indices.k_MMM];
omegaC = hasTurbulentViscosity ? vf::gpu::calculateOmega(omegaCoarse, turbulentViscosityCoarse[indices.k_000]) : omegaCoarse; omegaC = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omegaCoarse, turbulentViscosityCoarse[indices.k_000]) : omegaCoarse;
read(f_coarse, distCoarse, indices); read(f_coarse, distCoarse, indices);
momentsSet.calculateMMP(f_coarse, omegaC); momentsSet.calculateMMP(f_coarse, omegaC);
...@@ -403,7 +404,7 @@ template<bool hasTurbulentViscosity> __global__ void scaleCF_compressible( ...@@ -403,7 +404,7 @@ template<bool hasTurbulentViscosity> __global__ void scaleCF_compressible(
indices.k_0MM = indices.k_MMM; indices.k_0MM = indices.k_MMM;
indices.k_MMM = neighborXcoarse[indices.k_MMM]; indices.k_MMM = neighborXcoarse[indices.k_MMM];
omegaC = hasTurbulentViscosity ? vf::gpu::calculateOmega(omegaCoarse, turbulentViscosityCoarse[indices.k_000]) : omegaCoarse; omegaC = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omegaCoarse, turbulentViscosityCoarse[indices.k_000]) : omegaCoarse;
read(f_coarse, distCoarse, indices); read(f_coarse, distCoarse, indices);
momentsSet.calculatePMP(f_coarse, omegaC); momentsSet.calculatePMP(f_coarse, omegaC);
...@@ -421,7 +422,7 @@ template<bool hasTurbulentViscosity> __global__ void scaleCF_compressible( ...@@ -421,7 +422,7 @@ template<bool hasTurbulentViscosity> __global__ void scaleCF_compressible(
indices.k_0M0 = k_base_MM0; indices.k_0M0 = k_base_MM0;
indices.k_MM0 = neighborXcoarse[k_base_MM0]; indices.k_MM0 = neighborXcoarse[k_base_MM0];
omegaC = hasTurbulentViscosity ? vf::gpu::calculateOmega(omegaCoarse, turbulentViscosityCoarse[indices.k_000]) : omegaCoarse; omegaC = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omegaCoarse, turbulentViscosityCoarse[indices.k_000]) : omegaCoarse;
read(f_coarse, distCoarse, indices); read(f_coarse, distCoarse, indices);
momentsSet.calculatePMM(f_coarse, omegaC); momentsSet.calculatePMM(f_coarse, omegaC);
...@@ -449,7 +450,7 @@ template<bool hasTurbulentViscosity> __global__ void scaleCF_compressible( ...@@ -449,7 +450,7 @@ template<bool hasTurbulentViscosity> __global__ void scaleCF_compressible(
indices.k_0MM = k_base_0MM; indices.k_0MM = k_base_0MM;
indices.k_MMM = k_base_MMM; indices.k_MMM = k_base_MMM;
omegaC = hasTurbulentViscosity ? vf::gpu::calculateOmega(omegaCoarse, turbulentViscosityCoarse[indices.k_000]) : omegaCoarse; omegaC = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omegaCoarse, turbulentViscosityCoarse[indices.k_000]) : omegaCoarse;
read(f_coarse, distCoarse, indices); read(f_coarse, distCoarse, indices);
momentsSet.calculateMPM(f_coarse, omegaC); momentsSet.calculateMPM(f_coarse, omegaC);
...@@ -467,7 +468,7 @@ template<bool hasTurbulentViscosity> __global__ void scaleCF_compressible( ...@@ -467,7 +468,7 @@ template<bool hasTurbulentViscosity> __global__ void scaleCF_compressible(
indices.k_0MM = neighborZcoarse[indices.k_0MM]; indices.k_0MM = neighborZcoarse[indices.k_0MM];
indices.k_MMM = neighborZcoarse[indices.k_MMM]; indices.k_MMM = neighborZcoarse[indices.k_MMM];
omegaC = hasTurbulentViscosity ? vf::gpu::calculateOmega(omegaCoarse, turbulentViscosityCoarse[indices.k_000]) : omegaCoarse; omegaC = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omegaCoarse, turbulentViscosityCoarse[indices.k_000]) : omegaCoarse;
read(f_coarse, distCoarse, indices); read(f_coarse, distCoarse, indices);
momentsSet.calculateMPP(f_coarse, omegaC); momentsSet.calculateMPP(f_coarse, omegaC);
...@@ -484,7 +485,7 @@ template<bool hasTurbulentViscosity> __global__ void scaleCF_compressible( ...@@ -484,7 +485,7 @@ template<bool hasTurbulentViscosity> __global__ void scaleCF_compressible(
indices.k_0MM = indices.k_MMM; indices.k_0MM = indices.k_MMM;
indices.k_MMM = neighborXcoarse[indices.k_MMM]; indices.k_MMM = neighborXcoarse[indices.k_MMM];
omegaC = hasTurbulentViscosity ? vf::gpu::calculateOmega(omegaCoarse, turbulentViscosityCoarse[indices.k_000]) : omegaCoarse; omegaC = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omegaCoarse, turbulentViscosityCoarse[indices.k_000]) : omegaCoarse;
read(f_coarse, distCoarse, indices); read(f_coarse, distCoarse, indices);
momentsSet.calculatePPP(f_coarse, omegaC); momentsSet.calculatePPP(f_coarse, omegaC);
...@@ -501,7 +502,7 @@ template<bool hasTurbulentViscosity> __global__ void scaleCF_compressible( ...@@ -501,7 +502,7 @@ template<bool hasTurbulentViscosity> __global__ void scaleCF_compressible(
indices.k_0M0 = k_base_MM0; indices.k_0M0 = k_base_MM0;
indices.k_MM0 = neighborXcoarse[k_base_MM0]; indices.k_MM0 = neighborXcoarse[k_base_MM0];
omegaC = hasTurbulentViscosity ? vf::gpu::calculateOmega(omegaCoarse, turbulentViscosityCoarse[indices.k_000]) : omegaCoarse; omegaC = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omegaCoarse, turbulentViscosityCoarse[indices.k_000]) : omegaCoarse;
read(f_coarse, distCoarse, indices); read(f_coarse, distCoarse, indices);
momentsSet.calculatePPM(f_coarse, omegaC); momentsSet.calculatePPM(f_coarse, omegaC);
......
src/gpu/VirtualFluids_GPU/GPU/GridScaling/scaleFC_compressible.cu
+ 1
− 1
View file @ 848f64da
...@@ -60,7 +60,7 @@ template <bool hasTurbulentViscosity> __device__ void interpolate( ...@@ -60,7 +60,7 @@ template <bool hasTurbulentViscosity> __device__ void interpolate(
vf::gpu::ListIndices indices(indicesCoarse000[nodeIndex], neighborXcoarse, neighborYcoarse, neighborZcoarse); vf::gpu::ListIndices indices(indicesCoarse000[nodeIndex], neighborXcoarse, neighborYcoarse, neighborZcoarse);
const real epsilonNew = c2o1; // ratio of grid resolutions const real epsilonNew = c2o1; // ratio of grid resolutions
const real omegaCoarseNew = hasTurbulentViscosity ? vf::gpu::calculateOmega(omegaCoarse, turbulentViscosityCoarse[indices.k_000]) : omegaCoarse; const real omegaCoarseNew = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omegaCoarse, turbulentViscosityCoarse[indices.k_000]) : omegaCoarse;
real fCoarse[27]; real fCoarse[27];
vf::lbm::interpolateFC(fCoarse, epsilonNew, omegaCoarseNew, coefficients); vf::lbm::interpolateFC(fCoarse, epsilonNew, omegaCoarseNew, coefficients);
......
src/gpu/VirtualFluids_GPU/Kernel/Compressible/FluidFlow/K17/K17CompressibleNavierStokes.cu
+ 5
− 18
View file @ 848f64da
...@@ -60,7 +60,7 @@ void K17CompressibleNavierStokes<turbulenceModel>::run() ...@@ -60,7 +60,7 @@ void K17CompressibleNavierStokes<turbulenceModel>::run()
auto collision = [] __device__(vf::lbm::CollisionParameter & parameter, vf::lbm::MacroscopicValues & macroscopicValues, auto collision = [] __device__(vf::lbm::CollisionParameter & parameter, vf::lbm::MacroscopicValues & macroscopicValues,
vf::lbm::TurbulentViscosity & turbulentViscosity) { vf::lbm::TurbulentViscosity & turbulentViscosity) {
return vf::lbm::runK17CompressibleNavierStokes<turbulenceModel, false>(parameter, macroscopicValues, turbulentViscosity); return vf::lbm::runK17CompressibleNavierStokes<turbulenceModel>(parameter, macroscopicValues, turbulentViscosity);
}; };
vf::gpu::runCollision<decltype(collision), turbulenceModel, false, false><<<cudaGrid.grid, cudaGrid.threads>>>(collision, kernelParameter); vf::gpu::runCollision<decltype(collision), turbulenceModel, false, false><<<cudaGrid.grid, cudaGrid.threads>>>(collision, kernelParameter);
...@@ -75,24 +75,19 @@ void K17CompressibleNavierStokes<turbulenceModel>::runOnIndices(const unsigned i ...@@ -75,24 +75,19 @@ void K17CompressibleNavierStokes<turbulenceModel>::runOnIndices(const unsigned i
{ {
cudaStream_t stream = para->getStreamManager()->getStream(streamIndex); cudaStream_t stream = para->getStreamManager()->getStream(streamIndex);
auto collision = [] __device__(vf::lbm::CollisionParameter& parameter, vf::lbm::MacroscopicValues& macroscopicValues, vf::lbm::TurbulentViscosity& turbulentViscosity) {
return vf::lbm::runK17CompressibleNavierStokes<turbulenceModel>(parameter, macroscopicValues, turbulentViscosity);
};
switch (collisionTemplate) { switch (collisionTemplate) {
case CollisionTemplate::Default: { case CollisionTemplate::Default: {
vf::gpu::GPUCollisionParameter kernelParameter = getCollisionParameter(para, level, indices, size_indices); vf::gpu::GPUCollisionParameter kernelParameter = getCollisionParameter(para, level, indices, size_indices);
auto collision = [] __device__(vf::lbm::CollisionParameter& parameter, vf::lbm::MacroscopicValues& macroscopicValues, vf::lbm::TurbulentViscosity& turbulentViscosity) {
return vf::lbm::runK17CompressibleNavierStokes<turbulenceModel, false>(parameter, macroscopicValues, turbulentViscosity);
};
vf::gpu::runCollision<decltype(collision), turbulenceModel, false, false><<<cudaGrid.grid, cudaGrid.threads, 0, stream>>>(collision, kernelParameter); vf::gpu::runCollision<decltype(collision), turbulenceModel, false, false><<<cudaGrid.grid, cudaGrid.threads, 0, stream>>>(collision, kernelParameter);
break; break;
} }
case CollisionTemplate::WriteMacroVars: { case CollisionTemplate::WriteMacroVars: {
vf::gpu::GPUCollisionParameter kernelParameter = getCollisionParameter(para, level, indices, size_indices); vf::gpu::GPUCollisionParameter kernelParameter = getCollisionParameter(para, level, indices, size_indices);
auto collision = [] __device__(vf::lbm::CollisionParameter& parameter, vf::lbm::MacroscopicValues& macroscopicValues, vf::lbm::TurbulentViscosity& turbulentViscosity) {
return vf::lbm::runK17CompressibleNavierStokes<turbulenceModel, true>(parameter, macroscopicValues, turbulentViscosity);
};
vf::gpu::runCollision<decltype(collision), turbulenceModel, true, false><<<cudaGrid.grid, cudaGrid.threads, 0, stream>>>(collision, kernelParameter); vf::gpu::runCollision<decltype(collision), turbulenceModel, true, false><<<cudaGrid.grid, cudaGrid.threads, 0, stream>>>(collision, kernelParameter);
break; break;
...@@ -100,20 +95,12 @@ void K17CompressibleNavierStokes<turbulenceModel>::runOnIndices(const unsigned i ...@@ -100,20 +95,12 @@ void K17CompressibleNavierStokes<turbulenceModel>::runOnIndices(const unsigned i
case CollisionTemplate::SubDomainBorder: case CollisionTemplate::SubDomainBorder:
case CollisionTemplate::AllFeatures: { case CollisionTemplate::AllFeatures: {
vf::gpu::GPUCollisionParameter kernelParameter = getCollisionParameter(para, level, indices, size_indices); vf::gpu::GPUCollisionParameter kernelParameter = getCollisionParameter(para, level, indices, size_indices);
auto collision = [] __device__(vf::lbm::CollisionParameter& parameter, vf::lbm::MacroscopicValues& macroscopicValues, vf::lbm::TurbulentViscosity& turbulentViscosity) {
return vf::lbm::runK17CompressibleNavierStokes<turbulenceModel, true>(parameter, macroscopicValues, turbulentViscosity);
};
vf::gpu::runCollision<decltype(collision), turbulenceModel, true, true><<<cudaGrid.grid, cudaGrid.threads, 0, stream>>>(collision, kernelParameter); vf::gpu::runCollision<decltype(collision), turbulenceModel, true, true><<<cudaGrid.grid, cudaGrid.threads, 0, stream>>>(collision, kernelParameter);
break; break;
} }
case CollisionTemplate::ApplyBodyForce: { case CollisionTemplate::ApplyBodyForce: {
vf::gpu::GPUCollisionParameter kernelParameter = getCollisionParameter(para, level, indices, size_indices); vf::gpu::GPUCollisionParameter kernelParameter = getCollisionParameter(para, level, indices, size_indices);
auto collision = [] __device__(vf::lbm::CollisionParameter& parameter, vf::lbm::MacroscopicValues& macroscopicValues, vf::lbm::TurbulentViscosity& turbulentViscosity) {
return vf::lbm::runK17CompressibleNavierStokes<turbulenceModel, false>(parameter, macroscopicValues, turbulentViscosity);
};
vf::gpu::runCollision<decltype(collision), turbulenceModel, false, true><<<cudaGrid.grid, cudaGrid.threads, 0, stream>>>(collision, kernelParameter); vf::gpu::runCollision<decltype(collision), turbulenceModel, false, true><<<cudaGrid.grid, cudaGrid.threads, 0, stream>>>(collision, kernelParameter);
break; break;
......
src/gpu/VirtualFluids_GPU/LBM/GPUHelperFunctions/KernelUtilities.h
+ 0
− 5
View file @ 848f64da
...@@ -391,11 +391,6 @@ __inline__ __device__ void writeInverse(Distributions27& destination, const List ...@@ -391,11 +391,6 @@ __inline__ __device__ void writeInverse(Distributions27& destination, const List
(destination.f[DIR_MMM])[indices.k_MMM] = f[DIR_PPP]; (destination.f[DIR_MMM])[indices.k_MMM] = f[DIR_PPP];
} }
__inline__ __device__ real calculateOmega(const real omega_old, real turbulenceViscosity)
{
return omega_old / (c1o1 + c3o1 * omega_old * turbulenceViscosity);
}
} }
#endif #endif
src/gpu/VirtualFluids_GPU/LBM/GPUHelperFunctions/RunCollision.cuh
+ 4
− 4
View file @ 848f64da
...@@ -58,9 +58,9 @@ __global__ void runCollision(CollisionFunctor collision, GPUCollisionParameter c ...@@ -58,9 +58,9 @@ __global__ void runCollision(CollisionFunctor collision, GPUCollisionParameter c
para.forceZ = (collisionParameter.forces[2] + collisionParameter.bodyForceZ[k_000]) * c1o2 * collisionParameter.forceFactor; para.forceZ = (collisionParameter.forces[2] + collisionParameter.bodyForceZ[k_000]) * c1o2 * collisionParameter.forceFactor;
// Reset body force. To be used when not using round-off correction. // Reset body force. To be used when not using round-off correction.
collisionParameter.bodyForceX[k_000] = 0.0f; collisionParameter.bodyForceX[k_000] = c0o1;
collisionParameter.bodyForceX[k_000] = 0.0f; collisionParameter.bodyForceY[k_000] = c0o1;
collisionParameter.bodyForceX[k_000] = 0.0f; collisionParameter.bodyForceZ[k_000] = c0o1;
//////////////////////////////////////////////////////////////////////////////////// ////////////////////////////////////////////////////////////////////////////////////
//!> Round-off correction //!> Round-off correction
...@@ -97,7 +97,7 @@ __global__ void runCollision(CollisionFunctor collision, GPUCollisionParameter c ...@@ -97,7 +97,7 @@ __global__ void runCollision(CollisionFunctor collision, GPUCollisionParameter c
vf::lbm::MacroscopicValues macroscopicValues; vf::lbm::MacroscopicValues macroscopicValues;
collision(para, macroscopicValues, turbulentViscosity); collision(para, macroscopicValues, turbulentViscosity);
if (writeMacroscopicVariables) { if (writeMacroscopicVariables || turbulenceModel == vf::lbm::TurbulenceModel::AMD) {
collisionParameter.vx[k_000] = macroscopicValues.vx; collisionParameter.vx[k_000] = macroscopicValues.vx;
collisionParameter.vy[k_000] = macroscopicValues.vy; collisionParameter.vy[k_000] = macroscopicValues.vy;
collisionParameter.vz[k_000] = macroscopicValues.vz; collisionParameter.vz[k_000] = macroscopicValues.vz;
......
src/gpu/VirtualFluids_GPU/LBM/GPUHelperFunctions/ScalingUtilities.h
+ 10
− 9
View file @ 848f64da
...@@ -41,6 +41,7 @@ ...@@ -41,6 +41,7 @@
#include <basics/DataTypes.h> #include <basics/DataTypes.h>
#include <lbm/interpolation/InterpolationCoefficients.h> #include <lbm/interpolation/InterpolationCoefficients.h>
#include <lbm/collision/TurbulentViscosity.h>
using namespace vf::basics::constant; using namespace vf::basics::constant;
using namespace vf::lbm; using namespace vf::lbm;
...@@ -97,7 +98,7 @@ template<bool hasTurbulentViscosity> __device__ void calculateMomentSet( ...@@ -97,7 +98,7 @@ template<bool hasTurbulentViscosity> __device__ void calculateMomentSet(
indices.k_0MM = k_base_0MM; indices.k_0MM = k_base_0MM;
indices.k_MMM = k_base_MMM; indices.k_MMM = k_base_MMM;
omega_ = hasTurbulentViscosity ? calculateOmega(omega, turbulentViscosity[indices.k_000]) : omega; omega_ = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omega, turbulentViscosity[indices.k_000]) : omega;
real f_fine[27]; real f_fine[27];
...@@ -123,7 +124,7 @@ template<bool hasTurbulentViscosity> __device__ void calculateMomentSet( ...@@ -123,7 +124,7 @@ template<bool hasTurbulentViscosity> __device__ void calculateMomentSet(
indices.k_0MM = neighborZ[indices.k_0MM]; indices.k_0MM = neighborZ[indices.k_0MM];
indices.k_MMM = neighborZ[indices.k_MMM]; indices.k_MMM = neighborZ[indices.k_MMM];
omega_ = hasTurbulentViscosity ? calculateOmega(omega, turbulentViscosity[indices.k_000]) : omega; omega_ = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omega, turbulentViscosity[indices.k_000]) : omega;
read(f_fine, distFine, indices); read(f_fine, distFine, indices);
...@@ -146,7 +147,7 @@ template<bool hasTurbulentViscosity> __device__ void calculateMomentSet( ...@@ -146,7 +147,7 @@ template<bool hasTurbulentViscosity> __device__ void calculateMomentSet(
indices.k_0MM = indices.k_MMM; indices.k_0MM = indices.k_MMM;
indices.k_MMM = neighborX[indices.k_MMM]; indices.k_MMM = neighborX[indices.k_MMM];
omega_ = hasTurbulentViscosity ? calculateOmega(omega, turbulentViscosity[indices.k_000]) : omega; omega_ = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omega, turbulentViscosity[indices.k_000]) : omega;
read(f_fine, distFine, indices); read(f_fine, distFine, indices);
if (forcesForAllNodesX) if (forcesForAllNodesX)
...@@ -168,7 +169,7 @@ template<bool hasTurbulentViscosity> __device__ void calculateMomentSet( ...@@ -168,7 +169,7 @@ template<bool hasTurbulentViscosity> __device__ void calculateMomentSet(
indices.k_0M0 = k_base_MM0; indices.k_0M0 = k_base_MM0;
indices.k_MM0 = neighborX[k_base_MM0]; indices.k_MM0 = neighborX[k_base_MM0];
omega_ = hasTurbulentViscosity ? calculateOmega(omega, turbulentViscosity[indices.k_000]) : omega; omega_ = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omega, turbulentViscosity[indices.k_000]) : omega;
read(f_fine, distFine, indices); read(f_fine, distFine, indices);
if (forcesForAllNodesX) if (forcesForAllNodesX)
...@@ -201,7 +202,7 @@ template<bool hasTurbulentViscosity> __device__ void calculateMomentSet( ...@@ -201,7 +202,7 @@ template<bool hasTurbulentViscosity> __device__ void calculateMomentSet(
indices.k_0MM = k_base_0MM; indices.k_0MM = k_base_0MM;
indices.k_MMM = k_base_MMM; indices.k_MMM = k_base_MMM;
omega_ = hasTurbulentViscosity ? calculateOmega(omega, turbulentViscosity[indices.k_000]) : omega; omega_ = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omega, turbulentViscosity[indices.k_000]) : omega;
read(f_fine, distFine, indices); read(f_fine, distFine, indices);
if (forcesForAllNodesX) if (forcesForAllNodesX)
...@@ -223,7 +224,7 @@ template<bool hasTurbulentViscosity> __device__ void calculateMomentSet( ...@@ -223,7 +224,7 @@ template<bool hasTurbulentViscosity> __device__ void calculateMomentSet(
indices.k_0MM = neighborZ[indices.k_0MM]; indices.k_0MM = neighborZ[indices.k_0MM];
indices.k_MMM = neighborZ[indices.k_MMM]; indices.k_MMM = neighborZ[indices.k_MMM];
omega_ = hasTurbulentViscosity ? calculateOmega(omega, turbulentViscosity[indices.k_000]) : omega; omega_ = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omega, turbulentViscosity[indices.k_000]) : omega;
read(f_fine, distFine, indices); read(f_fine, distFine, indices);
if (forcesForAllNodesX) if (forcesForAllNodesX)
...@@ -245,7 +246,7 @@ template<bool hasTurbulentViscosity> __device__ void calculateMomentSet( ...@@ -245,7 +246,7 @@ template<bool hasTurbulentViscosity> __device__ void calculateMomentSet(
indices.k_0MM = indices.k_MMM; indices.k_0MM = indices.k_MMM;
indices.k_MMM = neighborX[indices.k_MMM]; indices.k_MMM = neighborX[indices.k_MMM];
omega_ = hasTurbulentViscosity ? calculateOmega(omega, turbulentViscosity[indices.k_000]) : omega; omega_ = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omega, turbulentViscosity[indices.k_000]) : omega;
read(f_fine, distFine, indices); read(f_fine, distFine, indices);
if (forcesForAllNodesX) if (forcesForAllNodesX)
...@@ -266,8 +267,8 @@ template<bool hasTurbulentViscosity> __device__ void calculateMomentSet( ...@@ -266,8 +267,8 @@ template<bool hasTurbulentViscosity> __device__ void calculateMomentSet(
indices.k_M00 = neighborX[k_base_M00]; indices.k_M00 = neighborX[k_base_M00];
indices.k_0M0 = k_base_MM0; indices.k_0M0 = k_base_MM0;
indices.k_MM0 = neighborX[k_base_MM0]; indices.k_MM0 = neighborX[k_base_MM0];
omega_ = hasTurbulentViscosity ? calculateOmega(omega, turbulentViscosity[indices.k_000]) : omega; omega_ = hasTurbulentViscosity ? vf::lbm::calculateOmegaWithturbulentViscosity(omega, turbulentViscosity[indices.k_000]) : omega;
read(f_fine, distFine, indices); read(f_fine, distFine, indices);
if (forcesForAllNodesX) if (forcesForAllNodesX)
......
src/lbm/collision/K17CompressibleNavierStokes.h
+ 7
− 13
View file @ 848f64da
...@@ -79,7 +79,7 @@ namespace vf::lbm ...@@ -79,7 +79,7 @@ namespace vf::lbm
//! ]</b></a> and \ref <a href="https://doi.org/10.1016/j.jcp.2017.07.004"><b>[ M. Geier et al. (2017), //! ]</b></a> and \ref <a href="https://doi.org/10.1016/j.jcp.2017.07.004"><b>[ M. Geier et al. (2017),
//! DOI:10.1016/j.jcp.2017.07.004 ]</b></a> //! DOI:10.1016/j.jcp.2017.07.004 ]</b></a>
//////////////////////////////////////////////////////////////////////////////// ////////////////////////////////////////////////////////////////////////////////
template <TurbulenceModel turbulenceModel, bool writeMacroscopicVariables> template <TurbulenceModel turbulenceModel>
__host__ __device__ void runK17CompressibleNavierStokes(CollisionParameter& parameter, MacroscopicValues& macroscopicValues, TurbulentViscosity& turbulentViscosity) __host__ __device__ void runK17CompressibleNavierStokes(CollisionParameter& parameter, MacroscopicValues& macroscopicValues, TurbulentViscosity& turbulentViscosity)
{ {
auto& distribution = parameter.distribution; auto& distribution = parameter.distribution;
...@@ -237,10 +237,7 @@ __host__ __device__ void runK17CompressibleNavierStokes(CollisionParameter& para ...@@ -237,10 +237,7 @@ __host__ __device__ void runK17CompressibleNavierStokes(CollisionParameter& para
//////////////////////////////////////////////////////////////////////////////////// ////////////////////////////////////////////////////////////////////////////////////
//! - Calculate modified omega with turbulent viscosity //! - Calculate modified omega with turbulent viscosity
//! //!
real omega = parameter.omega; const real omega = turbulenceModel == TurbulenceModel::None ? parameter.omega : vf::lbm::calculateOmegaWithturbulentViscosity(parameter.omega, turbulentViscosity.value);
if (turbulenceModel != TurbulenceModel::None) {
omega /= (c1o1 + c3o1 * parameter.omega * turbulentViscosity.value);
}
//////////////////////////////////////////////////////////// ////////////////////////////////////////////////////////////
// 2. // 2.
real OxxPyyPzz = c1o1; real OxxPyyPzz = c1o1;
...@@ -510,14 +507,6 @@ __host__ __device__ void runK17CompressibleNavierStokes(CollisionParameter& para ...@@ -510,14 +507,6 @@ __host__ __device__ void runK17CompressibleNavierStokes(CollisionParameter& para
m_010 = -m_010; m_010 = -m_010;
m_001 = -m_001; m_001 = -m_001;
// Write to array here to distribute read/write
if (writeMacroscopicVariables || turbulenceModel == TurbulenceModel::AMD) {
macroscopicValues.rho = drho;
macroscopicValues.vx = vvx;
macroscopicValues.vy = vvy;
macroscopicValues.vz = vvz;
}
//////////////////////////////////////////////////////////////////////////////////// ////////////////////////////////////////////////////////////////////////////////////
//! - Chimera transform from central moments to well conditioned distributions as defined in Appendix J in //! - Chimera transform from central moments to well conditioned distributions as defined in Appendix J in
//! <a href="https://doi.org/10.1016/j.camwa.2015.05.001"><b>[ M. Geier et al. (2015), //! <a href="https://doi.org/10.1016/j.camwa.2015.05.001"><b>[ M. Geier et al. (2015),
...@@ -588,6 +577,11 @@ __host__ __device__ void runK17CompressibleNavierStokes(CollisionParameter& para ...@@ -588,6 +577,11 @@ __host__ __device__ void runK17CompressibleNavierStokes(CollisionParameter& para
distribution[DIR_MMP] = f_MMP; distribution[DIR_MMP] = f_MMP;
distribution[DIR_MPM] = f_MPM; distribution[DIR_MPM] = f_MPM;
distribution[DIR_MMM] = f_MMM; distribution[DIR_MMM] = f_MMM;
macroscopicValues.rho = drho;
macroscopicValues.vx = vvx;
macroscopicValues.vy = vvy;
macroscopicValues.vz = vvz;
} }
} // namespace vf::lbm } // namespace vf::lbm
src/lbm/collision/TurbulentViscosity.h
+ 6
− 1
View file @ 848f64da
...@@ -71,7 +71,7 @@ inline __host__ __device__ real calcTurbulentViscositySmagorinsky(real Cs, real ...@@ -71,7 +71,7 @@ inline __host__ __device__ real calcTurbulentViscositySmagorinsky(real Cs, real
} }
template <typename T> template <typename T>
__host__ __device__ int max( T a, T b ) __host__ __device__ T max( T a, T b )
{ {
return ( a > b ) ? a : b; return ( a > b ) ? a : b;
} }
...@@ -87,6 +87,11 @@ inline __host__ __device__ real calcTurbulentViscosityQR(real C, real dxux, real ...@@ -87,6 +87,11 @@ inline __host__ __device__ real calcTurbulentViscosityQR(real C, real dxux, real
return C * max(R, c0o1) / Q; return C * max(R, c0o1) / Q;
} }
inline __host__ __device__ real calculateOmegaWithturbulentViscosity(real omega, real turbulenceViscosity)
{
return omega / (c1o1 + c3o1 * omega * turbulenceViscosity);
}
} // namespace vf::lbm } // namespace vf::lbm
#endif //TURBULENT_VISCOSITY_H #endif //TURBULENT_VISCOSITY_H
utilities/ci-regression-tests/regression-tests-ci.yml.j2
+ 5
− 0
View file @ 848f64da
...@@ -16,6 +16,11 @@ stages: ...@@ -16,6 +16,11 @@ stages:
- chmod +x ./regression-tests/* - chmod +x ./regression-tests/*
- pip install fieldcompare - pip install fieldcompare
artifacts:
expire_in: 1 hrs
paths:
- output/
{% for regression_test in regression_tests %} {% for regression_test in regression_tests %}
run-regression-test-{{ regression_test }}: run-regression-test-{{ regression_test }}:
extends: .regression-test extends: .regression-test
......
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment