From 087f7557f2e9d07ba21adfdffb2e991b27f7e4eb Mon Sep 17 00:00:00 2001
From: peters <peters@irmb.tu-bs.de>
Date: Wed, 29 Sep 2021 14:33:08 +0200
Subject: [PATCH] Update GKS code to newest version of VF. Also: Build GKS in
 CI

---
 .gitlab-ci.yml                                 |  3 ++-
 CMake/3rd/vtk.cmake                            |  8 ++++----
 CMakeLists.txt                                 |  5 ++---
 apps/gpu/GKS/Flame7cm/CMakeLists.txt           |  7 +++----
 apps/gpu/GKS/Flame7cm/Flame7cm.cpp             | 15 +++++++--------
 src/gpu/GksGpu/Analyzer/EnstrophyAnalyzer.cu   |  2 ++
 src/gpu/GksGpu/Analyzer/HeatFluxAnalyzer.cu    |  2 ++
 src/gpu/GksGpu/Analyzer/TurbulenceAnalyzer.cu  |  2 ++
 .../BoundaryConditions/CreepingMassFlux.cu     |  2 +-
 src/gpu/GksGpu/BoundaryConditions/HeatFlux.cu  |  2 +-
 src/gpu/GksGpu/BoundaryConditions/Inflow.cu    |  2 +-
 .../BoundaryConditions/MassCompensation.cu     |  2 +-
 src/gpu/GksGpu/CMakeLists.txt                  |  5 ++++-
 .../GksGpu/CellProperties/CellProperties.cuh   |  4 ++++
 src/gpu/GksGpu/CellUpdate/Reaction.cuh         | 12 ++++++------
 src/gpu/GksGpu/DataBase/DataBase.cpp           |  4 ++++
 .../GksGpu/FlowStateData/AccessDeviceData.cuh  |  4 ++++
 src/gpu/GksGpu/FlowStateData/FlowStateData.cuh |  9 ++++++++-
 .../FlowStateData/FlowStateDataConversion.cuh  |  4 ++++
 .../GksGpu/FlowStateData/HeatCapacities.cuh    | 16 ++++++++++------
 .../FlowStateData/ThermalDependencies.cuh      |  4 ++++
 .../GksGpu/FluxComputation/FluxComputation.cu  |  4 ++--
 src/gpu/GksGpu/Output/VtkWriter.cpp            |  4 ++++
 src/gpu/GksGpu/Output/VtkWriter.h              |  5 +++++
 src/gpu/GksMeshAdapter/CMakeLists.txt          |  2 +-
 src/gpu/GksMeshAdapter/GksMeshAdapter.cpp      | 18 +++++++++++-------
 src/lbm/CMakeLists.txt                         |  2 +-
 27 files changed, 100 insertions(+), 49 deletions(-)

diff --git a/.gitlab-ci.yml b/.gitlab-ci.yml
index d2423ba9d..d9cc93734 100644
--- a/.gitlab-ci.yml
+++ b/.gitlab-ci.yml
@@ -54,6 +54,7 @@ stages:
       --preset=all_make
       -DBUILD_WARNINGS_AS_ERRORS=ON
       -DCMAKE_CUDA_ARCHITECTURES=60
+      -DBUILD_VF_GKS=ON
     - make -j4
     - ccache --show-stats
 
@@ -101,7 +102,7 @@ msvc_16:
     - cd $CI_PROJECT_DIR
     - md -force $env:BUILD_FOLDER
     - cd $env:BUILD_FOLDER
-    - cmake .. --preset=all_msvc -DCMAKE_CUDA_ARCHITECTURES=61 -DBUILD_WARNINGS_AS_ERRORS=ON
+    - cmake .. --preset=all_msvc -DCMAKE_CUDA_ARCHITECTURES=61 -DBUILD_WARNINGS_AS_ERRORS=ON -DBUILD_VF_GKS=ON
     - MSBuild.exe VirtualFluids.sln /property:Configuration=$env:BUILD_CONFIGURATION /verbosity:minimal /maxcpucount:4
 
   cache:
diff --git a/CMake/3rd/vtk.cmake b/CMake/3rd/vtk.cmake
index 83cf22e88..6f5106428 100644
--- a/CMake/3rd/vtk.cmake
+++ b/CMake/3rd/vtk.cmake
@@ -3,9 +3,9 @@
 # VTK_DIR needs to bet set to the VTK build directory in the config file.
 #########################################################################
 find_package(VTK REQUIRED)
-    vf_get_library_name(library_name)
+vf_get_library_name(library_name)
 
-    include(${VTK_USE_FILE})
-    target_include_directories(${library_name} PRIVATE ${VTK_INCLUDE_DIRS})
+include(${VTK_USE_FILE})
+target_include_directories(${library_name} PRIVATE ${VTK_INCLUDE_DIRS})
 
-    target_link_libraries(${library_name} PRIVATE ${VTK_LIBRARIES})
+target_link_libraries(${library_name} PRIVATE ${VTK_LIBRARIES})
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 7796c73f8..9d619f78d 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -49,7 +49,6 @@ option(BUILD_USE_OPENMP "Build VirtualFluids with openmp" ON)
 option(BUILD_USE_BOOST "Build VirtualFluids with boost" OFF)
 
 # vf gpu
-option(BUILD_VF_GPU          "Build VirtualFluids GPU"     ON )
 option(BUILD_VF_GKS          "Build VirtualFluids GKS"     OFF )
 option(BUILD_VF_TRAFFIC      "Build VirtualFluids Traffic" OFF)
 option(BUILD_JSONCPP         "Builds json cpp "            OFF)
@@ -97,7 +96,7 @@ IF( BUILD_VF_DOUBLE_ACCURACY )
     list(APPEND VF_COMPILER_DEFINITION VF_DOUBLE_ACCURACY)
 ENDIF()
 
-if(BUILD_VF_GPU)
+if(BUILD_VF_GPU OR BUILD_VF_GKS)
     include(CheckLanguage)
     check_language(CUDA)
 
@@ -192,6 +191,6 @@ add_subdirectory(src/lbm)
 if (BUILD_VF_CPU)
     include (cpu.cmake)
 endif()
-if(BUILD_VF_GPU)
+if(BUILD_VF_GPU OR BUILD_VF_GKS)
     include (gpu.cmake)
 endif()
diff --git a/apps/gpu/GKS/Flame7cm/CMakeLists.txt b/apps/gpu/GKS/Flame7cm/CMakeLists.txt
index c3f57dcdd..75ca5fa4b 100644
--- a/apps/gpu/GKS/Flame7cm/CMakeLists.txt
+++ b/apps/gpu/GKS/Flame7cm/CMakeLists.txt
@@ -1,6 +1,5 @@
-PROJECT(Flame7cm)
+PROJECT(Flame7cm LANGUAGES CUDA CXX)
 
-vf_add_library(BUILDTYPE binary PRIVATE_LINK basics GridGenerator GksMeshAdapter GksVtkAdapter GksGpu FILES Flame7cm.cpp )
+vf_add_library(BUILDTYPE binary PRIVATE_LINK basics GridGenerator GksMeshAdapter GksVtkAdapter GksGpu MPI::MPI_CXX FILES Flame7cm.cpp )
 
-include (${VF_CMAKE_DIR}/3rd/cuda.cmake)
-include (${VF_CMAKE_DIR}/3rd/mpi.cmake)
+set_source_files_properties(Flame7cm.cpp PROPERTIES LANGUAGE CUDA)
diff --git a/apps/gpu/GKS/Flame7cm/Flame7cm.cpp b/apps/gpu/GKS/Flame7cm/Flame7cm.cpp
index e0b736dfe..4323ce5ae 100644
--- a/apps/gpu/GKS/Flame7cm/Flame7cm.cpp
+++ b/apps/gpu/GKS/Flame7cm/Flame7cm.cpp
@@ -152,11 +152,10 @@ void thermalCavity( std::string path, std::string simulationName, uint _gpuIndex
 
     ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
 
-    auto gridFactory = GridFactory::make();
-    gridFactory->setGridStrategy(Device::CPU);
-    gridFactory->setTriangularMeshDiscretizationMethod(TriangularMeshDiscretizationMethod::POINT_IN_OBJECT);
+    // auto gridFactory = GridFactory::make();
+    // gridFactory->setTriangularMeshDiscretizationMethod(TriangularMeshDiscretizationMethod::POINT_IN_OBJECT);
 
-    auto gridBuilder = MultipleGridBuilder::makeShared(gridFactory);
+    auto gridBuilder = MultipleGridBuilder::makeShared();
 
     ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
 
@@ -476,14 +475,14 @@ int main( int argc, char* argv[])
     {
         thermalCavity( path, simulationName, gpuIndex, nx, useTempLimiter, restartIter );
     }
-    catch (const std::exception& e)
-    {     
-        *logging::out << logging::Logger::LOGGER_ERROR << e.what() << "\n";
-    }
     catch (const std::bad_alloc& e)
     {  
         *logging::out << logging::Logger::LOGGER_ERROR << "Bad Alloc:" << e.what() << "\n";
     }
+    catch (const std::exception& e)
+    {     
+        *logging::out << logging::Logger::LOGGER_ERROR << e.what() << "\n";
+    }
     catch (...)
     {
         *logging::out << logging::Logger::LOGGER_ERROR << "Unknown exception!\n";
diff --git a/src/gpu/GksGpu/Analyzer/EnstrophyAnalyzer.cu b/src/gpu/GksGpu/Analyzer/EnstrophyAnalyzer.cu
index 5a3dc76db..346692bfd 100644
--- a/src/gpu/GksGpu/Analyzer/EnstrophyAnalyzer.cu
+++ b/src/gpu/GksGpu/Analyzer/EnstrophyAnalyzer.cu
@@ -59,6 +59,8 @@ bool EnstrophyAnalyzer::run(uint iter)
     this->enstrophyTimeSeries.push_back( EnstrophyTmp );
 
     //*logging::out << logging::Logger::INFO_HIGH << "EKin = " << EKin << "\n";
+
+    return true;
 }
 
 ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/GksGpu/Analyzer/HeatFluxAnalyzer.cu b/src/gpu/GksGpu/Analyzer/HeatFluxAnalyzer.cu
index 266d3b564..ed68f8d95 100644
--- a/src/gpu/GksGpu/Analyzer/HeatFluxAnalyzer.cu
+++ b/src/gpu/GksGpu/Analyzer/HeatFluxAnalyzer.cu
@@ -63,6 +63,8 @@ bool HeatFluxAnalyzer::run(uint iter, Parameters parameters)
     this->heatFluxTimeSeries.push_back( q / qIdeal );
 
     if( iter % this->outputIter == 0 ) *logging::out << logging::Logger::INFO_HIGH << "q = " << q / qIdeal << "\n";
+
+    return true;
 }
 
 //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/GksGpu/Analyzer/TurbulenceAnalyzer.cu b/src/gpu/GksGpu/Analyzer/TurbulenceAnalyzer.cu
index d05c47cd2..5e896e03e 100644
--- a/src/gpu/GksGpu/Analyzer/TurbulenceAnalyzer.cu
+++ b/src/gpu/GksGpu/Analyzer/TurbulenceAnalyzer.cu
@@ -51,6 +51,8 @@ bool TurbulenceAnalyzer::run(uint iter, Parameters parameters)
     getLastCudaError("TurbulenceAnalyzer::run(uint iter, Parameters parameters)");
 
     this->counter++;
+
+    return true;
 }
 
 ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/GksGpu/BoundaryConditions/CreepingMassFlux.cu b/src/gpu/GksGpu/BoundaryConditions/CreepingMassFlux.cu
index 2b8b8174f..fd5591824 100644
--- a/src/gpu/GksGpu/BoundaryConditions/CreepingMassFlux.cu
+++ b/src/gpu/GksGpu/BoundaryConditions/CreepingMassFlux.cu
@@ -91,7 +91,7 @@ __host__ __device__ inline void boundaryConditionFunction(const DataBaseStruct&
                                                           const uint startIndex,
                                                           const uint index)
 {
-    uint ghostCellIdx  = boundaryCondition.ghostCells [ startIndex + index ];
+    // uint ghostCellIdx  = boundaryCondition.ghostCells [ startIndex + index ];
     uint domainCellIdx = boundaryCondition.domainCells[ startIndex + index ];
 
     ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/GksGpu/BoundaryConditions/HeatFlux.cu b/src/gpu/GksGpu/BoundaryConditions/HeatFlux.cu
index 3ecd1b6cd..87f880bcf 100644
--- a/src/gpu/GksGpu/BoundaryConditions/HeatFlux.cu
+++ b/src/gpu/GksGpu/BoundaryConditions/HeatFlux.cu
@@ -91,7 +91,7 @@ __host__ __device__ inline void boundaryConditionFunction(const DataBaseStruct&
                                                           const uint startIndex,
                                                           const uint index)
 {
-    uint ghostCellIdx  = boundaryCondition.ghostCells [ startIndex + index ];
+    // uint ghostCellIdx  = boundaryCondition.ghostCells [ startIndex + index ];
     uint domainCellIdx = boundaryCondition.domainCells[ startIndex + index ];
 
     ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/GksGpu/BoundaryConditions/Inflow.cu b/src/gpu/GksGpu/BoundaryConditions/Inflow.cu
index 21ab98293..7f9b2777f 100644
--- a/src/gpu/GksGpu/BoundaryConditions/Inflow.cu
+++ b/src/gpu/GksGpu/BoundaryConditions/Inflow.cu
@@ -86,7 +86,7 @@ __host__ __device__ inline void boundaryConditionFunction(const DataBaseStruct&
 {
     uint ghostCellIdx  = boundaryCondition.ghostCells [ startIndex + index ];
     uint domainCellIdx = boundaryCondition.domainCells[ startIndex + index ];
-    uint secondCellIdx = boundaryCondition.secondCells[ startIndex + index ];
+    // uint secondCellIdx = boundaryCondition.secondCells[ startIndex + index ];
 
     PrimitiveVariables ghostCellPrim;
     {
diff --git a/src/gpu/GksGpu/BoundaryConditions/MassCompensation.cu b/src/gpu/GksGpu/BoundaryConditions/MassCompensation.cu
index 4aaf40634..f6e697426 100644
--- a/src/gpu/GksGpu/BoundaryConditions/MassCompensation.cu
+++ b/src/gpu/GksGpu/BoundaryConditions/MassCompensation.cu
@@ -91,7 +91,7 @@ __host__ __device__ inline void boundaryConditionFunction(const DataBaseStruct&
                                                           const uint startIndex,
                                                           const uint index)
 {
-    uint ghostCellIdx  = boundaryCondition.ghostCells [ startIndex + index ];
+    // uint ghostCellIdx  = boundaryCondition.ghostCells [ startIndex + index ];
     uint domainCellIdx = boundaryCondition.domainCells[ startIndex + index ];
 
     ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/GksGpu/CMakeLists.txt b/src/gpu/GksGpu/CMakeLists.txt
index da404e020..6da2a9ec8 100644
--- a/src/gpu/GksGpu/CMakeLists.txt
+++ b/src/gpu/GksGpu/CMakeLists.txt
@@ -1,3 +1,6 @@
 project(GksGpu LANGUAGES CUDA CXX)
 
-vf_add_library(PRIVATE_LINK basics GksMeshAdapter OpenMP::OpenMP_CXX MPI::MPI_CXX)
+vf_add_library(PRIVATE_LINK basics lbmCuda GksMeshAdapter OpenMP::OpenMP_CXX MPI::MPI_CXX)
+
+target_include_directories(GksGpu PRIVATE "${VF_THIRD_DIR}/cuda_samples/")
+target_compile_options(GksGpu PRIVATE "-fPIC")
\ No newline at end of file
diff --git a/src/gpu/GksGpu/CellProperties/CellProperties.cuh b/src/gpu/GksGpu/CellProperties/CellProperties.cuh
index 1ce36e85b..08731b9f5 100644
--- a/src/gpu/GksGpu/CellProperties/CellProperties.cuh
+++ b/src/gpu/GksGpu/CellProperties/CellProperties.cuh
@@ -4,9 +4,13 @@
 #ifdef __CUDACC__
 #include <cuda_runtime.h>
 #else
+#ifndef __host__
 #define __host__
+#endif
+#ifndef __device__
 #define __device__
 #endif
+#endif
 
 //////////////////////////////////////////////////////////////////////////
 
diff --git a/src/gpu/GksGpu/CellUpdate/Reaction.cuh b/src/gpu/GksGpu/CellUpdate/Reaction.cuh
index 4bf317b27..21ba61220 100644
--- a/src/gpu/GksGpu/CellUpdate/Reaction.cuh
+++ b/src/gpu/GksGpu/CellUpdate/Reaction.cuh
@@ -36,14 +36,14 @@ inline __host__ __device__ real getTurbulentViscosityDeardorff(const DataBaseStr
         real uHead = c1o2 * prim.U;
 
         {
-            uint neighborCellIndex = dataBase.cellToCell[CELL_TO_CELL(cellIndex, 0, dataBase.numberOfCells)];
+            // uint neighborCellIndex = dataBase.cellToCell[CELL_TO_CELL(cellIndex, 0, dataBase.numberOfCells)];
             readCellData(cellIndex, dataBase, neighborCons);
             neighborPrim = toPrimitiveVariables(neighborCons, parameters.K);
 
             uHead += c1o4 * neighborPrim.U;
         }
         {
-            uint neighborCellIndex = dataBase.cellToCell[CELL_TO_CELL(cellIndex, 1, dataBase.numberOfCells)];
+            // uint neighborCellIndex = dataBase.cellToCell[CELL_TO_CELL(cellIndex, 1, dataBase.numberOfCells)];
             readCellData(cellIndex, dataBase, neighborCons);
             neighborPrim = toPrimitiveVariables(neighborCons, parameters.K);
 
@@ -57,14 +57,14 @@ inline __host__ __device__ real getTurbulentViscosityDeardorff(const DataBaseStr
         real vHead = c1o2 * prim.V;
 
         {
-            uint neighborCellIndex = dataBase.cellToCell[CELL_TO_CELL(cellIndex, 2, dataBase.numberOfCells)];
+            // uint neighborCellIndex = dataBase.cellToCell[CELL_TO_CELL(cellIndex, 2, dataBase.numberOfCells)];
             readCellData(cellIndex, dataBase, neighborCons);
             neighborPrim = toPrimitiveVariables(neighborCons, parameters.K);
 
             vHead += c1o4 * neighborPrim.V;
         }
         {
-            uint neighborCellIndex = dataBase.cellToCell[CELL_TO_CELL(cellIndex, 3, dataBase.numberOfCells)];
+            // uint neighborCellIndex = dataBase.cellToCell[CELL_TO_CELL(cellIndex, 3, dataBase.numberOfCells)];
             readCellData(cellIndex, dataBase, neighborCons);
             neighborPrim = toPrimitiveVariables(neighborCons, parameters.K);
 
@@ -78,14 +78,14 @@ inline __host__ __device__ real getTurbulentViscosityDeardorff(const DataBaseStr
         real wHead = c1o2 * prim.W;
 
         {
-            uint neighborCellIndex = dataBase.cellToCell[CELL_TO_CELL(cellIndex, 4, dataBase.numberOfCells)];
+            // uint neighborCellIndex = dataBase.cellToCell[CELL_TO_CELL(cellIndex, 4, dataBase.numberOfCells)];
             readCellData(cellIndex, dataBase, neighborCons);
             neighborPrim = toPrimitiveVariables(neighborCons, parameters.K);
 
             wHead += c1o4 * neighborPrim.W;
         }
         {
-            uint neighborCellIndex = dataBase.cellToCell[CELL_TO_CELL(cellIndex, 5, dataBase.numberOfCells)];
+            // uint neighborCellIndex = dataBase.cellToCell[CELL_TO_CELL(cellIndex, 5, dataBase.numberOfCells)];
             readCellData(cellIndex, dataBase, neighborCons);
             neighborPrim = toPrimitiveVariables(neighborCons, parameters.K);
 
diff --git a/src/gpu/GksGpu/DataBase/DataBase.cpp b/src/gpu/GksGpu/DataBase/DataBase.cpp
index 21c51b757..46921a683 100644
--- a/src/gpu/GksGpu/DataBase/DataBase.cpp
+++ b/src/gpu/GksGpu/DataBase/DataBase.cpp
@@ -14,6 +14,10 @@
 #include "GksMeshAdapter/GksMeshAdapter.h"
 #include "Communication/Communicator.h"
 
+#include <lbm/constants/NumericConstants.h>
+
+using namespace vf::lbm::constant;
+
 namespace GksGpu {
 
 DataBase::DataBase( std::string type ) 
diff --git a/src/gpu/GksGpu/FlowStateData/AccessDeviceData.cuh b/src/gpu/GksGpu/FlowStateData/AccessDeviceData.cuh
index 3ff9848f1..2ad158173 100644
--- a/src/gpu/GksGpu/FlowStateData/AccessDeviceData.cuh
+++ b/src/gpu/GksGpu/FlowStateData/AccessDeviceData.cuh
@@ -4,9 +4,13 @@
 #ifdef __CUDACC__
 #include <cuda_runtime.h>
 #else
+#ifndef __host__
 #define __host__
+#endif
+#ifndef __device__
 #define __device__
 #endif
+#endif
 
 #include "Core/DataTypes.h"
 #include "Core/RealConstants.h"
diff --git a/src/gpu/GksGpu/FlowStateData/FlowStateData.cuh b/src/gpu/GksGpu/FlowStateData/FlowStateData.cuh
index 3ec3a7c9e..3b7929b39 100644
--- a/src/gpu/GksGpu/FlowStateData/FlowStateData.cuh
+++ b/src/gpu/GksGpu/FlowStateData/FlowStateData.cuh
@@ -4,15 +4,22 @@
 #ifdef __CUDACC__
 #include <cuda_runtime.h>
 #else
+#ifndef __host__
 #define __host__
+#endif
+#ifndef __device__
 #define __device__
 #endif
+#endif
 
 #include "Core/DataTypes.h"
-#include "Core/RealConstants.h"
 
 #include "Definitions/PassiveScalar.h"
 
+#include <lbm/constants/NumericConstants.h>
+
+using namespace vf::lbm::constant;
+
 namespace GksGpu {
 
 //////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/GksGpu/FlowStateData/FlowStateDataConversion.cuh b/src/gpu/GksGpu/FlowStateData/FlowStateDataConversion.cuh
index c33f02ea3..b7b759c99 100644
--- a/src/gpu/GksGpu/FlowStateData/FlowStateDataConversion.cuh
+++ b/src/gpu/GksGpu/FlowStateData/FlowStateDataConversion.cuh
@@ -4,9 +4,13 @@
 #ifdef __CUDACC__
 #include <cuda_runtime.h>
 #else
+#ifndef __host__
 #define __host__
+#endif
+#ifndef __device__
 #define __device__
 #endif
+#endif
 
 #include "Core/DataTypes.h"
 #include "Core/RealConstants.h"
diff --git a/src/gpu/GksGpu/FlowStateData/HeatCapacities.cuh b/src/gpu/GksGpu/FlowStateData/HeatCapacities.cuh
index 3002ed3dd..04a164aa3 100644
--- a/src/gpu/GksGpu/FlowStateData/HeatCapacities.cuh
+++ b/src/gpu/GksGpu/FlowStateData/HeatCapacities.cuh
@@ -1,12 +1,16 @@
 //#ifndef HeatCapacities_H
 //#define HeatCapacities_H
 //
-//#ifdef __CUDACC__
-//#include <cuda_runtime.h>
-//#else
-//#define __host__
-//#define __device__
-//#endif
+// #ifdef __CUDACC__
+// #include <cuda_runtime.h>
+// #else
+// #ifndef __host__
+// #define __host__
+// #endif
+// #ifndef __device__
+// #define __device__
+// #endif
+// #endif
 //
 //#include "Core/DataTypes.h"
 //#include "Core/RealConstants.h"
diff --git a/src/gpu/GksGpu/FlowStateData/ThermalDependencies.cuh b/src/gpu/GksGpu/FlowStateData/ThermalDependencies.cuh
index 9f3a268a5..47eb261a0 100644
--- a/src/gpu/GksGpu/FlowStateData/ThermalDependencies.cuh
+++ b/src/gpu/GksGpu/FlowStateData/ThermalDependencies.cuh
@@ -4,9 +4,13 @@
 #ifdef __CUDACC__
 #include <cuda_runtime.h>
 #else
+#ifndef __host__
 #define __host__
+#endif
+#ifndef __device__
 #define __device__
 #endif
+#endif
 
 #include <math.h>
 
diff --git a/src/gpu/GksGpu/FluxComputation/FluxComputation.cu b/src/gpu/GksGpu/FluxComputation/FluxComputation.cu
index 8c9358636..25ba5726b 100644
--- a/src/gpu/GksGpu/FluxComputation/FluxComputation.cu
+++ b/src/gpu/GksGpu/FluxComputation/FluxComputation.cu
@@ -152,7 +152,7 @@ __host__ __device__ inline void fluxFunction(DataBaseStruct dataBase, Parameters
     {
         if( parameters.spongeLayerIdx == 0 )
         {
-            real x = dataBase.faceCenter[VEC_X(faceIndex, dataBase.numberOfFaces)];
+            // real x = dataBase.faceCenter[VEC_X(faceIndex, dataBase.numberOfFaces)];
             real z = dataBase.faceCenter[VEC_Z(faceIndex, dataBase.numberOfFaces)];
 
             real muNew = parameters.mu;
@@ -168,7 +168,7 @@ __host__ __device__ inline void fluxFunction(DataBaseStruct dataBase, Parameters
         }
         if( parameters.spongeLayerIdx == 1 )
         {
-            real x = dataBase.faceCenter[VEC_X(faceIndex, dataBase.numberOfFaces)];
+            // real x = dataBase.faceCenter[VEC_X(faceIndex, dataBase.numberOfFaces)];
             real z = dataBase.faceCenter[VEC_Z(faceIndex, dataBase.numberOfFaces)];
 
             real muNew = parameters.mu;
diff --git a/src/gpu/GksGpu/Output/VtkWriter.cpp b/src/gpu/GksGpu/Output/VtkWriter.cpp
index a1a0ab9f6..234151c7d 100644
--- a/src/gpu/GksGpu/Output/VtkWriter.cpp
+++ b/src/gpu/GksGpu/Output/VtkWriter.cpp
@@ -47,6 +47,8 @@
 #include "FlowStateData/FlowStateDataConversion.cuh"
 #include "FlowStateData/AccessDeviceData.cuh"
 
+namespace GksGpu {
+
 void VtkWriter::write(std::shared_ptr<DataBase> dataBase, Parameters parameters, std::string filename)
 {
     *logging::out << logging::Logger::INFO_INTERMEDIATE << "Write " << filename << ".vtu" << " ... \n";
@@ -144,3 +146,5 @@ void VtkWriter::write(std::shared_ptr<DataBase> dataBase, Parameters parameters,
 
     *logging::out << logging::Logger::INFO_INTERMEDIATE << "done!\n";
 }
+
+}
diff --git a/src/gpu/GksGpu/Output/VtkWriter.h b/src/gpu/GksGpu/Output/VtkWriter.h
index 0596fc7bd..679fae55b 100644
--- a/src/gpu/GksGpu/Output/VtkWriter.h
+++ b/src/gpu/GksGpu/Output/VtkWriter.h
@@ -38,9 +38,12 @@
 
 #include "GksGpu_export.h"
 
+namespace GksGpu {
+
 struct DataBase;
 struct Parameters;
 
+
 class GKSGPU_EXPORT VtkWriter
 {
 public:
@@ -49,4 +52,6 @@ public:
                        std::string filename );
 };
 
+}
+
 #endif
\ No newline at end of file
diff --git a/src/gpu/GksMeshAdapter/CMakeLists.txt b/src/gpu/GksMeshAdapter/CMakeLists.txt
index cb00b3c01..b9a2d12df 100644
--- a/src/gpu/GksMeshAdapter/CMakeLists.txt
+++ b/src/gpu/GksMeshAdapter/CMakeLists.txt
@@ -1,3 +1,3 @@
 project(GksMeshAdapter LANGUAGES CUDA CXX)
 
-vf_add_library(PRIVATE_LINK basics GridGenerator)
+vf_add_library(PRIVATE_LINK basics GridGenerator lbmCuda)
diff --git a/src/gpu/GksMeshAdapter/GksMeshAdapter.cpp b/src/gpu/GksMeshAdapter/GksMeshAdapter.cpp
index 16f5c2085..8d032dfee 100644
--- a/src/gpu/GksMeshAdapter/GksMeshAdapter.cpp
+++ b/src/gpu/GksMeshAdapter/GksMeshAdapter.cpp
@@ -22,6 +22,10 @@
 #include "MeshCell.h"
 #include "MeshFace.h"
 
+#include <lbm/constants/NumericConstants.h>
+
+using namespace vf::lbm::constant;
+
 using namespace vf::gpu;
 
 GksMeshAdapter::GksMeshAdapter(SPtr<MultipleGridBuilder> gridBuilder)
@@ -518,7 +522,7 @@ void GksMeshAdapter::sortFaces()
     // sort into blocks
     ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
 
-    std::array<char, 3> orientations = {'x', 'y', 'z'};
+    // std::array<char, 3> orientations = {'x', 'y', 'z'};
 
     for( uint level = 0; level < this->gridBuilder->getNumberOfLevels(); level++ )
     {
@@ -527,17 +531,17 @@ void GksMeshAdapter::sortFaces()
             uint start =         this->startOfFacesPerLevelXYZ [ 3 * level + idx];
             uint end   = start + this->numberOfFacesPerLevelXYZ[ 3 * level + idx];
 
-            real xMax = (*std::max_element(this->faces.begin() + start, this->faces.begin() + end, [this](MeshFace lhs, MeshFace rhs) { return lhs.faceCenter.x < rhs.faceCenter.x; })).faceCenter.x;
-            real yMax = (*std::max_element(this->faces.begin() + start, this->faces.begin() + end, [this](MeshFace lhs, MeshFace rhs) { return lhs.faceCenter.y < rhs.faceCenter.y; })).faceCenter.y;
-            real zMax = (*std::max_element(this->faces.begin() + start, this->faces.begin() + end, [this](MeshFace lhs, MeshFace rhs) { return lhs.faceCenter.z < rhs.faceCenter.z; })).faceCenter.z;
+            // real xMax = (*std::max_element(this->faces.begin() + start, this->faces.begin() + end, [this](MeshFace lhs, MeshFace rhs) { return lhs.faceCenter.x < rhs.faceCenter.x; })).faceCenter.x;
+            // real yMax = (*std::max_element(this->faces.begin() + start, this->faces.begin() + end, [this](MeshFace lhs, MeshFace rhs) { return lhs.faceCenter.y < rhs.faceCenter.y; })).faceCenter.y;
+            // real zMax = (*std::max_element(this->faces.begin() + start, this->faces.begin() + end, [this](MeshFace lhs, MeshFace rhs) { return lhs.faceCenter.z < rhs.faceCenter.z; })).faceCenter.z;
 
             real xMin = (*std::min_element(this->faces.begin() + start, this->faces.begin() + end, [this](MeshFace lhs, MeshFace rhs) { return lhs.faceCenter.x < rhs.faceCenter.x; })).faceCenter.x;
             real yMin = (*std::min_element(this->faces.begin() + start, this->faces.begin() + end, [this](MeshFace lhs, MeshFace rhs) { return lhs.faceCenter.y < rhs.faceCenter.y; })).faceCenter.y;
             real zMin = (*std::min_element(this->faces.begin() + start, this->faces.begin() + end, [this](MeshFace lhs, MeshFace rhs) { return lhs.faceCenter.z < rhs.faceCenter.z; })).faceCenter.z;
 
-            real xRange = xMax - xMin;
-            real yRange = yMax - yMin;
-            real zRange = zMax - zMin;
+            // real xRange = xMax - xMin;
+            // real yRange = yMax - yMin;
+            // real zRange = zMax - zMin;
 
             uint blockDim = 8;
 
diff --git a/src/lbm/CMakeLists.txt b/src/lbm/CMakeLists.txt
index 56b03bded..afa90bdd3 100644
--- a/src/lbm/CMakeLists.txt
+++ b/src/lbm/CMakeLists.txt
@@ -7,6 +7,6 @@ if(BUILD_VF_CPU)
     vf_add_tests()
 endif()
 
-if(BUILD_VF_GPU)
+if(BUILD_VF_GPU OR BUILD_VF_GKS)
     add_subdirectory(cuda)
 endif()
-- 
GitLab