From a416c020aada38ec905127104b4244def2aa9425 Mon Sep 17 00:00:00 2001 From: Andrey Prokopenko Date: Tue, 24 Sep 2024 20:37:09 -0400 Subject: [PATCH 1/2] Switch examples to APIv2 --- .../example_cuda_access_traits.cpp | 5 ++- .../example_host_access_traits.cpp | 14 +++++--- examples/brute_force/example_brute_force.cpp | 4 ++- examples/callback/example_callback.cpp | 36 +++++++++---------- .../example_molecular_dynamics.cpp | 10 ++++-- examples/raytracing/example_raytracing.cpp | 14 +++++--- .../example_intersection.cpp | 5 +-- .../triangle_intersection.cpp | 5 +-- examples/viz/tree_visualization.cpp | 13 ++++++- src/geometry/ArborX_DetailsAlgorithms.hpp | 13 +++++-- 10 files changed, 79 insertions(+), 40 deletions(-) diff --git a/examples/access_traits/example_cuda_access_traits.cpp b/examples/access_traits/example_cuda_access_traits.cpp index e07d1c999..5d9fb7e4a 100644 --- a/examples/access_traits/example_cuda_access_traits.cpp +++ b/examples/access_traits/example_cuda_access_traits.cpp @@ -79,7 +79,10 @@ int main(int argc, char *argv[]) cudaMemcpyAsync(d_a, a.data(), sizeof(a), cudaMemcpyHostToDevice, stream); Kokkos::Cuda cuda{stream}; - ArborX::BVH bvh{cuda, PointCloud{d_a, d_a, d_a, N}}; + ArborX::BoundingVolumeHierarchy>> + bvh{cuda, + ArborX::Experimental::attach_indices(PointCloud{d_a, d_a, d_a, N})}; Kokkos::View indices("Example::indices", 0); Kokkos::View offset("Example::offset", 0); diff --git a/examples/access_traits/example_host_access_traits.cpp b/examples/access_traits/example_host_access_traits.cpp index eb99c29b0..048269636 100644 --- a/examples/access_traits/example_host_access_traits.cpp +++ b/examples/access_traits/example_host_access_traits.cpp @@ -41,14 +41,18 @@ int main(int argc, char *argv[]) }); // Pass directly the vector of points to use the access traits defined above - ArborX::BVH bvh{Kokkos::DefaultHostExecutionSpace{}, - points}; + ArborX::BoundingVolumeHierarchy> + bvh{Kokkos::DefaultHostExecutionSpace{}, + ArborX::Experimental::attach_indices(points)}; // As a supported alternative, wrap the vector in an unmanaged View - bvh = ArborX::BVH{ + bvh = ArborX::BoundingVolumeHierarchy>{ Kokkos::DefaultHostExecutionSpace{}, - Kokkos::View{ - points.data(), points.size()}}; + ArborX::Experimental::attach_indices( + Kokkos::View{ + points.data(), points.size()})}; return 0; } diff --git a/examples/brute_force/example_brute_force.cpp b/examples/brute_force/example_brute_force.cpp index 3c3c29251..ae876835c 100644 --- a/examples/brute_force/example_brute_force.cpp +++ b/examples/brute_force/example_brute_force.cpp @@ -74,7 +74,9 @@ int main(int argc, char *argv[]) unsigned int out_count; { - ArborX::BoundingVolumeHierarchy bvh{space, primitives}; + ArborX::BoundingVolumeHierarchy>> + bvh{space, ArborX::Experimental::attach_indices(primitives)}; Kokkos::View indices("Example::indices_ref", 0); Kokkos::View offset("Example::offset_ref", 0); diff --git a/examples/callback/example_callback.cpp b/examples/callback/example_callback.cpp index c738434da..6772c4b32 100644 --- a/examples/callback/example_callback.cpp +++ b/examples/callback/example_callback.cpp @@ -52,12 +52,13 @@ struct ArborX::AccessTraits struct PrintfCallback { - template - KOKKOS_FUNCTION void operator()(Predicate, int primitive, + template + KOKKOS_FUNCTION void operator()(Predicate, Value const &value, OutputFunctor const &out) const { - Kokkos::printf("Found %d from functor\n", primitive); - out(primitive); + auto const index = value.index; + Kokkos::printf("Found %d from functor\n", index); + out(index); } }; @@ -77,12 +78,13 @@ int main(int argc, char *argv[]) return Point{rd(), rd(), rd()}; }); - ArborX::BVH bvh{ - ExecutionSpace{}, - Kokkos::create_mirror_view_and_copy( - MemorySpace{}, - Kokkos::View( - points.data(), points.size()))}; + ArborX::BoundingVolumeHierarchy> + bvh{ExecutionSpace{}, + ArborX::Experimental::attach_indices( + Kokkos::create_mirror_view_and_copy( + MemorySpace{}, Kokkos::View( + points.data(), points.size())))}; { Kokkos::View values("Example::values", 0); @@ -92,9 +94,8 @@ int main(int argc, char *argv[]) #ifndef __NVCC__ ArborX::query( bvh, ExecutionSpace{}, FirstOctant{}, - KOKKOS_LAMBDA(auto /*predicate*/, int primitive, - auto /*output_functor*/) { - Kokkos::printf("Found %d from generic lambda\n", primitive); + KOKKOS_LAMBDA(auto /*predicate*/, auto value, auto /*output_functor*/) { + Kokkos::printf("Found %d from generic lambda\n", value.index); }, values, offsets); #endif @@ -109,9 +110,8 @@ int main(int argc, char *argv[]) #ifndef __NVCC__ ArborX::query( bvh, ExecutionSpace{}, NearestToOrigin{k}, - KOKKOS_LAMBDA(auto /*predicate*/, int primitive, - auto /*output_functor*/) { - Kokkos::printf("Found %d from generic lambda\n", primitive); + KOKKOS_LAMBDA(auto /*predicate*/, auto value, auto /*output_functor*/) { + Kokkos::printf("Found %d from generic lambda\n", value.index); }, values, offsets); #endif @@ -125,8 +125,8 @@ int main(int argc, char *argv[]) #ifndef __NVCC__ bvh.query( ExecutionSpace{}, FirstOctant{}, - KOKKOS_LAMBDA(auto /*predicate*/, int j) { - Kokkos::printf("%d %d %d\n", ++c(), -1, j); + KOKKOS_LAMBDA(auto /*predicate*/, auto value) { + Kokkos::printf("%d %d %d\n", ++c(), -1, value.index); }); #endif } diff --git a/examples/molecular_dynamics/example_molecular_dynamics.cpp b/examples/molecular_dynamics/example_molecular_dynamics.cpp index 413d407d0..f772c3805 100644 --- a/examples/molecular_dynamics/example_molecular_dynamics.cpp +++ b/examples/molecular_dynamics/example_molecular_dynamics.cpp @@ -40,10 +40,12 @@ struct ArborX::AccessTraits, ArborX::PredicatesTag> struct ExcludeSelfCollision { - template - KOKKOS_FUNCTION void operator()(Predicate const &predicate, int i, + template + KOKKOS_FUNCTION void operator()(Predicate const &predicate, + Value const &value, OutputFunctor const &out) const { + int const i = value.index; int const j = getData(predicate); if (i != j) { @@ -115,7 +117,9 @@ int main(int argc, char *argv[]) // TODO scale velocities Kokkos::Profiling::popRegion(); - ArborX::BVH index(execution_space, particles); + ArborX::BoundingVolumeHierarchy>> + index(execution_space, ArborX::Experimental::attach_indices(particles)); Kokkos::View indices("Example::indices", 0); Kokkos::View offsets("Example::offsets", 0); diff --git a/examples/raytracing/example_raytracing.cpp b/examples/raytracing/example_raytracing.cpp index ea54b77f2..111d52ed8 100644 --- a/examples/raytracing/example_raytracing.cpp +++ b/examples/raytracing/example_raytracing.cpp @@ -64,12 +64,13 @@ struct DepositEnergy Kokkos::View _ray_energy; Kokkos::View _energy; - template + template KOKKOS_FUNCTION void operator()(Predicate const &predicate, - int const primitive_index) const + Value const &value) const { float length; float entrylength; + int const primitive_index = value.index; auto const &ray = ArborX::getGeometry(predicate); auto const &box = _boxes(primitive_index); int const predicate_index = ArborX::getData(predicate); @@ -149,13 +150,14 @@ struct AccumulateRaySphereIntersections { Kokkos::View *, MemorySpace> _boxes; - template + template KOKKOS_FUNCTION void operator()(Predicate const &predicate, - int const primitive_index, + Value const &value, OutputFunctor const &out) const { float length; float entrylength; + int const primitive_index = value.index; auto const &ray = ArborX::getGeometry(predicate); auto const &box = _boxes(primitive_index); int const predicate_index = ArborX::getData(predicate); @@ -284,7 +286,9 @@ int main(int argc, char *argv[]) Kokkos::Profiling::popRegion(); // Construct BVH - ArborX::BVH bvh{exec_space, boxes}; + ArborX::BoundingVolumeHierarchy>> + bvh{exec_space, ArborX::Experimental::attach_indices(boxes)}; // OrderedIntersects-based approach Kokkos::View energy_ordered_intersects; diff --git a/examples/simple_intersection/example_intersection.cpp b/examples/simple_intersection/example_intersection.cpp index f0f8c12ef..c090238c6 100644 --- a/examples/simple_intersection/example_intersection.cpp +++ b/examples/simple_intersection/example_intersection.cpp @@ -50,8 +50,9 @@ int main(int argc, char *argv[]) ExecutionSpace space; - ArborX::BVH> const tree( - space, ArborX::Experimental::attach_indices(boxes)); + ArborX::BoundingVolumeHierarchy> const + tree(space, ArborX::Experimental::attach_indices(boxes)); // The query will resize indices and offsets accordingly Kokkos::View indices("Example::indices", 0); diff --git a/examples/triangle_intersection/triangle_intersection.cpp b/examples/triangle_intersection/triangle_intersection.cpp index f68866a2a..8e99b99d4 100644 --- a/examples/triangle_intersection/triangle_intersection.cpp +++ b/examples/triangle_intersection/triangle_intersection.cpp @@ -84,8 +84,9 @@ int main(int argc, char *argv[]) ExecutionSpace space; - ArborX::BVH> const tree( - space, ArborX::Experimental::attach_indices(triangles)); + ArborX::BoundingVolumeHierarchy> const + tree(space, ArborX::Experimental::attach_indices(triangles)); // The query will resize indices and offsets accordingly Kokkos::View indices("Example::indices", 0); diff --git a/examples/viz/tree_visualization.cpp b/examples/viz/tree_visualization.cpp index ec31d04b7..f16708e65 100644 --- a/examples/viz/tree_visualization.cpp +++ b/examples/viz/tree_visualization.cpp @@ -90,11 +90,22 @@ void viz(std::string const &prefix, std::string const &infile, int n_neighbors) using DeviceType = ExecutionSpace::device_type; using Point = ArborX::Point<3>; + using Box = ArborX::Box<3>; Kokkos::View points("Example::points", 0); loadPointCloud(infile, points); - ArborX::BVH bvh{ExecutionSpace{}, points}; + Kokkos::View boxes("Example::boxes", points.size()); + Kokkos::parallel_for( + "Example::copy_points_to_boxes", + Kokkos::RangePolicy(0, points.size()), + KOKKOS_LAMBDA(int i) { + boxes(i) = {points(i), points(i)}; + }); + + ArborX::BoundingVolumeHierarchy> + bvh{ExecutionSpace{}, ArborX::Experimental::attach_indices(boxes)}; using TreeVisualization = ArborX::Details::TreeVisualization; using TikZVisitor = typename TreeVisualization::TikZVisitor; diff --git a/src/geometry/ArborX_DetailsAlgorithms.hpp b/src/geometry/ArborX_DetailsAlgorithms.hpp index 0539d0db3..b3599adc9 100644 --- a/src/geometry/ArborX_DetailsAlgorithms.hpp +++ b/src/geometry/ArborX_DetailsAlgorithms.hpp @@ -521,15 +521,24 @@ template struct intersects { KOKKOS_FUNCTION static constexpr bool apply(Point const &point, - Box const &other) + Box const &box) { constexpr int DIM = GeometryTraits::dimension_v; for (int d = 0; d < DIM; ++d) - if (point[d] > other.maxCorner()[d] || point[d] < other.minCorner()[d]) + if (point[d] > box.maxCorner()[d] || point[d] < box.minCorner()[d]) return false; return true; } }; +template +struct intersects +{ + KOKKOS_FUNCTION static constexpr bool apply(Box const &box, + Point const &point) + { + return Details::intersects(point, box); + } +}; // check if a sphere intersects with an axis-aligned bounding box template From 79f734bd2bd7bfd9a5e49cc2913eb6485653df71 Mon Sep 17 00:00:00 2001 From: Andrey Prokopenko Date: Tue, 24 Sep 2024 20:43:55 -0400 Subject: [PATCH 2/2] Switch most of benchmarks to APIv2 --- .../execution_space_instances_driver.cpp | 18 +++++++++++------- .../triangulated_surface_distance.cpp | 2 +- 2 files changed, 12 insertions(+), 8 deletions(-) diff --git a/benchmarks/execution_space_instances/execution_space_instances_driver.cpp b/benchmarks/execution_space_instances/execution_space_instances_driver.cpp index 5197153b6..c49b825b6 100644 --- a/benchmarks/execution_space_instances/execution_space_instances_driver.cpp +++ b/benchmarks/execution_space_instances/execution_space_instances_driver.cpp @@ -80,8 +80,8 @@ struct CountCallback { Kokkos::View _counts; - template - KOKKOS_FUNCTION void operator()(Query const &query, int) const + template + KOKKOS_FUNCTION void operator()(Query const &query, Value const &) const { auto const i = ArborX::getData(query); Kokkos::atomic_increment(&_counts(i)); @@ -184,17 +184,21 @@ int main(int argc, char *argv[]) InstanceManager instance_manager(num_exec_spaces); auto const &instances = instance_manager.get_instances(); - std::vector> trees; + std::vector>> + trees; for (int p = 0; p < num_problems; ++p) { auto const &exec_space = instances[p % num_exec_spaces]; trees.emplace_back( - exec_space, Kokkos::subview(primitives, Kokkos::pair( - p * num_primitives, - (p + 1) * num_primitives))); + exec_space, + ArborX::Experimental::attach_indices(Kokkos::subview( + primitives, Kokkos::pair(p * num_primitives, + (p + 1) * num_primitives)))); } - ArborX::BVH tree(instances[0], primitives); + ArborX::BoundingVolumeHierarchy> + tree(instances[0], ArborX::Experimental::attach_indices(primitives)); Kokkos::View counts("Benchmark::counts", num_predicates * num_problems); diff --git a/benchmarks/triangulated_surface_distance/triangulated_surface_distance.cpp b/benchmarks/triangulated_surface_distance/triangulated_surface_distance.cpp index 526e19f89..a820aa7f2 100644 --- a/benchmarks/triangulated_surface_distance/triangulated_surface_distance.cpp +++ b/benchmarks/triangulated_surface_distance/triangulated_surface_distance.cpp @@ -166,7 +166,7 @@ int main(int argc, char *argv[]) std::cout << "#triangles : " << triangles.size() << '\n'; std::cout << "#queries : " << random_points.size() << '\n'; - ArborX::BVH index( + ArborX::BoundingVolumeHierarchy index( space, Triangles{vertices, triangles}); Kokkos::View offset("Benchmark::offsets", 0);