Skip to content

Commit

Permalink
universal vector
Browse files Browse the repository at this point in the history
  • Loading branch information
@pca006132
pca006132 committed May 14, 2022
1 parent fadad77 commit 933eade
Show file tree
Hide file tree
Showing 13 changed files with 153 additions and 195 deletions.
12 changes: 6 additions & 6 deletions collider/src/collider.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -252,7 +252,7 @@ Collider::Collider(const VecDH<Box>& 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);
Expand All @@ -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(
zip(querriesIn.cbeginD(), countAt(0)), querriesIn.size(),
thrust::device, zip(querriesIn.cbeginD(), countAt(0)), querriesIn.size(),
FindCollisions<T>({querryTri.ptrDpq(), nOverlapsD.ptrD(), maxOverlaps,
nodeBBox_.ptrD(), internalChildren_.ptrD()}));
nOverlaps = nOverlapsD.H()[0];
Expand Down Expand Up @@ -305,13 +305,13 @@ void Collider::UpdateBoxes(const VecDH<Box>& leafBB) {
// copy in leaf node Boxs
strided_range<VecDH<Box>::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<int> 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()}));
}
Expand All @@ -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;
Expand Down
21 changes: 11 additions & 10 deletions manifold/src/boolean3.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -14,6 +14,7 @@

#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 @@ -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();
Expand Down Expand Up @@ -246,7 +247,7 @@ std::tuple<VecDH<int>, VecDH<glm::vec4>> Shadow11(SparseIndices &p1q1,
VecDH<glm::vec4> 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,
Expand Down Expand Up @@ -343,7 +344,7 @@ std::tuple<VecDH<int>, VecDH<float>> 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(),
Expand Down Expand Up @@ -453,7 +454,7 @@ std::tuple<VecDH<int>, VecDH<glm::vec3>> Intersect12(
VecDH<glm::vec3> 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(),
Expand All @@ -471,19 +472,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(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<int> w03val(w03.size());
VecDH<int> 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<int>());
return w03;
};
Expand Down
51 changes: 26 additions & 25 deletions manifold/src/boolean_result.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -14,6 +14,7 @@

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

#include "boolean3.h"
#include "polygon.h"
Expand Down Expand Up @@ -83,45 +84,45 @@ std::tuple<VecDH<int>, VecDH<int>> 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<int> facePQ2R(inP.NumTri() + inQ.NumTri() + 1);
auto keepFace =
thrust::make_transform_iterator(sidesPerFacePQ.beginD(), NotZero());
thrust::inclusive_scan(keepFace, keepFace + sidesPerFacePQ.size(),
thrust::inclusive_scan(thrust::device, keepFace, keepFace + sidesPerFacePQ.size(),
facePQ2R.beginD() + 1);
int numFaceR = facePQ2R.H().back();
facePQ2R.resize(inP.NumTri() + inQ.NumTri());

outR.faceNormal_.resize(numFaceR);
auto next = thrust::copy_if(inP.faceNormal_.beginD(), inP.faceNormal_.endD(),
auto next = thrust::copy_if(thrust::device, inP.faceNormal_.beginD(), inP.faceNormal_.endD(),
keepFace, outR.faceNormal_.beginD(),
thrust::identity<bool>());
if (invertQ) {
auto start = thrust::make_transform_iterator(inQ.faceNormal_.beginD(),
thrust::negate<glm::vec3>());
auto end = thrust::make_transform_iterator(inQ.faceNormal_.endD(),
thrust::negate<glm::vec3>());
thrust::copy_if(start, end, keepFace + inP.NumTri(), next,
thrust::copy_if(thrust::device, start, end, keepFace + inP.NumTri(), next,
thrust::identity<bool>());
} 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<bool>());
}

auto newEnd =
thrust::remove(sidesPerFacePQ.beginD(), sidesPerFacePQ.endD(), 0);
thrust::remove(thrust::device, sidesPerFacePQ.beginD(), sidesPerFacePQ.endD(), 0);
VecDH<int> 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());

Expand Down Expand Up @@ -211,7 +212,7 @@ void AppendPartialEdges(Manifold::Impl &outR, VecH<bool> &wholeHalfedgeP,
std::map<int, std::vector<EdgePos>> &edgesP,
VecH<Ref> &halfedgeRef, const Manifold::Impl &inP,
const VecH<int> &i03, const VecH<int> &vP2R,
const thrust::host_vector<int>::const_iterator faceP2R,
const VecDH<int>::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),
Expand Down Expand Up @@ -406,7 +407,7 @@ void AppendWholeEdges(Manifold::Impl &outR, VecDH<int> &facePtrR,
const VecDH<int> &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(),
Expand Down Expand Up @@ -482,7 +483,7 @@ std::pair<VecDH<BaryRef>, VecDH<int>> CalculateMeshRelation(
VecDH<int> halfedgeBary(halfedgeRef.size());
VecDH<int> 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(
Expand Down Expand Up @@ -550,33 +551,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(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<int> vP2R(inP_.NumVert());
thrust::exclusive_scan(i03.beginD(), i03.endD(), vP2R.beginD(), 0, AbsSum());
thrust::exclusive_scan(thrust::device, i03.beginD(), i03.endD(), vP2R.beginD(), 0, AbsSum());
int numVertR = AbsSum()(vP2R.H().back(), i03.H().back());
const int nPv = numVertR;

VecDH<int> vQ2R(inQ_.NumVert());
thrust::exclusive_scan(i30.beginD(), i30.endD(), vQ2R.beginD(), numVertR,
thrust::exclusive_scan(thrust::device, i30.beginD(), i30.endD(), vQ2R.beginD(), numVertR,
AbsSum());
numVertR = AbsSum()(vQ2R.H().back(), i30.H().back());
const int nQv = numVertR - nPv;

VecDH<int> 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());
}
const int n12 = numVertR - nPv - nQv;

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

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

#include "graph.h"
#include "impl.h"
Expand Down Expand Up @@ -183,7 +184,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(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.
Expand Down Expand Up @@ -403,20 +404,20 @@ Manifold Manifold::Compose(const std::vector<Manifold>& 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}));

Expand Down Expand Up @@ -461,7 +462,7 @@ std::vector<Manifold> Manifold::Decompose() const {
VecDH<int> 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()),
Expand All @@ -470,11 +471,11 @@ std::vector<Manifold> Manifold::Decompose() const {
meshes[i].pImpl_->vertPos_.resize(nVert);

VecDH<int> 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);
Expand Down
7 changes: 4 additions & 3 deletions manifold/src/edge_op.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -13,6 +13,7 @@
// limitations under the License.

#include "impl.h"
#include <thrust/execution_policy.h>

namespace {
using namespace manifold;
Expand Down Expand Up @@ -121,7 +122,7 @@ void Manifold::Impl::SimplifyTopology() {
VecDH<int> 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);
Expand All @@ -131,7 +132,7 @@ void Manifold::Impl::SimplifyTopology() {
flaggedEdges.resize(halfedge_.size());
numFlagged =
thrust::copy_if(
countAt(0), countAt(halfedge_.size()), flaggedEdges.beginD(),
thrust::device, countAt(0), countAt(halfedge_.size()), flaggedEdges.beginD(),
FlagEdge({halfedge_.cptrD(), meshRelation_.triBary.cptrD()})) -
flaggedEdges.beginD();
flaggedEdges.resize(numFlagged);
Expand All @@ -140,7 +141,7 @@ void Manifold::Impl::SimplifyTopology() {

flaggedEdges.resize(halfedge_.size());
numFlagged = thrust::copy_if(
countAt(0), countAt(halfedge_.size()), flaggedEdges.beginD(),
thrust::device, countAt(0), countAt(halfedge_.size()), flaggedEdges.beginD(),
SwappableEdge({halfedge_.cptrD(), vertPos_.cptrD(),
faceNormal_.cptrD(), precision_})) -
flaggedEdges.beginD();
Expand Down
Loading

0 comments on commit 933eade

Please sign in to comment.