diff --git a/solvers/iterateGeometry/iterateGeometry.cpp b/solvers/iterateGeometry/iterateGeometry.cpp index d541a5bf..1279d1c6 100755 --- a/solvers/iterateGeometry/iterateGeometry.cpp +++ b/solvers/iterateGeometry/iterateGeometry.cpp @@ -35,12 +35,10 @@ Licence: #include "commandLine.hpp" //#include "readControlDict.hpp" -using namespace pFlow; - int main( int argc, char* argv[] ) { -commandLine cmds( +pFlow::commandLine cmds( "iterateGeometry", "Performs simulation without particles, only geometry is solved"); @@ -55,8 +53,8 @@ commandLine cmds( // this should be palced in each main -processors::initProcessors(argc, argv); -initialize_pFlowProcessors(); +pFlow::processors::initProcessors(argc, argv); +pFlow::initialize_pFlowProcessors(); #include "initialize_Control.hpp" #include "setProperty.hpp" @@ -71,7 +69,7 @@ initialize_pFlowProcessors(); // this should be palced in each main #include "finalize.hpp" -processors::finalizeProcessors(); +pFlow::processors::finalizeProcessors(); } diff --git a/solvers/iterateSphereParticles/createDEMComponents.hpp b/solvers/iterateSphereParticles/createDEMComponents.hpp index 85cb798a..e9975f44 100755 --- a/solvers/iterateSphereParticles/createDEMComponents.hpp +++ b/solvers/iterateSphereParticles/createDEMComponents.hpp @@ -20,6 +20,6 @@ Licence: // REPORT(0)<<"\nReading sphere particles . . ."<numSurfaces() != motionComponentName_.size() ) { fatalErrorInFunction<< - "Number of surfaces is not equal to number of motion component names"<numSurfaces() << + ") is not equal to number of motion component names("<< + motionComponentName_.size()<<")"<::findMotionIndex() return true; } +namespace pFlow::GMotion +{ + template + void moveGeometry( + real dt, + uint32 nPoints, + const ModelInterface& mModel, + const deviceViewType1D& pointMIndexD, + const deviceViewType1D& pointsD + ) + { + Kokkos::parallel_for( + "geometryMotion::movePoints", + deviceRPolicyStatic(0, nPoints), + LAMBDA_HD(uint32 i){ + auto newPos = mModel.transferPoint(pointMIndexD[i], pointsD[i], dt); + pointsD[i] = newPos; + }); + + Kokkos::fence(); + } +} template bool pFlow::geometryMotion::moveGeometry() @@ -84,8 +105,15 @@ template 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::movePoints", deviceRPolicyStatic(0, numPoints()), LAMBDA_HD(uint32 i){ @@ -93,7 +121,7 @@ template pointsD[i] = newPos; }); - Kokkos::fence(); + Kokkos::fence();*/ // move the motion components motionModel_.move(iter, t,dt); diff --git a/src/Interaction/CMakeLists.txt b/src/Interaction/CMakeLists.txt index c0265808..96197f6a 100644 --- a/src/Interaction/CMakeLists.txt +++ b/src/Interaction/CMakeLists.txt @@ -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 diff --git a/src/Interaction/contactLists/sortedContactList.hpp b/src/Interaction/contactLists/sortedContactList.hpp index 73f0ce4b..4f41caac 100644 --- a/src/Interaction/contactLists/sortedContactList.hpp +++ b/src/Interaction/contactLists/sortedContactList.hpp @@ -138,7 +138,7 @@ public: start, end, newPair); - idx0!=-1) + idx0!=static_cast(-1)) { values_[idx] = values0_[idx0]; } @@ -147,7 +147,7 @@ public: start, end, newPair); - idx0!=-1) + idx0!= static_cast(-1) ) { values_[idx] = values0_[idx0]; diff --git a/src/Interaction/contactLists/unsortedContactList.hpp b/src/Interaction/contactLists/unsortedContactList.hpp index 8a61b2ce..7aa98237 100644 --- a/src/Interaction/contactLists/unsortedContactList.hpp +++ b/src/Interaction/contactLists/unsortedContactList.hpp @@ -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(-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(-1)) { setValue(idx, val); return true;; @@ -156,7 +156,7 @@ public: { if( uint32 idx0 = container0_.find(this->getPair(idx)); - idx0!=-1 ) + idx0!= static_cast(-1) ) { values_[idx] = values0_[idx0]; } diff --git a/src/Interaction/contactLists/unsortedPairs.hpp b/src/Interaction/contactLists/unsortedPairs.hpp index fd4f117d..fb622bdb 100644 --- a/src/Interaction/contactLists/unsortedPairs.hpp +++ b/src/Interaction/contactLists/unsortedPairs.hpp @@ -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(-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(-1); else return insertResult.index(); @@ -154,7 +154,7 @@ public: idx != Kokkos::UnorderedMapInvalidIndex ) return idx; else - return -1; + return static_cast(-1); } INLINE_FUNCTION_HD diff --git a/src/Interaction/contactSearch/boundaries/periodicBoundaryContactSearch/ppwBndryContactSearchKernels.cpp b/src/Interaction/contactSearch/boundaries/periodicBoundaryContactSearch/ppwBndryContactSearchKernels.cpp index ea266e78..a0d47e78 100644 --- a/src/Interaction/contactSearch/boundaries/periodicBoundaryContactSearch/ppwBndryContactSearchKernels.cpp +++ b/src/Interaction/contactSearch/boundaries/periodicBoundaryContactSearch/ppwBndryContactSearchKernels.cpp @@ -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(-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++; } diff --git a/src/Interaction/contactSearch/boundaries/periodicBoundaryContactSearch/wallBoundaryContactSearch.cpp b/src/Interaction/contactSearch/boundaries/periodicBoundaryContactSearch/wallBoundaryContactSearch.cpp index 99bbb8ac..acfad226 100644 --- a/src/Interaction/contactSearch/boundaries/periodicBoundaryContactSearch/wallBoundaryContactSearch.cpp +++ b/src/Interaction/contactSearch/boundaries/periodicBoundaryContactSearch/wallBoundaryContactSearch.cpp @@ -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(-1)) { notInsertedUpdate++; } diff --git a/src/Interaction/contactSearch/methods/cellBased/NBS/NBSLoop.hpp b/src/Interaction/contactSearch/methods/cellBased/NBS/NBSLoop.hpp index ab3eeea7..ecbc63e5 100644 --- a/src/Interaction/contactSearch/methods/cellBased/NBS/NBSLoop.hpp +++ b/src/Interaction/contactSearch/methods/cellBased/NBS/NBSLoop.hpp @@ -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(-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(-1)) { getFullUpdate++; } diff --git a/src/Interaction/contactSearch/methods/cellBased/NBS/cellsWallLevel0.cpp b/src/Interaction/contactSearch/methods/cellBased/NBS/cellsWallLevel0.cpp index 00932281..e1b6e3a7 100644 --- a/src/Interaction/contactSearch/methods/cellBased/NBS/cellsWallLevel0.cpp +++ b/src/Interaction/contactSearch/methods/cellBased/NBS/cellsWallLevel0.cpp @@ -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(n), - static_cast(iTri) ) == -1 ) + static_cast(iTri) ) == static_cast(-1) + ) innerUpdate++; } n = particleMap.next(n); diff --git a/src/Interaction/sphereInteraction/sphereInteraction/sphereInteractionKernels.hpp b/src/Interaction/sphereInteraction/sphereInteraction/sphereInteractionKernels.hpp index c581d598..751c0fec 100644 --- a/src/Interaction/sphereInteraction/sphereInteraction/sphereInteractionKernels.hpp +++ b/src/Interaction/sphereInteraction/sphereInteraction/sphereInteractionKernels.hpp @@ -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 "<(-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"<(-1)); const auto& p = position_[n]; if(!searchBox_.isInside(p)) diff --git a/src/Particles/Insertion/insertion/insertion.cpp b/src/Particles/Insertion/insertion/insertion.cpp index d8d0881d..bb551a47 100644 --- a/src/Particles/Insertion/insertion/insertion.cpp +++ b/src/Particles/Insertion/insertion/insertion.cpp @@ -74,9 +74,6 @@ pFlow::insertion::pStruct() const return particles_.pStruct(); } - - - bool pFlow::insertion::readInsertionDict() { diff --git a/src/Particles/Insertion/insertion/insertion.hpp b/src/Particles/Insertion/insertion/insertion.hpp index f8419aa3..f39f5c35 100644 --- a/src/Particles/Insertion/insertion/insertion.hpp +++ b/src/Particles/Insertion/insertion/insertion.hpp @@ -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 ; }; diff --git a/src/Particles/SphereParticles/sphereParticles/sphereParticles.hpp b/src/Particles/SphereParticles/sphereParticles/sphereParticles.hpp index 0f78e498..d71e751a 100644 --- a/src/Particles/SphereParticles/sphereParticles/sphereParticles.hpp +++ b/src/Particles/SphereParticles/sphereParticles/sphereParticles.hpp @@ -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 { diff --git a/src/Particles/particles/regularParticleIdHandler/regularParticleIdHandler.cpp b/src/Particles/particles/regularParticleIdHandler/regularParticleIdHandler.cpp index 1916f430..b8a55aba 100644 --- a/src/Particles/particles/regularParticleIdHandler/regularParticleIdHandler.cpp +++ b/src/Particles/particles/regularParticleIdHandler/regularParticleIdHandler.cpp @@ -15,8 +15,11 @@ pFlow::regularParticleIdHandler::regularParticleIdHandler pFlow::Pair pFlow::regularParticleIdHandler::getIdRange(uint32 nNewParticles) { + + if(nNewParticles==0) return {0,0}; + uint32 startId; - if(maxId_==-1) + if(maxId_== static_cast(-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(-1)) { fillSequence(*this,0u); maxId_ = size()-1; diff --git a/src/Particles/particles/regularParticleIdHandler/regularParticleIdHandler.hpp b/src/Particles/particles/regularParticleIdHandler/regularParticleIdHandler.hpp index a237e0da..3e4f8b41 100644 --- a/src/Particles/particles/regularParticleIdHandler/regularParticleIdHandler.hpp +++ b/src/Particles/particles/regularParticleIdHandler/regularParticleIdHandler.hpp @@ -31,7 +31,7 @@ class regularParticleIdHandler { private: - uint32 maxId_ = -1; + uint32 maxId_ = static_cast(-1); bool initialIdCheck()override; public: diff --git a/src/Particles/particles/shape/baseShapeNames.hpp b/src/Particles/particles/shape/baseShapeNames.hpp index bc513b06..7e0a00c3 100644 --- a/src/Particles/particles/shape/baseShapeNames.hpp +++ b/src/Particles/particles/shape/baseShapeNames.hpp @@ -132,7 +132,7 @@ public: return true; } } - idx = -1; + idx = static_cast(-1); return false; } @@ -144,6 +144,8 @@ public: // - IO + using fileDictionary::write; + bool write(iOstream& os)const override; }; diff --git a/src/phasicFlow/CMakeLists.txt b/src/phasicFlow/CMakeLists.txt index 408fc042..7de9ceda 100644 --- a/src/phasicFlow/CMakeLists.txt +++ b/src/phasicFlow/CMakeLists.txt @@ -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 diff --git a/src/phasicFlow/Kokkos/ViewAlgorithms.hpp b/src/phasicFlow/Kokkos/ViewAlgorithms.hpp index db071f0b..28e32abd 100644 --- a/src/phasicFlow/Kokkos/ViewAlgorithms.hpp +++ b/src/phasicFlow/Kokkos/ViewAlgorithms.hpp @@ -421,10 +421,10 @@ binarySearch( ) { if (end <= start) - return -1; + return static_cast(-1); if (auto res = binarySearch_(view.data() + start, end - start, val); - res != -1) + res != static_cast(-1)) { return res + start; } diff --git a/src/phasicFlow/commandLine/CLI/TypeTools.hpp b/src/phasicFlow/commandLine/CLI/TypeTools.hpp index 2b87ec60..667d3eac 100644 --- a/src/phasicFlow/commandLine/CLI/TypeTools.hpp +++ b/src/phasicFlow/commandLine/CLI/TypeTools.hpp @@ -146,11 +146,11 @@ template 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() } #ifdef __CUDACC__ -#pragma diag_default 2361 +#pragma nv_diag_default 2361 #endif , std::is_move_assignable()); diff --git a/src/phasicFlow/containers/Field/Field.hpp b/src/phasicFlow/containers/Field/Field.hpp index 887b994d..738811bc 100644 --- a/src/phasicFlow/containers/Field/Field.hpp +++ b/src/phasicFlow/containers/Field/Field.hpp @@ -24,6 +24,7 @@ Licence: #include "types.hpp" #include "VectorSingle.hpp" +#include "wordVectorHost.hpp" #include "Vector.hpp" #include "streams.hpp" diff --git a/src/phasicFlow/containers/Field/Fields.cpp b/src/phasicFlow/containers/Field/Fields.cpp index b0cbad75..e76daa11 100644 --- a/src/phasicFlow/containers/Field/Fields.cpp +++ b/src/phasicFlow/containers/Field/Fields.cpp @@ -31,5 +31,4 @@ template class pFlow::Field; template class pFlow::Field; - - +template class pFlow::Field; diff --git a/src/phasicFlow/containers/List/ListPtr/ListPtrI.hpp b/src/phasicFlow/containers/List/ListPtr/ListPtrI.hpp index 77ef5a99..861239de 100644 --- a/src/phasicFlow/containers/List/ListPtr/ListPtrI.hpp +++ b/src/phasicFlow/containers/List/ListPtr/ListPtrI.hpp @@ -39,6 +39,7 @@ inline bool pFlow::ListPtr::copy(const ListPtrType& src) } template +inline T* pFlow::ListPtr::ptr(size_t i) { @@ -51,6 +52,7 @@ T* pFlow::ListPtr::ptr(size_t i) } template +inline const T* pFlow::ListPtr::ptr ( size_t i @@ -66,6 +68,7 @@ const T* pFlow::ListPtr::ptr } template +inline auto pFlow::ListPtr::pos ( size_t i @@ -84,6 +87,7 @@ auto pFlow::ListPtr::pos } template +inline auto pFlow::ListPtr::pos ( size_t i @@ -102,6 +106,7 @@ auto pFlow::ListPtr::pos } template +inline pFlow::ListPtr::ListPtr ( const ListPtrType& src @@ -119,6 +124,7 @@ pFlow::ListPtr::ListPtr template +inline pFlow::ListPtr& pFlow::ListPtr::operator= ( const ListPtrType& rhs @@ -144,6 +150,7 @@ pFlow::ListPtr& pFlow::ListPtr::operator= } template +inline pFlow::uniquePtr pFlow::ListPtr::set ( size_t i, T* ptr @@ -155,6 +162,7 @@ pFlow::uniquePtr pFlow::ListPtr::set template +inline pFlow::uniquePtr pFlow::ListPtr::set ( size_t i, @@ -179,6 +187,7 @@ pFlow::uniquePtr pFlow::ListPtr::set template template +inline pFlow::uniquePtr pFlow::ListPtr::setSafe ( size_t i, @@ -190,6 +199,7 @@ pFlow::uniquePtr pFlow::ListPtr::setSafe } template +inline void pFlow::ListPtr::push_back ( T* ptr @@ -199,6 +209,7 @@ void pFlow::ListPtr::push_back } template +inline void pFlow::ListPtr::push_back(uniquePtr&& ptr) { list_.push_back( ptr.release() ); @@ -206,13 +217,15 @@ void pFlow::ListPtr::push_back(uniquePtr&& ptr) template template +inline void pFlow::ListPtr::push_backSafe(Args&&... args) { - auto ptr=makeUnique(std::forward(args)...) ; + uniquePtr ptr = makeUnique(std::forward(args)...) ; push_back(ptr); } template +inline T& pFlow::ListPtr::operator[] ( size_t i @@ -231,6 +244,7 @@ T& pFlow::ListPtr::operator[] } template +inline const T& pFlow::ListPtr::operator[] ( size_t i @@ -248,18 +262,21 @@ const T& pFlow::ListPtr::operator[] } template +inline size_t pFlow::ListPtr::size()const { return list_.size(); } template +inline auto pFlow::ListPtr::empty() const { return list_.emtpy(); } template +inline pFlow::uniquePtr pFlow::ListPtr::release ( size_t i @@ -273,15 +290,14 @@ pFlow::uniquePtr pFlow::ListPtr::release template +inline void pFlow::ListPtr::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::clear() } template +inline void pFlow::ListPtr::clear ( size_t i diff --git a/src/phasicFlow/containers/VectorHD/VectorSingle.cpp b/src/phasicFlow/containers/VectorHD/VectorSingle.cpp index dd093270..0e6f02b2 100644 --- a/src/phasicFlow/containers/VectorHD/VectorSingle.cpp +++ b/src/phasicFlow/containers/VectorHD/VectorSingle.cpp @@ -55,30 +55,14 @@ void pFlow::VectorSingle::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::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 @@ -209,21 +184,7 @@ pFlow::VectorSingle::VectorSingle : VectorSingle(name, src.capacity(), src.size(), RESERVE()) { - if constexpr(isTriviallyCopyable_) - { - copy(deviceView(), src.deviceView()); - } - else if constexpr( isHostAccessible_) - { - for(auto i=0u; i @@ -235,21 +196,7 @@ pFlow::VectorSingle::VectorSingle : VectorSingle(name, src.size(), src.size(), RESERVE()) { - if constexpr(isTriviallyCopyable_) - { - copy(deviceView(), src); - } - else if constexpr( isHostAccessible_) - { - for(auto i=0u; i @@ -325,18 +272,7 @@ INLINE_FUNCTION_H auto pFlow::VectorSingle::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::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::assign changeSize(srcSize); } - if constexpr( isTriviallyCopyable_ ) - { - // - unmanaged view in the host - hostViewType1D temp(src.data(), srcSize ); - copy(deviceView(), temp); - } - else if constexpr( isHostAccessible_) - { - for(auto i=0u; i temp(src.data(), srcSize ); + copy(deviceView(), temp); } @@ -543,21 +453,8 @@ void pFlow::VectorSingle::assignFromHost(const VectorTypeHost& sr changeSize(srcSize); } - if constexpr(isTriviallyCopyable_) - { - copy(deviceView(), src.hostView()); - } - else if constexpr( isHostAccessible_) - { - for(auto i=0u; i @@ -580,25 +477,30 @@ void pFlow::VectorSingle::assign changeSize(srcSize); } - - if constexpr(isTriviallyCopyable_) - { - copy(deviceView(), src.deviceView()); - } - else if constexpr( isHostAccessible_) - { - for(auto i=0u; i +template +INLINE_FUNCTION_H +void pFlow::VectorSingle::assignFromDevice( + const VectorSingle& 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 INLINE_FUNCTION_H void pFlow::VectorSingle::append(const ViewType1D& appVec) @@ -615,21 +517,8 @@ void pFlow::VectorSingle::append(const ViewType1D view_, Kokkos::make_pair(oldS, newSize)); - if constexpr( isTriviallyCopyable_) - { - copy(appendView, appVec); - } - else if constexpr( isHostAccessible_) - { - for(auto i=0u; i @@ -650,22 +539,7 @@ void pFlow::VectorSingle::append hostViewType1D temp(appVec.data(), srcSize ); auto dest = Kokkos::subview(view_, Kokkos::make_pair(oldSize,newSize)); - if constexpr( isTriviallyCopyable_) - { - copy(dest, temp); - } - else if constexpr( isHostAccessible_) - { - for(auto i=0u; i @@ -687,21 +561,8 @@ void pFlow::VectorSingle::append view_, Kokkos::make_pair(oldS, newSize)); - if constexpr( isTriviallyCopyable_) - { - copy(appendView, appVec.deviceView()); - } - else if constexpr( isHostAccessible_) - { - for(auto i=0u; i @@ -936,22 +797,8 @@ bool pFlow::VectorSingle::reorderItems(const uint32IndexContainer setSize(newSize); - - if constexpr( isTriviallyCopyable_ ) - { - copy(deviceView(), sortedView); - } - else if constexpr( isHostAccessible_) - { - for(auto i=0u; i -class VectorSingle; template class VectorSingle @@ -101,6 +98,8 @@ private: static constexpr bool isTriviallyCopyable_ = std::is_trivially_copyable_v; + 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 INLINE_FUNCTION_H - void assignFromDevice(const VectorSingle& 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& src, bool srcCapacity = true); INLINE_FUNCTION_H void append(const ViewType1D& appVec); diff --git a/src/phasicFlow/containers/VectorHD/VectorSingles.hpp b/src/phasicFlow/containers/VectorHD/VectorSingles.hpp index d0c3acdf..2479715a 100644 --- a/src/phasicFlow/containers/VectorHD/VectorSingles.hpp +++ b/src/phasicFlow/containers/VectorHD/VectorSingles.hpp @@ -25,6 +25,7 @@ Licence: #include "types.hpp" #include "VectorSingle.hpp" +#include "wordVectorHost.hpp" namespace pFlow { @@ -77,6 +78,8 @@ typedef VectorSingle realx3x3Vector_D; typedef VectorSingle realx3x3Vector_H; +typedef VectorSingle wordVector_H; + } #endif \ No newline at end of file diff --git a/src/phasicFlow/containers/VectorHD/wordVectorHost.cpp b/src/phasicFlow/containers/VectorHD/wordVectorHost.cpp new file mode 100644 index 00000000..1941a239 --- /dev/null +++ b/src/phasicFlow/containers/VectorHD/wordVectorHost.cpp @@ -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" + diff --git a/src/phasicFlow/containers/VectorHD/wordVectorHost.hpp b/src/phasicFlow/containers/VectorHD/wordVectorHost.hpp new file mode 100644 index 00000000..e638bd27 --- /dev/null +++ b/src/phasicFlow/containers/VectorHD/wordVectorHost.hpp @@ -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 +{ +public: + + //- typedefs for accessing data + + using VectorType = VectorSingle; + + using VectorTypeHost = VectorSingle; + + 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; + + 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 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 & 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 & 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 clone() const + { + return makeUnique(*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(container_.data()), container_.capacity()); + return unManagedView_; + } + + /// Device view range [0, size) + INLINE_FUNCTION_H + auto deviceView()const + { + return viewType(const_cast(container_.data()), container_.size()); + } + + /// Return a view accessible on Host in range [0,capacity) + INLINE_FUNCTION_H + auto hostViewAll()const + { + return viewType(const_cast(container_.data()), container_.capacity()); + } + + + /// Return a view accessible on Host in range [0,size) + INLINE_FUNCTION_H + auto hostView()const + { + return viewType(const_cast(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& 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& 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 + INLINE_FUNCTION_H + void assignFromDevice(const VectorSingle& src, bool srcCapacity = true) + { + notImplementedFunction; + fatalExit; + } + + /*INLINE_FUNCTION_H + void append(const ViewType1D& appVec); + + INLINE_FUNCTION_H + void append(const std::vector& appVec); + + INLINE_FUNCTION_H + void append(const VectorType& appVec);*/ + + INLINE_FUNCTION_H + auto getSpan() + { + return span(container_.data(), container_.size()); + } + + INLINE_FUNCTION_H + auto getSpan()const + { + return span(const_cast(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& vals) + { + notImplementedFunction; + return false; + } + + INLINE_FUNCTION_H + bool insertSetElement( + const uint32IndexContainer& indices, + const ViewType1D 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 + 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 & ivec ) +{ + if( !ivec.read(is) ) + { + ioErrorInFile (is.name(), is.lineNumber()); + fatalExit; + } + return is; +} + +inline iOstream& operator << (iOstream& os, const VectorSingle& ovec ) +{ + + if( !ovec.write(os) ) + { + ioErrorInFile(os.name(), os.lineNumber()); + fatalExit; + } + + return os; +} + + +} // - pFlow + + + +#endif + + +/* +INLINE_FUNCTION_H + bool append(const deviceViewType1D& 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()); + }*/ \ No newline at end of file diff --git a/src/phasicFlow/containers/pointField/boundaryField/boundaryField/boundaryField.hpp b/src/phasicFlow/containers/pointField/boundaryField/boundaryField/boundaryField.hpp index 9570f86b..20267f33 100644 --- a/src/phasicFlow/containers/pointField/boundaryField/boundaryField/boundaryField.hpp +++ b/src/phasicFlow/containers/pointField/boundaryField/boundaryField/boundaryField.hpp @@ -135,18 +135,31 @@ public: FieldAccessType thisField()const { - return FieldAccessType( - this->size(), - this->indexList().deviceViewAll(), - internal_.deviceViewAll()); + if constexpr(isDeviceAccessible()) + 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()) + return FieldAccessType( + this->mirrorBoundary().size(), + this->mirrorBoundary().indexList().deviceViewAll(), + internal_.deviceViewAll()); + else + return FieldAccessType( + this->mirrorBoundary().size(), + this->mirrorBoundary().indexListHost().deviceViewAll(), + internal_.deviceViewAll()); + } virtual diff --git a/src/phasicFlow/containers/pointField/boundaryField/boundaryFieldList.hpp b/src/phasicFlow/containers/pointField/boundaryField/boundaryFieldList.hpp index 31a13661..9ef6fcb1 100644 --- a/src/phasicFlow/containers/pointField/boundaryField/boundaryFieldList.hpp +++ b/src/phasicFlow/containers/pointField/boundaryField/boundaryFieldList.hpp @@ -43,7 +43,7 @@ protected: const boundaryList& boundaries_; - uint32 slaveToMasterUpdateIter_ = -1; + uint32 slaveToMasterUpdateIter_ = static_cast(-1); public: @@ -95,7 +95,7 @@ public: bool slaveToMasterUpdateRequested()const { - return slaveToMasterUpdateIter_ != -1; + return slaveToMasterUpdateIter_ != static_cast(-1); } diff --git a/src/phasicFlow/dictionary/dictionary.hpp b/src/phasicFlow/dictionary/dictionary.hpp index 9baba109..1684987c 100644 --- a/src/phasicFlow/dictionary/dictionary.hpp +++ b/src/phasicFlow/dictionary/dictionary.hpp @@ -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; }; diff --git a/src/phasicFlow/dictionary/fileDictionary.hpp b/src/phasicFlow/dictionary/fileDictionary.hpp index b8685dfd..a74c989e 100644 --- a/src/phasicFlow/dictionary/fileDictionary.hpp +++ b/src/phasicFlow/dictionary/fileDictionary.hpp @@ -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; diff --git a/src/phasicFlow/globals/pFlowMacros.hpp b/src/phasicFlow/globals/pFlowMacros.hpp index 544fc150..8973b6ba 100755 --- a/src/phasicFlow/globals/pFlowMacros.hpp +++ b/src/phasicFlow/globals/pFlowMacros.hpp @@ -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__ diff --git a/src/phasicFlow/processors/localProcessors.hpp b/src/phasicFlow/processors/localProcessors.hpp index eab18c08..1e1d6d3e 100644 --- a/src/phasicFlow/processors/localProcessors.hpp +++ b/src/phasicFlow/processors/localProcessors.hpp @@ -145,6 +145,15 @@ public: } #endif + static + bool builtForMPI() + { + #ifdef pFlow_Build_MPI + return true; + #else + return false; + #endif + } }; } diff --git a/src/phasicFlow/processors/processors.cpp b/src/phasicFlow/processors/processors.cpp index f5eb066a..078c9a5d 100644 --- a/src/phasicFlow/processors/processors.cpp +++ b/src/phasicFlow/processors/processors.cpp @@ -29,7 +29,6 @@ Licence: #include "processors.hpp" #include "streams.hpp" -static int numVarsInitialized__ = 0; #ifdef pFlow_Build_MPI diff --git a/src/phasicFlow/repository/IOobject/IOfileHeader.cpp b/src/phasicFlow/repository/IOobject/IOfileHeader.cpp index 96a49f67..f67eec28 100644 --- a/src/phasicFlow/repository/IOobject/IOfileHeader.cpp +++ b/src/phasicFlow/repository/IOobject/IOfileHeader.cpp @@ -42,12 +42,22 @@ pFlow::uniquePtr pFlow::IOfileHeader::outStream()const return osPtr; } -pFlow::IOfileHeader::IOfileHeader -( - const objectFile& objf -) -: - objectFile(objf) +pFlow::uniquePtr pFlow::IOfileHeader::dummyOutStream() const +{ + auto osPtr = makeUnique( CWD()+word("dummyFile") , outFileBinary()); + + if(osPtr && owner()) + { + auto outPrecision = owner()->outFilePrecision(); + osPtr->precision(static_cast(outPrecision)); + } + + return osPtr; +} + +pFlow::IOfileHeader::IOfileHeader( + const objectFile &objf) + : objectFile(objf) {} pFlow::fileSystem pFlow::IOfileHeader::path() const diff --git a/src/phasicFlow/repository/IOobject/IOfileHeader.hpp b/src/phasicFlow/repository/IOobject/IOfileHeader.hpp index 1d2a1e4c..d4549c01 100644 --- a/src/phasicFlow/repository/IOobject/IOfileHeader.hpp +++ b/src/phasicFlow/repository/IOobject/IOfileHeader.hpp @@ -58,6 +58,8 @@ protected: // - ouput file stream uniquePtr outStream()const; + uniquePtr dummyOutStream()const; + public: // with owner diff --git a/src/phasicFlow/repository/IOobject/IOobject.cpp b/src/phasicFlow/repository/IOobject/IOobject.cpp index b5857ea1..6fae2f0a 100644 --- a/src/phasicFlow/repository/IOobject/IOobject.cpp +++ b/src/phasicFlow/repository/IOobject/IOobject.cpp @@ -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() <= 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; diff --git a/src/phasicFlow/smartPointers/uniquePtr.hpp b/src/phasicFlow/smartPointers/uniquePtr.hpp index c1363cde..7bb6a6b9 100644 --- a/src/phasicFlow/smartPointers/uniquePtr.hpp +++ b/src/phasicFlow/smartPointers/uniquePtr.hpp @@ -37,16 +37,15 @@ namespace pFlow template< - typename T, - typename Deleter = std::default_delete + typename T > class uniquePtr : - public std::unique_ptr + public std::unique_ptr { public: - using uniquePtrType = std::unique_ptr; + using uniquePtrType = std::unique_ptr; // using base constructors using uniquePtrType::unique_ptr; diff --git a/src/phasicFlow/streams/TStream/iTstream.cpp b/src/phasicFlow/streams/TStream/iTstream.cpp index 49b80cee..ac8e503e 100755 --- a/src/phasicFlow/streams/TStream/iTstream.cpp +++ b/src/phasicFlow/streams/TStream/iTstream.cpp @@ -325,7 +325,7 @@ void pFlow::iTstream::reset() size_t pFlow::iTstream::tell() { notImplementedFunction; - return -1; + return static_cast(-1); } const pFlow::tokenList& pFlow::iTstream::tokens()const diff --git a/src/phasicFlow/structuredData/boundaries/boundaryBase/scatteredFieldAccess.hpp b/src/phasicFlow/structuredData/boundaries/boundaryBase/scatteredFieldAccess.hpp index e051c803..c1967bc8 100644 --- a/src/phasicFlow/structuredData/boundaries/boundaryBase/scatteredFieldAccess.hpp +++ b/src/phasicFlow/structuredData/boundaries/boundaryBase/scatteredFieldAccess.hpp @@ -124,6 +124,7 @@ public: return indices_; } + INLINE_FUNCTION_HD uint32 index(uint32 i)const { return indices_[i]; diff --git a/src/phasicFlow/structuredData/boundaries/boundaryReflective/boundaryReflective.cpp b/src/phasicFlow/structuredData/boundaries/boundaryReflective/boundaryReflective.cpp index 6d818a23..74ee4d75 100644 --- a/src/phasicFlow/structuredData/boundaries/boundaryReflective/boundaryReflective.cpp +++ b/src/phasicFlow/structuredData/boundaries/boundaryReflective/boundaryReflective.cpp @@ -125,6 +125,8 @@ bool pFlow::boundaryReflective::afterIteration const auto& velocity = time().lookupObject(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()); } } ); diff --git a/src/phasicFlow/structuredData/cylinder/cylinder.cpp b/src/phasicFlow/structuredData/cylinder/cylinder.cpp index 8c706d3e..014462c2 100644 --- a/src/phasicFlow/structuredData/cylinder/cylinder.cpp +++ b/src/phasicFlow/structuredData/cylinder/cylinder.cpp @@ -18,7 +18,7 @@ Licence: -----------------------------------------------------------------------------*/ - +#include "math.hpp" #include "cylinder.hpp" #include "zAxis.hpp" #include "streams.hpp" diff --git a/src/phasicFlow/structuredData/pointStructure/internalPoints/internalPoints.cpp b/src/phasicFlow/structuredData/pointStructure/internalPoints/internalPoints.cpp index 4dc8ff0d..0ea527f6 100644 --- a/src/phasicFlow/structuredData/pointStructure/internalPoints/internalPoints.cpp +++ b/src/phasicFlow/structuredData/pointStructure/internalPoints/internalPoints.cpp @@ -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<< " "<( + msg.addAndName(message::ITEM_INSERT), + newIndices + ); } const auto& indices = varList.getObject( @@ -450,7 +465,7 @@ bool pFlow::internalPoints::insertPointsOnly( }// we should fill the scattered empty spots else { - pOutput<<"numNew to be inserted "<< numNew <( msg.addAndName(message::ITEM_INSERT), newIndices diff --git a/src/phasicFlow/structuredData/pointStructure/internalPoints/pointFlagKernels.hpp b/src/phasicFlow/structuredData/pointStructure/internalPoints/pointFlagKernels.hpp index de8940ba..e76533b9 100644 --- a/src/phasicFlow/structuredData/pointStructure/internalPoints/pointFlagKernels.hpp +++ b/src/phasicFlow/structuredData/pointStructure/internalPoints/pointFlagKernels.hpp @@ -63,12 +63,13 @@ pFlow::ViewType1D::memo ViewType1D 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::markOutOfBoxDelete { minRange = 0; maxRange = 0; - numDeleted == numActive_; + numDeleted = numActive_; } else { @@ -180,7 +181,8 @@ pFlow::pointFlag::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::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(minRange), Kokkos::Max(maxRange), @@ -247,7 +249,7 @@ bool pFlow::pointFlag::deletePoints if(numDeleted >= numActive_) { activeRange_ = {0, 0}; - numDeleted == numActive_; + numDeleted = numActive_; } numActive_ = numActive_ - numDeleted; @@ -290,7 +292,7 @@ bool pFlow::pointFlag::deletePoints if(numDeleted >= numActive_) { activeRange_ = {0, 0}; - numDeleted == numActive_; + numDeleted = numActive_; } numActive_ = numActive_ - numDeleted; @@ -309,13 +311,14 @@ bool pFlow::pointFlag::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(); diff --git a/src/phasicFlow/structuredData/zAxis/array2D.hpp b/src/phasicFlow/structuredData/zAxis/array2D.hpp index f76e6080..3a1bb059 100644 --- a/src/phasicFlow/structuredData/zAxis/array2D.hpp +++ b/src/phasicFlow/structuredData/zAxis/array2D.hpp @@ -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; } }; diff --git a/src/phasicFlow/triSurface/triangleFunctions.hpp b/src/phasicFlow/triSurface/triangleFunctions.hpp index 7345c07f..c6812d2f 100644 --- a/src/phasicFlow/triSurface/triangleFunctions.hpp +++ b/src/phasicFlow/triSurface/triangleFunctions.hpp @@ -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); } diff --git a/src/phasicFlow/types/basicTypes/math.hpp b/src/phasicFlow/types/basicTypes/math.hpp index ee140058..fbe3a6f1 100755 --- a/src/phasicFlow/types/basicTypes/math.hpp +++ b/src/phasicFlow/types/basicTypes/math.hpp @@ -21,15 +21,16 @@ Licence: #ifndef __math_hpp__ #define __math_hpp__ + +#include "builtinTypes.hpp" +#include "pFlowMacros.hpp" + #ifdef __CUDACC__ #include "math.h" #else #include #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 diff --git a/src/setHelpers/setProperty.hpp b/src/setHelpers/setProperty.hpp index 43ec5ef2..978662d4 100644 --- a/src/setHelpers/setProperty.hpp +++ b/src/setHelpers/setProperty.hpp @@ -23,6 +23,6 @@ Licence: REPORT(0)<<"\nReading proprties . . . "<(vertecies)); solidNames_.push_back(name); } @@ -283,7 +283,7 @@ void pFlow::stlFile::addSolid realx3x3Vector&& vertecies ) { - solids_.push_backSafe(std::move(vertecies)); + solids_.push_back(makeUnique(vertecies)); solidNames_.push_back(name); } @@ -325,7 +325,7 @@ bool pFlow::stlFile::write()const { oFstream os(file_); os.precision(8); - for(label i=0; i= size() ) @@ -369,7 +369,7 @@ const pFlow::realx3x3Vector& pFlow::stlFile::solid const pFlow::word& pFlow::stlFile::name ( - label i + size_t i )const { if(i >= size() ) diff --git a/src/phasicFlow/triSurface/stlFile.hpp b/utilities/Utilities/geometryPhasicFlow/stlWall/stlFile.hpp similarity index 97% rename from src/phasicFlow/triSurface/stlFile.hpp rename to utilities/Utilities/geometryPhasicFlow/stlWall/stlFile.hpp index 6f47fbd5..eea2f928 100644 --- a/src/phasicFlow/triSurface/stlFile.hpp +++ b/utilities/Utilities/geometryPhasicFlow/stlWall/stlFile.hpp @@ -106,10 +106,10 @@ public: size_t size()const; // - vertecies of ith solid - const realx3x3Vector& solid(uint64 i)const; + const realx3x3Vector& solid(size_t i)const; // - name of ith solid - const word& name(uint64 i)const; + const word& name(size_t i)const; }; diff --git a/utilities/checkPhasicFlow/checkPhasicFlow.cpp b/utilities/checkPhasicFlow/checkPhasicFlow.cpp index 3ffe4168..cbd36d03 100755 --- a/utilities/checkPhasicFlow/checkPhasicFlow.cpp +++ b/utilities/checkPhasicFlow/checkPhasicFlow.cpp @@ -20,6 +20,7 @@ Licence: #include "KokkosTypes.hpp" #include "systemControl.hpp" +#include "localProcessors.hpp" #include "commandLine.hpp" @@ -40,9 +41,11 @@ if(!cmds.parse(argc, argv)) return 0; #include "initialize.hpp" output< pStructPtr = nullptr; + pFlow::uniquePtr pStructPtr = nullptr; if(!setOnly) @@ -84,13 +82,13 @@ int main( int argc, char* argv[] ) // position particles based on the dict content REPORT(0)<< "Positioning points . . . \n"<(Control, finalPos); + pStructPtr = pFlow::makeUnique(Control, finalPos); REPORT(1)<< "Created pStruct with "<< pStructPtr().size() << " points and capacity "<< @@ -100,11 +98,11 @@ int main( int argc, char* argv[] ) else { // read the content of pStruct from 0/pStructure - pStructPtr = makeUnique(Control); + pStructPtr = pFlow::makeUnique(Control); } - List> allObjects; + pFlow::List> allObjects; if(!positionOnly) { @@ -113,7 +111,7 @@ int main( int argc, char* argv[] ) auto& sfDict = cpDict.subDict("setFields"); - setFieldList defValueList(sfDict.subDict("defaultValue")); + pFlow::setFieldList defValueList(sfDict.subDict("defaultValue")); for(auto& sfEntry: defValueList) { @@ -128,7 +126,7 @@ int main( int argc, char* argv[] ) } } - output<("shapeName"); + auto& shapeName = Control.time().template lookupObject("shapeName"); REPORT(0)<< "Converting shapeName field to shapeIndex field"<