diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Utilities/DistributionHelper.cu b/src/gpu/VirtualFluids_GPU/Kernel/Utilities/DistributionHelper.cu
index 08241e2a55ea9bb053fb8097757d706ae0345f6a..db61c77857eee5847af7e5c58cc9005f331cd42e 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Utilities/DistributionHelper.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Utilities/DistributionHelper.cu
@@ -10,74 +10,73 @@ namespace vf
namespace gpu
{
-__device__ __host__ Distributions27 getDistributions27(real *distributions, unsigned int size_Mat, bool isEvenTimestep)
+__device__ __host__ DistributionReferences27 getDistributionReferences27(real *distributions, unsigned int size_Mat, bool isEvenTimestep)
{
-
- Distributions27 dist;
+ DistributionReferences27 distribution_references;
if (isEvenTimestep) {
- dist.f[dirE] = &distributions[dirE * size_Mat];
- dist.f[dirW] = &distributions[dirW * size_Mat];
- dist.f[dirN] = &distributions[dirN * size_Mat];
- dist.f[dirS] = &distributions[dirS * size_Mat];
- dist.f[dirT] = &distributions[dirT * size_Mat];
- dist.f[dirB] = &distributions[dirB * size_Mat];
- dist.f[dirNE] = &distributions[dirNE * size_Mat];
- dist.f[dirSW] = &distributions[dirSW * size_Mat];
- dist.f[dirSE] = &distributions[dirSE * size_Mat];
- dist.f[dirNW] = &distributions[dirNW * size_Mat];
- dist.f[dirTE] = &distributions[dirTE * size_Mat];
- dist.f[dirBW] = &distributions[dirBW * size_Mat];
- dist.f[dirBE] = &distributions[dirBE * size_Mat];
- dist.f[dirTW] = &distributions[dirTW * size_Mat];
- dist.f[dirTN] = &distributions[dirTN * size_Mat];
- dist.f[dirBS] = &distributions[dirBS * size_Mat];
- dist.f[dirBN] = &distributions[dirBN * size_Mat];
- dist.f[dirTS] = &distributions[dirTS * size_Mat];
- dist.f[dirREST] = &distributions[dirREST * size_Mat];
- dist.f[dirTNE] = &distributions[dirTNE * size_Mat];
- dist.f[dirTSW] = &distributions[dirTSW * size_Mat];
- dist.f[dirTSE] = &distributions[dirTSE * size_Mat];
- dist.f[dirTNW] = &distributions[dirTNW * size_Mat];
- dist.f[dirBNE] = &distributions[dirBNE * size_Mat];
- dist.f[dirBSW] = &distributions[dirBSW * size_Mat];
- dist.f[dirBSE] = &distributions[dirBSE * size_Mat];
- dist.f[dirBNW] = &distributions[dirBNW * size_Mat];
+ distribution_references.f[dirE] = &distributions[dirE * size_Mat];
+ distribution_references.f[dirW] = &distributions[dirW * size_Mat];
+ distribution_references.f[dirN] = &distributions[dirN * size_Mat];
+ distribution_references.f[dirS] = &distributions[dirS * size_Mat];
+ distribution_references.f[dirT] = &distributions[dirT * size_Mat];
+ distribution_references.f[dirB] = &distributions[dirB * size_Mat];
+ distribution_references.f[dirNE] = &distributions[dirNE * size_Mat];
+ distribution_references.f[dirSW] = &distributions[dirSW * size_Mat];
+ distribution_references.f[dirSE] = &distributions[dirSE * size_Mat];
+ distribution_references.f[dirNW] = &distributions[dirNW * size_Mat];
+ distribution_references.f[dirTE] = &distributions[dirTE * size_Mat];
+ distribution_references.f[dirBW] = &distributions[dirBW * size_Mat];
+ distribution_references.f[dirBE] = &distributions[dirBE * size_Mat];
+ distribution_references.f[dirTW] = &distributions[dirTW * size_Mat];
+ distribution_references.f[dirTN] = &distributions[dirTN * size_Mat];
+ distribution_references.f[dirBS] = &distributions[dirBS * size_Mat];
+ distribution_references.f[dirBN] = &distributions[dirBN * size_Mat];
+ distribution_references.f[dirTS] = &distributions[dirTS * size_Mat];
+ distribution_references.f[dirREST] = &distributions[dirREST * size_Mat];
+ distribution_references.f[dirTNE] = &distributions[dirTNE * size_Mat];
+ distribution_references.f[dirTSW] = &distributions[dirTSW * size_Mat];
+ distribution_references.f[dirTSE] = &distributions[dirTSE * size_Mat];
+ distribution_references.f[dirTNW] = &distributions[dirTNW * size_Mat];
+ distribution_references.f[dirBNE] = &distributions[dirBNE * size_Mat];
+ distribution_references.f[dirBSW] = &distributions[dirBSW * size_Mat];
+ distribution_references.f[dirBSE] = &distributions[dirBSE * size_Mat];
+ distribution_references.f[dirBNW] = &distributions[dirBNW * size_Mat];
} else {
- dist.f[dirW] = &distributions[dirE * size_Mat];
- dist.f[dirE] = &distributions[dirW * size_Mat];
- dist.f[dirS] = &distributions[dirN * size_Mat];
- dist.f[dirN] = &distributions[dirS * size_Mat];
- dist.f[dirB] = &distributions[dirT * size_Mat];
- dist.f[dirT] = &distributions[dirB * size_Mat];
- dist.f[dirSW] = &distributions[dirNE * size_Mat];
- dist.f[dirNE] = &distributions[dirSW * size_Mat];
- dist.f[dirNW] = &distributions[dirSE * size_Mat];
- dist.f[dirSE] = &distributions[dirNW * size_Mat];
- dist.f[dirBW] = &distributions[dirTE * size_Mat];
- dist.f[dirTE] = &distributions[dirBW * size_Mat];
- dist.f[dirTW] = &distributions[dirBE * size_Mat];
- dist.f[dirBE] = &distributions[dirTW * size_Mat];
- dist.f[dirBS] = &distributions[dirTN * size_Mat];
- dist.f[dirTN] = &distributions[dirBS * size_Mat];
- dist.f[dirTS] = &distributions[dirBN * size_Mat];
- dist.f[dirBN] = &distributions[dirTS * size_Mat];
- dist.f[dirREST] = &distributions[dirREST * size_Mat];
- dist.f[dirBSW] = &distributions[dirTNE * size_Mat];
- dist.f[dirBNE] = &distributions[dirTSW * size_Mat];
- dist.f[dirBNW] = &distributions[dirTSE * size_Mat];
- dist.f[dirBSE] = &distributions[dirTNW * size_Mat];
- dist.f[dirTSW] = &distributions[dirBNE * size_Mat];
- dist.f[dirTNE] = &distributions[dirBSW * size_Mat];
- dist.f[dirTNW] = &distributions[dirBSE * size_Mat];
- dist.f[dirTSE] = &distributions[dirBNW * size_Mat];
+ distribution_references.f[dirW] = &distributions[dirE * size_Mat];
+ distribution_references.f[dirE] = &distributions[dirW * size_Mat];
+ distribution_references.f[dirS] = &distributions[dirN * size_Mat];
+ distribution_references.f[dirN] = &distributions[dirS * size_Mat];
+ distribution_references.f[dirB] = &distributions[dirT * size_Mat];
+ distribution_references.f[dirT] = &distributions[dirB * size_Mat];
+ distribution_references.f[dirSW] = &distributions[dirNE * size_Mat];
+ distribution_references.f[dirNE] = &distributions[dirSW * size_Mat];
+ distribution_references.f[dirNW] = &distributions[dirSE * size_Mat];
+ distribution_references.f[dirSE] = &distributions[dirNW * size_Mat];
+ distribution_references.f[dirBW] = &distributions[dirTE * size_Mat];
+ distribution_references.f[dirTE] = &distributions[dirBW * size_Mat];
+ distribution_references.f[dirTW] = &distributions[dirBE * size_Mat];
+ distribution_references.f[dirBE] = &distributions[dirTW * size_Mat];
+ distribution_references.f[dirBS] = &distributions[dirTN * size_Mat];
+ distribution_references.f[dirTN] = &distributions[dirBS * size_Mat];
+ distribution_references.f[dirTS] = &distributions[dirBN * size_Mat];
+ distribution_references.f[dirBN] = &distributions[dirTS * size_Mat];
+ distribution_references.f[dirREST] = &distributions[dirREST * size_Mat];
+ distribution_references.f[dirBSW] = &distributions[dirTNE * size_Mat];
+ distribution_references.f[dirBNE] = &distributions[dirTSW * size_Mat];
+ distribution_references.f[dirBNW] = &distributions[dirTSE * size_Mat];
+ distribution_references.f[dirBSE] = &distributions[dirTNW * size_Mat];
+ distribution_references.f[dirTSW] = &distributions[dirBNE * size_Mat];
+ distribution_references.f[dirTNE] = &distributions[dirBSW * size_Mat];
+ distribution_references.f[dirTNW] = &distributions[dirBSE * size_Mat];
+ distribution_references.f[dirTSE] = &distributions[dirBNW * size_Mat];
}
- return dist;
+ return distribution_references;
}
__device__ DistributionWrapper::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(neighborX[k]), ks(neighborY[k]),
+ : distribution_references(getDistributionReferences27(distributions, size_Mat, isEvenTimestep)), k(k), kw(neighborX[k]), ks(neighborY[k]),
kb(neighborZ[k]), ksw(neighborY[kw]), kbw(neighborZ[kw]), kbs(neighborZ[ks]), kbsw(neighborZ[ksw])
{
read();
@@ -85,64 +84,64 @@ __device__ DistributionWrapper::DistributionWrapper(real *distributions, unsigne
__device__ void DistributionWrapper::read()
{
- distribution.f[vf::lbm::dir::PZZ] = (dist.f[dirE])[k];
- distribution.f[vf::lbm::dir::MZZ] = (dist.f[dirW])[kw];
- distribution.f[vf::lbm::dir::ZPZ] = (dist.f[dirN])[k];
- distribution.f[vf::lbm::dir::ZMZ] = (dist.f[dirS])[ks];
- distribution.f[vf::lbm::dir::ZZP] = (dist.f[dirT])[k];
- distribution.f[vf::lbm::dir::ZZM] = (dist.f[dirB])[kb];
- distribution.f[vf::lbm::dir::PPZ] = (dist.f[dirNE])[k];
- distribution.f[vf::lbm::dir::MMZ] = (dist.f[dirSW])[ksw];
- distribution.f[vf::lbm::dir::PMZ] = (dist.f[dirSE])[ks];
- distribution.f[vf::lbm::dir::MPZ] = (dist.f[dirNW])[kw];
- distribution.f[vf::lbm::dir::PZP] = (dist.f[dirTE])[k];
- distribution.f[vf::lbm::dir::MZM] = (dist.f[dirBW])[kbw];
- distribution.f[vf::lbm::dir::PZM] = (dist.f[dirBE])[kb];
- distribution.f[vf::lbm::dir::MZP] = (dist.f[dirTW])[kw];
- distribution.f[vf::lbm::dir::ZPP] = (dist.f[dirTN])[k];
- distribution.f[vf::lbm::dir::ZMM] = (dist.f[dirBS])[kbs];
- distribution.f[vf::lbm::dir::ZPM] = (dist.f[dirBN])[kb];
- distribution.f[vf::lbm::dir::ZMP] = (dist.f[dirTS])[ks];
- distribution.f[vf::lbm::dir::PPP] = (dist.f[dirTNE])[k];
- distribution.f[vf::lbm::dir::MPP] = (dist.f[dirTNW])[kw];
- 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[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];
+ distribution.f[vf::lbm::dir::PZZ] = (distribution_references.f[dirE])[k];
+ distribution.f[vf::lbm::dir::MZZ] = (distribution_references.f[dirW])[kw];
+ distribution.f[vf::lbm::dir::ZPZ] = (distribution_references.f[dirN])[k];
+ distribution.f[vf::lbm::dir::ZMZ] = (distribution_references.f[dirS])[ks];
+ distribution.f[vf::lbm::dir::ZZP] = (distribution_references.f[dirT])[k];
+ distribution.f[vf::lbm::dir::ZZM] = (distribution_references.f[dirB])[kb];
+ distribution.f[vf::lbm::dir::PPZ] = (distribution_references.f[dirNE])[k];
+ distribution.f[vf::lbm::dir::MMZ] = (distribution_references.f[dirSW])[ksw];
+ distribution.f[vf::lbm::dir::PMZ] = (distribution_references.f[dirSE])[ks];
+ distribution.f[vf::lbm::dir::MPZ] = (distribution_references.f[dirNW])[kw];
+ distribution.f[vf::lbm::dir::PZP] = (distribution_references.f[dirTE])[k];
+ distribution.f[vf::lbm::dir::MZM] = (distribution_references.f[dirBW])[kbw];
+ distribution.f[vf::lbm::dir::PZM] = (distribution_references.f[dirBE])[kb];
+ distribution.f[vf::lbm::dir::MZP] = (distribution_references.f[dirTW])[kw];
+ distribution.f[vf::lbm::dir::ZPP] = (distribution_references.f[dirTN])[k];
+ distribution.f[vf::lbm::dir::ZMM] = (distribution_references.f[dirBS])[kbs];
+ distribution.f[vf::lbm::dir::ZPM] = (distribution_references.f[dirBN])[kb];
+ distribution.f[vf::lbm::dir::ZMP] = (distribution_references.f[dirTS])[ks];
+ distribution.f[vf::lbm::dir::PPP] = (distribution_references.f[dirTNE])[k];
+ distribution.f[vf::lbm::dir::MPP] = (distribution_references.f[dirTNW])[kw];
+ distribution.f[vf::lbm::dir::PMP] = (distribution_references.f[dirTSE])[ks];
+ distribution.f[vf::lbm::dir::MMP] = (distribution_references.f[dirTSW])[ksw];
+ distribution.f[vf::lbm::dir::PPM] = (distribution_references.f[dirBNE])[kb];
+ distribution.f[vf::lbm::dir::MPM] = (distribution_references.f[dirBNW])[kbw];
+ distribution.f[vf::lbm::dir::PMM] = (distribution_references.f[dirBSE])[kbs];
+ distribution.f[vf::lbm::dir::MMM] = (distribution_references.f[dirBSW])[kbsw];
+ distribution.f[vf::lbm::dir::ZZZ] = (distribution_references.f[dirREST])[k];
}
__device__ void DistributionWrapper::write()
{
- (dist.f[dirE])[k] = distribution.f[vf::lbm::dir::PZZ];
- (dist.f[dirW])[kw] = distribution.f[vf::lbm::dir::MZZ];
- (dist.f[dirN])[k] = distribution.f[vf::lbm::dir::ZPZ];
- (dist.f[dirS])[ks] = distribution.f[vf::lbm::dir::ZMZ];
- (dist.f[dirT])[k] = distribution.f[vf::lbm::dir::ZZP];
- (dist.f[dirB])[kb] = distribution.f[vf::lbm::dir::ZZM];
- (dist.f[dirNE])[k] = distribution.f[vf::lbm::dir::PPZ];
- (dist.f[dirSW])[ksw] = distribution.f[vf::lbm::dir::MMZ];
- (dist.f[dirSE])[ks] = distribution.f[vf::lbm::dir::PMZ];
- (dist.f[dirNW])[kw] = distribution.f[vf::lbm::dir::MPZ];
- (dist.f[dirTE])[k] = distribution.f[vf::lbm::dir::PZP];
- (dist.f[dirBW])[kbw] = distribution.f[vf::lbm::dir::MZM];
- (dist.f[dirBE])[kb] = distribution.f[vf::lbm::dir::PZM];
- (dist.f[dirTW])[kw] = distribution.f[vf::lbm::dir::MZP];
- (dist.f[dirTN])[k] = distribution.f[vf::lbm::dir::ZPP];
- (dist.f[dirBS])[kbs] = distribution.f[vf::lbm::dir::ZMM];
- (dist.f[dirBN])[kb] = distribution.f[vf::lbm::dir::ZPM];
- (dist.f[dirTS])[ks] = distribution.f[vf::lbm::dir::ZMP];
- (dist.f[dirTNE])[k] = distribution.f[vf::lbm::dir::PPP];
- (dist.f[dirTNW])[kw] = distribution.f[vf::lbm::dir::MPP];
- (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[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];
+ (distribution_references.f[dirE])[k] = distribution.f[vf::lbm::dir::PZZ];
+ (distribution_references.f[dirW])[kw] = distribution.f[vf::lbm::dir::MZZ];
+ (distribution_references.f[dirN])[k] = distribution.f[vf::lbm::dir::ZPZ];
+ (distribution_references.f[dirS])[ks] = distribution.f[vf::lbm::dir::ZMZ];
+ (distribution_references.f[dirT])[k] = distribution.f[vf::lbm::dir::ZZP];
+ (distribution_references.f[dirB])[kb] = distribution.f[vf::lbm::dir::ZZM];
+ (distribution_references.f[dirNE])[k] = distribution.f[vf::lbm::dir::PPZ];
+ (distribution_references.f[dirSW])[ksw] = distribution.f[vf::lbm::dir::MMZ];
+ (distribution_references.f[dirSE])[ks] = distribution.f[vf::lbm::dir::PMZ];
+ (distribution_references.f[dirNW])[kw] = distribution.f[vf::lbm::dir::MPZ];
+ (distribution_references.f[dirTE])[k] = distribution.f[vf::lbm::dir::PZP];
+ (distribution_references.f[dirBW])[kbw] = distribution.f[vf::lbm::dir::MZM];
+ (distribution_references.f[dirBE])[kb] = distribution.f[vf::lbm::dir::PZM];
+ (distribution_references.f[dirTW])[kw] = distribution.f[vf::lbm::dir::MZP];
+ (distribution_references.f[dirTN])[k] = distribution.f[vf::lbm::dir::ZPP];
+ (distribution_references.f[dirBS])[kbs] = distribution.f[vf::lbm::dir::ZMM];
+ (distribution_references.f[dirBN])[kb] = distribution.f[vf::lbm::dir::ZPM];
+ (distribution_references.f[dirTS])[ks] = distribution.f[vf::lbm::dir::ZMP];
+ (distribution_references.f[dirTNE])[k] = distribution.f[vf::lbm::dir::PPP];
+ (distribution_references.f[dirTNW])[kw] = distribution.f[vf::lbm::dir::MPP];
+ (distribution_references.f[dirTSE])[ks] = distribution.f[vf::lbm::dir::PMP];
+ (distribution_references.f[dirTSW])[ksw] = distribution.f[vf::lbm::dir::MMP];
+ (distribution_references.f[dirBNE])[kb] = distribution.f[vf::lbm::dir::PPM];
+ (distribution_references.f[dirBNW])[kbw] = distribution.f[vf::lbm::dir::MPM];
+ (distribution_references.f[dirBSE])[kbs] = distribution.f[vf::lbm::dir::PMM];
+ (distribution_references.f[dirBSW])[kbsw] = distribution.f[vf::lbm::dir::MMM];
+ (distribution_references.f[dirREST])[k] = distribution.f[vf::lbm::dir::ZZZ];
}
__device__ unsigned int getNodeIndex()
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Utilities/DistributionHelper.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Utilities/DistributionHelper.cuh
index b657bc2d6cc49e0e0b26dcc5f9a0385f3b753da2..bbcb869ee963ed2b13c65c061e9d38d13aaafccd 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Utilities/DistributionHelper.cuh
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Utilities/DistributionHelper.cuh
@@ -35,7 +35,6 @@
#include "LBM/LB.h"
-
#include <lbm/CumulantChimeraK17.h>
namespace vf
@@ -43,8 +42,21 @@ namespace vf
namespace gpu
{
-__device__ __host__ Distributions27 getDistributions27(real* distributions, unsigned int size_Mat, bool isEvenTimestep);
-
+/**
+* Getting references to the 27 directions.
+* @params distributions 1D real* array containing all data (number of elements = 27 * matrix_size)
+* @params matrix_size number of discretizations nodes
+* @params isEvenTimestep: stored data dependent on timestep is based on the esoteric twist algorithm
+* @return a data struct containing the addresses to the 27 directions within the 1D distribution array
+*/
+__device__ __host__ DistributionReferences27 getDistributionReferences27(real* distributions, unsigned int matrix_size, bool isEvenTimestep);
+
+
+/**
+* Holds the references to all directions and the concrete distributions for a single node.
+* After instantiation the distributions are read to the member "distribution" from "distribution_references".
+* After computation the data can be written back to "distribution_references".
+*/
struct DistributionWrapper
{
__device__ DistributionWrapper(
@@ -60,8 +72,10 @@ struct DistributionWrapper
__device__ void write();
- Distributions27 dist;
+ // origin distributions to read from and write to after computation
+ DistributionReferences27 distribution_references;
+ // distribution pass to kernel computation
vf::lbm::Distribution27 distribution;
const uint k;
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Utilities/DistributionHelperTests.cpp b/src/gpu/VirtualFluids_GPU/Kernel/Utilities/DistributionHelperTests.cpp
index 0d0e620ad83fd43d5654c450a9fc97936e1483b4..46d3fe890fd2f412251044cce6ade51f08ba0185 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Utilities/DistributionHelperTests.cpp
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Utilities/DistributionHelperTests.cpp
@@ -22,7 +22,7 @@ TEST(DistributionHelperTests, getPointerToDistribution_WhenEvenTimeStep_ShouldBe
const uint size_Mat = 1;
const bool isEvenTimeStep = true;
- Distributions27 distribution_out = vf::gpu::getDistributions27(distributions_in, size_Mat, isEvenTimeStep);
+ Distributions27 distribution_out = vf::gpu::getDistributionReferences27(distributions_in, size_Mat, isEvenTimeStep);
EXPECT_THAT(*distribution_out.f[dirE], RealEq(distributions_in[dirE]));
EXPECT_THAT(*distribution_out.f[dirW], RealEq(distributions_in[dirW]));
@@ -61,7 +61,7 @@ TEST(DistributionHelperTests, getPointerToDistribution_WhenOddTimeStep_ShouldBeS
const int size_Mat = 1;
const bool isEvenTimeStep = false;
- Distributions27 distribution_out = vf::gpu::getDistributions27(distributions_in, size_Mat, isEvenTimeStep);
+ Distributions27 distribution_out = vf::gpu::getDistributionReferences27(distributions_in, size_Mat, isEvenTimeStep);
EXPECT_THAT(*distribution_out.f[dirW], RealEq(distributions_in[dirE]));
EXPECT_THAT(*distribution_out.f[dirE], RealEq(distributions_in[dirW]));
diff --git a/src/gpu/VirtualFluids_GPU/LBM/LB.h b/src/gpu/VirtualFluids_GPU/LBM/LB.h
index f5f91738739b95f897a1192cb7c1d77a6cfea9cb..52f932c7b751b0d31e9bf733a4f3b5c439a6755f 100644
--- a/src/gpu/VirtualFluids_GPU/LBM/LB.h
+++ b/src/gpu/VirtualFluids_GPU/LBM/LB.h
@@ -144,7 +144,7 @@ typedef struct Distri19{
// Distribution functions f 27
typedef struct Distri27{
real* f[27];
-} Distributions27;
+} Distributions27, DistributionReferences27;
//Q for second order BCs
typedef struct QforBC{