From 6f9876ae14cedc4c4fc672c4e000896624ebcf4a Mon Sep 17 00:00:00 2001
From: peters <peters@irmb.tu-bs.de>
Date: Wed, 24 Mar 2021 15:45:27 +0100
Subject: [PATCH] Bugfixing & Refactoring. CumulantChimera is working on the
GPU.
---
.../VirtualFluids_GPU/GPU/Cumulant27chim.cu | 749 +++++++++++++++---
src/lbm/CumulantChimera.h | 155 ++--
2 files changed, 716 insertions(+), 188 deletions(-)
diff --git a/src/gpu/VirtualFluids_GPU/GPU/Cumulant27chim.cu b/src/gpu/VirtualFluids_GPU/GPU/Cumulant27chim.cu
index a8c4b7ce6..49e2c10a2 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/Cumulant27chim.cu
+++ b/src/gpu/VirtualFluids_GPU/GPU/Cumulant27chim.cu
@@ -107,27 +107,30 @@ __device__ Distributions27 getDistributions27(real* distributions, unsigned int
return dist;
}
-struct ABC
+struct DistributionWrapper
{
- __device__ ABC(uint k,
- uint kw,
- uint ks,
- uint kb,
- uint ksw,
- uint kbw,
- uint kbs,
- uint kbsw) :
+ __device__ DistributionWrapper(
+ real* distributions,
+ unsigned int size_Mat,
+ bool isEvenTimestep,
+ uint k,
+ uint* neighborX,
+ uint* neighborY,
+ uint* neighborZ) :
+ dist(getDistributions27(distributions, size_Mat, isEvenTimestep)),
k(k),
- kw (kw),
- ks (ks),
- kb (kb),
- ksw (ksw),
- kbw (kbw),
- kbs (kbs),
- kbsw (kbsw)
- {}
-
- __device__ void read(const Distributions27& dist)
+ kw (neighborX[k]),
+ ks (neighborY[k]),
+ kb (neighborZ[k]),
+ ksw (neighborY[kw]),
+ kbw (neighborZ[kw]),
+ kbs (neighborZ[ks]),
+ kbsw(neighborZ[ksw])
+ {
+ read();
+ }
+
+ __device__ void read()
{
distribution.f[VF::LBM::DIR::PZZ] = (dist.f[dirE ])[k];
distribution.f[VF::LBM::DIR::MZZ] = (dist.f[dirW ])[kw];
@@ -152,13 +155,13 @@ struct ABC
distribution.f[VF::LBM::DIR::PMP] = (dist.f[dirTSE ])[ks];
distribution.f[VF::LBM::DIR::MMP] = (dist.f[dirTSW ])[ksw];
distribution.f[VF::LBM::DIR::PPM] = (dist.f[dirBNE ])[kb];
- distribution.f[VF::LBM::DIR::MPM] = (dist.f[dirBNE ])[kb];
+ distribution.f[VF::LBM::DIR::MPM] = (dist.f[dirBNW ])[kbw];
distribution.f[VF::LBM::DIR::PMM] = (dist.f[dirBSE ])[kbs];
distribution.f[VF::LBM::DIR::MMM] = (dist.f[dirBSW ])[kbsw];
distribution.f[VF::LBM::DIR::ZZZ] = (dist.f[dirREST])[k];
}
- __device__ void write(Distributions27& dist)
+ __device__ void write()
{
(dist.f[dirE ])[k] = distribution.f[VF::LBM::DIR::PZZ];
(dist.f[dirW ])[kw] = distribution.f[VF::LBM::DIR::MZZ];
@@ -183,12 +186,14 @@ struct ABC
(dist.f[dirTSE ])[ks] = distribution.f[VF::LBM::DIR::PMP];
(dist.f[dirTSW ])[ksw] = distribution.f[VF::LBM::DIR::MMP];
(dist.f[dirBNE ])[kb] = distribution.f[VF::LBM::DIR::PPM];
- (dist.f[dirBNE ])[kb] = distribution.f[VF::LBM::DIR::MPM];
+ (dist.f[dirBNW ])[kbw] = distribution.f[VF::LBM::DIR::MPM];
(dist.f[dirBSE ])[kbs] = distribution.f[VF::LBM::DIR::PMM];
(dist.f[dirBSW ])[kbsw] = distribution.f[VF::LBM::DIR::MMM];
(dist.f[dirREST])[k] = distribution.f[VF::LBM::DIR::ZZZ];
}
+ Distributions27 dist;
+
VF::LBM::Distribution27 distribution;
const uint k;
@@ -201,7 +206,39 @@ struct ABC
const uint kbsw;
};
-////////////////////////////////////////////////////////////////////////////////
+__device__ unsigned int getNodeIndex()
+{
+ const unsigned x = threadIdx.x;
+ const unsigned y = blockIdx.x;
+ const unsigned z = blockIdx.y;
+
+ const unsigned nx = blockDim.x;
+ const unsigned ny = gridDim.x;
+
+ return nx*(ny*z + y) + x;
+}
+
+__device__ bool isValidFluidNode(uint k, int size_Mat, uint nodeType)
+{
+ return (k < size_Mat) && (nodeType == GEO_FLUID);
+}
+
+__device__ void getLevelForce(real fx, real fy, real fz, int level, real* forces)
+{
+ real fx_t {1.}, fy_t {1.}, fz_t {1.};
+ for (int i = 0; i < level; i++)
+ {
+ fx_t *= c2o1;
+ fy_t *= c2o1;
+ fz_t *= c2o1;
+ }
+
+ forces[0] = fx / fx_t;
+ forces[1] = fy / fy_t;
+ forces[2] = fz / fz_t;
+}
+
+
extern "C" __global__ void Cumulant_K17_LBM_Device_Kernel(
real omega,
uint* typeOfGridNode,
@@ -213,6 +250,39 @@ extern "C" __global__ void Cumulant_K17_LBM_Device_Kernel(
int level,
real* forces,
bool isEvenTimestep)
+{
+ const uint k = getNodeIndex();
+ const uint nodeType = typeOfGridNode[k];
+
+ if (isValidFluidNode(k, size_Mat, nodeType))
+ {
+ DistributionWrapper distributionWrapper {
+ distributions, size_Mat, isEvenTimestep, k, neighborX, neighborY, neighborZ
+ };
+
+ real level_forces[3];
+ getLevelForce(forces[0], forces[1], forces[2], level, level_forces);
+
+ VF::LBM::cumulantChimera(distributionWrapper.distribution, omega, level_forces);
+
+ distributionWrapper.write();
+ }
+}
+
+
+
+////////////////////////////////////////////////////////////////////////////////
+extern "C" __global__ void Cumulant_K17_LBM_Device_Kernel_old(
+ real omega,
+ uint* typeOfGridNode,
+ uint* neighborX,
+ uint* neighborY,
+ uint* neighborZ,
+ real* distributions,
+ int size_Mat,
+ int level,
+ real* forces,
+ bool isEvenTimestep)
{
//////////////////////////////////////////////////////////////////////////
//! Cumulant K17 Kernel is based on \ref
@@ -244,112 +314,555 @@ extern "C" __global__ void Cumulant_K17_LBM_Device_Kernel(
//!
Distributions27 dist = getDistributions27(distributions, size_Mat, isEvenTimestep);
- // ////////////////////////////////////////////////////////////////////////////////
- // //! - Set neighbor indices (necessary for indirect addressing)
- const uint kw = neighborX[k];
- const uint ks = neighborY[k];
- const uint kb = neighborZ[k];
- const uint ksw = neighborY[kw];
- const uint kbw = neighborZ[kw];
- const uint kbs = neighborZ[ks];
- const uint kbsw = neighborZ[ksw];
-
- ABC abc = {
- k,
- kw ,
- ks ,
- kb ,
- ksw ,
- kbw ,
- kbs ,
- kbsw
- };
+ ////////////////////////////////////////////////////////////////////////////////
+ //! - Set neighbor indices (necessary for indirect addressing)
+ uint kw = neighborX[k];
+ uint ks = neighborY[k];
+ uint kb = neighborZ[k];
+ uint ksw = neighborY[kw];
+ uint kbw = neighborZ[kw];
+ uint kbs = neighborZ[ks];
+ uint kbsw = neighborZ[ksw];
+ ////////////////////////////////////////////////////////////////////////////////////
+ //! - Set local distributions
+ //!
+ // real mfcbb = (dist.f[dirE ])[k];
+ // real mfabb = (dist.f[dirW ])[kw];
+ // real mfbcb = (dist.f[dirN ])[k];
+ // real mfbab = (dist.f[dirS ])[ks];
+ // real mfbbc = (dist.f[dirT ])[k];
+ // real mfbba = (dist.f[dirB ])[kb];
+ // real mfccb = (dist.f[dirNE ])[k];
+ // real mfaab = (dist.f[dirSW ])[ksw];
+ // real mfcab = (dist.f[dirSE ])[ks];
+ // real mfacb = (dist.f[dirNW ])[kw];
+ // real mfcbc = (dist.f[dirTE ])[k];
+ // real mfaba = (dist.f[dirBW ])[kbw];
+ // real mfcba = (dist.f[dirBE ])[kb];
+ // real mfabc = (dist.f[dirTW ])[kw];
+ // real mfbcc = (dist.f[dirTN ])[k];
+ // real mfbaa = (dist.f[dirBS ])[kbs];
+ // real mfbca = (dist.f[dirBN ])[kb];
+ // real mfbac = (dist.f[dirTS ])[ks];
+ // real mfbbb = (dist.f[dirREST])[k];
+ // real mfccc = (dist.f[dirTNE ])[k];
+ // real mfaac = (dist.f[dirTSW ])[ksw];
+ // real mfcac = (dist.f[dirTSE ])[ks];
+ // real mfacc = (dist.f[dirTNW ])[kw];
+ // real mfcca = (dist.f[dirBNE ])[kb];
+ // real mfaaa = (dist.f[dirBSW ])[kbsw];
+ // real mfcaa = (dist.f[dirBSE ])[kbs];
+ // real mfaca = (dist.f[dirBNW ])[kbw];
+
+ VF::LBM::Distribution27 distribution_in;
+
+ distribution_in.f[VF::LBM::DIR::PZZ] = (dist.f[dirE ])[k];
+ distribution_in.f[VF::LBM::DIR::MZZ] = (dist.f[dirW ])[kw];
+ distribution_in.f[VF::LBM::DIR::ZPZ] = (dist.f[dirN ])[k];
+ distribution_in.f[VF::LBM::DIR::ZMZ] = (dist.f[dirS ])[ks];
+ distribution_in.f[VF::LBM::DIR::ZZP] = (dist.f[dirT ])[k];
+ distribution_in.f[VF::LBM::DIR::ZZM] = (dist.f[dirB ])[kb];
+ distribution_in.f[VF::LBM::DIR::PPZ] = (dist.f[dirNE ])[k];
+ distribution_in.f[VF::LBM::DIR::MMZ] = (dist.f[dirSW ])[ksw];
+ distribution_in.f[VF::LBM::DIR::PMZ] = (dist.f[dirSE ])[ks];
+ distribution_in.f[VF::LBM::DIR::MPZ] = (dist.f[dirNW ])[kw];
+ distribution_in.f[VF::LBM::DIR::PZP] = (dist.f[dirTE ])[k];
+ distribution_in.f[VF::LBM::DIR::MZM] = (dist.f[dirBW ])[kbw];
+ distribution_in.f[VF::LBM::DIR::PZM] = (dist.f[dirBE ])[kb];
+ distribution_in.f[VF::LBM::DIR::MZP] = (dist.f[dirTW ])[kw];
+ distribution_in.f[VF::LBM::DIR::ZPP] = (dist.f[dirTN ])[k];
+ distribution_in.f[VF::LBM::DIR::ZMM] = (dist.f[dirBS ])[kbs];
+ distribution_in.f[VF::LBM::DIR::ZPM] = (dist.f[dirBN ])[kb];
+ distribution_in.f[VF::LBM::DIR::ZMP] = (dist.f[dirTS ])[ks];
+ distribution_in.f[VF::LBM::DIR::PPP] = (dist.f[dirTNE ])[k];
+ distribution_in.f[VF::LBM::DIR::MPP] = (dist.f[dirTNW ])[kw];
+ distribution_in.f[VF::LBM::DIR::PMP] = (dist.f[dirTSE ])[ks];
+ distribution_in.f[VF::LBM::DIR::MMP] = (dist.f[dirTSW ])[ksw];
+ distribution_in.f[VF::LBM::DIR::PPM] = (dist.f[dirBNE ])[kb];
+ distribution_in.f[VF::LBM::DIR::MPM] = (dist.f[dirBNW ])[kbw];
+ distribution_in.f[VF::LBM::DIR::PMM] = (dist.f[dirBSE ])[kbs];
+ distribution_in.f[VF::LBM::DIR::MMM] = (dist.f[dirBSW ])[kbsw];
+ distribution_in.f[VF::LBM::DIR::ZZZ] = (dist.f[dirREST])[k];
+
+ real mfcbb = distribution_in.f[VF::LBM::DIR::PZZ];
+ real mfabb = distribution_in.f[VF::LBM::DIR::MZZ];
+ real mfbcb = distribution_in.f[VF::LBM::DIR::ZPZ];
+ real mfbab = distribution_in.f[VF::LBM::DIR::ZMZ];
+ real mfbbc = distribution_in.f[VF::LBM::DIR::ZZP];
+ real mfbba = distribution_in.f[VF::LBM::DIR::ZZM];
+ real mfccb = distribution_in.f[VF::LBM::DIR::PPZ];
+ real mfaab = distribution_in.f[VF::LBM::DIR::MMZ];
+ real mfcab = distribution_in.f[VF::LBM::DIR::PMZ];
+ real mfacb = distribution_in.f[VF::LBM::DIR::MPZ];
+ real mfcbc = distribution_in.f[VF::LBM::DIR::PZP];
+ real mfaba = distribution_in.f[VF::LBM::DIR::MZM];
+ real mfcba = distribution_in.f[VF::LBM::DIR::PZM];
+ real mfabc = distribution_in.f[VF::LBM::DIR::MZP];
+ real mfbcc = distribution_in.f[VF::LBM::DIR::ZPP];
+ real mfbaa = distribution_in.f[VF::LBM::DIR::ZMM];
+ real mfbca = distribution_in.f[VF::LBM::DIR::ZPM];
+ real mfbac = distribution_in.f[VF::LBM::DIR::ZMP];
+ real mfccc = distribution_in.f[VF::LBM::DIR::PPP];
+ real mfacc = distribution_in.f[VF::LBM::DIR::MPP];
+ real mfcac = distribution_in.f[VF::LBM::DIR::PMP];
+ real mfaac = distribution_in.f[VF::LBM::DIR::MMP];
+ real mfcca = distribution_in.f[VF::LBM::DIR::PPM];
+ real mfaca = distribution_in.f[VF::LBM::DIR::MPM];
+ real mfcaa = distribution_in.f[VF::LBM::DIR::PMM];
+ real mfaaa = distribution_in.f[VF::LBM::DIR::MMM];
+ real mfbbb = distribution_in.f[VF::LBM::DIR::ZZZ];
+
+
+ ////////////////////////////////////////////////////////////////////////////////////
+ //! - Calculate density and velocity using pyramid summation for low round-off errors as in Eq. (J1)-(J3) \ref
+ //! <a href="https://doi.org/10.1016/j.camwa.2015.05.001"><b>[ M. Geier et al. (2015), DOI:10.1016/j.camwa 2015.05.001 ]</b></a>
+ //!
+ real drho =
+ ((((mfccc + mfaaa) + (mfaca + mfcac)) + ((mfacc + mfcaa) + (mfaac + mfcca))) +
+ (((mfbac + mfbca) + (mfbaa + mfbcc)) + ((mfabc + mfcba) + (mfaba + mfcbc)) + ((mfacb + mfcab) + (mfaab + mfccb))) +
+ ((mfabb + mfcbb) + (mfbab + mfbcb) + (mfbba + mfbbc))) + mfbbb;
+ real rho = c1o1 + drho;
+ real OOrho = c1o1 / rho;
+ real vvx =
+ ((((mfccc - mfaaa) + (mfcac - mfaca)) + ((mfcaa - mfacc) + (mfcca - mfaac))) +
+ (((mfcba - mfabc) + (mfcbc - mfaba)) + ((mfcab - mfacb) + (mfccb - mfaab))) +
+ (mfcbb - mfabb)) * OOrho;
+ real vvy =
+ ((((mfccc - mfaaa) + (mfaca - mfcac)) + ((mfacc - mfcaa) + (mfcca - mfaac))) +
+ (((mfbca - mfbac) + (mfbcc - mfbaa)) + ((mfacb - mfcab) + (mfccb - mfaab))) +
+ (mfbcb - mfbab)) * OOrho;
+ real vvz =
+ ((((mfccc - mfaaa) + (mfcac - mfaca)) + ((mfacc - mfcaa) + (mfaac - mfcca))) +
+ (((mfbac - mfbca) + (mfbcc - mfbaa)) + ((mfabc - mfcba) + (mfcbc - mfaba))) +
+ (mfbbc - mfbba)) * OOrho;
+ ////////////////////////////////////////////////////////////////////////////////////
+ //! - Add half of the acceleration (body force) to the velocity as in Eq. (42) \ref
+ //! <a href="https://doi.org/10.1016/j.camwa.2015.05.001"><b>[ M. Geier et al. (2015), DOI:10.1016/j.camwa 2015.05.001 ]</b></a>
+ //!
+ real fx = forces[0];
+ real fy = forces[1];
+ real fz = forces[2];
+
+ real fx_t {1.}, fy_t {1.}, fz_t {1.};
+ for (int i = 0; i < level; i++)
+ {
+ fx_t *= c2o1;
+ fy_t *= c2o1;
+ fz_t *= c2o1;
+ }
+
+ fx /= fx_t;
+ fy /= fy_t;
+ fz /= fz_t;
+ //real forces[3] {fx, fy, fz};
- // ////////////////////////////////////////////////////////////////////////////////////
- // //! - Set local distributions
- // //!
- // VF::LBM::Distribution27 distribution_in;
-
- // distribution_in.f[VF::LBM::DIR::PZZ] = (dist.f[dirE ])[k];
- // distribution_in.f[VF::LBM::DIR::MZZ] = (dist.f[dirW ])[kw];
- // distribution_in.f[VF::LBM::DIR::ZPZ] = (dist.f[dirN ])[k];
- // distribution_in.f[VF::LBM::DIR::ZMZ] = (dist.f[dirS ])[ks];
- // distribution_in.f[VF::LBM::DIR::ZZP] = (dist.f[dirT ])[k];
- // distribution_in.f[VF::LBM::DIR::ZZM] = (dist.f[dirB ])[kb];
- // distribution_in.f[VF::LBM::DIR::PPZ] = (dist.f[dirNE ])[k];
- // distribution_in.f[VF::LBM::DIR::MMZ] = (dist.f[dirSW ])[ksw];
- // distribution_in.f[VF::LBM::DIR::PMZ] = (dist.f[dirSE ])[ks];
- // distribution_in.f[VF::LBM::DIR::MPZ] = (dist.f[dirNW ])[kw];
- // distribution_in.f[VF::LBM::DIR::PZP] = (dist.f[dirTE ])[k];
- // distribution_in.f[VF::LBM::DIR::MZM] = (dist.f[dirBW ])[kbw];
- // distribution_in.f[VF::LBM::DIR::PZM] = (dist.f[dirBE ])[kb];
- // distribution_in.f[VF::LBM::DIR::MZP] = (dist.f[dirTW ])[kw];
- // distribution_in.f[VF::LBM::DIR::ZPP] = (dist.f[dirTN ])[k];
- // distribution_in.f[VF::LBM::DIR::ZMM] = (dist.f[dirBS ])[kbs];
- // distribution_in.f[VF::LBM::DIR::ZPM] = (dist.f[dirBN ])[kb];
- // distribution_in.f[VF::LBM::DIR::ZMP] = (dist.f[dirTS ])[ks];
- // distribution_in.f[VF::LBM::DIR::PPP] = (dist.f[dirTNE ])[k];
- // distribution_in.f[VF::LBM::DIR::MPP] = (dist.f[dirTNW ])[kw];
- // distribution_in.f[VF::LBM::DIR::PMP] = (dist.f[dirTSE ])[ks];
- // distribution_in.f[VF::LBM::DIR::MMP] = (dist.f[dirTSW ])[ksw];
- // distribution_in.f[VF::LBM::DIR::PPM] = (dist.f[dirBNE ])[kb];
- // distribution_in.f[VF::LBM::DIR::MPM] = (dist.f[dirBNE ])[kb];
- // distribution_in.f[VF::LBM::DIR::PMM] = (dist.f[dirBSE ])[kbs];
- // distribution_in.f[VF::LBM::DIR::MMM] = (dist.f[dirBSW ])[kbsw];
- // distribution_in.f[VF::LBM::DIR::ZZZ] = (dist.f[dirREST])[k];
-
- abc.read(dist);
-
- real vx = forces[0] / pow((double)c2o1, (double)level);
- real vy = forces[1] / pow((double)c2o1, (double)level);
- real vz = forces[2] / pow((double)c2o1, (double)level);
- real forces[3] {vx, vy, vz};
-
- // VF::LBM::Distribution27 distribution_out = VF::LBM::cumulantChimera(distribution_in, omega, forces);
-
- abc.distribution = VF::LBM::cumulantChimera(abc.distribution, omega, forces);
-
- abc.write(dist);
+ vvx += fx * c1o2;
+ vvy += fy * c1o2;
+ vvz += fz * c1o2;
+ ////////////////////////////////////////////////////////////////////////////////////
+ // calculate the square of velocities for this lattice node
+ real vx2 = vvx*vvx;
+ real vy2 = vvy*vvy;
+ real vz2 = vvz*vvz;
+ ////////////////////////////////////////////////////////////////////////////////////
+ //! - Set relaxation limiters for third order cumulants to default value \f$ \lambda=0.001 \f$ according to section 6 in \ref
+ //! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), DOI:10.1016/j.jcp.2017.05 040 ]</b></a>
+ //!
+ real wadjust;
+ real qudricLimitP = c1o100;
+ real qudricLimitM = c1o100;
+ real qudricLimitD = c1o100;
+ ////////////////////////////////////////////////////////////////////////////////////
+ //! - Chimera transform from well conditioned distributions to central moments as defined in Appendix J in \ref
+ //! <a href="https://doi.org/10.1016/j.camwa.2015.05.001"><b>[ M. Geier et al. (2015), DOI:10.1016/j.camwa 2015.05.001 ]</b></a>
+ //! see also Eq. (6)-(14) in \ref
+ //! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), DOI:10.1016/j.jcp.2017.05 040 ]</b></a>
+ //!
+ ////////////////////////////////////////////////////////////////////////////////////
+ // Z - Dir
+ VF::LBM::forwardInverseChimeraWithK(mfaaa, mfaab, mfaac, vvz, vz2, c36o1, c1o36);
+ VF::LBM::forwardInverseChimeraWithK(mfaba, mfabb, mfabc, vvz, vz2, c9o1, c1o9);
+ VF::LBM::forwardInverseChimeraWithK(mfaca, mfacb, mfacc, vvz, vz2, c36o1, c1o36);
+ VF::LBM::forwardInverseChimeraWithK(mfbaa, mfbab, mfbac, vvz, vz2, c9o1, c1o9);
+ VF::LBM::forwardInverseChimeraWithK(mfbba, mfbbb, mfbbc, vvz, vz2, c9o4, c4o9);
+ VF::LBM::forwardInverseChimeraWithK(mfbca, mfbcb, mfbcc, vvz, vz2, c9o1, c1o9);
+ VF::LBM::forwardInverseChimeraWithK(mfcaa, mfcab, mfcac, vvz, vz2, c36o1, c1o36);
+ VF::LBM::forwardInverseChimeraWithK(mfcba, mfcbb, mfcbc, vvz, vz2, c9o1, c1o9);
+ VF::LBM::forwardInverseChimeraWithK(mfcca, mfccb, mfccc, vvz, vz2, c36o1, c1o36);
+ ////////////////////////////////////////////////////////////////////////////////////
+ // Y - Dir
+ VF::LBM::forwardInverseChimeraWithK(mfaaa, mfaba, mfaca, vvy, vy2, c6o1, c1o6);
+ VF::LBM::forwardChimera( mfaab, mfabb, mfacb, vvy, vy2);
+ VF::LBM::forwardInverseChimeraWithK(mfaac, mfabc, mfacc, vvy, vy2, c18o1, c1o18);
+ VF::LBM::forwardInverseChimeraWithK(mfbaa, mfbba, mfbca, vvy, vy2, c3o2, c2o3);
+ VF::LBM::forwardChimera( mfbab, mfbbb, mfbcb, vvy, vy2);
+ VF::LBM::forwardInverseChimeraWithK(mfbac, mfbbc, mfbcc, vvy, vy2, c9o2, c2o9);
+ VF::LBM::forwardInverseChimeraWithK(mfcaa, mfcba, mfcca, vvy, vy2, c6o1, c1o6);
+ VF::LBM::forwardChimera( mfcab, mfcbb, mfccb, vvy, vy2);
+ VF::LBM::forwardInverseChimeraWithK(mfcac, mfcbc, mfccc, vvy, vy2, c18o1, c1o18);
+ ////////////////////////////////////////////////////////////////////////////////////
+ // X - Dir
+ VF::LBM::forwardInverseChimeraWithK(mfaaa, mfbaa, mfcaa, vvx, vx2, c1o1, c1o1);
+ VF::LBM::forwardChimera( mfaba, mfbba, mfcba, vvx, vx2);
+ VF::LBM::forwardInverseChimeraWithK(mfaca, mfbca, mfcca, vvx, vx2, c3o1, c1o3);
+ VF::LBM::forwardChimera( mfaab, mfbab, mfcab, vvx, vx2);
+ VF::LBM::forwardChimera( mfabb, mfbbb, mfcbb, vvx, vx2);
+ VF::LBM::forwardChimera( mfacb, mfbcb, mfccb, vvx, vx2);
+ VF::LBM::forwardInverseChimeraWithK(mfaac, mfbac, mfcac, vvx, vx2, c3o1, c1o3);
+ VF::LBM::forwardChimera( mfabc, mfbbc, mfcbc, vvx, vx2);
+ VF::LBM::forwardInverseChimeraWithK(mfacc, mfbcc, mfccc, vvx, vx2, c3o1, c1o9);
+ ////////////////////////////////////////////////////////////////////////////////////
+ //! - Setting relaxation rates for non-hydrodynamic cumulants (default values). Variable names and equations according to
+ //! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), DOI:10.1016/j.jcp.2017.05 040 ]</b></a>
+ //! => [NAME IN PAPER]=[NAME IN CODE]=[DEFAULT VALUE].
+ //! - Trace of second order cumulants \f$ C_{200}+C_{020}+C_{002} \f$ used to adjust bulk viscosity:\f$\omega_2=OxxPyyPzz=1.0 \f$.
+ //! - Third order cumulants \f$ C_{120}+C_{102}, C_{210}+C_{012}, C_{201}+C_{021} \f$: \f$ \omega_3=OxyyPxzz \f$ set according to Eq. (111) with simplifications assuming \f$ \omega_2=1.0\f$.
+ //! - Third order cumulants \f$ C_{120}-C_{102}, C_{210}-C_{012}, C_{201}-C_{021} \f$: \f$ \omega_4 = OxyyMxzz \f$ set according to Eq. (112) with simplifications assuming \f$ \omega_2 = 1.0\f$.
+ //! - Third order cumulants \f$ C_{111} \f$: \f$ \omega_5 = Oxyz \f$ set according to Eq. (113) with simplifications assuming \f$ \omega_2 = 1.0\f$ (modify for different bulk viscosity).
+ //! - Fourth order cumulants \f$ C_{220}, C_{202}, C_{022}, C_{211}, C_{121}, C_{112} \f$: for simplification all set to the same default value \f$ \omega_6=\omega_7=\omega_8=O4=1.0 \f$.
+ //! - Fifth order cumulants \f$ C_{221}, C_{212}, C_{122}\f$: \f$\omega_9=O5=1.0\f$.
+ //! - Sixth order cumulant \f$ C_{222}\f$: \f$\omega_{10}=O6=1.0\f$.
+ //!
+ ////////////////////////////////////////////////////////////
+ //2.
+ real OxxPyyPzz = c1o1;
+ ////////////////////////////////////////////////////////////
+ //3.
+ real OxyyPxzz = c8o1 * (-c2o1 + omega) * ( c1o1 + c2o1*omega) / (-c8o1 - c14o1*omega + c7o1*omega*omega);
+ real OxyyMxzz = c8o1 * (-c2o1 + omega) * (-c7o1 + c4o1*omega) / (c56o1 - c50o1*omega + c9o1*omega*omega);
+ real Oxyz = c24o1 * (-c2o1 + omega) * (-c2o1 - c7o1*omega + c3o1*omega*omega) / (c48o1 + c152o1*omega - c130o1*omega*omega + c29o1*omega*omega*omega);
+ ////////////////////////////////////////////////////////////
+ //4.
+ real O4 = c1o1;
+ ////////////////////////////////////////////////////////////
+ //5.
+ real O5 = c1o1;
+ ////////////////////////////////////////////////////////////
+ //6.
+ real O6 = c1o1;
+ ////////////////////////////////////////////////////////////////////////////////////
+ //! - A and B: parameters for fourth order convergence of the diffusion term according to Eq. (114) and (115)
+ //! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), DOI:10.1016/j.jcp.2017.05 040 ]</b></a>
+ //! with simplifications assuming \f$ \omega_2 = 1.0 \f$ (modify for different bulk viscosity).
+ //!
+ real A = (c4o1 + c2o1*omega - c3o1*omega*omega) / (c2o1 - c7o1*omega + c5o1*omega*omega);
+ real B = (c4o1 + c28o1*omega - c14o1*omega*omega) / (c6o1 - c21o1*omega + c15o1*omega*omega);
+ ////////////////////////////////////////////////////////////////////////////////////
+ //! - Compute cumulants from central moments according to Eq. (20)-(23) in
+ //! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), DOI:10.1016/j.jcp.2017.05 040 ]</b></a>
+ //!
+ ////////////////////////////////////////////////////////////
+ //4.
+ real CUMcbb = mfcbb - ((mfcaa + c1o3) * mfabb + c2o1 * mfbba * mfbab) * OOrho;
+ real CUMbcb = mfbcb - ((mfaca + c1o3) * mfbab + c2o1 * mfbba * mfabb) * OOrho;
+ real CUMbbc = mfbbc - ((mfaac + c1o3) * mfbba + c2o1 * mfbab * mfabb) * OOrho;
+ real CUMcca = mfcca - (((mfcaa * mfaca + c2o1 * mfbba * mfbba) + c1o3 * (mfcaa + mfaca)) * OOrho - c1o9*(drho * OOrho));
+ real CUMcac = mfcac - (((mfcaa * mfaac + c2o1 * mfbab * mfbab) + c1o3 * (mfcaa + mfaac)) * OOrho - c1o9*(drho * OOrho));
+ real CUMacc = mfacc - (((mfaac * mfaca + c2o1 * mfabb * mfabb) + c1o3 * (mfaac + mfaca)) * OOrho - c1o9*(drho * OOrho));
+ ////////////////////////////////////////////////////////////
+ //5.
+ real CUMbcc = mfbcc - ((mfaac * mfbca + mfaca * mfbac + c4o1 * mfabb * mfbbb + c2o1 * (mfbab * mfacb + mfbba * mfabc)) + c1o3 * (mfbca + mfbac)) * OOrho;
+ real CUMcbc = mfcbc - ((mfaac * mfcba + mfcaa * mfabc + c4o1 * mfbab * mfbbb + c2o1 * (mfabb * mfcab + mfbba * mfbac)) + c1o3 * (mfcba + mfabc)) * OOrho;
+ real CUMccb = mfccb - ((mfcaa * mfacb + mfaca * mfcab + c4o1 * mfbba * mfbbb + c2o1 * (mfbab * mfbca + mfabb * mfcba)) + c1o3 * (mfacb + mfcab)) * OOrho;
+ ////////////////////////////////////////////////////////////
+ //6.
+ real CUMccc = mfccc + ((-c4o1 * mfbbb * mfbbb
+ - (mfcaa * mfacc + mfaca * mfcac + mfaac * mfcca)
+ - c4o1 * (mfabb * mfcbb + mfbab * mfbcb + mfbba * mfbbc)
+ - c2o1 * (mfbca * mfbac + mfcba * mfabc + mfcab * mfacb)) * OOrho
+ + (c4o1 * (mfbab * mfbab * mfaca + mfabb * mfabb * mfcaa + mfbba * mfbba * mfaac)
+ + c2o1 * (mfcaa * mfaca * mfaac)
+ + c16o1 * mfbba * mfbab * mfabb) * OOrho * OOrho
+ - c1o3 * (mfacc + mfcac + mfcca) * OOrho
+ - c1o9 * (mfcaa + mfaca + mfaac) * OOrho
+ + (c2o1 * (mfbab * mfbab + mfabb * mfabb + mfbba * mfbba)
+ + (mfaac * mfaca + mfaac * mfcaa + mfaca * mfcaa) + c1o3 *(mfaac + mfaca + mfcaa)) * OOrho * OOrho * c2o3
+ + c1o27*((drho * drho - drho) * OOrho * OOrho));
+ ////////////////////////////////////////////////////////////////////////////////////
+ //! - Compute linear combinations of second and third order cumulants
+ //!
+ ////////////////////////////////////////////////////////////
+ //2.
+ real mxxPyyPzz = mfcaa + mfaca + mfaac;
+ real mxxMyy = mfcaa - mfaca;
+ real mxxMzz = mfcaa - mfaac;
+ ////////////////////////////////////////////////////////////
+ //3.
+ real mxxyPyzz = mfcba + mfabc;
+ real mxxyMyzz = mfcba - mfabc;
+ real mxxzPyyz = mfcab + mfacb;
+ real mxxzMyyz = mfcab - mfacb;
+ real mxyyPxzz = mfbca + mfbac;
+ real mxyyMxzz = mfbca - mfbac;
+ ////////////////////////////////////////////////////////////////////////////////////
+ //incl. correction
+ ////////////////////////////////////////////////////////////
+ //! - Compute velocity gradients from second order cumulants according to Eq. (27)-(32)
+ //! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), DOI:10.1016/j.jcp.2017.05 040 ]</b></a>
+ //! Further explanations of the correction in viscosity in Appendix H of
+ //! <a href="https://doi.org/10.1016/j.camwa.2015.05.001"><b>[ M. Geier et al. (2015), DOI:10.1016/j.camwa 2015.05.001 ]</b></a>
+ //! Note that the division by rho is omitted here as we need rho times the gradients later.
+ //!
+ real Dxy = -c3o1*omega*mfbba;
+ real Dxz = -c3o1*omega*mfbab;
+ real Dyz = -c3o1*omega*mfabb;
+ real dxux = c1o2 * (-omega) *(mxxMyy + mxxMzz) + c1o2 * OxxPyyPzz * (mfaaa - mxxPyyPzz);
+ real dyuy = dxux + omega * c3o2 * mxxMyy;
+ real dzuz = dxux + omega * c3o2 * mxxMzz;
+ ////////////////////////////////////////////////////////////
+ //! - Relaxation of second order cumulants with correction terms according to Eq. (33)-(35) in
+ //! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), DOI:10.1016/j.jcp.2017.05 040 ]</b></a>
+ //!
+ mxxPyyPzz += OxxPyyPzz*(mfaaa - mxxPyyPzz) - c3o1 * (c1o1 - c1o2 * OxxPyyPzz) * (vx2 * dxux + vy2 * dyuy + vz2 * dzuz);
+ mxxMyy += omega * (-mxxMyy) - c3o1 * (c1o1 + c1o2 * (-omega)) * (vx2 * dxux - vy2 * dyuy);
+ mxxMzz += omega * (-mxxMzz) - c3o1 * (c1o1 + c1o2 * (-omega)) * (vx2 * dxux - vz2 * dzuz);
+ ////////////////////////////////////////////////////////////////////////////////////
+ ////no correction
+ //mxxPyyPzz += OxxPyyPzz*(mfaaa - mxxPyyPzz);
+ //mxxMyy += -(-omega) * (-mxxMyy);
+ //mxxMzz += -(-omega) * (-mxxMzz);
+ //////////////////////////////////////////////////////////////////////////
+ mfabb += omega * (-mfabb);
+ mfbab += omega * (-mfbab);
+ mfbba += omega * (-mfbba);
+ ////////////////////////////////////////////////////////////////////////////////////
+ //relax
+ //////////////////////////////////////////////////////////////////////////
+ // incl. limiter
+ //! - Relaxation of third order cumulants including limiter according to Eq. (116)-(123)
+ //! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), DOI:10.1016/j.jcp.2017.05 040 ]</b></a>
+ //!
+ wadjust = Oxyz + (c1o1 - Oxyz)*abs(mfbbb) / (abs(mfbbb) + qudricLimitD);
+ mfbbb += wadjust * (-mfbbb);
+ wadjust = OxyyPxzz + (c1o1 - OxyyPxzz)*abs(mxxyPyzz) / (abs(mxxyPyzz) + qudricLimitP);
+ mxxyPyzz += wadjust * (-mxxyPyzz);
+ wadjust = OxyyMxzz + (c1o1 - OxyyMxzz)*abs(mxxyMyzz) / (abs(mxxyMyzz) + qudricLimitM);
+ mxxyMyzz += wadjust * (-mxxyMyzz);
+ wadjust = OxyyPxzz + (c1o1 - OxyyPxzz)*abs(mxxzPyyz) / (abs(mxxzPyyz) + qudricLimitP);
+ mxxzPyyz += wadjust * (-mxxzPyyz);
+ wadjust = OxyyMxzz + (c1o1 - OxyyMxzz)*abs(mxxzMyyz) / (abs(mxxzMyyz) + qudricLimitM);
+ mxxzMyyz += wadjust * (-mxxzMyyz);
+ wadjust = OxyyPxzz + (c1o1 - OxyyPxzz)*abs(mxyyPxzz) / (abs(mxyyPxzz) + qudricLimitP);
+ mxyyPxzz += wadjust * (-mxyyPxzz);
+ wadjust = OxyyMxzz + (c1o1 - OxyyMxzz)*abs(mxyyMxzz) / (abs(mxyyMxzz) + qudricLimitM);
+ mxyyMxzz += wadjust * (-mxyyMxzz);
+ //////////////////////////////////////////////////////////////////////////
+ // no limiter
+ //mfbbb += OxyyMxzz * (-mfbbb);
+ //mxxyPyzz += OxyyPxzz * (-mxxyPyzz);
+ //mxxyMyzz += OxyyMxzz * (-mxxyMyzz);
+ //mxxzPyyz += OxyyPxzz * (-mxxzPyyz);
+ //mxxzMyyz += OxyyMxzz * (-mxxzMyyz);
+ //mxyyPxzz += OxyyPxzz * (-mxyyPxzz);
+ //mxyyMxzz += OxyyMxzz * (-mxyyMxzz);
+ ////////////////////////////////////////////////////////////////////////////////////
+ //! - Compute inverse linear combinations of second and third order cumulants
+ //!
+ mfcaa = c1o3 * (mxxMyy + mxxMzz + mxxPyyPzz);
+ mfaca = c1o3 * (-c2o1* mxxMyy + mxxMzz + mxxPyyPzz);
+ mfaac = c1o3 * (mxxMyy - c2o1* mxxMzz + mxxPyyPzz);
+ mfcba = ( mxxyMyzz + mxxyPyzz) * c1o2;
+ mfabc = (-mxxyMyzz + mxxyPyzz) * c1o2;
+ mfcab = ( mxxzMyyz + mxxzPyyz) * c1o2;
+ mfacb = (-mxxzMyyz + mxxzPyyz) * c1o2;
+ mfbca = ( mxyyMxzz + mxyyPxzz) * c1o2;
+ mfbac = (-mxyyMxzz + mxyyPxzz) * c1o2;
+ //////////////////////////////////////////////////////////////////////////
+ //////////////////////////////////////////////////////////////////////////
+ //4.
+ // no limiter
+ //! - Relax fourth order cumulants to modified equilibrium for fourth order convergence of diffusion according to Eq. (43)-(48)
+ //! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), DOI:10.1016/j.jcp.2017.05 040 ]</b></a>
+ //!
+ CUMacc = -O4*(c1o1 / omega - c1o2) * (dyuy + dzuz) * c2o3 * A + (c1o1 - O4) * (CUMacc);
+ CUMcac = -O4*(c1o1 / omega - c1o2) * (dxux + dzuz) * c2o3 * A + (c1o1 - O4) * (CUMcac);
+ CUMcca = -O4*(c1o1 / omega - c1o2) * (dyuy + dxux) * c2o3 * A + (c1o1 - O4) * (CUMcca);
+ CUMbbc = -O4*(c1o1 / omega - c1o2) * Dxy * c1o3 * B + (c1o1 - O4) * (CUMbbc);
+ CUMbcb = -O4*(c1o1 / omega - c1o2) * Dxz * c1o3 * B + (c1o1 - O4) * (CUMbcb);
+ CUMcbb = -O4*(c1o1 / omega - c1o2) * Dyz * c1o3 * B + (c1o1 - O4) * (CUMcbb);
+ //////////////////////////////////////////////////////////////////////////
+ //5.
+ CUMbcc += O5 * (-CUMbcc);
+ CUMcbc += O5 * (-CUMcbc);
+ CUMccb += O5 * (-CUMccb);
+ //////////////////////////////////////////////////////////////////////////
+ //6.
+ CUMccc += O6 * (-CUMccc);
+ ////////////////////////////////////////////////////////////////////////////////////
+ //! - Compute central moments from post collision cumulants according to Eq. (53)-(56) in
+ //! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), DOI:10.1016/j.jcp.2017.05 040 ]</b></a>
+ //!
+ //////////////////////////////////////////////////////////////////////////
+ //4.
+ mfcbb = CUMcbb + c1o3*((c3o1*mfcaa + c1o1) * mfabb + c6o1 * mfbba * mfbab) * OOrho;
+ mfbcb = CUMbcb + c1o3*((c3o1*mfaca + c1o1) * mfbab + c6o1 * mfbba * mfabb) * OOrho;
+ mfbbc = CUMbbc + c1o3*((c3o1*mfaac + c1o1) * mfbba + c6o1 * mfbab * mfabb) * OOrho;
+ mfcca = CUMcca + (((mfcaa * mfaca + c2o1 * mfbba * mfbba)*c9o1 + c3o1 * (mfcaa + mfaca)) * OOrho - (drho * OOrho))*c1o9;
+ mfcac = CUMcac + (((mfcaa * mfaac + c2o1 * mfbab * mfbab)*c9o1 + c3o1 * (mfcaa + mfaac)) * OOrho - (drho * OOrho))*c1o9;
+ mfacc = CUMacc + (((mfaac * mfaca + c2o1 * mfabb * mfabb)*c9o1 + c3o1 * (mfaac + mfaca)) * OOrho - (drho * OOrho))*c1o9;
+ //////////////////////////////////////////////////////////////////////////
+ //5.
+ mfbcc = CUMbcc + c1o3 *(c3o1*(mfaac * mfbca + mfaca * mfbac + c4o1 * mfabb * mfbbb + c2o1 * (mfbab * mfacb + mfbba * mfabc)) + (mfbca + mfbac)) * OOrho;
+ mfcbc = CUMcbc + c1o3 *(c3o1*(mfaac * mfcba + mfcaa * mfabc + c4o1 * mfbab * mfbbb + c2o1 * (mfabb * mfcab + mfbba * mfbac)) + (mfcba + mfabc)) * OOrho;
+ mfccb = CUMccb + c1o3 *(c3o1*(mfcaa * mfacb + mfaca * mfcab + c4o1 * mfbba * mfbbb + c2o1 * (mfbab * mfbca + mfabb * mfcba)) + (mfacb + mfcab)) * OOrho;
+ //////////////////////////////////////////////////////////////////////////
+ //6.
+ mfccc = CUMccc - ((-c4o1 * mfbbb * mfbbb
+ - (mfcaa * mfacc + mfaca * mfcac + mfaac * mfcca)
+ - c4o1 * (mfabb * mfcbb + mfbab * mfbcb + mfbba * mfbbc)
+ - c2o1 * (mfbca * mfbac + mfcba * mfabc + mfcab * mfacb)) * OOrho
+ + (c4o1 * (mfbab * mfbab * mfaca + mfabb * mfabb * mfcaa + mfbba * mfbba * mfaac)
+ + c2o1 * (mfcaa * mfaca * mfaac)
+ + c16o1 * mfbba * mfbab * mfabb) * OOrho * OOrho
+ - c1o3 * (mfacc + mfcac + mfcca) * OOrho
+ - c1o9 * (mfcaa + mfaca + mfaac) * OOrho
+ + (c2o1 * (mfbab * mfbab + mfabb * mfabb + mfbba * mfbba)
+ + (mfaac * mfaca + mfaac * mfcaa + mfaca * mfcaa) + c1o3 *(mfaac + mfaca + mfcaa)) * OOrho * OOrho * c2o3
+ + c1o27*((drho * drho - drho) * OOrho * OOrho));
+ ////////////////////////////////////////////////////////////////////////////////////
+ //! - Add acceleration (body force) to first order cumulants according to Eq. (85)-(87) in
+ //! <a href="https://doi.org/10.1016/j.camwa.2015.05.001"><b>[ M. Geier et al. (2015), DOI:10.1016/j.camwa 2015.05.001 ]</b></a>
+ //!
+ mfbaa = -mfbaa;
+ mfaba = -mfaba;
+ mfaab = -mfaab;
+ ////////////////////////////////////////////////////////////////////////////////////
+ //! - 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), DOI:10.1016/j.camwa 2015.05.001 ]</b></a>
+ //! see also Eq. (88)-(96) in
+ //! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), DOI:10.1016/j.jcp.2017.05 040 ]</b></a>
+ //!
+ ////////////////////////////////////////////////////////////////////////////////////
+ // X - Dir
+ VF::LBM::backwardInverseChimeraWithK(mfaaa, mfbaa, mfcaa, vvx, vx2, c1o1, c1o1);
+ VF::LBM::backwardChimera( mfaba, mfbba, mfcba, vvx, vx2);
+ VF::LBM::backwardInverseChimeraWithK(mfaca, mfbca, mfcca, vvx, vx2, c3o1, c1o3);
+ VF::LBM::backwardChimera( mfaab, mfbab, mfcab, vvx, vx2);
+ VF::LBM::backwardChimera( mfabb, mfbbb, mfcbb, vvx, vx2);
+ VF::LBM::backwardChimera( mfacb, mfbcb, mfccb, vvx, vx2);
+ VF::LBM::backwardInverseChimeraWithK(mfaac, mfbac, mfcac, vvx, vx2, c3o1, c1o3);
+ VF::LBM::backwardChimera( mfabc, mfbbc, mfcbc, vvx, vx2);
+ VF::LBM::backwardInverseChimeraWithK(mfacc, mfbcc, mfccc, vvx, vx2, c9o1, c1o9);
+ ////////////////////////////////////////////////////////////////////////////////////
+ // Y - Dir
+ VF::LBM::backwardInverseChimeraWithK(mfaaa, mfaba, mfaca, vvy, vy2, c6o1, c1o6);
+ VF::LBM::backwardChimera( mfaab, mfabb, mfacb, vvy, vy2);
+ VF::LBM::backwardInverseChimeraWithK(mfaac, mfabc, mfacc, vvy, vy2, c18o1, c1o18);
+ VF::LBM::backwardInverseChimeraWithK(mfbaa, mfbba, mfbca, vvy, vy2, c3o2, c2o3);
+ VF::LBM::backwardChimera( mfbab, mfbbb, mfbcb, vvy, vy2);
+ VF::LBM::backwardInverseChimeraWithK(mfbac, mfbbc, mfbcc, vvy, vy2, c9o2, c2o9);
+ VF::LBM::backwardInverseChimeraWithK(mfcaa, mfcba, mfcca, vvy, vy2, c6o1, c1o6);
+ VF::LBM::backwardChimera( mfcab, mfcbb, mfccb, vvy, vy2);
+ VF::LBM::backwardInverseChimeraWithK(mfcac, mfcbc, mfccc, vvy, vy2, c18o1, c1o18);
+ ////////////////////////////////////////////////////////////////////////////////////
+ // Z - Dir
+ VF::LBM::backwardInverseChimeraWithK(mfaaa, mfaab, mfaac, vvz, vz2, c36o1, c1o36);
+ VF::LBM::backwardInverseChimeraWithK(mfaba, mfabb, mfabc, vvz, vz2, c9o1, c1o9);
+ VF::LBM::backwardInverseChimeraWithK(mfaca, mfacb, mfacc, vvz, vz2, c36o1, c1o36);
+ VF::LBM::backwardInverseChimeraWithK(mfbaa, mfbab, mfbac, vvz, vz2, c9o1, c1o9);
+ VF::LBM::backwardInverseChimeraWithK(mfbba, mfbbb, mfbbc, vvz, vz2, c9o4, c4o9);
+ VF::LBM::backwardInverseChimeraWithK(mfbca, mfbcb, mfbcc, vvz, vz2, c9o1, c1o9);
+ VF::LBM::backwardInverseChimeraWithK(mfcaa, mfcab, mfcac, vvz, vz2, c36o1, c1o36);
+ VF::LBM::backwardInverseChimeraWithK(mfcba, mfcbb, mfcbc, vvz, vz2, c9o1, c1o9);
+ VF::LBM::backwardInverseChimeraWithK(mfcca, mfccb, mfccc, vvz, vz2, c36o1, c1o36);
////////////////////////////////////////////////////////////////////////////////////
//! - Write distributions: style of reading and writing the distributions from/to
//! stored arrays dependent on timestep is based on the esoteric twist algorithm
//! <a href="https://doi.org/10.3390/computation5020019"><b>[ M. Geier et al. (2017), DOI:10.3390/computation5020019 ]</b></a>
//!
- // (dist.f[dirE ])[k] = distribution_out.f[VF::LBM::DIR::PZZ];
- // (dist.f[dirW ])[kw] = distribution_out.f[VF::LBM::DIR::MZZ];
- // (dist.f[dirN ])[k] = distribution_out.f[VF::LBM::DIR::ZPZ];
- // (dist.f[dirS ])[ks] = distribution_out.f[VF::LBM::DIR::ZMZ];
- // (dist.f[dirT ])[k] = distribution_out.f[VF::LBM::DIR::ZZP];
- // (dist.f[dirB ])[kb] = distribution_out.f[VF::LBM::DIR::ZZM];
- // (dist.f[dirNE ])[k] = distribution_out.f[VF::LBM::DIR::PPZ];
- // (dist.f[dirSW ])[ksw] = distribution_out.f[VF::LBM::DIR::MMZ];
- // (dist.f[dirSE ])[ks] = distribution_out.f[VF::LBM::DIR::PMZ];
- // (dist.f[dirNW ])[kw] = distribution_out.f[VF::LBM::DIR::MPZ];
- // (dist.f[dirTE ])[k] = distribution_out.f[VF::LBM::DIR::PZP];
- // (dist.f[dirBW ])[kbw] = distribution_out.f[VF::LBM::DIR::MZM];
- // (dist.f[dirBE ])[kb] = distribution_out.f[VF::LBM::DIR::PZM];
- // (dist.f[dirTW ])[kw] = distribution_out.f[VF::LBM::DIR::MZP];
- // (dist.f[dirTN ])[k] = distribution_out.f[VF::LBM::DIR::ZPP];
- // (dist.f[dirBS ])[kbs] = distribution_out.f[VF::LBM::DIR::ZMM];
- // (dist.f[dirBN ])[kb] = distribution_out.f[VF::LBM::DIR::ZPM];
- // (dist.f[dirTS ])[ks] = distribution_out.f[VF::LBM::DIR::ZMP];
- // (dist.f[dirTNE ])[k] = distribution_out.f[VF::LBM::DIR::PPP];
- // (dist.f[dirTNW ])[kw] = distribution_out.f[VF::LBM::DIR::MPP];
- // (dist.f[dirTSE ])[ks] = distribution_out.f[VF::LBM::DIR::PMP];
- // (dist.f[dirTSW ])[ksw] = distribution_out.f[VF::LBM::DIR::MMP];
- // (dist.f[dirBNE ])[kb] = distribution_out.f[VF::LBM::DIR::PPM];
- // (dist.f[dirBNE ])[kb] = distribution_out.f[VF::LBM::DIR::MPM];
- // (dist.f[dirBSE ])[kbs] = distribution_out.f[VF::LBM::DIR::PMM];
- // (dist.f[dirBSW ])[kbsw] = distribution_out.f[VF::LBM::DIR::MMM];
- // (dist.f[dirREST])[k] = distribution_out.f[VF::LBM::DIR::ZZZ];
+ // (dist.f[dirE ])[k ] = mfabb;
+ // (dist.f[dirW ])[kw ] = mfcbb;
+ // (dist.f[dirN ])[k ] = mfbab;
+ // (dist.f[dirS ])[ks ] = mfbcb;
+ // (dist.f[dirT ])[k ] = mfbba;
+ // (dist.f[dirB ])[kb ] = mfbbc;
+ // (dist.f[dirNE ])[k ] = mfaab;
+ // (dist.f[dirSW ])[ksw ] = mfccb;
+ // (dist.f[dirSE ])[ks ] = mfacb;
+ // (dist.f[dirNW ])[kw ] = mfcab;
+ // (dist.f[dirTE ])[k ] = mfaba;
+ // (dist.f[dirBW ])[kbw ] = mfcbc;
+ // (dist.f[dirBE ])[kb ] = mfabc;
+ // (dist.f[dirTW ])[kw ] = mfcba;
+ // (dist.f[dirTN ])[k ] = mfbaa;
+ // (dist.f[dirBS ])[kbs ] = mfbcc;
+ // (dist.f[dirBN ])[kb ] = mfbac;
+ // (dist.f[dirTS ])[ks ] = mfbca;
+ // (dist.f[dirREST])[k ] = mfbbb;
+ // (dist.f[dirTNE ])[k ] = mfaaa;
+ // (dist.f[dirTSE ])[ks ] = mfaca;
+ // (dist.f[dirBNE ])[kb ] = mfaac;
+ // (dist.f[dirBSE ])[kbs ] = mfacc;
+ // (dist.f[dirTNW ])[kw ] = mfcaa;
+ // (dist.f[dirTSW ])[ksw ] = mfcca;
+ // (dist.f[dirBNW ])[kbw ] = mfcac;
+ // (dist.f[dirBSW ])[kbsw] = mfccc;
+
+ VF::LBM::Distribution27 distribution_out;
+
+ distribution_out.f[VF::LBM::DIR::MZZ] = mfcbb;
+ distribution_out.f[VF::LBM::DIR::PZZ] = mfabb;
+ distribution_out.f[VF::LBM::DIR::ZMZ] = mfbcb;
+ distribution_out.f[VF::LBM::DIR::ZPZ] = mfbab;
+ distribution_out.f[VF::LBM::DIR::ZZM] = mfbbc;
+ distribution_out.f[VF::LBM::DIR::ZZP] = mfbba;
+ distribution_out.f[VF::LBM::DIR::MMZ] = mfccb;
+ distribution_out.f[VF::LBM::DIR::PPZ] = mfaab;
+ distribution_out.f[VF::LBM::DIR::MPZ] = mfcab;
+ distribution_out.f[VF::LBM::DIR::PMZ] = mfacb;
+ distribution_out.f[VF::LBM::DIR::MZM] = mfcbc;
+ distribution_out.f[VF::LBM::DIR::PZP] = mfaba;
+ distribution_out.f[VF::LBM::DIR::MZP] = mfcba;
+ distribution_out.f[VF::LBM::DIR::PZM] = mfabc;
+ distribution_out.f[VF::LBM::DIR::ZMM] = mfbcc;
+ distribution_out.f[VF::LBM::DIR::ZPP] = mfbaa;
+ distribution_out.f[VF::LBM::DIR::ZMP] = mfbca;
+ distribution_out.f[VF::LBM::DIR::ZPM] = mfbac;
+ distribution_out.f[VF::LBM::DIR::MMM] = mfccc;
+ distribution_out.f[VF::LBM::DIR::PMM] = mfacc;
+ distribution_out.f[VF::LBM::DIR::MPM] = mfcac;
+ distribution_out.f[VF::LBM::DIR::PPM] = mfaac;
+ distribution_out.f[VF::LBM::DIR::MMP] = mfcca;
+ distribution_out.f[VF::LBM::DIR::PMP] = mfaca;
+ distribution_out.f[VF::LBM::DIR::MPP] = mfcaa;
+ distribution_out.f[VF::LBM::DIR::PPP] = mfaaa;
+ distribution_out.f[VF::LBM::DIR::ZZZ] = mfbbb;
+
+
+ (dist.f[dirE ])[k] = distribution_out.f[VF::LBM::DIR::PZZ];
+ (dist.f[dirW ])[kw] = distribution_out.f[VF::LBM::DIR::MZZ];
+ (dist.f[dirN ])[k] = distribution_out.f[VF::LBM::DIR::ZPZ];
+ (dist.f[dirS ])[ks] = distribution_out.f[VF::LBM::DIR::ZMZ];
+ (dist.f[dirT ])[k] = distribution_out.f[VF::LBM::DIR::ZZP];
+ (dist.f[dirB ])[kb] = distribution_out.f[VF::LBM::DIR::ZZM];
+ (dist.f[dirNE ])[k] = distribution_out.f[VF::LBM::DIR::PPZ];
+ (dist.f[dirSW ])[ksw] = distribution_out.f[VF::LBM::DIR::MMZ];
+ (dist.f[dirSE ])[ks] = distribution_out.f[VF::LBM::DIR::PMZ];
+ (dist.f[dirNW ])[kw] = distribution_out.f[VF::LBM::DIR::MPZ];
+ (dist.f[dirTE ])[k] = distribution_out.f[VF::LBM::DIR::PZP];
+ (dist.f[dirBW ])[kbw] = distribution_out.f[VF::LBM::DIR::MZM];
+ (dist.f[dirBE ])[kb] = distribution_out.f[VF::LBM::DIR::PZM];
+ (dist.f[dirTW ])[kw] = distribution_out.f[VF::LBM::DIR::MZP];
+ (dist.f[dirTN ])[k] = distribution_out.f[VF::LBM::DIR::ZPP];
+ (dist.f[dirBS ])[kbs] = distribution_out.f[VF::LBM::DIR::ZMM];
+ (dist.f[dirBN ])[kb] = distribution_out.f[VF::LBM::DIR::ZPM];
+ (dist.f[dirTS ])[ks] = distribution_out.f[VF::LBM::DIR::ZMP];
+ (dist.f[dirTNE ])[k] = distribution_out.f[VF::LBM::DIR::PPP];
+ (dist.f[dirTNW ])[kw] = distribution_out.f[VF::LBM::DIR::MPP];
+ (dist.f[dirTSE ])[ks] = distribution_out.f[VF::LBM::DIR::PMP];
+ (dist.f[dirTSW ])[ksw] = distribution_out.f[VF::LBM::DIR::MMP];
+ (dist.f[dirBNE ])[kb] = distribution_out.f[VF::LBM::DIR::PPM];
+ (dist.f[dirBNW ])[kbw] = distribution_out.f[VF::LBM::DIR::MPM];
+ (dist.f[dirBSE ])[kbs] = distribution_out.f[VF::LBM::DIR::PMM];
+ (dist.f[dirBSW ])[kbsw] = distribution_out.f[VF::LBM::DIR::MMM];
+ (dist.f[dirREST])[k] = distribution_out.f[VF::LBM::DIR::ZZZ];
+
}
}
////////////////////////////////////////////////////////////////////////////////
-
-////////////////////////////////////////////////////////////////////////////////
-extern "C" __global__ void Cumulant_K17_LBM_Device_Kernel_old(
+extern "C" __global__ void Cumulant_K17_LBM_Device_Kernel_old_origin(
real omega,
uint* typeOfGridNode,
uint* neighborX,
diff --git a/src/lbm/CumulantChimera.h b/src/lbm/CumulantChimera.h
index 90c3d53cd..690ff227d 100644
--- a/src/lbm/CumulantChimera.h
+++ b/src/lbm/CumulantChimera.h
@@ -23,16 +23,28 @@ namespace LBM
struct Distribution27
{
- real f[27];
+ real f[27];
- inline __host__ __device__ real getDensity_() const
- {
- return getDensity(f);
- }
+ inline __host__ __device__ real getDensity_() const
+ {
+ return getDensity(f);
+ }
};
-inline __host__ __device__ Distribution27 cumulantChimera(const Distribution27& distribution, real omega, const real* forces)
+
+//////////////////////////////////////////////////////////////////////////
+//! Cumulant K17 Kernel is based on \ref
+//! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), DOI:10.1016/j.jcp.2017.05.040 ]</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>
+//////////////////////////////////////////////////////////////////////////
+inline __host__ __device__ void cumulantChimera(Distribution27& distribution, real omega, real* forces)
{
+ ////////////////////////////////////////////////////////////////////////////////////
+ //! - Read distributions: style of reading and writing the distributions from/to
+ //! stored arrays dependent on timestep is based on the esoteric twist algorithm
+ //! <a href="https://doi.org/10.3390/computation5020019"><b>[ M. Geier et al. (2017), DOI:10.3390/computation5020019 ]</b></a>
+ //!
real mfcbb = distribution.f[DIR::PZZ];
real mfabb = distribution.f[DIR::MZZ];
real mfbcb = distribution.f[DIR::ZPZ];
@@ -66,23 +78,23 @@ inline __host__ __device__ Distribution27 cumulantChimera(const Distribution27&
//! <a href="https://doi.org/10.1016/j.camwa.2015.05.001"><b>[ M. Geier et al. (2015), DOI:10.1016/j.camwa 2015.05.001 ]</b></a>
//!
real drho =
- ((((mfccc + mfaaa) + (mfaca + mfcac)) + ((mfacc + mfcaa) + (mfaac + mfcca))) +
- (((mfbac + mfbca) + (mfbaa + mfbcc)) + ((mfabc + mfcba) + (mfaba + mfcbc)) + ((mfacb + mfcab) + (mfaab + mfccb))) +
- ((mfabb + mfcbb) + (mfbab + mfbcb) + (mfbba + mfbbc))) + mfbbb;
+ ((((mfccc + mfaaa) + (mfaca + mfcac)) + ((mfacc + mfcaa) + (mfaac + mfcca))) +
+ (((mfbac + mfbca) + (mfbaa + mfbcc)) + ((mfabc + mfcba) + (mfaba + mfcbc)) + ((mfacb + mfcab) + (mfaab + mfccb))) +
+ ((mfabb + mfcbb) + (mfbab + mfbcb) + (mfbba + mfbbc))) + mfbbb;
real rho = c1o1 + drho;
real OOrho = c1o1 / rho;
real vvx =
- ((((mfccc - mfaaa) + (mfcac - mfaca)) + ((mfcaa - mfacc) + (mfcca - mfaac))) +
- (((mfcba - mfabc) + (mfcbc - mfaba)) + ((mfcab - mfacb) + (mfccb - mfaab))) +
- (mfcbb - mfabb)) * OOrho;
+ ((((mfccc - mfaaa) + (mfcac - mfaca)) + ((mfcaa - mfacc) + (mfcca - mfaac))) +
+ (((mfcba - mfabc) + (mfcbc - mfaba)) + ((mfcab - mfacb) + (mfccb - mfaab))) +
+ (mfcbb - mfabb)) * OOrho;
real vvy =
- ((((mfccc - mfaaa) + (mfaca - mfcac)) + ((mfacc - mfcaa) + (mfcca - mfaac))) +
- (((mfbca - mfbac) + (mfbcc - mfbaa)) + ((mfacb - mfcab) + (mfccb - mfaab))) +
- (mfbcb - mfbab)) * OOrho;
+ ((((mfccc - mfaaa) + (mfaca - mfcac)) + ((mfacc - mfcaa) + (mfcca - mfaac))) +
+ (((mfbca - mfbac) + (mfbcc - mfbaa)) + ((mfacb - mfcab) + (mfccb - mfaab))) +
+ (mfbcb - mfbab)) * OOrho;
real vvz =
- ((((mfccc - mfaaa) + (mfcac - mfaca)) + ((mfacc - mfcaa) + (mfaac - mfcca))) +
- (((mfbac - mfbca) + (mfbcc - mfbaa)) + ((mfabc - mfcba) + (mfcbc - mfaba))) +
- (mfbbc - mfbba)) * OOrho;
+ ((((mfccc - mfaaa) + (mfcac - mfaca)) + ((mfacc - mfcaa) + (mfaac - mfcca))) +
+ (((mfbac - mfbca) + (mfbcc - mfbaa)) + ((mfabc - mfcba) + (mfcbc - mfaba))) +
+ (mfbbc - mfbba)) * OOrho;
////////////////////////////////////////////////////////////////////////////////////
//! - Add half of the acceleration (body force) to the velocity as in Eq. (42) \ref
//! <a href="https://doi.org/10.1016/j.camwa.2015.05.001"><b>[ M. Geier et al. (2015), DOI:10.1016/j.camwa 2015.05.001 ]</b></a>
@@ -198,17 +210,17 @@ inline __host__ __device__ Distribution27 cumulantChimera(const Distribution27&
////////////////////////////////////////////////////////////
//6.
real CUMccc = mfccc + ((-c4o1 * mfbbb * mfbbb
- - (mfcaa * mfacc + mfaca * mfcac + mfaac * mfcca)
- - c4o1 * (mfabb * mfcbb + mfbab * mfbcb + mfbba * mfbbc)
- - c2o1 * (mfbca * mfbac + mfcba * mfabc + mfcab * mfacb)) * OOrho
- + (c4o1 * (mfbab * mfbab * mfaca + mfabb * mfabb * mfcaa + mfbba * mfbba * mfaac)
- + c2o1 * (mfcaa * mfaca * mfaac)
- + c16o1 * mfbba * mfbab * mfabb) * OOrho * OOrho
- - c1o3 * (mfacc + mfcac + mfcca) * OOrho
- - c1o9 * (mfcaa + mfaca + mfaac) * OOrho
- + (c2o1 * (mfbab * mfbab + mfabb * mfabb + mfbba * mfbba)
- + (mfaac * mfaca + mfaac * mfcaa + mfaca * mfcaa) + c1o3 *(mfaac + mfaca + mfcaa)) * OOrho * OOrho * c2o3
- + c1o27*((drho * drho - drho) * OOrho * OOrho));
+ - (mfcaa * mfacc + mfaca * mfcac + mfaac * mfcca)
+ - c4o1 * (mfabb * mfcbb + mfbab * mfbcb + mfbba * mfbbc)
+ - c2o1 * (mfbca * mfbac + mfcba * mfabc + mfcab * mfacb)) * OOrho
+ + (c4o1 * (mfbab * mfbab * mfaca + mfabb * mfabb * mfcaa + mfbba * mfbba * mfaac)
+ + c2o1 * (mfcaa * mfaca * mfaac)
+ + c16o1 * mfbba * mfbab * mfabb) * OOrho * OOrho
+ - c1o3 * (mfacc + mfcac + mfcca) * OOrho
+ - c1o9 * (mfcaa + mfaca + mfaac) * OOrho
+ + (c2o1 * (mfbab * mfbab + mfabb * mfabb + mfbba * mfbba)
+ + (mfaac * mfaca + mfaac * mfcaa + mfaca * mfcaa) + c1o3 *(mfaac + mfaca + mfcaa)) * OOrho * OOrho * c2o3
+ + c1o27*((drho * drho - drho) * OOrho * OOrho));
////////////////////////////////////////////////////////////////////////////////////
//! - Compute linear combinations of second and third order cumulants
//!
@@ -339,17 +351,17 @@ inline __host__ __device__ Distribution27 cumulantChimera(const Distribution27&
//////////////////////////////////////////////////////////////////////////
//6.
mfccc = CUMccc - ((-c4o1 * mfbbb * mfbbb
- - (mfcaa * mfacc + mfaca * mfcac + mfaac * mfcca)
- - c4o1 * (mfabb * mfcbb + mfbab * mfbcb + mfbba * mfbbc)
- - c2o1 * (mfbca * mfbac + mfcba * mfabc + mfcab * mfacb)) * OOrho
- + (c4o1 * (mfbab * mfbab * mfaca + mfabb * mfabb * mfcaa + mfbba * mfbba * mfaac)
- + c2o1 * (mfcaa * mfaca * mfaac)
- + c16o1 * mfbba * mfbab * mfabb) * OOrho * OOrho
- - c1o3 * (mfacc + mfcac + mfcca) * OOrho
- - c1o9 * (mfcaa + mfaca + mfaac) * OOrho
- + (c2o1 * (mfbab * mfbab + mfabb * mfabb + mfbba * mfbba)
- + (mfaac * mfaca + mfaac * mfcaa + mfaca * mfcaa) + c1o3 *(mfaac + mfaca + mfcaa)) * OOrho * OOrho * c2o3
- + c1o27*((drho * drho - drho) * OOrho * OOrho));
+ - (mfcaa * mfacc + mfaca * mfcac + mfaac * mfcca)
+ - c4o1 * (mfabb * mfcbb + mfbab * mfbcb + mfbba * mfbbc)
+ - c2o1 * (mfbca * mfbac + mfcba * mfabc + mfcab * mfacb)) * OOrho
+ + (c4o1 * (mfbab * mfbab * mfaca + mfabb * mfabb * mfcaa + mfbba * mfbba * mfaac)
+ + c2o1 * (mfcaa * mfaca * mfaac)
+ + c16o1 * mfbba * mfbab * mfabb) * OOrho * OOrho
+ - c1o3 * (mfacc + mfcac + mfcca) * OOrho
+ - c1o9 * (mfcaa + mfaca + mfaac) * OOrho
+ + (c2o1 * (mfbab * mfbab + mfabb * mfabb + mfbba * mfbba)
+ + (mfaac * mfaca + mfaac * mfcaa + mfaca * mfcaa) + c1o3 *(mfaac + mfaca + mfcaa)) * OOrho * OOrho * c2o3
+ + c1o27*((drho * drho - drho) * OOrho * OOrho));
////////////////////////////////////////////////////////////////////////////////////
//! - Add acceleration (body force) to first order cumulants according to Eq. (85)-(87) in
//! <a href="https://doi.org/10.1016/j.camwa.2015.05.001"><b>[ M. Geier et al. (2015), DOI:10.1016/j.camwa 2015.05.001 ]</b></a>
@@ -397,39 +409,42 @@ inline __host__ __device__ Distribution27 cumulantChimera(const Distribution27&
VF::LBM::backwardInverseChimeraWithK(mfcba, mfcbb, mfcbc, vvz, vz2, c9o1, c1o9);
VF::LBM::backwardInverseChimeraWithK(mfcca, mfccb, mfccc, vvz, vz2, c36o1, c1o36);
- Distribution27 distribution_calculated;
-
- distribution_calculated.f[DIR::PZZ] = mfcbb;
- distribution_calculated.f[DIR::MZZ] = mfabb;
- distribution_calculated.f[DIR::ZPZ] = mfbcb;
- distribution_calculated.f[DIR::ZMZ] = mfbab;
- distribution_calculated.f[DIR::ZZP] = mfbbc;
- distribution_calculated.f[DIR::ZZM] = mfbba;
- distribution_calculated.f[DIR::PPZ] = mfccb;
- distribution_calculated.f[DIR::MMZ] = mfaab;
- distribution_calculated.f[DIR::PMZ] = mfcab;
- distribution_calculated.f[DIR::MPZ] = mfacb;
- distribution_calculated.f[DIR::PZP] = mfcbc;
- distribution_calculated.f[DIR::MZM] = mfaba;
- distribution_calculated.f[DIR::PZM] = mfcba;
- distribution_calculated.f[DIR::MZP] = mfabc;
- distribution_calculated.f[DIR::ZPP] = mfbcc;
- distribution_calculated.f[DIR::ZMM] = mfbaa;
- distribution_calculated.f[DIR::ZPM] = mfbca;
- distribution_calculated.f[DIR::ZMP] = mfbac;
- distribution_calculated.f[DIR::PPP] = mfccc;
- distribution_calculated.f[DIR::MPP] = mfacc;
- distribution_calculated.f[DIR::PMP] = mfcac;
- distribution_calculated.f[DIR::MMP] = mfaac;
- distribution_calculated.f[DIR::PPM] = mfcca;
- distribution_calculated.f[DIR::MPM] = mfaca;
- distribution_calculated.f[DIR::PMM] = mfcaa;
- distribution_calculated.f[DIR::MMM] = mfaaa;
- distribution_calculated.f[DIR::ZZZ] = mfbbb;
- return distribution_calculated;
+ ////////////////////////////////////////////////////////////////////////////////////
+ //! - Write distributions: style of reading and writing the distributions from/to
+ //! stored arrays dependent on timestep is based on the esoteric twist algorithm
+ //! <a href="https://doi.org/10.3390/computation5020019"><b>[ M. Geier et al. (2017), DOI:10.3390/computation5020019 ]</b></a>
+ //!
+ distribution.f[VF::LBM::DIR::MZZ] = mfcbb;
+ distribution.f[VF::LBM::DIR::PZZ] = mfabb;
+ distribution.f[VF::LBM::DIR::ZMZ] = mfbcb;
+ distribution.f[VF::LBM::DIR::ZPZ] = mfbab;
+ distribution.f[VF::LBM::DIR::ZZM] = mfbbc;
+ distribution.f[VF::LBM::DIR::ZZP] = mfbba;
+ distribution.f[VF::LBM::DIR::MMZ] = mfccb;
+ distribution.f[VF::LBM::DIR::PPZ] = mfaab;
+ distribution.f[VF::LBM::DIR::MPZ] = mfcab;
+ distribution.f[VF::LBM::DIR::PMZ] = mfacb;
+ distribution.f[VF::LBM::DIR::MZM] = mfcbc;
+ distribution.f[VF::LBM::DIR::PZP] = mfaba;
+ distribution.f[VF::LBM::DIR::MZP] = mfcba;
+ distribution.f[VF::LBM::DIR::PZM] = mfabc;
+ distribution.f[VF::LBM::DIR::ZMM] = mfbcc;
+ distribution.f[VF::LBM::DIR::ZPP] = mfbaa;
+ distribution.f[VF::LBM::DIR::ZMP] = mfbca;
+ distribution.f[VF::LBM::DIR::ZPM] = mfbac;
+ distribution.f[VF::LBM::DIR::MMM] = mfccc;
+ distribution.f[VF::LBM::DIR::PMM] = mfacc;
+ distribution.f[VF::LBM::DIR::MPM] = mfcac;
+ distribution.f[VF::LBM::DIR::PPM] = mfaac;
+ distribution.f[VF::LBM::DIR::MMP] = mfcca;
+ distribution.f[VF::LBM::DIR::PMP] = mfaca;
+ distribution.f[VF::LBM::DIR::MPP] = mfcaa;
+ distribution.f[VF::LBM::DIR::PPP] = mfaaa;
+ distribution.f[VF::LBM::DIR::ZZZ] = mfbbb;
}
+
}
}
#endif
--
GitLab