diff --git a/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.cpp b/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.cpp index 3151e6bedeb6a96666f11f0040de2c95b20cc42c..e6960cc18ff7905fdcc351f3396a7a0948243dbf 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.cpp +++ b/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.cpp @@ -27,6 +27,25 @@ void KernelImp::setCheckParameterStrategy(std::shared_ptr<CheckParameterStrategy this->checkStrategy = strategy; } + + KernelImp::KernelImp(std::shared_ptr<Parameter> para, int level) : para(para), level(level) {} -KernelImp::KernelImp() {} \ No newline at end of file +KernelImp::KernelImp() {} + +std::unique_ptr<std::pair<dim3, dim3>> KernelImp::calcGridDimensions(unsigned int size_Mat, int numberOfThreads) +{ + 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); + std::pair<dim3, dim3> dimensions(grid, threads); + return std::make_unique<std::pair<dim3, dim3>>(dimensions); +} diff --git a/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h b/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h index c5215dbdd00e41a76b49c6c67808a6e7c44bb9a8..ca41116e40608e970b771e583858ec0631651b1c 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h +++ b/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h @@ -33,6 +33,8 @@ protected: KernelGroup myKernelGroup; vf::gpu::CudaGrid cudaGrid; + + std::unique_ptr<std::pair<dim3, dim3>> calcGridDimensions(unsigned int size_Mat, int numberOfThreads); }; #endif diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17/CumulantK17Comp.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17/CumulantK17Comp.cu index 86704c6143c2d48ee070dad58e12f49036dea43d..1caf0340c11cfc86214781ad89e2c6318730b5a6 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17/CumulantK17Comp.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17/CumulantK17Comp.cu @@ -10,23 +10,9 @@ std::shared_ptr<CumulantK17Comp> CumulantK17Comp::getNewInstance(std::shared_ptr void CumulantK17Comp::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); + dim3 grid, threads; + std::tie(grid, threads) = + *calcGridDimensions(para->getParD(level)->size_Mat_SP, para->getParD(level)->numberofthreads); LB_Kernel_CumulantK17Comp <<< grid, threads >>>(para->getParD(level)->omega, para->getParD(level)->geoSP, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chim/CumulantK17CompChim.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chim/CumulantK17CompChim.cu index 38df36655e3705d26d3ee9614970cde9d2fdc87b..a31505f33bcf8d1fda5f4feb28af5262baf1c4ac 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chim/CumulantK17CompChim.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chim/CumulantK17CompChim.cu @@ -10,23 +10,9 @@ std::shared_ptr<CumulantK17CompChim> CumulantK17CompChim::getNewInstance(std::sh void CumulantK17CompChim::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); + dim3 grid, threads; + std::tie(grid, threads) = + *calcGridDimensions(para->getParD(level)->size_Mat_SP, para->getParD(level)->numberofthreads); LB_Kernel_CumulantK17CompChim <<< grid, threads >>>( para->getParD(level)->omega, 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 3f493de191473307193230843bc366f516b444e3..b2ba5fdb1c7db76dccf3d0da1c494618ad6afaaa 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 @@ -14,7 +14,8 @@ std::shared_ptr<CumulantK17CompChimSparse> CumulantK17CompChimSparse::getNewInst void CumulantK17CompChimSparse::run() { dim3 grid, threads; - std::tie(grid, threads) = *calcGridDimensions(para->getParD(level)->numberOfFluidNodes); + std::tie(grid, threads) = + *calcGridDimensions(para->getParD(level)->numberOfFluidNodes, para->getParD(level)->numberofthreads); LB_Kernel_CumulantK17CompChimSparse <<< grid, threads >>>( para->getParD(level)->omega, @@ -35,7 +36,8 @@ void CumulantK17CompChimSparse::run() 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); + std::tie(grid, threads) = + *calcGridDimensions(para->getParD(level)->numberOfFluidNodes, para->getParD(level)->numberofthreads); cudaStream_t stream = (streamIndex == -1) ? CU_STREAM_LEGACY : para->getStreamManager().getStream(streamIndex); @@ -66,21 +68,3 @@ CumulantK17CompChimSparse::CumulantK17CompChimSparse(std::shared_ptr<Parameter> myKernelGroup = BasicKernel; } -std::unique_ptr<std::pair<dim3, dim3>> CumulantK17CompChimSparse::calcGridDimensions(unsigned int size_Mat) -{ - int numberOfThreads = para->getParD(level)->numberofthreads; - - 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); - std::pair<dim3, dim3> dimensions(grid, threads); - return std::make_unique<std::pair<dim3, dim3>>(dimensions); -} 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 906c880337d1747f96bc7a0f5d5b7988d6235ee5..8233e9b9b2ff98b17222cb55328ba5a7cd84cc3c 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 @@ -13,7 +13,6 @@ public: private: CumulantK17CompChimSparse(); CumulantK17CompChimSparse(std::shared_ptr<Parameter> para, int level); - std::unique_ptr<std::pair<dim3, dim3>> calcGridDimensions(unsigned int size_Mat); }; #endif