Skip to content

Commit

Permalink
Merge pull request #121 from pca006132/universal-vector
Browse files Browse the repository at this point in the history 
universal vector
  • Loading branch information
@elalish
elalish authored May 17, 2022
2 parents 8f21bc5 + d229f5f commit bfc8e5c
Show file tree
Hide file tree
Showing 16 changed files with 578 additions and 498 deletions.
18 changes: 9 additions & 9 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,14 +274,14 @@ 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.cbegin(), countAt(0)), querriesIn.size(),
FindCollisions<T>({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<float>(querriesIn.size()) / lastQuery;
if (ratio > 1000) // do not trust the ratio if it is too large
maxOverlaps *= 2;
Expand All @@ -303,15 +303,15 @@ 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(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(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(
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_.begin(), nodeBBox_.end(),
TransformBox({transform}));
}
return axisAligned;
Expand Down
28 changes: 14 additions & 14 deletions manifold/src/boolean3.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -85,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(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(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 @@ -246,7 +246,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.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 @@ -343,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(
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 @@ -453,8 +453,8 @@ 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),
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 @@ -471,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(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.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(p0q2.beginD(reverse), p0q2.endD(reverse),
s02.beginD(), w03vert.beginD(), w03val.beginD());
thrust::scatter(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(w03.beginD(), w03.endD(), w03.beginD(),
thrust::transform(thrust::device, w03.begin(), w03.end(), w03.begin(),
thrust::negate<int>());
return w03;
};
Expand Down
Loading

0 comments on commit bfc8e5c

Please sign in to comment.