Skip to content

Commit

Permalink
vec_dh cleanup
Browse files Browse the repository at this point in the history
  • Loading branch information
@pca006132
pca006132 committed May 16, 2022
1 parent c316836 commit dbbcdbb
Show file tree
Hide file tree
Showing 15 changed files with 271 additions and 332 deletions.
10 changes: 5 additions & 5 deletions collider/src/collider.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -274,7 +274,7 @@ SparseIndices Collider::Collisions(const VecDH<T>& querriesIn) const {
VecDH<int> 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<T>({querryTri.ptrDpq(), nOverlapsD.ptrD(), maxOverlaps,
nodeBBox_.ptrD(), internalChildren_.ptrD()}));
nOverlaps = nOverlapsD[0];
Expand Down Expand Up @@ -303,12 +303,12 @@ void Collider::UpdateBoxes(const VecDH<Box>& leafBB) {
ALWAYS_ASSERT(leafBB.size() == NumLeaves(), userErr,
"must have the same number of updated boxes as original");
// copy in leaf node Boxs
strided_range<VecDH<Box>::IterD> leaves(nodeBBox_.beginD(), nodeBBox_.endD(),
strided_range<VecDH<Box>::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<int> 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(),
Expand All @@ -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;
Expand Down
29 changes: 14 additions & 15 deletions manifold/src/boolean3.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,6 @@

#include "boolean3.h"
#include <limits>
#include <thrust/execution_policy.h>

// TODO: make this runtime configurable for quicker debug
constexpr bool kVerbose = false;
Expand Down Expand Up @@ -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();
Expand Down Expand Up @@ -247,7 +246,7 @@ std::tuple<VecDH<int>, VecDH<glm::vec4>> Shadow11(SparseIndices &p1q1,
VecDH<glm::vec4> 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,
Expand Down Expand Up @@ -344,8 +343,8 @@ std::tuple<VecDH<int>, VecDH<float>> 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}));
Expand Down Expand Up @@ -454,8 +453,8 @@ std::tuple<VecDH<int>, VecDH<glm::vec3>> Intersect12(
VecDH<glm::vec3> 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(),
Expand All @@ -472,19 +471,19 @@ VecDH<int> Winding03(const Manifold::Impl &inP, SparseIndices &p0q2,
// verts that are not shadowed (not in p0q2) have winding number zero.
VecDH<int> 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<int> w03val(w03.size());
VecDH<int> 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<int>());
return w03;
};
Expand Down
63 changes: 31 additions & 32 deletions manifold/src/boolean_result.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,6 @@

#include <algorithm>
#include <map>
#include <thrust/execution_policy.h>

#include "boolean3.h"
#include "polygon.h"
Expand Down Expand Up @@ -84,46 +83,46 @@ std::tuple<VecDH<int>, VecDH<int>> 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<int> 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<bool>());
if (invertQ) {
auto start = thrust::make_transform_iterator(inQ.faceNormal_.beginD(),
auto start = thrust::make_transform_iterator(inQ.faceNormal_.begin(),
thrust::negate<glm::vec3>());
auto end = thrust::make_transform_iterator(inQ.faceNormal_.endD(),
auto end = thrust::make_transform_iterator(inQ.faceNormal_.end(),
thrust::negate<glm::vec3>());
thrust::copy_if(thrust::device, start, end, keepFace + inP.NumTri(), next,
thrust::identity<bool>());
} 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<bool>());
}

auto newEnd =
thrust::remove(thrust::device, sidesPerFacePQ.beginD(), sidesPerFacePQ.endD(), 0);
VecDH<int> 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<int> 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);
Expand Down Expand Up @@ -212,7 +211,7 @@ void AppendPartialEdges(Manifold::Impl &outR, VecDH<char> &wholeHalfedgeP,
std::map<int, std::vector<EdgePos>> &edgesP,
VecDH<Ref> &halfedgeRef, const Manifold::Impl &inP,
const VecDH<int> &i03, const VecDH<int> &vP2R,
const VecDH<int>::IterHc faceP2R,
const VecDH<int>::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),
Expand Down Expand Up @@ -407,7 +406,7 @@ void AppendWholeEdges(Manifold::Impl &outR, VecDH<int> &facePtrR,
const VecDH<int> &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(),
Expand Down Expand Up @@ -483,8 +482,8 @@ std::pair<VecDH<BaryRef>, VecDH<int>> CalculateMeshRelation(
VecDH<int> halfedgeBary(halfedgeRef.size());
VecDH<int> 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(),
Expand Down Expand Up @@ -551,33 +550,33 @@ Manifold::Impl Boolean3::Result(Manifold::OpType op) const {
VecDH<int> i21(x21_.size());
VecDH<int> i03(w03_.size());
VecDH<int> 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<int> 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<int> 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<int> 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());
}
const int n12 = numVertR - nPv - nQv;

VecDH<int> 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());
}
Expand All @@ -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) {
Expand Down
47 changes: 23 additions & 24 deletions manifold/src/constructors.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,6 @@
// limitations under the License.

#include <thrust/sequence.h>
#include <thrust/execution_policy.h>

#include "graph.h"
#include "impl.h"
Expand Down Expand Up @@ -184,7 +183,7 @@ Manifold Manifold::Sphere(float radius, int circularSegments) {
Manifold sphere;
sphere.pImpl_ = std::make_unique<Impl>(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.
Expand Down Expand Up @@ -404,21 +403,21 @@ Manifold Manifold::Compose(const std::vector<Manifold>& 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();
Expand Down Expand Up @@ -462,22 +461,22 @@ std::vector<Manifold> Manifold::Decompose() const {
VecDH<int> 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<int> 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);
Expand Down
Loading

0 comments on commit dbbcdbb

Please sign in to comment.