From 558dc74becf387279ff8a09016506c511a579a12 Mon Sep 17 00:00:00 2001
From: HenrikAsmuth <henrik.asmuth@geo.uu.se>
Date: Thu, 13 Oct 2022 13:11:11 +0200
Subject: [PATCH] merge fixes
---
.../Calculation/CollisisionStrategy.cpp | 5 +++--
.../VirtualFluids_GPU/Calculation/UpdateGrid27.h | 2 +-
.../CumulantK17Almighty/CumulantK17Almighty.cu | 16 +++++++++-------
.../CumulantK17Almighty/CumulantK17Almighty.h | 2 +-
4 files changed, 14 insertions(+), 11 deletions(-)
diff --git a/src/gpu/VirtualFluids_GPU/Calculation/CollisisionStrategy.cpp b/src/gpu/VirtualFluids_GPU/Calculation/CollisisionStrategy.cpp
index 2c1d3e74e..819119176 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 ac597bcb4..6f5db02ae 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 449876899..8ddaa6381 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 b49b34e63..1819537c0 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();
--
GitLab