From 45a90f0930efa12d8d84ac568618fabf1b5115cc Mon Sep 17 00:00:00 2001 From: Andrey Prokopenko Date: Mon, 14 Oct 2024 10:39:56 -0400 Subject: [PATCH 1/5] Add deduction guides for BVH and BruteForce taking in indexable getter --- src/spatial/ArborX_BruteForce.hpp | 11 +++++++++++ src/spatial/ArborX_LinearBVH.hpp | 12 ++++++++++++ 2 files changed, 23 insertions(+) diff --git a/src/spatial/ArborX_BruteForce.hpp b/src/spatial/ArborX_BruteForce.hpp index eefd8c4ea..09dd2b52b 100644 --- a/src/spatial/ArborX_BruteForce.hpp +++ b/src/spatial/ArborX_BruteForce.hpp @@ -136,6 +136,17 @@ KOKKOS_FUNCTION typename Details::AccessValues::memory_space, typename Details::AccessValues::value_type>; +template +#if KOKKOS_VERSION >= 40400 +KOKKOS_DEDUCTION_GUIDE +#else +KOKKOS_FUNCTION +#endif + BruteForce(ExecutionSpace, Values, IndexableGetter) -> BruteForce< + typename Details::AccessValues::memory_space, + typename Details::AccessValues::value_type, + IndexableGetter>; + template template diff --git a/src/spatial/ArborX_LinearBVH.hpp b/src/spatial/ArborX_LinearBVH.hpp index e6f2cafff..29fa18fd4 100644 --- a/src/spatial/ArborX_LinearBVH.hpp +++ b/src/spatial/ArborX_LinearBVH.hpp @@ -178,6 +178,18 @@ KOKKOS_FUNCTION typename Details::AccessValues::memory_space, typename Details::AccessValues::value_type>; +template +#if KOKKOS_VERSION >= 40400 +KOKKOS_DEDUCTION_GUIDE +#else +KOKKOS_FUNCTION +#endif + BoundingVolumeHierarchy(ExecutionSpace, Values, IndexableGetter) + -> BoundingVolumeHierarchy< + typename Details::AccessValues::memory_space, + typename Details::AccessValues::value_type, + IndexableGetter>; + template Date: Mon, 14 Oct 2024 10:42:15 -0400 Subject: [PATCH 2/5] Update code to eliminate usage of APIv1 in v2 mode --- .../brute_force_vs_bvh_timpl.hpp | 22 ++++----- .../bvh_driver/benchmark_registration.hpp | 32 ++++++++++++- benchmarks/bvh_driver/bvh_driver.cpp | 7 +-- .../dbscan/ArborX_DBSCANVerification.hpp | 17 +++++-- .../example_cuda_access_traits.cpp | 13 ++--- examples/brute_force/example_brute_force.cpp | 47 +++++++++++++++---- .../example_intersection.cpp | 20 ++++---- .../detail/ArborX_DistributedTreeNearest.hpp | 6 ++- .../detail/ArborX_DistributedTreeSpatial.hpp | 6 ++- .../detail/ArborX_DistributedTreeUtils.hpp | 12 +++++ test/tstQueryTreeIntersectsKDOP.cpp | 38 ++++++++++++--- 11 files changed, 166 insertions(+), 54 deletions(-) diff --git a/benchmarks/brute_force_vs_bvh/brute_force_vs_bvh_timpl.hpp b/benchmarks/brute_force_vs_bvh/brute_force_vs_bvh_timpl.hpp index db95d0bd6..6440f0de9 100644 --- a/benchmarks/brute_force_vs_bvh/brute_force_vs_bvh_timpl.hpp +++ b/benchmarks/brute_force_vs_bvh/brute_force_vs_bvh_timpl.hpp @@ -75,43 +75,43 @@ template static void run_fp(int nprimitives, int nqueries, int nrepeats) { ExecutionSpace space{}; + Placeholder primitives{nprimitives}; Placeholder predicates{nqueries}; + using Point = ArborX::Point; for (int i = 0; i < nrepeats; i++) { [[maybe_unused]] unsigned int out_count; { Kokkos::Timer timer; - ArborX::BoundingVolumeHierarchy bvh{ - space, ArborX::Experimental::attach_indices(primitives)}; + ArborX::BoundingVolumeHierarchy bvh{space, primitives}; - Kokkos::View indices("Benchmark::indices_ref", 0); + Kokkos::View values("Benchmark::values_ref", 0); Kokkos::View offset("Benchmark::offset_ref", 0); - bvh.query(space, predicates, indices, offset); + bvh.query(space, predicates, values, offset); space.fence(); double time = timer.seconds(); if (i == 0) printf("Collisions: %.5f\n", - (float)(indices.extent(0)) / (nprimitives * nqueries)); + (float)(values.extent(0)) / (nprimitives * nqueries)); printf("Time BVH : %lf\n", time); - out_count = indices.extent(0); + out_count = values.extent(0); } { Kokkos::Timer timer; - ArborX::BruteForce brute{ - space, ArborX::Experimental::attach_indices(primitives)}; + ArborX::BruteForce brute{space, primitives}; - Kokkos::View indices("Benchmark::indices", 0); + Kokkos::View values("Benchmark::values", 0); Kokkos::View offset("Benchmark::offset", 0); - brute.query(space, predicates, indices, offset); + brute.query(space, predicates, values, offset); space.fence(); double time = timer.seconds(); printf("Time BF : %lf\n", time); - assert(out_count == indices.extent(0)); + assert(out_count == values.extent(0)); } } } diff --git a/benchmarks/bvh_driver/benchmark_registration.hpp b/benchmarks/bvh_driver/benchmark_registration.hpp index a228a78ce..179f831fd 100644 --- a/benchmarks/bvh_driver/benchmark_registration.hpp +++ b/benchmarks/bvh_driver/benchmark_registration.hpp @@ -32,6 +32,34 @@ struct is_boost_rtree> : std::true_type template inline constexpr bool is_boost_rtree_v = is_boost_rtree::value; +template +struct Iota +{ + using memory_space = MemorySpace; + using index_type = Index; + + size_t _n; + + template >> + Iota(T n) + : _n(n) + {} +}; + +template +struct ArborX::AccessTraits, ArborX::PrimitivesTag> +{ + using Self = Iota; + + using memory_space = typename Self::memory_space; + static KOKKOS_FUNCTION size_t size(Self const &self) { return self._n; } + static KOKKOS_FUNCTION auto get(Self const &, size_t i) + { + return (typename Self::index_type)i; + } +}; + struct Spec { using PointCloudType = ArborXBenchmark::PointCloudType; @@ -135,7 +163,9 @@ auto makeTree(ExecutionSpace const &space, Primitives const &primitives) if constexpr (is_boost_rtree_v) return TreeType(space, primitives); else - return TreeType(space, ArborX::Experimental::attach_indices(primitives)); + return TreeType(space, + Iota{primitives.size()}, + primitives); } template diff --git a/benchmarks/bvh_driver/bvh_driver.cpp b/benchmarks/bvh_driver/bvh_driver.cpp index 2fa521b69..1741ff764 100644 --- a/benchmarks/bvh_driver/bvh_driver.cpp +++ b/benchmarks/bvh_driver/bvh_driver.cpp @@ -51,9 +51,10 @@ struct BenchmarkRegistration template using BVHBenchmarkRegistration = BenchmarkRegistration< - ExecutionSpace, - ArborX::BoundingVolumeHierarchy>>>; + ExecutionSpace, ArborX::BoundingVolumeHierarchy< + typename ExecutionSpace::memory_space, int, + Kokkos::View *, + typename ExecutionSpace::memory_space>>>; void register_bvh_benchmarks(Spec const &spec) { diff --git a/benchmarks/dbscan/ArborX_DBSCANVerification.hpp b/benchmarks/dbscan/ArborX_DBSCANVerification.hpp index 78fa1eff9..d98bf4970 100644 --- a/benchmarks/dbscan/ArborX_DBSCANVerification.hpp +++ b/benchmarks/dbscan/ArborX_DBSCANVerification.hpp @@ -276,6 +276,16 @@ bool verifyClusters(ExecutionSpace const &exec_space, IndicesView indices, }); } +struct IndexOnlyCallback +{ + template + KOKKOS_FUNCTION auto operator()(Query const &, Value const &value, + Output const &out) const + { + out(value.index); + } +}; + template bool verifyDBSCAN(ExecutionSpace exec_space, Primitives const &primitives, float eps, int core_min_size, LabelsView const &labels) @@ -301,12 +311,10 @@ bool verifyDBSCAN(ExecutionSpace exec_space, Primitives const &primitives, ArborX::BoundingVolumeHierarchy bvh( exec_space, ArborX::Experimental::attach_indices(points)); - auto const predicates = ArborX::Experimental::attach_indices( - ArborX::Experimental::make_intersects(points, eps)); - Kokkos::View indices("ArborX::DBSCAN::indices", 0); Kokkos::View offset("ArborX::DBSCAN::offset", 0); - ArborX::query(bvh, exec_space, predicates, indices, offset); + bvh.query(exec_space, ArborX::Experimental::make_intersects(points, eps), + IndexOnlyCallback{}, indices, offset); auto passed = Details::verifyClusters(exec_space, indices, offset, labels, core_min_size); @@ -314,6 +322,7 @@ bool verifyDBSCAN(ExecutionSpace exec_space, Primitives const &primitives, return passed; } + } // namespace Details } // namespace ArborX diff --git a/examples/access_traits/example_cuda_access_traits.cpp b/examples/access_traits/example_cuda_access_traits.cpp index e323f03ac..2aa275a15 100644 --- a/examples/access_traits/example_cuda_access_traits.cpp +++ b/examples/access_traits/example_cuda_access_traits.cpp @@ -79,19 +79,20 @@ int main(int argc, char *argv[]) cudaMemcpyAsync(d_a, a.data(), sizeof(a), cudaMemcpyHostToDevice, stream); Kokkos::Cuda cuda{stream}; - ArborX::BoundingVolumeHierarchy bvh{ - cuda, ArborX::Experimental::attach_indices(PointCloud{d_a, d_a, d_a, N})}; + ArborX::BoundingVolumeHierarchy bvh{cuda, PointCloud{d_a, d_a, d_a, N}}; - Kokkos::View indices("Example::indices", 0); + Kokkos::View *, Kokkos::CudaSpace> points("Example::points", + 0); Kokkos::View offset("Example::offset", 0); - ArborX::query(bvh, cuda, Spheres{d_a, d_a, d_a, d_a, N}, indices, offset); + ArborX::query(bvh, cuda, Spheres{d_a, d_a, d_a, d_a, N}, points, offset); Kokkos::parallel_for( - "Example::print_indices", Kokkos::RangePolicy(cuda, 0, N), + "Example::print_points", Kokkos::RangePolicy(cuda, 0, N), KOKKOS_LAMBDA(int i) { for (int j = offset(i); j < offset(i + 1); ++j) { - printf("%i %i\n", i, indices(j)); + printf("%i: (%.1f, %.1f, %.1f)\n", i, points(j)[0], points(j)[1], + points(j)[2]); } }); diff --git a/examples/brute_force/example_brute_force.cpp b/examples/brute_force/example_brute_force.cpp index 6bf5543d2..46cfb65a3 100644 --- a/examples/brute_force/example_brute_force.cpp +++ b/examples/brute_force/example_brute_force.cpp @@ -23,13 +23,41 @@ struct Dummy using ExecutionSpace = Kokkos::DefaultExecutionSpace; using MemorySpace = ExecutionSpace::memory_space; -template <> -struct ArborX::AccessTraits +template +struct Iota { using memory_space = MemorySpace; - using size_type = typename MemorySpace::size_type; - static KOKKOS_FUNCTION size_type size(Dummy const &d) { return d.count; } - static KOKKOS_FUNCTION auto get(Dummy const &, size_type i) + using index_type = Index; + + size_t _n; + + template >> + Iota(T n) + : _n(n) + {} +}; + +template +struct ArborX::AccessTraits, ArborX::PrimitivesTag> +{ + using Self = Iota; + + using memory_space = typename Self::memory_space; + static KOKKOS_FUNCTION size_t size(Self const &self) { return self._n; } + static KOKKOS_FUNCTION auto get(Self const &, size_t i) + { + return (typename Self::index_type)i; + } +}; + +struct DummyIndexableGetter +{ + int count; + + using memory_space = MemorySpace; + KOKKOS_FUNCTION auto size() const { return count; } + KOKKOS_FUNCTION auto operator()(int i) const { return ArborX::Point{(float)i, (float)i, (float)i}; } @@ -69,13 +97,13 @@ int main(int argc, char *argv[]) int nprimitives = 5; int npredicates = 5; - Dummy primitives{nprimitives}; + Iota primitives{nprimitives}; + DummyIndexableGetter indexable_getter{nprimitives}; Dummy predicates{npredicates}; unsigned int out_count; { - ArborX::BoundingVolumeHierarchy bvh{ - space, ArborX::Experimental::attach_indices(primitives)}; + ArborX::BoundingVolumeHierarchy bvh{space, primitives, indexable_getter}; Kokkos::View indices("Example::indices_ref", 0); Kokkos::View offset("Example::offset_ref", 0); @@ -88,8 +116,7 @@ int main(int argc, char *argv[]) } { - ArborX::BruteForce brute{space, - ArborX::Experimental::attach_indices(primitives)}; + ArborX::BruteForce brute{space, primitives, indexable_getter}; Kokkos::View indices("Example::indices", 0); Kokkos::View offset("Example::offset", 0); diff --git a/examples/simple_intersection/example_intersection.cpp b/examples/simple_intersection/example_intersection.cpp index e0d03be79..1386b6020 100644 --- a/examples/simple_intersection/example_intersection.cpp +++ b/examples/simple_intersection/example_intersection.cpp @@ -50,18 +50,20 @@ int main(int argc, char *argv[]) ExecutionSpace space; + using Value = ArborX::PairValueIndex; + ArborX::BoundingVolumeHierarchy const tree( space, ArborX::Experimental::attach_indices(boxes)); - // The query will resize indices and offsets accordingly - Kokkos::View indices("Example::indices", 0); + // The query will resize values and offsets accordingly + Kokkos::View values("Example::values", 0); Kokkos::View offsets("Example::offsets", 0); - tree.query(space, queries, indices, offsets); + tree.query(space, queries, values, offsets); auto offsets_host = Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace{}, offsets); - auto indices_host = - Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace{}, indices); + auto values_host = + Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace{}, values); // Expected output: // offsets: 0 1 2 6 @@ -70,10 +72,10 @@ int main(int argc, char *argv[]) std::cout << "offsets: "; std::copy(offsets_host.data(), offsets_host.data() + offsets.size(), std::ostream_iterator(std::cout, " ")); - std::cout << "\nindices: "; - std::copy(indices_host.data(), indices_host.data() + indices.size(), - std::ostream_iterator(std::cout, " ")); - std::cout << "\n"; + std::cout << "\nindices:"; + for (int i = 0; i < values_host.extent_int(0); ++i) + std::cout << " " << values_host(i).index; + std::cout << '\n'; return 0; } diff --git a/src/distributed/detail/ArborX_DistributedTreeNearest.hpp b/src/distributed/detail/ArborX_DistributedTreeNearest.hpp index c8c0d0a43..5aa1a7dc0 100644 --- a/src/distributed/detail/ArborX_DistributedTreeNearest.hpp +++ b/src/distributed/detail/ArborX_DistributedTreeNearest.hpp @@ -52,7 +52,8 @@ void DistributedTreeImpl::phaseI(ExecutionSpace const &space, Tree const &tree, // Find the k nearest local trees. Kokkos::View offset(prefix + "::offset", 0); Kokkos::View nearest_ranks(prefix + "::nearest_ranks", 0); - tree._top_tree.query(space, predicates, nearest_ranks, offset); + tree._top_tree.query(space, predicates, DistributedTree::IndexOnlyCallback{}, + nearest_ranks, offset); // Accumulate total leave count in the local trees until it reaches k which // is the number of neighbors queried for. Stop if local trees get @@ -139,7 +140,8 @@ void DistributedTreeImpl::phaseII(ExecutionSpace const &space, Tree const &tree, tree._top_tree.query(space, WithinDistanceFromPredicates{ predicates, distances}, - nearest_ranks, offset); + DistributedTree::IndexOnlyCallback{}, nearest_ranks, + offset); auto const &bottom_tree = tree._bottom_tree; using BottomTree = std::decay_t; diff --git a/src/distributed/detail/ArborX_DistributedTreeSpatial.hpp b/src/distributed/detail/ArborX_DistributedTreeSpatial.hpp index 74c4d099e..ab37583e7 100644 --- a/src/distributed/detail/ArborX_DistributedTreeSpatial.hpp +++ b/src/distributed/detail/ArborX_DistributedTreeSpatial.hpp @@ -50,7 +50,8 @@ DistributedTreeImpl::queryDispatch(SpatialPredicateTag, Tree const &tree, Kokkos::View intersected_ranks( "ArborX::DistributedTree::query::spatial::intersected_ranks", 0); - top_tree.query(space, predicates, intersected_ranks, offset); + top_tree.query(space, predicates, DistributedTree::IndexOnlyCallback{}, + intersected_ranks, offset); DistributedTree::forwardQueriesAndCommunicateResults( tree.getComm(), space, tree._bottom_tree, predicates, callback, @@ -81,7 +82,8 @@ void DistributedTreeImpl::queryDispatch(SpatialPredicateTag, Tree const &tree, Kokkos::View intersected_ranks( prefix + "::intersected_ranks", 0); Kokkos::View offset(prefix + "::offset", 0); - top_tree.query(space, predicates, intersected_ranks, offset); + top_tree.query(space, predicates, DistributedTree::IndexOnlyCallback{}, + intersected_ranks, offset); using Query = typename Predicates::value_type; Kokkos::View fwd_predicates(prefix + "::fwd_predicates", diff --git a/src/distributed/detail/ArborX_DistributedTreeUtils.hpp b/src/distributed/detail/ArborX_DistributedTreeUtils.hpp index f61465337..683444e09 100644 --- a/src/distributed/detail/ArborX_DistributedTreeUtils.hpp +++ b/src/distributed/detail/ArborX_DistributedTreeUtils.hpp @@ -14,6 +14,7 @@ #include #include +#include #include #include #include @@ -335,6 +336,17 @@ void filterResults(ExecutionSpace const &space, Predicates const &queries, offset = new_offset; } +struct IndexOnlyCallback +{ + template + KOKKOS_FUNCTION auto operator()(Query const &, + PairValueIndex const &value, + Output const &out) const + { + out(value.index); + } +}; + } // namespace ArborX::Details::DistributedTree #endif diff --git a/test/tstQueryTreeIntersectsKDOP.cpp b/test/tstQueryTreeIntersectsKDOP.cpp index f0b7cf193..a62c885e9 100644 --- a/test/tstQueryTreeIntersectsKDOP.cpp +++ b/test/tstQueryTreeIntersectsKDOP.cpp @@ -18,6 +18,34 @@ #include +template +struct Iota +{ + using memory_space = MemorySpace; + using index_type = Index; + + size_t _n; + + template >> + Iota(T n) + : _n(n) + {} +}; + +template +struct ArborX::AccessTraits, ArborX::PrimitivesTag> +{ + using Self = Iota; + + using memory_space = typename Self::memory_space; + static KOKKOS_FUNCTION size_t size(Self const &self) { return self._n; } + static KOKKOS_FUNCTION auto get(Self const &, size_t i) + { + return (typename Self::index_type)i; + } +}; + #include #include "Search_UnitTestHelpers.hpp" @@ -28,8 +56,6 @@ BOOST_AUTO_TEST_CASE_TEMPLATE(intersects_kdop, DeviceType, ARBORX_DEVICE_TYPES) using MemorySpace = typename DeviceType::memory_space; using Point = ArborX::Point<3>; - using Tree = ArborX::BoundingVolumeHierarchy>; std::vector primitives = { {{0, 0, 0}}, // 0 @@ -46,12 +72,12 @@ BOOST_AUTO_TEST_CASE_TEMPLATE(intersects_kdop, DeviceType, ARBORX_DEVICE_TYPES) {{0, 0, 2}}, // 11 {{0, 0, 3}}, // 12 }; - Tree const tree( - ExecutionSpace{}, - ArborX::Experimental::attach_indices(Kokkos::create_mirror_view_and_copy( + ArborX::BoundingVolumeHierarchy const tree( + ExecutionSpace{}, Iota{primitives.size()}, + Kokkos::create_mirror_view_and_copy( MemorySpace{}, Kokkos::View( - primitives.data(), primitives.size())))); + primitives.data(), primitives.size()))); // (0,0,0)->(1,2,3) box with (0,0,0)--(0,0,3) edge cut off ArborX::Experimental::KDOP<3, 18> x; From 095a4a06084f69a5262709a513bf33b88a6dbd80 Mon Sep 17 00:00:00 2001 From: Andrey Prokopenko Date: Mon, 14 Oct 2024 10:42:50 -0400 Subject: [PATCH 3/5] Remove APIv1 convenience shortcuts from BVH and BruteForce --- src/spatial/ArborX_BruteForce.hpp | 35 ++++--------------------------- src/spatial/ArborX_LinearBVH.hpp | 35 ++++--------------------------- 2 files changed, 8 insertions(+), 62 deletions(-) diff --git a/src/spatial/ArborX_BruteForce.hpp b/src/spatial/ArborX_BruteForce.hpp index 09dd2b52b..0cdc3d5e7 100644 --- a/src/spatial/ArborX_BruteForce.hpp +++ b/src/spatial/ArborX_BruteForce.hpp @@ -81,37 +81,10 @@ class BruteForce using Predicates = Details::AccessValues; using Tag = typename Predicates::value_type::Tag; - // Automatically add LegacyDefaultCallback if - // 1. A user does not provide a callback - // 2. The index is constructed on PairValueIndex - // 3. The output value_type is an integral type - constexpr bool use_convenient_shortcut = []() { - if constexpr (!Kokkos::is_view_v>) - return false; - else if constexpr (!Details::is_pair_value_index_v) - return false; - else - return std::is_integral_v< - typename std::decay_t::value_type>; - }(); - - if constexpr (use_convenient_shortcut) - { - // Simplified way to get APIv1 result using APIv2 interface - Details::CrsGraphWrapperImpl::queryDispatch( - Tag{}, *this, space, Predicates{user_predicates}, - Details::LegacyDefaultCallback{}, // inject legacy callback arg - std::forward(callback_or_view), - std::forward(view), std::forward(args)...); - return; - } - else - { - Details::CrsGraphWrapperImpl::queryDispatch( - Tag{}, *this, space, Predicates{user_predicates}, - std::forward(callback_or_view), - std::forward(view), std::forward(args)...); - } + Details::CrsGraphWrapperImpl::queryDispatch( + Tag{}, *this, space, Predicates{user_predicates}, + std::forward(callback_or_view), + std::forward(view), std::forward(args)...); } KOKKOS_FUNCTION auto const &indexable_get() const diff --git a/src/spatial/ArborX_LinearBVH.hpp b/src/spatial/ArborX_LinearBVH.hpp index 29fa18fd4..b89b54478 100644 --- a/src/spatial/ArborX_LinearBVH.hpp +++ b/src/spatial/ArborX_LinearBVH.hpp @@ -103,37 +103,10 @@ class BoundingVolumeHierarchy using Predicates = Details::AccessValues; using Tag = typename Predicates::value_type::Tag; - // Automatically add LegacyDefaultCallback if - // 1. A user does not provide a callback - // 2. The index is constructed on PairValueIndex - // 3. The output value_type is an integral type - constexpr bool use_convenient_shortcut = []() { - if constexpr (!Kokkos::is_view_v>) - return false; - else if constexpr (!Details::is_pair_value_index_v) - return false; - else - return std::is_integral_v< - typename std::decay_t::value_type>; - }(); - - if constexpr (use_convenient_shortcut) - { - // Simplified way to get APIv1 result using APIv2 interface - Details::CrsGraphWrapperImpl::queryDispatch( - Tag{}, *this, space, Predicates{user_predicates}, - Details::LegacyDefaultCallback{}, // inject legacy callback arg - std::forward(callback_or_view), - std::forward(view), std::forward(args)...); - return; - } - else - { - Details::CrsGraphWrapperImpl::queryDispatch( - Tag{}, *this, space, Predicates{user_predicates}, - std::forward(callback_or_view), - std::forward(view), std::forward(args)...); - } + Details::CrsGraphWrapperImpl::queryDispatch( + Tag{}, *this, space, Predicates{user_predicates}, + std::forward(callback_or_view), + std::forward(view), std::forward(args)...); } template From 1492fdfbdbafe0d0694ff675422ee164bc58a75b Mon Sep 17 00:00:00 2001 From: Andrey Prokopenko Date: Tue, 15 Oct 2024 09:54:58 -0400 Subject: [PATCH 4/5] Remove is_pair_v --- src/spatial/detail/ArborX_PairValueIndex.hpp | 15 --------------- 1 file changed, 15 deletions(-) diff --git a/src/spatial/detail/ArborX_PairValueIndex.hpp b/src/spatial/detail/ArborX_PairValueIndex.hpp index 28941c6fd..2c9c69807 100644 --- a/src/spatial/detail/ArborX_PairValueIndex.hpp +++ b/src/spatial/detail/ArborX_PairValueIndex.hpp @@ -31,21 +31,6 @@ struct PairValueIndex Index index; }; -namespace Details -{ -template -struct is_pair_value_index : public std::false_type -{}; - -template -struct is_pair_value_index> : public std::true_type -{}; - -template -inline constexpr bool is_pair_value_index_v = is_pair_value_index::value; - -} // namespace Details - } // namespace ArborX #endif From e1250ce16f5097793837230c9b63e925a7c43762 Mon Sep 17 00:00:00 2001 From: Andrey Prokopenko Date: Mon, 11 Nov 2024 10:35:05 -0500 Subject: [PATCH 5/5] Simplify Iota struct usage --- .../bvh_driver/benchmark_registration.hpp | 21 ++++++------------- examples/brute_force/example_brute_force.cpp | 18 ++++------------ test/tstQueryTreeIntersectsKDOP.cpp | 20 +++++------------- 3 files changed, 15 insertions(+), 44 deletions(-) diff --git a/benchmarks/bvh_driver/benchmark_registration.hpp b/benchmarks/bvh_driver/benchmark_registration.hpp index 179f831fd..f22991d83 100644 --- a/benchmarks/bvh_driver/benchmark_registration.hpp +++ b/benchmarks/bvh_driver/benchmark_registration.hpp @@ -32,19 +32,12 @@ struct is_boost_rtree> : std::true_type template inline constexpr bool is_boost_rtree_v = is_boost_rtree::value; -template +template struct Iota { + static_assert(Kokkos::is_memory_space_v); using memory_space = MemorySpace; - using index_type = Index; - - size_t _n; - - template >> - Iota(T n) - : _n(n) - {} + int _n; }; template @@ -54,10 +47,7 @@ struct ArborX::AccessTraits, ArborX::PrimitivesTag> using memory_space = typename Self::memory_space; static KOKKOS_FUNCTION size_t size(Self const &self) { return self._n; } - static KOKKOS_FUNCTION auto get(Self const &, size_t i) - { - return (typename Self::index_type)i; - } + static KOKKOS_FUNCTION auto get(Self const &, int i) { return i; } }; struct Spec @@ -164,7 +154,8 @@ auto makeTree(ExecutionSpace const &space, Primitives const &primitives) return TreeType(space, primitives); else return TreeType(space, - Iota{primitives.size()}, + Iota{ + static_cast(primitives.size())}, primitives); } diff --git a/examples/brute_force/example_brute_force.cpp b/examples/brute_force/example_brute_force.cpp index 46cfb65a3..22445f79a 100644 --- a/examples/brute_force/example_brute_force.cpp +++ b/examples/brute_force/example_brute_force.cpp @@ -23,19 +23,12 @@ struct Dummy using ExecutionSpace = Kokkos::DefaultExecutionSpace; using MemorySpace = ExecutionSpace::memory_space; -template +template struct Iota { + static_assert(Kokkos::is_memory_space_v); using memory_space = MemorySpace; - using index_type = Index; - - size_t _n; - - template >> - Iota(T n) - : _n(n) - {} + int _n; }; template @@ -45,10 +38,7 @@ struct ArborX::AccessTraits, ArborX::PrimitivesTag> using memory_space = typename Self::memory_space; static KOKKOS_FUNCTION size_t size(Self const &self) { return self._n; } - static KOKKOS_FUNCTION auto get(Self const &, size_t i) - { - return (typename Self::index_type)i; - } + static KOKKOS_FUNCTION auto get(Self const &, int i) { return i; } }; struct DummyIndexableGetter diff --git a/test/tstQueryTreeIntersectsKDOP.cpp b/test/tstQueryTreeIntersectsKDOP.cpp index a62c885e9..21dac3737 100644 --- a/test/tstQueryTreeIntersectsKDOP.cpp +++ b/test/tstQueryTreeIntersectsKDOP.cpp @@ -18,19 +18,12 @@ #include -template +template struct Iota { + static_assert(Kokkos::is_memory_space_v); using memory_space = MemorySpace; - using index_type = Index; - - size_t _n; - - template >> - Iota(T n) - : _n(n) - {} + int _n; }; template @@ -40,10 +33,7 @@ struct ArborX::AccessTraits, ArborX::PrimitivesTag> using memory_space = typename Self::memory_space; static KOKKOS_FUNCTION size_t size(Self const &self) { return self._n; } - static KOKKOS_FUNCTION auto get(Self const &, size_t i) - { - return (typename Self::index_type)i; - } + static KOKKOS_FUNCTION auto get(Self const &, int i) { return i; } }; #include @@ -73,7 +63,7 @@ BOOST_AUTO_TEST_CASE_TEMPLATE(intersects_kdop, DeviceType, ARBORX_DEVICE_TYPES) {{0, 0, 3}}, // 12 }; ArborX::BoundingVolumeHierarchy const tree( - ExecutionSpace{}, Iota{primitives.size()}, + ExecutionSpace{}, Iota{static_cast(primitives.size())}, Kokkos::create_mirror_view_and_copy( MemorySpace{}, Kokkos::View(