From c31683618e13d76385b50d50b293a9f07e30e819 Mon Sep 17 00:00:00 2001 From: pca006132 Date: Sat, 14 May 2022 22:20:38 +0800 Subject: [PATCH 1/3] universal vector --- collider/src/collider.cpp | 16 +- manifold/src/boolean3.cpp | 21 +-- manifold/src/boolean_result.cpp | 125 ++++++------- manifold/src/constructors.cpp | 29 +-- manifold/src/edge_op.cpp | 89 +++++----- manifold/src/face_op.cpp | 49 +++--- manifold/src/impl.cpp | 50 +++--- manifold/src/impl.h | 2 +- manifold/src/manifold.cpp | 11 +- manifold/src/properties.cpp | 24 +-- manifold/src/shared.h | 4 +- manifold/src/smoothing.cpp | 20 +-- manifold/src/sort.cpp | 34 ++-- utilities/include/sparse.h | 17 +- utilities/include/structs.h | 24 +-- utilities/include/vec_dh.h | 300 +++++++++++++++++++++++--------- 16 files changed, 478 insertions(+), 337 deletions(-) diff --git a/collider/src/collider.cpp b/collider/src/collider.cpp index dc86e0e12..ff8eadd48 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,14 +274,14 @@ 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]; + nOverlaps = nOverlapsD[0]; if (nOverlaps <= maxOverlaps) break; else { // if not enough memory was allocated, guess how much will be needed - int lastQuery = querryTri.Get(0).H().back(); + int lastQuery = querryTri.Get(0).back(); float ratio = static_cast(querriesIn.size()) / lastQuery; if (ratio > 1000) // do not trust the ratio if it is too large maxOverlaps *= 2; @@ -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..3e8cc17cd 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(); + int numFaceR = facePQ2R.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,19 +112,19 @@ 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()); + outR.halfedge_.resize(faceEdge.back()); return std::make_tuple(faceEdge, facePQ2R); } @@ -137,15 +138,15 @@ struct EdgePos { void AddNewEdgeVerts( std::map> &edgesP, std::map, std::vector> &edgesNew, - const SparseIndices &p1q2, const VecH &i12, const VecH &v12R, - const VecH &halfedgeP, bool forward) { + const SparseIndices &p1q2, const VecDH &i12, const VecDH &v12R, + const VecDH &halfedgeP, bool forward) { // For each edge of P that intersects a face of Q (p1q2), add this vertex to // P's corresponding edge vector and to the two new edges, which are // intersections between the face of Q and the two faces of P attached to the // edge. The direction and duplicity are given by i12, while v12R remaps to // the output vert index. When forward is false, all is reversed. - const VecH &p1 = p1q2.Get(!forward).H(); - const VecH &q2 = p1q2.Get(forward).H(); + const VecDH &p1 = p1q2.Get(!forward); + const VecDH &q2 = p1q2.Get(forward); for (int i = 0; i < p1q2.size(); ++i) { const int edgeP = p1[i]; const int faceQ = q2[i]; @@ -206,21 +207,21 @@ struct Ref { int PQ, tri, vert; }; -void AppendPartialEdges(Manifold::Impl &outR, VecH &wholeHalfedgeP, - VecH &facePtrR, +void AppendPartialEdges(Manifold::Impl &outR, VecDH &wholeHalfedgeP, + VecDH &facePtrR, std::map> &edgesP, - VecH &halfedgeRef, const Manifold::Impl &inP, - const VecH &i03, const VecH &vP2R, - const thrust::host_vector::const_iterator faceP2R, + VecDH &halfedgeRef, const Manifold::Impl &inP, + const VecDH &i03, const VecDH &vP2R, + 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), // while remaping them to the output using vP2R. Use the verts position // projected along the edge vector to pair them up, then distribute these // edges to their faces. - VecH &halfedgeR = outR.halfedge_.H(); - const VecH &vertPosP = inP.vertPos_.H(); - const VecH &halfedgeP = inP.halfedge_.H(); + VecDH &halfedgeR = outR.halfedge_; + const VecDH &vertPosP = inP.vertPos_; + const VecDH &halfedgeP = inP.halfedge_; for (auto &value : edgesP) { const int edgeP = value.first; @@ -235,13 +236,13 @@ void AppendPartialEdges(Manifold::Impl &outR, VecH &wholeHalfedgeP, const glm::vec3 edgeVec = vertPosP[vEnd] - vertPosP[vStart]; // Fill in the edge positions of the old points. for (EdgePos &edge : edgePosP) { - edge.edgePos = glm::dot(outR.vertPos_.H()[edge.vert], edgeVec); + edge.edgePos = glm::dot(outR.vertPos_[edge.vert], edgeVec); } int inclusion = i03[vStart]; bool reversed = inclusion < 0; EdgePos edgePos = {vP2R[vStart], - glm::dot(outR.vertPos_.H()[vP2R[vStart]], edgeVec), + glm::dot(outR.vertPos_[vP2R[vStart]], edgeVec), inclusion > 0}; for (int j = 0; j < glm::abs(inclusion); ++j) { edgePosP.push_back(edgePos); @@ -250,7 +251,7 @@ void AppendPartialEdges(Manifold::Impl &outR, VecH &wholeHalfedgeP, inclusion = i03[vEnd]; reversed |= inclusion < 0; - edgePos = {vP2R[vEnd], glm::dot(outR.vertPos_.H()[vP2R[vEnd]], edgeVec), + edgePos = {vP2R[vEnd], glm::dot(outR.vertPos_[vP2R[vEnd]], edgeVec), inclusion < 0}; for (int j = 0; j < glm::abs(inclusion); ++j) { edgePosP.push_back(edgePos); @@ -294,12 +295,12 @@ void AppendPartialEdges(Manifold::Impl &outR, VecH &wholeHalfedgeP, } void AppendNewEdges( - Manifold::Impl &outR, VecH &facePtrR, + Manifold::Impl &outR, VecDH &facePtrR, std::map, std::vector> &edgesNew, - VecH &halfedgeRef, const VecH &facePQ2R, const int numFaceP) { + VecDH &halfedgeRef, const VecDH &facePQ2R, const int numFaceP) { // Pair up each edge's verts and distribute to faces based on indices in key. - VecH &halfedgeR = outR.halfedge_.H(); - VecH &vertPosR = outR.vertPos_.H(); + VecDH &halfedgeR = outR.halfedge_; + VecDH &vertPosR = outR.vertPos_; for (auto &value : edgesNew) { const int faceP = value.first.first; @@ -402,11 +403,11 @@ struct DuplicateHalfedges { void AppendWholeEdges(Manifold::Impl &outR, VecDH &facePtrR, VecDH &halfedgeRef, const Manifold::Impl &inP, - const VecDH wholeHalfedgeP, const VecDH &i03, + const VecDH wholeHalfedgeP, const VecDH &i03, 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( @@ -492,7 +493,7 @@ std::pair, VecDH> CalculateMeshRelation( inP.meshRelation_.triBary.cptrD(), inQ.meshRelation_.triBary.cptrD(), inP.meshRelation_.barycentric.cptrD(), inQ.meshRelation_.barycentric.cptrD(), invertQ, outR.precision_})); - outR.meshRelation_.barycentric.resize(idx.H()[0]); + outR.meshRelation_.barycentric.resize(idx[0]); return std::make_pair(faceRef, halfedgeBary); } } // namespace @@ -550,35 +551,35 @@ 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()); - int numVertR = AbsSum()(vP2R.H().back(), i03.H().back()); + thrust::exclusive_scan(thrust::device, i03.beginD(), i03.endD(), vP2R.beginD(), 0, AbsSum()); + int numVertR = AbsSum()(vP2R.back(), i03.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()); + numVertR = AbsSum()(vQ2R.back(), i30.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()); + numVertR = AbsSum()(v12R.back(), i12.back()); } const int n12 = numVertR - nPv - nQv; 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()); + numVertR = AbsSum()(v21R.back(), i21.back()); } const int n21 = numVertR - nPv - nQv - n12; @@ -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) { @@ -620,10 +621,10 @@ Manifold::Impl Boolean3::Result(Manifold::OpType op) const { // This key is the face index of std::map, std::vector> edgesNew; - AddNewEdgeVerts(edgesP, edgesNew, p1q2_, i12.H(), v12R.H(), - inP_.halfedge_.H(), true); - AddNewEdgeVerts(edgesQ, edgesNew, p2q1_, i21.H(), v21R.H(), - inQ_.halfedge_.H(), false); + AddNewEdgeVerts(edgesP, edgesNew, p1q2_, i12, v12R, + inP_.halfedge_, true); + AddNewEdgeVerts(edgesQ, edgesNew, p2q1_, i21, v21R, + inQ_.halfedge_, false); // Level 4 VecDH faceEdge; @@ -636,20 +637,20 @@ Manifold::Impl Boolean3::Result(Manifold::OpType op) const { // next one knows where to slot in. VecDH facePtrR = faceEdge; // Intersected halfedges are marked false. - VecDH wholeHalfedgeP(inP_.halfedge_.size(), true); - VecDH wholeHalfedgeQ(inQ_.halfedge_.size(), true); + VecDH wholeHalfedgeP(inP_.halfedge_.size(), true); + VecDH wholeHalfedgeQ(inQ_.halfedge_.size(), true); // The halfedgeRef contains the data that will become triBary once the faces // are triangulated. VecDH halfedgeRef(2 * outR.NumEdge()); - AppendPartialEdges(outR, wholeHalfedgeP.H(), facePtrR.H(), edgesP, - halfedgeRef.H(), inP_, i03.H(), vP2R.H(), facePQ2R.begin(), + AppendPartialEdges(outR, wholeHalfedgeP, facePtrR, edgesP, + halfedgeRef, inP_, i03, vP2R, facePQ2R.begin(), true); - AppendPartialEdges(outR, wholeHalfedgeQ.H(), facePtrR.H(), edgesQ, - halfedgeRef.H(), inQ_, i30.H(), vQ2R.H(), + AppendPartialEdges(outR, wholeHalfedgeQ, facePtrR, edgesQ, + halfedgeRef, inQ_, i30, vQ2R, facePQ2R.begin() + inP_.NumTri(), false); - AppendNewEdges(outR, facePtrR.H(), edgesNew, halfedgeRef.H(), facePQ2R.H(), + AppendNewEdges(outR, facePtrR, edgesNew, halfedgeRef, facePQ2R, inP_.NumTri()); AppendWholeEdges(outR, facePtrR, halfedgeRef, inP_, wholeHalfedgeP, i03, vP2R, diff --git a/manifold/src/constructors.cpp b/manifold/src/constructors.cpp index 10e61457b..cdc1dfe35 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. @@ -212,9 +213,9 @@ Manifold Manifold::Extrude(Polygons crossSection, float height, int nDivisions, "scale values cannot be negative"); Manifold extrusion; ++nDivisions; - auto& vertPos = extrusion.pImpl_->vertPos_.H(); + auto& vertPos = extrusion.pImpl_->vertPos_; VecDH triVertsDH; - auto& triVerts = triVertsDH.H(); + auto& triVerts = triVertsDH; int nCrossSection = 0; bool isCone = scaleTop.x == 0.0 && scaleTop.y == 0.0; int idx = 0; @@ -289,9 +290,9 @@ Manifold Manifold::Revolve(const Polygons& crossSection, int circularSegments) { int nDivisions = circularSegments > 2 ? circularSegments : GetCircularSegments(radius); Manifold revoloid; - auto& vertPos = revoloid.pImpl_->vertPos_.H(); + auto& vertPos = revoloid.pImpl_->vertPos_; VecDH triVertsDH; - auto& triVerts = triVertsDH.H(); + auto& triVerts = triVertsDH; float dPhi = 360.0f / nDivisions; for (const auto& poly : crossSection) { int start = -1; @@ -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..0926db2f3 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,52 +122,50 @@ 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); - for (const int edge : flaggedEdges.H()) CollapseEdge(edge); + for (const int edge : flaggedEdges) CollapseEdge(edge); 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); - for (const int edge : flaggedEdges.H()) CollapseEdge(edge); + for (const int edge : flaggedEdges) CollapseEdge(edge); 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(); flaggedEdges.resize(numFlagged); - for (const int edge : flaggedEdges.H()) { + for (const int edge : flaggedEdges) { RecursiveEdgeSwap(edge); } } void Manifold::Impl::PairUp(int edge0, int edge1) { - VecH& halfedge = halfedge_.H(); - halfedge[edge0].pairedHalfedge = edge1; - halfedge[edge1].pairedHalfedge = edge0; + halfedge_[edge0].pairedHalfedge = edge1; + halfedge_[edge1].pairedHalfedge = edge0; } // Traverses CW around startEdge.endVert from startEdge to endEdge // (edgeEdge.endVert must == startEdge.endVert), updating each edge to point // to vert instead. void Manifold::Impl::UpdateVert(int vert, int startEdge, int endEdge) { - VecH& halfedge = halfedge_.H(); while (startEdge != endEdge) { - halfedge[startEdge].endVert = vert; + halfedge_[startEdge].endVert = vert; startEdge = NextHalfedge(startEdge); - halfedge[startEdge].startVert = vert; - startEdge = halfedge[startEdge].pairedHalfedge; + halfedge_[startEdge].startVert = vert; + startEdge = halfedge_[startEdge].pairedHalfedge; } } @@ -174,58 +173,52 @@ void Manifold::Impl::UpdateVert(int vert, int startEdge, int endEdge) { // instead we duplicate the two verts and attach the manifolds the other way // across this edge. void Manifold::Impl::FormLoop(int current, int end) { - VecH& halfedge = halfedge_.H(); - VecH& vertPos = vertPos_.H(); + int startVert = vertPos_.size(); + vertPos_.push_back(vertPos_[halfedge_[current].startVert]); + int endVert = vertPos_.size(); + vertPos_.push_back(vertPos_[halfedge_[current].endVert]); - int startVert = vertPos.size(); - vertPos.push_back(vertPos[halfedge[current].startVert]); - int endVert = vertPos.size(); - vertPos.push_back(vertPos[halfedge[current].endVert]); - - int oldMatch = halfedge[current].pairedHalfedge; - int newMatch = halfedge[end].pairedHalfedge; + int oldMatch = halfedge_[current].pairedHalfedge; + int newMatch = halfedge_[end].pairedHalfedge; UpdateVert(startVert, oldMatch, newMatch); UpdateVert(endVert, end, current); - halfedge[current].pairedHalfedge = newMatch; - halfedge[newMatch].pairedHalfedge = current; - halfedge[end].pairedHalfedge = oldMatch; - halfedge[oldMatch].pairedHalfedge = end; + halfedge_[current].pairedHalfedge = newMatch; + halfedge_[newMatch].pairedHalfedge = current; + halfedge_[end].pairedHalfedge = oldMatch; + halfedge_[oldMatch].pairedHalfedge = end; RemoveIfFolded(end); } void Manifold::Impl::CollapseTri(const glm::ivec3& triEdge) { - VecH& halfedge = halfedge_.H(); - int pair1 = halfedge[triEdge[1]].pairedHalfedge; - int pair2 = halfedge[triEdge[2]].pairedHalfedge; - halfedge[pair1].pairedHalfedge = pair2; - halfedge[pair2].pairedHalfedge = pair1; + int pair1 = halfedge_[triEdge[1]].pairedHalfedge; + int pair2 = halfedge_[triEdge[2]].pairedHalfedge; + halfedge_[pair1].pairedHalfedge = pair2; + halfedge_[pair2].pairedHalfedge = pair1; for (int i : {0, 1, 2}) { - halfedge[triEdge[i]] = {-1, -1, -1, -1}; + halfedge_[triEdge[i]] = {-1, -1, -1, -1}; } } void Manifold::Impl::RemoveIfFolded(int edge) { - VecH& halfedge = halfedge_.H(); - VecH& vertPos = vertPos_.H(); const glm::ivec3 tri0edge = TriOf(edge); - const glm::ivec3 tri1edge = TriOf(halfedge[edge].pairedHalfedge); - if (halfedge[tri0edge[1]].endVert == halfedge[tri1edge[1]].endVert) { + const glm::ivec3 tri1edge = TriOf(halfedge_[edge].pairedHalfedge); + if (halfedge_[tri0edge[1]].endVert == halfedge_[tri1edge[1]].endVert) { for (int i : {0, 1, 2}) { - vertPos[halfedge[tri0edge[i]].startVert] = glm::vec3(NAN); - halfedge[tri0edge[i]] = {-1, -1, -1, -1}; - halfedge[tri1edge[i]] = {-1, -1, -1, -1}; + vertPos_[halfedge_[tri0edge[i]].startVert] = glm::vec3(NAN); + halfedge_[tri0edge[i]] = {-1, -1, -1, -1}; + halfedge_[tri1edge[i]] = {-1, -1, -1, -1}; } } } void Manifold::Impl::CollapseEdge(const int edge) { - VecH& halfedge = halfedge_.H(); - VecH& vertPos = vertPos_.H(); - VecH& triNormal = faceNormal_.H(); - VecH& triBary = meshRelation_.triBary.H(); + VecDH& halfedge = halfedge_; + VecDH& vertPos = vertPos_; + VecDH& triNormal = faceNormal_; + VecDH& triBary = meshRelation_.triBary; const Halfedge toRemove = halfedge[edge]; if (toRemove.pairedHalfedge < 0) return; @@ -315,10 +308,10 @@ void Manifold::Impl::CollapseEdge(const int edge) { } void Manifold::Impl::RecursiveEdgeSwap(const int edge) { - const VecH& vertPos = vertPos_.H(); - VecH& halfedge = halfedge_.H(); - VecH& triNormal = faceNormal_.H(); - VecH& triBary = meshRelation_.triBary.H(); + const VecDH& vertPos = vertPos_; + VecDH& halfedge = halfedge_; + VecDH& triNormal = faceNormal_; + VecDH& triBary = meshRelation_.triBary; if (halfedge[edge].pairedHalfedge < 0) return; @@ -371,7 +364,7 @@ void Manifold::Impl::RecursiveEdgeSwap(const int edge) { const glm::vec3 uvw2 = a * uvw0 + (1 - a) * uvw1; // And assign it. const int newBary = meshRelation_.barycentric.size(); - meshRelation_.barycentric.H().push_back(uvw2); + meshRelation_.barycentric.push_back(uvw2); triBary[tri1].vertBary[perm1[0]] = newBary; triBary[tri0].vertBary[perm0[2]] = newBary; diff --git a/manifold/src/face_op.cpp b/manifold/src/face_op.cpp index cc061a492..fa5df3d2a 100644 --- a/manifold/src/face_op.cpp +++ b/manifold/src/face_op.cpp @@ -32,19 +32,22 @@ namespace manifold { void Manifold::Impl::Face2Tri(const VecDH& faceEdge, const VecDH& faceRef, const VecDH& halfedgeBary) { - VecDH triVertsOut; - VecDH triNormalOut; - - VecH& triVerts = triVertsOut.H(); - VecH& triNormal = triNormalOut.H(); - const VecH& vertPos = vertPos_.H(); - const VecH& faceEdgeH = faceEdge.H(); - const VecH& halfedge = halfedge_.H(); - const VecH& faceNormal = faceNormal_.H(); - meshRelation_.triBary.resize(0); + VecDH triVerts; + VecDH triNormal; + VecDH &triBary = meshRelation_.triBary; + triBary.resize(0); triVerts.reserve(faceEdge.size()); triNormal.reserve(faceEdge.size()); - meshRelation_.triBary.H().reserve(faceEdge.size()*3); + triBary.reserve(faceEdge.size()*3); + + const VecDH& vertPos = vertPos_; + const VecDH& faceEdgeH = faceEdge; + const VecDH& halfedge = halfedge_; + const VecDH& faceNormal = faceNormal_; + // meshRelation_.triBary.resize(0); + // std::vector triVerts; + // std::vector triNormal; + // std::vector triBary; for (int face = 0; face < faceEdgeH.size() - 1; ++face) { const int firstEdge = faceEdgeH[face]; @@ -79,10 +82,10 @@ void Manifold::Impl::Face2Tri(const VecDH& faceEdge, triVerts.push_back(tri); triNormal.push_back(normal); - meshRelation_.triBary.H().push_back(faceRef.H()[face]); + triBary.push_back(faceRef[face]); for (int k : {0, 1, 2}) { int index = linearSearch(mapping, tri[k]); - meshRelation_.triBary.H().back().vertBary[k] = halfedgeBary.H()[firstEdge + index]; + triBary.back().vertBary[k] = halfedgeBary[firstEdge + index]; } } else if (numEdge == 4) { // Pair of triangles int mapping[4] = {halfedge[firstEdge].startVert, @@ -131,10 +134,10 @@ void Manifold::Impl::Face2Tri(const VecDH& faceEdge, for (auto tri : { tri0, tri1 }) { triVerts.push_back(tri); triNormal.push_back(normal); - meshRelation_.triBary.H().push_back(faceRef.H()[face]); + triBary.push_back(faceRef[face]); for (int k : {0, 1, 2}) { int index = linearSearch(mapping, tri[k]); - meshRelation_.triBary.H().back().vertBary[k] = halfedgeBary.H()[firstEdge + index]; + triBary.back().vertBary[k] = halfedgeBary[firstEdge + index]; } } } else { // General triangulation @@ -142,7 +145,7 @@ void Manifold::Impl::Face2Tri(const VecDH& faceEdge, std::map vertBary; for (int j = firstEdge; j < lastEdge; ++j) - vertBary[halfedge[j].startVert] = halfedgeBary.H()[j]; + vertBary[halfedge[j].startVert] = halfedgeBary[j]; Polygons polys; try { @@ -160,16 +163,16 @@ void Manifold::Impl::Face2Tri(const VecDH& faceEdge, for (auto tri : newTris) { triVerts.push_back(tri); triNormal.push_back(normal); - meshRelation_.triBary.H().push_back(faceRef.H()[face]); + triBary.push_back(faceRef[face]); for (int k : {0, 1, 2}) { - meshRelation_.triBary.H().back().vertBary[k] = + triBary.back().vertBary[k] = vertBary[tri[k]]; } } } } - faceNormal_ = std::move(triNormalOut); - CreateAndFixHalfedges(triVertsOut); + faceNormal_ = std::move(triNormal); + CreateAndFixHalfedges(triVerts); } /** @@ -177,9 +180,9 @@ void Manifold::Impl::Face2Tri(const VecDH& faceEdge, * projection of the vertices. */ Polygons Manifold::Impl::Face2Polygons(int face, glm::mat3x2 projection, - const VecH& faceEdge) const { - const VecH& vertPos = vertPos_.H(); - const VecH& halfedge = halfedge_.H(); + const VecDH& faceEdge) const { + const VecDH& vertPos = vertPos_; + const VecDH& halfedge = halfedge_; const int firstEdge = faceEdge[face]; const int lastEdge = faceEdge[face + 1]; diff --git a/manifold/src/impl.cpp b/manifold/src/impl.cpp index 1bca96287..b35e5fef7 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_})); @@ -421,7 +421,7 @@ int Manifold::Impl::InitializeNewReference( graph.add_nodes(i); } for (int i = 0; i < face2face.size(); ++i) { - const thrust::pair edge = face2face.H()[i]; + const thrust::pair edge = face2face[i]; if (edge.first < 0) continue; graph.add_edge(edge.first, edge.second); } @@ -433,13 +433,13 @@ int Manifold::Impl::InitializeNewReference( for (int tri = 0; tri < NumTri(); ++tri) { const int comp = components[tri]; const int current = comp2tri[comp]; - if (current < 0 || triArea.H()[tri] > triArea.H()[current]) { + if (current < 0 || triArea[tri] > triArea[current]) { comp2tri[comp] = tri; - triArea.H()[comp] = triArea.H()[tri]; + triArea[comp] = triArea[tri]; } } - VecH& triBary = meshRelation_.triBary.H(); + VecDH& triBary = meshRelation_.triBary; std::map, int> triVert2bary; for (int tri = 0; tri < NumTri(); ++tri) { @@ -448,25 +448,25 @@ int Manifold::Impl::InitializeNewReference( glm::mat3 triPos; for (int i : {0, 1, 2}) { - const int vert = halfedge_.H()[3 * refTri + i].startVert; - triPos[i] = vertPos_.H()[vert]; + const int vert = halfedge_[3 * refTri + i].startVert; + triPos[i] = vertPos_[vert]; triVert2bary[{refTri, vert}] = i - 3; } glm::ivec3 vertBary; bool coplanar = true; for (int i : {0, 1, 2}) { - const int vert = halfedge_.H()[3 * tri + i].startVert; + const int vert = halfedge_[3 * tri + i].startVert; if (triVert2bary.find({refTri, vert}) == triVert2bary.end()) { const glm::vec3 uvw = - GetBarycentric(vertPos_.H()[vert], triPos, precision_); + GetBarycentric(vertPos_[vert], triPos, precision_); if (isnan(uvw[0])) { coplanar = false; triVert2bary[{refTri, vert}] = -4; break; } triVert2bary[{refTri, vert}] = meshRelation_.barycentric.size(); - meshRelation_.barycentric.H().push_back(uvw); + meshRelation_.barycentric.push_back(uvw); } const int bary = triVert2bary[{refTri, vert}]; if (bary < -3) { @@ -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/impl.h b/manifold/src/impl.h index d1bcf8af7..27c21b177 100644 --- a/manifold/src/impl.h +++ b/manifold/src/impl.h @@ -94,7 +94,7 @@ struct Manifold::Impl { void Face2Tri(const VecDH& faceEdge, const VecDH& faceRef, const VecDH& halfedgeBary); Polygons Face2Polygons(int face, glm::mat3x2 projection, - const VecH& faceEdge) const; + const VecDH& faceEdge) const; // edge_op.cu void SimplifyTopology(); diff --git a/manifold/src/manifold.cpp b/manifold/src/manifold.cpp index ee6bb7838..35f28796f 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; @@ -439,7 +440,7 @@ Manifold& Manifold::Transform(const glm::mat4x3& m) { */ Manifold& Manifold::Warp(std::function warpFunc) { pImpl_->ApplyTransform(); - thrust::for_each_n(pImpl_->vertPos_.begin(), NumVert(), warpFunc); + thrust::for_each_n(thrust::host, pImpl_->vertPos_.begin(), NumVert(), warpFunc); pImpl_->Update(); pImpl_->faceNormal_.resize(0); // force recalculation of triNormal pImpl_->CalculateNormals(); 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..c8c996bf5 100644 --- a/manifold/src/smoothing.cpp +++ b/manifold/src/smoothing.cpp @@ -356,14 +356,14 @@ 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()})); if (!sharpenedEdges.empty()) { - const VecH& halfedge = halfedge_.H(); - const VecH& triBary = meshRelation_.triBary.H(); + const VecDH& halfedge = halfedge_; + const VecDH& triBary = meshRelation_.triBary; // sharpenedEdges are referenced to the input Mesh, but the triangles have // been sorted in creating the Manifold, so the indices are converted using @@ -396,7 +396,7 @@ void Manifold::Impl::CreateTangents( {edge.second, edge.first}); } - VecH& tangent = halfedgeTangent_.H(); + VecDH& tangent = halfedgeTangent_; for (const auto& value : vertTangents) { const std::vector& vert = value.second; // Sharp edges that end are smooth at their terminal vert. @@ -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..3b18b065e 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); @@ -140,19 +141,19 @@ class SparseIndices { "Different number of values than indicies!"); size_t size = pqEnd - pqBegin; VecDH result(size); - VecDH found(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; } void Dump() const { - const auto& p = Get(0).H(); - const auto& q = Get(1).H(); + const auto& p = Get(0); + const auto& q = Get(1); std::cout << "SparseIndices = " << std::endl; for (int i = 0; i < size(); ++i) { std::cout << i << ", p = " << p[i] << ", q = " << q[i] << std::endl; diff --git a/utilities/include/structs.h b/utilities/include/structs.h index c84d835f4..78fc9a999 100644 --- a/utilities/include/structs.h +++ b/utilities/include/structs.h @@ -414,17 +414,17 @@ struct Box { } }; -/** - * Print the contents of this vector to standard output. - */ -template -void Dump(const std::vector& vec) { - std::cout << "Vec = " << std::endl; - for (int i = 0; i < vec.size(); ++i) { - std::cout << i << ", " << vec[i] << ", " << std::endl; - } - std::cout << std::endl; -} +// /** +// * Print the contents of this vector to standard output. +// */ +// template +// void Dump(const std::vector& vec) { +// std::cout << "Vec = " << std::endl; +// for (int i = 0; i < vec.size(); ++i) { +// std::cout << i << ", " << vec[i] << ", " << std::endl; +// } +// std::cout << std::endl; +// } inline std::ostream& operator<<(std::ostream& stream, const Box& box) { return stream << "min: " << box.min.x << ", " << box.min.y << ", " @@ -470,4 +470,4 @@ inline std::ostream& operator<<(std::ostream& stream, const BaryRef& ref) { /** @} */ } // namespace manifold -#undef HOST_DEVICE \ No newline at end of file +#undef HOST_DEVICE diff --git a/utilities/include/vec_dh.h b/utilities/include/vec_dh.h index 38191b72f..81efff582 100644 --- a/utilities/include/vec_dh.h +++ b/utilities/include/vec_dh.h @@ -13,8 +13,8 @@ // limitations under the License. #pragma once -#include -#include +#include +#include namespace manifold { @@ -22,103 +22,196 @@ namespace manifold { * @{ */ template -using VecH = thrust::host_vector; - -template -void Dump(const VecH& vec) { - std::cout << "VecDH = " << std::endl; +void Dump(const std::vector& vec) { + std::cout << "Vec = " << std::endl; for (int i = 0; i < vec.size(); ++i) { std::cout << i << ", " << vec[i] << ", " << std::endl; } std::cout << std::endl; } +/* + * Host and device vector implementation. This uses `thrust::universal_vector` for + * storage, so data can be moved by the hardware on demand, allows using more + * memory than the available GPU memory, reduce memory overhead and provide + * speedup due to less synchronization. + * + * Due to https://github.com/NVIDIA/thrust/issues/1690 , `push_back` operations + * on universal vectors are *VERY* slow, so a `std::vector` is used as a cache. + * The cache will be created when we perform `push_back` or `reserve` operations + * on the `VecDH`, and destroyed when we try to access device iterator/pointer. + * For better performance, please avoid interspersing `push_back` between device + * memory accesses, as that will cause repeated synchronization and hurts + * performance. + * Note that it is *NOT SAFE* to first obtain a host(device) pointer, perform + * some device(host) modification, and then read the host(device) pointer again + * (on the same vector). The memory will be inconsistent in that case. + */ template class VecDH { public: - VecDH() {} + VecDH() { + impl_ = thrust::universal_vector(); + } VecDH(int size, T val = T()) { - device_.resize(size, val); - host_valid_ = false; + impl_.resize(size, val); + implModified = true; + cacheModified = false; } VecDH(const std::vector& vec) { - host_ = vec; - device_valid_ = false; + cache = vec; + cacheModified = true; + implModified = false; + } + + VecDH(const VecDH& other) { + if (!other.cacheModified) { + if (other.impl_.size() > 0) + impl_ = other.impl_; + else + impl_.clear(); + implModified = true; + cacheModified = false; + } else { + cache = other.cache; + implModified = false; + cacheModified = true; + } + } + + VecDH(VecDH&& other) { + if (!other.cacheModified) { + if (other.impl_.size() > 0) + impl_ = std::move(other.impl_); + else + impl_.clear(); + implModified = true; + cacheModified = false; + } else { + cache = std::move(other.cache); + cacheModified = true; + implModified = false; + } } - int size() const { return device_valid_ ? device_.size() : host_.size(); } + VecDH& operator=(const VecDH& other) { + if (!other.cacheModified) { + if (other.impl_.size() > 0) + impl_ = other.impl_; + else + impl_.clear(); + implModified = true; + cacheModified = false; + } else { + cache = other.cache; + cacheModified = true; + implModified = false; + } + return *this; + } + + VecDH& operator=(VecDH&& other) { + if (!other.cacheModified) { + if (other.impl_.size() > 0) + impl_ = std::move(other.impl_); + else + impl_.clear(); + implModified = true; + cacheModified = false; + } else { + cache = std::move(other.cache); + cacheModified = true; + implModified = false; + } + return *this; + } + + int size() const { + if (!cacheModified) + return impl_.size(); + return cache.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(); + if (cacheModified) { + cache.resize(newSize, val); + if (shrink) cache.shrink_to_fit(); + } else { + impl_.resize(newSize, val); + if (shrink) impl_.shrink_to_fit(); + implModified = true; } } 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_); + if (!cacheModified && !other.cacheModified) { + implModified = true; + other.implModified = true; + impl_.swap(other.impl_); + } else { + syncCache(); + other.syncCache(); + cacheModified = true; + implModified = false; + other.cacheModified = true; + other.implModified = false; + cache.swap(other.cache); + } } - 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(); + syncImpl(); + implModified = true; + return impl_.begin(); } IterH end() { - RefreshHost(); - device_valid_ = false; - return host_.end(); + syncImpl(); + implModified = true; + return impl_.end(); } IterHc cbegin() const { - RefreshHost(); - return host_.cbegin(); + syncImpl(); + return impl_.cbegin(); } IterHc cend() const { - RefreshHost(); - return host_.cend(); + syncImpl(); + return impl_.cend(); } IterHc begin() const { return cbegin(); } IterHc end() const { return cend(); } IterD beginD() { - RefreshDevice(); - host_valid_ = false; - return device_.begin(); + syncImpl(); + implModified = true; + return impl_.begin(); } IterD endD() { - RefreshDevice(); - host_valid_ = false; - return device_.end(); + syncImpl(); + implModified = true; + return impl_.end(); } IterDc cbeginD() const { - RefreshDevice(); - return device_.cbegin(); + syncImpl(); + return impl_.cbegin(); } IterDc cendD() const { - RefreshDevice(); - return device_.cend(); + syncImpl(); + return impl_.cend(); } IterDc beginD() const { return cbeginD(); } @@ -126,65 +219,112 @@ class VecDH { T* ptrD() { if (size() == 0) return nullptr; - RefreshDevice(); - host_valid_ = false; - return device_.data().get(); + syncImpl(); + implModified = true; + return impl_.data().get(); } const T* cptrD() const { if (size() == 0) return nullptr; - RefreshDevice(); - return device_.data().get(); + syncImpl(); + return impl_.data().get(); } const T* ptrD() const { return cptrD(); } T* ptrH() { if (size() == 0) return nullptr; - RefreshHost(); - device_valid_ = false; - return host_.data(); + if (cacheModified) { + return cache.data(); + } else { + implModified = true; + return impl_.data().get(); + } } const T* cptrH() const { if (size() == 0) return nullptr; - RefreshHost(); - return host_.data(); + if (cacheModified) { + return cache.data(); + } else { + return impl_.data().get(); + } } const T* ptrH() const { return cptrH(); } - const VecH& H() const { - RefreshHost(); - return host_; + T& operator[](int i) { + if (!cacheModified) { + implModified = true; + return impl_[i]; + } else { + cacheModified = true; + return cache[i]; + } + } + + const T& operator[](int i) const { + if (!cacheModified) { + return impl_[i]; + } else { + return cache[i]; + } + } + + T& back() { + if (!cacheModified) { + implModified = true; + return impl_.back(); + } else { + return cache.back(); + } } - VecH& H() { - RefreshHost(); - device_valid_ = false; - return host_; + const T& back() const { + if (!cacheModified) { + return impl_.back(); + } else { + return cache.back(); + } } - void Dump() const { manifold::Dump(H()); } + void push_back(const T& val) { + syncCache(); + cacheModified = true; + cache.push_back(val); + } + + void reserve(int n) { + syncCache(); + cacheModified = true; + cache.reserve(n); + } + + void Dump() const { + syncCache(); + manifold::Dump(cache); + } 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; + mutable thrust::universal_vector impl_; + + mutable bool implModified = false; + mutable bool cacheModified = false; + mutable std::vector cache; + + void syncImpl() const { + if (cacheModified) { + impl_ = cache; + cache.clear(); } + cacheModified = false; } - void RefreshDevice() const { - if (!device_valid_) { - device_ = host_; - device_valid_ = true; + void syncCache() const { + if (implModified) { + cache = std::vector(impl_.begin(), impl_.end()); } + implModified = false; } }; @@ -201,4 +341,4 @@ class VecD { const int size_; }; /** @} */ -} // namespace manifold \ No newline at end of file +} // namespace manifold From dbbcdbbd49e4594429c60da66f7046c2f4d3a836 Mon Sep 17 00:00:00 2001 From: pca006132 Date: Mon, 16 May 2022 14:30:19 +0800 Subject: [PATCH 2/3] vec_dh cleanup --- collider/src/collider.cpp | 10 ++-- manifold/src/boolean3.cpp | 29 +++++----- manifold/src/boolean_result.cpp | 63 +++++++++++----------- manifold/src/constructors.cpp | 47 ++++++++-------- manifold/src/edge_op.cpp | 95 +++++++++++++++------------------ manifold/src/face_op.cpp | 83 +++++++++++++--------------- manifold/src/impl.cpp | 31 ++++++----- manifold/src/manifold.cpp | 7 ++- manifold/src/properties.cpp | 19 ++++--- manifold/src/shared.h | 11 ++-- manifold/src/smoothing.cpp | 27 +++++----- manifold/src/sort.cpp | 54 +++++++++---------- utilities/include/sparse.h | 48 ++++++++--------- utilities/include/structs.h | 22 ++++---- utilities/include/vec_dh.h | 57 ++++---------------- 15 files changed, 271 insertions(+), 332 deletions(-) diff --git a/collider/src/collider.cpp b/collider/src/collider.cpp index ff8eadd48..f02227ad8 100644 --- a/collider/src/collider.cpp +++ b/collider/src/collider.cpp @@ -274,7 +274,7 @@ SparseIndices Collider::Collisions(const VecDH& querriesIn) const { VecDH nOverlapsD(1, 0); // calculate Bounding Box overlaps thrust::for_each_n( - thrust::device, zip(querriesIn.cbeginD(), countAt(0)), querriesIn.size(), + thrust::device, zip(querriesIn.cbegin(), countAt(0)), querriesIn.size(), FindCollisions({querryTri.ptrDpq(), nOverlapsD.ptrD(), maxOverlaps, nodeBBox_.ptrD(), internalChildren_.ptrD()})); nOverlaps = nOverlapsD[0]; @@ -303,12 +303,12 @@ void Collider::UpdateBoxes(const VecDH& leafBB) { ALWAYS_ASSERT(leafBB.size() == NumLeaves(), userErr, "must have the same number of updated boxes as original"); // copy in leaf node Boxs - strided_range::IterD> leaves(nodeBBox_.beginD(), nodeBBox_.endD(), + strided_range::Iter> leaves(nodeBBox_.begin(), nodeBBox_.end(), 2); - thrust::copy(thrust::device, leafBB.cbeginD(), leafBB.cendD(), leaves.begin()); + thrust::copy(thrust::device, leafBB.cbegin(), leafBB.cend(), leaves.begin()); // create global counters VecDH counter(NumInternal()); - thrust::fill(thrust::device, counter.beginD(), counter.endD(), 0); + thrust::fill(thrust::device, counter.begin(), counter.end(), 0); // kernel over leaves to save internal Boxs thrust::for_each_n( thrust::device, countAt(0), NumLeaves(), @@ -330,7 +330,7 @@ bool Collider::Transform(glm::mat4x3 transform) { if (count != 2) axisAligned = false; } if (axisAligned) { - thrust::for_each(thrust::device, nodeBBox_.beginD(), nodeBBox_.endD(), + thrust::for_each(thrust::device, nodeBBox_.begin(), nodeBBox_.end(), TransformBox({transform})); } return axisAligned; diff --git a/manifold/src/boolean3.cpp b/manifold/src/boolean3.cpp index 248a4ac18..5bd239332 100644 --- a/manifold/src/boolean3.cpp +++ b/manifold/src/boolean3.cpp @@ -14,7 +14,6 @@ #include "boolean3.h" #include -#include // TODO: make this runtime configurable for quicker debug constexpr bool kVerbose = false; @@ -86,12 +85,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(thrust::device, zip(countAt(0), p1q2.beginD(0), p1q2.beginD(1)), + thrust::for_each_n(thrust::device, zip(countAt(0), p1q2.begin(0), p1q2.begin(1)), p1q2.size(), CopyFaceEdges({p1q1.ptrDpq(), inQ.halfedge_.cptrD()})); p1q1.SwapPQ(); - thrust::for_each_n(thrust::device, zip(countAt(p1q2.size()), p2q1.beginD(1), p2q1.beginD(0)), + thrust::for_each_n(thrust::device, zip(countAt(p1q2.size()), p2q1.begin(1), p2q1.begin(0)), p2q1.size(), CopyFaceEdges({p1q1.ptrDpq(), inP.halfedge_.cptrD()})); p1q1.SwapPQ(); @@ -247,7 +246,7 @@ std::tuple, VecDH> Shadow11(SparseIndices &p1q1, VecDH xyzz11(p1q1.size()); thrust::for_each_n( - thrust::device, zip(xyzz11.beginD(), s11.beginD(), p1q1.beginD(0), p1q1.beginD(1)), + thrust::device, zip(xyzz11.begin(), s11.begin(), p1q1.begin(0), p1q1.begin(1)), p1q1.size(), Kernel11({inP.vertPos_.cptrD(), inQ.vertPos_.cptrD(), inP.halfedge_.cptrD(), inQ.halfedge_.cptrD(), expandP, @@ -344,8 +343,8 @@ std::tuple, VecDH> Shadow02(const Manifold::Impl &inP, auto vertNormalP = forward ? inP.vertNormal_.cptrD() : inQ.vertNormal_.cptrD(); thrust::for_each_n( - thrust::device, zip(s02.beginD(), z02.beginD(), p0q2.beginD(!forward), - p0q2.beginD(forward)), + thrust::device, zip(s02.begin(), z02.begin(), p0q2.begin(!forward), + p0q2.begin(forward)), p0q2.size(), Kernel02({inP.vertPos_.cptrD(), inQ.halfedge_.cptrD(), inQ.vertPos_.cptrD(), forward, expandP, vertNormalP})); @@ -454,8 +453,8 @@ std::tuple, VecDH> Intersect12( VecDH v12(p1q2.size()); thrust::for_each_n( - thrust::device, zip(x12.beginD(), v12.beginD(), p1q2.beginD(!forward), - p1q2.beginD(forward)), + thrust::device, zip(x12.begin(), v12.begin(), p1q2.begin(!forward), + p1q2.begin(forward)), p1q2.size(), Kernel12({p0q2.ptrDpq(), s02.ptrD(), z02.cptrD(), p0q2.size(), p1q1.ptrDpq(), s11.ptrD(), xyzz11.cptrD(), p1q1.size(), @@ -472,19 +471,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(thrust::device, p0q2.beginD(reverse), p0q2.endD(reverse))) - thrust::sort_by_key(thrust::device, p0q2.beginD(reverse), p0q2.endD(reverse), s02.beginD()); + if (!thrust::is_sorted(thrust::device, p0q2.begin(reverse), p0q2.end(reverse))) + thrust::sort_by_key(thrust::device, p0q2.begin(reverse), p0q2.end(reverse), s02.begin()); VecDH w03val(w03.size()); VecDH w03vert(w03.size()); // sum known s02 values into w03 (winding number) auto endPair = - thrust::reduce_by_key(thrust::device, p0q2.beginD(reverse), p0q2.endD(reverse), - s02.beginD(), w03vert.beginD(), w03val.beginD()); - thrust::scatter(thrust::device, w03val.beginD(), endPair.second, w03vert.beginD(), - w03.beginD()); + thrust::reduce_by_key(thrust::device, p0q2.begin(reverse), p0q2.end(reverse), + s02.begin(), w03vert.begin(), w03val.begin()); + thrust::scatter(thrust::device, w03val.begin(), endPair.second, w03vert.begin(), + w03.begin()); if (reverse) - thrust::transform(thrust::device, w03.beginD(), w03.endD(), w03.beginD(), + thrust::transform(thrust::device, w03.begin(), w03.end(), w03.begin(), thrust::negate()); return w03; }; diff --git a/manifold/src/boolean_result.cpp b/manifold/src/boolean_result.cpp index 3e8cc17cd..807727fc0 100644 --- a/manifold/src/boolean_result.cpp +++ b/manifold/src/boolean_result.cpp @@ -14,7 +14,6 @@ #include #include -#include #include "boolean3.h" #include "polygon.h" @@ -84,46 +83,46 @@ std::tuple, VecDH> SizeOutput( auto sidesPerFaceP = sidesPerFacePQ.ptrD(); auto sidesPerFaceQ = sidesPerFacePQ.ptrD() + inP.NumTri(); - thrust::for_each(thrust::device, inP.halfedge_.beginD(), inP.halfedge_.endD(), + thrust::for_each(thrust::device, inP.halfedge_.begin(), inP.halfedge_.end(), CountVerts({sidesPerFaceP, i03.cptrD()})); - thrust::for_each(thrust::device, inQ.halfedge_.beginD(), inQ.halfedge_.endD(), + thrust::for_each(thrust::device, inQ.halfedge_.begin(), inQ.halfedge_.end(), CountVerts({sidesPerFaceQ, i30.cptrD()})); thrust::for_each_n( - thrust::device, zip(p1q2.beginD(0), p1q2.beginD(1), i12.beginD()), i12.size(), + thrust::device, zip(p1q2.begin(0), p1q2.begin(1), i12.begin()), i12.size(), CountNewVerts({sidesPerFaceP, sidesPerFaceQ, inP.halfedge_.cptrD()})); thrust::for_each_n( - thrust::device, zip(p2q1.beginD(1), p2q1.beginD(0), i21.beginD()), i21.size(), + thrust::device, zip(p2q1.begin(1), p2q1.begin(0), i21.begin()), 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::make_transform_iterator(sidesPerFacePQ.begin(), NotZero()); thrust::inclusive_scan(thrust::device, keepFace, keepFace + sidesPerFacePQ.size(), - facePQ2R.beginD() + 1); + facePQ2R.begin() + 1); int numFaceR = facePQ2R.back(); facePQ2R.resize(inP.NumTri() + inQ.NumTri()); outR.faceNormal_.resize(numFaceR); - auto next = thrust::copy_if(thrust::device, inP.faceNormal_.beginD(), inP.faceNormal_.endD(), - keepFace, outR.faceNormal_.beginD(), + auto next = thrust::copy_if(thrust::device, inP.faceNormal_.begin(), inP.faceNormal_.end(), + keepFace, outR.faceNormal_.begin(), thrust::identity()); if (invertQ) { - auto start = thrust::make_transform_iterator(inQ.faceNormal_.beginD(), + auto start = thrust::make_transform_iterator(inQ.faceNormal_.begin(), thrust::negate()); - auto end = thrust::make_transform_iterator(inQ.faceNormal_.endD(), + auto end = thrust::make_transform_iterator(inQ.faceNormal_.end(), thrust::negate()); thrust::copy_if(thrust::device, start, end, keepFace + inP.NumTri(), next, thrust::identity()); } else { - thrust::copy_if(thrust::device, inQ.faceNormal_.beginD(), inQ.faceNormal_.endD(), + thrust::copy_if(thrust::device, inQ.faceNormal_.begin(), inQ.faceNormal_.end(), keepFace + inP.NumTri(), next, thrust::identity()); } auto newEnd = - thrust::remove(thrust::device, sidesPerFacePQ.beginD(), sidesPerFacePQ.endD(), 0); - VecDH faceEdge(newEnd - sidesPerFacePQ.beginD() + 1); - thrust::inclusive_scan(thrust::device, sidesPerFacePQ.beginD(), newEnd, - faceEdge.beginD() + 1); + thrust::remove(thrust::device, sidesPerFacePQ.begin(), sidesPerFacePQ.end(), 0); + VecDH faceEdge(newEnd - sidesPerFacePQ.begin() + 1); + thrust::inclusive_scan(thrust::device, sidesPerFacePQ.begin(), newEnd, + faceEdge.begin() + 1); outR.halfedge_.resize(faceEdge.back()); return std::make_tuple(faceEdge, facePQ2R); @@ -212,7 +211,7 @@ void AppendPartialEdges(Manifold::Impl &outR, VecDH &wholeHalfedgeP, std::map> &edgesP, VecDH &halfedgeRef, const Manifold::Impl &inP, const VecDH &i03, const VecDH &vP2R, - const VecDH::IterHc faceP2R, + const VecDH::IterC 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), @@ -407,7 +406,7 @@ void AppendWholeEdges(Manifold::Impl &outR, VecDH &facePtrR, const VecDH &vP2R, const int *faceP2R, bool forward) { thrust::for_each_n( - thrust::device, zip(wholeHalfedgeP.beginD(), inP.halfedge_.beginD(), countAt(0)), + thrust::device, zip(wholeHalfedgeP.begin(), inP.halfedge_.begin(), countAt(0)), inP.halfedge_.size(), DuplicateHalfedges({outR.halfedge_.ptrD(), halfedgeRef.ptrD(), facePtrR.ptrD(), inP.halfedge_.cptrD(), i03.cptrD(), @@ -483,8 +482,8 @@ std::pair, VecDH> CalculateMeshRelation( VecDH halfedgeBary(halfedgeRef.size()); VecDH idx(1, 0); thrust::for_each_n( - thrust::device, zip(halfedgeBary.beginD(), halfedgeRef.beginD(), - outR.halfedge_.cbeginD()), + thrust::device, zip(halfedgeBary.begin(), halfedgeRef.begin(), + outR.halfedge_.cbegin()), halfedgeRef.size(), CreateBarycentric( {outR.meshRelation_.barycentric.ptrD(), faceRef.ptrD(), idx.ptrD(), @@ -551,25 +550,25 @@ Manifold::Impl Boolean3::Result(Manifold::OpType op) const { VecDH i21(x21_.size()); VecDH i03(w03_.size()); VecDH i30(w30_.size()); - 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); + thrust::transform(thrust::device, x12_.begin(), x12_.end(), i12.begin(), c3 * _1); + thrust::transform(thrust::device, x21_.begin(), x21_.end(), i21.begin(), c3 * _1); + thrust::transform(thrust::device, w03_.begin(), w03_.end(), i03.begin(), c1 + c3 * _1); + thrust::transform(thrust::device, w30_.begin(), w30_.end(), i30.begin(), c2 + c3 * _1); VecDH vP2R(inP_.NumVert()); - thrust::exclusive_scan(thrust::device, i03.beginD(), i03.endD(), vP2R.beginD(), 0, AbsSum()); + thrust::exclusive_scan(thrust::device, i03.begin(), i03.end(), vP2R.begin(), 0, AbsSum()); int numVertR = AbsSum()(vP2R.back(), i03.back()); const int nPv = numVertR; VecDH vQ2R(inQ_.NumVert()); - thrust::exclusive_scan(thrust::device, i30.beginD(), i30.endD(), vQ2R.beginD(), numVertR, + thrust::exclusive_scan(thrust::device, i30.begin(), i30.end(), vQ2R.begin(), numVertR, AbsSum()); numVertR = AbsSum()(vQ2R.back(), i30.back()); const int nQv = numVertR - nPv; VecDH v12R(v12_.size()); if (v12_.size() > 0) { - thrust::exclusive_scan(thrust::device, i12.beginD(), i12.endD(), v12R.beginD(), numVertR, + thrust::exclusive_scan(thrust::device, i12.begin(), i12.end(), v12R.begin(), numVertR, AbsSum()); numVertR = AbsSum()(v12R.back(), i12.back()); } @@ -577,7 +576,7 @@ Manifold::Impl Boolean3::Result(Manifold::OpType op) const { VecDH v21R(v21_.size()); if (v21_.size() > 0) { - thrust::exclusive_scan(thrust::device, i21.beginD(), i21.endD(), v21R.beginD(), numVertR, + thrust::exclusive_scan(thrust::device, i21.begin(), i21.end(), v21R.begin(), numVertR, AbsSum()); numVertR = AbsSum()(v21R.back(), i21.back()); } @@ -593,14 +592,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(thrust::device, zip(i03.beginD(), vP2R.beginD(), inP_.vertPos_.beginD()), + thrust::for_each_n(thrust::device, zip(i03.begin(), vP2R.begin(), inP_.vertPos_.begin()), inP_.NumVert(), DuplicateVerts({outR.vertPos_.ptrD()})); - thrust::for_each_n(thrust::device, zip(i30.beginD(), vQ2R.beginD(), inQ_.vertPos_.beginD()), + thrust::for_each_n(thrust::device, zip(i30.begin(), vQ2R.begin(), inQ_.vertPos_.begin()), inQ_.NumVert(), DuplicateVerts({outR.vertPos_.ptrD()})); // New vertices created from intersections: - thrust::for_each_n(thrust::device, zip(i12.beginD(), v12R.beginD(), v12_.beginD()), + thrust::for_each_n(thrust::device, zip(i12.begin(), v12R.begin(), v12_.begin()), i12.size(), DuplicateVerts({outR.vertPos_.ptrD()})); - thrust::for_each_n(thrust::device, zip(i21.beginD(), v21R.beginD(), v21_.beginD()), + thrust::for_each_n(thrust::device, zip(i21.begin(), v21R.begin(), v21_.begin()), i21.size(), DuplicateVerts({outR.vertPos_.ptrD()})); if (kVerbose) { diff --git a/manifold/src/constructors.cpp b/manifold/src/constructors.cpp index cdc1dfe35..e0be1550a 100644 --- a/manifold/src/constructors.cpp +++ b/manifold/src/constructors.cpp @@ -13,7 +13,6 @@ // limitations under the License. #include -#include #include "graph.h" #include "impl.h" @@ -184,7 +183,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(thrust::device, sphere.pImpl_->vertPos_.beginD(), sphere.NumVert(), + thrust::for_each_n(thrust::device, sphere.pImpl_->vertPos_.begin(), sphere.NumVert(), ToSphere({radius})); sphere.pImpl_->Finish(); // Ignore preceding octahedron. @@ -404,21 +403,21 @@ Manifold Manifold::Compose(const std::vector& manifolds) { const Impl& impl = *(manifold.pImpl_); impl.ApplyTransform(); - thrust::copy(thrust::device, impl.vertPos_.beginD(), impl.vertPos_.endD(), - combined.vertPos_.beginD() + nextVert); - thrust::copy(thrust::device, impl.faceNormal_.beginD(), impl.faceNormal_.endD(), - combined.faceNormal_.beginD() + nextTri); - thrust::copy(thrust::device, impl.halfedgeTangent_.beginD(), impl.halfedgeTangent_.endD(), - combined.halfedgeTangent_.beginD() + nextEdge); - thrust::copy(thrust::device, impl.meshRelation_.barycentric.beginD(), - impl.meshRelation_.barycentric.endD(), - combined.meshRelation_.barycentric.beginD() + nextBary); - thrust::transform(thrust::device, impl.meshRelation_.triBary.beginD(), - impl.meshRelation_.triBary.endD(), - combined.meshRelation_.triBary.beginD() + nextTri, + thrust::copy(thrust::device, impl.vertPos_.begin(), impl.vertPos_.end(), + combined.vertPos_.begin() + nextVert); + thrust::copy(thrust::device, impl.faceNormal_.begin(), impl.faceNormal_.end(), + combined.faceNormal_.begin() + nextTri); + thrust::copy(thrust::device, impl.halfedgeTangent_.begin(), impl.halfedgeTangent_.end(), + combined.halfedgeTangent_.begin() + nextEdge); + thrust::copy(thrust::device, impl.meshRelation_.barycentric.begin(), + impl.meshRelation_.barycentric.end(), + combined.meshRelation_.barycentric.begin() + nextBary); + thrust::transform(thrust::device, impl.meshRelation_.triBary.begin(), + impl.meshRelation_.triBary.end(), + combined.meshRelation_.triBary.begin() + nextTri, UpdateTriBary({nextBary})); - thrust::transform(thrust::device, impl.halfedge_.beginD(), impl.halfedge_.endD(), - combined.halfedge_.beginD() + nextEdge, + thrust::transform(thrust::device, impl.halfedge_.begin(), impl.halfedge_.end(), + combined.halfedge_.begin() + nextEdge, UpdateHalfedge({nextVert, nextEdge, nextTri})); nextVert += manifold.NumVert(); @@ -462,22 +461,22 @@ std::vector Manifold::Decompose() const { VecDH vertNew2Old(NumVert()); int nVert = thrust::copy_if( - thrust::device, zip(pImpl_->vertPos_.beginD(), countAt(0)), - zip(pImpl_->vertPos_.endD(), countAt(NumVert())), - vertLabel.beginD(), - zip(meshes[i].pImpl_->vertPos_.beginD(), vertNew2Old.beginD()), + thrust::device, zip(pImpl_->vertPos_.begin(), countAt(0)), + zip(pImpl_->vertPos_.end(), countAt(NumVert())), + vertLabel.begin(), + zip(meshes[i].pImpl_->vertPos_.begin(), vertNew2Old.begin()), Equals({i})) - - zip(meshes[i].pImpl_->vertPos_.beginD(), countAt(0)); + zip(meshes[i].pImpl_->vertPos_.begin(), countAt(0)); meshes[i].pImpl_->vertPos_.resize(nVert); VecDH faceNew2Old(NumTri()); - thrust::sequence(thrust::device, faceNew2Old.beginD(), faceNew2Old.endD()); + thrust::sequence(thrust::device, faceNew2Old.begin(), faceNew2Old.end()); int nFace = thrust::remove_if( - thrust::device, faceNew2Old.beginD(), faceNew2Old.endD(), + thrust::device, faceNew2Old.begin(), faceNew2Old.end(), RemoveFace({pImpl_->halfedge_.cptrD(), vertLabel.cptrD(), i})) - - faceNew2Old.beginD(); + faceNew2Old.begin(); faceNew2Old.resize(nFace); meshes[i].pImpl_->GatherFaces(*pImpl_, faceNew2Old); diff --git a/manifold/src/edge_op.cpp b/manifold/src/edge_op.cpp index 0926db2f3..eecf01568 100644 --- a/manifold/src/edge_op.cpp +++ b/manifold/src/edge_op.cpp @@ -13,7 +13,6 @@ // limitations under the License. #include "impl.h" -#include namespace { using namespace manifold; @@ -122,9 +121,9 @@ void Manifold::Impl::SimplifyTopology() { VecDH flaggedEdges(halfedge_.size()); int numFlagged = thrust::copy_if( - thrust::device, countAt(0), countAt(halfedge_.size()), flaggedEdges.beginD(), + thrust::device, countAt(0), countAt(halfedge_.size()), flaggedEdges.begin(), ShortEdge({halfedge_.cptrD(), vertPos_.cptrD(), precision_})) - - flaggedEdges.beginD(); + flaggedEdges.begin(); flaggedEdges.resize(numFlagged); for (const int edge : flaggedEdges) CollapseEdge(edge); @@ -132,19 +131,19 @@ void Manifold::Impl::SimplifyTopology() { flaggedEdges.resize(halfedge_.size()); numFlagged = thrust::copy_if( - thrust::device, countAt(0), countAt(halfedge_.size()), flaggedEdges.beginD(), + thrust::device, countAt(0), countAt(halfedge_.size()), flaggedEdges.begin(), FlagEdge({halfedge_.cptrD(), meshRelation_.triBary.cptrD()})) - - flaggedEdges.beginD(); + flaggedEdges.begin(); flaggedEdges.resize(numFlagged); for (const int edge : flaggedEdges) CollapseEdge(edge); flaggedEdges.resize(halfedge_.size()); numFlagged = thrust::copy_if( - thrust::device, countAt(0), countAt(halfedge_.size()), flaggedEdges.beginD(), + thrust::device, countAt(0), countAt(halfedge_.size()), flaggedEdges.begin(), SwappableEdge({halfedge_.cptrD(), vertPos_.cptrD(), faceNormal_.cptrD(), precision_})) - - flaggedEdges.beginD(); + flaggedEdges.begin(); flaggedEdges.resize(numFlagged); for (const int edge : flaggedEdges) { @@ -215,42 +214,39 @@ void Manifold::Impl::RemoveIfFolded(int edge) { } void Manifold::Impl::CollapseEdge(const int edge) { - VecDH& halfedge = halfedge_; - VecDH& vertPos = vertPos_; - VecDH& triNormal = faceNormal_; VecDH& triBary = meshRelation_.triBary; - const Halfedge toRemove = halfedge[edge]; + const Halfedge toRemove = halfedge_[edge]; if (toRemove.pairedHalfedge < 0) return; const int endVert = toRemove.endVert; const glm::ivec3 tri0edge = TriOf(edge); const glm::ivec3 tri1edge = TriOf(toRemove.pairedHalfedge); - const glm::vec3 pNew = vertPos[endVert]; - const glm::vec3 pOld = vertPos[toRemove.startVert]; + const glm::vec3 pNew = vertPos_[endVert]; + const glm::vec3 pOld = vertPos_[toRemove.startVert]; const glm::vec3 delta = pNew - pOld; const bool shortEdge = glm::dot(delta, delta) < precision_ * precision_; std::vector edges; // Orbit endVert - int current = halfedge[tri0edge[1]].pairedHalfedge; + int current = halfedge_[tri0edge[1]].pairedHalfedge; while (current != tri1edge[2]) { current = NextHalfedge(current); edges.push_back(current); - current = halfedge[current].pairedHalfedge; + current = halfedge_[current].pairedHalfedge; } // Orbit startVert - int start = halfedge[tri1edge[1]].pairedHalfedge; + int start = halfedge_[tri1edge[1]].pairedHalfedge; const BaryRef ref0 = triBary[edge / 3]; const BaryRef ref1 = triBary[toRemove.pairedHalfedge / 3]; if (!shortEdge) { current = start; - glm::vec3 pLast = vertPos[halfedge[tri1edge[1]].endVert]; + glm::vec3 pLast = vertPos_[halfedge_[tri1edge[1]].endVert]; while (current != tri0edge[2]) { current = NextHalfedge(current); - glm::vec3 pNext = vertPos[halfedge[current].endVert]; + glm::vec3 pNext = vertPos_[halfedge_[current].endVert]; const int tri = current / 3; const BaryRef ref = triBary[tri]; // Don't collapse if the edge is not redundant (this may have changed due @@ -260,18 +256,18 @@ void Manifold::Impl::CollapseEdge(const int edge) { return; // Don't collapse edge if it would cause a triangle to invert. - const glm::mat3x2 projection = GetAxisAlignedProjection(triNormal[tri]); + const glm::mat3x2 projection = GetAxisAlignedProjection(faceNormal_[tri]); if (CCW(projection * pNext, projection * pLast, projection * pNew, precision_) < 0) return; pLast = pNext; - current = halfedge[current].pairedHalfedge; + current = halfedge_[current].pairedHalfedge; } } // Remove toRemove.startVert and replace with endVert. - vertPos[toRemove.startVert] = glm::vec3(NAN); + vertPos_[toRemove.startVert] = glm::vec3(NAN); CollapseTri(tri1edge); // Orbit startVert @@ -289,10 +285,10 @@ void Manifold::Impl::CollapseEdge(const int edge) { : ref1.vertBary[toRemove.pairedHalfedge % 3]; } - const int vert = halfedge[current].endVert; - const int next = halfedge[current].pairedHalfedge; + const int vert = halfedge_[current].endVert; + const int next = halfedge_[current].pairedHalfedge; for (int i = 0; i < edges.size(); ++i) { - if (vert == halfedge[edges[i]].endVert) { + if (vert == halfedge_[edges[i]].endVert) { FormLoop(edges[i], current); start = next; edges.resize(i); @@ -308,48 +304,45 @@ void Manifold::Impl::CollapseEdge(const int edge) { } void Manifold::Impl::RecursiveEdgeSwap(const int edge) { - const VecDH& vertPos = vertPos_; - VecDH& halfedge = halfedge_; - VecDH& triNormal = faceNormal_; VecDH& triBary = meshRelation_.triBary; - if (halfedge[edge].pairedHalfedge < 0) return; + if (halfedge_[edge].pairedHalfedge < 0) return; - const int pair = halfedge[edge].pairedHalfedge; + const int pair = halfedge_[edge].pairedHalfedge; const glm::ivec3 tri0edge = TriOf(edge); const glm::ivec3 tri1edge = TriOf(pair); const glm::ivec3 perm0 = TriOf(edge % 3); const glm::ivec3 perm1 = TriOf(pair % 3); - glm::mat3x2 projection = GetAxisAlignedProjection(triNormal[edge / 3]); + glm::mat3x2 projection = GetAxisAlignedProjection(faceNormal_[edge / 3]); glm::vec2 v[4]; for (int i : {0, 1, 2}) - v[i] = projection * vertPos[halfedge[tri0edge[i]].startVert]; + v[i] = projection * vertPos_[halfedge_[tri0edge[i]].startVert]; // Only operate on the long edge of a degenerate triangle. if (CCW(v[0], v[1], v[2], precision_) > 0 || !Is01Longest(v[0], v[1], v[2])) return; // Switch to neighbor's projection. - projection = GetAxisAlignedProjection(triNormal[halfedge[pair].face]); + projection = GetAxisAlignedProjection(faceNormal_[halfedge_[pair].face]); for (int i : {0, 1, 2}) - v[i] = projection * vertPos[halfedge[tri0edge[i]].startVert]; - v[3] = projection * vertPos[halfedge[tri1edge[2]].startVert]; + v[i] = projection * vertPos_[halfedge_[tri0edge[i]].startVert]; + v[3] = projection * vertPos_[halfedge_[tri1edge[2]].startVert]; auto SwapEdge = [&]() { // The 0-verts are swapped to the opposite 2-verts. - const int v0 = halfedge[tri0edge[2]].startVert; - const int v1 = halfedge[tri1edge[2]].startVert; - halfedge[tri0edge[0]].startVert = v1; - halfedge[tri0edge[2]].endVert = v1; - halfedge[tri1edge[0]].startVert = v0; - halfedge[tri1edge[2]].endVert = v0; - PairUp(tri0edge[0], halfedge[tri1edge[2]].pairedHalfedge); - PairUp(tri1edge[0], halfedge[tri0edge[2]].pairedHalfedge); + const int v0 = halfedge_[tri0edge[2]].startVert; + const int v1 = halfedge_[tri1edge[2]].startVert; + halfedge_[tri0edge[0]].startVert = v1; + halfedge_[tri0edge[2]].endVert = v1; + halfedge_[tri1edge[0]].startVert = v0; + halfedge_[tri1edge[2]].endVert = v0; + PairUp(tri0edge[0], halfedge_[tri1edge[2]].pairedHalfedge); + PairUp(tri1edge[0], halfedge_[tri0edge[2]].pairedHalfedge); PairUp(tri0edge[2], tri1edge[2]); // Both triangles are now subsets of the neighboring triangle. - const int tri0 = halfedge[tri0edge[0]].face; - const int tri1 = halfedge[tri1edge[0]].face; - triNormal[tri0] = triNormal[tri1]; + const int tri0 = halfedge_[tri0edge[0]].face; + const int tri1 = halfedge_[tri1edge[0]].face; + faceNormal_[tri0] = faceNormal_[tri1]; triBary[tri0] = triBary[tri1]; triBary[tri0].vertBary[perm0[1]] = triBary[tri1].vertBary[perm1[0]]; triBary[tri0].vertBary[perm0[0]] = triBary[tri1].vertBary[perm1[2]]; @@ -369,16 +362,16 @@ void Manifold::Impl::RecursiveEdgeSwap(const int edge) { triBary[tri0].vertBary[perm0[2]] = newBary; // if the new edge already exists, duplicate the verts and split the mesh. - int current = halfedge[tri1edge[0]].pairedHalfedge; - const int endVert = halfedge[tri1edge[1]].endVert; + int current = halfedge_[tri1edge[0]].pairedHalfedge; + const int endVert = halfedge_[tri1edge[1]].endVert; while (current != tri0edge[1]) { current = NextHalfedge(current); - if (halfedge[current].endVert == endVert) { + if (halfedge_[current].endVert == endVert) { FormLoop(tri0edge[2], current); RemoveIfFolded(tri0edge[2]); return; } - current = halfedge[current].pairedHalfedge; + current = halfedge_[current].pairedHalfedge; } }; @@ -403,7 +396,7 @@ void Manifold::Impl::RecursiveEdgeSwap(const int edge) { } // Normal path SwapEdge(); - RecursiveEdgeSwap(halfedge[tri0edge[1]].pairedHalfedge); - RecursiveEdgeSwap(halfedge[tri1edge[0]].pairedHalfedge); + RecursiveEdgeSwap(halfedge_[tri0edge[1]].pairedHalfedge); + RecursiveEdgeSwap(halfedge_[tri1edge[0]].pairedHalfedge); } } // namespace manifold diff --git a/manifold/src/face_op.cpp b/manifold/src/face_op.cpp index fa5df3d2a..d90ea6492 100644 --- a/manifold/src/face_op.cpp +++ b/manifold/src/face_op.cpp @@ -40,21 +40,12 @@ void Manifold::Impl::Face2Tri(const VecDH& faceEdge, triNormal.reserve(faceEdge.size()); triBary.reserve(faceEdge.size()*3); - const VecDH& vertPos = vertPos_; - const VecDH& faceEdgeH = faceEdge; - const VecDH& halfedge = halfedge_; - const VecDH& faceNormal = faceNormal_; - // meshRelation_.triBary.resize(0); - // std::vector triVerts; - // std::vector triNormal; - // std::vector triBary; - - for (int face = 0; face < faceEdgeH.size() - 1; ++face) { - const int firstEdge = faceEdgeH[face]; - const int lastEdge = faceEdgeH[face + 1]; + for (int face = 0; face < faceEdge.size() - 1; ++face) { + const int firstEdge = faceEdge[face]; + const int lastEdge = faceEdge[face + 1]; const int numEdge = lastEdge - firstEdge; ALWAYS_ASSERT(numEdge >= 3, topologyErr, "face has less than three edges."); - const glm::vec3 normal = faceNormal[face]; + const glm::vec3 normal = faceNormal_[face]; auto linearSearch = [](const int* mapping, int value) { int i = 0; @@ -64,15 +55,15 @@ void Manifold::Impl::Face2Tri(const VecDH& faceEdge, }; if (numEdge == 3) { // Single triangle - int mapping[3] = {halfedge[firstEdge].startVert, - halfedge[firstEdge + 1].startVert, - halfedge[firstEdge + 2].startVert}; - glm::ivec3 tri(halfedge[firstEdge].startVert, - halfedge[firstEdge + 1].startVert, - halfedge[firstEdge + 2].startVert); - glm::ivec3 ends(halfedge[firstEdge].endVert, - halfedge[firstEdge + 1].endVert, - halfedge[firstEdge + 2].endVert); + int mapping[3] = {halfedge_[firstEdge].startVert, + halfedge_[firstEdge + 1].startVert, + halfedge_[firstEdge + 2].startVert}; + glm::ivec3 tri(halfedge_[firstEdge].startVert, + halfedge_[firstEdge + 1].startVert, + halfedge_[firstEdge + 2].startVert); + glm::ivec3 ends(halfedge_[firstEdge].endVert, + halfedge_[firstEdge + 1].endVert, + halfedge_[firstEdge + 2].endVert); if (ends[0] == tri[2]) { std::swap(tri[1], tri[2]); std::swap(ends[1], ends[2]); @@ -88,26 +79,26 @@ void Manifold::Impl::Face2Tri(const VecDH& faceEdge, triBary.back().vertBary[k] = halfedgeBary[firstEdge + index]; } } else if (numEdge == 4) { // Pair of triangles - int mapping[4] = {halfedge[firstEdge].startVert, - halfedge[firstEdge + 1].startVert, - halfedge[firstEdge + 2].startVert, - halfedge[firstEdge + 3].startVert}; + int mapping[4] = {halfedge_[firstEdge].startVert, + halfedge_[firstEdge + 1].startVert, + halfedge_[firstEdge + 2].startVert, + halfedge_[firstEdge + 3].startVert}; const glm::mat3x2 projection = GetAxisAlignedProjection(normal); - auto triCCW = [&projection, &vertPos, this](const glm::ivec3 tri) { - return CCW(projection * vertPos[tri[0]], projection * vertPos[tri[1]], - projection * vertPos[tri[2]], precision_) >= 0; + auto triCCW = [&projection, this](const glm::ivec3 tri) { + return CCW(projection * this->vertPos_[tri[0]], projection * this->vertPos_[tri[1]], + projection * this->vertPos_[tri[2]], precision_) >= 0; }; - glm::ivec3 tri0(halfedge[firstEdge].startVert, - halfedge[firstEdge].endVert, -1); + glm::ivec3 tri0(halfedge_[firstEdge].startVert, + halfedge_[firstEdge].endVert, -1); glm::ivec3 tri1(-1, -1, tri0[0]); for (const int i : {1, 2, 3}) { - if (halfedge[firstEdge + i].startVert == tri0[1]) { - tri0[2] = halfedge[firstEdge + i].endVert; + if (halfedge_[firstEdge + i].startVert == tri0[1]) { + tri0[2] = halfedge_[firstEdge + i].endVert; tri1[0] = tri0[2]; } - if (halfedge[firstEdge + i].endVert == tri0[0]) { - tri1[1] = halfedge[firstEdge + i].startVert; + if (halfedge_[firstEdge + i].endVert == tri0[0]) { + tri1[1] = halfedge_[firstEdge + i].startVert; } } ALWAYS_ASSERT(glm::all(glm::greaterThanEqual(tri0, glm::ivec3(0))) && @@ -122,8 +113,8 @@ void Manifold::Impl::Face2Tri(const VecDH& faceEdge, tri0[2] = tri1[0]; tri1[2] = tri0[0]; } else if (firstValid) { - glm::vec3 firstCross = vertPos[tri0[0]] - vertPos[tri1[0]]; - glm::vec3 secondCross = vertPos[tri0[1]] - vertPos[tri1[1]]; + glm::vec3 firstCross = vertPos_[tri0[0]] - vertPos_[tri1[0]]; + glm::vec3 secondCross = vertPos_[tri0[1]] - vertPos_[tri1[1]]; if (glm::dot(firstCross, firstCross) < glm::dot(secondCross, secondCross)) { tri0[2] = tri1[0]; @@ -145,15 +136,15 @@ void Manifold::Impl::Face2Tri(const VecDH& faceEdge, std::map vertBary; for (int j = firstEdge; j < lastEdge; ++j) - vertBary[halfedge[j].startVert] = halfedgeBary[j]; + vertBary[halfedge_[j].startVert] = halfedgeBary[j]; Polygons polys; try { - polys = Face2Polygons(face, projection, faceEdgeH); + polys = Face2Polygons(face, projection, faceEdge); } catch (const std::exception& e) { std::cout << e.what() << std::endl; - for (int edge = faceEdgeH[face]; edge < faceEdgeH[face + 1]; ++edge) - std::cout << "halfedge: " << edge << ", " << halfedge[edge] + for (int edge = faceEdge[face]; edge < faceEdge[face + 1]; ++edge) + std::cout << "halfedge: " << edge << ", " << halfedge_[edge] << std::endl; throw; } @@ -181,15 +172,13 @@ void Manifold::Impl::Face2Tri(const VecDH& faceEdge, */ Polygons Manifold::Impl::Face2Polygons(int face, glm::mat3x2 projection, const VecDH& faceEdge) const { - const VecDH& vertPos = vertPos_; - const VecDH& halfedge = halfedge_; const int firstEdge = faceEdge[face]; const int lastEdge = faceEdge[face + 1]; std::map vert_edge; for (int edge = firstEdge; edge < lastEdge; ++edge) { ALWAYS_ASSERT( - vert_edge.emplace(std::make_pair(halfedge[edge].startVert, edge)) + vert_edge.emplace(std::make_pair(halfedge_[edge].startVert, edge)) .second, topologyErr, "face has duplicate vertices."); } @@ -204,9 +193,9 @@ Polygons Manifold::Impl::Face2Polygons(int face, glm::mat3x2 projection, thisEdge = startEdge; polys.push_back({}); } - int vert = halfedge[thisEdge].startVert; - polys.back().push_back({projection * vertPos[vert], vert}); - const auto result = vert_edge.find(halfedge[thisEdge].endVert); + int vert = halfedge_[thisEdge].startVert; + polys.back().push_back({projection * vertPos_[vert], vert}); + const auto result = vert_edge.find(halfedge_[thisEdge].endVert); ALWAYS_ASSERT(result != vert_edge.end(), topologyErr, "nonmanifold edge"); thisEdge = result->second; vert_edge.erase(result); diff --git a/manifold/src/impl.cpp b/manifold/src/impl.cpp index b35e5fef7..9b5513bd7 100644 --- a/manifold/src/impl.cpp +++ b/manifold/src/impl.cpp @@ -12,7 +12,6 @@ // See the License for the specific language governing permissions and // limitations under the License. -#include #include #include @@ -374,7 +373,7 @@ void Manifold::Impl::DuplicateMeshIDs() { } void Manifold::Impl::ReinitializeReference(int meshID) { - thrust::for_each_n(thrust::device, zip(meshRelation_.triBary.beginD(), countAt(0)), NumTri(), + thrust::for_each_n(thrust::device, zip(meshRelation_.triBary.begin(), countAt(0)), NumTri(), InitializeBaryRef({meshID, halfedge_.cptrD()})); } @@ -402,7 +401,7 @@ int Manifold::Impl::InitializeNewReference( "propertyTolerance."); const int numSets = properties.size() / numProps; - ALWAYS_ASSERT(thrust::all_of(thrust::device, triPropertiesD.beginD(), triPropertiesD.endD(), + ALWAYS_ASSERT(thrust::all_of(thrust::device, triPropertiesD.begin(), triPropertiesD.end(), CheckProperties({numSets})), userErr, "triProperties value is outside the properties range."); @@ -411,7 +410,7 @@ int Manifold::Impl::InitializeNewReference( VecDH> face2face(halfedge_.size(), {-1, -1}); VecDH triArea(NumTri()); thrust::for_each_n( - thrust::device, zip(face2face.beginD(), countAt(0)), halfedge_.size(), + thrust::device, zip(face2face.begin(), countAt(0)), halfedge_.size(), CoplanarEdge({triArea.ptrD(), halfedge_.cptrD(), vertPos_.cptrD(), triPropertiesD.cptrD(), propertiesD.cptrD(), propertyToleranceD.cptrD(), numProps, precision_})); @@ -493,9 +492,9 @@ void Manifold::Impl::CreateHalfedges(const VecDH& triVerts) { const int numTri = triVerts.size(); halfedge_.resize(3 * numTri); VecDH edge(3 * numTri); - thrust::for_each_n(thrust::device, zip(countAt(0), triVerts.beginD()), numTri, + thrust::for_each_n(thrust::device, zip(countAt(0), triVerts.begin()), numTri, Tri2Halfedges({halfedge_.ptrD(), edge.ptrD()})); - thrust::sort(thrust::device, edge.beginD(), edge.endD()); + thrust::sort(thrust::device, edge.begin(), edge.end()); thrust::for_each_n(thrust::device, countAt(0), halfedge_.size() / 2, LinkHalfedges({halfedge_.ptrD(), edge.cptrD()})); } @@ -511,7 +510,7 @@ void Manifold::Impl::CreateAndFixHalfedges(const VecDH& triVerts) { halfedge_.resize(0); halfedge_.resize(3 * numTri); VecDH edge(3 * numTri); - thrust::for_each_n(thrust::device, zip(countAt(0), triVerts.beginD()), numTri, + thrust::for_each_n(thrust::device, zip(countAt(0), triVerts.begin()), 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 +518,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(thrust::device, edge.beginD(), edge.endD()); + thrust::stable_sort(thrust::device, edge.begin(), edge.end()); 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 +550,14 @@ void Manifold::Impl::ApplyTransform() const { */ void Manifold::Impl::ApplyTransform() { if (transform_ == glm::mat4x3(1.0f)) return; - thrust::for_each(thrust::device, vertPos_.beginD(), vertPos_.endD(), + thrust::for_each(thrust::device, vertPos_.begin(), vertPos_.end(), Transform4x3({transform_})); glm::mat3 normalTransform = glm::inverse(glm::transpose(glm::mat3(transform_))); - thrust::for_each(thrust::device, faceNormal_.beginD(), faceNormal_.endD(), + thrust::for_each(thrust::device, faceNormal_.begin(), faceNormal_.end(), TransformNormals({normalTransform})); - thrust::for_each(thrust::device, vertNormal_.beginD(), vertNormal_.endD(), + thrust::for_each(thrust::device, vertNormal_.begin(), vertNormal_.end(), TransformNormals({normalTransform})); // This optimization does a cheap collider update if the transform is // axis-aligned. @@ -601,17 +600,17 @@ void Manifold::Impl::SetPrecision(float minPrecision) { */ void Manifold::Impl::CalculateNormals() { vertNormal_.resize(NumVert()); - thrust::fill(thrust::device, vertNormal_.beginD(), vertNormal_.endD(), glm::vec3(0)); + thrust::fill(thrust::device, vertNormal_.begin(), vertNormal_.end(), glm::vec3(0)); bool calculateTriNormal = false; if (faceNormal_.size() != NumTri()) { faceNormal_.resize(NumTri()); calculateTriNormal = true; } thrust::for_each_n( - thrust::device, zip(faceNormal_.beginD(), countAt(0)), NumTri(), + thrust::device, zip(faceNormal_.begin(), countAt(0)), NumTri(), AssignNormals({vertNormal_.ptrD(), vertPos_.cptrD(), halfedge_.cptrD(), precision_, calculateTriNormal})); - thrust::for_each(thrust::device, vertNormal_.beginD(), vertNormal_.endD(), Normalize()); + thrust::for_each(thrust::device, vertNormal_.begin(), vertNormal_.end(), Normalize()); } /** @@ -623,12 +622,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(thrust::device, zip(QedgeBB.beginD(), edges.cbeginD()), numEdge, + thrust::for_each_n(thrust::device, zip(QedgeBB.begin(), edges.cbegin()), numEdge, EdgeBox({Q.vertPos_.cptrD()})); SparseIndices q1p2 = collider_.Collisions(QedgeBB); - thrust::for_each(thrust::device, q1p2.beginD(0), q1p2.endD(0), ReindexEdge({edges.cptrD()})); + thrust::for_each(thrust::device, q1p2.begin(0), q1p2.end(0), ReindexEdge({edges.cptrD()})); return q1p2; } diff --git a/manifold/src/manifold.cpp b/manifold/src/manifold.cpp index 35f28796f..642d2acdf 100644 --- a/manifold/src/manifold.cpp +++ b/manifold/src/manifold.cpp @@ -14,7 +14,6 @@ #include "boolean3.h" #include "impl.h" -#include namespace { using namespace manifold; @@ -289,11 +288,11 @@ MeshRelation Manifold::GetMeshRelation() const { std::vector Manifold::GetMeshIDs() const { VecDH meshIDs(NumTri()); thrust::for_each_n( - thrust::device, zip(meshIDs.beginD(), pImpl_->meshRelation_.triBary.beginD()), NumTri(), + thrust::device, zip(meshIDs.begin(), pImpl_->meshRelation_.triBary.begin()), NumTri(), GetMeshID()); - thrust::sort(thrust::device, meshIDs.beginD(), meshIDs.endD()); - int n = thrust::unique(thrust::device, meshIDs.beginD(), meshIDs.endD()) - meshIDs.beginD(); + thrust::sort(thrust::device, meshIDs.begin(), meshIDs.end()); + int n = thrust::unique(thrust::device, meshIDs.begin(), meshIDs.end()) - meshIDs.begin(); meshIDs.resize(n); std::vector out; diff --git a/manifold/src/properties.cpp b/manifold/src/properties.cpp index 672b61f03..3862fec2e 100644 --- a/manifold/src/properties.cpp +++ b/manifold/src/properties.cpp @@ -13,7 +13,6 @@ // limitations under the License. #include -#include #include #include #include @@ -219,7 +218,7 @@ bool Manifold::Impl::IsManifold() const { CheckManifold({halfedge_.cptrD()})); VecDH halfedge(halfedge_); - thrust::sort(thrust::device, halfedge.beginD(), halfedge.endD()); + thrust::sort(thrust::device, halfedge.begin(), halfedge.end()); isManifold &= thrust::all_of(thrust::device, countAt(0), countAt(2 * NumEdge() - 1), NoDuplicates({halfedge.cptrD()})); return isManifold; @@ -269,20 +268,20 @@ Curvature Manifold::Impl::GetCurvature() const { vertArea.ptrD(), degree.ptrD(), halfedge_.cptrD(), vertPos_.cptrD(), faceNormal_.cptrD()})); thrust::for_each_n( - thrust::device, zip(vertMeanCurvature.beginD(), vertGaussianCurvature.beginD(), - vertArea.beginD(), degree.beginD()), + thrust::device, zip(vertMeanCurvature.begin(), vertGaussianCurvature.begin(), + vertArea.begin(), degree.begin()), NumVert(), NormalizeCurvature()); result.minMeanCurvature = - thrust::reduce(thrust::device, vertMeanCurvature.beginD(), vertMeanCurvature.endD(), + thrust::reduce(thrust::device, vertMeanCurvature.begin(), vertMeanCurvature.end(), std::numeric_limits::infinity(), thrust::minimum()); result.maxMeanCurvature = - thrust::reduce(thrust::device, vertMeanCurvature.beginD(), vertMeanCurvature.endD(), + thrust::reduce(thrust::device, vertMeanCurvature.begin(), vertMeanCurvature.end(), -std::numeric_limits::infinity(), thrust::maximum()); result.minGaussianCurvature = thrust::reduce( - thrust::device, vertGaussianCurvature.beginD(), vertGaussianCurvature.endD(), std::numeric_limits::infinity(), + thrust::device, vertGaussianCurvature.begin(), vertGaussianCurvature.end(), std::numeric_limits::infinity(), thrust::minimum()); result.maxGaussianCurvature = thrust::reduce( - thrust::device, vertGaussianCurvature.beginD(), vertGaussianCurvature.endD(), + thrust::device, vertGaussianCurvature.begin(), vertGaussianCurvature.end(), -std::numeric_limits::infinity(), thrust::maximum()); result.vertMeanCurvature.insert(result.vertMeanCurvature.end(), vertMeanCurvature.begin(), @@ -299,9 +298,9 @@ Curvature Manifold::Impl::GetCurvature() const { * range for Morton code calculation. */ void Manifold::Impl::CalculateBBox() { - bBox_.min = thrust::reduce(thrust::device, vertPos_.beginD(), vertPos_.endD(), + bBox_.min = thrust::reduce(thrust::device, vertPos_.begin(), vertPos_.end(), glm::vec3(std::numeric_limits::infinity()), PosMin()); - bBox_.max = thrust::reduce(thrust::device, vertPos_.beginD(), vertPos_.endD(), + bBox_.max = thrust::reduce(thrust::device, vertPos_.begin(), vertPos_.end(), glm::vec3(-std::numeric_limits::infinity()), PosMax()); } } // namespace manifold diff --git a/manifold/src/shared.h b/manifold/src/shared.h index 295810a4f..f78b7ada8 100644 --- a/manifold/src/shared.h +++ b/manifold/src/shared.h @@ -14,6 +14,9 @@ #pragma once +#include +#include +#include "utils.h" #include "vec_dh.h" namespace manifold { @@ -23,7 +26,7 @@ namespace manifold { */ __host__ __device__ inline glm::vec3 SafeNormalize(glm::vec3 v) { v = glm::normalize(v); - return isfinite(v.x) ? v : glm::vec3(0); + return glm::isfinite(v.x) ? v : glm::vec3(0); } __host__ __device__ inline int NextHalfedge(int current) { @@ -146,10 +149,10 @@ struct TmpInvalid { VecDH inline CreateTmpEdges(const VecDH& halfedge) { VecDH edges(halfedge.size()); - thrust::for_each_n(thrust::device, zip(edges.beginD(), halfedge.beginD(), countAt(0)), + thrust::for_each_n(thrust::device, zip(edges.begin(), halfedge.begin(), countAt(0)), edges.size(), Halfedge2Tmp()); - int numEdge = thrust::remove_if(thrust::device, edges.beginD(), edges.endD(), TmpInvalid()) - - edges.beginD(); + int numEdge = thrust::remove_if(thrust::device, edges.begin(), edges.end(), TmpInvalid()) - + edges.begin(); ALWAYS_ASSERT(numEdge == halfedge.size() / 2, topologyErr, "Not oriented!"); edges.resize(numEdge); return edges; diff --git a/manifold/src/smoothing.cpp b/manifold/src/smoothing.cpp index c8c996bf5..7a41d0753 100644 --- a/manifold/src/smoothing.cpp +++ b/manifold/src/smoothing.cpp @@ -356,19 +356,18 @@ void Manifold::Impl::CreateTangents( const int numHalfedge = halfedge_.size(); halfedgeTangent_.resize(numHalfedge); - thrust::for_each_n(thrust::device, zip(halfedgeTangent_.beginD(), halfedge_.cbeginD()), + thrust::for_each_n(thrust::device, zip(halfedgeTangent_.begin(), halfedge_.cbegin()), numHalfedge, SmoothBezier({vertPos_.cptrD(), faceNormal_.cptrD(), vertNormal_.cptrD(), halfedge_.cptrD()})); if (!sharpenedEdges.empty()) { - const VecDH& halfedge = halfedge_; const VecDH& triBary = meshRelation_.triBary; // sharpenedEdges are referenced to the input Mesh, but the triangles have // been sorted in creating the Manifold, so the indices are converted using // meshRelation_. - std::vector oldHalfedge2New(halfedge.size()); + std::vector oldHalfedge2New(halfedge_.size()); for (int tri = 0; tri < NumTri(); ++tri) { int oldTri = triBary[tri].tri; for (int i : {0, 1, 2}) oldHalfedge2New[3 * oldTri + i] = 3 * tri + i; @@ -380,7 +379,7 @@ void Manifold::Impl::CreateTangents( for (Smoothness edge : sharpenedEdges) { if (edge.smoothness == 1) continue; edge.halfedge = oldHalfedge2New[edge.halfedge]; - int pair = halfedge[edge.halfedge].pairedHalfedge; + int pair = halfedge_[edge.halfedge].pairedHalfedge; if (edges.find(pair) == edges.end()) { edges[edge.halfedge] = {edge, {pair, 1}}; } else { @@ -391,8 +390,8 @@ void Manifold::Impl::CreateTangents( std::map> vertTangents; for (const auto &value : edges) { const Pair edge = value.second; - vertTangents[halfedge[edge.first.halfedge].startVert].push_back(edge); - vertTangents[halfedge[edge.second.halfedge].startVert].push_back( + vertTangents[halfedge_[edge.first.halfedge].startVert].push_back(edge); + vertTangents[halfedge_[edge.second.halfedge].startVert].push_back( {edge.second, edge.first}); } @@ -414,7 +413,7 @@ void Manifold::Impl::CreateTangents( tangent[second].w); auto SmoothHalf = [&](int first, int last, float smoothness) { - int current = NextHalfedge(halfedge[first].pairedHalfedge); + int current = NextHalfedge(halfedge_[first].pairedHalfedge); while (current != last) { const float cosBeta = glm::dot( newTangent, glm::normalize(glm::vec3(tangent[current]))); @@ -422,7 +421,7 @@ void Manifold::Impl::CreateTangents( (1 - smoothness) * cosBeta * cosBeta + smoothness; tangent[current] = glm::vec4(factor * glm::vec3(tangent[current]), tangent[current].w); - current = NextHalfedge(halfedge[current].pairedHalfedge); + current = NextHalfedge(halfedge_[current].pairedHalfedge); } }; @@ -444,7 +443,7 @@ void Manifold::Impl::CreateTangents( do { tangent[current] = glm::vec4(smoothness * glm::vec3(tangent[current]), tangent[current].w); - current = NextHalfedge(halfedge[current].pairedHalfedge); + current = NextHalfedge(halfedge_[current].pairedHalfedge); } while (current != start); } } @@ -476,12 +475,12 @@ Manifold::Impl::MeshRelationD Manifold::Impl::Subdivide(int n) { VecDH edges = CreateTmpEdges(halfedge_); VecDH half2Edge(2 * numEdge); - thrust::for_each_n(thrust::device, zip(countAt(0), edges.beginD()), numEdge, + thrust::for_each_n(thrust::device, zip(countAt(0), edges.begin()), numEdge, ReindexHalfedge({half2Edge.ptrD()})); - thrust::for_each_n(thrust::device, zip(countAt(0), edges.beginD()), numEdge, + thrust::for_each_n(thrust::device, zip(countAt(0), edges.begin()), numEdge, EdgeVerts({vertPos_.ptrD(), numVert, n})); thrust::for_each_n( - thrust::device, zip(countAt(0), oldMeshRelation.triBary.beginD()), numTri, + thrust::device, zip(countAt(0), oldMeshRelation.triBary.begin()), numTri, InteriorVerts({vertPos_.ptrD(), relation.barycentric.ptrD(), relation.triBary.ptrD(), meshRelation_.barycentric.ptrD(), meshRelation_.triBary.ptrD(), @@ -504,12 +503,12 @@ void Manifold::Impl::Refine(int n) { VecDH vertBary(NumVert()); VecDH lock(NumVert(), 0); thrust::for_each_n( - thrust::device, zip(relation.triBary.beginD(), countAt(0)), NumTri(), + thrust::device, zip(relation.triBary.begin(), countAt(0)), NumTri(), TriBary2Vert({vertBary.ptrD(), lock.ptrD(), relation.barycentric.cptrD(), halfedge_.cptrD()})); thrust::for_each_n( - thrust::device, zip(vertPos_.beginD(), vertBary.beginD()), NumVert(), + thrust::device, zip(vertPos_.begin(), vertBary.begin()), NumVert(), InterpTri({old.halfedge_.cptrD(), old.halfedgeTangent_.cptrD(), old.vertPos_.cptrD()})); } diff --git a/manifold/src/sort.cpp b/manifold/src/sort.cpp index 94b0fef09..7a0a67af4 100644 --- a/manifold/src/sort.cpp +++ b/manifold/src/sort.cpp @@ -122,8 +122,8 @@ template void Permute(VecDH& inOut, const VecDH& new2Old) { VecDH tmp(inOut); inOut.resize(new2Old.size()); - thrust::gather(new2Old.beginD(), new2Old.endD(), tmp.beginD(), - inOut.beginD()); + thrust::gather(new2Old.begin(), new2Old.end(), tmp.begin(), + inOut.begin()); } template void Permute(VecDH&, const VecDH&); @@ -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(thrust::device, halfedge_.beginD(), halfedge_.endD(), extrema, Extrema()); + thrust::reduce(thrust::device, halfedge_.begin(), halfedge_.end(), extrema, Extrema()); ALWAYS_ASSERT(extrema.startVert >= 0, topologyErr, "Vertex index is negative!"); @@ -210,21 +210,21 @@ void Manifold::Impl::Finish() { */ void Manifold::Impl::SortVerts() { VecDH vertMorton(NumVert()); - thrust::for_each_n(thrust::device, zip(vertMorton.beginD(), vertPos_.cbeginD()), NumVert(), + thrust::for_each_n(thrust::device, zip(vertMorton.begin(), vertPos_.cbegin()), NumVert(), Morton({bBox_})); VecDH vertNew2Old(NumVert()); - thrust::sequence(thrust::device, vertNew2Old.beginD(), vertNew2Old.endD()); - thrust::sort_by_key(thrust::device, vertMorton.beginD(), vertMorton.endD(), - zip(vertPos_.beginD(), vertNew2Old.beginD())); + thrust::sequence(thrust::device, vertNew2Old.begin(), vertNew2Old.end()); + thrust::sort_by_key(thrust::device, vertMorton.begin(), vertMorton.end(), + zip(vertPos_.begin(), vertNew2Old.begin())); ReindexVerts(vertNew2Old, NumVert()); // 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(thrust::device, vertMorton.beginD(), vertMorton.endD(), kNoCode) - - vertMorton.beginD(); + thrust::find(thrust::device, vertMorton.begin(), vertMorton.end(), kNoCode) - + vertMorton.begin(); 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(thrust::device, countAt(0), countAt(NumVert()), vertNew2Old.beginD(), - vertOld2New.beginD()); - thrust::for_each(thrust::device, halfedge_.beginD(), halfedge_.endD(), + thrust::scatter(thrust::device, countAt(0), countAt(NumVert()), vertNew2Old.begin(), + vertOld2New.begin()); + thrust::for_each(thrust::device, halfedge_.begin(), halfedge_.end(), Reindex({vertOld2New.cptrD()})); } @@ -252,7 +252,7 @@ void Manifold::Impl::GetFaceBoxMorton(VecDH& faceBox, faceBox.resize(NumTri()); faceMorton.resize(NumTri()); thrust::for_each_n( - thrust::device, zip(faceMorton.beginD(), faceBox.beginD(), countAt(0)), NumTri(), + thrust::device, zip(faceMorton.begin(), faceBox.begin(), countAt(0)), NumTri(), FaceMortonBox({halfedge_.cptrD(), vertPos_.cptrD(), bBox_})); } @@ -263,16 +263,16 @@ void Manifold::Impl::GetFaceBoxMorton(VecDH& faceBox, void Manifold::Impl::SortFaces(VecDH& faceBox, VecDH& faceMorton) { VecDH faceNew2Old(NumTri()); - thrust::sequence(thrust::device, faceNew2Old.beginD(), faceNew2Old.endD()); + thrust::sequence(thrust::device, faceNew2Old.begin(), faceNew2Old.end()); - thrust::sort_by_key(thrust::device, faceMorton.beginD(), faceMorton.endD(), - zip(faceBox.beginD(), faceNew2Old.beginD())); + thrust::sort_by_key(thrust::device, faceMorton.begin(), faceMorton.end(), + zip(faceBox.begin(), faceNew2Old.begin())); // 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(thrust::device, faceMorton.beginD(), faceMorton.endD(), kNoCode) - - faceMorton.beginD(); + thrust::find(thrust::device, faceMorton.begin(), faceMorton.end(), kNoCode) - + faceMorton.begin(); faceBox.resize(newNumTri); faceMorton.resize(newNumTri); faceNew2Old.resize(newNumTri); @@ -295,8 +295,8 @@ void Manifold::Impl::GatherFaces(const VecDH& faceNew2Old) { VecDH oldHalfedge(std::move(halfedge_)); VecDH oldHalfedgeTangent(std::move(halfedgeTangent_)); VecDH faceOld2New(oldHalfedge.size() / 3); - thrust::scatter(thrust::device, countAt(0), countAt(numTri), faceNew2Old.beginD(), - faceOld2New.beginD()); + thrust::scatter(thrust::device, countAt(0), countAt(numTri), faceNew2Old.begin(), + faceOld2New.begin()); halfedge_.resize(3 * numTri); if (oldHalfedgeTangent.size() != 0) halfedgeTangent_.resize(3 * numTri); @@ -311,21 +311,21 @@ void Manifold::Impl::GatherFaces(const Impl& old, const VecDH& faceNew2Old) { const int numTri = faceNew2Old.size(); meshRelation_.triBary.resize(numTri); - thrust::gather(thrust::device, faceNew2Old.beginD(), faceNew2Old.endD(), - old.meshRelation_.triBary.beginD(), - meshRelation_.triBary.beginD()); + thrust::gather(thrust::device, faceNew2Old.begin(), faceNew2Old.end(), + old.meshRelation_.triBary.begin(), + meshRelation_.triBary.begin()); meshRelation_.barycentric = old.meshRelation_.barycentric; DuplicateMeshIDs(); if (old.faceNormal_.size() == old.NumTri()) { faceNormal_.resize(numTri); - thrust::gather(thrust::device, faceNew2Old.beginD(), faceNew2Old.endD(), - old.faceNormal_.beginD(), faceNormal_.beginD()); + thrust::gather(thrust::device, faceNew2Old.begin(), faceNew2Old.end(), + old.faceNormal_.begin(), faceNormal_.begin()); } VecDH faceOld2New(old.NumTri()); - thrust::scatter(thrust::device, countAt(0), countAt(numTri), faceNew2Old.beginD(), - faceOld2New.beginD()); + thrust::scatter(thrust::device, countAt(0), countAt(numTri), faceNew2Old.begin(), + faceOld2New.begin()); halfedge_.resize(3 * numTri); if (old.halfedgeTangent_.size() != 0) halfedgeTangent_.resize(3 * numTri); diff --git a/utilities/include/sparse.h b/utilities/include/sparse.h index 3b18b065e..585840e36 100644 --- a/utilities/include/sparse.h +++ b/utilities/include/sparse.h @@ -34,12 +34,12 @@ class SparseIndices { // stored in vectors separate from this class, but having the same length. public: SparseIndices(int size = 0) : p(size), q(size) {} - typedef typename VecDH::IterD Iter; + typedef typename VecDH::Iter Iter; typedef typename thrust::zip_iterator> Zip; - Zip beginDpq() { return zip(p.beginD(), q.beginD()); } - Zip endDpq() { return zip(p.endD(), q.endD()); } - Iter beginD(bool use_q) { return use_q ? q.beginD() : p.beginD(); } - Iter endD(bool use_q) { return use_q ? q.endD() : p.endD(); } + Zip beginPQ() { return zip(p.begin(), q.begin()); } + Zip endPQ() { return zip(p.end(), q.end()); } + Iter begin(bool use_q) { return use_q ? q.begin() : p.begin(); } + Iter end(bool use_q) { return use_q ? q.end() : p.end(); } int* ptrD(bool use_q) { return use_q ? q.ptrD() : p.ptrD(); } thrust::pair ptrDpq(int idx = 0) { return thrust::make_pair(p.ptrD() + idx, q.ptrD() + idx); @@ -53,23 +53,21 @@ class SparseIndices { return out; } - typedef typename VecDH::IterDc IterC; + typedef typename VecDH::IterC IterC; typedef typename thrust::zip_iterator> ZipC; - ZipC beginDpq() const { return zip(p.beginD(), q.beginD()); } - ZipC endDpq() const { return zip(p.endD(), q.endD()); } - IterC beginD(bool use_q) const { return use_q ? q.beginD() : p.beginD(); } - IterC endD(bool use_q) const { return use_q ? q.endD() : p.endD(); } + ZipC beginPQ() const { return zip(p.begin(), q.begin()); } + ZipC endPQ() const { return zip(p.end(), q.end()); } + IterC begin(bool use_q) const { return use_q ? q.begin() : p.begin(); } + IterC end(bool use_q) const { return use_q ? q.end() : p.end(); } const int* ptrD(bool use_q) const { return use_q ? q.ptrD() : p.ptrD(); } - typedef typename VecDH::IterHc IterHC; - typedef typename thrust::zip_iterator> ZipHC; - ZipHC beginHpq() const { return zip(p.begin(), q.begin()); } - ZipHC endHpq() const { return zip(p.end(), q.end()); } + ZipC beginHpq() const { return zip(p.begin(), q.begin()); } + ZipC endHpq() const { return zip(p.end(), q.end()); } int size() const { return p.size(); } void SwapPQ() { p.swap(q); } - void Sort() { thrust::sort(beginDpq(), endDpq()); } + void Sort() { thrust::sort(beginPQ(), endPQ()); } void Resize(int size) { p.resize(size, -1); @@ -78,7 +76,7 @@ class SparseIndices { void Unique() { Sort(); - int newSize = thrust::unique(beginDpq(), endDpq()) - beginDpq(); + int newSize = thrust::unique(beginPQ(), endPQ()) - beginPQ(); Resize(newSize); } @@ -91,8 +89,8 @@ class SparseIndices { size_t RemoveZeros(VecDH& S) { ALWAYS_ASSERT(S.size() == p.size(), userErr, "Different number of values than indicies!"); - auto zBegin = zip(S.beginD(), beginD(false), beginD(true)); - auto zEnd = zip(S.endD(), endD(false), endD(true)); + auto zBegin = zip(S.begin(), begin(false), begin(true)); + auto zEnd = zip(S.end(), end(false), end(true)); size_t size = thrust::remove_if(zBegin, zEnd, firstZero()) - zBegin; S.resize(size, -1); p.resize(size, -1); @@ -124,8 +122,8 @@ class SparseIndices { size_t KeepFinite(VecDH& v, VecDH& x) { ALWAYS_ASSERT(x.size() == p.size(), userErr, "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)); + auto zBegin = zip(v.begin(), x.begin(), begin(false), begin(true)); + auto zEnd = zip(v.end(), x.end(), end(false), end(true)); size_t size = thrust::remove_if(thrust::device, zBegin, zEnd, firstNonFinite()) - zBegin; v.resize(size); x.resize(size, -1); @@ -143,11 +141,11 @@ class SparseIndices { VecDH result(size); VecDH found(size); VecDH temp(size); - 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()); + thrust::fill(thrust::device, result.begin(), result.end(), missingVal); + thrust::binary_search(thrust::device, beginPQ(), endPQ(), pqBegin, pqEnd, found.begin()); + thrust::lower_bound(thrust::device, beginPQ(), endPQ(), pqBegin, pqEnd, temp.begin()); + thrust::gather_if(thrust::device, temp.begin(), temp.end(), found.begin(), val.begin(), + result.begin()); return result; } diff --git a/utilities/include/structs.h b/utilities/include/structs.h index 78fc9a999..ae3bc5785 100644 --- a/utilities/include/structs.h +++ b/utilities/include/structs.h @@ -414,17 +414,17 @@ struct Box { } }; -// /** -// * Print the contents of this vector to standard output. -// */ -// template -// void Dump(const std::vector& vec) { -// std::cout << "Vec = " << std::endl; -// for (int i = 0; i < vec.size(); ++i) { -// std::cout << i << ", " << vec[i] << ", " << std::endl; -// } -// std::cout << std::endl; -// } +/** + * Print the contents of this vector to standard output. + */ +template +void Dump(const std::vector& vec) { + std::cout << "Vec = " << std::endl; + for (int i = 0; i < vec.size(); ++i) { + std::cout << i << ", " << vec[i] << ", " << std::endl; + } + std::cout << std::endl; +} inline std::ostream& operator<<(std::ostream& stream, const Box& box) { return stream << "min: " << box.min.x << ", " << box.min.y << ", " diff --git a/utilities/include/vec_dh.h b/utilities/include/vec_dh.h index 81efff582..98aa4f242 100644 --- a/utilities/include/vec_dh.h +++ b/utilities/include/vec_dh.h @@ -15,21 +15,11 @@ #pragma once #include #include +#include +#include "structs.h" namespace manifold { -/** @addtogroup Private - * @{ - */ -template -void Dump(const std::vector& vec) { - std::cout << "Vec = " << std::endl; - for (int i = 0; i < vec.size(); ++i) { - std::cout << i << ", " << vec[i] << ", " << std::endl; - } - std::cout << std::endl; -} - /* * Host and device vector implementation. This uses `thrust::universal_vector` for * storage, so data can be moved by the hardware on demand, allows using more @@ -162,60 +152,33 @@ class VecDH { } } - 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() { - syncImpl(); - implModified = true; - return impl_.begin(); - } - - IterH end() { - syncImpl(); - implModified = true; - return impl_.end(); - } - - IterHc cbegin() const { - syncImpl(); - return impl_.cbegin(); - } - - IterHc cend() const { - syncImpl(); - return impl_.cend(); - } - - IterHc begin() const { return cbegin(); } - IterHc end() const { return cend(); } + using Iter = typename thrust::universal_vector::iterator; + using IterC = typename thrust::universal_vector::const_iterator; - IterD beginD() { + Iter begin() { syncImpl(); implModified = true; return impl_.begin(); } - IterD endD() { + Iter end() { syncImpl(); implModified = true; return impl_.end(); } - IterDc cbeginD() const { + IterC cbegin() const { syncImpl(); return impl_.cbegin(); } - IterDc cendD() const { + IterC cend() const { syncImpl(); return impl_.cend(); } - IterDc beginD() const { return cbeginD(); } - IterDc endD() const { return cendD(); } + IterC begin() const { return cbegin(); } + IterC end() const { return cend(); } T* ptrD() { if (size() == 0) return nullptr; From d229f5f2f6118af33b6254eae168018b4b204547 Mon Sep 17 00:00:00 2001 From: pca006132 Date: Tue, 17 May 2022 13:47:00 +0800 Subject: [PATCH 3/3] fixed removed cache --- utilities/include/vec_dh.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/utilities/include/vec_dh.h b/utilities/include/vec_dh.h index 98aa4f242..2eabed682 100644 --- a/utilities/include/vec_dh.h +++ b/utilities/include/vec_dh.h @@ -284,7 +284,7 @@ class VecDH { } void syncCache() const { - if (implModified) { + if (implModified || cache.empty()) { cache = std::vector(impl_.begin(), impl_.end()); } implModified = false;