diff --git a/CMake/VirtualFluidsMacros.cmake b/CMake/VirtualFluidsMacros.cmake
index dfef8f98478eb90f5f8b6760f52b61b169e19e77..d0dba357409d1251595eb22b6fc646e5c1c7942b 100644
--- a/CMake/VirtualFluidsMacros.cmake
+++ b/CMake/VirtualFluidsMacros.cmake
@@ -152,6 +152,17 @@ function(vf_add_library)
ARCHIVE_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib"
PDB_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/bin")
+ # link time optimization
+ include(CheckIPOSupported)
+ check_ipo_supported(RESULT ipo_supported OUTPUT ipo_error)
+
+ if( ipo_supported )
+ status_lib("IPO / LTO enabled")
+ set_target_properties(${library_name} PROPERTIES INTERPROCEDURAL_OPTIMIZATION TRUE)
+ else()
+ status_lib("IPO / LTO not supported: <${ipo_error}>")
+ endif()
+
# clang-tidy
if(BUILD_VF_CLANG_TIDY)
find_program(CLANG_TIDY_PROGRAM NAMES clang-tidy)
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 7ac14b98432bc0951b2cbf1b5a4e6a608f8fb160..e6eb2b0020d2f60d27bfb0152989bc43957dec50 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -104,6 +104,7 @@ if(BUILD_VF_GPU)
message("CUDA Architecture: ${CMAKE_CUDA_ARCHITECTURES}")
endif()
+
#################################################################################
# COMMON LIBRARIES
#################################################################################
diff --git a/src/cpu/VirtualFluidsCore/LBM/CumulantK17LBMKernel.cpp b/src/cpu/VirtualFluidsCore/LBM/CumulantK17LBMKernel.cpp
index daed493b9cc1afddbd92acabcd551da0f463ea26..6fe89ee7a0727781db2af58be7050ed4b2d11b50 100644
--- a/src/cpu/VirtualFluidsCore/LBM/CumulantK17LBMKernel.cpp
+++ b/src/cpu/VirtualFluidsCore/LBM/CumulantK17LBMKernel.cpp
@@ -39,6 +39,8 @@
#include "Block3D.h"
#include "BCArray3D.h"
+#include <lbm/BackwardChimera.h>
+
#define PROOF_CORRECTNESS
using namespace UbMath;
@@ -256,39 +258,39 @@ void CumulantK17LBMKernel::calculate(int step)
//!
////////////////////////////////////////////////////////////////////////////////////
// Z - Dir
- forwardInverseChimeraWithK(mfaaa, mfaab, mfaac, vvz, vz2, c36, c1o36);
- forwardInverseChimeraWithK(mfaba, mfabb, mfabc, vvz, vz2, c9, c1o9);
- forwardInverseChimeraWithK(mfaca, mfacb, mfacc, vvz, vz2, c36, c1o36);
- forwardInverseChimeraWithK(mfbaa, mfbab, mfbac, vvz, vz2, c9, c1o9);
- forwardInverseChimeraWithK(mfbba, mfbbb, mfbbc, vvz, vz2, c9o4, c4o9);
- forwardInverseChimeraWithK(mfbca, mfbcb, mfbcc, vvz, vz2, c9, c1o9);
- forwardInverseChimeraWithK(mfcaa, mfcab, mfcac, vvz, vz2, c36, c1o36);
- forwardInverseChimeraWithK(mfcba, mfcbb, mfcbc, vvz, vz2, c9, c1o9);
- forwardInverseChimeraWithK(mfcca, mfccb, mfccc, vvz, vz2, c36, c1o36);
+ VF::LBM::forwardInverseChimeraWithK(mfaaa, mfaab, mfaac, vvz, vz2, c36, c1o36);
+ VF::LBM::forwardInverseChimeraWithK(mfaba, mfabb, mfabc, vvz, vz2, c9, c1o9);
+ VF::LBM::forwardInverseChimeraWithK(mfaca, mfacb, mfacc, vvz, vz2, c36, c1o36);
+ VF::LBM::forwardInverseChimeraWithK(mfbaa, mfbab, mfbac, vvz, vz2, c9, c1o9);
+ VF::LBM::forwardInverseChimeraWithK(mfbba, mfbbb, mfbbc, vvz, vz2, c9o4, c4o9);
+ VF::LBM::forwardInverseChimeraWithK(mfbca, mfbcb, mfbcc, vvz, vz2, c9, c1o9);
+ VF::LBM::forwardInverseChimeraWithK(mfcaa, mfcab, mfcac, vvz, vz2, c36, c1o36);
+ VF::LBM::forwardInverseChimeraWithK(mfcba, mfcbb, mfcbc, vvz, vz2, c9, c1o9);
+ VF::LBM::forwardInverseChimeraWithK(mfcca, mfccb, mfccc, vvz, vz2, c36, c1o36);
////////////////////////////////////////////////////////////////////////////////////
// Y - Dir
- forwardInverseChimeraWithK(mfaaa, mfaba, mfaca, vvy, vy2, c6, c1o6);
- forwardChimera(mfaab, mfabb, mfacb, vvy, vy2);
- forwardInverseChimeraWithK(mfaac, mfabc, mfacc, vvy, vy2, c18, c1o18);
- forwardInverseChimeraWithK(mfbaa, mfbba, mfbca, vvy, vy2, c3o2, c2o3);
- forwardChimera(mfbab, mfbbb, mfbcb, vvy, vy2);
- forwardInverseChimeraWithK(mfbac, mfbbc, mfbcc, vvy, vy2, c9o2, c2o9);
- forwardInverseChimeraWithK(mfcaa, mfcba, mfcca, vvy, vy2, c6, c1o6);
- forwardChimera(mfcab, mfcbb, mfccb, vvy, vy2);
- forwardInverseChimeraWithK(mfcac, mfcbc, mfccc, vvy, vy2, c18, c1o18);
+ VF::LBM::forwardInverseChimeraWithK(mfaaa, mfaba, mfaca, vvy, vy2, c6, c1o6);
+ VF::LBM::forwardChimera(mfaab, mfabb, mfacb, vvy, vy2);
+ VF::LBM::forwardInverseChimeraWithK(mfaac, mfabc, mfacc, vvy, vy2, c18, 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, c6, c1o6);
+ VF::LBM::forwardChimera(mfcab, mfcbb, mfccb, vvy, vy2);
+ VF::LBM::forwardInverseChimeraWithK(mfcac, mfcbc, mfccc, vvy, vy2, c18, c1o18);
////////////////////////////////////////////////////////////////////////////////////
// X - Dir
- forwardInverseChimeraWithK(mfaaa, mfbaa, mfcaa, vvx, vx2, c1, c1);
- forwardChimera(mfaba, mfbba, mfcba, vvx, vx2);
- forwardInverseChimeraWithK(mfaca, mfbca, mfcca, vvx, vx2, c3, c1o3);
- forwardChimera(mfaab, mfbab, mfcab, vvx, vx2);
- forwardChimera(mfabb, mfbbb, mfcbb, vvx, vx2);
- forwardChimera(mfacb, mfbcb, mfccb, vvx, vx2);
- forwardInverseChimeraWithK(mfaac, mfbac, mfcac, vvx, vx2, c3, c1o3);
- forwardChimera(mfabc, mfbbc, mfcbc, vvx, vx2);
- forwardInverseChimeraWithK(mfacc, mfbcc, mfccc, vvx, vx2, c9, c1o9);
+ VF::LBM::forwardInverseChimeraWithK(mfaaa, mfbaa, mfcaa, vvx, vx2, c1, c1);
+ VF::LBM::forwardChimera(mfaba, mfbba, mfcba, vvx, vx2);
+ VF::LBM::forwardInverseChimeraWithK(mfaca, mfbca, mfcca, vvx, vx2, c3, 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, c3, c1o3);
+ VF::LBM::forwardChimera(mfabc, mfbbc, mfcbc, vvx, vx2);
+ VF::LBM::forwardInverseChimeraWithK(mfacc, mfbcc, mfccc, vvx, vx2, c9, c1o9);
////////////////////////////////////////////////////////////////////////////////////
//! - Setting relaxation rates for non-hydrodynamic cumulants (default values). Variable names and equations according to
@@ -536,39 +538,39 @@ void CumulantK17LBMKernel::calculate(int step)
//!
////////////////////////////////////////////////////////////////////////////////////
// X - Dir
- backwardInverseChimeraWithK(mfaaa, mfbaa, mfcaa, vvx, vx2, c1, c1);
- backwardChimera(mfaba, mfbba, mfcba, vvx, vx2);
- backwardInverseChimeraWithK(mfaca, mfbca, mfcca, vvx, vx2, c3, c1o3);
- backwardChimera(mfaab, mfbab, mfcab, vvx, vx2);
- backwardChimera(mfabb, mfbbb, mfcbb, vvx, vx2);
- backwardChimera(mfacb, mfbcb, mfccb, vvx, vx2);
- backwardInverseChimeraWithK(mfaac, mfbac, mfcac, vvx, vx2, c3, c1o3);
- backwardChimera(mfabc, mfbbc, mfcbc, vvx, vx2);
- backwardInverseChimeraWithK(mfacc, mfbcc, mfccc, vvx, vx2, c9, c1o9);
+ VF::LBM::backwardInverseChimeraWithK(mfaaa, mfbaa, mfcaa, vvx, vx2, c1, c1);
+ VF::LBM::backwardChimera(mfaba, mfbba, mfcba, vvx, vx2);
+ VF::LBM::backwardInverseChimeraWithK(mfaca, mfbca, mfcca, vvx, vx2, c3, 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, c3, c1o3);
+ VF::LBM::backwardChimera(mfabc, mfbbc, mfcbc, vvx, vx2);
+ VF::LBM::backwardInverseChimeraWithK(mfacc, mfbcc, mfccc, vvx, vx2, c9, c1o9);
////////////////////////////////////////////////////////////////////////////////////
// Y - Dir
- backwardInverseChimeraWithK(mfaaa, mfaba, mfaca, vvy, vy2, c6, c1o6);
- backwardChimera(mfaab, mfabb, mfacb, vvy, vy2);
- backwardInverseChimeraWithK(mfaac, mfabc, mfacc, vvy, vy2, c18, c1o18);
- backwardInverseChimeraWithK(mfbaa, mfbba, mfbca, vvy, vy2, c3o2, c2o3);
- backwardChimera(mfbab, mfbbb, mfbcb, vvy, vy2);
- backwardInverseChimeraWithK(mfbac, mfbbc, mfbcc, vvy, vy2, c9o2, c2o9);
- backwardInverseChimeraWithK(mfcaa, mfcba, mfcca, vvy, vy2, c6, c1o6);
- backwardChimera(mfcab, mfcbb, mfccb, vvy, vy2);
- backwardInverseChimeraWithK(mfcac, mfcbc, mfccc, vvy, vy2, c18, c1o18);
+ VF::LBM::backwardInverseChimeraWithK(mfaaa, mfaba, mfaca, vvy, vy2, c6, c1o6);
+ VF::LBM::backwardChimera(mfaab, mfabb, mfacb, vvy, vy2);
+ VF::LBM::backwardInverseChimeraWithK(mfaac, mfabc, mfacc, vvy, vy2, c18, 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, c6, c1o6);
+ VF::LBM::backwardChimera(mfcab, mfcbb, mfccb, vvy, vy2);
+ VF::LBM::backwardInverseChimeraWithK(mfcac, mfcbc, mfccc, vvy, vy2, c18, c1o18);
////////////////////////////////////////////////////////////////////////////////////
// Z - Dir
- backwardInverseChimeraWithK(mfaaa, mfaab, mfaac, vvz, vz2, c36, c1o36);
- backwardInverseChimeraWithK(mfaba, mfabb, mfabc, vvz, vz2, c9, c1o9);
- backwardInverseChimeraWithK(mfaca, mfacb, mfacc, vvz, vz2, c36, c1o36);
- backwardInverseChimeraWithK(mfbaa, mfbab, mfbac, vvz, vz2, c9, c1o9);
- backwardInverseChimeraWithK(mfbba, mfbbb, mfbbc, vvz, vz2, c9o4, c4o9);
- backwardInverseChimeraWithK(mfbca, mfbcb, mfbcc, vvz, vz2, c9, c1o9);
- backwardInverseChimeraWithK(mfcaa, mfcab, mfcac, vvz, vz2, c36, c1o36);
- backwardInverseChimeraWithK(mfcba, mfcbb, mfcbc, vvz, vz2, c9, c1o9);
- backwardInverseChimeraWithK(mfcca, mfccb, mfccc, vvz, vz2, c36, c1o36);
+ VF::LBM::backwardInverseChimeraWithK(mfaaa, mfaab, mfaac, vvz, vz2, c36, c1o36);
+ VF::LBM::backwardInverseChimeraWithK(mfaba, mfabb, mfabc, vvz, vz2, c9, c1o9);
+ VF::LBM::backwardInverseChimeraWithK(mfaca, mfacb, mfacc, vvz, vz2, c36, c1o36);
+ VF::LBM::backwardInverseChimeraWithK(mfbaa, mfbab, mfbac, vvz, vz2, c9, c1o9);
+ VF::LBM::backwardInverseChimeraWithK(mfbba, mfbbb, mfbbc, vvz, vz2, c9o4, c4o9);
+ VF::LBM::backwardInverseChimeraWithK(mfbca, mfbcb, mfbcc, vvz, vz2, c9, c1o9);
+ VF::LBM::backwardInverseChimeraWithK(mfcaa, mfcab, mfcac, vvz, vz2, c36, c1o36);
+ VF::LBM::backwardInverseChimeraWithK(mfcba, mfcbb, mfcbc, vvz, vz2, c9, c1o9);
+ VF::LBM::backwardInverseChimeraWithK(mfcca, mfccb, mfccc, vvz, vz2, c36, c1o36);
////////////////////////////////////////////////////////////////////////////////////
//////////////////////////////////////////////////////////////////////////
diff --git a/src/cpu/VirtualFluidsCore/LBM/LBMKernel.h b/src/cpu/VirtualFluidsCore/LBM/LBMKernel.h
index 5c3e067656e2fc5cb06046e28c0442b15c180050..f4d96d946a9f3bab367e7bc935b54352e09063a6 100644
--- a/src/cpu/VirtualFluidsCore/LBM/LBMKernel.h
+++ b/src/cpu/VirtualFluidsCore/LBM/LBMKernel.h
@@ -53,6 +53,7 @@ public:
public:
LBMKernel();
+ virtual ~LBMKernel() = default;
virtual SPtr<LBMKernel> clone() = 0;
diff --git a/src/cpu/simulationconfig/include/simulationconfig/AbstractLBMSystem.h b/src/cpu/simulationconfig/include/simulationconfig/AbstractLBMSystem.h
index 0c3ac72db2d3cb3488860227171f2fd27600c028..156649c6fd3c060117e8247efab2765f56f7c77b 100644
--- a/src/cpu/simulationconfig/include/simulationconfig/AbstractLBMSystem.h
+++ b/src/cpu/simulationconfig/include/simulationconfig/AbstractLBMSystem.h
@@ -7,6 +7,8 @@
class AbstractLBMSystem {
public:
+ virtual ~AbstractLBMSystem() = default;
+
virtual int getNumberOfDirections() = 0;
virtual std::shared_ptr<Interactor3D> makeInteractor() = 0;
diff --git a/src/cpu/simulationconfig/include/simulationconfig/KernelFactory.h b/src/cpu/simulationconfig/include/simulationconfig/KernelFactory.h
index 6b151d96f9c538bcfbf3aeb693afd56e36e6cf35..6f984849bea91a832d0d778107feacb3cfdbdca6 100644
--- a/src/cpu/simulationconfig/include/simulationconfig/KernelFactory.h
+++ b/src/cpu/simulationconfig/include/simulationconfig/KernelFactory.h
@@ -20,6 +20,7 @@ public:
};
KernelFactory() = default;
+ virtual ~KernelFactory() = default;
std::shared_ptr<LBMKernel> makeKernel(KernelType kernelType);
diff --git a/src/cpu/simulationconfig/src/Simulation.cpp b/src/cpu/simulationconfig/src/Simulation.cpp
index 6d80a1cfa2cc62d2575d1e8f67193b879f217b90..e876f65af4f7388eaeaa9d087b14e398b0452bbd 100644
--- a/src/cpu/simulationconfig/src/Simulation.cpp
+++ b/src/cpu/simulationconfig/src/Simulation.cpp
@@ -99,7 +99,7 @@ void Simulation::run()
grid->setPeriodicX2(gridParameters->periodicBoundaryInX2);
grid->setPeriodicX3(gridParameters->periodicBoundaryInX3);
- int &numberOfNodesInReferenceDirection = gridParameters->numberOfNodesPerDirection[gridParameters->referenceDirectionIndex];
+ //int &numberOfNodesInReferenceDirection = gridParameters->numberOfNodesPerDirection[gridParameters->referenceDirectionIndex];
std::shared_ptr<LBMUnitConverter> converter = makeLBMUnitConverter();
int &nodesInX1 = gridParameters->numberOfNodesPerDirection[0];
@@ -132,7 +132,7 @@ void Simulation::run()
grid->accept(kernelVisitor);
intHelper.setBC();
- double bulkViscosity = physicalParameters->latticeViscosity * physicalParameters->bulkViscosityFactor;
+ //double bulkViscosity = physicalParameters->latticeViscosity * physicalParameters->bulkViscosityFactor;
//auto iProcessor = std::make_shared<CompressibleOffsetMomentsInterpolationProcessor>();
//iProcessor->setBulkViscosity(physicalParameters->latticeViscosity, bulkViscosity);
diff --git a/src/lbm/BackwardChimera.cpp b/src/lbm/BackwardChimera.cpp
new file mode 100644
index 0000000000000000000000000000000000000000..38b81322906dcf7099d1392074d57767fd87d6ca
--- /dev/null
+++ b/src/lbm/BackwardChimera.cpp
@@ -0,0 +1,49 @@
+#include "BackwardChimera.h
+
+#include <basics/Core/RealConstants.h>
+
+
+void VF::LBM::forwardInverseChimeraWithK(real &mfa, real &mfb, real &mfc, real vv,
+ real v2, real Kinverse, real K)
+{
+ const real m2 = mfa + mfc;
+ const real m1 = mfc - mfa;
+ const real m0 = m2 + mfb;
+
+ mfa = m0;
+ m0 *= Kinverse;
+ m0 += c1;
+ mfb = (m1 * Kinverse - m0 * vv) * K;
+ mfc = ((m2 - c2 * m1 * vv) * Kinverse + v2 * m0) * K;
+}
+
+void VF::LBM::backwardInverseChimeraWithK(real &mfa, real &mfb, real &mfc, real vv,
+ real v2, real Kinverse, real K)
+{
+ const real m0 = (((mfc - mfb) * c1o2 + mfb * vv) * Kinverse + (mfa * Kinverse + c1) * (v2 - vv) * c1o2) * K;
+ const real m1 = (((mfa - mfc) - c2 * mfb * vv) * Kinverse + (mfa * Kinverse + c1) * (-v2)) * K;
+
+ mfc = (((mfc + mfb) * c1o2 + mfb * vv) * Kinverse + (mfa * Kinverse + c1) * (v2 + vv) * c1o2) * K;
+ mfa = m0;
+ mfb = m1;
+}
+
+void VF::LBM::forwardChimera(real &mfa, real &mfb, real &mfc, real vv, real v2)
+{
+ const real m1 = (mfa + mfc) + mfb;
+ const real m2 = mfc - mfa;
+
+ mfc = (mfc + mfa) + (v2 * m1 - c2 * vv * m2);
+ mfb = m2 - vv * m1;
+ mfa = m1;
+}
+
+void VF::LBM::backwardChimera(real &mfa, real &mfb, real &mfc, real vv, real v2)
+{
+ const real ma = (mfc + mfa * (v2 - vv)) * c1o2 + mfb * (vv - c1o2);
+ const real mb = ((mfa - mfc) - mfa * v2) - c2 * mfb * vv;
+
+ mfc = (mfc + mfa * (v2 + vv)) * c1o2 + mfb * (vv + c1o2);
+ mfb = mb;
+ mfa = ma;
+}
diff --git a/src/lbm/BackwardChimera.h b/src/lbm/BackwardChimera.h
new file mode 100644
index 0000000000000000000000000000000000000000..9243fcbfeb29e4a18a5ec75de4d8d23dd9c7b236
--- /dev/null
+++ b/src/lbm/BackwardChimera.h
@@ -0,0 +1,52 @@
+#ifndef LBM_BACKWARDCHIMERA_H
+#define LBM_BACKWARDCHIMERA_H
+
+#ifdef __CUDACC__
+#include <cuda_runtime.h>
+#else
+#ifndef __host__
+#define __host__
+#endif
+#ifndef __device__
+#define __device__
+#endif
+#endif
+
+#include "Core/DataTypes.h"
+
+namespace VF
+{
+namespace LBM
+{
+
+////////////////////////////////////////////////////////////////////////////////
+//! \brief forward chimera transformation \ref forwardInverseChimeraWithK
+//! Transformation from distributions to central moments according to Eq. (6)-(14) 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> Modified for lower round-off errors.
+////////////////////////////////////////////////////////////////////////////////
+__host__ __device__ void forwardInverseChimeraWithK(real &mfa, real &mfb, real &mfc, real vv,
+ real v2, real Kinverse, real K);
+////////////////////////////////////////////////////////////////////////////////
+//! \brief backward chimera transformation \ref backwardInverseChimeraWithK
+//! Transformation from central moments to distributions according to Eq. (57)-(65) 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> ] Modified for lower round-off errors.
+////////////////////////////////////////////////////////////////////////////////
+__host__ __device__ void backwardInverseChimeraWithK(real &mfa, real &mfb, real &mfc, real vv,
+ real v2, real Kinverse, real K);
+////////////////////////////////////////////////////////////////////////////////
+//! \brief forward chimera transformation \ref forwardChimera
+//! Transformation from distributions to central moments according to Eq. (6)-(14) 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> for \f$ K_{abc}=0 \f$. This is to avoid unnessary floating point operations. Modified for lower round-off errors.
+////////////////////////////////////////////////////////////////////////////////
+__host__ __device__ void forwardChimera(real &mfa, real &mfb, real &mfc, real vv, real v2);
+////////////////////////////////////////////////////////////////////////////////
+//! \brief backward chimera transformation \ref backwardChimera
+//! Transformation from central moments to distributions according to Eq. (57)-(65) 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> for \f$ K_{abc}=0 \f$. This is to avoid unnessary floating point operations. Modified for lower round-off errors.
+////////////////////////////////////////////////////////////////////////////////
+__host__ __device__ void backwardChimera(real &mfa, real &mfb, real &mfc, real vv, real v2);
+
+}
+}
+
+#endif
diff --git a/src/lbm/MacroscopicQuantities.cpp b/src/lbm/MacroscopicQuantities.cpp
index 84813be617e8c3d50207a4ad3e1cec2d2c0e6b12..3a465d5f6b971461cd99035b9dd8d23159d73ac5 100644
--- a/src/lbm/MacroscopicQuantities.cpp
+++ b/src/lbm/MacroscopicQuantities.cpp
@@ -3,7 +3,7 @@
#include "D3Q27.h"
-__host__ __device__ real VF::LBM::getDensity(const real *const &f /*[27]*/)
+real VF::LBM::getDensity(const real *const &f /*[27]*/)
{
return ((f[DIR::TNE] + f[DIR::BSW]) + (f[DIR::TSE] + f[DIR::BNW])) + ((f[DIR::BSE] + f[DIR::TNW]) + (f[DIR::TSW] + f[DIR::BNE])) +
(((f[DIR::NE] + f[DIR::SW]) + (f[DIR::SE] + f[DIR::NW])) + ((f[DIR::TE] + f[DIR::BW]) + (f[DIR::BE] + f[DIR::TW])) +
diff --git a/src/lbm/cuda/CMakeLists.txt b/src/lbm/cuda/CMakeLists.txt
index f024461cf3d17d15f511b82b5518c979675acf28..6a5d42e83869299b5d627c26369daecaf7082a52 100644
--- a/src/lbm/cuda/CMakeLists.txt
+++ b/src/lbm/cuda/CMakeLists.txt
@@ -8,3 +8,4 @@ set_target_properties(lbmCuda PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
set_source_files_properties(../MacroscopicQuantities.cpp PROPERTIES LANGUAGE CUDA)
+set_source_files_properties(../BackwardChimera.cpp PROPERTIES LANGUAGE CUDA)
\ No newline at end of file