From e7b94822770dca689ca7e3d29f1f3f4f640ea950 Mon Sep 17 00:00:00 2001
From: Henry <henry.korb@geo.uu.se>
Date: Thu, 13 Oct 2022 17:05:15 +0200
Subject: [PATCH] fix merge bugs and remove BCstruct header

---
 .../Calculation/PlaneCalculations.cpp         |  1 -
 .../Calculation/UpdateGrid27.cpp              | 36 +++++-----
 .../Communication/ExchangeData27.cpp          | 20 +++---
 .../GridReaderFiles/GridReader.h              |  1 -
 .../GridReaderGenerator/GridGenerator.h       |  1 -
 .../Factories/BoundaryConditionFactory.h      |  1 -
 .../Factories/BoundaryConditionStructs.cuh    | 72 -------------------
 src/gpu/VirtualFluids_GPU/FindQ/FindQ.h       |  1 -
 .../GPU/CudaMemoryManager.cpp                 |  2 +-
 src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h |  1 -
 .../VirtualFluids_GPU/Init/PositionReader.cpp |  1 -
 .../KernelManager/BCKernelManager.h           |  1 -
 src/gpu/VirtualFluids_GPU/LBM/LB.h            | 68 ++++++++++++++++++
 .../VirtualFluids_GPU/Output/QDebugWriter.hpp |  1 -
 .../Output/UnstructuredGridWriter.hpp         |  1 -
 .../Parameter/CudaStreamManager.h             |  3 +-
 .../VirtualFluids_GPU/Parameter/Parameter.h   |  1 -
 17 files changed, 99 insertions(+), 113 deletions(-)
 delete mode 100644 src/gpu/VirtualFluids_GPU/Factories/BoundaryConditionStructs.cuh

diff --git a/src/gpu/VirtualFluids_GPU/Calculation/PlaneCalculations.cpp b/src/gpu/VirtualFluids_GPU/Calculation/PlaneCalculations.cpp
index ca1bba908..13b6bd662 100644
--- a/src/gpu/VirtualFluids_GPU/Calculation/PlaneCalculations.cpp
+++ b/src/gpu/VirtualFluids_GPU/Calculation/PlaneCalculations.cpp
@@ -10,7 +10,6 @@
 //using namespace std;
 //////////////////////////////////////////////////////////////////////////
 #include "Core/StringUtilities/StringUtil.h"
-#include "BoundaryConditions/BoundaryConditionStructs.cuh"
 
 void setSizeOfPlane(Parameter* para, int lev, unsigned int z)
 {
diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
index 32efd4133..f9fa76890 100644
--- a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
+++ b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
@@ -120,29 +120,29 @@ void UpdateGrid27::collisionAdvectionDiffusion(int level)
 
 void UpdateGrid27::prepareExchangeMultiGPU(int level, CudaStreamIndex streamIndex)
 {
-    prepareExchangeCollDataXGPU27AllNodes(para.get(), level);
-    prepareExchangeCollDataYGPU27AllNodes(para.get(), level);
-    prepareExchangeCollDataZGPU27AllNodes(para.get(), level);
+    prepareExchangeCollDataXGPU27AllNodes(para.get(), level, streamIndex);
+    prepareExchangeCollDataYGPU27AllNodes(para.get(), level, streamIndex);
+    prepareExchangeCollDataZGPU27AllNodes(para.get(), level, streamIndex);
 }
 
 void UpdateGrid27::prepareExchangeMultiGPUAfterFtoC(int level, CudaStreamIndex streamIndex)
 {
-    prepareExchangeCollDataXGPU27AfterFtoC(para.get(), level);
-    prepareExchangeCollDataYGPU27AfterFtoC(para.get(), level);
-    prepareExchangeCollDataZGPU27AfterFtoC(para.get(), level);
+    prepareExchangeCollDataXGPU27AfterFtoC(para.get(), level, streamIndex);
+    prepareExchangeCollDataYGPU27AfterFtoC(para.get(), level, streamIndex);
+    prepareExchangeCollDataZGPU27AfterFtoC(para.get(), level, streamIndex);
 }
 
 void UpdateGrid27::exchangeMultiGPU(int level, CudaStreamIndex streamIndex)
 {
     //////////////////////////////////////////////////////////////////////////
     // 3D domain decomposition
-    exchangeCollDataXGPU27AllNodes(para.get(), comm, cudaMemoryManager.get(), level);
-    exchangeCollDataYGPU27AllNodes(para.get(), comm, cudaMemoryManager.get(), level);
-    exchangeCollDataZGPU27AllNodes(para.get(), comm, cudaMemoryManager.get(), level);
+    exchangeCollDataXGPU27AllNodes(para.get(), comm, cudaMemoryManager.get(), level, streamIndex);
+    exchangeCollDataYGPU27AllNodes(para.get(), comm, cudaMemoryManager.get(), level, streamIndex);
+    exchangeCollDataZGPU27AllNodes(para.get(), comm, cudaMemoryManager.get(), level, streamIndex);
 
-    scatterNodesFromRecvBufferXGPU27AllNodes(para.get(), level);
-    scatterNodesFromRecvBufferYGPU27AllNodes(para.get(), level);
-    scatterNodesFromRecvBufferZGPU27AllNodes(para.get(), level);
+    scatterNodesFromRecvBufferXGPU27AllNodes(para.get(), level, streamIndex);
+    scatterNodesFromRecvBufferYGPU27AllNodes(para.get(), level, streamIndex);
+    scatterNodesFromRecvBufferZGPU27AllNodes(para.get(), level, streamIndex);
 
     //////////////////////////////////////////////////////////////////////////
     // 3D domain decomposition convection diffusion
@@ -208,13 +208,13 @@ void UpdateGrid27::exchangeMultiGPUAfterFtoC(int level, CudaStreamIndex streamIn
 {
     //////////////////////////////////////////////////////////////////////////
     // 3D domain decomposition
-    exchangeCollDataXGPU27AfterFtoC(para.get(), comm, cudaMemoryManager.get(), level);
-    exchangeCollDataYGPU27AfterFtoC(para.get(), comm, cudaMemoryManager.get(), level);
-    exchangeCollDataZGPU27AfterFtoC(para.get(), comm, cudaMemoryManager.get(), level);
+    exchangeCollDataXGPU27AfterFtoC(para.get(), comm, cudaMemoryManager.get(), level, streamIndex);
+    exchangeCollDataYGPU27AfterFtoC(para.get(), comm, cudaMemoryManager.get(), level, streamIndex);
+    exchangeCollDataZGPU27AfterFtoC(para.get(), comm, cudaMemoryManager.get(), level, streamIndex);
 
-    scatterNodesFromRecvBufferXGPU27AfterFtoC(para.get(), level);
-    scatterNodesFromRecvBufferYGPU27AfterFtoC(para.get(), level);
-    scatterNodesFromRecvBufferZGPU27AfterFtoC(para.get(), level);
+    scatterNodesFromRecvBufferXGPU27AfterFtoC(para.get(), level, streamIndex);
+    scatterNodesFromRecvBufferYGPU27AfterFtoC(para.get(), level, streamIndex);
+    scatterNodesFromRecvBufferZGPU27AfterFtoC(para.get(), level, streamIndex);
 
     //////////////////////////////////////////////////////////////////////////
     // 3D domain decomposition convection diffusion
diff --git a/src/gpu/VirtualFluids_GPU/Communication/ExchangeData27.cpp b/src/gpu/VirtualFluids_GPU/Communication/ExchangeData27.cpp
index 177f7699f..7b8950a93 100644
--- a/src/gpu/VirtualFluids_GPU/Communication/ExchangeData27.cpp
+++ b/src/gpu/VirtualFluids_GPU/Communication/ExchangeData27.cpp
@@ -107,13 +107,13 @@ void copyEdgeNodes(std::vector<LBMSimulationParameter::EdgeNodePositions> &edgeN
 ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
 void prepareExchangeCollDataXGPU27AllNodes(Parameter *para, int level, CudaStreamIndex streamIndex)
 {
-    collectNodesInSendBufferGPU(para, level, &para->getParD(level)->sendProcessNeighborX,
+    collectNodesInSendBufferGPU(para, level, streamIndex, &para->getParD(level)->sendProcessNeighborX,
                                 (unsigned int)(para->getNumberOfProcessNeighborsX(level, "send")));
 }
 
 void prepareExchangeCollDataXGPU27AfterFtoC(Parameter *para, int level, CudaStreamIndex streamIndex)
 {
-    collectNodesInSendBufferGPU(para, level, &para->getParD(level)->sendProcessNeighborsAfterFtoCX,
+    collectNodesInSendBufferGPU(para, level, streamIndex, &para->getParD(level)->sendProcessNeighborsAfterFtoCX,
                                 (unsigned int)(para->getNumberOfProcessNeighborsX(level, "send")));
 }
 
@@ -192,13 +192,13 @@ void exchangeCollDataXGPU27(Parameter *para, vf::gpu::Communicator &comm, CudaMe
 ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
 void prepareExchangeCollDataYGPU27AllNodes(Parameter *para, int level, CudaStreamIndex streamIndex)
 {
-    collectNodesInSendBufferGPU(para, level, &para->getParD(level)->sendProcessNeighborY,
+    collectNodesInSendBufferGPU(para, level, streamIndex, &para->getParD(level)->sendProcessNeighborY,
                                 (unsigned int)(para->getNumberOfProcessNeighborsY(level, "send")));
 }
 
 void prepareExchangeCollDataYGPU27AfterFtoC(Parameter *para, int level, CudaStreamIndex streamIndex)
 {
-    collectNodesInSendBufferGPU(para, level, &para->getParD(level)->sendProcessNeighborsAfterFtoCY,
+    collectNodesInSendBufferGPU(para, level, streamIndex, &para->getParD(level)->sendProcessNeighborsAfterFtoCY,
                                 (unsigned int)(para->getNumberOfProcessNeighborsY(level, "send")));
 }
 
@@ -224,13 +224,13 @@ void exchangeCollDataYGPU27AfterFtoC(Parameter *para, vf::gpu::Communicator &com
 
 void scatterNodesFromRecvBufferYGPU27AllNodes(Parameter *para, int level, CudaStreamIndex streamIndex)
 {
-    scatterNodesFromRecvBufferGPU(para, level, &para->getParD(level)->recvProcessNeighborY,
+    scatterNodesFromRecvBufferGPU(para, level, streamIndex, &para->getParD(level)->recvProcessNeighborY,
                                   (unsigned int)(para->getNumberOfProcessNeighborsY(level, "send")));
 }
 
 void scatterNodesFromRecvBufferYGPU27AfterFtoC(Parameter *para, int level, CudaStreamIndex streamIndex)
 {
-    scatterNodesFromRecvBufferGPU(para, level, &para->getParD(level)->recvProcessNeighborsAfterFtoCY,
+    scatterNodesFromRecvBufferGPU(para, level, streamIndex, &para->getParD(level)->recvProcessNeighborsAfterFtoCY,
                                   (unsigned int)(para->getNumberOfProcessNeighborsY(level, "send")));
 }
 
@@ -289,13 +289,13 @@ void exchangeCollDataYGPU27(Parameter *para, vf::gpu::Communicator &comm, CudaMe
 ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
 void prepareExchangeCollDataZGPU27AllNodes(Parameter *para, int level, CudaStreamIndex streamIndex)
 {
-    collectNodesInSendBufferGPU(para, level, &para->getParD(level)->sendProcessNeighborZ,
+    collectNodesInSendBufferGPU(para, level, streamIndex, &para->getParD(level)->sendProcessNeighborZ,
                                 (unsigned int)(para->getNumberOfProcessNeighborsZ(level, "send")));
 }
 
 void prepareExchangeCollDataZGPU27AfterFtoC(Parameter *para, int level, CudaStreamIndex streamIndex)
 {
-    collectNodesInSendBufferGPU(para, level, &para->getParD(level)->sendProcessNeighborsAfterFtoCZ,
+    collectNodesInSendBufferGPU(para, level, streamIndex, &para->getParD(level)->sendProcessNeighborsAfterFtoCZ,
                                 (unsigned int)(para->getNumberOfProcessNeighborsZ(level, "send")));
 }
 
@@ -320,13 +320,13 @@ void exchangeCollDataZGPU27AfterFtoC(Parameter *para, vf::gpu::Communicator &com
 
 void scatterNodesFromRecvBufferZGPU27AllNodes(Parameter *para, int level, CudaStreamIndex streamIndex)
 {
-    scatterNodesFromRecvBufferGPU(para, level, &para->getParD(level)->recvProcessNeighborZ,
+    scatterNodesFromRecvBufferGPU(para, level, streamIndex, &para->getParD(level)->recvProcessNeighborZ,
                                   (unsigned int)(para->getNumberOfProcessNeighborsZ(level, "send")));
 }
 
 void scatterNodesFromRecvBufferZGPU27AfterFtoC(Parameter *para, int level, CudaStreamIndex streamIndex)
 {
-    scatterNodesFromRecvBufferGPU(para, level, &para->getParD(level)->recvProcessNeighborsAfterFtoCZ,
+    scatterNodesFromRecvBufferGPU(para, level, streamIndex, &para->getParD(level)->recvProcessNeighborsAfterFtoCZ,
                                   (unsigned int)(para->getNumberOfProcessNeighborsZ(level, "send")));
 }
 
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.h b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.h
index ce08a8537..18efb6a78 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.h
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.h
@@ -8,7 +8,6 @@
 #include <memory>
 
 #include "LBM/LB.h"
-#include "BoundaryConditions/BoundaryConditionStructs.cuh"
 
 //#include "VirtualFluids_GPU_export.h"
 
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.h b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.h
index b9ca1c284..41b922560 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.h
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.h
@@ -40,7 +40,6 @@
 #include <memory>
 
 #include "LBM/LB.h"
-#include "BoundaryConditions/BoundaryConditionStructs.cuh"
 
 
 class Parameter;
diff --git a/src/gpu/VirtualFluids_GPU/Factories/BoundaryConditionFactory.h b/src/gpu/VirtualFluids_GPU/Factories/BoundaryConditionFactory.h
index e9bf8bcb7..f7cd7cc1a 100644
--- a/src/gpu/VirtualFluids_GPU/Factories/BoundaryConditionFactory.h
+++ b/src/gpu/VirtualFluids_GPU/Factories/BoundaryConditionFactory.h
@@ -39,7 +39,6 @@
 #include <variant>
 
 #include "LBM/LB.h"
-#include "BoundaryConditionStructs.cuh"
 #include "Parameter/Parameter.h"
 #include "gpu/GridGenerator/grid/BoundaryConditions/Side.h"
 
diff --git a/src/gpu/VirtualFluids_GPU/Factories/BoundaryConditionStructs.cuh b/src/gpu/VirtualFluids_GPU/Factories/BoundaryConditionStructs.cuh
deleted file mode 100644
index 88f65bd68..000000000
--- a/src/gpu/VirtualFluids_GPU/Factories/BoundaryConditionStructs.cuh
+++ /dev/null
@@ -1,72 +0,0 @@
-#ifndef BoundaryConditionStructs_H
-#define BoundaryConditionStructs_H
-#include "Core/DataTypes.h"
-
-//Q for second order BCs
-//! \struct to manage sub-grid-distances (q) for second order Boundary Conditions (BCs)
-typedef struct QforBC{
-   int* k;
-   int* kN;
-   long long* valueQ;
-   real* qread;
-   real* q27[27];
-   real* q19[19];
-   unsigned int numberOfBCnodes=0;
-   int kArray;
-   real *Vx,      *Vy,      *Vz;
-   real *Vx1,     *Vy1,     *Vz1;
-   real *deltaVz, *RhoBC;
-   real *normalX, *normalY, *normalZ;
-}QforBoundaryConditions;
-
-typedef struct QforPrecursorBC{
-   int* k;
-   int numberOfBCnodes=0;
-   int sizeQ;
-   int numberOfPrecursorNodes=0;
-   uint nPrecursorReads=0;
-   uint nTRead;
-   size_t numberOfQuantities;
-   real* q27[27];
-   uint* planeNeighborNT, *planeNeighborNB, *planeNeighborST, *planeNeighborSB;
-   real* weightsNT, *weightsNB, *weightsST,  *weightsSB;
-   real* last, *current, *next;
-   real velocityX, velocityY, velocityZ;
-}QforPrecursorBoundaryConditions;
-
-//BCTemp
-typedef struct TempforBC{
-   int* k;
-   real* temp;
-   int kTemp=0;
-}TempforBoundaryConditions;
-
-//BCTempVel
-typedef struct TempVelforBC{
-   int* k;
-   real* temp;
-   real* tempPulse;
-   real* velo;
-   int kTemp=0;
-}TempVelforBoundaryConditions;
-
-//BCTempPress
-typedef struct TempPressforBC{
-   int* k;
-   real* temp;
-   real* velo;
-   int kTemp=0;
-}TempPressforBoundaryConditions;
-
-// Settings for wall model used in StressBC
-typedef struct WMparas{
-   real* z0;
-   int* samplingOffset;
-   bool hasMonitor;
-   real* u_star;
-   real* Fx;
-   real* Fy;
-   real* Fz;
-}WallModelParameters;
-
-#endif
\ No newline at end of file
diff --git a/src/gpu/VirtualFluids_GPU/FindQ/FindQ.h b/src/gpu/VirtualFluids_GPU/FindQ/FindQ.h
index 39b77dc87..d1e8a08a7 100644
--- a/src/gpu/VirtualFluids_GPU/FindQ/FindQ.h
+++ b/src/gpu/VirtualFluids_GPU/FindQ/FindQ.h
@@ -4,7 +4,6 @@
 #include "LBM/LB.h"
 #include "lbm/constants/D3Q27.h"
 #include "Parameter/Parameter.h"
-#include "BoundaryConditions/BoundaryConditionStructs.cuh"
 
 
 void findQ(Parameter* para, int lev);
diff --git a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
index ec79d1c55..b0a2ea595 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
+++ b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
@@ -1736,7 +1736,7 @@ void CudaMemoryManager::cudaCopyPrecursorBC(int lev)
 void CudaMemoryManager::cudaCopyPrecursorData(int lev)
 {
     auto prec = &parameter->getParH(lev)->precursorBC;
-    auto precStream = parameter->getStreamManager()->getStream(CudaStreamManager::StreamIndex::precursorStream);
+    auto precStream = parameter->getStreamManager()->getStream(CudaStreamIndex::Precursor);
     size_t memSize = prec->numberOfPrecursorNodes*sizeof(real)*prec->numberOfQuantities;
     checkCudaErrors( cudaStreamSynchronize(precStream) );
     checkCudaErrors( cudaMemcpyAsync(parameter->getParD(lev)->precursorBC.next, prec->next, memSize, cudaMemcpyHostToDevice, precStream) );
diff --git a/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h b/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h
index 283218cd1..ceb70fb12 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h
+++ b/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h
@@ -9,7 +9,6 @@
 #define GPU_INTERFACE_H
 
 #include "LBM/LB.h"
-#include "BoundaryConditions/BoundaryConditionStructs.cuh"
 
 #include <cuda_runtime.h>
 #include <curand.h>
diff --git a/src/gpu/VirtualFluids_GPU/Init/PositionReader.cpp b/src/gpu/VirtualFluids_GPU/Init/PositionReader.cpp
index 2ea931014..04052b38c 100644
--- a/src/gpu/VirtualFluids_GPU/Init/PositionReader.cpp
+++ b/src/gpu/VirtualFluids_GPU/Init/PositionReader.cpp
@@ -3,7 +3,6 @@
 #include "Parameter/Parameter.h"
 
 #include <basics/utilities/UbFileInputASCII.h>
-#include "BoundaryConditions/BoundaryConditionStructs.cuh"
 
 using namespace vf::lbm::dir;
 
diff --git a/src/gpu/VirtualFluids_GPU/KernelManager/BCKernelManager.h b/src/gpu/VirtualFluids_GPU/KernelManager/BCKernelManager.h
index 386cf95fc..339100e6b 100644
--- a/src/gpu/VirtualFluids_GPU/KernelManager/BCKernelManager.h
+++ b/src/gpu/VirtualFluids_GPU/KernelManager/BCKernelManager.h
@@ -40,7 +40,6 @@
 #include "LBM/LB.h"
 #include "PointerDefinitions.h"
 #include "VirtualFluids_GPU_export.h"
-#include "BoundaryConditions/BoundaryConditionStructs.cuh"
 
 
 class CudaMemoryManager;
diff --git a/src/gpu/VirtualFluids_GPU/LBM/LB.h b/src/gpu/VirtualFluids_GPU/LBM/LB.h
index 734e15c33..cedf1cda9 100644
--- a/src/gpu/VirtualFluids_GPU/LBM/LB.h
+++ b/src/gpu/VirtualFluids_GPU/LBM/LB.h
@@ -198,6 +198,74 @@ typedef struct SubgridDist27{
    real* q[27];
 } SubgridDistances27;
 
+//Q for second order BCs
+//! \struct to manage sub-grid-distances (q) for second order Boundary Conditions (BCs)
+typedef struct QforBC{
+   int* k;
+   int* kN;
+   long long* valueQ;
+   real* qread;
+   real* q27[27];
+   real* q19[19];
+   unsigned int numberOfBCnodes=0;
+   int kArray;
+   real *Vx,      *Vy,      *Vz;
+   real *Vx1,     *Vy1,     *Vz1;
+   real *deltaVz, *RhoBC;
+   real *normalX, *normalY, *normalZ;
+}QforBoundaryConditions;
+
+typedef struct QforPrecursorBC{
+   int* k;
+   int numberOfBCnodes=0;
+   int sizeQ;
+   int numberOfPrecursorNodes=0;
+   uint nPrecursorReads=0;
+   uint nTRead;
+   size_t numberOfQuantities;
+   real* q27[27];
+   uint* planeNeighborNT, *planeNeighborNB, *planeNeighborST, *planeNeighborSB;
+   real* weightsNT, *weightsNB, *weightsST,  *weightsSB;
+   real* last, *current, *next;
+   real velocityX, velocityY, velocityZ;
+}QforPrecursorBoundaryConditions;
+
+//BCTemp
+typedef struct TempforBC{
+   int* k;
+   real* temp;
+   int kTemp=0;
+}TempforBoundaryConditions;
+
+//BCTempVel
+typedef struct TempVelforBC{
+   int* k;
+   real* temp;
+   real* tempPulse;
+   real* velo;
+   int kTemp=0;
+}TempVelforBoundaryConditions;
+
+//BCTempPress
+typedef struct TempPressforBC{
+   int* k;
+   real* temp;
+   real* velo;
+   int kTemp=0;
+}TempPressforBoundaryConditions;
+
+// Settings for wall model used in StressBC
+typedef struct WMparas{
+   real* z0;
+   int* samplingOffset;
+   bool hasMonitor;
+   real* u_star;
+   real* Fx;
+   real* Fy;
+   real* Fz;
+}WallModelParameters;
+
+
 //measurePoints
 typedef struct MeasP{
 	std::string name;
diff --git a/src/gpu/VirtualFluids_GPU/Output/QDebugWriter.hpp b/src/gpu/VirtualFluids_GPU/Output/QDebugWriter.hpp
index 946810d5c..d00663657 100644
--- a/src/gpu/VirtualFluids_GPU/Output/QDebugWriter.hpp
+++ b/src/gpu/VirtualFluids_GPU/Output/QDebugWriter.hpp
@@ -12,7 +12,6 @@
 #include "basics/utilities/UbSystem.h"
 #include <basics/writer/WbWriterVtkXmlBinary.h>
 #include "Core/StringUtilities/StringUtil.h"
-#include "BoundaryConditions/BoundaryConditionStructs.cuh"
 
 //using namespace std;
 
diff --git a/src/gpu/VirtualFluids_GPU/Output/UnstructuredGridWriter.hpp b/src/gpu/VirtualFluids_GPU/Output/UnstructuredGridWriter.hpp
index 2beb1d25c..81f2c028a 100644
--- a/src/gpu/VirtualFluids_GPU/Output/UnstructuredGridWriter.hpp
+++ b/src/gpu/VirtualFluids_GPU/Output/UnstructuredGridWriter.hpp
@@ -9,7 +9,6 @@
 #include "LBM/LB.h"
 #include "lbm/constants/D3Q27.h"
 #include "Parameter/Parameter.h"
-#include "BoundaryConditions/BoundaryConditionStructs.cuh"
 #include "basics/utilities/UbSystem.h"
 #include <basics/writer/WbWriterVtkXmlBinary.h>
 #include <basics/writer/WbWriterVtkXmlASCII.h>
diff --git a/src/gpu/VirtualFluids_GPU/Parameter/CudaStreamManager.h b/src/gpu/VirtualFluids_GPU/Parameter/CudaStreamManager.h
index 33d09a6b9..f567df97b 100644
--- a/src/gpu/VirtualFluids_GPU/Parameter/CudaStreamManager.h
+++ b/src/gpu/VirtualFluids_GPU/Parameter/CudaStreamManager.h
@@ -37,7 +37,8 @@ enum class CudaStreamIndex
     {
         Legacy,
         Bulk,
-        Border
+        Border,
+        Precursor
     };
 class CudaStreamManager
 {
diff --git a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
index 1384f101f..aff100584 100644
--- a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
+++ b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
@@ -43,7 +43,6 @@
 #include "LBM/LB.h"
 #include "PreCollisionInteractor/PreCollisionInteractor.h"
 #include "TurbulenceModels/TurbulenceModelFactory.h"
-#include "BoundaryConditions/BoundaryConditionStructs.cuh"
 
 #include "VirtualFluids_GPU_export.h"
 
-- 
GitLab