merge fixes

src/gpu/VirtualFluids_GPU/Calculation/CollisisionStrategy.cpp

@@ -42,7 +42,8 @@ void CollisionAndExchange_noStreams_indexKernel::operator()(UpdateGrid27 *update
     updateGrid->collisionUsingIndices(  level, t, 
                                         para->getParD(level)->taggedFluidNodeIndices[CollisionTemplate::Default],
                                         para->getParD(level)->numberOfTaggedFluidNodes[CollisionTemplate::Default],  
-                                        CollisionTemplate::Default, -1);
+                                        CollisionTemplate::Default, 
+                                        CudaStreamIndex::Legacy);
 
     //! 2. exchange information between GPUs
     updateGrid->exchangeMultiGPU_noStreams_withPrepare(level, false);
@@ -81,7 +82,7 @@ void CollisionAndExchange_streams::operator()(UpdateGrid27 *updateGrid, Paramete
 
     //! 3. launch the collision kernel for bulk nodes
     //!
-    para->getStreamManager()->waitOnStartBulkKernelEvent(bulkStreamIndex);
+    para->getStreamManager()->waitOnStartBulkKernelEvent(CudaStreamIndex::Bulk);
     updateGrid->collisionUsingIndices(  level, t, 
                                         para->getParD(level)->taggedFluidNodeIndices[CollisionTemplate::Default],
                                         para->getParD(level)->numberOfTaggedFluidNodes[CollisionTemplate::Default], 

src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h

@@ -31,7 +31,7 @@ public:
 
 private:
     void collisionAllNodes(int level, unsigned int t);
-    void collisionUsingIndices(int level, unsigned int t, uint *indices, uint numberOfIndices = 0, CollisionTemplate collisionTemplate, CudaStreamIndex streamIndex=CudaStreamIndex::Legacy);
+    void collisionUsingIndices(int level, unsigned int t, uint *indices, uint numberOfIndices, CollisionTemplate collisionTemplate, CudaStreamIndex streamIndex=CudaStreamIndex::Legacy);
     void collisionAdvectionDiffusion(int level);
 
     void postCollisionBC(int level);

src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Almighty/CumulantK17Almighty.cu

@@ -37,14 +37,13 @@ void CumulantK17Almighty<turbulenceModel>::run()
 }
 
 template<TurbulenceModel turbulenceModel>
-void CumulantK17Almighty<turbulenceModel>::runOnIndices( const unsigned int *indices, unsigned int size_indices, CollisionTemplate collisionTemplate, int streamIndex )
+void CumulantK17Almighty<turbulenceModel>::runOnIndices( const unsigned int *indices, unsigned int size_indices, CollisionTemplate collisionTemplate, CudaStreamIndex streamIndex)
 {
-	cudaStream_t stream = (streamIndex == -1) ? CU_STREAM_LEGACY : para->getStreamManager()->getStream(streamIndex);
-
 	switch (collisionTemplate)
 	{
 		case CollisionTemplate::Default:
-			LB_Kernel_CumulantK17Almighty < turbulenceModel, false, false  > <<< cudaGrid.grid, cudaGrid.threads, 0, stream >>>(   	para->getParD(level)->omega, 	
+			LB_Kernel_CumulantK17Almighty < turbulenceModel, false, false  > <<< cudaGrid.grid, cudaGrid.threads, 0, para->getStreamManager()->getStream(streamIndex) >>>(   
+																																	para->getParD(level)->omega, 	
 																																	para->getParD(level)->typeOfGridNode, 										
 																																	para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ,	
 																																	para->getParD(level)->distributions.f[0],	
@@ -64,7 +63,8 @@ void CumulantK17Almighty<turbulenceModel>::runOnIndices( const unsigned int *ind
 			break;
 		
 		case CollisionTemplate::WriteMacroVars:
-			LB_Kernel_CumulantK17Almighty < turbulenceModel, true, false  > <<< cudaGrid.grid, cudaGrid.threads, 0, stream >>>( para->getParD(level)->omega, 	
+			LB_Kernel_CumulantK17Almighty < turbulenceModel, true, false  > <<< cudaGrid.grid, cudaGrid.threads, 0, para->getStreamManager()->getStream(streamIndex) >>>( 
+																																para->getParD(level)->omega, 	
 																																para->getParD(level)->typeOfGridNode, 										
 																																para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ,	
 																																para->getParD(level)->distributions.f[0],	
@@ -85,7 +85,8 @@ void CumulantK17Almighty<turbulenceModel>::runOnIndices( const unsigned int *ind
 		
 		case CollisionTemplate::Border:
 		case CollisionTemplate::AllFeatures:
-			LB_Kernel_CumulantK17Almighty < turbulenceModel, true, true  > <<< cudaGrid.grid, cudaGrid.threads, 0, stream >>>(  para->getParD(level)->omega, 	
+			LB_Kernel_CumulantK17Almighty < turbulenceModel, true, true  > <<< cudaGrid.grid, cudaGrid.threads, 0, para->getStreamManager()->getStream(streamIndex) >>>(  
+																																para->getParD(level)->omega, 	
 																																para->getParD(level)->typeOfGridNode, 										
 																																para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ,	
 																																para->getParD(level)->distributions.f[0],	
@@ -104,7 +105,8 @@ void CumulantK17Almighty<turbulenceModel>::runOnIndices( const unsigned int *ind
 																																size_indices);
 			break;
 		case CollisionTemplate::ApplyBodyForce:
-			LB_Kernel_CumulantK17Almighty < turbulenceModel, false, true  > <<< cudaGrid.grid, cudaGrid.threads, 0, stream >>>( para->getParD(level)->omega, 	
+			LB_Kernel_CumulantK17Almighty < turbulenceModel, false, true  > <<< cudaGrid.grid, cudaGrid.threads, 0, para->getStreamManager()->getStream(streamIndex) >>>( 
+																																para->getParD(level)->omega, 	
 																																para->getParD(level)->typeOfGridNode, 										
 																																para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ,	
 																																para->getParD(level)->distributions.f[0],	

src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Almighty/CumulantK17Almighty.h

@@ -10,7 +10,7 @@ class CumulantK17Almighty : public KernelImp
 public:
 	static std::shared_ptr< CumulantK17Almighty<turbulenceModel> > getNewInstance(std::shared_ptr< Parameter> para, int level);
 	void run() override;
-    void runOnIndices(const unsigned int *indices, unsigned int size_indices, CollisionTemplate collisionTemplate, int streamIndex = -1) override;
+    void runOnIndices(const unsigned int *indices, unsigned int size_indices, CollisionTemplate collisionTemplate, CudaStreamIndex streamIndex) override;
 
 private:
     CumulantK17Almighty();