xfygogo/phasicFlow
1
0
Fork 0
mirror of https://github.com/PhasicFlow/phasicFlow.git synced 2025-06-22 16:28:30 +00:00
Code Issues Packages Projects Releases Wiki Activity

bug remove for GPU run after CPU MPI parallelization

Browse Source 
- specialization of VectorSingle for word
- dummyFile creation to solve write to file in MPI mode
This commit is contained in:
Hamidreza Norouzi
2024-05-18 18:40:25 +03:30
parent 614b2f732e
commit 4e8b921514
69 changed files with 1124 additions and 490 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
src/Geometry/geometry/geometry.cpp
View File

@ -249,7 +249,9 @@ pFlow::geometry::geometry
if( this->numSurfaces() != motionComponentName_.size() )
{
fatalErrorInFunction<<
"Number of surfaces is not equal to number of motion component names"<<endl;
"Number of surfaces ("<< this->numSurfaces() <<
") is not equal to number of motion component names("<<
motionComponentName_.size()<<")"<<endl;
fatalExit;
}

34
src/Geometry/geometryMotion/geometryMotion.cpp
View File

@ -1,4 +1,3 @@
#include "geometryMotion.hpp"
/*------------------------------- phasicFlow ---------------------------------
O C enter of
O O E ngineering and
@ -70,6 +69,28 @@ bool pFlow::geometryMotion<MotionModel>::findMotionIndex()
return true;
}
namespace pFlow::GMotion
{
template<typename ModelInterface>
void moveGeometry(
real dt,
uint32 nPoints,
const ModelInterface& mModel,
const deviceViewType1D<uint32>& pointMIndexD,
const deviceViewType1D<realx3>& pointsD
)
{
Kokkos::parallel_for(
"geometryMotion<MotionModel>::movePoints",
deviceRPolicyStatic(0, nPoints),
LAMBDA_HD(uint32 i){
auto newPos = mModel.transferPoint(pointMIndexD[i], pointsD[i], dt);
pointsD[i] = newPos;
});
Kokkos::fence();
}
}
template<typename MotionModel>
bool pFlow::geometryMotion<MotionModel>::moveGeometry()
@ -84,8 +105,15 @@ template<typename MotionModel>
auto& pointMIndexD= pointMotionIndex_.deviceViewAll();
auto& pointsD = points().deviceViewAll();
pFlow::GMotion::moveGeometry(
dt,
numPoints(),
motionModel_.getModelInterface(iter, t, dt),
pointMotionIndex_.deviceViewAll(),
points().deviceViewAll()
);
Kokkos::parallel_for(
/*Kokkos::parallel_for(
"geometryMotion<MotionModel>::movePoints",
deviceRPolicyStatic(0, numPoints()),
LAMBDA_HD(uint32 i){
@ -93,7 +121,7 @@ template<typename MotionModel>
pointsD[i] = newPos;
});
Kokkos::fence();
Kokkos::fence();*/
// move the motion components
motionModel_.move(iter, t,dt);

4
src/Interaction/CMakeLists.txt
View File

@ -7,8 +7,8 @@ contactSearch/methods/cellBased/NBS/NBS.cpp
contactSearch/methods/cellBased/NBS/cellsWallLevel0.cpp
contactSearch/boundaries/boundaryContactSearch/boundaryContactSearch.cpp
contactSearch/boundaries/twoPartContactSearch/twoPartContactSearchKernels.cpp
contactSearch/boundaries/twoPartContactSearch/twoPartContactSearch.cpp
#contactSearch/boundaries/twoPartContactSearch/twoPartContactSearchKernels.cpp
#contactSearch/boundaries/twoPartContactSearch/twoPartContactSearch.cpp
contactSearch/boundaries/periodicBoundaryContactSearch/ppwBndryContactSearchKernels.cpp
contactSearch/boundaries/periodicBoundaryContactSearch/ppwBndryContactSearch.cpp
contactSearch/boundaries/periodicBoundaryContactSearch/wallBoundaryContactSearch.cpp

4
src/Interaction/contactLists/sortedContactList.hpp
View File

@ -138,7 +138,7 @@ public:
start,
end,
newPair);
idx0!=-1)
idx0!=static_cast<uint32>(-1))
{
values_[idx] = values0_[idx0];
}
@ -147,7 +147,7 @@ public:
start,
end,
newPair);
idx0!=-1)
idx0!= static_cast<uint32>(-1) )
{
values_[idx] = values0_[idx0];

6
src/Interaction/contactLists/unsortedContactList.hpp
View File

@ -124,7 +124,7 @@ public:
INLINE_FUNCTION_HD
bool getValue(const PairType& p, ValueType& val)const
{
if(auto idx = this->find(p); idx!=-1)
if(auto idx = this->find(p); idx!=static_cast<uint32>(-1))
{
val = getValue(idx);
return true;
@ -141,7 +141,7 @@ public:
INLINE_FUNCTION_HD
bool setValue(const PairType& p, const ValueType& val)const
{
if(uint32 idx = this->find(p); idx!=-1)
if(uint32 idx = this->find(p); idx!=static_cast<uint32>(-1))
{
setValue(idx, val);
return true;;
@ -156,7 +156,7 @@ public:
{
if( uint32 idx0 =
container0_.find(this->getPair(idx));
idx0!=-1 )
idx0!= static_cast<uint32>(-1) )
{
values_[idx] = values0_[idx0];
}

6
src/Interaction/contactLists/unsortedPairs.hpp
View File

@ -105,7 +105,7 @@ public:
uint32 insert(idType i, idType j)const
{
if(auto insertResult = container_.insert(PairType(i,j)); insertResult.failed())
return -1;
return static_cast<uint32>(-1);
else
return insertResult.index();
@ -115,7 +115,7 @@ public:
uint32 insert(const PairType& p)const
{
if(auto insertResult = container_.insert(p); insertResult.failed())
return -1;
return static_cast<uint32>(-1);
else
return insertResult.index();
@ -154,7 +154,7 @@ public:
idx != Kokkos::UnorderedMapInvalidIndex )
return idx;
else
return -1;
return static_cast<uint32>(-1);
}
INLINE_FUNCTION_HD

4
src/Interaction/contactSearch/boundaries/periodicBoundaryContactSearch/ppwBndryContactSearchKernels.cpp
View File

@ -80,14 +80,14 @@ pFlow::uint32 pFlow::pweBndryContactSearchKernels::broadSearchPP
if(!searchCells.inCellRange(ind))continue;
uint32 thisI = head(ind.x(),ind.y(),ind.z());
while (thisI!=-1)
while (thisI!=static_cast<uint32>(-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<uint32>(-1))
{
getFullUpdate++;
}

13
src/Interaction/contactSearch/boundaries/periodicBoundaryContactSearch/wallBoundaryContactSearch.cpp
View File

@ -113,7 +113,10 @@ pFlow::uint32 pFlow::wallBoundaryContactSearch::findPairsElementRangeCount
uint32 nNotInserted = 0;
uint32 nThis = pPoints.size();
const auto& numElements = numElements_;
const auto& elementBox = elementBox_;
const auto& validBox = validBox_;
Kokkos::parallel_reduce(
"pFlow::wallBoundaryContactSearch::findPairsElementRangeCount",
deviceRPolicyDynamic(0,nThis),
@ -123,11 +126,11 @@ pFlow::uint32 pFlow::wallBoundaryContactSearch::findPairsElementRangeCount
int32x3 ind;
if( searchCells.pointIndexInDomain(p, ind) )
{
for(uint32 nTri=0; nTri<numElements_; nTri++)
for(uint32 nTri=0; nTri<numElements; nTri++)
{
if( validBox_[nTri]== 0)continue;
if( elementBox_[nTri].isInside(ind)&&
pairs.insert(i,nTri+baseTriIndex) == -1)
if( validBox[nTri]== 0)continue;
if( elementBox[nTri].isInside(ind)&&
pairs.insert(i,nTri+baseTriIndex) == static_cast<uint32>(-1))
{
notInsertedUpdate++;
}

4
src/Interaction/contactSearch/methods/cellBased/NBS/NBSLoop.hpp
View File

@ -43,7 +43,7 @@ while( m != mapperNBS::NoPos)
auto lm = m;
if(lm>ln) Swap(lm,ln);
if( pairs.insert(lm,ln) == -1)
if( pairs.insert(lm,ln) == static_cast<uint32>(-1))
{
getFullUpdate++;
}
@ -86,7 +86,7 @@ while( m != mapperNBS::NoPos)
auto ln = n;
auto lm = m;
if(lm>ln) Swap(lm,ln);
if( pairs.insert(lm,ln) == -1)
if( pairs.insert(lm,ln) == static_cast<uint32>(-1))
{
getFullUpdate++;
}

37
src/Interaction/contactSearch/methods/cellBased/NBS/cellsWallLevel0.cpp
View File

@ -85,21 +85,26 @@ bool pFlow::cellsWallLevel0::broadSearch
bool pFlow::cellsWallLevel0::build(const cells & searchBox)
{
const auto& points = points_;
const auto& vertices = vertices_;
const auto& elementBox = elementBox_;
const auto cellExtent = cellExtent_;
Kokkos::parallel_for(
"pFlow::cellsWallLevel0::build",
deviceRPolicyStatic(0,numElements_),
CLASS_LAMBDA_HD(uint32 i)
LAMBDA_HD(uint32 i)
{
auto v = vertices_[i];
auto p1 = points_[v.x()];
auto p2 = points_[v.y()];
auto p3 = points_[v.z()];
auto v = vertices[i];
auto p1 = points[v.x()];
auto p2 = points[v.y()];
auto p3 = points[v.z()];
realx3 minP;
realx3 maxP;
searchBox.extendBox(p1, p2, p3, cellExtent_, minP, maxP);
elementBox_[i] = iBoxType(searchBox.pointIndex(minP), searchBox.pointIndex(maxP));
searchBox.extendBox(p1, p2, p3, cellExtent, minP, maxP);
elementBox[i] = iBoxType(searchBox.pointIndex(minP), searchBox.pointIndex(maxP));
});
Kokkos::fence();
@ -153,7 +158,12 @@ pFlow::int32 pFlow::cellsWallLevel0::findPairsElementRangeCount
{
uint32 getFull =0;
const auto& elementBox = elementBox_;
const auto& normals = normals_;
const auto& points = points_;
const auto& vertices = vertices_;
const auto cellExtent = cellExtent_;
Kokkos::parallel_reduce(
"pFlow::cellsWallLevel0::findPairsElementRangeCount",
tpPWContactSearch(numElements_, Kokkos::AUTO),
@ -163,10 +173,10 @@ pFlow::int32 pFlow::cellsWallLevel0::findPairsElementRangeCount
const uint32 iTri = teamMember.league_rank();
const auto triBox = elementBox_[iTri];
const auto triBox = elementBox[iTri];
const auto triPlane = infinitePlane(
normals_[iTri],
points_[vertices_[iTri].x()]);
normals[iTri],
points[vertices[iTri].x()]);
uint32 getFull2 = 0;
@ -186,11 +196,12 @@ pFlow::int32 pFlow::cellsWallLevel0::findPairsElementRangeCount
while( n != particleMap.NoPos)
{
// id is wall id the pair is (particle id, wall id)
if( abs(triPlane.pointFromPlane(pPoints[n]))< pDiams[n]*sizeRatio*cellExtent_)
if( abs(triPlane.pointFromPlane(pPoints[n]))< pDiams[n]*sizeRatio*cellExtent)
{
if( pairs.insert(
static_cast<csIdType>(n),
static_cast<csIdType>(iTri) ) == -1 )
static_cast<csIdType>(iTri) ) == static_cast<csIdType>(-1)
)
innerUpdate++;
}
n = particleMap.next(n);

2
src/Interaction/sphereInteraction/sphereInteraction/sphereInteractionKernels.hpp
View File

@ -261,7 +261,7 @@ struct pwInteractionFunctor
int32 propId_i = propId_[i];
int32 wPropId_j = wPropId_[tj];
realx3 FCn, FCt, Mri, Mrj, Mij, Mji;
realx3 FCn, FCt, Mri, Mrj, Mij;
//output<< "before "<<history.overlap_t_<<endl;
// calculates contact force
forceModel_.contactForce(

2
src/MotionModel/entities/stationary/stationary.hpp
View File

@ -62,7 +62,7 @@ public:
INLINE_FUNCTION_HD
realx3 linVelocityPoint(const realx3 &)const
{
return zero3;
return realx3(0);
}
INLINE_FUNCTION_HD

6
src/MotionModel/rotatingAxisMotion/rotatingAxisMotion.hpp
View File

@ -76,8 +76,6 @@ protected:
{
return true;
}
void impl_setTime(uint32 iter, real t, real dt)const;
public:
@ -90,6 +88,7 @@ public:
const dictionary& dict,
repository* owner);
using fileDictionary::write;
bool write(iOstream& os, const IOPattern& iop)const override;
@ -98,6 +97,9 @@ public:
{
return rotatingAxis({0,0,0}, {1,0,0}, 0.0);
}
// TODO: make this method protected
void impl_setTime(uint32 iter, real t, real dt)const;
};
} // pFlow

2
src/MotionModel/stationaryWall/stationaryWall.hpp
View File

@ -84,6 +84,8 @@ public:
repository* owner);
using fileDictionary::write;
bool write(iOstream& os, const IOPattern& iop)const override;
static

6
src/MotionModel/vibratingMotion/vibratingMotion.hpp
View File

@ -82,7 +82,6 @@ protected:
return true;
}
void impl_setTime(uint32 iter, real t, real dt)const;
public:
@ -99,15 +98,20 @@ public:
/// Destructor
~vibratingMotion()override = default;
using fileDictionary::write;
bool write(iOstream& os, const IOPattern& iop)const override;
static
auto noneComponent()
{
return vibrating();
}
// TODO: make this protected
void impl_setTime(uint32 iter, real t, real dt)const;
};
} // pFlow

4
src/Particles/Insertion/collisionCheck/collisionCheck.cpp
View File

@ -60,7 +60,7 @@ pFlow::collisionCheck::checkPoint(const realx3& p, const real d) const
{
uint32 n = head_(i, j, k);
while( n != -1)
while( n != static_cast<uint32>(-1))
{
if( ((position_[n]-p).length() - 0.5*(diameters_[n]+d )) <= 0.0 )
{
@ -85,7 +85,7 @@ pFlow::collisionCheck::mapLastAddedParticle()
"size mismatch of next and position"<<endl;
return false;
}
next_.push_back(-1);
next_.push_back(static_cast<uint32>(-1));
const auto& p = position_[n];
if(!searchBox_.isInside(p))

3
src/Particles/Insertion/insertion/insertion.cpp
View File

@ -74,9 +74,6 @@ pFlow::insertion::pStruct() const
return particles_.pStruct();
}
bool
pFlow::insertion::readInsertionDict()
{

6
src/Particles/Insertion/insertion/insertion.hpp
View File

@ -32,7 +32,9 @@ class pointStructure;
/**
* Base class for particle insertion
*/
class insertion : public fileDictionary
class insertion
:
public fileDictionary
{
private:
@ -118,6 +120,8 @@ public:
/*/// read from iIstream
virtual bool read(iIstream& is) = 0;*/
using fileDictionary::write;
/// write to iOstream
bool write(iOstream& os, const IOPattern& iop)const override ;
};

7
src/Particles/SphereParticles/sphereParticles/sphereParticles.hpp
View File

@ -82,7 +82,7 @@ private:
Timer fieldUpdateTimer_;
private:
bool initializeParticles();
bool getParticlesInfoFromShape(
const wordVector& shapeNames,
@ -119,11 +119,14 @@ public:
*/
/*bool insertParticles
(
const realx3Vector& position,
const realx3Vector& position,
const wordVector& shapes,
const setFieldList& setField
) override ;*/
// TODO: make this method private later
bool initializeParticles();
/// const reference to shapes object
const auto& spheres() const
{

7
src/Particles/particles/regularParticleIdHandler/regularParticleIdHandler.cpp
View File

@ -15,8 +15,11 @@ pFlow::regularParticleIdHandler::regularParticleIdHandler
pFlow::Pair<pFlow::uint32, pFlow::uint32>
pFlow::regularParticleIdHandler::getIdRange(uint32 nNewParticles)
{
if(nNewParticles==0) return {0,0};
uint32 startId;
if(maxId_==-1)
if(maxId_== static_cast<uint32>(-1))
{
startId = 0;
}
@ -37,7 +40,7 @@ bool pFlow::regularParticleIdHandler::initialIdCheck()
uint32 maxId = max( *this );
/// particles should get ids from 0 to size-1
if(maxId == -1)
if(maxId == static_cast<uint32>(-1))
{
fillSequence(*this,0u);
maxId_ = size()-1;

2
src/Particles/particles/regularParticleIdHandler/regularParticleIdHandler.hpp
View File

@ -31,7 +31,7 @@ class regularParticleIdHandler
{
private:
uint32 maxId_ = -1;
uint32 maxId_ = static_cast<uint32>(-1);
bool initialIdCheck()override;
public:

4
src/Particles/particles/shape/baseShapeNames.hpp
View File

@ -132,7 +132,7 @@ public:
return true;
}
}
idx = -1;
idx = static_cast<uint32>(-1);
return false;
}
@ -144,6 +144,8 @@ public:
// - IO
using fileDictionary::write;
bool write(iOstream& os)const override;
};

1
src/phasicFlow/CMakeLists.txt
View File

@ -59,6 +59,7 @@ Timer/Timers.cpp
containers/Vector/Vectors.cpp
containers/VectorHD/wordVectorHost.cpp
containers/VectorHD/VectorSingles.cpp
containers/Field/Fields.cpp
containers/symArrayHD/symArrays.cpp

4
src/phasicFlow/Kokkos/ViewAlgorithms.hpp
View File

@ -421,10 +421,10 @@ binarySearch(
)
{
if (end <= start)
return -1;
return static_cast<uint32>(-1);
if (auto res = binarySearch_(view.data() + start, end - start, val);
res != -1)
res != static_cast<uint32>(-1))
{
return res + start;
}

4
src/phasicFlow/commandLine/CLI/TypeTools.hpp
View File

@ -146,11 +146,11 @@ template <typename T, typename C> class is_direct_constructible {
static auto test(int, std::true_type) -> decltype(
// NVCC warns about narrowing conversions here
#ifdef __CUDACC__
#pragma diag_suppress 2361
#pragma nv_diag_suppress 2361
#endif
TT { std::declval<CC>() }
#ifdef __CUDACC__
#pragma diag_default 2361
#pragma nv_diag_default 2361
#endif
,
std::is_move_assignable<TT>());

1
src/phasicFlow/containers/Field/Field.hpp
View File

@ -24,6 +24,7 @@ Licence:
#include "types.hpp"
#include "VectorSingle.hpp"
#include "wordVectorHost.hpp"
#include "Vector.hpp"
#include "streams.hpp"

3
src/phasicFlow/containers/Field/Fields.cpp
View File

@ -31,5 +31,4 @@ template class pFlow::Field<pFlow::real>;
template class pFlow::Field<pFlow::realx3>;
template class pFlow::Field<pFlow::word, pFlow::HostSpace>;

25
src/phasicFlow/containers/List/ListPtr/ListPtrI.hpp
View File

@ -39,6 +39,7 @@ inline bool pFlow::ListPtr<T>::copy(const ListPtrType& src)
}
template<typename T>
inline
T* pFlow::ListPtr<T>::ptr(size_t i)
{
@ -51,6 +52,7 @@ T* pFlow::ListPtr<T>::ptr(size_t i)
}
template<typename T>
inline
const T* pFlow::ListPtr<T>::ptr
(
size_t i
@ -66,6 +68,7 @@ const T* pFlow::ListPtr<T>::ptr
}
template<typename T>
inline
auto pFlow::ListPtr<T>::pos
(
size_t i
@ -84,6 +87,7 @@ auto pFlow::ListPtr<T>::pos
}
template<typename T>
inline
auto pFlow::ListPtr<T>::pos
(
size_t i
@ -102,6 +106,7 @@ auto pFlow::ListPtr<T>::pos
}
template<typename T>
inline
pFlow::ListPtr<T>::ListPtr
(
const ListPtrType& src
@ -119,6 +124,7 @@ pFlow::ListPtr<T>::ListPtr
template<typename T>
inline
pFlow::ListPtr<T>& pFlow::ListPtr<T>::operator=
(
const ListPtrType& rhs
@ -144,6 +150,7 @@ pFlow::ListPtr<T>& pFlow::ListPtr<T>::operator=
}
template<typename T>
inline
pFlow::uniquePtr<T> pFlow::ListPtr<T>::set
(
size_t i, T* ptr
@ -155,6 +162,7 @@ pFlow::uniquePtr<T> pFlow::ListPtr<T>::set
template<typename T>
inline
pFlow::uniquePtr<T> pFlow::ListPtr<T>::set
(
size_t i,
@ -179,6 +187,7 @@ pFlow::uniquePtr<T> pFlow::ListPtr<T>::set
template<typename T>
template<typename... Args>
inline
pFlow::uniquePtr<T> pFlow::ListPtr<T>::setSafe
(
size_t i,
@ -190,6 +199,7 @@ pFlow::uniquePtr<T> pFlow::ListPtr<T>::setSafe
}
template<typename T>
inline
void pFlow::ListPtr<T>::push_back
(
T* ptr
@ -199,6 +209,7 @@ void pFlow::ListPtr<T>::push_back
}
template<typename T>
inline
void pFlow::ListPtr<T>::push_back(uniquePtr<T>&& ptr)
{
list_.push_back( ptr.release() );
@ -206,13 +217,15 @@ void pFlow::ListPtr<T>::push_back(uniquePtr<T>&& ptr)
template<typename T>
template<typename... Args>
inline
void pFlow::ListPtr<T>::push_backSafe(Args&&... args)
{
auto ptr=makeUnique<T>(std::forward<Args>(args)...) ;
uniquePtr<T> ptr = makeUnique<T>(std::forward<Args>(args)...) ;
push_back(ptr);
}
template<typename T>
inline
T& pFlow::ListPtr<T>::operator[]
(
size_t i
@ -231,6 +244,7 @@ T& pFlow::ListPtr<T>::operator[]
}
template<typename T>
inline
const T& pFlow::ListPtr<T>::operator[]
(
size_t i
@ -248,18 +262,21 @@ const T& pFlow::ListPtr<T>::operator[]
}
template<typename T>
inline
size_t pFlow::ListPtr<T>::size()const
{
return list_.size();
}
template<typename T>
inline
auto pFlow::ListPtr<T>::empty() const
{
return list_.emtpy();
}
template<typename T>
inline
pFlow::uniquePtr<T> pFlow::ListPtr<T>::release
(
size_t i
@ -273,15 +290,14 @@ pFlow::uniquePtr<T> pFlow::ListPtr<T>::release
template<typename T>
inline
void pFlow::ListPtr<T>::clear()
{
int i =0;
for( auto iter = list_.begin(); iter != list_.end(); ++iter )
{
if(*iter != nullptr)
{
{
delete *iter;
*iter = nullptr;
}
@ -291,6 +307,7 @@ void pFlow::ListPtr<T>::clear()
}
template<typename T>
inline
void pFlow::ListPtr<T>::clear
(
size_t i

239
src/phasicFlow/containers/VectorHD/VectorSingle.cpp
View File

@ -55,30 +55,14 @@ void pFlow::VectorSingle<T,MemorySpace>::changeCapacity
bool withInit
)
{
if constexpr( isTriviallyCopyable_ )
if(withInit)
{
if(withInit)
{
resizeInit(view_, actualCap);
}
else
{
resizeNoInit(view_, actualCap);
}
}
else if constexpr( isHostAccessible_ )
{
viewType newView(view_.label(), actualCap);
for(auto i=0u; i<min(size_,actualCap); i++)
{
newView(i) = view_(i);
}
view_ = newView;
resizeInit(view_, actualCap);
}
else
{
static_assert("changeCapacity is not a valid operation for non-trivially-copyable data type on device memory");
resizeNoInit(view_, actualCap);
}
}
@ -91,17 +75,8 @@ pFlow::uint32 pFlow::VectorSingle<T,MemorySpace>::reallocateCapacitySize
uint32 s
)
{
if constexpr (isTriviallyCopyable_)
{
reallocNoInit(view_, cap);
}
else
{
viewType newView(view_.label(), cap);
view_ = newView;
}
return setSize(s);
reallocNoInit(view_, cap);
return setSize(s);
}
template<typename T, typename MemorySpace>
@ -209,21 +184,7 @@ pFlow::VectorSingle<T,MemorySpace>::VectorSingle
:
VectorSingle(name, src.capacity(), src.size(), RESERVE())
{
if constexpr(isTriviallyCopyable_)
{
copy(deviceView(), src.deviceView());
}
else if constexpr( isHostAccessible_)
{
for(auto i=0u; i<src.size(); i++)
{
view_[i] = src.view_[i];
}
}
else
{
static_assert("This constructor is not valid for non-trivially copyable data type on device memory");
}
copy(deviceView(), src.deviceView());
}
template<typename T, typename MemorySpace>
@ -235,21 +196,7 @@ pFlow::VectorSingle<T,MemorySpace>::VectorSingle
:
VectorSingle(name, src.size(), src.size(), RESERVE())
{
if constexpr(isTriviallyCopyable_)
{
copy(deviceView(), src);
}
else if constexpr( isHostAccessible_)
{
for(auto i=0u; i<size(); i++)
{
view_[i] = src[i];
}
}
else
{
static_assert("This constructor is not valid for non-trivially copyable data type on device memory");
}
copy(deviceView(), src);
}
template<typename T, typename MemorySpace>
@ -325,18 +272,7 @@ INLINE_FUNCTION_H
auto pFlow::VectorSingle<T,MemorySpace>::hostViewAll()const
{
auto hView = Kokkos::create_mirror_view(view_);
if constexpr(isTriviallyCopyable_)
{
copy(hView, view_);
}
else if constexpr( isHostAccessible_ )
{
// nothing to be done, since it is already a host memory
}
else
{
static_assert("hostViewAll is not valid for non-trivially copyable data type on device memory");
}
copy(hView, view_);
return hView;
}
@ -345,19 +281,7 @@ INLINE_FUNCTION_H
auto pFlow::VectorSingle<T,MemorySpace>::hostView()const
{
auto hView = Kokkos::create_mirror_view(deviceView());
if constexpr(isTriviallyCopyable_)
{
copy(hView, deviceView());
}
else if constexpr( isHostAccessible_ )
{
// nothing to be done, since it is already a host memory
}
else
{
static_assert("hostView is not valid for non-trivially copyable data type on device memory");
}
copy(hView, deviceView());
return hView;
}
@ -497,23 +421,9 @@ void pFlow::VectorSingle<T,MemorySpace>::assign
changeSize(srcSize);
}
if constexpr( isTriviallyCopyable_ )
{
// - unmanaged view in the host
hostViewType1D<const T> temp(src.data(), srcSize );
copy(deviceView(), temp);
}
else if constexpr( isHostAccessible_)
{
for(auto i=0u; i<srcSize; i++)
{
view_[i] = src[i];
}
}
else
{
static_assert("Not a valid operation for this data type on memory device");
}
// - unmanaged view in the host
hostViewType1D<const T> temp(src.data(), srcSize );
copy(deviceView(), temp);
}
@ -543,21 +453,8 @@ void pFlow::VectorSingle<T,MemorySpace>::assignFromHost(const VectorTypeHost& sr
changeSize(srcSize);
}
if constexpr(isTriviallyCopyable_)
{
copy(deviceView(), src.hostView());
}
else if constexpr( isHostAccessible_)
{
for(auto i=0u; i<src.size(); i++)
{
view_[i] = src.view_[i];
}
}
else
{
static_assert("Not a valid operation for this data type on device memory");
}
copy(deviceView(), src.hostView());
}
template<typename T, typename MemorySpace>
@ -580,25 +477,30 @@ void pFlow::VectorSingle<T,MemorySpace>::assign
changeSize(srcSize);
}
if constexpr(isTriviallyCopyable_)
{
copy(deviceView(), src.deviceView());
}
else if constexpr( isHostAccessible_)
{
for(auto i=0u; i<src.size(); i++)
{
view_[i] = src.view_[i];
}
}
else
{
static_assert("Not a valid operation for this data type on device memory");
}
copy(deviceView(), src.deviceView());
}
template<typename T, typename MemorySpace>
template<typename MSpace>
INLINE_FUNCTION_H
void pFlow::VectorSingle<T,MemorySpace>::assignFromDevice(
const VectorSingle<T, MSpace>& src,
bool srcCapacity
)
{
uint32 srcSize = src.size();
uint32 srcCap = src.capacity();
if(srcCapacity && srcCap != capacity()){
reallocateCapacitySize(srcCap, srcSize);
}
else {
changeSize(srcSize);
}
copy(deviceView(), src.deviceView());
}
template <typename T, typename MemorySpace>
INLINE_FUNCTION_H
void pFlow::VectorSingle<T, MemorySpace>::append(const ViewType1D<T,MemorySpace>& appVec)
@ -615,21 +517,8 @@ void pFlow::VectorSingle<T, MemorySpace>::append(const ViewType1D<T,MemorySpace>
view_,
Kokkos::make_pair<uint32>(oldS, newSize));
if constexpr( isTriviallyCopyable_)
{
copy(appendView, appVec);
}
else if constexpr( isHostAccessible_)
{
for(auto i=0u; i<appVec.size(); i++)
{
appendView[i] = appVec[i];
}
}
else
{
static_assert("not a valid operation for this data type on device memory");
}
copy(appendView, appVec);
}
template <typename T, typename MemorySpace>
@ -650,22 +539,7 @@ void pFlow::VectorSingle<T, MemorySpace>::append
hostViewType1D<const T> temp(appVec.data(), srcSize );
auto dest = Kokkos::subview(view_, Kokkos::make_pair<uint32>(oldSize,newSize));
if constexpr( isTriviallyCopyable_)
{
copy(dest, temp);
}
else if constexpr( isHostAccessible_)
{
for(auto i=0u; i<appVec.size(); i++)
{
dest[i] = appVec[i];
}
}
else
{
static_assert("not a valid operation for this data type on device memory");
}
copy(dest, temp);
}
template <typename T, typename MemorySpace>
@ -687,21 +561,8 @@ void pFlow::VectorSingle<T, MemorySpace>::append
view_,
Kokkos::make_pair<uint32>(oldS, newSize));
if constexpr( isTriviallyCopyable_)
{
copy(appendView, appVec.deviceView());
}
else if constexpr( isHostAccessible_)
{
for(auto i=0u; i<appVec.size(); i++)
{
appendView[i] = appVec.view_[i];
}
}
else
{
static_assert("not a valid operation for this data type on device memory");
}
copy(appendView, appVec.deviceView());
}
template <typename T, typename MemorySpace>
@ -936,22 +797,8 @@ bool pFlow::VectorSingle<T,MemorySpace>::reorderItems(const uint32IndexContainer
setSize(newSize);
if constexpr( isTriviallyCopyable_ )
{
copy(deviceView(), sortedView);
}
else if constexpr( isHostAccessible_)
{
for(auto i=0u; i<newSize; i++)
{
view_[i] = sortedView[i];
}
}
else
{
static_assert("Not a valid operation for this data type on memory device");
}
copy(deviceView(), sortedView);
return true;
}

25
src/phasicFlow/containers/VectorHD/VectorSingle.hpp
View File

@ -39,9 +39,6 @@ Licence:
namespace pFlow
{
//- Forward
template<typename T, typename MemorySpace>
class VectorSingle;
template<typename T, typename MemorySpace=void>
class VectorSingle
@ -101,6 +98,8 @@ private:
static constexpr
bool isTriviallyCopyable_ = std::is_trivially_copyable_v<T>;
static_assert(isTriviallyCopyable_, "This type is not trivially copyable");
/// Evaluate capacity based on the input size
static INLINE_FUNCTION_H uint32 evalCapacity(uint32 n)
{
@ -290,25 +289,7 @@ public:
template<typename MSpace>
INLINE_FUNCTION_H
void assignFromDevice(const VectorSingle<T, MSpace>& src, bool srcCapacity = true)
{
uint32 srcSize = src.size();
uint32 srcCap = src.capacity();
if(srcCapacity && srcCap != capacity()){
reallocateCapacitySize(srcCap, srcSize);
}
else {
changeSize(srcSize);
}
if constexpr(isTriviallyCopyable_){
copy(deviceView(), src.deviceView());
}
else{
static_assert("Not a valid operation for this data type ");
}
}
void assignFromDevice(const VectorSingle<T, MSpace>& src, bool srcCapacity = true);
INLINE_FUNCTION_H
void append(const ViewType1D<T,MemorySpace>& appVec);

3
src/phasicFlow/containers/VectorHD/VectorSingles.hpp
View File

@ -25,6 +25,7 @@ Licence:
#include "types.hpp"
#include "VectorSingle.hpp"
#include "wordVectorHost.hpp"
namespace pFlow
{
@ -77,6 +78,8 @@ typedef VectorSingle<realx3x3> realx3x3Vector_D;
typedef VectorSingle<realx3x3, HostSpace> realx3x3Vector_H;
typedef VectorSingle<word, HostSpace> wordVector_H;
}
#endif

23
src/phasicFlow/containers/VectorHD/wordVectorHost.cpp Normal file
View File

@ -0,0 +1,23 @@
/*------------------------------- 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 "wordVectorHost.hpp"

558
src/phasicFlow/containers/VectorHD/wordVectorHost.hpp Normal file
View File

@ -0,0 +1,558 @@
/*------------------------------- 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.
-----------------------------------------------------------------------------*/
#ifndef __wordVectorHost_hpp__
#define __wordVectorHost_hpp__
#include "VectorSingle.hpp"
#include "Vector.hpp"
namespace pFlow
{
template<>
class VectorSingle<word, HostSpace>
{
public:
//- typedefs for accessing data
using VectorType = VectorSingle<word, HostSpace>;
using VectorTypeHost = VectorSingle<word, HostSpace>;
using iterator = word*;
using const_iterator = const word*;
using reference = word&;
using const_reference = const word&;
using value_type = word;
using pointer = word*;
using const_pointer = const word*;
//- typedefs related to memory management
using viewType = ViewType1D<word, HostSpace>;
using device_type = typename viewType::device_type;
using memory_space = typename viewType::memory_space;
using execution_space = typename viewType::execution_space;
private:
// - Data members
Vector<word> container_;
mutable viewType unManagedView_;
// - protected members and methods
/// Is the memory of this vector accessible from Host
static constexpr
bool isHostAccessible_ = true;
/// Is the memory of this vector accessiple from Divce
static constexpr
bool isDeviceAccessible_ = false;
static constexpr
bool isTriviallyCopyable_ = false;
public:
/// Type info
TypeInfoTemplateNV111("VectorSingle", word , memoerySpaceName());
//// - Constructors
/// Empty vector
VectorSingle() = default;
/// Empty vector with a name (capacity = 2)
explicit VectorSingle(const word& name)
:
container_(name)
{}
/// Vector with name and size n
VectorSingle(const word& name, uint32 n)
:
container_(name, n)
{}
/// Vector with name, size and value
VectorSingle(const word& name, uint32 n, const word& val)
:
container_(name, n, val)
{}
/// Vector with name, size (n) and reserved capacity
VectorSingle(const word& name, uint32 cap, uint32 n, const RESERVE& r )
:
container_(name, cap, n, r)
{}
/// Construct with a name and form std::vector (host memory)
VectorSingle(const word& name, const std::vector<word> & src)
:
container_(name, src)
{}
/// Construct with a name and form std::vector (host memory) and with a desired capacity.
VectorSingle(const word& name, const std::vector<word> & src, uint32 cap)
:
container_(name, src, cap)
{}
/// Copy construct (performs deep copy)
VectorSingle(const VectorSingle& src)=default;
/// Copy construct with a new name (perform deep copy)
VectorSingle(const word& name, const VectorSingle& src)
:
container_(name, src.container_)
{}
/// Copy assignment (perform deep copy from rhs to *this)
VectorSingle& operator = (const VectorSingle& rhs)=default;
/// Move construct
VectorSingle(VectorSingle&&) = default;
/// Move assignment
VectorSingle& operator= (VectorSingle&&) = default;
/// @brief Descructor
/// This may not destroy the underlying memory, sice view is
/// shared_ptr and maybe referenced by another object too
~VectorSingle() = default;
/// Clone as a uniquePtr (perform deep copy)
INLINE_FUNCTION_H
uniquePtr<VectorSingle> clone() const
{
return makeUnique<VectorSingle>(*this);
}
//// - Methods
/// Return *this
INLINE_FUNCTION_H
VectorType& VectorField()
{
return *this;
}
/// Return *this
INLINE_FUNCTION_H
const VectorType& VectorField()const
{
return *this;
}
/// Device view range [0,capcity)
INLINE_FUNCTION_H
auto& deviceViewAll()
{
// return un-managed view
unManagedView_ = viewType(container_.data(), container_.capacity());
return unManagedView_;
}
/// Device view range [0,capcity)
INLINE_FUNCTION_H
const auto& deviceViewAll() const
{
// return un-managed view
unManagedView_ = viewType(const_cast<word*>(container_.data()), container_.capacity());
return unManagedView_;
}
/// Device view range [0, size)
INLINE_FUNCTION_H
auto deviceView()const
{
return viewType(const_cast<word*>(container_.data()), container_.size());
}
/// Return a view accessible on Host in range [0,capacity)
INLINE_FUNCTION_H
auto hostViewAll()const
{
return viewType(const_cast<word*>(container_.data()), container_.capacity());
}
/// Return a view accessible on Host in range [0,size)
INLINE_FUNCTION_H
auto hostView()const
{
return viewType(const_cast<word*>(container_.data()), container_.size());
}
/// Name of the vector
INLINE_FUNCTION_H
word name()const
{
return container_.name();
}
/// Size of the vector
INLINE_FUNCTION_H
uint32 size()const
{
return container_.size();
}
// Capcity of the vector
INLINE_FUNCTION_H
uint32 capacity()const
{
return container_.capacity();
}
/// If vector is empty
INLINE_FUNCTION_H
bool empty()const
{
return container_.size()==0uL;
}
/// Reserve capacity for vector
/// Preserve the content.
INLINE_FUNCTION_H
void reserve(uint32 cap)
{
container_.reserve(cap);
}
/// Reallocate memory to new cap and set size to 0.
/*INLINE_FUNCTION_H
void reallocate(uint32 cap);
/// Reallocate memory to new cap and set size to newSize.
/// Do not preserve the content
INLINE_FUNCTION_H
void reallocate(uint32 cap, uint32 newSize);*/
/// Resize the vector and preserve the content
INLINE_FUNCTION_H
void resize(uint32 n)
{
container_.resize(n);
}
/// Resize the vector and assign the value to it.
INLINE_FUNCTION_H
void resize(uint32 n, const word& val)
{
container_.resize(n, val);
}
/// Clear the vector, but keep the allocated memory unchanged
INLINE_FUNCTION_H
void clear()
{
container_.clear();
}
/// Fill the range [0,size) with val
INLINE_FUNCTION_H
void fill(const word& val)
{
container_.fill(val);
}
INLINE_FUNCTION_H
void fill(rangeU32 r, const word& val)
{
container_.fill(r.start(), r.end(), val);
}
/// Change size of the vector and assign val to vector and
INLINE_FUNCTION_H
void assign(size_t n, const word& val)
{
container_.assign(n, val);
}
/// Assign source vector with specified capacity.
/// The size of *this becomes the size of src.
INLINE_FUNCTION_H
void assign(const std::vector<word>& src, uint32 cap)
{
container_.reserve(cap);
this->assign(src);
}
/// Assign source vector.
/// The size of *this becomes the size of src.
/// The capacity of *this becomes the capacity of src.
INLINE_FUNCTION_H
void assign(const std::vector<word>& src)
{
container_.assign(src.begin(), src.end());
}
/// Assign source vector from host side.
/// The size of *this becomes the size of src.
/// The capacity of *this becomes the capacity of src.
INLINE_FUNCTION_H
void assignFromHost(const VectorTypeHost& src)
{
notImplementedFunction;
fatalExit;
}
INLINE_FUNCTION_H
void assign(const VectorType& src, bool srcCapacity = true)
{
notImplementedFunction;
fatalExit;
}
template<typename MSpace>
INLINE_FUNCTION_H
void assignFromDevice(const VectorSingle<word, MSpace>& src, bool srcCapacity = true)
{
notImplementedFunction;
fatalExit;
}
/*INLINE_FUNCTION_H
void append(const ViewType1D<T,MemorySpace>& appVec);
INLINE_FUNCTION_H
void append(const std::vector<T>& appVec);
INLINE_FUNCTION_H
void append(const VectorType& appVec);*/
INLINE_FUNCTION_H
auto getSpan()
{
return span<word>(container_.data(), container_.size());
}
INLINE_FUNCTION_H
auto getSpan()const
{
return span<word>(const_cast<word*>(container_.data()), container_.size());
}
INLINE_FUNCTION_H
bool insertSetElement(const uint32IndexContainer& indices, const word& val)
{
notImplementedFunction;
return false;
}
INLINE_FUNCTION_H
bool insertSetElement(const uint32IndexContainer& indices, const std::vector<word>& vals)
{
notImplementedFunction;
return false;
}
INLINE_FUNCTION_H
bool insertSetElement(
const uint32IndexContainer& indices,
const ViewType1D<word, memory_space> vals)
{
notImplementedFunction;
return false;
}
INLINE_FUNCTION_H
bool reorderItems(const uint32IndexContainer& indices)
{
notImplementedFunction;
return false;
}
/// @brief push a new element at the end (host call only)
/// resize if necessary and works on host accessible vector.
void push_back(const word& val)
{
container_.push_back(val);
}
INLINE_FUNCTION_H
pointer data(){
return container_.data();
}
INLINE_FUNCTION_H
const_pointer data()const{
return container_.data();
}
/// Return begin iterator. It works when devices is host accessible.
auto
begin(){
return container_.begin();
}
/// Return begin iterator. it works when host is accessible.
const auto
begin()const {
return container_.begin();
}
auto end(){
return container_.end();
}
/// Return end iterator. it works when host is accessible.
const auto end()const{
return container_.end();
}
/// Return reference to element i. it works when host is accessible.
word& operator[](size_t i){
return container_[i];
}
const word& operator[](size_t i)const{
return container_[i];
}
//// - IO operations
/// Read vector from stream
FUNCTION_H
bool read(iIstream& is)
{
return container_.read(is);
}
/// Read vector from stream
FUNCTION_H
bool read(iIstream& is, const IOPattern& iop)
{
return container_.read(is, iop);
}
/// Write the vector to os
FUNCTION_H
bool write(iOstream& os, const IOPattern& iop)const
{
return container_.write(os, iop);
}
FUNCTION_H
bool write(iOstream& os)const
{
return container_.write(os);
}
template<typename HostMask>
FUNCTION_H
bool write(iOstream& os, const IOPattern& iop, const HostMask& mask)const
{
notImplementedFunction;
return false;
}
/// Name of the memory space
static
constexpr const char* memoerySpaceName()
{
return memory_space::name();
}
}; // class wordVectorHost
inline iIstream& operator >> (iIstream & is, VectorSingle<word, HostSpace> & ivec )
{
if( !ivec.read(is) )
{
ioErrorInFile (is.name(), is.lineNumber());
fatalExit;
}
return is;
}
inline iOstream& operator << (iOstream& os, const VectorSingle<word, HostSpace>& ovec )
{
if( !ovec.write(os) )
{
ioErrorInFile(os.name(), os.lineNumber());
fatalExit;
}
return os;
}
} // - pFlow
#endif
/*
INLINE_FUNCTION_H
bool append(const deviceViewType1D<T>& dVec, size_t numElems)
{
if(numElems == 0 )return true;
auto oldSize = size_;
auto newSize = size_ + numElems;
if(this->empty() || newSize > capacity() )
{
resize(newSize);
}
else
{
size_ = size_+numElems;
}
auto dSubView = Kokkos::subview(view_, Kokkos::make_pair(oldSize, newSize));
copy(dSubView, dVec);
return true;
}
INLINE_FUNCTION_H
bool append(const VectorSingle& Vec)
{
return append(Vec.deviceView(), Vec.size());
}*/

29
src/phasicFlow/containers/pointField/boundaryField/boundaryField/boundaryField.hpp
View File

@ -135,18 +135,31 @@ public:
FieldAccessType thisField()const
{
return FieldAccessType(
this->size(),
this->indexList().deviceViewAll(),
internal_.deviceViewAll());
if constexpr(isDeviceAccessible<execution_space>())
return FieldAccessType(
this->size(),
this->indexList().deviceViewAll(),
internal_.deviceViewAll());
else
return FieldAccessType(
this->size(),
this->boundary().indexListHost().deviceViewAll(),
internal_.deviceViewAll());
}
FieldAccessType mirrorField()const
{
return FieldAccessType(
this->mirrorBoundary().size(),
this->mirrorBoundary().indexList().deviceViewAll(),
internal_.deviceViewAll());
if constexpr(isDeviceAccessible<execution_space>())
return FieldAccessType(
this->mirrorBoundary().size(),
this->mirrorBoundary().indexList().deviceViewAll(),
internal_.deviceViewAll());
else
return FieldAccessType(
this->mirrorBoundary().size(),
this->mirrorBoundary().indexListHost().deviceViewAll(),
internal_.deviceViewAll());
}
virtual

4
src/phasicFlow/containers/pointField/boundaryField/boundaryFieldList.hpp
View File

@ -43,7 +43,7 @@ protected:
const boundaryList& boundaries_;
uint32 slaveToMasterUpdateIter_ = -1;
uint32 slaveToMasterUpdateIter_ = static_cast<uint32>(-1);
public:
@ -95,7 +95,7 @@ public:
bool slaveToMasterUpdateRequested()const
{
return slaveToMasterUpdateIter_ != -1;
return slaveToMasterUpdateIter_ != static_cast<uint32>(-1);
}

4
src/phasicFlow/dictionary/dictionary.hpp
View File

@ -302,10 +302,10 @@ public:
//// - IO operations
/// read from stream
virtual bool read(iIstream& is);
bool read(iIstream& is) override;
/// write to stream
virtual bool write(iOstream& os) const;
bool write(iOstream& os) const override;
};

3
src/phasicFlow/dictionary/fileDictionary.hpp
View File

@ -46,6 +46,9 @@ public:
const dictionary& dict,
repository* owner=nullptr);
using dictionary::read;
using dictionary::write;
/// read from stream
bool read(iIstream& is, const IOPattern& iop) override;

4
src/phasicFlow/globals/pFlowMacros.hpp
View File

@ -37,6 +37,10 @@ Licence:
#define CONSUME_PARAM(x) (void)(x);
#if defined(pFlow_Build_Cuda) && !defined(__CUDACC__)
#define __CUDACC__
#endif
#ifdef __CUDACC__
#define INLINE_FUNCTION_HD inline __host__ __device__
#define INLINE_FUNCTION_D inline __device__

9
src/phasicFlow/processors/localProcessors.hpp
View File

@ -145,6 +145,15 @@ public:
}
#endif
static
bool builtForMPI()
{
#ifdef pFlow_Build_MPI
return true;
#else
return false;
#endif
}
};
}

1
src/phasicFlow/processors/processors.cpp
View File

@ -29,7 +29,6 @@ Licence:
#include "processors.hpp"
#include "streams.hpp"
static int numVarsInitialized__ = 0;
#ifdef pFlow_Build_MPI

22
src/phasicFlow/repository/IOobject/IOfileHeader.cpp
View File

@ -42,12 +42,22 @@ pFlow::uniquePtr<pFlow::oFstream> pFlow::IOfileHeader::outStream()const
return osPtr;
}
pFlow::IOfileHeader::IOfileHeader
(
const objectFile& objf
)
:
objectFile(objf)
pFlow::uniquePtr<pFlow::oFstream> pFlow::IOfileHeader::dummyOutStream() const
{
auto osPtr = makeUnique<oFstream>( CWD()+word("dummyFile") , outFileBinary());
if(osPtr && owner())
{
auto outPrecision = owner()->outFilePrecision();
osPtr->precision(static_cast<int>(outPrecision));
}
return osPtr;
}
pFlow::IOfileHeader::IOfileHeader(
const objectFile &objf)
: objectFile(objf)
{}
pFlow::fileSystem pFlow::IOfileHeader::path() const

2
src/phasicFlow/repository/IOobject/IOfileHeader.hpp
View File

@ -58,6 +58,8 @@ protected:
// - ouput file stream
uniquePtr<oFstream> outStream()const;
uniquePtr<oFstream> dummyOutStream()const;
public:
// with owner

29
src/phasicFlow/repository/IOobject/IOobject.cpp
View File

@ -122,18 +122,35 @@ bool pFlow::IOobject::writeObject() const
if(ioPattern().thisCallWrite())
{
if(auto ptrOS = outStream(); ptrOS )
if( ioPattern().thisProcWriteData())
{
return writeObject(ptrOS());
if(auto ptrOS = outStream(); ptrOS )
{
return writeObject(ptrOS());
}
else
{
warningInFunction<<
"error in opening file "<< path() <<endl;
return false;
}
}
else
{
warningInFunction<<
"error in opening file "<< path() <<endl;
return false;
if(auto ptrOS = dummyOutStream(); ptrOS )
{
return writeObject(ptrOS());
}
else
{
warningInFunction<<
"error in opening file "<< path() <<endl;
return false;
}
}
}
return true;

2
src/phasicFlow/repository/Time/baseTimeControl.cpp
View File

@ -68,7 +68,7 @@ pFlow::baseTimeControl::baseTimeControl
pFlow::baseTimeControl::baseTimeControl(int32 start, int32 end, int32 stride, const word &intervalPrefix)
:
isTimeStep_(true),
iRange_(start, end, max(stride,1)),
iRange_(start, end, std::max(stride,1)),
intervalPrefix_(
intervalPrefix.size()==0uL? word("interval"): intervalPrefix+"Interval"
)

12
src/phasicFlow/repository/Time/timeControl.cpp
View File

@ -18,7 +18,7 @@ Licence:
-----------------------------------------------------------------------------*/
#include "math.hpp"
#include "timeControl.hpp"
#include "dictionary.hpp"
@ -137,14 +137,14 @@ pFlow::word pFlow::timeControl::timeName()const
bool pFlow::timeControl::finalTime()const
{
if( currentTime_ >= endTime_ ) return true;
if( abs(currentTime_-endTime_) < 0.5*dt_ )return true;
if( std::abs(currentTime_-endTime_) < 0.5*dt_ )return true;
return false;
}
bool pFlow::timeControl::reachedStopAt()const
{
if( currentTime_ >= stopAt_ ) return true;
if( abs(currentTime_-stopAt_) < 0.5*dt_ )return true;
if( std::abs(currentTime_-stopAt_) < 0.5*dt_ )return true;
return false;
}
@ -154,7 +154,7 @@ void pFlow::timeControl::checkForOutputToFile()
bool save = false;
if(managedExternaly_)
{
if( abs(currentTime_-writeTime_) < 0.5*dt_)
if( std::abs(currentTime_-writeTime_) < 0.5*dt_)
{
save = true;
lastSaved_ = currentTime_;
@ -162,12 +162,12 @@ void pFlow::timeControl::checkForOutputToFile()
}
else
{
if ( abs(currentTime_ - lastSaved_ - saveInterval_) < 0.5 * dt_ )
if ( std::abs(currentTime_ - lastSaved_ - saveInterval_) < 0.5 * dt_ )
{
lastSaved_ = currentTime_;
save = true;
}
else if( abs(currentTime_ - lastSaved_) < min( pow(10.0,-1.0*timePrecision_), 0.5 *dt_) )
else if( std::abs(currentTime_ - lastSaved_) < std::min( pow(10.0,-1.0*timePrecision_), 0.5 *dt_) )
{
lastSaved_ = currentTime_;
save = true;

7
src/phasicFlow/smartPointers/uniquePtr.hpp
View File

@ -37,16 +37,15 @@ namespace pFlow
template<
typename T,
typename Deleter = std::default_delete<T>
typename T
>
class uniquePtr
:
public std::unique_ptr<T, Deleter>
public std::unique_ptr<T>
{
public:
using uniquePtrType = std::unique_ptr<T, Deleter>;
using uniquePtrType = std::unique_ptr<T>;
// using base constructors
using uniquePtrType::unique_ptr;

2
src/phasicFlow/streams/TStream/iTstream.cpp
View File

@ -325,7 +325,7 @@ void pFlow::iTstream::reset()
size_t pFlow::iTstream::tell()
{
notImplementedFunction;
return -1;
return static_cast<size_t>(-1);
}
const pFlow::tokenList& pFlow::iTstream::tokens()const

1
src/phasicFlow/structuredData/boundaries/boundaryBase/scatteredFieldAccess.hpp
View File

@ -124,6 +124,7 @@ public:
return indices_;
}
INLINE_FUNCTION_HD
uint32 index(uint32 i)const
{
return indices_[i];

4
src/phasicFlow/structuredData/boundaries/boundaryReflective/boundaryReflective.cpp
View File

@ -125,6 +125,8 @@ bool pFlow::boundaryReflective::afterIteration
const auto& velocity = time().lookupObject<realx3PointField_D>(velocityName_);
const auto& velocityD = velocity.deviceViewAll();
const auto restitution = restitution_;
Kokkos::parallel_for(
"pFlow::boundaryReflective::velocityChange",
@ -137,7 +139,7 @@ bool pFlow::boundaryReflective::afterIteration
if(vn < 0)
{
realx3 vt = vel - vn*p.normal();
vel = restitution_*(vt - vn*p.normal());
vel = restitution*(vt - vn*p.normal());
}
}
);

2
src/phasicFlow/structuredData/cylinder/cylinder.cpp
View File

@ -18,7 +18,7 @@ Licence:
-----------------------------------------------------------------------------*/
#include "math.hpp"
#include "cylinder.hpp"
#include "zAxis.hpp"
#include "streams.hpp"

23
src/phasicFlow/structuredData/pointStructure/internalPoints/internalPoints.cpp
View File

@ -317,6 +317,9 @@ pFlow::internalPoints::insertPoints(
auto aRange = pFlagsD_.activeRange();
uint32 emptySpots = pFlagsD_.capacity() - pFlagsD_.numActive();
if(emptySpots!= 0) emptySpots--;
message msg;
if( numNew > emptySpots )
@ -348,8 +351,20 @@ pFlow::internalPoints::insertPoints(
// we should fill the scattered empty spots
else
{
notImplementedFunction;
return false;
auto newIndices = pFlagsD_.getEmptyPoints(numNew);
if(numNew != newIndices.size())
{
fatalErrorInFunction<<"not enough empty points in pointFlag"<<
numNew<< " "<<newIndices.size() <<endl;
pOutput<< pFlagsD_.capacity()<<endl;
pOutput<< pFlagsD_.numActive()<<endl;
return false;
}
varList.emplaceBack<uint32IndexContainer>(
msg.addAndName(message::ITEM_INSERT),
newIndices
);
}
const auto& indices = varList.getObject<uint32IndexContainer>(
@ -450,7 +465,7 @@ bool pFlow::internalPoints::insertPointsOnly(
}// we should fill the scattered empty spots
else
{
pOutput<<"numNew to be inserted "<< numNew <<endl;
auto newIndices = pFlagsD_.getEmptyPoints(numNew);
if(numNew != newIndices.size())
{
@ -460,7 +475,7 @@ bool pFlow::internalPoints::insertPointsOnly(
pOutput<< pFlagsD_.numActive()<<endl;
return false;
}
pOutput<<newIndices<<endl;
varList.emplaceBack<uint32IndexContainer>(
msg.addAndName(message::ITEM_INSERT),
newIndices

21
src/phasicFlow/structuredData/pointStructure/internalPoints/pointFlagKernels.hpp
View File

@ -63,12 +63,13 @@ pFlow::ViewType1D<pFlow::uint32, typename pFlow::pointFlag<ExecutionSpace>::memo
ViewType1D<uint32,memory_space> emptyPoints("emptyPoints", numToGet);
auto flags = flags_;
Kokkos::parallel_for(
"getEmptyPoints",
rPolicy(0, cap),
LAMBDA_HD(uint32 i){
indices(i) = flags_[i] == DELETED;
indices(i) = flags[i] == DELETED;
});
Kokkos::fence();
@ -151,7 +152,7 @@ pFlow::uint32 pFlow::pointFlag<ExecutionSpace>::markOutOfBoxDelete
{
minRange = 0;
maxRange = 0;
numDeleted == numActive_;
numDeleted = numActive_;
}
else
{
@ -180,7 +181,8 @@ pFlow::pointFlag<ExecutionSpace>::addInternalPoints
uint32 maxRange;
uint32 numAdded = 0;
const auto& flag = flags_;
Kokkos::parallel_reduce(
"pointFlagKernels::addInternalPoints",
deviceRPolicyStatic(0, points.extent(0)),
@ -191,10 +193,10 @@ pFlow::pointFlag<ExecutionSpace>::addInternalPoints
uint32& addToUpdate)
{
uint32 idx = points(i);
if( flags_[idx] <= DELETED) addToUpdate ++;
if( flag[idx] <= DELETED) addToUpdate ++;
minUpdate = min(minUpdate,idx);
maxUpdate = max(maxUpdate,idx);
flags_[idx] = INTERNAL;
flag[idx] = INTERNAL;
},
Kokkos::Min<uint32>(minRange),
Kokkos::Max<uint32>(maxRange),
@ -247,7 +249,7 @@ bool pFlow::pointFlag<ExecutionSpace>::deletePoints
if(numDeleted >= numActive_)
{
activeRange_ = {0, 0};
numDeleted == numActive_;
numDeleted = numActive_;
}
numActive_ = numActive_ - numDeleted;
@ -290,7 +292,7 @@ bool pFlow::pointFlag<ExecutionSpace>::deletePoints
if(numDeleted >= numActive_)
{
activeRange_ = {0, 0};
numDeleted == numActive_;
numDeleted = numActive_;
}
numActive_ = numActive_ - numDeleted;
@ -309,13 +311,14 @@ bool pFlow::pointFlag<ExecutionSpace>::changeFlags
)
{
auto flg = getBoundaryFlag(boundaryIndex);
const auto& flags = flags_;
Kokkos::parallel_for
(
"pointFlag::changeFlags",
rPolicy(0, changePoints.size()),
LAMBDA_HD(uint32 i)
{
flags_[changePoints(i)] = flg;
flags[changePoints(i)] = flg;
}
);
Kokkos::fence();

4
src/phasicFlow/structuredData/zAxis/array2D.hpp
View File

@ -33,12 +33,12 @@ struct array2D
constexpr size_t nCols()const noexcept
{
return nCols;
return nCol;
}
constexpr size_t nRows()const noexcept
{
return nRows;
return nRow;
}
};

384
src/phasicFlow/triSurface/stlFile.cpp
View File

@ -1,384 +0,0 @@
/*------------------------------- 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 "stlFile.hpp"
#include "iFstream.hpp"
#include "oFstream.hpp"
#include "error.hpp"
namespace pFlow
{
inline bool badInput(iIstream& is, token& tok )
{
return is.bad() || !tok.good();
}
inline bool checkWordToken(iIstream& is, token& tok, const word& check)
{
if( badInput(is, tok) || is.eof() || !tok.isWord() || tok.wordToken() != check ) return false;
return true;
}
inline bool checkNumberToken(iIstream& is, token& tok, real& val)
{
if(badInput(is, tok) || is.eof() || !tok.isNumber() )return false;
val = tok.number();
return true;
}
}
bool pFlow::stlFile::readSolid
(
iIstream& is,
realx3x3Vector & vertecies,
word & name
)
{
token tok;
is>> tok;
if(!checkWordToken(is, tok, "solid")) return false;
// check if there is a name associated with solid
name = "";
int32 nWords =0;
bool reachedFacet = false;
is >> tok;
while (nWords < 20 )
{
if( badInput(is, tok) ) return false;
//if(!tok.isWord()) return false;
nWords++;
if(tok.isWord() && tok.wordToken() != "facet" )
{
name += tok.wordToken();
}
else if( tok.isNumber())
{
auto val = tok.number();
name += real2Word(val);
}
else if( tok.isPunctuation())
{
name += tok.pToken();
}
else if (tok.isWord() && tok.wordToken() == "facet")
{
is.putBack(tok);
reachedFacet = true;
break;
}
else
{
return false;
}
is >> tok;
}
if(!reachedFacet) return false;
vertecies.clear();
while(true )
{
is>>tok;
if( badInput(is,tok) || !tok.isWord() )return false;
word wTok = tok.wordToken();
if( wTok == "endsolid" ) return true; // end of solid
if( wTok != "facet" ) return false;
// read facet
is.putBack(tok);
realx3x3 tri;
if( !readFacet(is, tri) ) return false;
vertecies.push_back(tri);
}
return true;
}
bool pFlow::stlFile::readFacet
(
iIstream& is,
realx3x3& tri
)
{
token tok;
is>>tok;
if( !checkWordToken(is, tok, "facet") ) return false;
is >> tok;
if( !checkWordToken(is, tok , "normal") ) return false;
real val;
for(uint32 i=0; i<3; i++ )
{
is>>tok;
if( !checkNumberToken(is, tok, val))return false;
}
is>> tok;
if( !checkWordToken(is, tok, "outer")) return false;
is>> tok;
if(!checkWordToken(is, tok, "loop")) return false;
realx3 v;
for(uint32 i=0; i<3; i++)
{
is>>tok;
if(!checkWordToken(is, tok, "vertex")) return false;
is>>tok;
if(!checkNumberToken(is, tok, v.x()))return false;
is>>tok;
if(!checkNumberToken(is, tok, v.y()))return false;
is>>tok;
if(!checkNumberToken(is, tok, v.z()))return false;
if( i==0 ) tri.x() = v;
if( i==1 ) tri.y() = v;
if( i==2) tri.z() = v;
}
is>> tok;
if(!checkWordToken(is, tok, "endloop")) return false;
is>> tok;
if(!checkWordToken(is, tok, "endfacet")) return false;
return true;
}
bool pFlow::stlFile::writeFacet
(
iOstream& os,
const realx3x3& tri
)const
{
realx3 n = cross( tri.y() - tri.x(), tri.z()-tri.x());
n.normalize();
os.incrIndent();
indent(os) << "facet" << spaceToken()
<< "normal"<<spaceToken()
<< n.x() << spaceToken()
<< n.y() << spaceToken()
<< n.z() << endl;
os.incrIndent();
indent(os) << "outer loop"<<endl;
os.incrIndent();
indent(os) << "vertex"<< spaceToken()
<< tri.x().x() << spaceToken()
<< tri.x().y() << spaceToken()
<< tri.x().z() << endl;
indent(os) << "vertex"<< spaceToken()
<< tri.y().x() << spaceToken()
<< tri.y().y() << spaceToken()
<< tri.y().z() << endl;
indent(os) << "vertex"<< spaceToken()
<< tri.z().x() << spaceToken()
<< tri.z().y() << spaceToken()
<< tri.z().z() << endl;
os.decrIndent();
indent(os) << "endloop"<<endl;
os.decrIndent();
indent(os)<< "endfacet"<<endl;
os.decrIndent();
return true;
}
bool pFlow::stlFile::writeSolid
(
iOstream& os,
const realx3x3Vector& vertecies,
const word& name
)const
{
os<< "solid"<<spaceToken()<<name<<endl;
for(const auto& tri: vertecies)
{
writeFacet(os, tri);
}
os<< "endsolid"<<endl;
return true;
}
pFlow::stlFile::stlFile( fileSystem file )
:
file_(file)
{
}
pFlow::stlFile::stlFile
(
fileSystem file,
const word& name,
const realx3x3Vector& vertecies
)
:
stlFile(file)
{
addSolid(name, vertecies);
}
pFlow::stlFile::stlFile
(
fileSystem file,
const word& name,
realx3x3Vector&& vertecies
)
:
stlFile(file)
{
addSolid(name, vertecies);
}
void pFlow::stlFile::addSolid
(
const word& name,
const realx3x3Vector& vertecies
)
{
solids_.push_backSafe(vertecies);
solidNames_.push_back(name);
}
void pFlow::stlFile::addSolid
(
const word& name,
realx3x3Vector&& vertecies
)
{
solids_.push_backSafe(std::move(vertecies));
solidNames_.push_back(name);
}
bool pFlow::stlFile::read()
{
solids_.clear();
solidNames_.clear();
// open file
iFstream is(file_);
token tok;
while (true)
{
realx3x3Vector vertecies;
word name;
if(!readSolid(is, vertecies, name))
{
ioErrorInFile(is.name(), is.lineNumber());
return false;
}
addSolid(name, std::move(vertecies));
is >> tok;
if( is.eof() || !tok.good())return true;
is.putBack(tok);
}
return true;
}
bool pFlow::stlFile::write()const
{
oFstream os(file_);
os.precision(8);
for(label i=0; i<size(); i++)
{
writeSolid(os, solids_[i], solidNames_[i]);
}
return true;
}
void pFlow::stlFile::setFile
(
fileSystem file
) const
{
file_ = file;
}
const pFlow::wordList& pFlow::stlFile::names()const
{
return solidNames_;
}
size_t pFlow::stlFile::size()const
{
return solids_.size();
}
const pFlow::realx3x3Vector& pFlow::stlFile::solid
(
label i
)const
{
if(i >= size() )
{
fatalErrorInFunction<<
"requested out of range solid from stlFile "<<
file_<<endl;
fatalExit;
}
return solids_[i];
}
const pFlow::word& pFlow::stlFile::name
(
label i
)const
{
if(i >= size() )
{
fatalErrorInFunction<<
"requested out of range solid name from stlFile "<<
file_<<endl;
fatalExit;
}
return solidNames_[i];
}

119
src/phasicFlow/triSurface/stlFile.hpp
View File

@ -1,119 +0,0 @@
/*------------------------------- 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.
-----------------------------------------------------------------------------*/
#ifndef __stlFile_hpp__
#define __stlFile_hpp__
#include "types.hpp"
#include "Vectors.hpp"
#include "Lists.hpp"
#include "fileSystem.hpp"
namespace pFlow
{
class iIstream;
class stlFile
{
protected:
//// - data memebrs
// - list of verticies of all solids
ListPtr<realx3x3Vector> solids_;
// - list of names of all solids
wordList solidNames_;
// - file name of stl file (used for reading and writing)
mutable fileSystem file_;
// - protected members
bool readSolid(iIstream& is, realx3x3Vector & vertecies, word & name);
bool readFacet(iIstream& is, realx3x3& tri);
bool writeSolid(iOstream& os, const realx3x3Vector& vertecies, const word& name)const;
bool writeFacet(iOstream& os, const realx3x3& tri)const;
public:
//// - Constructors
// - construct with file name
// an empty stlFile
stlFile( fileSystem file );
// - construct with file name and one solid (copy)
stlFile( fileSystem file, const word& name, const realx3x3Vector& vertecies);
// - construct with file name and one solid (copy)
stlFile( fileSystem file, const word& name, realx3x3Vector&& vertecies);
// - copy construct
stlFile(const stlFile&) = default;
// - move construct
stlFile(stlFile&&) = default;
~stlFile() = default;
//// - Methods
// - add a solid at the end of list, with copy operation
void addSolid(const word& name, const realx3x3Vector& vertecies);
// - add a solid at the end of list, with move operation
void addSolid(const word& name, realx3x3Vector&& vertecies);
// - clear current content and read from file
bool read();
// - write the current contnet to file
bool write()const;
// - set stl file path
void setFile(fileSystem file) const;
// - name of solids
const wordList& names()const;
// - number of solids
size_t size()const;
// - vertecies of ith solid
const realx3x3Vector& solid(uint64 i)const;
// - name of ith solid
const word& name(uint64 i)const;
};
} // pFlow
#endif //__stlFile_hpp__

2
src/phasicFlow/triSurface/triangleFunctions.hpp
View File

@ -39,7 +39,7 @@ realx3 normal(const realx3& p1, const realx3& p2, const realx3& p3)
{
auto n = cross(p2-p1, p3-p1);
if( equal(n.length(), 0.0) )
return zero3;
return realx3(0);
else
return normalize(n);
}

191
src/phasicFlow/types/basicTypes/math.hpp
View File

@ -21,15 +21,16 @@ Licence:
#ifndef __math_hpp__
#define __math_hpp__
#include "builtinTypes.hpp"
#include "pFlowMacros.hpp"
#ifdef __CUDACC__
#include "math.h"
#else
#include <cmath>
#endif
#include "builtinTypes.hpp"
#include "pFlowMacros.hpp"
//* * * * * * * * * * * List of functinos * * * * * * * * //
// abs, mod, exp, log, log10, pow, sqrt, cbrt
// sin, cos, tan, asin, acos, atan, atan2
@ -51,22 +52,26 @@ abs(real x)
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD
int64
abs(int64 x)
{
#ifdef __CUDACC__
return ::abs(x);
#else
return std::abs(x);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD int32
abs(int32 x)
{
#ifdef __CUDACC__
return ::abs(x);
#else
return std::abs(x);
}
#endif
}
INLINE_FUNCTION_HD real
mod(real x, real y)
@ -78,6 +83,7 @@ mod(real x, real y)
#endif
}
INLINE_FUNCTION_HD int64
mod(int64 x, int64 y)
{
@ -102,147 +108,188 @@ mod(uint32 x, uint32 y)
return x % y;
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD real
remainder(real x, real y)
{
#ifdef __CUDACC__
return ::remainder(x,y);
#else
return std::remainder(x, y);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_H
INLINE_FUNCTION_HD
real
exp(real x)
{
#ifdef __CUDACC__
return ::exp(x);
#else
return std::exp(x);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD
real
log(real x)
{
#ifdef __CUDACC__
return ::log(x);
#else
return std::log(x);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD
real
log10(real x)
{
#ifdef __CUDACC__
return ::log10(x);
#else
return std::log10(x);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD
real
pow(real x, real y)
{
#ifdef __CUDACC__
return ::pow(x,y);
#else
return std::pow(x, y);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD
real
sqrt(real x)
{
#ifdef __CUDACC__
return ::sqrt(x);
#else
return std::sqrt(x);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD
real
cbrt(real x)
{
#ifdef __CUDACC__
return ::cbrt(x);
#else
return std::cbrt(x);
#endif
}
#endif
#ifndef __CUDACC__
INLINE_FUNCTION_HD
real
sin(real x)
{
#ifdef __CUDACC__
return ::sin(x);
#else
return std::sin(x);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD
real
cos(real x)
{
#ifdef __CUDACC__
return ::cos(x);
#else
return std::cos(x);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD
real
tan(real x)
{
#ifdef __CUDACC__
return ::tan(x);
#else
return std::tan(x);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD
real
asin(real x)
{
#ifdef __CUDACC__
return ::asin(x);
#else
return std::asin(x);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD
real
acos(real x)
{
#ifdef __CUDACC__
return ::acos(x);
#else
return std::acos(x);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD
real
atan(real x)
{
#ifdef __CUDACC__
return ::atan(x);
#else
return std::atan(x);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD real
atan2(real y, real x)
{
#ifdef __CUDACC__
return ::atan2(y,x);
#else
return std::atan2(y, x);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD
real
sinh(real x)
{
#ifdef __CUDACC__
return ::sinh(x);
#else
return std::sinh(x);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD
real
cosh(real x)
{
#ifdef __CUDACC__
return ::cosh(x);
#else
return std::cosh(x);
}
#endif
}
INLINE_FUNCTION_HD real
tanh(real x)
@ -274,14 +321,17 @@ acosh(real x)
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD
real
atanh(real x)
{
#ifdef __CUDACC__
return ::atanh(x);
#else
return std::atanh(x);
}
#endif
}
INLINE_FUNCTION_HD real
min(real x, real y)
@ -293,40 +343,54 @@ min(real x, real y)
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD
int64
min(int32 x, int32 y)
{
#ifdef __CUDACC__
return ::min(x,y);
#else
return std::min(x, y);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD
int64
min(int64 x, int64 y)
{
#ifdef __CUDACC__
return ::min(x,y);
#else
return std::min(x, y);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD
uint64
min(uint64 x, uint64 y)
{
#ifdef __CUDACC__
return ::min(x,y);
#else
return std::min(x, y);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD uint32
min(uint32 x, uint32 y)
{
#ifdef __CUDACC__
return ::min(x,y);
#else
return std::min(x, y);
}
#endif
}
INLINE_FUNCTION_HD real
max(real x, real y)
@ -338,37 +402,48 @@ max(real x, real y)
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD int64
max(int64 x, int64 y)
{
#ifdef __CUDACC__
return ::max(x, y);
#else
return std::max(x, y);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD int32
max(int32 x, int32 y)
{
#ifdef __CUDACC__
return ::max(x, y);
#else
return std::max(x, y);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD uint64
max(uint64 x, uint64 y)
{
#ifdef __CUDACC__
return ::max(x, y);
#else
return std::max(x, y);
}
#endif
}
#ifndef __CUDACC__
INLINE_FUNCTION_HD uint32
max(uint32 x, uint32 y)
{
#ifdef __CUDACC__
return ::max(x, y);
#else
return std::max(x, y);
}
#endif
}
} // pFlow

2
src/setHelpers/setProperty.hpp
View File

@ -23,6 +23,6 @@ Licence:
REPORT(0)<<"\nReading proprties . . . "<<END_REPORT;
auto proprties = property(propertyFile__, Control.caseSetup().path());
auto proprties = pFlow::property(pFlow::propertyFile__, Control.caseSetup().path());
#endif // __setProperty_hpp__