move precursor Stream to streammanager

src/gpu/VirtualFluids_GPU/BoundaryConditions/BoundaryConditionStructs.cuh

 #ifndef BoundaryConditionStructs_H
 #define BoundaryConditionStructs_H
 #include "Core/DataTypes.h"
-#include <cuda_runtime.h>
 
 //Q for second order BCs
 //! \struct to manage sub-grid-distances (q) for second order Boundary Conditions (BCs)
@@ -33,7 +32,6 @@ typedef struct QforPrecursorBC{
    real* weightsNT, *weightsNB, *weightsST,  *weightsSB;
    real* last, *current, *next;
    real velocityX, velocityY, velocityZ;
-   cudaStream_t stream;
 }QforPrecursorBoundaryConditions;
 
 //BCTemp

src/gpu/VirtualFluids_GPU/CMakeLists.txt

@@ -22,8 +22,5 @@ if(BUILD_VF_UNIT_TESTS)
 	set_source_files_properties(DataStructureInitializer/GridReaderGenerator/IndexRearrangementForStreamsTest.cpp PROPERTIES LANGUAGE CUDA)
     set_source_files_properties(Communication/ExchangeData27Test.cpp PROPERTIES LANGUAGE CUDA)
     set_source_files_properties(BoundaryConditions/BoundaryConditionFactoryTest.cpp PROPERTIES LANGUAGE CUDA)
-    set_source_files_properties(Parameter/ParameterTest.cpp PROPERTIES LANGUAGE CUDA)
-    set_source_files_properties(KernelManager/BCKernelManagerTest.cpp PROPERTIES LANGUAGE CUDA)
-    set_source_files_properties(Parameter/EdgeNodeFinderTest.cpp PROPERTIES LANGUAGE CUDA)
     target_include_directories(VirtualFluids_GPUTests PRIVATE "${VF_THIRD_DIR}/cuda_samples/")
 endif()

src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp

@@ -1704,7 +1704,6 @@ void CudaMemoryManager::cudaAllocPrecursorData(int lev)
 {
     size_t size = parameter->getParH(lev)->precursorBC.numberOfPrecursorNodes*sizeof(real)*parameter->getParH(lev)->precursorBC.numberOfQuantities;
 
-    checkCudaErrors( cudaStreamCreate(&parameter->getParH(lev)->precursorBC.stream) );
     checkCudaErrors( cudaMallocHost((void**) &parameter->getParH(lev)->precursorBC.last, size));
     checkCudaErrors( cudaMallocHost((void**) &parameter->getParH(lev)->precursorBC.current, size));
     checkCudaErrors( cudaMallocHost((void**) &parameter->getParH(lev)->precursorBC.next, size));
@@ -1739,9 +1738,10 @@ void CudaMemoryManager::cudaCopyPrecursorBC(int lev)
 void CudaMemoryManager::cudaCopyPrecursorData(int lev)
 {
     auto prec = &parameter->getParH(lev)->precursorBC;
+    auto precStream = parameter->getStreamManager()->getStream(parameter->getStreamManager()->getPrecursorStreamIndex());
     size_t memSize = prec->numberOfPrecursorNodes*sizeof(real)*prec->numberOfQuantities;
-    checkCudaErrors( cudaStreamSynchronize(prec->stream) );
-    checkCudaErrors( cudaMemcpyAsync(parameter->getParD(lev)->precursorBC.next, prec->next, memSize, cudaMemcpyHostToDevice, prec->stream)) ;
+    checkCudaErrors( cudaStreamSynchronize(precStream) );
+    checkCudaErrors( cudaMemcpyAsync(parameter->getParD(lev)->precursorBC.next, prec->next, memSize, cudaMemcpyHostToDevice, precStream) );
 }
 
 

src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp

@@ -115,7 +115,7 @@ void Simulation::init(GridProvider &gridProvider, BoundaryConditionFactory *bcFa
     //////////////////////////////////////////////////////////////////////////
     // CUDA streams
     if (para->getUseStreams()) {
-        para->getStreamManager()->launchStreams(2u);
+        para->getStreamManager()->launchStreams(3u);
         para->getStreamManager()->createCudaEvents();
     }
     //////////////////////////////////////////////////////////////////////////

src/gpu/VirtualFluids_GPU/Parameter/CudaStreamManager.cpp

@@ -51,6 +51,8 @@ int CudaStreamManager::getBorderStreamIndex() { return borderStreamIndex; }
 
 int CudaStreamManager::getBulkStreamIndex() { return bulkStreamIndex; }
 
+int CudaStreamManager::getPrecursorStreamIndex() { return precursorStreamIndex; }
+
 void CudaStreamManager::createCudaEvents()
 {
     checkCudaErrors(cudaEventCreateWithFlags(&startBulkKernel, cudaEventDisableTiming));

src/gpu/VirtualFluids_GPU/Parameter/CudaStreamManager.h

@@ -40,6 +40,7 @@ class CudaStreamManager
 private:
     std::vector<cudaStream_t> cudaStreams;
     cudaEvent_t startBulkKernel = NULL;
+    const int precursorStreamIndex    = 2;
     const int borderStreamIndex       = 1;
     const int bulkStreamIndex         = 0;
 
@@ -47,9 +48,9 @@ public:
     void launchStreams(uint numberOfStreams);
     void terminateStreams();
     cudaStream_t &getStream(uint streamIndex);
-
     int getBorderStreamIndex();
     int getBulkStreamIndex();
+    int getPrecursorStreamIndex();
 
     // Events
     void createCudaEvents();