Add test kernel for Cuda Streams

apps/gpu/LBM/MusselOyster/MusselOyster.cpp

@@ -144,7 +144,7 @@ void multipleLevel(const std::string& configPath)
     // para->setMaxLevel(2);
 
     // para->setMainKernel("CumulantK15Comp");
-    para->setMainKernel("CumulantK17CompChim");
+    para->setMainKernel("CumulantK17CompChimSparse");
 
     if (useMultiGPU) {
         para->setDevices(std::vector<uint>{ (uint)0, (uint)1 });
@@ -265,8 +265,6 @@ void multipleLevel(const std::string& configPath)
 
             gridBuilder->setVelocityBoundaryCondition(SideType::GEOMETRY, 0.0, 0.0, 0.0);
             //////////////////////////////////////////////////////////////////////////
-            SPtr<Grid> grid = gridBuilder->getGrid(gridBuilder->getNumberOfLevels() - 1);
-            //////////////////////////////////////////////////////////////////////////
 
             gridBuilder->writeGridsToVtk("E:/temp/MusselOyster/" + bivalveType + "/grid/");
             // gridBuilder->writeArrows ("E:/temp/MusselOyster/" + bivalveType + "/arrow");

src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse.cu

+#include "CumulantK17CompChimSparse.h"
+
+#include "Parameter/Parameter.h"
+#include "CumulantK17CompChimSparse_Device.cuh"
+
+std::shared_ptr<CumulantK17CompChimSparse> CumulantK17CompChimSparse::getNewInstance(std::shared_ptr<Parameter> para,
+                                                                               int level)
+{
+    return std::shared_ptr<CumulantK17CompChimSparse>(new CumulantK17CompChimSparse(para, level));
+}
+
+void CumulantK17CompChimSparse::run()
+{
+	int numberOfThreads = para->getParD(level)->numberofthreads;
+	int size_Mat = para->getParD(level)->size_Mat_SP;
+
+	int Grid = (size_Mat / numberOfThreads) + 1;
+	int Grid1, Grid2;
+	if (Grid>512)
+	{
+		Grid1 = 512;
+		Grid2 = (Grid / Grid1) + 1;
+	}
+	else
+	{
+		Grid1 = 1;
+		Grid2 = Grid;
+	}
+	dim3 grid(Grid1, Grid2);
+	dim3 threads(numberOfThreads, 1, 1);
+
+	LB_Kernel_CumulantK17CompChimSparse<<<grid, threads>>>(
+		para->getParD(level)->omega,
+		para->getParD(level)->geoSP,
+		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->getForcesDev(),
+        para->getQuadricLimitersDev(),
+		para->getParD(level)->evenOrOdd);
+	getLastCudaError("LB_Kernel_CumulantK17CompChim execution failed");
+}
+
+CumulantK17CompChimSparse::CumulantK17CompChimSparse(std::shared_ptr<Parameter> para, int level)
+{
+	this->para = para;
+	this->level = level;
+
+	myPreProcessorTypes.push_back(InitCompSP27);
+
+	myKernelGroup = BasicKernel;
+}
\ No newline at end of file

src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse.h

+#ifndef CUMULANT_K17_COMP_CHIM_SPARSE_H
+#define CUMULANT_K17_COMP_CHIM_SPARSE_H
+
+#include "Kernel/KernelImp.h"
+
+class CumulantK17CompChimSparse : public KernelImp
+{
+public:
+    static std::shared_ptr<CumulantK17CompChimSparse> getNewInstance(std::shared_ptr<Parameter> para, int level);
+	void run();
+
+private:
+    CumulantK17CompChimSparse();
+    CumulantK17CompChimSparse(std::shared_ptr<Parameter> para, int level);
+};
+
+#endif 

src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse_Device.cuh

+#ifndef LB_Kernel_CUMULANT_K17_COMP_CHIM_SPARSE_H
+#define LB_Kernel_CUMULANT_K17_COMP_CHIM_SPARSE_H
+
+#include <DataTypes.h>
+#include <curand.h>
+
+extern "C" __global__ void LB_Kernel_CumulantK17CompChimSparse(
+	real omega,
+	uint* typeOfGridNode,
+	uint* neighborX,
+	uint* neighborY,
+	uint* neighborZ,
+	real* distributions,
+	int size_Mat,
+	int level,
+	real* forces,
+	real* quadricLimiters,
+	bool isEvenTimestep);
+#endif

src/gpu/VirtualFluids_GPU/Kernel/Utilities/KernelFactory/KernelFactoryImp.cpp

@@ -11,6 +11,7 @@
 #include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17/CumulantK17Comp.h"
 #include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Unified/CumulantK17Unified.h"
 #include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chim/CumulantK17CompChim.h"
+#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse.h"
 #include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Bulk/CumulantK17BulkComp.h"
 #include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27.h"
 #include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp.h"
@@ -136,6 +137,9 @@ std::shared_ptr<Kernel> KernelFactoryImp::makeKernel(std::shared_ptr<Parameter>
     } else if (kernel == "CumulantK17CompChim") {
         newKernel     = CumulantK17CompChim::getNewInstance(para, level);
         checkStrategy = FluidFlowCompStrategy::getInstance();
+    } else if (kernel == "CumulantK17CompChimSparse") {
+        newKernel     = CumulantK17CompChimSparse::getNewInstance(para, level);
+        checkStrategy = FluidFlowCompStrategy::getInstance();
     } else if (kernel == "CumulantAll4CompSP27") {
         newKernel     = CumulantAll4CompSP27::getNewInstance(para, level);
         checkStrategy = FluidFlowCompStrategy::getInstance();