From 2c3f7d070d74aa022ca75823826be4dfca93a346 Mon Sep 17 00:00:00 2001
From: Anna Wellmann <a.wellmann@tu-braunschweig.de>
Date: Thu, 9 Nov 2023 14:56:08 +0000
Subject: [PATCH] Fix bugs after merge

---
 apps/gpu/RotatingGrid/CMakeLists.txt          |  2 +-
 apps/gpu/RotatingGrid/RotatingGrid.cpp        | 22 ++---
 src/gpu/core/GPU/CalcMac27.cu                 |  2 +-
 .../interpolateRotatingToStatic.cu            | 56 ++++++-------
 .../interpolateStaticToRotating.cu            | 56 ++++++-------
 .../LBM/GPUHelperFunctions/KernelUtilities.h  |  2 +-
 src/gpu/core/Output/NeighborDebugWriter.hpp   | 56 -------------
 src/gpu/core/Output/WriterUtilitiesTest.cpp   |  2 +-
 src/lbm/MacroscopicQuantities.h               | 84 +++++++++----------
 9 files changed, 113 insertions(+), 169 deletions(-)
 delete mode 100644 src/gpu/core/Output/NeighborDebugWriter.hpp

diff --git a/apps/gpu/RotatingGrid/CMakeLists.txt b/apps/gpu/RotatingGrid/CMakeLists.txt
index bd64b62d4..8e307a095 100644
--- a/apps/gpu/RotatingGrid/CMakeLists.txt
+++ b/apps/gpu/RotatingGrid/CMakeLists.txt
@@ -2,4 +2,4 @@ PROJECT(DrivenCavity LANGUAGES CXX)
 
 #LIST(APPEND CS_COMPILER_FLAGS_CXX "-DOMPI_SKIP_MPICXX" )
 
-vf_add_library(BUILDTYPE binary PRIVATE_LINK basics VirtualFluids_GPU GridGenerator FILES RotatingGrid.cpp)
+vf_add_library(BUILDTYPE binary PRIVATE_LINK basics gpu_core GridGenerator FILES RotatingGrid.cpp)
diff --git a/apps/gpu/RotatingGrid/RotatingGrid.cpp b/apps/gpu/RotatingGrid/RotatingGrid.cpp
index 2d95dd3fd..84b96bc81 100644
--- a/apps/gpu/RotatingGrid/RotatingGrid.cpp
+++ b/apps/gpu/RotatingGrid/RotatingGrid.cpp
@@ -56,17 +56,17 @@
 
 //////////////////////////////////////////////////////////////////////////
 
-#include "VirtualFluids_GPU/DataStructureInitializer/GridProvider.h"
-#include "VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.h"
-#include "VirtualFluids_GPU/Factories/BoundaryConditionFactory.h"
-#include "VirtualFluids_GPU/Factories/GridScalingFactory.h"
-#include "VirtualFluids_GPU/GPU/CudaMemoryManager.h"
-#include "VirtualFluids_GPU/Kernel/Utilities/KernelTypes.h"
-#include "VirtualFluids_GPU/LBM/Simulation.h"
-#include "VirtualFluids_GPU/Output/FileWriter.h"
-#include "VirtualFluids_GPU/Output/NeighborDebugWriter.h"
-#include "VirtualFluids_GPU/Parameter/Parameter.h"
-#include "VirtualFluids_GPU/Parameter/ParameterRotatingGrid.h"
+#include "core/DataStructureInitializer/GridProvider.h"
+#include "core/DataStructureInitializer/GridReaderGenerator/GridGenerator.h"
+#include "core/Factories/BoundaryConditionFactory.h"
+#include "core/Factories/GridScalingFactory.h"
+#include "core/GPU/CudaMemoryManager.h"
+#include "core/Kernel/KernelTypes.h"
+#include "core/LBM/Simulation.h"
+#include "core/Output/FileWriter.h"
+#include "core/Output/NeighborDebugWriter.h"
+#include "core/Parameter/Parameter.h"
+#include "core/Parameter/ParameterRotatingGrid.h"
 
 //////////////////////////////////////////////////////////////////////////
 
diff --git a/src/gpu/core/GPU/CalcMac27.cu b/src/gpu/core/GPU/CalcMac27.cu
index 9c9502d50..046362319 100644
--- a/src/gpu/core/GPU/CalcMac27.cu
+++ b/src/gpu/core/GPU/CalcMac27.cu
@@ -291,7 +291,7 @@ __inline__ __device__ void calcMacCompSP27ForNode(real* vxD, real* vyD, real* vz
     vf::gpu::ListIndices listIndices(nodeIndex, neighborX, neighborY, neighborZ);
 
     real distribution[27];
-    vf::gpu::read(distribution, dist, listIndices);
+    vf::gpu::getPreCollisionDistribution(distribution, dist, listIndices);
 
     rhoD[nodeIndex] = vf::lbm::getDensity(distribution);
     vxD[nodeIndex] = vf::lbm::getCompressibleVelocityX1(distribution, rhoD[nodeIndex]);
diff --git a/src/gpu/core/GPU/GridScaling/interpolateRotatingToStatic.cu b/src/gpu/core/GPU/GridScaling/interpolateRotatingToStatic.cu
index b2d86b05d..ff673acb5 100644
--- a/src/gpu/core/GPU/GridScaling/interpolateRotatingToStatic.cu
+++ b/src/gpu/core/GPU/GridScaling/interpolateRotatingToStatic.cu
@@ -475,33 +475,33 @@ __global__ void interpolateRotatingToStatic(
     // ////////////////////////////////////////////////////////////////////////////////
 
     real fStatic[27];
-    fStatic[DIR_000] = f000;
-    fStatic[DIR_P00] = fP00;
-    fStatic[DIR_M00] = fM00;
-    fStatic[DIR_0P0] = f0P0;
-    fStatic[DIR_0M0] = f0M0;
-    fStatic[DIR_00P] = f00P;
-    fStatic[DIR_00M] = f00M;
-    fStatic[DIR_PP0] = fPP0;
-    fStatic[DIR_MM0] = fMM0;
-    fStatic[DIR_PM0] = fPM0;
-    fStatic[DIR_MP0] = fMP0;
-    fStatic[DIR_P0P] = fP0P;
-    fStatic[DIR_M0M] = fM0M;
-    fStatic[DIR_P0M] = fP0M;
-    fStatic[DIR_M0P] = fM0P;
-    fStatic[DIR_0PP] = f0PP;
-    fStatic[DIR_0MM] = f0MM;
-    fStatic[DIR_0PM] = f0PM;
-    fStatic[DIR_0MP] = f0MP;
-    fStatic[DIR_PPP] = fPPP;
-    fStatic[DIR_MPP] = fMPP;
-    fStatic[DIR_PMP] = fPMP;
-    fStatic[DIR_MMP] = fMMP;
-    fStatic[DIR_PPM] = fPPM;
-    fStatic[DIR_MPM] = fMPM;
-    fStatic[DIR_PMM] = fPMM;
-    fStatic[DIR_MMM] = fMMM;
+    fStatic[d000] = f000;
+    fStatic[dP00] = fP00;
+    fStatic[dM00] = fM00;
+    fStatic[d0P0] = f0P0;
+    fStatic[d0M0] = f0M0;
+    fStatic[d00P] = f00P;
+    fStatic[d00M] = f00M;
+    fStatic[dPP0] = fPP0;
+    fStatic[dMM0] = fMM0;
+    fStatic[dPM0] = fPM0;
+    fStatic[dMP0] = fMP0;
+    fStatic[dP0P] = fP0P;
+    fStatic[dM0M] = fM0M;
+    fStatic[dP0M] = fP0M;
+    fStatic[dM0P] = fM0P;
+    fStatic[d0PP] = f0PP;
+    fStatic[d0MM] = f0MM;
+    fStatic[d0PM] = f0PM;
+    fStatic[d0MP] = f0MP;
+    fStatic[dPPP] = fPPP;
+    fStatic[dMPP] = fMPP;
+    fStatic[dPMP] = fPMP;
+    fStatic[dMMP] = fMMP;
+    fStatic[dPPM] = fPPM;
+    fStatic[dMPM] = fMPM;
+    fStatic[dPMM] = fPMM;
+    fStatic[dMMM] = fMMM;
 
     // get distribution pointers for destination node
     Distributions27 distStatic;
@@ -509,7 +509,7 @@ __global__ void interpolateRotatingToStatic(
 
     // write
     vf::gpu::ListIndices indicesStaticForWriting(destinationIndex, neighborXstatic, neighborYstatic, neighborZstatic);
-    vf::gpu::write(distStatic, indicesStaticForWriting, fStatic);
+    vf::gpu::setPreCollisionDistribution(distStatic, indicesStaticForWriting, fStatic);
 }
 
 __global__ void updateGlobalCoordinates(
diff --git a/src/gpu/core/GPU/GridScaling/interpolateStaticToRotating.cu b/src/gpu/core/GPU/GridScaling/interpolateStaticToRotating.cu
index 6e80893b4..508826e3d 100644
--- a/src/gpu/core/GPU/GridScaling/interpolateStaticToRotating.cu
+++ b/src/gpu/core/GPU/GridScaling/interpolateStaticToRotating.cu
@@ -346,33 +346,33 @@ __global__ void interpolateStaticToRotating(
     // ////////////////////////////////////////////////////////////////////////////////
 
     real fRotating[27];
-    fRotating[DIR_000] = f000;
-    fRotating[DIR_P00] = fP00;
-    fRotating[DIR_M00] = fM00;
-    fRotating[DIR_0P0] = f0P0;
-    fRotating[DIR_0M0] = f0M0;
-    fRotating[DIR_00P] = f00P;
-    fRotating[DIR_00M] = f00M;
-    fRotating[DIR_PP0] = fPP0;
-    fRotating[DIR_MM0] = fMM0;
-    fRotating[DIR_PM0] = fPM0;
-    fRotating[DIR_MP0] = fMP0;
-    fRotating[DIR_P0P] = fP0P;
-    fRotating[DIR_M0M] = fM0M;
-    fRotating[DIR_P0M] = fP0M;
-    fRotating[DIR_M0P] = fM0P;
-    fRotating[DIR_0PP] = f0PP;
-    fRotating[DIR_0MM] = f0MM;
-    fRotating[DIR_0PM] = f0PM;
-    fRotating[DIR_0MP] = f0MP;
-    fRotating[DIR_PPP] = fPPP;
-    fRotating[DIR_MPP] = fMPP;
-    fRotating[DIR_PMP] = fPMP;
-    fRotating[DIR_MMP] = fMMP;
-    fRotating[DIR_PPM] = fPPM;
-    fRotating[DIR_MPM] = fMPM;
-    fRotating[DIR_PMM] = fPMM;
-    fRotating[DIR_MMM] = fMMM;
+    fRotating[d000] = f000;
+    fRotating[dP00] = fP00;
+    fRotating[dM00] = fM00;
+    fRotating[d0P0] = f0P0;
+    fRotating[d0M0] = f0M0;
+    fRotating[d00P] = f00P;
+    fRotating[d00M] = f00M;
+    fRotating[dPP0] = fPP0;
+    fRotating[dMM0] = fMM0;
+    fRotating[dPM0] = fPM0;
+    fRotating[dMP0] = fMP0;
+    fRotating[dP0P] = fP0P;
+    fRotating[dM0M] = fM0M;
+    fRotating[dP0M] = fP0M;
+    fRotating[dM0P] = fM0P;
+    fRotating[d0PP] = f0PP;
+    fRotating[d0MM] = f0MM;
+    fRotating[d0PM] = f0PM;
+    fRotating[d0MP] = f0MP;
+    fRotating[dPPP] = fPPP;
+    fRotating[dMPP] = fMPP;
+    fRotating[dPMP] = fPMP;
+    fRotating[dMMP] = fMMP;
+    fRotating[dPPM] = fPPM;
+    fRotating[dMPM] = fMPM;
+    fRotating[dPMM] = fPMM;
+    fRotating[dMMM] = fMMM;
 
     // get distribution pointers for destination node
     Distributions27 distRoating;
@@ -381,7 +381,7 @@ __global__ void interpolateStaticToRotating(
     // write
     vf::gpu::ListIndices indicesRotatingForWriting(destinationIndex, neighborXrotating, neighborYrotating,
                                                    neighborZrotating);
-    vf::gpu::write(distRoating, indicesRotatingForWriting, fRotating);
+    vf::gpu::setPreCollisionDistribution(distRoating, indicesRotatingForWriting, fRotating);
 }
 
 __global__ void traverseStaticToRotating(
diff --git a/src/gpu/core/LBM/GPUHelperFunctions/KernelUtilities.h b/src/gpu/core/LBM/GPUHelperFunctions/KernelUtilities.h
index ad75d197e..f0d64d1dc 100644
--- a/src/gpu/core/LBM/GPUHelperFunctions/KernelUtilities.h
+++ b/src/gpu/core/LBM/GPUHelperFunctions/KernelUtilities.h
@@ -231,7 +231,7 @@ __inline__ __device__ bool isValidFluidNode(uint nodeType)
 
 struct ListIndices
 {
-    __device__ ListIndices() = default;
+    __device__ ListIndices() {};
     __device__ ListIndices(unsigned int index, const unsigned int* neighborX, const unsigned int* neighborY,
                            const unsigned int* neighborZ)
     {
diff --git a/src/gpu/core/Output/NeighborDebugWriter.hpp b/src/gpu/core/Output/NeighborDebugWriter.hpp
deleted file mode 100644
index 71ee3eb1c..000000000
--- a/src/gpu/core/Output/NeighborDebugWriter.hpp
+++ /dev/null
@@ -1,56 +0,0 @@
-#include "NeighborDebugWriter.h"
-
-#include "LBM/LB.h"
-#include "Logger.h"
-#include "Parameter/Parameter.h"
-#include "basics/utilities/UbSystem.h"
-#include "gpu/GridGenerator/grid/NodeValues.h"
-#include "lbm/constants/D3Q27.h"
-#include <basics/writer/WbWriterVtkXmlBinary.h>
-
-#include "StringUtilities/StringUtil.h"
-#include "Utilities/FindNeighbors.h"
-
-void NeighborDebugWriter::writeNeighborLinkLines(LBMSimulationParameter *parH, int direction, const std::string &name,
-                                   WbWriter *writer)
-{
-    VF_LOG_INFO("Write node links in direction {}.", direction);
-
-    const unsigned long long numberOfNodes = parH->numberOfNodes;
-    std::vector<UbTupleFloat3> nodes;
-    nodes.reserve(numberOfNodes);
-    std::vector<UbTupleInt2> cells;
-    cells.reserve(numberOfNodes/2);
-
-    for (size_t position = 0; position < numberOfNodes; position++) {
-        if (parH->typeOfGridNode[position] != GEO_FLUID)
-            continue;
-
-        const double x1 = parH->coordinateX[position];
-        const double x2 = parH->coordinateY[position];
-        const double x3 = parH->coordinateZ[position];
-
-        const uint positionNeighbor = getNeighborIndex(parH, (uint)position, direction);
-
-        const double x1Neighbor = parH->coordinateX[positionNeighbor];
-        const double x2Neighbor = parH->coordinateY[positionNeighbor];
-        const double x3Neighbor = parH->coordinateZ[positionNeighbor];
-
-        nodes.emplace_back(float(x1), float(x2), float(x3));
-        nodes.emplace_back(float(x1Neighbor), float(x2Neighbor), float(x3Neighbor));
-
-        cells.emplace_back((int)nodes.size() - 2, (int)nodes.size() - 1);
-    }
-    writer->writeLines(name, nodes, cells);
-}
-
-void NeighborDebugWriter::writeNeighborLinkLinesDebug(Parameter *para)
-{
-    for (int level = 0; level <= para->getMaxLevel(); level++) {
-        for (size_t direction = vf::lbm::dir::STARTDIR; direction <= vf::lbm::dir::ENDDIR; direction++) {
-            const std::string fileName = para->getFName() + "_" + StringUtil::toString<int>(level) + "_Link_" +
-                                         std::to_string(direction) + "_Debug.vtk";
-            writeNeighborLinkLines(para->getParH(level).get(), (int)direction, fileName, WbWriterVtkXmlBinary::getInstance());
-        }
-    }
-}
\ No newline at end of file
diff --git a/src/gpu/core/Output/WriterUtilitiesTest.cpp b/src/gpu/core/Output/WriterUtilitiesTest.cpp
index bf39c44de..4f9032817 100644
--- a/src/gpu/core/Output/WriterUtilitiesTest.cpp
+++ b/src/gpu/core/Output/WriterUtilitiesTest.cpp
@@ -1,6 +1,6 @@
 
 #include "WriterUtilities.h"
-#include "gpu/VirtualFluids_GPU/Utilities/testUtilitiesGPU.h"
+#include "gpu/core/Utilities/testUtilitiesGPU.h"
 #include <gmock/gmock.h>
 #include <numeric>
 
diff --git a/src/lbm/MacroscopicQuantities.h b/src/lbm/MacroscopicQuantities.h
index 269bbe325..f76c8b73a 100644
--- a/src/lbm/MacroscopicQuantities.h
+++ b/src/lbm/MacroscopicQuantities.h
@@ -8,7 +8,7 @@
 #define __device__
 #endif
 
-#include <gpu/VirtualFluids_GPU/LBM/LB.h>
+#include <gpu/core/LBM/LB.h>
 #include <basics/DataTypes.h>
 #include <basics/constants/NumericConstants.h>
 
@@ -104,53 +104,53 @@ inline __host__ __device__ void getCompressibleMacroscopicValues(const real *con
 
 inline __host__ __device__ void getCompressibleMacroscopicValues(const Distributions27 &dist /*[27] for all nodes*/, uint nodeIndex, real &drho, real &vx1, real &vx2, real &vx3)
 {
-    drho = ((dist.f[dir::DIR_PPP][nodeIndex] + dist.f[dir::DIR_MMM][nodeIndex]) +
-            (dist.f[dir::DIR_PMP][nodeIndex] + dist.f[dir::DIR_MPM][nodeIndex])) +
-           ((dist.f[dir::DIR_PMM][nodeIndex] + dist.f[dir::DIR_MPP][nodeIndex]) +
-            (dist.f[dir::DIR_MMP][nodeIndex] + dist.f[dir::DIR_PPM][nodeIndex])) +
-           (((dist.f[dir::DIR_PP0][nodeIndex] + dist.f[dir::DIR_MM0][nodeIndex]) +
-             (dist.f[dir::DIR_PM0][nodeIndex] + dist.f[dir::DIR_MP0][nodeIndex])) +
-            ((dist.f[dir::DIR_P0P][nodeIndex] + dist.f[dir::DIR_M0M][nodeIndex]) +
-             (dist.f[dir::DIR_P0M][nodeIndex] + dist.f[dir::DIR_M0P][nodeIndex])) +
-            ((dist.f[dir::DIR_0PM][nodeIndex] + dist.f[dir::DIR_0MP][nodeIndex]) +
-             (dist.f[dir::DIR_0PP][nodeIndex] + dist.f[dir::DIR_0MM][nodeIndex]))) +
-           ((dist.f[dir::DIR_P00][nodeIndex] + dist.f[dir::DIR_M00][nodeIndex]) +
-            (dist.f[dir::DIR_0P0][nodeIndex] + dist.f[dir::DIR_0M0][nodeIndex]) +
-            (dist.f[dir::DIR_00P][nodeIndex] + dist.f[dir::DIR_00M][nodeIndex])) +
-           dist.f[dir::DIR_000][nodeIndex];
+    drho = ((dist.f[dir::dPPP][nodeIndex] + dist.f[dir::dMMM][nodeIndex]) +
+            (dist.f[dir::dPMP][nodeIndex] + dist.f[dir::dMPM][nodeIndex])) +
+           ((dist.f[dir::dPMM][nodeIndex] + dist.f[dir::dMPP][nodeIndex]) +
+            (dist.f[dir::dMMP][nodeIndex] + dist.f[dir::dPPM][nodeIndex])) +
+           (((dist.f[dir::dPP0][nodeIndex] + dist.f[dir::dMM0][nodeIndex]) +
+             (dist.f[dir::dPM0][nodeIndex] + dist.f[dir::dMP0][nodeIndex])) +
+            ((dist.f[dir::dP0P][nodeIndex] + dist.f[dir::dM0M][nodeIndex]) +
+             (dist.f[dir::dP0M][nodeIndex] + dist.f[dir::dM0P][nodeIndex])) +
+            ((dist.f[dir::d0PM][nodeIndex] + dist.f[dir::d0MP][nodeIndex]) +
+             (dist.f[dir::d0PP][nodeIndex] + dist.f[dir::d0MM][nodeIndex]))) +
+           ((dist.f[dir::dP00][nodeIndex] + dist.f[dir::dM00][nodeIndex]) +
+            (dist.f[dir::d0P0][nodeIndex] + dist.f[dir::d0M0][nodeIndex]) +
+            (dist.f[dir::d00P][nodeIndex] + dist.f[dir::d00M][nodeIndex])) +
+           dist.f[dir::d000][nodeIndex];
     real oneOverRho = vf::basics::constant::c1o1 / (drho + vf::basics::constant::c1o1);
 
-    vx1 = ((((dist.f[dir::DIR_PPP][nodeIndex] - dist.f[dir::DIR_MMM][nodeIndex]) +
-             (dist.f[dir::DIR_PMP][nodeIndex] - dist.f[dir::DIR_MPM][nodeIndex])) +
-            ((dist.f[dir::DIR_PMM][nodeIndex] - dist.f[dir::DIR_MPP][nodeIndex]) +
-             (dist.f[dir::DIR_PPM][nodeIndex] - dist.f[dir::DIR_MMP][nodeIndex]))) +
-           (((dist.f[dir::DIR_P0M][nodeIndex] - dist.f[dir::DIR_M0P][nodeIndex]) +
-             (dist.f[dir::DIR_P0P][nodeIndex] - dist.f[dir::DIR_M0M][nodeIndex])) +
-            ((dist.f[dir::DIR_PM0][nodeIndex] - dist.f[dir::DIR_MP0][nodeIndex]) +
-             (dist.f[dir::DIR_PP0][nodeIndex] - dist.f[dir::DIR_MM0][nodeIndex]))) +
-           (dist.f[dir::DIR_P00][nodeIndex] - dist.f[dir::DIR_M00][nodeIndex]));
+    vx1 = ((((dist.f[dir::dPPP][nodeIndex] - dist.f[dir::dMMM][nodeIndex]) +
+             (dist.f[dir::dPMP][nodeIndex] - dist.f[dir::dMPM][nodeIndex])) +
+            ((dist.f[dir::dPMM][nodeIndex] - dist.f[dir::dMPP][nodeIndex]) +
+             (dist.f[dir::dPPM][nodeIndex] - dist.f[dir::dMMP][nodeIndex]))) +
+           (((dist.f[dir::dP0M][nodeIndex] - dist.f[dir::dM0P][nodeIndex]) +
+             (dist.f[dir::dP0P][nodeIndex] - dist.f[dir::dM0M][nodeIndex])) +
+            ((dist.f[dir::dPM0][nodeIndex] - dist.f[dir::dMP0][nodeIndex]) +
+             (dist.f[dir::dPP0][nodeIndex] - dist.f[dir::dMM0][nodeIndex]))) +
+           (dist.f[dir::dP00][nodeIndex] - dist.f[dir::dM00][nodeIndex]));
     vx1 *= oneOverRho;
 
-    vx2 = ((((dist.f[dir::DIR_PPP][nodeIndex] - dist.f[dir::DIR_MMM][nodeIndex]) +
-             (dist.f[dir::DIR_MPM][nodeIndex] - dist.f[dir::DIR_PMP][nodeIndex])) +
-            ((dist.f[dir::DIR_MPP][nodeIndex] - dist.f[dir::DIR_PMM][nodeIndex]) +
-             (dist.f[dir::DIR_PPM][nodeIndex] - dist.f[dir::DIR_MMP][nodeIndex]))) +
-           (((dist.f[dir::DIR_0PM][nodeIndex] - dist.f[dir::DIR_0MP][nodeIndex]) +
-             (dist.f[dir::DIR_0PP][nodeIndex] - dist.f[dir::DIR_0MM][nodeIndex])) +
-            ((dist.f[dir::DIR_MP0][nodeIndex] - dist.f[dir::DIR_PM0][nodeIndex]) +
-             (dist.f[dir::DIR_PP0][nodeIndex] - dist.f[dir::DIR_MM0][nodeIndex]))) +
-           (dist.f[dir::DIR_0P0][nodeIndex] - dist.f[dir::DIR_0M0][nodeIndex]));
+    vx2 = ((((dist.f[dir::dPPP][nodeIndex] - dist.f[dir::dMMM][nodeIndex]) +
+             (dist.f[dir::dMPM][nodeIndex] - dist.f[dir::dPMP][nodeIndex])) +
+            ((dist.f[dir::dMPP][nodeIndex] - dist.f[dir::dPMM][nodeIndex]) +
+             (dist.f[dir::dPPM][nodeIndex] - dist.f[dir::dMMP][nodeIndex]))) +
+           (((dist.f[dir::d0PM][nodeIndex] - dist.f[dir::d0MP][nodeIndex]) +
+             (dist.f[dir::d0PP][nodeIndex] - dist.f[dir::d0MM][nodeIndex])) +
+            ((dist.f[dir::dMP0][nodeIndex] - dist.f[dir::dPM0][nodeIndex]) +
+             (dist.f[dir::dPP0][nodeIndex] - dist.f[dir::dMM0][nodeIndex]))) +
+           (dist.f[dir::d0P0][nodeIndex] - dist.f[dir::d0M0][nodeIndex]));
     vx2 *= oneOverRho;
 
-    vx3 = ((((dist.f[dir::DIR_PPP][nodeIndex] - dist.f[dir::DIR_MMM][nodeIndex]) +
-             (dist.f[dir::DIR_PMP][nodeIndex] - dist.f[dir::DIR_MPM][nodeIndex])) +
-            ((dist.f[dir::DIR_MPP][nodeIndex] - dist.f[dir::DIR_PMM][nodeIndex]) +
-             (dist.f[dir::DIR_MMP][nodeIndex] - dist.f[dir::DIR_PPM][nodeIndex]))) +
-           (((dist.f[dir::DIR_0MP][nodeIndex] - dist.f[dir::DIR_0PM][nodeIndex]) +
-             (dist.f[dir::DIR_0PP][nodeIndex] - dist.f[dir::DIR_0MM][nodeIndex])) +
-            ((dist.f[dir::DIR_M0P][nodeIndex] - dist.f[dir::DIR_P0M][nodeIndex]) +
-             (dist.f[dir::DIR_P0P][nodeIndex] - dist.f[dir::DIR_M0M][nodeIndex]))) +
-           (dist.f[dir::DIR_00P][nodeIndex] - dist.f[dir::DIR_00M][nodeIndex]));
+    vx3 = ((((dist.f[dir::dPPP][nodeIndex] - dist.f[dir::dMMM][nodeIndex]) +
+             (dist.f[dir::dPMP][nodeIndex] - dist.f[dir::dMPM][nodeIndex])) +
+            ((dist.f[dir::dMPP][nodeIndex] - dist.f[dir::dPMM][nodeIndex]) +
+             (dist.f[dir::dMMP][nodeIndex] - dist.f[dir::dPPM][nodeIndex]))) +
+           (((dist.f[dir::d0MP][nodeIndex] - dist.f[dir::d0PM][nodeIndex]) +
+             (dist.f[dir::d0PP][nodeIndex] - dist.f[dir::d0MM][nodeIndex])) +
+            ((dist.f[dir::dM0P][nodeIndex] - dist.f[dir::dP0M][nodeIndex]) +
+             (dist.f[dir::dP0P][nodeIndex] - dist.f[dir::dM0M][nodeIndex]))) +
+           (dist.f[dir::d00P][nodeIndex] - dist.f[dir::d00M][nodeIndex]));
     vx3 *= oneOverRho;
 }
 
-- 
GitLab