From 94b9097ab6dd1b45316652717a5eae49cdbd5149 Mon Sep 17 00:00:00 2001 From: Anna Wellmann <a.wellmann@tu-braunschweig.de> Date: Mon, 6 Sep 2021 17:13:48 +0200 Subject: [PATCH] Fix some bugs in f to c --- apps/gpu/LBM/MusselOyster/MusselOyster.cpp | 4 ++-- .../VirtualFluids_GPU/Calculation/UpdateGrid27.cpp | 2 +- src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h | 1 + src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu | 11 ++++++----- 4 files changed, 10 insertions(+), 8 deletions(-) diff --git a/apps/gpu/LBM/MusselOyster/MusselOyster.cpp b/apps/gpu/LBM/MusselOyster/MusselOyster.cpp index 709778989..12ef13eba 100644 --- a/apps/gpu/LBM/MusselOyster/MusselOyster.cpp +++ b/apps/gpu/LBM/MusselOyster/MusselOyster.cpp @@ -118,7 +118,7 @@ void multipleLevel(const std::string& configPath) std::string bivalveType = "MUSSEL"; // "MUSSEL" "OYSTER" std::string gridPath(gridPathParent + bivalveType); // only for GridGenerator, for GridReader the gridPath needs to be set in the config file - real dxGrid = (real)1.0; + real dxGrid = (real)2.0; real vxLB = (real)0.051; // LB units real Re = (real)300.0; real viscosityLB = (vxLB * dxGrid) / Re; @@ -155,9 +155,9 @@ void multipleLevel(const std::string& configPath) para->setMaxLevel(1); - //para->setMainKernel("CumulantK17CompChim"); if (useStreams) para->setUseStreams(); + //para->setMainKernel("CumulantK17CompChim"); para->setMainKernel("CumulantK17CompChimStream"); *logging::out << logging::Logger::INFO_HIGH << "Kernel: " << para->getMainKernel() << "\n"; diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp index af0511662..4da3c6745 100644 --- a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp +++ b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp @@ -1151,7 +1151,7 @@ void fineToCoarseUsingIndex(Parameter *para, int level, uint *fluidNodeIndices, para->getParD(level)->intFC.ICellFCC, para->getParD(level)->intFC.ICellFCF, para->getParD(level)->K_FC, para->getParD(level)->omega, para->getParD(level + 1)->omega, para->getParD(level)->vis, para->getParD(level)->nx, para->getParD(level)->ny, para->getParD(level + 1)->nx, para->getParD(level + 1)->ny, - para->getParD(level)->offFC, fluidNodeIndices, numberOfFluidNodes, stream); + para->getParD(level)->numberofthreads, para->getParD(level)->offFC, fluidNodeIndices, numberOfFluidNodes, stream); getLastCudaError("ScaleFC27_RhoSq_comp_Stream execution failed"); ////////////////////////////////////////////////////////////////////////// diff --git a/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h b/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h index 80f683ea0..1220b5ef4 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h +++ b/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h @@ -2115,6 +2115,7 @@ extern "C" void ScaleFC_RhoSq_comp_27_Stream( real* DC, unsigned int nyC, unsigned int nxF, unsigned int nyF, + unsigned int numberOfThreads, OffFC offFC, unsigned int *fluidNodeIndices, unsigned int numberOfFluidNodes, diff --git a/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu b/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu index b2a848e98..feab617a2 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu +++ b/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu @@ -6242,7 +6242,7 @@ extern "C" void ScaleFC_RhoSq_comp_27(real* DC, offFC); getLastCudaError("scaleFC_RhoSq_27 execution failed"); } -extern "C" void ScaleFC_RhoSq_comp_27_Stream(real *DC, +extern "C" void ScaleFC_RhoSq_comp_27_Stream(real * DC, real * DF, unsigned int * neighborCX, unsigned int * neighborCY, @@ -6263,12 +6263,13 @@ extern "C" void ScaleFC_RhoSq_comp_27_Stream(real *DC, unsigned int nyC, unsigned int nxF, unsigned int nyF, + unsigned int numberOfThreads, OffFC offFC, unsigned int *fluidNodeIndices, unsigned int numberOfFluidNodes, CUstream_st* stream) { - int Grid = (kFC / numberOfFluidNodes) + 1; + int Grid = (kFC / numberOfThreads) + 1; int Grid1, Grid2; if (Grid > 512) { Grid1 = 512; @@ -6278,11 +6279,11 @@ extern "C" void ScaleFC_RhoSq_comp_27_Stream(real *DC, Grid2 = Grid; } dim3 gridINT_FC(Grid1, Grid2); - dim3 threads(numberOfFluidNodes, 1, 1); + dim3 threads(numberOfThreads, 1, 1); - scaleFC_RhoSq_comp_27_Stream<<<gridINT_FC, threads, 0, stream>>>(DC, DF, neighborCX, neighborCY, neighborCZ, neighborFX, neighborFY, neighborFZ, size_MatC, size_MatF, evenOrOdd, + scaleFC_RhoSq_comp_27_Stream<<<gridINT_FC, threads, 0 , stream>>>(DC, DF, neighborCX, neighborCY, neighborCZ, neighborFX, neighborFY, neighborFZ, size_MatC, size_MatF, evenOrOdd, posC, posFSWB, kFC, omCoarse, omFine, nu, nxC, nyC, nxF, nyF, offFC, fluidNodeIndices, numberOfFluidNodes); - getLastCudaError("scaleFC_RhoSq_27_Stream execution failed"); + getLastCudaError("scaleFC_RhoSq_comp_27_Stream execution failed"); } ////////////////////////////////////////////////////////////////////////// extern "C" void ScaleFC_RhoSq_3rdMom_comp_27( real* DC, -- GitLab