Merge pull request #1159 from aprokop/apiv2_benchmarks_examples

Switch examples and most benchmarks to APIv2

benchmarks/execution_space_instances/execution_space_instances_driver.cpp

@@ -80,8 +80,8 @@ struct CountCallback
 {
   Kokkos::View<int *, MemorySpace> _counts;
 
-  template <typename Query>
-  KOKKOS_FUNCTION void operator()(Query const &query, int) const
+  template <typename Query, typename Value>
+  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<ExecutionSpace> instance_manager(num_exec_spaces);
   auto const &instances = instance_manager.get_instances();
 
-  std::vector<ArborX::BVH<MemorySpace>> trees;
+  std::vector<ArborX::BoundingVolumeHierarchy<MemorySpace,
+                                              ArborX::PairValueIndex<Point>>>
+      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<int, int>(
-                                                    p * num_primitives,
-                                                    (p + 1) * num_primitives)));
+        exec_space,
+        ArborX::Experimental::attach_indices(Kokkos::subview(
+            primitives, Kokkos::pair<int, int>(p * num_primitives,
+                                               (p + 1) * num_primitives))));
   }
-  ArborX::BVH<MemorySpace> tree(instances[0], primitives);
+  ArborX::BoundingVolumeHierarchy<MemorySpace, ArborX::PairValueIndex<Point>>
+      tree(instances[0], ArborX::Experimental::attach_indices(primitives));
 
   Kokkos::View<int *, MemorySpace> counts("Benchmark::counts",
                                           num_predicates * num_problems);

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<MemorySpace, Triangle> index(
+  ArborX::BoundingVolumeHierarchy<MemorySpace, Triangle> index(
       space, Triangles<MemorySpace>{vertices, triangles});
 
   Kokkos::View<int *, MemorySpace> offset("Benchmark::offsets", 0);

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<Kokkos::CudaSpace> bvh{cuda, PointCloud{d_a, d_a, d_a, N}};
+  ArborX::BoundingVolumeHierarchy<Kokkos::CudaSpace,
+                                  ArborX::PairValueIndex<ArborX::Point<3>>>
+      bvh{cuda,
+          ArborX::Experimental::attach_indices(PointCloud{d_a, d_a, d_a, N})};
 
   Kokkos::View<int *, Kokkos::CudaSpace> indices("Example::indices", 0);
   Kokkos::View<int *, Kokkos::CudaSpace> offset("Example::offset", 0);

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<Kokkos::HostSpace> bvh{Kokkos::DefaultHostExecutionSpace{},
-                                     points};
+  ArborX::BoundingVolumeHierarchy<Kokkos::HostSpace,
+                                  ArborX::PairValueIndex<Point>>
+      bvh{Kokkos::DefaultHostExecutionSpace{},
+          ArborX::Experimental::attach_indices(points)};
 
   // As a supported alternative, wrap the vector in an unmanaged View
-  bvh = ArborX::BVH<Kokkos::HostSpace>{
+  bvh = ArborX::BoundingVolumeHierarchy<Kokkos::HostSpace,
+                                        ArborX::PairValueIndex<Point>>{
       Kokkos::DefaultHostExecutionSpace{},
-      Kokkos::View<Point *, Kokkos::HostSpace, Kokkos::MemoryUnmanaged>{
-          points.data(), points.size()}};
+      ArborX::Experimental::attach_indices(
+          Kokkos::View<Point *, Kokkos::HostSpace, Kokkos::MemoryUnmanaged>{
+              points.data(), points.size()})};
 
   return 0;
 }

examples/brute_force/example_brute_force.cpp

@@ -74,7 +74,9 @@ int main(int argc, char *argv[])
 
   unsigned int out_count;
   {
-    ArborX::BoundingVolumeHierarchy<MemorySpace> bvh{space, primitives};
+    ArborX::BoundingVolumeHierarchy<MemorySpace,
+                                    ArborX::PairValueIndex<ArborX::Point<3>>>
+        bvh{space, ArborX::Experimental::attach_indices(primitives)};
 
     Kokkos::View<int *, ExecutionSpace> indices("Example::indices_ref", 0);
     Kokkos::View<int *, ExecutionSpace> offset("Example::offset_ref", 0);

examples/callback/example_callback.cpp

@@ -52,12 +52,13 @@ struct ArborX::AccessTraits<NearestToOrigin, ArborX::PredicatesTag>
 
 struct PrintfCallback
 {
-  template <typename Predicate, typename OutputFunctor>
-  KOKKOS_FUNCTION void operator()(Predicate, int primitive,
+  template <typename Predicate, typename Value, typename OutputFunctor>
+  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<MemorySpace> bvh{
-      ExecutionSpace{},
-      Kokkos::create_mirror_view_and_copy(
-          MemorySpace{},
-          Kokkos::View<Point *, Kokkos::HostSpace, Kokkos::MemoryUnmanaged>(
-              points.data(), points.size()))};
+  ArborX::BoundingVolumeHierarchy<MemorySpace, ArborX::PairValueIndex<Point>>
+      bvh{ExecutionSpace{},
+          ArborX::Experimental::attach_indices(
+              Kokkos::create_mirror_view_and_copy(
+                  MemorySpace{}, Kokkos::View<Point *, Kokkos::HostSpace,
+                                              Kokkos::MemoryUnmanaged>(
+                                     points.data(), points.size())))};
 
   {
     Kokkos::View<int *, MemorySpace> 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
   }

examples/molecular_dynamics/example_molecular_dynamics.cpp

@@ -40,10 +40,12 @@ struct ArborX::AccessTraits<Neighbors<MemorySpace>, ArborX::PredicatesTag>
 
 struct ExcludeSelfCollision
 {
-  template <class Predicate, class OutputFunctor>
-  KOKKOS_FUNCTION void operator()(Predicate const &predicate, int i,
+  template <class Predicate, typename Value, class OutputFunctor>
+  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<MemorySpace> index(execution_space, particles);
+  ArborX::BoundingVolumeHierarchy<MemorySpace,
+                                  ArborX::PairValueIndex<ArborX::Point<3>>>
+      index(execution_space, ArborX::Experimental::attach_indices(particles));
 
   Kokkos::View<int *, MemorySpace> indices("Example::indices", 0);
   Kokkos::View<int *, MemorySpace> offsets("Example::offsets", 0);

examples/raytracing/example_raytracing.cpp

@@ -64,12 +64,13 @@ struct DepositEnergy
   Kokkos::View<float *, MemorySpace> _ray_energy;
   Kokkos::View<float *, MemorySpace> _energy;
 
-  template <typename Predicate>
+  template <typename Predicate, typename Value>
   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<ArborX::Box<3> *, MemorySpace> _boxes;
 
-  template <typename Predicate, typename OutputFunctor>
+  template <typename Predicate, typename Value, typename OutputFunctor>
   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<MemorySpace> bvh{exec_space, boxes};
+  ArborX::BoundingVolumeHierarchy<MemorySpace,
+                                  ArborX::PairValueIndex<ArborX::Box<3>>>
+      bvh{exec_space, ArborX::Experimental::attach_indices(boxes)};
 
   // OrderedIntersects-based approach
   Kokkos::View<float *, MemorySpace> energy_ordered_intersects;

examples/simple_intersection/example_intersection.cpp

@@ -50,8 +50,9 @@ int main(int argc, char *argv[])
 
   ExecutionSpace space;
 
-  ArborX::BVH<MemorySpace, ArborX::PairValueIndex<Box>> const tree(
-      space, ArborX::Experimental::attach_indices(boxes));
+  ArborX::BoundingVolumeHierarchy<MemorySpace,
+                                  ArborX::PairValueIndex<Box>> const
+      tree(space, ArborX::Experimental::attach_indices(boxes));
 
   // The query will resize indices and offsets accordingly
   Kokkos::View<int *, MemorySpace> indices("Example::indices", 0);

examples/triangle_intersection/triangle_intersection.cpp

@@ -84,8 +84,9 @@ int main(int argc, char *argv[])
 
   ExecutionSpace space;
 
-  ArborX::BVH<MemorySpace, ArborX::PairValueIndex<Triangle>> const tree(
-      space, ArborX::Experimental::attach_indices(triangles));
+  ArborX::BoundingVolumeHierarchy<MemorySpace,
+                                  ArborX::PairValueIndex<Triangle>> const
+      tree(space, ArborX::Experimental::attach_indices(triangles));
 
   // The query will resize indices and offsets accordingly
   Kokkos::View<unsigned *, MemorySpace> indices("Example::indices", 0);

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<Point *, DeviceType> points("Example::points", 0);
   loadPointCloud(infile, points);
 
-  ArborX::BVH<Kokkos::HostSpace> bvh{ExecutionSpace{}, points};
+  Kokkos::View<Box *, DeviceType> boxes("Example::boxes", points.size());
+  Kokkos::parallel_for(
+      "Example::copy_points_to_boxes",
+      Kokkos::RangePolicy<ExecutionSpace>(0, points.size()),
+      KOKKOS_LAMBDA(int i) {
+        boxes(i) = {points(i), points(i)};
+      });
+
+  ArborX::BoundingVolumeHierarchy<Kokkos::HostSpace,
+                                  ArborX::PairValueIndex<Box>>
+      bvh{ExecutionSpace{}, ArborX::Experimental::attach_indices(boxes)};
 
   using TreeVisualization = ArborX::Details::TreeVisualization;
   using TikZVisitor = typename TreeVisualization::TikZVisitor;

src/geometry/ArborX_DetailsAlgorithms.hpp

@@ -535,15 +535,24 @@ template <typename Point, typename Box>
 struct intersects<PointTag, BoxTag, Point, Box>
 {
   KOKKOS_FUNCTION static constexpr bool apply(Point const &point,
-                                              Box const &other)
+                                              Box const &box)
   {
     constexpr int DIM = GeometryTraits::dimension_v<Point>;
     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 <typename Box, typename Point>
+struct intersects<BoxTag, PointTag, Box, Point>
+{
+  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 <typename Sphere, typename Box>