bug fix to work with CUDA run

This commit is contained in:
Hamidreza Norouzi
2024-05-22 09:52:48 +03:30
parent e05bd2c350
commit d2003b35e6
6 changed files with 100 additions and 28 deletions

View File

@ -87,14 +87,14 @@ pFlow::twoPartContactSearchKernels::broadSearchPP(
continue; continue;
uint32 thisI = head(ind.x(), ind.y(), ind.z()); uint32 thisI = head(ind.x(), ind.y(), ind.z());
while (thisI != -1) while (thisI != static_cast<uint32>(-1))
{ {
auto d_n = sizeRatio * diams[thisI]; auto d_n = sizeRatio * diams[thisI];
// first item is for this boundary and second itme, // first item is for this boundary and second itme,
// for mirror // for mirror
if(sphereSphereCheckB(p_m, points[thisI], d_m, d_n)&& if(sphereSphereCheckB(p_m, points[thisI], d_m, d_n)&&
ppPairs.insert(thisI,mrrI) == -1) ppPairs.insert(thisI,mrrI) == static_cast<uint32>(-1))
{ {
getFullUpdate++; getFullUpdate++;
} }
@ -161,14 +161,14 @@ pFlow::twoPartContactSearchKernels::broadSearchPP(
} }
uint32 i1 = head(ind.x(), ind.y(), ind.z()); uint32 i1 = head(ind.x(), ind.y(), ind.z());
while (i1 != -1) while (i1 != static_cast<uint32>(-1))
{ {
auto d_n = sizeRatio * diams1[i1]; auto d_n = sizeRatio * diams1[i1];
// first item is for this boundary and second itme, // first item is for this boundary and second itme,
// for mirror // for mirror
if(sphereSphereCheckB(p_m, points1[i1], d_m, d_n)&& if(sphereSphereCheckB(p_m, points1[i1], d_m, d_n)&&
ppPairs.insert(i1,i2) == -1) ppPairs.insert(i1,i2) == static_cast<uint32>(-1))
{ {
getFullUpdate++; getFullUpdate++;
} }

View File

@ -41,9 +41,12 @@ bool pFlow::MPI::MPISimulationDomain::createBoundaryDicts()
auto& mpiBoundaries = this->subDict("MPIBoundaries"); auto& mpiBoundaries = this->subDict("MPIBoundaries");
real neighborLength = boundaries.getVal<real>("neighborLength"); real neighborLength = boundaries.getVal<real>("neighborLength");
auto boundaryExtntionLengthRatio = auto boundaryExtntionLengthRatio = max(
boundaries.getValOrSet<real>("boundaryExtntionLengthRatio", 0.1); boundaries.getValOrSet<real>("boundaryExtntionLengthRatio", 0.1),
auto updateIntercal = boundaries.getValOrSet<uint32>("updateInterval", 1u); 0.0);
auto updateIntercal = max(
boundaries.getValOrSet<uint32>("updateInterval", 1u),
1u);
auto neighbors = findPlaneNeighbors(); auto neighbors = findPlaneNeighbors();

View File

@ -136,12 +136,28 @@ bool pFlow::MPI::processorBoundaryField<T, MemorySpace>::hearChanges(
message::eventName(message::BNDR_PROCTRANSFER_SEND) message::eventName(message::BNDR_PROCTRANSFER_SEND)
); );
FieldAccessType transferData( if constexpr( isDeviceAccessible<execution_space>())
indices.size(), {
indices.deviceViewAll(), FieldAccessType transferData(
this->internal().deviceViewAll() indices.size(),
); indices.deviceViewAll(),
sender_.sendData(pFlowProcessors(),transferData); 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)) else if(msg.equivalentTo(message::BNDR_PROCTRANSFER_RECIEVE))
{ {

View File

@ -19,6 +19,7 @@ Licence:
-----------------------------------------------------------------------------*/ -----------------------------------------------------------------------------*/
#include "boundaryProcessor.hpp" #include "boundaryProcessor.hpp"
#include "boundaryProcessorKernels.hpp"
#include "dictionary.hpp" #include "dictionary.hpp"
#include "mpiCommunication.hpp" #include "mpiCommunication.hpp"
#include "boundaryBaseKernels.hpp" #include "boundaryBaseKernels.hpp"
@ -69,8 +70,6 @@ pFlow::MPI::boundaryProcessor::beforeIteration(uint32 iterNum, real t, real dt)
thisNumPoints_ = size(); thisNumPoints_ = size();
uint32 oldNeighborProcNumPoints = neighborProcNumPoints_;
auto req = MPI_REQUEST_NULL; auto req = MPI_REQUEST_NULL;
MPI_Isend( MPI_Isend(
&thisNumPoints_, &thisNumPoints_,
@ -103,7 +102,6 @@ pFlow::MPI::boundaryProcessor::beforeIteration(uint32 iterNum, real t, real dt)
return false; return false;
} }
return true; return true;
} }
@ -154,23 +152,21 @@ bool pFlow::MPI::boundaryProcessor::transferData(uint32 iter, int step)
transferFlags.fill(0u); transferFlags.fill(0u);
const auto& transferD = transferFlags.deviceViewAll(); const auto& transferD = transferFlags.deviceViewAll();
auto points = thisPoints(); deviceScatteredFieldAccess<realx3> points = thisPoints();
auto p = boundaryPlane().infPlane(); auto p = boundaryPlane().infPlane();
numToTransfer_ = 0; numToTransfer_ = 0;
Kokkos::parallel_reduce
Kokkos::parallel_reduce
( (
"boundaryProcessor::afterIteration", "boundaryProcessor::afterIteration",
deviceRPolicyStatic(0,s), deviceRPolicyStatic(0,s),
LAMBDA_HD(uint32 i, uint32& transferToUpdate) boundaryProcessorKernels::markNegative(
{ boundaryPlane().infPlane(),
if(p.pointInNegativeSide(points(i))) transferFlags.deviceViewAll(),
{ thisPoints()
transferD(i)=1; ),
transferToUpdate++;
}
},
numToTransfer_ numToTransfer_
); );
@ -206,7 +202,7 @@ bool pFlow::MPI::boundaryProcessor::transferData(uint32 iter, int step)
thisBoundaryIndex(), thisBoundaryIndex(),
pFlowProcessors().localCommunicator(), pFlowProcessors().localCommunicator(),
&req), true ); &req), true );
//pOutput<<"sent "<< numToTransfer_<<endl;
CheckMPI(recv( CheckMPI(recv(
numToRecieve_, numToRecieve_,
neighborProcessorNo(), neighborProcessorNo(),
@ -214,6 +210,8 @@ bool pFlow::MPI::boundaryProcessor::transferData(uint32 iter, int step)
pFlowProcessors().localCommunicator(), pFlowProcessors().localCommunicator(),
StatusesIgnore), true); StatusesIgnore), true);
//pOutput<<"recieved "<<numToRecieve_<<endl;
MPI_Request_free(&req); MPI_Request_free(&req);
return true; return true;
} }

View File

@ -0,0 +1,56 @@
/*------------------------------- phasicFlow ---------------------------------
O C enter of
O O E ngineering and
O O M ultiscale modeling of
OOOOOOO F luid flow
------------------------------------------------------------------------------
Copyright (C): www.cemf.ir
email: hamid.r.norouzi AT gmail.com
------------------------------------------------------------------------------
Licence:
This file is part of phasicFlow code. It is a free software for simulating
granular and multiphase flows. You can redistribute it and/or modify it under
the terms of GNU General Public License v3 or any other later versions.
phasicFlow is distributed to help others in their research in the field of
granular and multiphase flows, but WITHOUT ANY WARRANTY; without even the
implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
-----------------------------------------------------------------------------*/
#include "phasicFlowKokkos.hpp"
#include "infinitePlane.hpp"
#include "scatteredFieldAccess.hpp"
namespace pFlow::boundaryProcessorKernels
{
struct markNegative
{
markNegative(const infinitePlane& pl,
const deviceViewType1D<uint32>& f,
const deviceScatteredFieldAccess<realx3>& p
)
:
plane_(pl),
flags_(f),
points_(p)
{}
infinitePlane plane_;
deviceViewType1D<uint32> flags_;
deviceScatteredFieldAccess<realx3> points_;
INLINE_FUNCTION_HD
void operator()(uint32 i, uint32& transferToUpdate)const
{
if(plane_.pointInNegativeSide(points_(i)))
{
flags_(i)=1;
transferToUpdate++;
}
}
};
}

View File

@ -116,7 +116,6 @@ public:
buffer_.clear(); buffer_.clear();
buffer_.resize(numToRecieve); buffer_.resize(numToRecieve);
Status status;
CheckMPI( CheckMPI(
Irecv( Irecv(
buffer_.getSpan(), buffer_.getSpan(),