From ba7d8052a00b8774c5f1aa0128d2259b4692a455 Mon Sep 17 00:00:00 2001 From: Damien L-G Date: Fri, 29 May 2020 09:06:38 -0400 Subject: [PATCH 01/10] Add is_tagged_post_callback class template --- src/details/ArborX_Callbacks.hpp | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/src/details/ArborX_Callbacks.hpp b/src/details/ArborX_Callbacks.hpp index a52e67fde..68c557493 100644 --- a/src/details/ArborX_Callbacks.hpp +++ b/src/details/ArborX_Callbacks.hpp @@ -79,6 +79,13 @@ using SpatialPredicateInlineCallbackArchetypeExpression = template using CallbackTagArchetypeAlias = typename Callback::tag; +template +struct is_tagged_post_callback + : std::is_same, + PostCallbackTag>::type +{ +}; + // output functor to pass to the callback during detection template struct Sink @@ -106,7 +113,7 @@ void check_valid_callback(Callback const &, Predicates const &, using Predicate = typename Traits::Helper::type; // FIXME - constexpr bool short_circuit = std::is_same{}; + constexpr bool short_circuit = is_tagged_post_callback{}; static_assert( short_circuit || (std::is_same{} && From a817653a723dae527cf62aa870fc0cbc37460546 Mon Sep 17 00:00:00 2001 From: Damien L-G Date: Fri, 29 May 2020 09:08:54 -0400 Subject: [PATCH 02/10] Do not require callback to be tagegd --- src/details/ArborX_Callbacks.hpp | 8 -------- 1 file changed, 8 deletions(-) diff --git a/src/details/ArborX_Callbacks.hpp b/src/details/ArborX_Callbacks.hpp index 68c557493..7dde92f52 100644 --- a/src/details/ArborX_Callbacks.hpp +++ b/src/details/ArborX_Callbacks.hpp @@ -100,14 +100,6 @@ template void check_valid_callback(Callback const &, Predicates const &, OutputView const &) { - static_assert(is_detected{}, - "Callback must define 'tag' member type"); - - using CallbackTag = detected_t; - static_assert(std::is_same{} || - std::is_same{}, - "Tag must be either 'InlineCallbackTag' or 'PostCallbackTag'"); - using Access = Traits::Access; using PredicateTag = typename Traits::Helper::tag; using Predicate = typename Traits::Helper::type; From d966cbeb8021de15a75d7f16120088903af77d9c Mon Sep 17 00:00:00 2001 From: Damien L-G Date: Fri, 29 May 2020 09:25:53 -0400 Subject: [PATCH 03/10] Rework queryDispatch overload set --- .../ArborX_DetailsBoundingVolumeHierarchyImpl.hpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/src/details/ArborX_DetailsBoundingVolumeHierarchyImpl.hpp b/src/details/ArborX_DetailsBoundingVolumeHierarchyImpl.hpp index 0df4856db..dc770780e 100644 --- a/src/details/ArborX_DetailsBoundingVolumeHierarchyImpl.hpp +++ b/src/details/ArborX_DetailsBoundingVolumeHierarchyImpl.hpp @@ -117,7 +117,8 @@ namespace BoundingVolumeHierarchyImpl // is called. template -std::enable_if_t::value> +std::enable_if_t{} && + Kokkos::is_view{} && Kokkos::is_view{}> queryDispatch(SpatialPredicateTag, BVH const &bvh, ExecutionSpace const &space, Predicates const &predicates, Callback const &callback, OutputView &out, OffsetView &offset, @@ -189,8 +190,7 @@ queryDispatch(SpatialPredicateTag, BVH const &bvh, ExecutionSpace const &space, template -inline std::enable_if_t< - std::is_same::value> +inline std::enable_if_t{}> queryDispatch(SpatialPredicateTag, BVH const &bvh, ExecutionSpace const &space, Predicates const &predicates, Callback const &callback, OutputView &out, OffsetView &offset, @@ -206,7 +206,8 @@ queryDispatch(SpatialPredicateTag, BVH const &bvh, ExecutionSpace const &space, template -std::enable_if_t::value> +std::enable_if_t{} && + Kokkos::is_view{} && Kokkos::is_view{}> queryDispatch(NearestPredicateTag, BVH const &bvh, ExecutionSpace const &space, Predicates const &predicates, Callback const &callback, OutputView &out, OffsetView &offset, @@ -266,8 +267,7 @@ queryDispatch(NearestPredicateTag, BVH const &bvh, ExecutionSpace const &space, template -inline std::enable_if_t< - std::is_same::value> +inline std::enable_if_t{}> queryDispatch(NearestPredicateTag, BVH const &bvh, ExecutionSpace const &space, Predicates const &predicates, Callback const &callback, OutputView &out, OffsetView &offset, From 61ea2893dbe70e96116f849814a30ebe49cd55a8 Mon Sep 17 00:00:00 2001 From: dalg24 Date: Fri, 29 May 2020 10:40:52 -0400 Subject: [PATCH 04/10] Make static assertion fail if using generic lambda with nvcc --- src/details/ArborX_Callbacks.hpp | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/src/details/ArborX_Callbacks.hpp b/src/details/ArborX_Callbacks.hpp index 7dde92f52..b1efec41e 100644 --- a/src/details/ArborX_Callbacks.hpp +++ b/src/details/ArborX_Callbacks.hpp @@ -100,6 +100,13 @@ template void check_valid_callback(Callback const &, Predicates const &, OutputView const &) { +#ifdef __NVCC__ + // Without it would get a segmentation fault and no diagnostic whatsoever + static_assert( + !__nv_is_extended_host_device_lambda_closure_type(Callback), + "__host__ __device__ extended lambdas cannot be generic lambdas"); +#endif + using Access = Traits::Access; using PredicateTag = typename Traits::Helper::tag; using Predicate = typename Traits::Helper::type; From 3fd48398a9eeffb05f0d5b02307748fd2646e6c3 Mon Sep 17 00:00:00 2001 From: Damien L-G Date: Fri, 29 May 2020 09:43:58 -0400 Subject: [PATCH 05/10] Update compile-only callback unit test --- test/tstCallbacks.cpp | 18 ++++++++++++++---- 1 file changed, 14 insertions(+), 4 deletions(-) diff --git a/test/tstCallbacks.cpp b/test/tstCallbacks.cpp index bbabfa04c..969e23300 100644 --- a/test/tstCallbacks.cpp +++ b/test/tstCallbacks.cpp @@ -70,9 +70,19 @@ int main() // Uncomment to see error messages - // check_valid_callback(SpatialPredicateCallbackMissingTag{}, - // SpatialPredicates{}, v); + check_valid_callback(SpatialPredicateCallbackMissingTag{}, + SpatialPredicates{}, v); + + check_valid_callback(NearestPredicateCallbackMissingTag{}, + NearestPredicates{}, v); - // check_valid_callback(NearestPredicateCallbackMissingTag{}, - // NearestPredicates{}, v); +#ifndef __NVCC__ + check_valid_callback( + [](auto const &predicate, int primitive, auto const &out) {}, + SpatialPredicates{}, v); + + check_valid_callback([](auto const &predicate, int primitive, float distance, + auto const &out) {}, + NearestPredicates{}, v); +#endif } From e11b3076c321f0a3035d0cc020945c5e0b43d1fb Mon Sep 17 00:00:00 2001 From: Damien L-G Date: Fri, 29 May 2020 13:53:06 -0400 Subject: [PATCH 06/10] Add callback example --- examples/CMakeLists.txt | 2 + examples/callback/CMakeLists.txt | 3 + examples/callback/example_callback.cpp | 138 +++++++++++++++++++++++++ 3 files changed, 143 insertions(+) create mode 100644 examples/callback/CMakeLists.txt create mode 100644 examples/callback/example_callback.cpp diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index aef663c35..c2c1a1a54 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -11,6 +11,8 @@ endif() add_subdirectory(access_traits) +add_subdirectory(callback) + find_package(Boost COMPONENTS program_options) if(Boost_FOUND) add_subdirectory(viz) diff --git a/examples/callback/CMakeLists.txt b/examples/callback/CMakeLists.txt new file mode 100644 index 000000000..eb1990c4a --- /dev/null +++ b/examples/callback/CMakeLists.txt @@ -0,0 +1,3 @@ +add_executable(ArborX_Callback.exe example_callback.cpp) +target_link_libraries(ArborX_Callback.exe ${ArborX_TARGET}) +add_test(NAME ArborX_Callback_Example COMMAND ./ArborX_Callback.exe) diff --git a/examples/callback/example_callback.cpp b/examples/callback/example_callback.cpp new file mode 100644 index 000000000..61be5691a --- /dev/null +++ b/examples/callback/example_callback.cpp @@ -0,0 +1,138 @@ +/**************************************************************************** + * Copyright (c) 2012-2020 by the ArborX authors * + * All rights reserved. * + * * + * This file is part of the ArborX library. ArborX is * + * distributed under a BSD 3-clause license. For the licensing terms see * + * the LICENSE file in the top-level directory. * + * * + * SPDX-License-Identifier: BSD-3-Clause * + ****************************************************************************/ + +#include + +#include + +#include +#include +#include + +using ExecutionSpace = Kokkos::DefaultExecutionSpace; +using MemorySpace = ExecutionSpace::memory_space; + +struct FirstOctant +{ +}; + +struct NearestToOrigin +{ + int k; +}; + +namespace ArborX +{ +namespace Traits +{ +template <> +struct Access +{ + KOKKOS_FUNCTION static std::size_t size(FirstOctant) { return 1; } + KOKKOS_FUNCTION static auto get(FirstOctant, std::size_t) + { + return intersects(Box{{{0, 0, 0}}, {{1, 1, 1}}}); + } + using memory_space = MemorySpace; +}; +template <> +struct Access +{ + KOKKOS_FUNCTION static std::size_t size(NearestToOrigin) { return 1; } + KOKKOS_FUNCTION static auto get(NearestToOrigin d, std::size_t) + { + return nearest(Point{0, 0, 0}, d.k); + } + using memory_space = MemorySpace; +}; +} // namespace Traits +} // namespace ArborX + +struct PairIndexDistance +{ + int index; + float distance; +}; + +struct PrintfCallback +{ + template + KOKKOS_FUNCTION void operator()(Predicate, int primitive, + OutputFunctor const &out) const + { + printf("Found %d from functor\n", primitive); + out(primitive); + } + template + KOKKOS_FUNCTION void operator()(Predicate, int primitive, float distance, + OutputFunctor const &out) const + { + printf("Found %d with distance %.3f from functor\n", primitive, distance); + out({primitive, distance}); + } +}; + +int main(int argc, char *argv[]) +{ + Kokkos::ScopeGuard guard(argc, argv); + + int const n = 100; + std::vector points; + // Fill vector with random points in [-1, 1]^3 + std::uniform_real_distribution dis{-1., 1.}; + std::default_random_engine gen; + auto rd = [&]() { return dis(gen); }; + std::generate_n(std::back_inserter(points), n, [&]() { + return ArborX::Point{rd(), rd(), rd()}; + }); + + ArborX::BVH bvh{ + ExecutionSpace{}, + Kokkos::create_mirror_view_and_copy( + MemorySpace{}, + Kokkos::View(points.data(), points.size()))}; + + { + Kokkos::View values("values", 0); + Kokkos::View offsets("offsets", 0); + bvh.query(ExecutionSpace{}, FirstOctant{}, PrintfCallback{}, values, + offsets); +#ifndef __NVCC__ + bvh.query( + ExecutionSpace{}, FirstOctant{}, + KOKKOS_LAMBDA(auto /*predicate*/, int primitive, + auto /*output_functor*/) { + printf("Found %d from generic lambda\n", primitive); + }, + values, offsets); +#endif + } + + { + int const k = 10; + Kokkos::View values("values", 0); + Kokkos::View offsets("offsets", 0); + bvh.query(ExecutionSpace{}, NearestToOrigin{k}, PrintfCallback{}, values, + offsets); +#ifndef __NVCC__ + bvh.query(ExecutionSpace{}, NearestToOrigin{k}, + KOKKOS_LAMBDA(auto /*predicate*/, int primitive, float distance, + auto /*output_functor*/) { + printf("Found %d with distance %.3f from generic lambda\n", + primitive, distance); + }, + values, offsets); +#endif + } + + return 0; +} From 421bceaf718b5bf1fd63ad3c80cc25fd24cc5f37 Mon Sep 17 00:00:00 2001 From: Damien L-G Date: Fri, 29 May 2020 13:53:36 -0400 Subject: [PATCH 07/10] Add experimental tree traversal with callback to example --- examples/callback/example_callback.cpp | 25 ++++++++++++++++++------- 1 file changed, 18 insertions(+), 7 deletions(-) diff --git a/examples/callback/example_callback.cpp b/examples/callback/example_callback.cpp index 61be5691a..372a805bf 100644 --- a/examples/callback/example_callback.cpp +++ b/examples/callback/example_callback.cpp @@ -107,13 +107,12 @@ int main(int argc, char *argv[]) bvh.query(ExecutionSpace{}, FirstOctant{}, PrintfCallback{}, values, offsets); #ifndef __NVCC__ - bvh.query( - ExecutionSpace{}, FirstOctant{}, - KOKKOS_LAMBDA(auto /*predicate*/, int primitive, - auto /*output_functor*/) { - printf("Found %d from generic lambda\n", primitive); - }, - values, offsets); + bvh.query(ExecutionSpace{}, FirstOctant{}, + KOKKOS_LAMBDA(auto /*predicate*/, int primitive, + auto /*output_functor*/) { + printf("Found %d from generic lambda\n", primitive); + }, + values, offsets); #endif } @@ -134,5 +133,17 @@ int main(int argc, char *argv[]) #endif } + { + // EXPERIMENTAL + // TODO replace with BVH::query(ExecutionSpace, Predicates, Callback) when + // new overload is added + Kokkos::View> c( + "counter"); + + ArborX::Details::traverse( + ExecutionSpace{}, bvh, FirstOctant{}, + KOKKOS_LAMBDA(int i, int j) { printf("%d %d %d\n", ++c(), i, j); }); + } + return 0; } From 8a1f3b96b952741ebf182da73f92a6f98e4e3bb0 Mon Sep 17 00:00:00 2001 From: Damien L-G Date: Fri, 29 May 2020 14:03:01 -0400 Subject: [PATCH 08/10] Treat "post" callback detection separately --- src/details/ArborX_Callbacks.hpp | 10 +++------- .../ArborX_DetailsBoundingVolumeHierarchyImpl.hpp | 13 ++++++++++++- 2 files changed, 15 insertions(+), 8 deletions(-) diff --git a/src/details/ArborX_Callbacks.hpp b/src/details/ArborX_Callbacks.hpp index b1efec41e..f2ba99b11 100644 --- a/src/details/ArborX_Callbacks.hpp +++ b/src/details/ArborX_Callbacks.hpp @@ -111,14 +111,10 @@ void check_valid_callback(Callback const &, Predicates const &, using PredicateTag = typename Traits::Helper::tag; using Predicate = typename Traits::Helper::type; - // FIXME - constexpr bool short_circuit = is_tagged_post_callback{}; static_assert( - short_circuit || - (std::is_same{} && - is_detected>{}) || + (std::is_same{} && + is_detected>{}) || (std::is_same{} && is_detected>{}), diff --git a/src/details/ArborX_DetailsBoundingVolumeHierarchyImpl.hpp b/src/details/ArborX_DetailsBoundingVolumeHierarchyImpl.hpp index dc770780e..50e691fdb 100644 --- a/src/details/ArborX_DetailsBoundingVolumeHierarchyImpl.hpp +++ b/src/details/ArborX_DetailsBoundingVolumeHierarchyImpl.hpp @@ -325,7 +325,8 @@ queryDispatch(NearestPredicateTag, BVH const &bvh, ExecutionSpace const &space, } // namespace BoundingVolumeHierarchyImpl template -std::enable_if_t{}> +std::enable_if_t{} && + !is_tagged_post_callback{}> check_valid_callback_if_first_argument_is_not_a_view( Callback const &callback, Predicates const &predicates, OutputView const &out) @@ -333,6 +334,16 @@ check_valid_callback_if_first_argument_is_not_a_view( check_valid_callback(callback, predicates, out); } +template +std::enable_if_t{} && + is_tagged_post_callback{}> +check_valid_callback_if_first_argument_is_not_a_view( + Callback const &callback, Predicates const &predicates, + OutputView const &out) +{ + // TODO +} + template std::enable_if_t{}> check_valid_callback_if_first_argument_is_not_a_view(View const &, From 816a422d33ecfc643d6d27453bcff0cf2cc431ba Mon Sep 17 00:00:00 2001 From: Damien L-G Date: Fri, 29 May 2020 14:23:46 -0400 Subject: [PATCH 09/10] Cleanup callback unit tests --- test/tstCallbacks.cpp | 37 ++++++++++++++++++++++++++++++------- 1 file changed, 30 insertions(+), 7 deletions(-) diff --git a/test/tstCallbacks.cpp b/test/tstCallbacks.cpp index 969e23300..51ca9c7c2 100644 --- a/test/tstCallbacks.cpp +++ b/test/tstCallbacks.cpp @@ -51,6 +51,18 @@ struct NearestPredicateCallbackMissingTag } }; +struct Wrong +{ +}; + +struct SpatialPredicateCallbackDoesNotTakeCorrectArgument +{ + template + void operator()(Wrong, int, OutputFunctor const &) const + { + } +}; + int main() { using ArborX::Details::check_valid_callback; @@ -68,21 +80,32 @@ int main() ArborX::Details::CallbackDefaultNearestPredicateWithDistance{}, NearestPredicates{}, v); - // Uncomment to see error messages - + // not required to tag inline callbacks any more check_valid_callback(SpatialPredicateCallbackMissingTag{}, SpatialPredicates{}, v); check_valid_callback(NearestPredicateCallbackMissingTag{}, NearestPredicates{}, v); + // generic lambdas are supported if not using NVCC #ifndef __NVCC__ - check_valid_callback( - [](auto const &predicate, int primitive, auto const &out) {}, - SpatialPredicates{}, v); + check_valid_callback([](auto const & /*predicate*/, int /*primitive*/, + auto const & /*out*/) {}, + SpatialPredicates{}, v); - check_valid_callback([](auto const &predicate, int primitive, float distance, - auto const &out) {}, + check_valid_callback([](auto const & /*predicate*/, int /*primitive*/, + float /*distance*/, auto const & /*out*/) {}, NearestPredicates{}, v); #endif + + // Uncomment to see error messages + + // check_valid_callback(SpatialPredicateCallbackDoesNotTakeCorrectArgument{}, + // SpatialPredicates{}, v); + + // check_valid_callback(ArborX::Details::CallbackDefaultSpatialPredicate{}, + // NearestPredicates{}, v); + + // check_valid_callback(ArborX::Details::CallbackDefaultNearestPredicate{}, + // SpatialPredicates{}, v); } From 67f0e609617733108fd798e0ea8cb46fa5e060b3 Mon Sep 17 00:00:00 2001 From: Damien L-G Date: Fri, 29 May 2020 16:53:21 -0400 Subject: [PATCH 10/10] Fixup unused parameter warnings --- src/details/ArborX_DetailsBoundingVolumeHierarchyImpl.hpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/src/details/ArborX_DetailsBoundingVolumeHierarchyImpl.hpp b/src/details/ArborX_DetailsBoundingVolumeHierarchyImpl.hpp index 50e691fdb..a2b66b437 100644 --- a/src/details/ArborX_DetailsBoundingVolumeHierarchyImpl.hpp +++ b/src/details/ArborX_DetailsBoundingVolumeHierarchyImpl.hpp @@ -337,9 +337,9 @@ check_valid_callback_if_first_argument_is_not_a_view( template std::enable_if_t{} && is_tagged_post_callback{}> -check_valid_callback_if_first_argument_is_not_a_view( - Callback const &callback, Predicates const &predicates, - OutputView const &out) +check_valid_callback_if_first_argument_is_not_a_view(Callback const &, + Predicates const &, + OutputView const &) { // TODO }