From d2003b35e6149aedecaf96a49cac2217f9616aab Mon Sep 17 00:00:00 2001 From: Hamidreza Norouzi Date: Wed, 22 May 2024 09:52:48 +0330 Subject: [PATCH] bug fix to work with CUDA run --- .../twoPartContactSearchKernels.cpp | 8 +-- .../domain/MPISimulationDomain.cpp | 9 ++- .../pointField/processorBoundaryField.cpp | 28 ++++++++-- .../boundaries/boundaryProcessor.cpp | 26 ++++----- .../boundaries/boundaryProcessorKernels.hpp | 56 +++++++++++++++++++ .../pointStructure/boundaries/dataSender.hpp | 1 - 6 files changed, 100 insertions(+), 28 deletions(-) create mode 100644 src/phasicFlow/MPIParallelization/pointStructure/boundaries/boundaryProcessorKernels.hpp diff --git a/src/Interaction/contactSearch/boundaries/twoPartContactSearch/twoPartContactSearchKernels.cpp b/src/Interaction/contactSearch/boundaries/twoPartContactSearch/twoPartContactSearchKernels.cpp index 515e5af1..9faa80fc 100644 --- a/src/Interaction/contactSearch/boundaries/twoPartContactSearch/twoPartContactSearchKernels.cpp +++ b/src/Interaction/contactSearch/boundaries/twoPartContactSearch/twoPartContactSearchKernels.cpp @@ -87,14 +87,14 @@ pFlow::twoPartContactSearchKernels::broadSearchPP( continue; uint32 thisI = head(ind.x(), ind.y(), ind.z()); - while (thisI != -1) + while (thisI != static_cast(-1)) { auto d_n = sizeRatio * diams[thisI]; // first item is for this boundary and second itme, // for mirror if(sphereSphereCheckB(p_m, points[thisI], d_m, d_n)&& - ppPairs.insert(thisI,mrrI) == -1) + ppPairs.insert(thisI,mrrI) == static_cast(-1)) { getFullUpdate++; } @@ -161,14 +161,14 @@ pFlow::twoPartContactSearchKernels::broadSearchPP( } uint32 i1 = head(ind.x(), ind.y(), ind.z()); - while (i1 != -1) + while (i1 != static_cast(-1)) { auto d_n = sizeRatio * diams1[i1]; // first item is for this boundary and second itme, // for mirror if(sphereSphereCheckB(p_m, points1[i1], d_m, d_n)&& - ppPairs.insert(i1,i2) == -1) + ppPairs.insert(i1,i2) == static_cast(-1)) { getFullUpdate++; } diff --git a/src/phasicFlow/MPIParallelization/domain/MPISimulationDomain.cpp b/src/phasicFlow/MPIParallelization/domain/MPISimulationDomain.cpp index 87b050eb..9a5ee76e 100644 --- a/src/phasicFlow/MPIParallelization/domain/MPISimulationDomain.cpp +++ b/src/phasicFlow/MPIParallelization/domain/MPISimulationDomain.cpp @@ -41,9 +41,12 @@ bool pFlow::MPI::MPISimulationDomain::createBoundaryDicts() auto& mpiBoundaries = this->subDict("MPIBoundaries"); real neighborLength = boundaries.getVal("neighborLength"); - auto boundaryExtntionLengthRatio = - boundaries.getValOrSet("boundaryExtntionLengthRatio", 0.1); - auto updateIntercal = boundaries.getValOrSet("updateInterval", 1u); + auto boundaryExtntionLengthRatio = max( + boundaries.getValOrSet("boundaryExtntionLengthRatio", 0.1), + 0.0); + auto updateIntercal = max( + boundaries.getValOrSet("updateInterval", 1u), + 1u); auto neighbors = findPlaneNeighbors(); diff --git a/src/phasicFlow/MPIParallelization/pointField/processorBoundaryField.cpp b/src/phasicFlow/MPIParallelization/pointField/processorBoundaryField.cpp index 820831b4..b7348a2a 100644 --- a/src/phasicFlow/MPIParallelization/pointField/processorBoundaryField.cpp +++ b/src/phasicFlow/MPIParallelization/pointField/processorBoundaryField.cpp @@ -136,12 +136,28 @@ bool pFlow::MPI::processorBoundaryField::hearChanges( message::eventName(message::BNDR_PROCTRANSFER_SEND) ); - FieldAccessType transferData( - indices.size(), - indices.deviceViewAll(), - this->internal().deviceViewAll() - ); - sender_.sendData(pFlowProcessors(),transferData); + if constexpr( isDeviceAccessible()) + { + FieldAccessType transferData( + indices.size(), + indices.deviceViewAll(), + this->internal().deviceViewAll() + ); + + sender_.sendData(pFlowProcessors(),transferData); + } + else + { + FieldAccessType transferData( + indices.size(), + indices.hostViewAll(), + this->internal().deviceViewAll() + ); + + sender_.sendData(pFlowProcessors(),transferData); + } + + } else if(msg.equivalentTo(message::BNDR_PROCTRANSFER_RECIEVE)) { diff --git a/src/phasicFlow/MPIParallelization/pointStructure/boundaries/boundaryProcessor.cpp b/src/phasicFlow/MPIParallelization/pointStructure/boundaries/boundaryProcessor.cpp index 76be7508..54bed6ad 100644 --- a/src/phasicFlow/MPIParallelization/pointStructure/boundaries/boundaryProcessor.cpp +++ b/src/phasicFlow/MPIParallelization/pointStructure/boundaries/boundaryProcessor.cpp @@ -19,6 +19,7 @@ Licence: -----------------------------------------------------------------------------*/ #include "boundaryProcessor.hpp" +#include "boundaryProcessorKernels.hpp" #include "dictionary.hpp" #include "mpiCommunication.hpp" #include "boundaryBaseKernels.hpp" @@ -69,8 +70,6 @@ pFlow::MPI::boundaryProcessor::beforeIteration(uint32 iterNum, real t, real dt) thisNumPoints_ = size(); - uint32 oldNeighborProcNumPoints = neighborProcNumPoints_; - auto req = MPI_REQUEST_NULL; MPI_Isend( &thisNumPoints_, @@ -103,7 +102,6 @@ pFlow::MPI::boundaryProcessor::beforeIteration(uint32 iterNum, real t, real dt) return false; } - return true; } @@ -154,23 +152,21 @@ bool pFlow::MPI::boundaryProcessor::transferData(uint32 iter, int step) transferFlags.fill(0u); const auto& transferD = transferFlags.deviceViewAll(); - auto points = thisPoints(); + deviceScatteredFieldAccess points = thisPoints(); auto p = boundaryPlane().infPlane(); numToTransfer_ = 0; - Kokkos::parallel_reduce + + Kokkos::parallel_reduce ( "boundaryProcessor::afterIteration", deviceRPolicyStatic(0,s), - LAMBDA_HD(uint32 i, uint32& transferToUpdate) - { - if(p.pointInNegativeSide(points(i))) - { - transferD(i)=1; - transferToUpdate++; - } - }, + boundaryProcessorKernels::markNegative( + boundaryPlane().infPlane(), + transferFlags.deviceViewAll(), + thisPoints() + ), numToTransfer_ ); @@ -206,13 +202,15 @@ bool pFlow::MPI::boundaryProcessor::transferData(uint32 iter, int step) thisBoundaryIndex(), pFlowProcessors().localCommunicator(), &req), true ); - + //pOutput<<"sent "<< numToTransfer_<& f, + const deviceScatteredFieldAccess& p + ) + : + plane_(pl), + flags_(f), + points_(p) + {} + + infinitePlane plane_; + deviceViewType1D flags_; + deviceScatteredFieldAccess points_; + + INLINE_FUNCTION_HD + void operator()(uint32 i, uint32& transferToUpdate)const + { + if(plane_.pointInNegativeSide(points_(i))) + { + flags_(i)=1; + transferToUpdate++; + } + } + +}; + +} \ No newline at end of file diff --git a/src/phasicFlow/MPIParallelization/pointStructure/boundaries/dataSender.hpp b/src/phasicFlow/MPIParallelization/pointStructure/boundaries/dataSender.hpp index 18b907c8..a4c5d39b 100644 --- a/src/phasicFlow/MPIParallelization/pointStructure/boundaries/dataSender.hpp +++ b/src/phasicFlow/MPIParallelization/pointStructure/boundaries/dataSender.hpp @@ -116,7 +116,6 @@ public: buffer_.clear(); buffer_.resize(numToRecieve); - Status status; CheckMPI( Irecv( buffer_.getSpan(),