From 933eade4f5625a3455c719a1082968ec8c7ecb73 Mon Sep 17 00:00:00 2001 From: pca006132 Date: Sat, 14 May 2022 22:20:38 +0800 Subject: [PATCH] universal vector --- collider/src/collider.cpp | 12 ++-- manifold/src/boolean3.cpp | 21 ++++--- manifold/src/boolean_result.cpp | 51 +++++++-------- manifold/src/constructors.cpp | 21 ++++--- manifold/src/edge_op.cpp | 7 ++- manifold/src/impl.cpp | 32 +++++----- manifold/src/manifold.cpp | 9 +-- manifold/src/properties.cpp | 24 +++---- manifold/src/shared.h | 4 +- manifold/src/smoothing.cpp | 14 ++--- manifold/src/sort.cpp | 34 +++++----- utilities/include/sparse.h | 11 ++-- utilities/include/vec_dh.h | 108 +++++++++----------------------- 13 files changed, 153 insertions(+), 195 deletions(-) diff --git a/collider/src/collider.cpp b/collider/src/collider.cpp index dc86e0e12..781430e98 100644 --- a/collider/src/collider.cpp +++ b/collider/src/collider.cpp @@ -252,7 +252,7 @@ Collider::Collider(const VecDH& leafBB, nodeParent_.resize(num_nodes, -1); internalChildren_.resize(leafBB.size() - 1, thrust::make_pair(-1, -1)); // organize tree - thrust::for_each_n(countAt(0), NumInternal(), + thrust::for_each_n(thrust::device, countAt(0), NumInternal(), CreateRadixTree({nodeParent_.ptrD(), internalChildren_.ptrD(), leafMorton})); UpdateBoxes(leafBB); @@ -274,7 +274,7 @@ SparseIndices Collider::Collisions(const VecDH& querriesIn) const { VecDH nOverlapsD(1, 0); // calculate Bounding Box overlaps thrust::for_each_n( - zip(querriesIn.cbeginD(), countAt(0)), querriesIn.size(), + thrust::device, zip(querriesIn.cbeginD(), countAt(0)), querriesIn.size(), FindCollisions({querryTri.ptrDpq(), nOverlapsD.ptrD(), maxOverlaps, nodeBBox_.ptrD(), internalChildren_.ptrD()})); nOverlaps = nOverlapsD.H()[0]; @@ -305,13 +305,13 @@ void Collider::UpdateBoxes(const VecDH& leafBB) { // copy in leaf node Boxs strided_range::IterD> leaves(nodeBBox_.beginD(), nodeBBox_.endD(), 2); - thrust::copy(leafBB.cbeginD(), leafBB.cendD(), leaves.begin()); + thrust::copy(thrust::device, leafBB.cbeginD(), leafBB.cendD(), leaves.begin()); // create global counters VecDH counter(NumInternal()); - thrust::fill(counter.beginD(), counter.endD(), 0); + thrust::fill(thrust::device, counter.beginD(), counter.endD(), 0); // kernel over leaves to save internal Boxs thrust::for_each_n( - countAt(0), NumLeaves(), + thrust::device, countAt(0), NumLeaves(), BuildInternalBoxes({nodeBBox_.ptrD(), counter.ptrD(), nodeParent_.ptrD(), internalChildren_.ptrD()})); } @@ -330,7 +330,7 @@ bool Collider::Transform(glm::mat4x3 transform) { if (count != 2) axisAligned = false; } if (axisAligned) { - thrust::for_each(nodeBBox_.beginD(), nodeBBox_.endD(), + thrust::for_each(thrust::device, nodeBBox_.beginD(), nodeBBox_.endD(), TransformBox({transform})); } return axisAligned; diff --git a/manifold/src/boolean3.cpp b/manifold/src/boolean3.cpp index 1387c1efc..248a4ac18 100644 --- a/manifold/src/boolean3.cpp +++ b/manifold/src/boolean3.cpp @@ -14,6 +14,7 @@ #include "boolean3.h" #include +#include // TODO: make this runtime configurable for quicker debug constexpr bool kVerbose = false; @@ -85,12 +86,12 @@ struct CopyFaceEdges { SparseIndices Filter11(const Manifold::Impl &inP, const Manifold::Impl &inQ, const SparseIndices &p1q2, const SparseIndices &p2q1) { SparseIndices p1q1(3 * p1q2.size() + 3 * p2q1.size()); - thrust::for_each_n(zip(countAt(0), p1q2.beginD(0), p1q2.beginD(1)), + thrust::for_each_n(thrust::device, zip(countAt(0), p1q2.beginD(0), p1q2.beginD(1)), p1q2.size(), CopyFaceEdges({p1q1.ptrDpq(), inQ.halfedge_.cptrD()})); p1q1.SwapPQ(); - thrust::for_each_n(zip(countAt(p1q2.size()), p2q1.beginD(1), p2q1.beginD(0)), + thrust::for_each_n(thrust::device, zip(countAt(p1q2.size()), p2q1.beginD(1), p2q1.beginD(0)), p2q1.size(), CopyFaceEdges({p1q1.ptrDpq(), inP.halfedge_.cptrD()})); p1q1.SwapPQ(); @@ -246,7 +247,7 @@ std::tuple, VecDH> Shadow11(SparseIndices &p1q1, VecDH xyzz11(p1q1.size()); thrust::for_each_n( - zip(xyzz11.beginD(), s11.beginD(), p1q1.beginD(0), p1q1.beginD(1)), + thrust::device, zip(xyzz11.beginD(), s11.beginD(), p1q1.beginD(0), p1q1.beginD(1)), p1q1.size(), Kernel11({inP.vertPos_.cptrD(), inQ.vertPos_.cptrD(), inP.halfedge_.cptrD(), inQ.halfedge_.cptrD(), expandP, @@ -343,7 +344,7 @@ std::tuple, VecDH> Shadow02(const Manifold::Impl &inP, auto vertNormalP = forward ? inP.vertNormal_.cptrD() : inQ.vertNormal_.cptrD(); thrust::for_each_n( - zip(s02.beginD(), z02.beginD(), p0q2.beginD(!forward), + thrust::device, zip(s02.beginD(), z02.beginD(), p0q2.beginD(!forward), p0q2.beginD(forward)), p0q2.size(), Kernel02({inP.vertPos_.cptrD(), inQ.halfedge_.cptrD(), @@ -453,7 +454,7 @@ std::tuple, VecDH> Intersect12( VecDH v12(p1q2.size()); thrust::for_each_n( - zip(x12.beginD(), v12.beginD(), p1q2.beginD(!forward), + thrust::device, zip(x12.beginD(), v12.beginD(), p1q2.beginD(!forward), p1q2.beginD(forward)), p1q2.size(), Kernel12({p0q2.ptrDpq(), s02.ptrD(), z02.cptrD(), p0q2.size(), @@ -471,19 +472,19 @@ VecDH Winding03(const Manifold::Impl &inP, SparseIndices &p0q2, // verts that are not shadowed (not in p0q2) have winding number zero. VecDH w03(inP.NumVert(), 0); - if (!thrust::is_sorted(p0q2.beginD(reverse), p0q2.endD(reverse))) - thrust::sort_by_key(p0q2.beginD(reverse), p0q2.endD(reverse), s02.beginD()); + if (!thrust::is_sorted(thrust::device, p0q2.beginD(reverse), p0q2.endD(reverse))) + thrust::sort_by_key(thrust::device, p0q2.beginD(reverse), p0q2.endD(reverse), s02.beginD()); VecDH w03val(w03.size()); VecDH w03vert(w03.size()); // sum known s02 values into w03 (winding number) auto endPair = - thrust::reduce_by_key(p0q2.beginD(reverse), p0q2.endD(reverse), + thrust::reduce_by_key(thrust::device, p0q2.beginD(reverse), p0q2.endD(reverse), s02.beginD(), w03vert.beginD(), w03val.beginD()); - thrust::scatter(w03val.beginD(), endPair.second, w03vert.beginD(), + thrust::scatter(thrust::device, w03val.beginD(), endPair.second, w03vert.beginD(), w03.beginD()); if (reverse) - thrust::transform(w03.beginD(), w03.endD(), w03.beginD(), + thrust::transform(thrust::device, w03.beginD(), w03.endD(), w03.beginD(), thrust::negate()); return w03; }; diff --git a/manifold/src/boolean_result.cpp b/manifold/src/boolean_result.cpp index a98b3635c..6acfb5e2e 100644 --- a/manifold/src/boolean_result.cpp +++ b/manifold/src/boolean_result.cpp @@ -14,6 +14,7 @@ #include #include +#include #include "boolean3.h" #include "polygon.h" @@ -83,27 +84,27 @@ std::tuple, VecDH> SizeOutput( auto sidesPerFaceP = sidesPerFacePQ.ptrD(); auto sidesPerFaceQ = sidesPerFacePQ.ptrD() + inP.NumTri(); - thrust::for_each(inP.halfedge_.beginD(), inP.halfedge_.endD(), + thrust::for_each(thrust::device, inP.halfedge_.beginD(), inP.halfedge_.endD(), CountVerts({sidesPerFaceP, i03.cptrD()})); - thrust::for_each(inQ.halfedge_.beginD(), inQ.halfedge_.endD(), + thrust::for_each(thrust::device, inQ.halfedge_.beginD(), inQ.halfedge_.endD(), CountVerts({sidesPerFaceQ, i30.cptrD()})); thrust::for_each_n( - zip(p1q2.beginD(0), p1q2.beginD(1), i12.beginD()), i12.size(), + thrust::device, zip(p1q2.beginD(0), p1q2.beginD(1), i12.beginD()), i12.size(), CountNewVerts({sidesPerFaceP, sidesPerFaceQ, inP.halfedge_.cptrD()})); thrust::for_each_n( - zip(p2q1.beginD(1), p2q1.beginD(0), i21.beginD()), i21.size(), + thrust::device, zip(p2q1.beginD(1), p2q1.beginD(0), i21.beginD()), i21.size(), CountNewVerts({sidesPerFaceQ, sidesPerFaceP, inQ.halfedge_.cptrD()})); VecDH facePQ2R(inP.NumTri() + inQ.NumTri() + 1); auto keepFace = thrust::make_transform_iterator(sidesPerFacePQ.beginD(), NotZero()); - thrust::inclusive_scan(keepFace, keepFace + sidesPerFacePQ.size(), + thrust::inclusive_scan(thrust::device, keepFace, keepFace + sidesPerFacePQ.size(), facePQ2R.beginD() + 1); int numFaceR = facePQ2R.H().back(); facePQ2R.resize(inP.NumTri() + inQ.NumTri()); outR.faceNormal_.resize(numFaceR); - auto next = thrust::copy_if(inP.faceNormal_.beginD(), inP.faceNormal_.endD(), + auto next = thrust::copy_if(thrust::device, inP.faceNormal_.beginD(), inP.faceNormal_.endD(), keepFace, outR.faceNormal_.beginD(), thrust::identity()); if (invertQ) { @@ -111,17 +112,17 @@ std::tuple, VecDH> SizeOutput( thrust::negate()); auto end = thrust::make_transform_iterator(inQ.faceNormal_.endD(), thrust::negate()); - thrust::copy_if(start, end, keepFace + inP.NumTri(), next, + thrust::copy_if(thrust::device, start, end, keepFace + inP.NumTri(), next, thrust::identity()); } else { - thrust::copy_if(inQ.faceNormal_.beginD(), inQ.faceNormal_.endD(), + thrust::copy_if(thrust::device, inQ.faceNormal_.beginD(), inQ.faceNormal_.endD(), keepFace + inP.NumTri(), next, thrust::identity()); } auto newEnd = - thrust::remove(sidesPerFacePQ.beginD(), sidesPerFacePQ.endD(), 0); + thrust::remove(thrust::device, sidesPerFacePQ.beginD(), sidesPerFacePQ.endD(), 0); VecDH faceEdge(newEnd - sidesPerFacePQ.beginD() + 1); - thrust::inclusive_scan(sidesPerFacePQ.beginD(), newEnd, + thrust::inclusive_scan(thrust::device, sidesPerFacePQ.beginD(), newEnd, faceEdge.beginD() + 1); outR.halfedge_.resize(faceEdge.H().back()); @@ -211,7 +212,7 @@ void AppendPartialEdges(Manifold::Impl &outR, VecH &wholeHalfedgeP, std::map> &edgesP, VecH &halfedgeRef, const Manifold::Impl &inP, const VecH &i03, const VecH &vP2R, - const thrust::host_vector::const_iterator faceP2R, + const VecDH::IterHc faceP2R, bool forward) { // Each edge in the map is partially retained; for each of these, look up // their original verts and include them based on their winding number (i03), @@ -406,7 +407,7 @@ void AppendWholeEdges(Manifold::Impl &outR, VecDH &facePtrR, const VecDH &vP2R, const int *faceP2R, bool forward) { thrust::for_each_n( - zip(wholeHalfedgeP.beginD(), inP.halfedge_.beginD(), countAt(0)), + thrust::device, zip(wholeHalfedgeP.beginD(), inP.halfedge_.beginD(), countAt(0)), inP.halfedge_.size(), DuplicateHalfedges({outR.halfedge_.ptrD(), halfedgeRef.ptrD(), facePtrR.ptrD(), inP.halfedge_.cptrD(), i03.cptrD(), @@ -482,7 +483,7 @@ std::pair, VecDH> CalculateMeshRelation( VecDH halfedgeBary(halfedgeRef.size()); VecDH idx(1, 0); thrust::for_each_n( - zip(halfedgeBary.beginD(), halfedgeRef.beginD(), + thrust::device, zip(halfedgeBary.beginD(), halfedgeRef.beginD(), outR.halfedge_.cbeginD()), halfedgeRef.size(), CreateBarycentric( @@ -550,25 +551,25 @@ Manifold::Impl Boolean3::Result(Manifold::OpType op) const { VecDH i21(x21_.size()); VecDH i03(w03_.size()); VecDH i30(w30_.size()); - thrust::transform(x12_.beginD(), x12_.endD(), i12.beginD(), c3 * _1); - thrust::transform(x21_.beginD(), x21_.endD(), i21.beginD(), c3 * _1); - thrust::transform(w03_.beginD(), w03_.endD(), i03.beginD(), c1 + c3 * _1); - thrust::transform(w30_.beginD(), w30_.endD(), i30.beginD(), c2 + c3 * _1); + thrust::transform(thrust::device, x12_.beginD(), x12_.endD(), i12.beginD(), c3 * _1); + thrust::transform(thrust::device, x21_.beginD(), x21_.endD(), i21.beginD(), c3 * _1); + thrust::transform(thrust::device, w03_.beginD(), w03_.endD(), i03.beginD(), c1 + c3 * _1); + thrust::transform(thrust::device, w30_.beginD(), w30_.endD(), i30.beginD(), c2 + c3 * _1); VecDH vP2R(inP_.NumVert()); - thrust::exclusive_scan(i03.beginD(), i03.endD(), vP2R.beginD(), 0, AbsSum()); + thrust::exclusive_scan(thrust::device, i03.beginD(), i03.endD(), vP2R.beginD(), 0, AbsSum()); int numVertR = AbsSum()(vP2R.H().back(), i03.H().back()); const int nPv = numVertR; VecDH vQ2R(inQ_.NumVert()); - thrust::exclusive_scan(i30.beginD(), i30.endD(), vQ2R.beginD(), numVertR, + thrust::exclusive_scan(thrust::device, i30.beginD(), i30.endD(), vQ2R.beginD(), numVertR, AbsSum()); numVertR = AbsSum()(vQ2R.H().back(), i30.H().back()); const int nQv = numVertR - nPv; VecDH v12R(v12_.size()); if (v12_.size() > 0) { - thrust::exclusive_scan(i12.beginD(), i12.endD(), v12R.beginD(), numVertR, + thrust::exclusive_scan(thrust::device, i12.beginD(), i12.endD(), v12R.beginD(), numVertR, AbsSum()); numVertR = AbsSum()(v12R.H().back(), i12.H().back()); } @@ -576,7 +577,7 @@ Manifold::Impl Boolean3::Result(Manifold::OpType op) const { VecDH v21R(v21_.size()); if (v21_.size() > 0) { - thrust::exclusive_scan(i21.beginD(), i21.endD(), v21R.beginD(), numVertR, + thrust::exclusive_scan(thrust::device, i21.beginD(), i21.endD(), v21R.beginD(), numVertR, AbsSum()); numVertR = AbsSum()(v21R.H().back(), i21.H().back()); } @@ -592,14 +593,14 @@ Manifold::Impl Boolean3::Result(Manifold::OpType op) const { outR.vertPos_.resize(numVertR); // Add vertices, duplicating for inclusion numbers not in [-1, 1]. // Retained vertices from P and Q: - thrust::for_each_n(zip(i03.beginD(), vP2R.beginD(), inP_.vertPos_.beginD()), + thrust::for_each_n(thrust::device, zip(i03.beginD(), vP2R.beginD(), inP_.vertPos_.beginD()), inP_.NumVert(), DuplicateVerts({outR.vertPos_.ptrD()})); - thrust::for_each_n(zip(i30.beginD(), vQ2R.beginD(), inQ_.vertPos_.beginD()), + thrust::for_each_n(thrust::device, zip(i30.beginD(), vQ2R.beginD(), inQ_.vertPos_.beginD()), inQ_.NumVert(), DuplicateVerts({outR.vertPos_.ptrD()})); // New vertices created from intersections: - thrust::for_each_n(zip(i12.beginD(), v12R.beginD(), v12_.beginD()), + thrust::for_each_n(thrust::device, zip(i12.beginD(), v12R.beginD(), v12_.beginD()), i12.size(), DuplicateVerts({outR.vertPos_.ptrD()})); - thrust::for_each_n(zip(i21.beginD(), v21R.beginD(), v21_.beginD()), + thrust::for_each_n(thrust::device, zip(i21.beginD(), v21R.beginD(), v21_.beginD()), i21.size(), DuplicateVerts({outR.vertPos_.ptrD()})); if (kVerbose) { diff --git a/manifold/src/constructors.cpp b/manifold/src/constructors.cpp index 10e61457b..5641fcdb5 100644 --- a/manifold/src/constructors.cpp +++ b/manifold/src/constructors.cpp @@ -13,6 +13,7 @@ // limitations under the License. #include +#include #include "graph.h" #include "impl.h" @@ -183,7 +184,7 @@ Manifold Manifold::Sphere(float radius, int circularSegments) { Manifold sphere; sphere.pImpl_ = std::make_unique(Impl::Shape::OCTAHEDRON); sphere.pImpl_->Subdivide(n); - thrust::for_each_n(sphere.pImpl_->vertPos_.beginD(), sphere.NumVert(), + thrust::for_each_n(thrust::device, sphere.pImpl_->vertPos_.beginD(), sphere.NumVert(), ToSphere({radius})); sphere.pImpl_->Finish(); // Ignore preceding octahedron. @@ -403,20 +404,20 @@ Manifold Manifold::Compose(const std::vector& manifolds) { const Impl& impl = *(manifold.pImpl_); impl.ApplyTransform(); - thrust::copy(impl.vertPos_.beginD(), impl.vertPos_.endD(), + thrust::copy(thrust::device, impl.vertPos_.beginD(), impl.vertPos_.endD(), combined.vertPos_.beginD() + nextVert); - thrust::copy(impl.faceNormal_.beginD(), impl.faceNormal_.endD(), + thrust::copy(thrust::device, impl.faceNormal_.beginD(), impl.faceNormal_.endD(), combined.faceNormal_.beginD() + nextTri); - thrust::copy(impl.halfedgeTangent_.beginD(), impl.halfedgeTangent_.endD(), + thrust::copy(thrust::device, impl.halfedgeTangent_.beginD(), impl.halfedgeTangent_.endD(), combined.halfedgeTangent_.beginD() + nextEdge); - thrust::copy(impl.meshRelation_.barycentric.beginD(), + thrust::copy(thrust::device, impl.meshRelation_.barycentric.beginD(), impl.meshRelation_.barycentric.endD(), combined.meshRelation_.barycentric.beginD() + nextBary); - thrust::transform(impl.meshRelation_.triBary.beginD(), + thrust::transform(thrust::device, impl.meshRelation_.triBary.beginD(), impl.meshRelation_.triBary.endD(), combined.meshRelation_.triBary.beginD() + nextTri, UpdateTriBary({nextBary})); - thrust::transform(impl.halfedge_.beginD(), impl.halfedge_.endD(), + thrust::transform(thrust::device, impl.halfedge_.beginD(), impl.halfedge_.endD(), combined.halfedge_.beginD() + nextEdge, UpdateHalfedge({nextVert, nextEdge, nextTri})); @@ -461,7 +462,7 @@ std::vector Manifold::Decompose() const { VecDH vertNew2Old(NumVert()); int nVert = thrust::copy_if( - zip(pImpl_->vertPos_.beginD(), countAt(0)), + thrust::device, zip(pImpl_->vertPos_.beginD(), countAt(0)), zip(pImpl_->vertPos_.endD(), countAt(NumVert())), vertLabel.beginD(), zip(meshes[i].pImpl_->vertPos_.beginD(), vertNew2Old.beginD()), @@ -470,11 +471,11 @@ std::vector Manifold::Decompose() const { meshes[i].pImpl_->vertPos_.resize(nVert); VecDH faceNew2Old(NumTri()); - thrust::sequence(faceNew2Old.beginD(), faceNew2Old.endD()); + thrust::sequence(thrust::device, faceNew2Old.beginD(), faceNew2Old.endD()); int nFace = thrust::remove_if( - faceNew2Old.beginD(), faceNew2Old.endD(), + thrust::device, faceNew2Old.beginD(), faceNew2Old.endD(), RemoveFace({pImpl_->halfedge_.cptrD(), vertLabel.cptrD(), i})) - faceNew2Old.beginD(); faceNew2Old.resize(nFace); diff --git a/manifold/src/edge_op.cpp b/manifold/src/edge_op.cpp index 20878e1f0..1f8a82029 100644 --- a/manifold/src/edge_op.cpp +++ b/manifold/src/edge_op.cpp @@ -13,6 +13,7 @@ // limitations under the License. #include "impl.h" +#include namespace { using namespace manifold; @@ -121,7 +122,7 @@ void Manifold::Impl::SimplifyTopology() { VecDH flaggedEdges(halfedge_.size()); int numFlagged = thrust::copy_if( - countAt(0), countAt(halfedge_.size()), flaggedEdges.beginD(), + thrust::device, countAt(0), countAt(halfedge_.size()), flaggedEdges.beginD(), ShortEdge({halfedge_.cptrD(), vertPos_.cptrD(), precision_})) - flaggedEdges.beginD(); flaggedEdges.resize(numFlagged); @@ -131,7 +132,7 @@ void Manifold::Impl::SimplifyTopology() { flaggedEdges.resize(halfedge_.size()); numFlagged = thrust::copy_if( - countAt(0), countAt(halfedge_.size()), flaggedEdges.beginD(), + thrust::device, countAt(0), countAt(halfedge_.size()), flaggedEdges.beginD(), FlagEdge({halfedge_.cptrD(), meshRelation_.triBary.cptrD()})) - flaggedEdges.beginD(); flaggedEdges.resize(numFlagged); @@ -140,7 +141,7 @@ void Manifold::Impl::SimplifyTopology() { flaggedEdges.resize(halfedge_.size()); numFlagged = thrust::copy_if( - countAt(0), countAt(halfedge_.size()), flaggedEdges.beginD(), + thrust::device, countAt(0), countAt(halfedge_.size()), flaggedEdges.beginD(), SwappableEdge({halfedge_.cptrD(), vertPos_.cptrD(), faceNormal_.cptrD(), precision_})) - flaggedEdges.beginD(); diff --git a/manifold/src/impl.cpp b/manifold/src/impl.cpp index 1bca96287..a6e57b3a1 100644 --- a/manifold/src/impl.cpp +++ b/manifold/src/impl.cpp @@ -374,7 +374,7 @@ void Manifold::Impl::DuplicateMeshIDs() { } void Manifold::Impl::ReinitializeReference(int meshID) { - thrust::for_each_n(zip(meshRelation_.triBary.beginD(), countAt(0)), NumTri(), + thrust::for_each_n(thrust::device, zip(meshRelation_.triBary.beginD(), countAt(0)), NumTri(), InitializeBaryRef({meshID, halfedge_.cptrD()})); } @@ -402,7 +402,7 @@ int Manifold::Impl::InitializeNewReference( "propertyTolerance."); const int numSets = properties.size() / numProps; - ALWAYS_ASSERT(thrust::all_of(triPropertiesD.beginD(), triPropertiesD.endD(), + ALWAYS_ASSERT(thrust::all_of(thrust::device, triPropertiesD.beginD(), triPropertiesD.endD(), CheckProperties({numSets})), userErr, "triProperties value is outside the properties range."); @@ -411,7 +411,7 @@ int Manifold::Impl::InitializeNewReference( VecDH> face2face(halfedge_.size(), {-1, -1}); VecDH triArea(NumTri()); thrust::for_each_n( - zip(face2face.beginD(), countAt(0)), halfedge_.size(), + thrust::device, zip(face2face.beginD(), countAt(0)), halfedge_.size(), CoplanarEdge({triArea.ptrD(), halfedge_.cptrD(), vertPos_.cptrD(), triPropertiesD.cptrD(), propertiesD.cptrD(), propertyToleranceD.cptrD(), numProps, precision_})); @@ -493,10 +493,10 @@ void Manifold::Impl::CreateHalfedges(const VecDH& triVerts) { const int numTri = triVerts.size(); halfedge_.resize(3 * numTri); VecDH edge(3 * numTri); - thrust::for_each_n(zip(countAt(0), triVerts.beginD()), numTri, + thrust::for_each_n(thrust::device, zip(countAt(0), triVerts.beginD()), numTri, Tri2Halfedges({halfedge_.ptrD(), edge.ptrD()})); - thrust::sort(edge.beginD(), edge.endD()); - thrust::for_each_n(countAt(0), halfedge_.size() / 2, + thrust::sort(thrust::device, edge.beginD(), edge.endD()); + thrust::for_each_n(thrust::device, countAt(0), halfedge_.size() / 2, LinkHalfedges({halfedge_.ptrD(), edge.cptrD()})); } @@ -511,7 +511,7 @@ void Manifold::Impl::CreateAndFixHalfedges(const VecDH& triVerts) { halfedge_.resize(0); halfedge_.resize(3 * numTri); VecDH edge(3 * numTri); - thrust::for_each_n(zip(countAt(0), triVerts.beginD()), numTri, + thrust::for_each_n(thrust::device, zip(countAt(0), triVerts.beginD()), numTri, Tri2Halfedges({halfedge_.ptrD(), edge.ptrD()})); // Stable sort is required here so that halfedges from the same face are // paired together (the triangles were created in face order). In some @@ -519,7 +519,7 @@ void Manifold::Impl::CreateAndFixHalfedges(const VecDH& triVerts) { // two different faces, causing this edge to not be 2-manifold. We detect this // and fix it by swapping one of the identical edges, so it is important that // we have the edges paired according to their face. - thrust::stable_sort(edge.beginD(), edge.endD()); + thrust::stable_sort(thrust::device, edge.beginD(), edge.endD()); thrust::for_each_n(thrust::host, countAt(0), halfedge_.size() / 2, LinkHalfedges({halfedge_.ptrH(), edge.cptrH()})); thrust::for_each(thrust::host, countAt(1), countAt(halfedge_.size() / 2), @@ -551,14 +551,14 @@ void Manifold::Impl::ApplyTransform() const { */ void Manifold::Impl::ApplyTransform() { if (transform_ == glm::mat4x3(1.0f)) return; - thrust::for_each(vertPos_.beginD(), vertPos_.endD(), + thrust::for_each(thrust::device, vertPos_.beginD(), vertPos_.endD(), Transform4x3({transform_})); glm::mat3 normalTransform = glm::inverse(glm::transpose(glm::mat3(transform_))); - thrust::for_each(faceNormal_.beginD(), faceNormal_.endD(), + thrust::for_each(thrust::device, faceNormal_.beginD(), faceNormal_.endD(), TransformNormals({normalTransform})); - thrust::for_each(vertNormal_.beginD(), vertNormal_.endD(), + thrust::for_each(thrust::device, vertNormal_.beginD(), vertNormal_.endD(), TransformNormals({normalTransform})); // This optimization does a cheap collider update if the transform is // axis-aligned. @@ -601,17 +601,17 @@ void Manifold::Impl::SetPrecision(float minPrecision) { */ void Manifold::Impl::CalculateNormals() { vertNormal_.resize(NumVert()); - thrust::fill(vertNormal_.beginD(), vertNormal_.endD(), glm::vec3(0)); + thrust::fill(thrust::device, vertNormal_.beginD(), vertNormal_.endD(), glm::vec3(0)); bool calculateTriNormal = false; if (faceNormal_.size() != NumTri()) { faceNormal_.resize(NumTri()); calculateTriNormal = true; } thrust::for_each_n( - zip(faceNormal_.beginD(), countAt(0)), NumTri(), + thrust::device, zip(faceNormal_.beginD(), countAt(0)), NumTri(), AssignNormals({vertNormal_.ptrD(), vertPos_.cptrD(), halfedge_.cptrD(), precision_, calculateTriNormal})); - thrust::for_each(vertNormal_.beginD(), vertNormal_.endD(), Normalize()); + thrust::for_each(thrust::device, vertNormal_.beginD(), vertNormal_.endD(), Normalize()); } /** @@ -623,12 +623,12 @@ SparseIndices Manifold::Impl::EdgeCollisions(const Impl& Q) const { VecDH edges = CreateTmpEdges(Q.halfedge_); const int numEdge = edges.size(); VecDH QedgeBB(numEdge); - thrust::for_each_n(zip(QedgeBB.beginD(), edges.cbeginD()), numEdge, + thrust::for_each_n(thrust::device, zip(QedgeBB.beginD(), edges.cbeginD()), numEdge, EdgeBox({Q.vertPos_.cptrD()})); SparseIndices q1p2 = collider_.Collisions(QedgeBB); - thrust::for_each(q1p2.beginD(0), q1p2.endD(0), ReindexEdge({edges.cptrD()})); + thrust::for_each(thrust::device, q1p2.beginD(0), q1p2.endD(0), ReindexEdge({edges.cptrD()})); return q1p2; } diff --git a/manifold/src/manifold.cpp b/manifold/src/manifold.cpp index ee6bb7838..137f715b8 100644 --- a/manifold/src/manifold.cpp +++ b/manifold/src/manifold.cpp @@ -14,6 +14,7 @@ #include "boolean3.h" #include "impl.h" +#include namespace { using namespace manifold; @@ -124,7 +125,7 @@ Mesh Manifold::GetMesh() const { pImpl_->halfedgeTangent_.end()); result.triVerts.resize(NumTri()); - thrust::for_each_n(zip(result.triVerts.begin(), countAt(0)), NumTri(), + thrust::for_each_n(thrust::host, zip(result.triVerts.begin(), countAt(0)), NumTri(), MakeTri({pImpl_->halfedge_.cptrH()})); return result; @@ -288,11 +289,11 @@ MeshRelation Manifold::GetMeshRelation() const { std::vector Manifold::GetMeshIDs() const { VecDH meshIDs(NumTri()); thrust::for_each_n( - zip(meshIDs.beginD(), pImpl_->meshRelation_.triBary.beginD()), NumTri(), + thrust::device, zip(meshIDs.beginD(), pImpl_->meshRelation_.triBary.beginD()), NumTri(), GetMeshID()); - thrust::sort(meshIDs.beginD(), meshIDs.endD()); - int n = thrust::unique(meshIDs.beginD(), meshIDs.endD()) - meshIDs.beginD(); + thrust::sort(thrust::device, meshIDs.beginD(), meshIDs.endD()); + int n = thrust::unique(thrust::device, meshIDs.beginD(), meshIDs.endD()) - meshIDs.beginD(); meshIDs.resize(n); std::vector out; diff --git a/manifold/src/properties.cpp b/manifold/src/properties.cpp index 414e13105..672b61f03 100644 --- a/manifold/src/properties.cpp +++ b/manifold/src/properties.cpp @@ -215,12 +215,12 @@ namespace manifold { */ bool Manifold::Impl::IsManifold() const { if (halfedge_.size() == 0) return true; - bool isManifold = thrust::all_of(countAt(0), countAt(halfedge_.size()), + bool isManifold = thrust::all_of(thrust::device, countAt(0), countAt(halfedge_.size()), CheckManifold({halfedge_.cptrD()})); VecDH halfedge(halfedge_); - thrust::sort(halfedge.beginD(), halfedge.endD()); - isManifold &= thrust::all_of(countAt(0), countAt(2 * NumEdge() - 1), + thrust::sort(thrust::device, halfedge.beginD(), halfedge.endD()); + isManifold &= thrust::all_of(thrust::device, countAt(0), countAt(2 * NumEdge() - 1), NoDuplicates({halfedge.cptrD()})); return isManifold; } @@ -249,7 +249,7 @@ Properties Manifold::Impl::GetProperties() const { if (IsEmpty()) return {0, 0}; ApplyTransform(); thrust::pair areaVolume = thrust::transform_reduce( - countAt(0), countAt(NumTri()), + thrust::device, countAt(0), countAt(NumTri()), FaceAreaVolume({halfedge_.cptrD(), vertPos_.cptrD(), precision_}), thrust::make_pair(0.0f, 0.0f), SumPair()); return {areaVolume.first, areaVolume.second}; @@ -264,25 +264,25 @@ Curvature Manifold::Impl::GetCurvature() const { VecDH vertArea(NumVert(), 0); VecDH degree(NumVert(), 0); thrust::for_each( - countAt(0), countAt(NumTri()), + thrust::device, countAt(0), countAt(NumTri()), CurvatureAngles({vertMeanCurvature.ptrD(), vertGaussianCurvature.ptrD(), vertArea.ptrD(), degree.ptrD(), halfedge_.cptrD(), vertPos_.cptrD(), faceNormal_.cptrD()})); thrust::for_each_n( - zip(vertMeanCurvature.beginD(), vertGaussianCurvature.beginD(), + thrust::device, zip(vertMeanCurvature.beginD(), vertGaussianCurvature.beginD(), vertArea.beginD(), degree.beginD()), NumVert(), NormalizeCurvature()); result.minMeanCurvature = - thrust::reduce(vertMeanCurvature.beginD(), vertMeanCurvature.endD(), + thrust::reduce(thrust::device, vertMeanCurvature.beginD(), vertMeanCurvature.endD(), std::numeric_limits::infinity(), thrust::minimum()); result.maxMeanCurvature = - thrust::reduce(vertMeanCurvature.beginD(), vertMeanCurvature.endD(), + thrust::reduce(thrust::device, vertMeanCurvature.beginD(), vertMeanCurvature.endD(), -std::numeric_limits::infinity(), thrust::maximum()); result.minGaussianCurvature = thrust::reduce( - vertGaussianCurvature.beginD(), vertGaussianCurvature.endD(), std::numeric_limits::infinity(), + thrust::device, vertGaussianCurvature.beginD(), vertGaussianCurvature.endD(), std::numeric_limits::infinity(), thrust::minimum()); result.maxGaussianCurvature = thrust::reduce( - vertGaussianCurvature.beginD(), vertGaussianCurvature.endD(), + thrust::device, vertGaussianCurvature.beginD(), vertGaussianCurvature.endD(), -std::numeric_limits::infinity(), thrust::maximum()); result.vertMeanCurvature.insert(result.vertMeanCurvature.end(), vertMeanCurvature.begin(), @@ -299,9 +299,9 @@ Curvature Manifold::Impl::GetCurvature() const { * range for Morton code calculation. */ void Manifold::Impl::CalculateBBox() { - bBox_.min = thrust::reduce(vertPos_.beginD(), vertPos_.endD(), + bBox_.min = thrust::reduce(thrust::device, vertPos_.beginD(), vertPos_.endD(), glm::vec3(std::numeric_limits::infinity()), PosMin()); - bBox_.max = thrust::reduce(vertPos_.beginD(), vertPos_.endD(), + bBox_.max = thrust::reduce(thrust::device, vertPos_.beginD(), vertPos_.endD(), glm::vec3(-std::numeric_limits::infinity()), PosMax()); } } // namespace manifold diff --git a/manifold/src/shared.h b/manifold/src/shared.h index 317f2341b..295810a4f 100644 --- a/manifold/src/shared.h +++ b/manifold/src/shared.h @@ -146,9 +146,9 @@ struct TmpInvalid { VecDH inline CreateTmpEdges(const VecDH& halfedge) { VecDH edges(halfedge.size()); - thrust::for_each_n(zip(edges.beginD(), halfedge.beginD(), countAt(0)), + thrust::for_each_n(thrust::device, zip(edges.beginD(), halfedge.beginD(), countAt(0)), edges.size(), Halfedge2Tmp()); - int numEdge = thrust::remove_if(edges.beginD(), edges.endD(), TmpInvalid()) - + int numEdge = thrust::remove_if(thrust::device, edges.beginD(), edges.endD(), TmpInvalid()) - edges.beginD(); ALWAYS_ASSERT(numEdge == halfedge.size() / 2, topologyErr, "Not oriented!"); edges.resize(numEdge); diff --git a/manifold/src/smoothing.cpp b/manifold/src/smoothing.cpp index 6c81bb4e2..b86268d36 100644 --- a/manifold/src/smoothing.cpp +++ b/manifold/src/smoothing.cpp @@ -356,7 +356,7 @@ void Manifold::Impl::CreateTangents( const int numHalfedge = halfedge_.size(); halfedgeTangent_.resize(numHalfedge); - thrust::for_each_n(zip(halfedgeTangent_.beginD(), halfedge_.cbeginD()), + thrust::for_each_n(thrust::device, zip(halfedgeTangent_.beginD(), halfedge_.cbeginD()), numHalfedge, SmoothBezier({vertPos_.cptrD(), faceNormal_.cptrD(), vertNormal_.cptrD(), halfedge_.cptrD()})); @@ -476,12 +476,12 @@ Manifold::Impl::MeshRelationD Manifold::Impl::Subdivide(int n) { VecDH edges = CreateTmpEdges(halfedge_); VecDH half2Edge(2 * numEdge); - thrust::for_each_n(zip(countAt(0), edges.beginD()), numEdge, + thrust::for_each_n(thrust::device, zip(countAt(0), edges.beginD()), numEdge, ReindexHalfedge({half2Edge.ptrD()})); - thrust::for_each_n(zip(countAt(0), edges.beginD()), numEdge, + thrust::for_each_n(thrust::device, zip(countAt(0), edges.beginD()), numEdge, EdgeVerts({vertPos_.ptrD(), numVert, n})); thrust::for_each_n( - zip(countAt(0), oldMeshRelation.triBary.beginD()), numTri, + thrust::device, zip(countAt(0), oldMeshRelation.triBary.beginD()), numTri, InteriorVerts({vertPos_.ptrD(), relation.barycentric.ptrD(), relation.triBary.ptrD(), meshRelation_.barycentric.ptrD(), meshRelation_.triBary.ptrD(), @@ -489,7 +489,7 @@ Manifold::Impl::MeshRelationD Manifold::Impl::Subdivide(int n) { halfedge_.ptrD()})); // Create subtriangles VecDH triVerts(n * n * numTri); - thrust::for_each_n(countAt(0), numTri, + thrust::for_each_n(thrust::device, countAt(0), numTri, SplitTris({triVerts.ptrD(), halfedge_.cptrD(), half2Edge.cptrD(), numVert, triVertStart, n})); CreateHalfedges(triVerts); @@ -504,12 +504,12 @@ void Manifold::Impl::Refine(int n) { VecDH vertBary(NumVert()); VecDH lock(NumVert(), 0); thrust::for_each_n( - zip(relation.triBary.beginD(), countAt(0)), NumTri(), + thrust::device, zip(relation.triBary.beginD(), countAt(0)), NumTri(), TriBary2Vert({vertBary.ptrD(), lock.ptrD(), relation.barycentric.cptrD(), halfedge_.cptrD()})); thrust::for_each_n( - zip(vertPos_.beginD(), vertBary.beginD()), NumVert(), + thrust::device, zip(vertPos_.beginD(), vertBary.beginD()), NumVert(), InterpTri({old.halfedge_.cptrD(), old.halfedgeTangent_.cptrD(), old.vertPos_.cptrD()})); } diff --git a/manifold/src/sort.cpp b/manifold/src/sort.cpp index dd327082d..94b0fef09 100644 --- a/manifold/src/sort.cpp +++ b/manifold/src/sort.cpp @@ -187,7 +187,7 @@ void Manifold::Impl::Finish() { "Not an even number of faces after sorting faces!"); Halfedge extrema = {0, 0, 0, 0}; extrema = - thrust::reduce(halfedge_.beginD(), halfedge_.endD(), extrema, Extrema()); + thrust::reduce(thrust::device, halfedge_.beginD(), halfedge_.endD(), extrema, Extrema()); ALWAYS_ASSERT(extrema.startVert >= 0, topologyErr, "Vertex index is negative!"); @@ -210,12 +210,12 @@ void Manifold::Impl::Finish() { */ void Manifold::Impl::SortVerts() { VecDH vertMorton(NumVert()); - thrust::for_each_n(zip(vertMorton.beginD(), vertPos_.cbeginD()), NumVert(), + thrust::for_each_n(thrust::device, zip(vertMorton.beginD(), vertPos_.cbeginD()), NumVert(), Morton({bBox_})); VecDH vertNew2Old(NumVert()); - thrust::sequence(vertNew2Old.beginD(), vertNew2Old.endD()); - thrust::sort_by_key(vertMorton.beginD(), vertMorton.endD(), + thrust::sequence(thrust::device, vertNew2Old.beginD(), vertNew2Old.endD()); + thrust::sort_by_key(thrust::device, vertMorton.beginD(), vertMorton.endD(), zip(vertPos_.beginD(), vertNew2Old.beginD())); ReindexVerts(vertNew2Old, NumVert()); @@ -223,7 +223,7 @@ void Manifold::Impl::SortVerts() { // Verts were flagged for removal with NaNs and assigned kNoCode to sort them // to the end, which allows them to be removed. const int newNumVert = - thrust::find(vertMorton.beginD(), vertMorton.endD(), kNoCode) - + thrust::find(thrust::device, vertMorton.beginD(), vertMorton.endD(), kNoCode) - vertMorton.beginD(); vertPos_.resize(newNumVert); } @@ -236,9 +236,9 @@ void Manifold::Impl::SortVerts() { void Manifold::Impl::ReindexVerts(const VecDH& vertNew2Old, int oldNumVert) { VecDH vertOld2New(oldNumVert); - thrust::scatter(countAt(0), countAt(NumVert()), vertNew2Old.beginD(), + thrust::scatter(thrust::device, countAt(0), countAt(NumVert()), vertNew2Old.beginD(), vertOld2New.beginD()); - thrust::for_each(halfedge_.beginD(), halfedge_.endD(), + thrust::for_each(thrust::device, halfedge_.beginD(), halfedge_.endD(), Reindex({vertOld2New.cptrD()})); } @@ -252,7 +252,7 @@ void Manifold::Impl::GetFaceBoxMorton(VecDH& faceBox, faceBox.resize(NumTri()); faceMorton.resize(NumTri()); thrust::for_each_n( - zip(faceMorton.beginD(), faceBox.beginD(), countAt(0)), NumTri(), + thrust::device, zip(faceMorton.beginD(), faceBox.beginD(), countAt(0)), NumTri(), FaceMortonBox({halfedge_.cptrD(), vertPos_.cptrD(), bBox_})); } @@ -263,15 +263,15 @@ void Manifold::Impl::GetFaceBoxMorton(VecDH& faceBox, void Manifold::Impl::SortFaces(VecDH& faceBox, VecDH& faceMorton) { VecDH faceNew2Old(NumTri()); - thrust::sequence(faceNew2Old.beginD(), faceNew2Old.endD()); + thrust::sequence(thrust::device, faceNew2Old.beginD(), faceNew2Old.endD()); - thrust::sort_by_key(faceMorton.beginD(), faceMorton.endD(), + thrust::sort_by_key(thrust::device, faceMorton.beginD(), faceMorton.endD(), zip(faceBox.beginD(), faceNew2Old.beginD())); // Tris were flagged for removal with pairedHalfedge = -1 and assigned kNoCode // to sort them to the end, which allows them to be removed. const int newNumTri = - thrust::find(faceMorton.beginD(), faceMorton.endD(), kNoCode) - + thrust::find(thrust::device, faceMorton.beginD(), faceMorton.endD(), kNoCode) - faceMorton.beginD(); faceBox.resize(newNumTri); faceMorton.resize(newNumTri); @@ -295,13 +295,13 @@ void Manifold::Impl::GatherFaces(const VecDH& faceNew2Old) { VecDH oldHalfedge(std::move(halfedge_)); VecDH oldHalfedgeTangent(std::move(halfedgeTangent_)); VecDH faceOld2New(oldHalfedge.size() / 3); - thrust::scatter(countAt(0), countAt(numTri), faceNew2Old.beginD(), + thrust::scatter(thrust::device, countAt(0), countAt(numTri), faceNew2Old.beginD(), faceOld2New.beginD()); halfedge_.resize(3 * numTri); if (oldHalfedgeTangent.size() != 0) halfedgeTangent_.resize(3 * numTri); thrust::for_each_n( - countAt(0), numTri, + thrust::device, countAt(0), numTri, ReindexFace({halfedge_.ptrD(), halfedgeTangent_.ptrD(), oldHalfedge.cptrD(), oldHalfedgeTangent.cptrD(), faceNew2Old.cptrD(), faceOld2New.cptrD()})); @@ -311,7 +311,7 @@ void Manifold::Impl::GatherFaces(const Impl& old, const VecDH& faceNew2Old) { const int numTri = faceNew2Old.size(); meshRelation_.triBary.resize(numTri); - thrust::gather(faceNew2Old.beginD(), faceNew2Old.endD(), + thrust::gather(thrust::device, faceNew2Old.beginD(), faceNew2Old.endD(), old.meshRelation_.triBary.beginD(), meshRelation_.triBary.beginD()); meshRelation_.barycentric = old.meshRelation_.barycentric; @@ -319,18 +319,18 @@ void Manifold::Impl::GatherFaces(const Impl& old, if (old.faceNormal_.size() == old.NumTri()) { faceNormal_.resize(numTri); - thrust::gather(faceNew2Old.beginD(), faceNew2Old.endD(), + thrust::gather(thrust::device, faceNew2Old.beginD(), faceNew2Old.endD(), old.faceNormal_.beginD(), faceNormal_.beginD()); } VecDH faceOld2New(old.NumTri()); - thrust::scatter(countAt(0), countAt(numTri), faceNew2Old.beginD(), + thrust::scatter(thrust::device, countAt(0), countAt(numTri), faceNew2Old.beginD(), faceOld2New.beginD()); halfedge_.resize(3 * numTri); if (old.halfedgeTangent_.size() != 0) halfedgeTangent_.resize(3 * numTri); thrust::for_each_n( - countAt(0), numTri, + thrust::device, countAt(0), numTri, ReindexFace({halfedge_.ptrD(), halfedgeTangent_.ptrD(), old.halfedge_.cptrD(), old.halfedgeTangent_.cptrD(), faceNew2Old.cptrD(), faceOld2New.cptrD()})); diff --git a/utilities/include/sparse.h b/utilities/include/sparse.h index 6e3397e66..953f03305 100644 --- a/utilities/include/sparse.h +++ b/utilities/include/sparse.h @@ -18,6 +18,7 @@ #include #include #include +#include #include "structs.h" #include "utils.h" @@ -125,7 +126,7 @@ class SparseIndices { "Different number of values than indicies!"); auto zBegin = zip(v.beginD(), x.beginD(), beginD(false), beginD(true)); auto zEnd = zip(v.endD(), x.endD(), endD(false), endD(true)); - size_t size = thrust::remove_if(zBegin, zEnd, firstNonFinite()) - zBegin; + size_t size = thrust::remove_if(thrust::device, zBegin, zEnd, firstNonFinite()) - zBegin; v.resize(size); x.resize(size, -1); p.resize(size, -1); @@ -142,10 +143,10 @@ class SparseIndices { VecDH result(size); VecDH found(size); VecDH temp(size); - thrust::fill(result.beginD(), result.endD(), missingVal); - thrust::binary_search(beginDpq(), endDpq(), pqBegin, pqEnd, found.beginD()); - thrust::lower_bound(beginDpq(), endDpq(), pqBegin, pqEnd, temp.beginD()); - thrust::gather_if(temp.beginD(), temp.endD(), found.beginD(), val.beginD(), + thrust::fill(thrust::device, result.beginD(), result.endD(), missingVal); + thrust::binary_search(thrust::device, beginDpq(), endDpq(), pqBegin, pqEnd, found.beginD()); + thrust::lower_bound(thrust::device, beginDpq(), endDpq(), pqBegin, pqEnd, temp.beginD()); + thrust::gather_if(thrust::device, temp.beginD(), temp.endD(), found.beginD(), val.beginD(), result.beginD()); return result; } diff --git a/utilities/include/vec_dh.h b/utilities/include/vec_dh.h index 38191b72f..ccb713baf 100644 --- a/utilities/include/vec_dh.h +++ b/utilities/include/vec_dh.h @@ -13,8 +13,9 @@ // limitations under the License. #pragma once -#include -#include +#include +#include +#include namespace manifold { @@ -22,7 +23,7 @@ namespace manifold { * @{ */ template -using VecH = thrust::host_vector; +using VecH = thrust::universal_vector; template void Dump(const VecH& vec) { @@ -39,86 +40,63 @@ class VecDH { VecDH() {} VecDH(int size, T val = T()) { - device_.resize(size, val); - host_valid_ = false; + impl_.resize(size, val); } VecDH(const std::vector& vec) { - host_ = vec; - device_valid_ = false; + impl_ = vec; } - int size() const { return device_valid_ ? device_.size() : host_.size(); } + int size() const { return impl_.size(); } void resize(int newSize, T val = T()) { bool shrink = size() > 2 * newSize; - if (device_valid_) { - device_.resize(newSize, val); - if (shrink) device_.shrink_to_fit(); - } - if (host_valid_) { - host_.resize(newSize, val); - if (shrink) host_.shrink_to_fit(); - } + impl_.resize(newSize, val); + if (shrink) impl_.shrink_to_fit(); } void swap(VecDH& other) { - host_.swap(other.host_); - device_.swap(other.device_); - thrust::swap(host_valid_, other.host_valid_); - thrust::swap(device_valid_, other.device_valid_); + impl_.swap(other.impl_); } - using IterD = typename thrust::device_vector::iterator; - using IterH = typename thrust::host_vector::iterator; - using IterDc = typename thrust::device_vector::const_iterator; - using IterHc = typename thrust::host_vector::const_iterator; + using IterD = typename thrust::universal_vector::iterator; + using IterH = typename thrust::universal_vector::iterator; + using IterDc = typename thrust::universal_vector::const_iterator; + using IterHc = typename thrust::universal_vector::const_iterator; IterH begin() { - RefreshHost(); - device_valid_ = false; - return host_.begin(); + return impl_.begin(); } IterH end() { - RefreshHost(); - device_valid_ = false; - return host_.end(); + return impl_.end(); } IterHc cbegin() const { - RefreshHost(); - return host_.cbegin(); + return impl_.cbegin(); } IterHc cend() const { - RefreshHost(); - return host_.cend(); + return impl_.cend(); } IterHc begin() const { return cbegin(); } IterHc end() const { return cend(); } IterD beginD() { - RefreshDevice(); - host_valid_ = false; - return device_.begin(); + return impl_.begin(); } IterD endD() { - RefreshDevice(); - host_valid_ = false; - return device_.end(); + return impl_.end(); } IterDc cbeginD() const { - RefreshDevice(); - return device_.cbegin(); + return impl_.cbegin(); } IterDc cendD() const { - RefreshDevice(); - return device_.cend(); + return impl_.cend(); } IterDc beginD() const { return cbeginD(); } @@ -126,66 +104,40 @@ class VecDH { T* ptrD() { if (size() == 0) return nullptr; - RefreshDevice(); - host_valid_ = false; - return device_.data().get(); + return impl_.data().get(); } const T* cptrD() const { if (size() == 0) return nullptr; - RefreshDevice(); - return device_.data().get(); + return impl_.data().get(); } const T* ptrD() const { return cptrD(); } T* ptrH() { if (size() == 0) return nullptr; - RefreshHost(); - device_valid_ = false; - return host_.data(); + return impl_.data().get(); } const T* cptrH() const { if (size() == 0) return nullptr; - RefreshHost(); - return host_.data(); + return impl_.data().get(); } const T* ptrH() const { return cptrH(); } const VecH& H() const { - RefreshHost(); - return host_; + return impl_; } VecH& H() { - RefreshHost(); - device_valid_ = false; - return host_; + return impl_; } void Dump() const { manifold::Dump(H()); } private: - mutable bool host_valid_ = true; - mutable bool device_valid_ = true; - mutable thrust::host_vector host_; - mutable thrust::device_vector device_; - - void RefreshHost() const { - if (!host_valid_) { - host_ = device_; - host_valid_ = true; - } - } - - void RefreshDevice() const { - if (!device_valid_) { - device_ = host_; - device_valid_ = true; - } - } + mutable thrust::universal_vector impl_; }; template @@ -201,4 +153,4 @@ class VecD { const int size_; }; /** @} */ -} // namespace manifold \ No newline at end of file +} // namespace manifold