Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

universal vector #121

Merged
merged 4 commits into from
May 17, 2022
Merged

universal vector #121

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
c316836
universal vector
pca006132 May 14, 2022
dbbcdbb
vec_dh cleanup
pca006132 May 16, 2022
2a86198
Merge branch 'master' into universal-vector
pca006132 May 16, 2022
d229f5f
fixed removed cache
pca006132 May 17, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
16 changes: 8 additions & 8 deletions collider/src/collider.cpp
View file
Open in desktop
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.cbeginD(), 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 @@ -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
Open in desktop
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>
pca006132 marked this conversation as resolved.
Show resolved Hide resolved

// 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
Loading