From c550628975a58ffa9a73bb59efa8eb2ede4c0296 Mon Sep 17 00:00:00 2001
From: Anna Wellmann <a.wellmann@tu-braunschweig.de>
Date: Thu, 5 Aug 2021 17:40:41 +0200
Subject: [PATCH] Pass stream to collision kernel
---
src/gpu/GridGenerator/grid/Grid.h | 2 +-
src/gpu/GridGenerator/grid/GridImp.cu | 2 +-
.../Calculation/UpdateGrid27.cpp | 37 +++++++++++++------
.../Calculation/UpdateGrid27.h | 4 +-
src/gpu/VirtualFluids_GPU/Kernel/Kernel.h | 2 +-
.../VirtualFluids_GPU/Kernel/KernelImp.cpp | 2 +-
src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h | 2 +-
.../CumulantK17CompChimSparse.cu | 29 ++++++++++-----
.../CumulantK17CompChimSparse.h | 4 +-
src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp | 2 +-
.../VirtualFluids_GPU/Parameter/Parameter.h | 2 +-
11 files changed, 57 insertions(+), 31 deletions(-)
diff --git a/src/gpu/GridGenerator/grid/Grid.h b/src/gpu/GridGenerator/grid/Grid.h
index 32f571791..d4621562a 100644
--- a/src/gpu/GridGenerator/grid/Grid.h
+++ b/src/gpu/GridGenerator/grid/Grid.h
@@ -143,7 +143,7 @@ public:
// needed for CUDA Streams
virtual void findFluidNodeIndices(bool onlyBulk) = 0;
- virtual uint getNumberOfFluidNodes() const = 0;;
+ virtual uint getNumberOfFluidNodes() const = 0;
virtual void getFluidNodeIndices(uint *fluidNodeIndices) const = 0;
virtual uint getNumberOfFluidNodesBorder() const = 0;
diff --git a/src/gpu/GridGenerator/grid/GridImp.cu b/src/gpu/GridGenerator/grid/GridImp.cu
index ee54a8b02..3cb312617 100644
--- a/src/gpu/GridGenerator/grid/GridImp.cu
+++ b/src/gpu/GridGenerator/grid/GridImp.cu
@@ -1974,7 +1974,7 @@ CUDA_HOST void GridImp::getFluidNodeIndices(uint *fluidNodeIndices) const
uint GridImp::getNumberOfFluidNodesBorder() const
{
- return this->fluidNodeIndicesBorder.size();
+ return (uint)this->fluidNodeIndicesBorder.size();
}
void GridImp::getFluidNodeIndicesBorder(uint *fluidNodeIndicesBorder) const
diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
index 445244ac7..9d3bd9030 100644
--- a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
+++ b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
@@ -27,10 +27,10 @@ void updateGrid27(Parameter* para,
//////////////////////////////////////////////////////////////////////////
if (para->useStreams) {
- collision(para, pm, level, t, kernels, para->getParD(level)->fluidNodeIndices,
- para->getParD(level)->numberOfFluidNodes);
- collision(para, pm, level, t, kernels, para->getParD(level)->fluidNodeIndicesBorder,
- para->getParD(level)->numberOffluidNodesBorder);
+ collisionUsingIndex(para, pm, level, t, kernels, para->getParD(level)->fluidNodeIndices,
+ para->getParD(level)->numberOfFluidNodes, 0);
+ collisionUsingIndex(para, pm, level, t, kernels, para->getParD(level)->fluidNodeIndicesBorder,
+ para->getParD(level)->numberOffluidNodesBorder, 1);
}
else
collision(para, pm, level, t, kernels);
@@ -68,15 +68,30 @@ void updateGrid27(Parameter* para,
}
}
-void collision(Parameter* para, std::vector<std::shared_ptr<PorousMedia>>& pm, int level, unsigned int t, std::vector < SPtr< Kernel>>& kernels, uint* fluidNodeIndices, uint numberOfFluidNodes)
+void collision(Parameter* para, std::vector<std::shared_ptr<PorousMedia>>& pm, int level, unsigned int t, std::vector < SPtr< Kernel>>& kernels)
{
- if (para->useStreams)
- if (fluidNodeIndices != nullptr && numberOfFluidNodes != 0)
- kernels.at(level)->runOnIndices(fluidNodeIndices, numberOfFluidNodes);
- else
- std::cout << "in collision: fluidNodeIndices or numberOfFluidNodes not definded" << std::endl; // better use logger
+ kernels.at(level)->run();
+
+ //////////////////////////////////////////////////////////////////////////
+
+ if (para->getSimulatePorousMedia())
+ collisionPorousMedia(para, pm, level);
+
+ //////////////////////////////////////////////////////////////////////////
+
+ if (para->getDiffOn())
+ collisionAdvectionDiffusion(para, level);
+}
+
+void collisionUsingIndex(Parameter *para, std::vector<std::shared_ptr<PorousMedia>> &pm, int level, unsigned int t,
+ std::vector<SPtr<Kernel>> &kernels,
+ uint *fluidNodeIndices, uint numberOfFluidNodes, int stream)
+{
+ if (fluidNodeIndices != nullptr && numberOfFluidNodes != 0)
+ kernels.at(level)->runOnIndices(fluidNodeIndices, numberOfFluidNodes, stream);
else
- kernels.at(level)->run();
+ std::cout << "in collision: fluidNodeIndices or numberOfFluidNodes not definded"
+ << std::endl;
//////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h
index 0ce9383b9..44b02d36f 100644
--- a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h
+++ b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h
@@ -18,7 +18,9 @@ extern "C" void updateGrid27(Parameter* para,
unsigned int t,
std::vector < SPtr< Kernel>>& kernels);
-extern "C" void collision(Parameter *para, std::vector<std::shared_ptr<PorousMedia>> &pm, int level, unsigned int t, std::vector<SPtr<Kernel>> &kernels, uint *fluidNodeIndices = nullptr, uint numberOfFluidNodes = 0);
+extern "C" void collision(Parameter *para, std::vector<std::shared_ptr<PorousMedia>> &pm, int level, unsigned int t, std::vector<SPtr<Kernel>> &kernels);
+
+extern "C" void collisionUsingIndex(Parameter *para, std::vector<std::shared_ptr<PorousMedia>> &pm, int level, unsigned int t, std::vector<SPtr<Kernel>> &kernels, uint *fluidNodeIndices = nullptr, uint numberOfFluidNodes = 0, int stream = -1);
extern "C" void collisionPorousMedia(Parameter* para, std::vector<std::shared_ptr<PorousMedia>>& pm, int level);
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernel.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernel.h
index ceb1638ad..9f9f7539b 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernel.h
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernel.h
@@ -13,7 +13,7 @@ class Kernel
public:
virtual ~Kernel() = default;
virtual void run() = 0;
- virtual void runOnIndices(const unsigned int *indices, unsigned int size_indices) = 0;
+ virtual void runOnIndices(const unsigned int *indices, unsigned int size_indices, int stream = -1) = 0; //if stream == -1: run on default stream
virtual bool checkParameter() = 0;
virtual std::vector<PreProcessorType> getPreProcessorTypes() = 0;
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.cpp b/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.cpp
index cd06a7c94..3151e6bed 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.cpp
+++ b/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.cpp
@@ -3,7 +3,7 @@
#include "Kernel/Utilities/CheckParameterStrategy/CheckParameterStrategy.h"
-void KernelImp::runOnIndices(const unsigned int *indices, unsigned int size_indices)
+void KernelImp::runOnIndices(const unsigned int *indices, unsigned int size_indices, int stream)
{
printf("Method not implemented for this Kernel \n");
}
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h b/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h
index 0986e7b97..c5215dbdd 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h
+++ b/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h
@@ -14,7 +14,7 @@ class KernelImp : public Kernel
{
public:
virtual void run() = 0;
- virtual void runOnIndices(const unsigned int *indices, unsigned int size_indices);
+ virtual void runOnIndices(const unsigned int *indices, unsigned int size_indices, int stream = -1);
bool checkParameter();
std::vector<PreProcessorType> getPreProcessorTypes();
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse.cu
index 6c3555c8f..afad31181 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse.cu
@@ -3,6 +3,8 @@
#include "Parameter/Parameter.h"
#include "CumulantK17CompChimSparse_Device.cuh"
+#include <cuda.h>
+
std::shared_ptr<CumulantK17CompChimSparse> CumulantK17CompChimSparse::getNewInstance(std::shared_ptr<Parameter> para,
int level)
{
@@ -30,25 +32,32 @@ void CumulantK17CompChimSparse::run()
getLastCudaError("LB_Kernel_CumulantK17CompChim execution failed");
}
-void CumulantK17CompChimSparse::runOnIndices(const unsigned int *indices, unsigned int size_indices)
+void CumulantK17CompChimSparse::runOnIndices(const unsigned int *indices, unsigned int size_indices, int streamIndex)
{
dim3 grid, threads;
std::tie(grid, threads) = *calcGridDimensions(para->getParD(level)->numberOfFluidNodes);
- LB_Kernel_CumulantK17CompChimSparse<<<grid, threads, 0, para->getStream(0)>>>(
+ cudaStream_t stream;
+ if (streamIndex == -1)
+ stream = CU_STREAM_LEGACY;
+ else
+ stream = para->getStream(streamIndex);
+
+ LB_Kernel_CumulantK17CompChimSparse<<<grid, threads, 0, stream>>>(
para->getParD(level)->omega,
- para->getParD(level)->neighborX_SP,
- para->getParD(level)->neighborY_SP,
+ para->getParD(level)->neighborX_SP,
+ para->getParD(level)->neighborY_SP,
para->getParD(level)->neighborZ_SP,
- para->getParD(level)->d0SP.f[0],
- para->getParD(level)->size_Mat_SP,
- level,
+ para->getParD(level)->d0SP.f[0],
+ para->getParD(level)->size_Mat_SP,
+ level,
para->getForcesDev(),
- para->getQuadricLimitersDev(),
- para->getParD(level)->evenOrOdd,
+ para->getQuadricLimitersDev(),
+ para->getParD(level)->evenOrOdd,
indices,
- size_indices);
+ size_indices);
getLastCudaError("LB_Kernel_CumulantK17CompChim execution failed");
+
}
CumulantK17CompChimSparse::CumulantK17CompChimSparse(std::shared_ptr<Parameter> para, int level)
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse.h
index f1645cfae..906c88033 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse.h
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse.h
@@ -7,8 +7,8 @@ class CumulantK17CompChimSparse : public KernelImp
{
public:
static std::shared_ptr<CumulantK17CompChimSparse> getNewInstance(std::shared_ptr<Parameter> para, int level);
- void run();
- void runOnIndices(const unsigned int *indices, unsigned int size_indices) override;
+ void run() override;
+ void runOnIndices(const unsigned int *indices, unsigned int size_indices, int stream = -1) override;
private:
CumulantK17CompChimSparse();
diff --git a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
index 34c45493d..9bbfa66b6 100644
--- a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
+++ b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
@@ -106,7 +106,7 @@ void Simulation::init(SPtr<Parameter> para, SPtr<GridProvider> gridProvider, std
//////////////////////////////////////////////////////////////////////////
// CUDA streams
if(para->useStreams)
- para->launchStreams((uint)1);
+ para->launchStreams((uint)2);
//////////////////////////////////////////////////////////////////////////
//
//output << para->getNeedInterface().at(0) << "\n";
diff --git a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
index 9f70266e4..5d24abe4f 100644
--- a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
+++ b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
@@ -43,8 +43,8 @@
#include "VirtualFluids_GPU_export.h"
+#include <cuda.h>
#include <cuda_runtime.h>
-#include <helper_cuda.h>
struct curandStateXORWOW;
typedef struct curandStateXORWOW curandState;
--
GitLab