diff --git a/src/gpu/VirtualFluids_GPU/Calculation/CollisisionStrategy.cpp b/src/gpu/VirtualFluids_GPU/Calculation/CollisisionStrategy.cpp index 2c1d3e74e8cb3507f364e8e40c80f72648b665da..819119176f131d50a569dbf63c863d346149585c 100644 --- a/src/gpu/VirtualFluids_GPU/Calculation/CollisisionStrategy.cpp +++ b/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], diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h index ac597bcb463c9e4b5959b92cf8ad4e04c9ce3eb8..6f5db02aee211021994fea4257e4b19feaa56c48 100644 --- a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h +++ b/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); diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Almighty/CumulantK17Almighty.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Almighty/CumulantK17Almighty.cu index 4498768999cdd5db826b852121760ac588fd1d04..8ddaa6381d65707936b580509bfd277b3f990bda 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Almighty/CumulantK17Almighty.cu +++ b/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], diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Almighty/CumulantK17Almighty.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Almighty/CumulantK17Almighty.h index b49b34e632712d9910551e9ecb46e11a18a9ee86..1819537c095c0106b707ab551b798c0ed62c0e41 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Almighty/CumulantK17Almighty.h +++ b/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();