From d114f16d71a85ab29165a911c769bd2b77153e50 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Mon, 24 Jul 2023 19:00:16 -0400 Subject: [PATCH 01/23] Allow TreeTraversal to run individual queries in non-batch mode --- .../bvh_driver/benchmark_registration.hpp | 26 +++++++++++++++++++ src/details/ArborX_DetailsTreeTraversal.hpp | 23 +++++++++++++++- 2 files changed, 48 insertions(+), 1 deletion(-) diff --git a/benchmarks/bvh_driver/benchmark_registration.hpp b/benchmarks/bvh_driver/benchmark_registration.hpp index bcdfd6a1d..2bc5c68d0 100644 --- a/benchmarks/bvh_driver/benchmark_registration.hpp +++ b/benchmarks/bvh_driver/benchmark_registration.hpp @@ -259,6 +259,9 @@ void BM_radius_search(benchmark::State &state, Spec const &spec) spec.n_queries, benchmark::Counter::kIsIterationInvariantRate); } +struct Dummy +{}; + template void BM_radius_callback_search(benchmark::State &state, Spec const &spec) { @@ -284,9 +287,32 @@ void BM_radius_callback_search(benchmark::State &state, Spec const &spec) exec_space.fence(); auto const start = std::chrono::high_resolution_clock::now(); +// #define ARBORX_BENCHMARK_USE_BATCHES +#ifdef ARBORX_BENCHMARK_USE_BATCHES index.query(exec_space, queries, callback, ArborX::Experimental::TraversalPolicy().setPredicateSorting( spec.sort_predicates)); +#else + ArborX::Details::LegacyCallbackWrapper< + CountCallback, + ArborX::Details::PairIndexVolume> + wrapped_callback{callback}; + ArborX::Details::TreeTraversal + tree_traversal(index, wrapped_callback); + + auto const n = queries_no_index.extent(0); + + Kokkos::parallel_for( + "ArborX::Benchmarks::RadiusCallbackSearch", + Kokkos::RangePolicy(exec_space, 0, n), + KOKKOS_LAMBDA(int i) { + const auto &query = + ArborX::AccessTraits::get(queries, i); + tree_traversal.search(query); + }); +#endif exec_space.fence(); auto const end = std::chrono::high_resolution_clock::now(); diff --git a/src/details/ArborX_DetailsTreeTraversal.hpp b/src/details/ArborX_DetailsTreeTraversal.hpp index ffc6e0c2e..5fed15c16 100644 --- a/src/details/ArborX_DetailsTreeTraversal.hpp +++ b/src/details/ArborX_DetailsTreeTraversal.hpp @@ -70,6 +70,11 @@ struct TreeTraversal } } + TreeTraversal(BVH const &bvh, Callback const &callback) + : _bvh{bvh} + , _callback{callback} + {} + struct OneLeafTree {}; @@ -88,7 +93,12 @@ struct TreeTraversal KOKKOS_FUNCTION void operator()(int queryIndex) const { auto const &predicate = Access::get(_predicates, queryIndex); + search(predicate); + } + template + KOKKOS_FUNCTION void search(Query const &predicate) const + { int node = HappyTreeFriends::getRoot(_bvh); // start with root do { @@ -392,7 +402,7 @@ struct TreeTraversal; - template + template TreeTraversal(ExecutionSpace const &space, BVH const &bvh, Predicates const &predicates, Callback const &callback) : _bvh{bvh} @@ -422,6 +432,11 @@ struct TreeTraversal + KOKKOS_FUNCTION void search(Query const &predicate) const + { using ArborX::Details::HappyTreeFriends; using distance_type = decltype(predicate.distance( From a35a218d4f62813a870bc8775346b3276abd4b71 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Wed, 26 Jul 2023 10:55:27 -0400 Subject: [PATCH 02/23] Introduce ArborX::kernel_query --- benchmarks/bvh_driver/benchmark_registration.hpp | 7 +------ src/ArborX_CrsGraphWrapper.hpp | 11 +++++++++++ src/details/ArborX_DetailsTreeTraversal.hpp | 4 ++-- 3 files changed, 14 insertions(+), 8 deletions(-) diff --git a/benchmarks/bvh_driver/benchmark_registration.hpp b/benchmarks/bvh_driver/benchmark_registration.hpp index 2bc5c68d0..a27ad01d6 100644 --- a/benchmarks/bvh_driver/benchmark_registration.hpp +++ b/benchmarks/bvh_driver/benchmark_registration.hpp @@ -297,12 +297,7 @@ void BM_radius_callback_search(benchmark::State &state, Spec const &spec) CountCallback, ArborX::Details::PairIndexVolume> wrapped_callback{callback}; - ArborX::Details::TreeTraversal - tree_traversal(index, wrapped_callback); - auto const n = queries_no_index.extent(0); - Kokkos::parallel_for( "ArborX::Benchmarks::RadiusCallbackSearch", Kokkos::RangePolicy(exec_space, 0, n), @@ -310,7 +305,7 @@ void BM_radius_callback_search(benchmark::State &state, Spec const &spec) const auto &query = ArborX::AccessTraits::get(queries, i); - tree_traversal.search(query); + ArborX::kernel_query(index, query, wrapped_callback); }); #endif diff --git a/src/ArborX_CrsGraphWrapper.hpp b/src/ArborX_CrsGraphWrapper.hpp index a7c923419..bb56d26b8 100644 --- a/src/ArborX_CrsGraphWrapper.hpp +++ b/src/ArborX_CrsGraphWrapper.hpp @@ -13,6 +13,7 @@ #define ARBORX_CRS_GRAPH_WRAPPER_HPP #include "ArborX_DetailsCrsGraphWrapperImpl.hpp" +#include "ArborX_DetailsTreeTraversal.hpp" namespace ArborX { @@ -41,6 +42,16 @@ inline void query(Tree const &tree, ExecutionSpace const &space, Kokkos::Profiling::popRegion(); } +template +KOKKOS_FUNCTION void kernel_query(Tree const &tree, Predicate const &predicate, + Callback const &callback) +{ + ArborX::Details::TreeTraversal + tree_traversal(tree, callback); + tree_traversal.search(predicate); +} + } // namespace ArborX #endif diff --git a/src/details/ArborX_DetailsTreeTraversal.hpp b/src/details/ArborX_DetailsTreeTraversal.hpp index 5fed15c16..cc31b3829 100644 --- a/src/details/ArborX_DetailsTreeTraversal.hpp +++ b/src/details/ArborX_DetailsTreeTraversal.hpp @@ -70,7 +70,7 @@ struct TreeTraversal } } - TreeTraversal(BVH const &bvh, Callback const &callback) + KOKKOS_FUNCTION TreeTraversal(BVH const &bvh, Callback const &callback) : _bvh{bvh} , _callback{callback} {} @@ -432,7 +432,7 @@ struct TreeTraversal Date: Wed, 28 Jun 2023 15:17:41 -0400 Subject: [PATCH 03/23] Add Tokamak benchmark --- benchmarks/CMakeLists.txt | 1 + benchmarks/tokamak/CMakeLists.txt | 5 + benchmarks/tokamak/tokamak.cpp | 370 ++++++++++++++++++ src/details/ArborX_Callbacks.hpp | 6 +- src/details/ArborX_DetailsTreeTraversal.hpp | 80 ++-- .../ArborX_DetailsTreeVisualization.hpp | 4 +- 6 files changed, 439 insertions(+), 27 deletions(-) create mode 100644 benchmarks/tokamak/CMakeLists.txt create mode 100644 benchmarks/tokamak/tokamak.cpp diff --git a/benchmarks/CMakeLists.txt b/benchmarks/CMakeLists.txt index 0dfe84d14..cd0ccdc97 100644 --- a/benchmarks/CMakeLists.txt +++ b/benchmarks/CMakeLists.txt @@ -8,6 +8,7 @@ add_subdirectory(bvh_driver) add_subdirectory(dbscan) add_subdirectory(develop) add_subdirectory(execution_space_instances) +add_subdirectory(tokamak) add_subdirectory(union_find) if (ARBORX_ENABLE_MPI) diff --git a/benchmarks/tokamak/CMakeLists.txt b/benchmarks/tokamak/CMakeLists.txt new file mode 100644 index 000000000..cef5f3390 --- /dev/null +++ b/benchmarks/tokamak/CMakeLists.txt @@ -0,0 +1,5 @@ +add_executable(ArborX_Benchmark_Tokamak.exe + tokamak.cpp +) +target_link_libraries(ArborX_Benchmark_Tokamak.exe ArborX::ArborX) +add_test(NAME ArborX_Benchmark_Tokamak.exe COMMAND ./ArborX_Benchmark_Tokamak.exe) diff --git a/benchmarks/tokamak/tokamak.cpp b/benchmarks/tokamak/tokamak.cpp new file mode 100644 index 000000000..78d5628af --- /dev/null +++ b/benchmarks/tokamak/tokamak.cpp @@ -0,0 +1,370 @@ +/**************************************************************************** + * Copyright (c) 2017-2021 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 + +template +struct Triangle +{ + ArborX::ExperimentalHyperGeometry::Point a, b, c; +}; + +struct Mapping +{ + ArborX::ExperimentalHyperGeometry::Point<2> alpha; + ArborX::ExperimentalHyperGeometry::Point<2> beta; + ArborX::ExperimentalHyperGeometry::Point<2> p0; + + ArborX::Point get_coeff(ArborX::ExperimentalHyperGeometry::Point<2> p) const + { + float alpha_coeff = alpha[0] * (p[0] - p0[0]) + alpha[1] * (p[1] - p0[1]); + float beta_coeff = beta[0] * (p[0] - p0[0]) + beta[1] * (p[1] - p0[1]); + return {1 - alpha_coeff - beta_coeff, alpha_coeff, beta_coeff}; + } + + // x = a + alpha * (b - a) + beta * (c - a) + // = (1-beta-alpha) * a + alpha * b + beta * c + void compute(Triangle<2> const &triangle) + { + auto const &a = triangle.a; + auto const &b = triangle.b; + auto const &c = triangle.c; + + ArborX::ExperimentalHyperGeometry::Point<2> u = {b[0] - a[0], b[1] - a[1]}; + ArborX::ExperimentalHyperGeometry::Point<2> v = {c[0] - a[0], c[1] - a[1]}; + + float const inv_det = 1. / (v[1] * u[0] - v[0] * u[1]); + + alpha = ArborX::ExperimentalHyperGeometry::Point<2>{v[1] * inv_det, + -v[0] * inv_det}; + beta = ArborX::ExperimentalHyperGeometry::Point<2>{-u[1] * inv_det, + u[0] * inv_det}; + p0 = a; + } + + Triangle<2> get_triangle() const + { + float const inv_det = 1. / (alpha[0] * beta[1] - alpha[1] * beta[0]); + ArborX::ExperimentalHyperGeometry::Point<2> a = p0; + ArborX::ExperimentalHyperGeometry::Point<2> b = { + {p0[0] + inv_det * beta[1], p0[1] - inv_det * beta[0]}}; + ArborX::ExperimentalHyperGeometry::Point<2> c = { + {p0[0] - inv_det * alpha[1], p0[1] + inv_det * alpha[0]}}; + return {a, b, c}; + } +}; + +template +struct Points +{ + KOKKOS_FUNCTION auto const &get_point(int i) const { return points_(i, 0); } + + KOKKOS_FUNCTION auto size() const { return points_.extent(0); } + + Kokkos::View **, + Kokkos::LayoutLeft, typename DeviceType::memory_space> + points_; +}; + +template +struct Triangles +{ + // Return the number of triangles. + KOKKOS_FUNCTION int size() const { return triangles_.size(); } + + // Return the triangle with index i. + KOKKOS_FUNCTION Triangle<2> const &get_triangle(int i) const + { + return triangles_(i); + } + + KOKKOS_FUNCTION Mapping const &get_mapping(int i) const + { + return mappings_(i); + } + + Kokkos::View *, typename DeviceType::memory_space> triangles_; + Kokkos::View mappings_; +}; + +// For creating the bounding volume hierarchy given a Triangles object, we +// need to define the memory space, how to get the total number of objects, +// and how to access a specific box. Since there are corresponding functions in +// the Triangles class, we just resort to them. +template +struct ArborX::AccessTraits, ArborX::PrimitivesTag> +{ + using memory_space = typename DeviceType::memory_space; + static KOKKOS_FUNCTION int size(Triangles const &triangles) + { + return triangles.size(); + } + static KOKKOS_FUNCTION auto get(Triangles const &triangles, int i) + { + auto const &triangle = triangles.get_triangle(i); + ArborX::ExperimentalHyperGeometry::Box<2> box{}; + box += triangle.a; + box += triangle.b; + box += triangle.c; + return box; + } +}; + +template +class TriangleIntersectionCallback +{ +public: + TriangleIntersectionCallback(Triangles triangles) + : triangles_(triangles) + {} + + template + KOKKOS_FUNCTION void operator()( + Query const &query, + ArborX::Details::PairIndexVolume< + ArborX::ExperimentalHyperGeometry::Box<2>> const &predicate) const + { + ArborX::ExperimentalHyperGeometry::Point<2> const &point = + getGeometry(getPredicate(query)); + auto const &attachment = ArborX::getData(query); + + auto const &triangle = triangles_.get_triangle(predicate.index); + + auto const coeffs = + triangles_.get_mapping(predicate.index).get_coeff(point); + bool intersects = coeffs[0] >= 0 && coeffs[1] >= 0 && coeffs[2] >= 0; + + if (intersects) + { + attachment.triangle_index = predicate.index; + attachment.coeffs = coeffs; + } + } + +private: + Triangles triangles_; +}; + +template +Triangles +parse_stl(typename DeviceType::execution_space const &execution_space) +{ + std::vector> triangles_host; + std::vector mappings_host; + std::ifstream stl_file("RZGrid.stl"); + if (!stl_file.good()) + throw std::runtime_error("Cannot open file"); + std::string line; + std::istringstream in; + Mapping mapping; + Triangle<2> triangle; + std::string dummy; + while (std::getline(stl_file >> std::ws, line)) + { + if (line.find("outer loop") == std::string::npos) + continue; + + std::getline(stl_file >> std::ws, line); + in.str(line); + in >> dummy >> triangle.a[0] >> triangle.a[1]; + + std::getline(stl_file >> std::ws, line); + in.str(line); + in >> dummy >> triangle.b[0] >> triangle.b[1]; + + std::getline(stl_file >> std::ws, line); + in.str(line); + in >> dummy >> triangle.c[0] >> triangle.c[1]; + + mapping.compute(triangle); + + if (triangles_host.size() == 0) + { + std::cout << triangle.a[0] << ' ' << triangle.a[1] << '\n' + << triangle.b[0] << ' ' << triangle.b[1] << '\n' + << triangle.c[0] << ' ' << triangle.c[1] << '\n'; + std::cout << mapping.alpha[0] << ' ' << mapping.alpha[1] << '\n' + << mapping.beta[0] << ' ' << mapping.beta[1] << '\n' + << mapping.p0[0] << ' ' << mapping.p0[1] << '\n'; + } + + triangles_host.push_back(triangle); + mappings_host.push_back(mapping); + } + + std::cout << "Read " << triangles_host.size() << " Triangles\n"; + + Kokkos::View *, typename DeviceType::memory_space> triangles( + Kokkos::view_alloc(Kokkos::WithoutInitializing, "triangles"), + triangles_host.size()); + Kokkos::deep_copy(execution_space, triangles, + Kokkos::View *, Kokkos::HostSpace>( + triangles_host.data(), triangles_host.size())); + + Kokkos::View mappings( + Kokkos::view_alloc(Kokkos::WithoutInitializing, "mappings"), + mappings_host.size()); + Kokkos::deep_copy(execution_space, mappings, + Kokkos::View( + mappings_host.data(), mappings_host.size())); + + return {triangles, mappings}; +} + +template +Points +parse_points(typename DeviceType::execution_space const &execution_space) +{ + std::vector> points_host; + std::ifstream step_file("RK4Steps.txt"); + if (!step_file.good()) + throw std::runtime_error("Cannot open file"); + Mapping mapping; + ArborX::ExperimentalHyperGeometry::Point<2> point; + int id = 0; + int size_per_id = -1; + std::string line; + while (std::getline(step_file, line)) + { + std::stringstream in_line, in_word; + std::string word; + in_line.str(line); + int old_id = id; + for (unsigned int i = 0; i < 3; ++i) + { + std::getline(in_line, word, ','); + in_word << word; + } + in_word >> id >> point[0] >> point[1]; + + if (old_id != id) + { + if (size_per_id == -1) + size_per_id = points_host.size(); + else if (points_host.size() % size_per_id != 0) + { + std::cout << points_host.size() << ' ' << size_per_id << std::endl; + Kokkos::abort("different sizes!"); + } + } + points_host.push_back(point); + if (points_host.size() < 10) + std::cout << id << ' ' << point[0] << ' ' << point[1] << std::endl; + } + + std::cout << "Read " << points_host.size() / size_per_id << " ids which " + << size_per_id << " elements each.\n"; + Kokkos::View **, + Kokkos::LayoutRight, Kokkos::HostSpace> + points_host_view(points_host.data(), points_host.size() / size_per_id, + size_per_id); + std::cout << "id 0, point 1: " << points_host_view(0, 1)[0] << ' ' + << points_host_view(0, 1)[1] << std::endl; + std::cout << "id 1, point 0: " << points_host_view(1, 0)[0] << ' ' + << points_host_view(1, 0)[1] << std::endl; + + Kokkos::View **, + Kokkos::LayoutLeft, typename DeviceType::memory_space> + points(Kokkos::view_alloc(Kokkos::WithoutInitializing, "points"), + points_host.size() / size_per_id, size_per_id); + Kokkos::deep_copy(execution_space, points, points_host_view); + + return {points}; +} + +// Now that we have encapsulated the objects and queries to be used within the +// Triangles class, we can continue with performing the actual search. +int main() +{ + Kokkos::initialize(); + { + using ExecutionSpace = Kokkos::DefaultExecutionSpace; + using MemorySpace = typename ExecutionSpace::memory_space; + using DeviceType = Kokkos::Device; + ExecutionSpace execution_space; + + std::cout << "Create grid with triangles.\n"; + auto triangles = parse_stl(execution_space); + auto points = parse_points(execution_space); + + std::cout << "Creating BVH tree.\n"; + ArborX::BasicBoundingVolumeHierarchy< + MemorySpace, ArborX::ExperimentalHyperGeometry::Box<2>> const + tree(execution_space, triangles); + std::cout << "BVH tree set up.\n"; + + std::cout << "Starting the queries.\n"; + int const n = points.size(); + Kokkos::View offsets("offsets", n); + Kokkos::View coefficients("coefficients", n); + + struct Dummy + {}; + + struct Attachment + { + int &triangle_index; + ArborX::Point &coeffs; + }; + + ArborX::Details::TreeTraversal< + ArborX::BasicBoundingVolumeHierarchy< + MemorySpace, ArborX::ExperimentalHyperGeometry::Box<2>>, + Dummy, TriangleIntersectionCallback, + ArborX::Details::SpatialPredicateTag, + decltype(ArborX::attach( + ArborX::intersects(ArborX::ExperimentalHyperGeometry::Point<2>{}), + std::declval()))> + tree_traversal(tree, + TriangleIntersectionCallback{triangles}); + + std::cout << "n: " << n << std::endl; + + Kokkos::parallel_for( + "ArborX::TreeTraversal::spatial", + Kokkos::RangePolicy(execution_space, 0, n), + KOKKOS_LAMBDA(int i) { + tree_traversal.search( + ArborX::attach(ArborX::intersects(points.get_point(i)), + Attachment{offsets(i), coefficients(i)})); + }); + + std::cout << "Queries done.\n"; + + std::cout << "Starting checking results.\n"; + auto offsets_host = + Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace{}, offsets); + auto coeffs_host = + Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace{}, coefficients); + + for (int i = 0; i < n; ++i) + { + auto const &c = coeffs_host(i); + auto const &t = triangles.get_triangle(offsets_host(i)); + auto const &p_h = points.get_point(i); + auto const p = ArborX::ExperimentalHyperGeometry::Point<2>{ + c[0] * t.a[0] + c[1] * t.b[0] + c[2] * t.c[0], + c[0] * t.a[1] + c[1] * t.b[1] + c[2] * t.c[1]}; + if ((std::abs(p[0] - p_h[0]) > .1) || std::abs(p[1] - p_h[1]) > .1) + { + std::cout << "coeffs for point " << i << " are wrong!\n"; + } + } + + std::cout << "Checking results successful.\n"; + } + + Kokkos::finalize(); +} diff --git a/src/details/ArborX_Callbacks.hpp b/src/details/ArborX_Callbacks.hpp index 3e893fc0d..cca5d108d 100644 --- a/src/details/ArborX_Callbacks.hpp +++ b/src/details/ArborX_Callbacks.hpp @@ -149,8 +149,8 @@ KOKKOS_INLINE_FUNCTION Predicate &&predicate, Primitive &&primitive) { - return ((Callback &&) callback)((Predicate &&) predicate, - (Primitive &&) primitive) == + return ((Callback &&)callback)((Predicate &&)predicate, + (Primitive &&)primitive) == CallbackTreeTraversalControl::early_exit; } @@ -166,7 +166,7 @@ KOKKOS_INLINE_FUNCTION Predicate &&predicate, Primitive &&primitive) { - ((Callback &&) callback)((Predicate &&) predicate, (Primitive &&) primitive); + ((Callback &&)callback)((Predicate &&)predicate, (Primitive &&)primitive); return false; } diff --git a/src/details/ArborX_DetailsTreeTraversal.hpp b/src/details/ArborX_DetailsTreeTraversal.hpp index ffc6e0c2e..b6b11423b 100644 --- a/src/details/ArborX_DetailsTreeTraversal.hpp +++ b/src/details/ArborX_DetailsTreeTraversal.hpp @@ -29,12 +29,15 @@ namespace ArborX namespace Details { -template +template ::get( + std::declval(), 0))> struct TreeTraversal {}; -template -struct TreeTraversal +template +struct TreeTraversal { BVH _bvh; Predicates _predicates; @@ -42,12 +45,22 @@ struct TreeTraversal using Access = AccessTraits; - template - TreeTraversal(ExecutionSpace const &space, BVH const &bvh, - Predicates const &predicates, Callback const &callback) + TreeTraversal() = default; + + TreeTraversal(BVH const &bvh, Predicates const &predicates, + Callback const &callback) : _bvh{bvh} , _predicates{predicates} , _callback{callback} + {} + + TreeTraversal(BVH const &bvh, Callback const &callback) + : _bvh{bvh} + , _callback{callback} + {} + + template + void run(ExecutionSpace const &space) const { if (_bvh.empty()) { @@ -58,14 +71,14 @@ struct TreeTraversal Kokkos::parallel_for( "ArborX::TreeTraversal::spatial::degenerated_one_leaf_tree", Kokkos::RangePolicy( - space, 0, Access::size(predicates)), + space, 0, Access::size(_predicates)), *this); } else { Kokkos::parallel_for("ArborX::TreeTraversal::spatial", Kokkos::RangePolicy( - space, 0, Access::size(predicates)), + space, 0, Access::size(_predicates)), *this); } } @@ -88,7 +101,11 @@ struct TreeTraversal KOKKOS_FUNCTION void operator()(int queryIndex) const { auto const &predicate = Access::get(_predicates, queryIndex); + search(predicate); + } + KOKKOS_FUNCTION void search(Query const &predicate) const + { int node = HappyTreeFriends::getRoot(_bvh); // start with root do { @@ -174,12 +191,15 @@ struct TreeTraversal _buffer = BufferProvider{buffer, offset}; } - template - TreeTraversal(ExecutionSpace const &space, BVH const &bvh, - Predicates const &predicates, Callback const &callback) + TreeTraversal(BVH const &bvh, Predicates const &predicates, + Callback const &callback) : _bvh{bvh} , _predicates{predicates} , _callback{callback} + {} + + template + void run(ExecutionSpace const &space) { if (_bvh.empty()) { @@ -190,7 +210,7 @@ struct TreeTraversal Kokkos::parallel_for( "ArborX::TreeTraversal::nearest::degenerated_one_leaf_tree", Kokkos::RangePolicy( - space, 0, Access::size(predicates)), + space, 0, Access::size(_predicates)), *this); } else @@ -199,7 +219,7 @@ struct TreeTraversal Kokkos::parallel_for("ArborX::TreeTraversal::nearest", Kokkos::RangePolicy( - space, 0, Access::size(predicates)), + space, 0, Access::size(_predicates)), *this); } } @@ -382,9 +402,9 @@ struct TreeTraversal } }; -template +template struct TreeTraversal + Experimental::OrderedSpatialPredicateTag, Query> { BVH _bvh; Predicates _predicates; @@ -392,12 +412,22 @@ struct TreeTraversal; - template - TreeTraversal(ExecutionSpace const &space, BVH const &bvh, - Predicates const &predicates, Callback const &callback) + TreeTraversal() = default; + + TreeTraversal(BVH const &bvh, Predicates const &predicates, + Callback const &callback) : _bvh{bvh} , _predicates{predicates} , _callback{callback} + {} + + TreeTraversal(BVH const &bvh, Callback const &callback) + : _bvh{bvh} + , _callback{callback} + {} + + template + void run(ExecutionSpace const &space) const { if (_bvh.empty()) { @@ -409,7 +439,7 @@ struct TreeTraversal( - space, 0, Access::size(predicates)), + space, 0, Access::size(_predicates)), *this); } else @@ -417,7 +447,7 @@ struct TreeTraversal(space, 0, - Access::size(predicates)), + Access::size(_predicates)), *this); } } @@ -444,6 +474,11 @@ struct TreeTraversal; using Tag = typename AccessTraitsHelper::tag; - TreeTraversal(space, bvh, predicates, - callback); + TreeTraversal tree_traversal(bvh, predicates, + callback); + tree_traversal.run(space); } } // namespace Details diff --git a/src/details/ArborX_DetailsTreeVisualization.hpp b/src/details/ArborX_DetailsTreeVisualization.hpp index d4561da67..6eb7d6e09 100644 --- a/src/details/ArborX_DetailsTreeVisualization.hpp +++ b/src/details/ArborX_DetailsTreeVisualization.hpp @@ -200,8 +200,8 @@ struct TreeVisualization predicates(0) = pred; TreeTraversal - tree_traversal(ExecutionSpace{}, tree, predicates, - Callback{tree, visitor, permute}); + tree_traversal(tree, predicates, Callback{tree, visitor, permute}); + tree_traversal.run(ExecutionSpace{}); #endif } }; From bbc7dee935bdb58e6406e0e9a3f0e075acb521d3 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Thu, 29 Jun 2023 13:50:09 -0400 Subject: [PATCH 04/23] Perform multiple queries per thread, clean up benchmark --- benchmarks/tokamak/tokamak.cpp | 71 ++++++++++++++-------------------- 1 file changed, 29 insertions(+), 42 deletions(-) diff --git a/benchmarks/tokamak/tokamak.cpp b/benchmarks/tokamak/tokamak.cpp index 78d5628af..e9a6fcafe 100644 --- a/benchmarks/tokamak/tokamak.cpp +++ b/benchmarks/tokamak/tokamak.cpp @@ -66,18 +66,6 @@ struct Mapping } }; -template -struct Points -{ - KOKKOS_FUNCTION auto const &get_point(int i) const { return points_(i, 0); } - - KOKKOS_FUNCTION auto size() const { return points_.extent(0); } - - Kokkos::View **, - Kokkos::LayoutLeft, typename DeviceType::memory_space> - points_; -}; - template struct Triangles { @@ -224,7 +212,8 @@ parse_stl(typename DeviceType::execution_space const &execution_space) } template -Points +Kokkos::View **, Kokkos::LayoutLeft, + typename DeviceType::memory_space> parse_points(typename DeviceType::execution_space const &execution_space) { std::vector> points_host; @@ -281,7 +270,7 @@ parse_points(typename DeviceType::execution_space const &execution_space) points_host.size() / size_per_id, size_per_id); Kokkos::deep_copy(execution_space, points, points_host_view); - return {points}; + return points; } // Now that we have encapsulated the objects and queries to be used within the @@ -306,9 +295,7 @@ int main() std::cout << "BVH tree set up.\n"; std::cout << "Starting the queries.\n"; - int const n = points.size(); - Kokkos::View offsets("offsets", n); - Kokkos::View coefficients("coefficients", n); + int const n = points.extent(0); struct Dummy {}; @@ -336,34 +323,34 @@ int main() "ArborX::TreeTraversal::spatial", Kokkos::RangePolicy(execution_space, 0, n), KOKKOS_LAMBDA(int i) { - tree_traversal.search( - ArborX::attach(ArborX::intersects(points.get_point(i)), - Attachment{offsets(i), coefficients(i)})); + int triangle_index = 0; + ArborX::Point coefficients{}; + for (unsigned int j = 0; j < points.extent(1); ++j) + { + auto const &point = points(i, j); + auto const &triangle = triangles.get_triangle(triangle_index); + auto const &test_coeffs = + triangles.get_mapping(triangle_index).get_coeff(point); + bool intersects = test_coeffs[0] >= 0 && test_coeffs[1] >= 0 && + test_coeffs[2] >= 0; + if (intersects) + { + coefficients = test_coeffs; + KOKKOS_IMPL_DO_NOT_USE_PRINTF("%d, %d: same triangle\n", i, j); + } + else + { + tree_traversal.search( + ArborX::attach(ArborX::intersects(point), + Attachment{triangle_index, coefficients})); + KOKKOS_IMPL_DO_NOT_USE_PRINTF("%d, %d: %d %f %f %f\n", i, j, + triangle_index, coefficients[0], + coefficients[1], coefficients[2]); + } + } }); std::cout << "Queries done.\n"; - - std::cout << "Starting checking results.\n"; - auto offsets_host = - Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace{}, offsets); - auto coeffs_host = - Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace{}, coefficients); - - for (int i = 0; i < n; ++i) - { - auto const &c = coeffs_host(i); - auto const &t = triangles.get_triangle(offsets_host(i)); - auto const &p_h = points.get_point(i); - auto const p = ArborX::ExperimentalHyperGeometry::Point<2>{ - c[0] * t.a[0] + c[1] * t.b[0] + c[2] * t.c[0], - c[0] * t.a[1] + c[1] * t.b[1] + c[2] * t.c[1]}; - if ((std::abs(p[0] - p_h[0]) > .1) || std::abs(p[1] - p_h[1]) > .1) - { - std::cout << "coeffs for point " << i << " are wrong!\n"; - } - } - - std::cout << "Checking results successful.\n"; } Kokkos::finalize(); From bb0ff33b3c1bc89f3c34ab2cbac1c3e052ba6910 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Thu, 29 Jun 2023 14:09:34 -0400 Subject: [PATCH 05/23] Clean up ArborX --- benchmarks/tokamak/tokamak.cpp | 5 +---- src/details/ArborX_DetailsTreeTraversal.hpp | 15 +++++++-------- 2 files changed, 8 insertions(+), 12 deletions(-) diff --git a/benchmarks/tokamak/tokamak.cpp b/benchmarks/tokamak/tokamak.cpp index e9a6fcafe..5662d02ac 100644 --- a/benchmarks/tokamak/tokamak.cpp +++ b/benchmarks/tokamak/tokamak.cpp @@ -310,10 +310,7 @@ int main() ArborX::BasicBoundingVolumeHierarchy< MemorySpace, ArborX::ExperimentalHyperGeometry::Box<2>>, Dummy, TriangleIntersectionCallback, - ArborX::Details::SpatialPredicateTag, - decltype(ArborX::attach( - ArborX::intersects(ArborX::ExperimentalHyperGeometry::Point<2>{}), - std::declval()))> + ArborX::Details::SpatialPredicateTag> tree_traversal(tree, TriangleIntersectionCallback{triangles}); diff --git a/src/details/ArborX_DetailsTreeTraversal.hpp b/src/details/ArborX_DetailsTreeTraversal.hpp index b6b11423b..8250bda2c 100644 --- a/src/details/ArborX_DetailsTreeTraversal.hpp +++ b/src/details/ArborX_DetailsTreeTraversal.hpp @@ -29,15 +29,12 @@ namespace ArborX namespace Details { -template ::get( - std::declval(), 0))> +template struct TreeTraversal {}; -template -struct TreeTraversal +template +struct TreeTraversal { BVH _bvh; Predicates _predicates; @@ -104,6 +101,7 @@ struct TreeTraversal search(predicate); } + template KOKKOS_FUNCTION void search(Query const &predicate) const { int node = HappyTreeFriends::getRoot(_bvh); // start with root @@ -402,9 +400,9 @@ struct TreeTraversal } }; -template +template struct TreeTraversal + Experimental::OrderedSpatialPredicateTag> { BVH _bvh; Predicates _predicates; @@ -477,6 +475,7 @@ struct TreeTraversal KOKKOS_FUNCTION void search(Query const &predicate) const { using ArborX::Details::HappyTreeFriends; From ea2b7502bbea07eb7b16a633c5afed18f6bce7f4 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Mon, 3 Jul 2023 17:10:02 +0000 Subject: [PATCH 06/23] Fix for Cuda --- benchmarks/tokamak/tokamak.cpp | 11 +++++++---- 1 file changed, 7 insertions(+), 4 deletions(-) diff --git a/benchmarks/tokamak/tokamak.cpp b/benchmarks/tokamak/tokamak.cpp index 5662d02ac..ff60d2836 100644 --- a/benchmarks/tokamak/tokamak.cpp +++ b/benchmarks/tokamak/tokamak.cpp @@ -27,6 +27,7 @@ struct Mapping ArborX::ExperimentalHyperGeometry::Point<2> beta; ArborX::ExperimentalHyperGeometry::Point<2> p0; + KOKKOS_FUNCTION ArborX::Point get_coeff(ArborX::ExperimentalHyperGeometry::Point<2> p) const { float alpha_coeff = alpha[0] * (p[0] - p0[0]) + alpha[1] * (p[1] - p0[1]); @@ -268,11 +269,16 @@ parse_points(typename DeviceType::execution_space const &execution_space) Kokkos::LayoutLeft, typename DeviceType::memory_space> points(Kokkos::view_alloc(Kokkos::WithoutInitializing, "points"), points_host.size() / size_per_id, size_per_id); - Kokkos::deep_copy(execution_space, points, points_host_view); + auto points_tmp_view = Kokkos::create_mirror_view_and_copy( + typename DeviceType::memory_space{}, points_host_view); + Kokkos::deep_copy(execution_space, points, points_tmp_view); return points; } +struct Dummy +{}; + // Now that we have encapsulated the objects and queries to be used within the // Triangles class, we can continue with performing the actual search. int main() @@ -297,9 +303,6 @@ int main() std::cout << "Starting the queries.\n"; int const n = points.extent(0); - struct Dummy - {}; - struct Attachment { int &triangle_index; From 100818993376ae01b3e65d49b819b78e637bd9fd Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Fri, 4 Aug 2023 08:44:57 -0500 Subject: [PATCH 07/23] Use HyperTriangle --- benchmarks/tokamak/tokamak.cpp | 48 +++++++++++++++++----------------- 1 file changed, 24 insertions(+), 24 deletions(-) diff --git a/benchmarks/tokamak/tokamak.cpp b/benchmarks/tokamak/tokamak.cpp index ff60d2836..afc65f2ed 100644 --- a/benchmarks/tokamak/tokamak.cpp +++ b/benchmarks/tokamak/tokamak.cpp @@ -10,17 +10,12 @@ ****************************************************************************/ #include +#include #include #include -template -struct Triangle -{ - ArborX::ExperimentalHyperGeometry::Point a, b, c; -}; - struct Mapping { ArborX::ExperimentalHyperGeometry::Point<2> alpha; @@ -37,7 +32,7 @@ struct Mapping // x = a + alpha * (b - a) + beta * (c - a) // = (1-beta-alpha) * a + alpha * b + beta * c - void compute(Triangle<2> const &triangle) + void compute(ArborX::ExperimentalHyperGeometry::Triangle<2> const &triangle) { auto const &a = triangle.a; auto const &b = triangle.b; @@ -55,7 +50,7 @@ struct Mapping p0 = a; } - Triangle<2> get_triangle() const + ArborX::ExperimentalHyperGeometry::Triangle<2> get_triangle() const { float const inv_det = 1. / (alpha[0] * beta[1] - alpha[1] * beta[0]); ArborX::ExperimentalHyperGeometry::Point<2> a = p0; @@ -74,7 +69,8 @@ struct Triangles KOKKOS_FUNCTION int size() const { return triangles_.size(); } // Return the triangle with index i. - KOKKOS_FUNCTION Triangle<2> const &get_triangle(int i) const + KOKKOS_FUNCTION ArborX::ExperimentalHyperGeometry::Triangle<2> const & + get_triangle(int i) const { return triangles_(i); } @@ -84,7 +80,9 @@ struct Triangles return mappings_(i); } - Kokkos::View *, typename DeviceType::memory_space> triangles_; + Kokkos::View *, + typename DeviceType::memory_space> + triangles_; Kokkos::View mappings_; }; @@ -150,7 +148,7 @@ template Triangles parse_stl(typename DeviceType::execution_space const &execution_space) { - std::vector> triangles_host; + std::vector> triangles_host; std::vector mappings_host; std::ifstream stl_file("RZGrid.stl"); if (!stl_file.good()) @@ -158,7 +156,7 @@ parse_stl(typename DeviceType::execution_space const &execution_space) std::string line; std::istringstream in; Mapping mapping; - Triangle<2> triangle; + ArborX::ExperimentalHyperGeometry::Triangle<2> triangle; std::string dummy; while (std::getline(stl_file >> std::ws, line)) { @@ -195,12 +193,15 @@ parse_stl(typename DeviceType::execution_space const &execution_space) std::cout << "Read " << triangles_host.size() << " Triangles\n"; - Kokkos::View *, typename DeviceType::memory_space> triangles( - Kokkos::view_alloc(Kokkos::WithoutInitializing, "triangles"), - triangles_host.size()); - Kokkos::deep_copy(execution_space, triangles, - Kokkos::View *, Kokkos::HostSpace>( - triangles_host.data(), triangles_host.size())); + Kokkos::View *, + typename DeviceType::memory_space> + triangles(Kokkos::view_alloc(Kokkos::WithoutInitializing, "triangles"), + triangles_host.size()); + Kokkos::deep_copy( + execution_space, triangles, + Kokkos::View *, + Kokkos::HostSpace>(triangles_host.data(), + triangles_host.size())); Kokkos::View mappings( Kokkos::view_alloc(Kokkos::WithoutInitializing, "mappings"), @@ -296,7 +297,8 @@ int main() std::cout << "Creating BVH tree.\n"; ArborX::BasicBoundingVolumeHierarchy< - MemorySpace, ArborX::ExperimentalHyperGeometry::Box<2>> const + MemorySpace, ArborX::Details::PairIndexVolume< + ArborX::ExperimentalHyperGeometry::Box<2>>> const tree(execution_space, triangles); std::cout << "BVH tree set up.\n"; @@ -309,11 +311,9 @@ int main() ArborX::Point &coeffs; }; - ArborX::Details::TreeTraversal< - ArborX::BasicBoundingVolumeHierarchy< - MemorySpace, ArborX::ExperimentalHyperGeometry::Box<2>>, - Dummy, TriangleIntersectionCallback, - ArborX::Details::SpatialPredicateTag> + ArborX::Details::TreeTraversal, + ArborX::Details::SpatialPredicateTag> tree_traversal(tree, TriangleIntersectionCallback{triangles}); From 40b8ad9d2fae65a4f6e0a7e82b97420c615e9ba7 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Mon, 7 Aug 2023 16:40:23 -0400 Subject: [PATCH 08/23] Use ArborX::kernel_query --- benchmarks/tokamak/tokamak.cpp | 12 ++++-------- 1 file changed, 4 insertions(+), 8 deletions(-) diff --git a/benchmarks/tokamak/tokamak.cpp b/benchmarks/tokamak/tokamak.cpp index afc65f2ed..53e49e747 100644 --- a/benchmarks/tokamak/tokamak.cpp +++ b/benchmarks/tokamak/tokamak.cpp @@ -311,12 +311,6 @@ int main() ArborX::Point &coeffs; }; - ArborX::Details::TreeTraversal, - ArborX::Details::SpatialPredicateTag> - tree_traversal(tree, - TriangleIntersectionCallback{triangles}); - std::cout << "n: " << n << std::endl; Kokkos::parallel_for( @@ -340,9 +334,11 @@ int main() } else { - tree_traversal.search( + ArborX::kernel_query( + tree, ArborX::attach(ArborX::intersects(point), - Attachment{triangle_index, coefficients})); + Attachment{triangle_index, coefficients}), + TriangleIntersectionCallback{triangles}); KOKKOS_IMPL_DO_NOT_USE_PRINTF("%d, %d: %d %f %f %f\n", i, j, triangle_index, coefficients[0], coefficients[1], coefficients[2]); From 70e92fc47503fd185802d99198c8253959503129 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Mon, 7 Aug 2023 16:46:25 -0400 Subject: [PATCH 09/23] Use ArborX::CallbackTreeTraversalControl::early_exit --- benchmarks/tokamak/tokamak.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/benchmarks/tokamak/tokamak.cpp b/benchmarks/tokamak/tokamak.cpp index 53e49e747..0516f0b2a 100644 --- a/benchmarks/tokamak/tokamak.cpp +++ b/benchmarks/tokamak/tokamak.cpp @@ -118,7 +118,7 @@ class TriangleIntersectionCallback {} template - KOKKOS_FUNCTION void operator()( + KOKKOS_FUNCTION auto operator()( Query const &query, ArborX::Details::PairIndexVolume< ArborX::ExperimentalHyperGeometry::Box<2>> const &predicate) const @@ -137,7 +137,9 @@ class TriangleIntersectionCallback { attachment.triangle_index = predicate.index; attachment.coeffs = coeffs; + return ArborX::CallbackTreeTraversalControl::early_exit; } + return ArborX::CallbackTreeTraversalControl::normal_continuation; } private: From 005fd807f6d6234c538187a8f521ff4291c17c11 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Tue, 19 Sep 2023 16:07:51 -0400 Subject: [PATCH 10/23] Clean up tokamak example --- benchmarks/tokamak/tokamak.cpp | 185 +++++++++++++-------------------- 1 file changed, 70 insertions(+), 115 deletions(-) diff --git a/benchmarks/tokamak/tokamak.cpp b/benchmarks/tokamak/tokamak.cpp index 0516f0b2a..f675ea924 100644 --- a/benchmarks/tokamak/tokamak.cpp +++ b/benchmarks/tokamak/tokamak.cpp @@ -16,74 +16,23 @@ #include -struct Mapping -{ - ArborX::ExperimentalHyperGeometry::Point<2> alpha; - ArborX::ExperimentalHyperGeometry::Point<2> beta; - ArborX::ExperimentalHyperGeometry::Point<2> p0; - - KOKKOS_FUNCTION - ArborX::Point get_coeff(ArborX::ExperimentalHyperGeometry::Point<2> p) const - { - float alpha_coeff = alpha[0] * (p[0] - p0[0]) + alpha[1] * (p[1] - p0[1]); - float beta_coeff = beta[0] * (p[0] - p0[0]) + beta[1] * (p[1] - p0[1]); - return {1 - alpha_coeff - beta_coeff, alpha_coeff, beta_coeff}; - } - - // x = a + alpha * (b - a) + beta * (c - a) - // = (1-beta-alpha) * a + alpha * b + beta * c - void compute(ArborX::ExperimentalHyperGeometry::Triangle<2> const &triangle) - { - auto const &a = triangle.a; - auto const &b = triangle.b; - auto const &c = triangle.c; - - ArborX::ExperimentalHyperGeometry::Point<2> u = {b[0] - a[0], b[1] - a[1]}; - ArborX::ExperimentalHyperGeometry::Point<2> v = {c[0] - a[0], c[1] - a[1]}; - - float const inv_det = 1. / (v[1] * u[0] - v[0] * u[1]); - - alpha = ArborX::ExperimentalHyperGeometry::Point<2>{v[1] * inv_det, - -v[0] * inv_det}; - beta = ArborX::ExperimentalHyperGeometry::Point<2>{-u[1] * inv_det, - u[0] * inv_det}; - p0 = a; - } - - ArborX::ExperimentalHyperGeometry::Triangle<2> get_triangle() const - { - float const inv_det = 1. / (alpha[0] * beta[1] - alpha[1] * beta[0]); - ArborX::ExperimentalHyperGeometry::Point<2> a = p0; - ArborX::ExperimentalHyperGeometry::Point<2> b = { - {p0[0] + inv_det * beta[1], p0[1] - inv_det * beta[0]}}; - ArborX::ExperimentalHyperGeometry::Point<2> c = { - {p0[0] - inv_det * alpha[1], p0[1] + inv_det * alpha[0]}}; - return {a, b, c}; - } -}; +using Point = ArborX::ExperimentalHyperGeometry::Point<2>; +using Triangle = ArborX::ExperimentalHyperGeometry::Triangle<2>; -template +template struct Triangles { // Return the number of triangles. - KOKKOS_FUNCTION int size() const { return triangles_.size(); } + KOKKOS_FUNCTION int size() const { return _triangles.size(); } // Return the triangle with index i. - KOKKOS_FUNCTION ArborX::ExperimentalHyperGeometry::Triangle<2> const & - get_triangle(int i) const + KOKKOS_FUNCTION Triangle const &operator()(int i) const { - return triangles_(i); + return _triangles(i); } - KOKKOS_FUNCTION Mapping const &get_mapping(int i) const - { - return mappings_(i); - } - - Kokkos::View *, - typename DeviceType::memory_space> - triangles_; - Kokkos::View mappings_; + Kokkos::View *, MemorySpace> + _triangles; }; // For creating the bounding volume hierarchy given a Triangles object, we @@ -100,7 +49,7 @@ struct ArborX::AccessTraits, ArborX::PrimitivesTag> } static KOKKOS_FUNCTION auto get(Triangles const &triangles, int i) { - auto const &triangle = triangles.get_triangle(i); + auto const &triangle = triangles(i); ArborX::ExperimentalHyperGeometry::Box<2> box{}; box += triangle.a; box += triangle.b; @@ -109,33 +58,61 @@ struct ArborX::AccessTraits, ArborX::PrimitivesTag> } }; +KOKKOS_FUNCTION +ArborX::Point compute_barycentric_coordinates(Triangle const &triangle, + Point const &point) +{ + auto const &a = triangle.a; + auto const &b = triangle.b; + auto const &c = triangle.c; + + // Find coefficients alpha and beta such that + // x = a + alpha * (b - a) + beta * (c - a) + // = (1 - alpha - beta) * a + alpha * b + beta * c + // recognizing the linear system + // ((b - a) (c - a)) (alpha beta)^T = (x - a) + float u[2] = {b[0] - a[0], b[1] - a[1]}; + float v[2] = {c[0] - a[0], c[1] - a[1]}; + float const det = v[1] * u[0] - v[0] * u[1]; + if (det == 0) + Kokkos::abort("Degenerate triangles are not supported!"); + float const inv_det = 1.f / det; + + float alpha[2] = {v[1] * inv_det, -v[0] * inv_det}; + float beta[2] = {-u[1] * inv_det, u[0] * inv_det}; + + float alpha_coeff = + alpha[0] * (point[0] - a[0]) + alpha[1] * (point[1] - a[1]); + float beta_coeff = beta[0] * (point[0] - a[0]) + beta[1] * (point[1] - a[1]); + + return {1 - alpha_coeff - beta_coeff, alpha_coeff, beta_coeff}; +} + template class TriangleIntersectionCallback { public: - TriangleIntersectionCallback(Triangles triangles) - : triangles_(triangles) + TriangleIntersectionCallback( + Kokkos::View triangles) + : _triangles(triangles) {} - template - KOKKOS_FUNCTION auto operator()( - Query const &query, - ArborX::Details::PairIndexVolume< - ArborX::ExperimentalHyperGeometry::Box<2>> const &predicate) const + template + KOKKOS_FUNCTION auto operator()(Query const &query, + Primitive const &primitive) const { - ArborX::ExperimentalHyperGeometry::Point<2> const &point = - getGeometry(getPredicate(query)); + Point const &point = getGeometry(getPredicate(query)); auto const &attachment = ArborX::getData(query); + auto triangle_index = primitive.index; + Triangle const &triangle = _triangles(triangle_index); - auto const &triangle = triangles_.get_triangle(predicate.index); + ArborX::Point coeffs = compute_barycentric_coordinates(triangle, point); - auto const coeffs = - triangles_.get_mapping(predicate.index).get_coeff(point); bool intersects = coeffs[0] >= 0 && coeffs[1] >= 0 && coeffs[2] >= 0; if (intersects) { - attachment.triangle_index = predicate.index; + attachment.triangle_index = triangle_index; attachment.coeffs = coeffs; return ArborX::CallbackTreeTraversalControl::early_exit; } @@ -143,22 +120,20 @@ class TriangleIntersectionCallback } private: - Triangles triangles_; + Kokkos::View _triangles; }; template -Triangles +Kokkos::View parse_stl(typename DeviceType::execution_space const &execution_space) { - std::vector> triangles_host; - std::vector mappings_host; + std::vector triangles_host; std::ifstream stl_file("RZGrid.stl"); if (!stl_file.good()) throw std::runtime_error("Cannot open file"); std::string line; std::istringstream in; - Mapping mapping; - ArborX::ExperimentalHyperGeometry::Triangle<2> triangle; + Triangle triangle; std::string dummy; while (std::getline(stl_file >> std::ws, line)) { @@ -177,55 +152,37 @@ parse_stl(typename DeviceType::execution_space const &execution_space) in.str(line); in >> dummy >> triangle.c[0] >> triangle.c[1]; - mapping.compute(triangle); - if (triangles_host.size() == 0) { std::cout << triangle.a[0] << ' ' << triangle.a[1] << '\n' << triangle.b[0] << ' ' << triangle.b[1] << '\n' << triangle.c[0] << ' ' << triangle.c[1] << '\n'; - std::cout << mapping.alpha[0] << ' ' << mapping.alpha[1] << '\n' - << mapping.beta[0] << ' ' << mapping.beta[1] << '\n' - << mapping.p0[0] << ' ' << mapping.p0[1] << '\n'; } triangles_host.push_back(triangle); - mappings_host.push_back(mapping); } std::cout << "Read " << triangles_host.size() << " Triangles\n"; - Kokkos::View *, - typename DeviceType::memory_space> - triangles(Kokkos::view_alloc(Kokkos::WithoutInitializing, "triangles"), - triangles_host.size()); - Kokkos::deep_copy( - execution_space, triangles, - Kokkos::View *, - Kokkos::HostSpace>(triangles_host.data(), - triangles_host.size())); - - Kokkos::View mappings( - Kokkos::view_alloc(Kokkos::WithoutInitializing, "mappings"), - mappings_host.size()); - Kokkos::deep_copy(execution_space, mappings, - Kokkos::View( - mappings_host.data(), mappings_host.size())); - - return {triangles, mappings}; + Kokkos::View triangles( + Kokkos::view_alloc(Kokkos::WithoutInitializing, "triangles"), + triangles_host.size()); + Kokkos::deep_copy(execution_space, triangles, + Kokkos::View( + triangles_host.data(), triangles_host.size())); + + return {triangles}; } template -Kokkos::View **, Kokkos::LayoutLeft, - typename DeviceType::memory_space> +Kokkos::View parse_points(typename DeviceType::execution_space const &execution_space) { - std::vector> points_host; + std::vector points_host; std::ifstream step_file("RK4Steps.txt"); if (!step_file.good()) throw std::runtime_error("Cannot open file"); - Mapping mapping; - ArborX::ExperimentalHyperGeometry::Point<2> point; + Point point; int id = 0; int size_per_id = -1; std::string line; @@ -259,8 +216,7 @@ parse_points(typename DeviceType::execution_space const &execution_space) std::cout << "Read " << points_host.size() / size_per_id << " ids which " << size_per_id << " elements each.\n"; - Kokkos::View **, - Kokkos::LayoutRight, Kokkos::HostSpace> + Kokkos::View points_host_view(points_host.data(), points_host.size() / size_per_id, size_per_id); std::cout << "id 0, point 1: " << points_host_view(0, 1)[0] << ' ' @@ -268,8 +224,7 @@ parse_points(typename DeviceType::execution_space const &execution_space) std::cout << "id 1, point 0: " << points_host_view(1, 0)[0] << ' ' << points_host_view(1, 0)[1] << std::endl; - Kokkos::View **, - Kokkos::LayoutLeft, typename DeviceType::memory_space> + Kokkos::View points(Kokkos::view_alloc(Kokkos::WithoutInitializing, "points"), points_host.size() / size_per_id, size_per_id); auto points_tmp_view = Kokkos::create_mirror_view_and_copy( @@ -301,7 +256,7 @@ int main() ArborX::BasicBoundingVolumeHierarchy< MemorySpace, ArborX::Details::PairIndexVolume< ArborX::ExperimentalHyperGeometry::Box<2>>> const - tree(execution_space, triangles); + tree(execution_space, Triangles{triangles}); std::cout << "BVH tree set up.\n"; std::cout << "Starting the queries.\n"; @@ -324,9 +279,9 @@ int main() for (unsigned int j = 0; j < points.extent(1); ++j) { auto const &point = points(i, j); - auto const &triangle = triangles.get_triangle(triangle_index); + auto const &triangle = triangles(triangle_index); auto const &test_coeffs = - triangles.get_mapping(triangle_index).get_coeff(point); + compute_barycentric_coordinates(triangle, point); bool intersects = test_coeffs[0] >= 0 && test_coeffs[1] >= 0 && test_coeffs[2] >= 0; if (intersects) From 7670e3fb1cdd63089d92a5fcf5d80d43990d76b5 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Mon, 25 Sep 2023 15:11:56 -0400 Subject: [PATCH 11/23] Add dummy inputs --- build/benchmarks/tokamak/RK4Steps.txt | 20 ++++++++++++ build/benchmarks/tokamak/RZGrid.stl | 44 +++++++++++++++++++++++++++ 2 files changed, 64 insertions(+) create mode 100644 build/benchmarks/tokamak/RK4Steps.txt create mode 100644 build/benchmarks/tokamak/RZGrid.stl diff --git a/build/benchmarks/tokamak/RK4Steps.txt b/build/benchmarks/tokamak/RK4Steps.txt new file mode 100644 index 000000000..2ecbd177d --- /dev/null +++ b/build/benchmarks/tokamak/RK4Steps.txt @@ -0,0 +1,20 @@ +0 0.24, 0.24 +0 0.74, 0.24 +0 0.24, 0.74 +0 0.74, 0.74 +1 0.74, 0.24 +1 1.24, 0.24 +1 0.74, 0.74 +1 1.24, 0.74 +2 1.24, 0.24 +2 1.74, 0.24 +2 1.24, 0.74 +2 1.74, 0.74 +3 1.74, 0.24 +3 2.24, 0.24 +3 1.74, 0.74 +3 2.24, 0.74 +4 2.24, 0.24 +4 2.74, 0.24 +4 2.24, 0.74 +4 2.74, 0.74 diff --git a/build/benchmarks/tokamak/RZGrid.stl b/build/benchmarks/tokamak/RZGrid.stl new file mode 100644 index 000000000..450c48784 --- /dev/null +++ b/build/benchmarks/tokamak/RZGrid.stl @@ -0,0 +1,44 @@ +solid ascii + facet normal 0 -0 -1 + outer loop + vertex 0 0 0 + vertex 1 0 0 + vertex 0 1 0 + endloop + endfacet + facet normal 0 0 -1 + outer loop + vertex 1 0 0 + vertex 1 1 0 + vertex 0 1 0 + endloop + endfacet + facet normal 0 0 -1 + outer loop + vertex 1 0 0 + vertex 2 0 0 + vertex 1 1 0 + endloop + endfacet + facet normal 0 0 -1 + outer loop + vertex 2 0 0 + vertex 2 1 0 + vertex 1 1 0 + endloop + endfacet + facet normal 0 0 -1 + outer loop + vertex 2 0 0 + vertex 3 0 0 + vertex 2 1 0 + endloop + endfacet + facet normal 0 0 -1 + outer loop + vertex 3 0 0 + vertex 3 1 0 + vertex 2 1 0 + endloop + endfacet +endsolid From 45af564290d0ae5fde93357ef5ff14f1aa89a19b Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Mon, 25 Sep 2023 15:31:48 -0400 Subject: [PATCH 12/23] Print point and triangle --- benchmarks/tokamak/tokamak.cpp | 9 +++++---- src/details/ArborX_Callbacks.hpp | 6 +++--- src/details/ArborX_DetailsTreeVisualization.hpp | 2 +- 3 files changed, 9 insertions(+), 8 deletions(-) diff --git a/benchmarks/tokamak/tokamak.cpp b/benchmarks/tokamak/tokamak.cpp index 4c09a9533..b2fc486b8 100644 --- a/benchmarks/tokamak/tokamak.cpp +++ b/benchmarks/tokamak/tokamak.cpp @@ -287,7 +287,9 @@ int main() if (intersects) { coefficients = test_coeffs; - KOKKOS_IMPL_DO_NOT_USE_PRINTF("%d, %d: same triangle\n", i, j); + KOKKOS_IMPL_DO_NOT_USE_PRINTF("%d, %d: %f %f in %d (same)\n", i, + j, point[0], point[1], + triangle_index); } else { @@ -296,9 +298,8 @@ int main() ArborX::attach(ArborX::intersects(point), Attachment{triangle_index, coefficients}), TriangleIntersectionCallback{triangles}); - KOKKOS_IMPL_DO_NOT_USE_PRINTF("%d, %d: %d %f %f %f\n", i, j, - triangle_index, coefficients[0], - coefficients[1], coefficients[2]); + KOKKOS_IMPL_DO_NOT_USE_PRINTF("%d, %d: %f %f in %d\n", i, j, + point[0], point[1], triangle_index); } } }); diff --git a/src/details/ArborX_Callbacks.hpp b/src/details/ArborX_Callbacks.hpp index e885665f9..91b29de8f 100644 --- a/src/details/ArborX_Callbacks.hpp +++ b/src/details/ArborX_Callbacks.hpp @@ -145,8 +145,8 @@ KOKKOS_INLINE_FUNCTION Predicate &&predicate, Primitive &&primitive) { - return ((Callback &&)callback)((Predicate &&)predicate, - (Primitive &&)primitive) == + return ((Callback &&) callback)((Predicate &&) predicate, + (Primitive &&) primitive) == CallbackTreeTraversalControl::early_exit; } @@ -162,7 +162,7 @@ KOKKOS_INLINE_FUNCTION Predicate &&predicate, Primitive &&primitive) { - ((Callback &&)callback)((Predicate &&)predicate, (Primitive &&)primitive); + ((Callback &&) callback)((Predicate &&) predicate, (Primitive &&) primitive); return false; } diff --git a/src/details/ArborX_DetailsTreeVisualization.hpp b/src/details/ArborX_DetailsTreeVisualization.hpp index 96019b987..d4561da67 100644 --- a/src/details/ArborX_DetailsTreeVisualization.hpp +++ b/src/details/ArborX_DetailsTreeVisualization.hpp @@ -200,7 +200,7 @@ struct TreeVisualization predicates(0) = pred; TreeTraversal - tree_traversal(space, tree, predicates, + tree_traversal(ExecutionSpace{}, tree, predicates, Callback{tree, visitor, permute}); #endif } From ffcf399662677603fd62737b937a10502bf0acab Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Mon, 25 Sep 2023 15:48:44 -0400 Subject: [PATCH 13/23] Also print coefficients --- benchmarks/tokamak/tokamak.cpp | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) diff --git a/benchmarks/tokamak/tokamak.cpp b/benchmarks/tokamak/tokamak.cpp index b2fc486b8..5b80165ec 100644 --- a/benchmarks/tokamak/tokamak.cpp +++ b/benchmarks/tokamak/tokamak.cpp @@ -287,9 +287,10 @@ int main() if (intersects) { coefficients = test_coeffs; - KOKKOS_IMPL_DO_NOT_USE_PRINTF("%d, %d: %f %f in %d (same)\n", i, - j, point[0], point[1], - triangle_index); + KOKKOS_IMPL_DO_NOT_USE_PRINTF( + "%d, %d: %f %f in %d (same), coefficients: %f, %f, %f\n", i, + j, point[0], point[1], triangle_index, coefficients[0], + coefficients[1], coefficients[2]); } else { @@ -298,8 +299,10 @@ int main() ArborX::attach(ArborX::intersects(point), Attachment{triangle_index, coefficients}), TriangleIntersectionCallback{triangles}); - KOKKOS_IMPL_DO_NOT_USE_PRINTF("%d, %d: %f %f in %d\n", i, j, - point[0], point[1], triangle_index); + KOKKOS_IMPL_DO_NOT_USE_PRINTF( + "%d, %d: %f %f in %d, coefficients: %f, %f, %f\n", i, j, + point[0], point[1], triangle_index, coefficients[0], + coefficients[1], coefficients[2]); } } }); From c54db089e40d3094602a437af47fadede69c38ce Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Mon, 25 Sep 2023 16:32:53 -0400 Subject: [PATCH 14/23] Copy the dummy inout files into the correct directory --- benchmarks/tokamak/CMakeLists.txt | 3 +++ {build/benchmarks => benchmarks}/tokamak/RK4Steps.txt | 0 {build/benchmarks => benchmarks}/tokamak/RZGrid.stl | 0 3 files changed, 3 insertions(+) rename {build/benchmarks => benchmarks}/tokamak/RK4Steps.txt (100%) rename {build/benchmarks => benchmarks}/tokamak/RZGrid.stl (100%) diff --git a/benchmarks/tokamak/CMakeLists.txt b/benchmarks/tokamak/CMakeLists.txt index cef5f3390..50d3635c0 100644 --- a/benchmarks/tokamak/CMakeLists.txt +++ b/benchmarks/tokamak/CMakeLists.txt @@ -2,4 +2,7 @@ add_executable(ArborX_Benchmark_Tokamak.exe tokamak.cpp ) target_link_libraries(ArborX_Benchmark_Tokamak.exe ArborX::ArborX) +file(COPY ${CMAKE_CURRENT_SOURCE_DIR}/RK4Steps.txt + ${CMAKE_CURRENT_SOURCE_DIR}/RZGrid.stl + DESTINATION ${CMAKE_CURRENT_BINARY_DIR}) add_test(NAME ArborX_Benchmark_Tokamak.exe COMMAND ./ArborX_Benchmark_Tokamak.exe) diff --git a/build/benchmarks/tokamak/RK4Steps.txt b/benchmarks/tokamak/RK4Steps.txt similarity index 100% rename from build/benchmarks/tokamak/RK4Steps.txt rename to benchmarks/tokamak/RK4Steps.txt diff --git a/build/benchmarks/tokamak/RZGrid.stl b/benchmarks/tokamak/RZGrid.stl similarity index 100% rename from build/benchmarks/tokamak/RZGrid.stl rename to benchmarks/tokamak/RZGrid.stl From a170c589ea8c31fee15970a0ca52844bb90d0ae2 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Tue, 26 Sep 2023 10:18:28 -0400 Subject: [PATCH 15/23] TriangleIntersectionCallback needs to be device-constructible --- benchmarks/tokamak/tokamak.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/benchmarks/tokamak/tokamak.cpp b/benchmarks/tokamak/tokamak.cpp index 5b80165ec..498ff3ebb 100644 --- a/benchmarks/tokamak/tokamak.cpp +++ b/benchmarks/tokamak/tokamak.cpp @@ -92,7 +92,7 @@ template class TriangleIntersectionCallback { public: - TriangleIntersectionCallback( + KOKKOS_FUNCTION TriangleIntersectionCallback( Kokkos::View triangles) : _triangles(triangles) {} From 051ff3cc86701b4c58eaf28c98fcba371f2fdb93 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Mon, 30 Oct 2023 13:29:02 -0400 Subject: [PATCH 16/23] Fix iostream --- benchmarks/tokamak/tokamak.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/benchmarks/tokamak/tokamak.cpp b/benchmarks/tokamak/tokamak.cpp index 498ff3ebb..2f9cae363 100644 --- a/benchmarks/tokamak/tokamak.cpp +++ b/benchmarks/tokamak/tokamak.cpp @@ -14,6 +14,7 @@ #include +#include #include using Point = ArborX::ExperimentalHyperGeometry::Point<2>; From a9d65b1098ae2c743b809303b6b188617b1b39b8 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Mon, 30 Oct 2023 14:24:57 -0400 Subject: [PATCH 17/23] Fix compiling with the new interface --- benchmarks/tokamak/tokamak.cpp | 12 +++++++----- 1 file changed, 7 insertions(+), 5 deletions(-) diff --git a/benchmarks/tokamak/tokamak.cpp b/benchmarks/tokamak/tokamak.cpp index 2f9cae363..e99a6180a 100644 --- a/benchmarks/tokamak/tokamak.cpp +++ b/benchmarks/tokamak/tokamak.cpp @@ -14,9 +14,10 @@ #include -#include #include +#include +using Box = ArborX::ExperimentalHyperGeometry::Box<2>; using Point = ArborX::ExperimentalHyperGeometry::Point<2>; using Triangle = ArborX::ExperimentalHyperGeometry::Triangle<2>; @@ -51,7 +52,7 @@ struct ArborX::AccessTraits, ArborX::PrimitivesTag> static KOKKOS_FUNCTION auto get(Triangles const &triangles, int i) { auto const &triangle = triangles(i); - ArborX::ExperimentalHyperGeometry::Box<2> box{}; + Box box{}; box += triangle.a; box += triangle.b; box += triangle.c; @@ -255,9 +256,10 @@ int main() std::cout << "Creating BVH tree.\n"; ArborX::BasicBoundingVolumeHierarchy< - MemorySpace, ArborX::Details::PairIndexVolume< - ArborX::ExperimentalHyperGeometry::Box<2>>> const - tree(execution_space, Triangles{triangles}); + MemorySpace, ArborX::Details::PairIndexVolume> const + tree( + execution_space, + ArborX::Details::LegacyValues{triangles}); std::cout << "BVH tree set up.\n"; std::cout << "Starting the queries.\n"; From 16cf2ba64da1775a5c36e1dd54605bf7f997cd5e Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Mon, 30 Oct 2023 15:14:59 -0400 Subject: [PATCH 18/23] Convert to Triangle leaf nodes --- benchmarks/tokamak/tokamak.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/benchmarks/tokamak/tokamak.cpp b/benchmarks/tokamak/tokamak.cpp index e99a6180a..d901b599e 100644 --- a/benchmarks/tokamak/tokamak.cpp +++ b/benchmarks/tokamak/tokamak.cpp @@ -256,10 +256,10 @@ int main() std::cout << "Creating BVH tree.\n"; ArborX::BasicBoundingVolumeHierarchy< - MemorySpace, ArborX::Details::PairIndexVolume> const - tree( - execution_space, - ArborX::Details::LegacyValues{triangles}); + MemorySpace, ArborX::Details::PairIndexVolume> const + tree(execution_space, + ArborX::Details::LegacyValues{ + triangles}); std::cout << "BVH tree set up.\n"; std::cout << "Starting the queries.\n"; From 61f55e22de213b237885ba7933131f19ce206636 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Mon, 30 Oct 2023 16:41:54 -0400 Subject: [PATCH 19/23] Remove Triangle struct --- benchmarks/tokamak/tokamak.cpp | 39 ---------------------------------- 1 file changed, 39 deletions(-) diff --git a/benchmarks/tokamak/tokamak.cpp b/benchmarks/tokamak/tokamak.cpp index d901b599e..2be4fda0b 100644 --- a/benchmarks/tokamak/tokamak.cpp +++ b/benchmarks/tokamak/tokamak.cpp @@ -21,45 +21,6 @@ using Box = ArborX::ExperimentalHyperGeometry::Box<2>; using Point = ArborX::ExperimentalHyperGeometry::Point<2>; using Triangle = ArborX::ExperimentalHyperGeometry::Triangle<2>; -template -struct Triangles -{ - // Return the number of triangles. - KOKKOS_FUNCTION int size() const { return _triangles.size(); } - - // Return the triangle with index i. - KOKKOS_FUNCTION Triangle const &operator()(int i) const - { - return _triangles(i); - } - - Kokkos::View *, MemorySpace> - _triangles; -}; - -// For creating the bounding volume hierarchy given a Triangles object, we -// need to define the memory space, how to get the total number of objects, -// and how to access a specific box. Since there are corresponding functions in -// the Triangles class, we just resort to them. -template -struct ArborX::AccessTraits, ArborX::PrimitivesTag> -{ - using memory_space = typename DeviceType::memory_space; - static KOKKOS_FUNCTION int size(Triangles const &triangles) - { - return triangles.size(); - } - static KOKKOS_FUNCTION auto get(Triangles const &triangles, int i) - { - auto const &triangle = triangles(i); - Box box{}; - box += triangle.a; - box += triangle.b; - box += triangle.c; - return box; - } -}; - KOKKOS_FUNCTION ArborX::Point compute_barycentric_coordinates(Triangle const &triangle, Point const &point) From eb6731bf291c43f8c09f5429f76665ff3a4372ef Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Mon, 30 Oct 2023 16:53:56 -0400 Subject: [PATCH 20/23] Move tokamak benchmark to example --- benchmarks/CMakeLists.txt | 1 - examples/CMakeLists.txt | 1 + {benchmarks => examples}/tokamak/CMakeLists.txt | 0 {benchmarks => examples}/tokamak/RK4Steps.txt | 0 {benchmarks => examples}/tokamak/RZGrid.stl | 0 {benchmarks => examples}/tokamak/tokamak.cpp | 0 6 files changed, 1 insertion(+), 1 deletion(-) rename {benchmarks => examples}/tokamak/CMakeLists.txt (100%) rename {benchmarks => examples}/tokamak/RK4Steps.txt (100%) rename {benchmarks => examples}/tokamak/RZGrid.stl (100%) rename {benchmarks => examples}/tokamak/tokamak.cpp (100%) diff --git a/benchmarks/CMakeLists.txt b/benchmarks/CMakeLists.txt index 5973e8638..65fa7aa1a 100644 --- a/benchmarks/CMakeLists.txt +++ b/benchmarks/CMakeLists.txt @@ -6,7 +6,6 @@ message(STATUS "Found benchmark: ${benchmark_DIR} (version \"${benchmark_VERSION add_subdirectory(brute_force_vs_bvh) add_subdirectory(dbscan) add_subdirectory(execution_space_instances) -add_subdirectory(tokamak) if(NOT WIN32) # FIXME: for now, skip the benchmarks using Google benchmark # when building for Windows, as we have trouble linking it diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index a387fc0d6..ba16c3a1b 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -11,6 +11,7 @@ add_subdirectory(dbscan) add_subdirectory(molecular_dynamics) add_subdirectory(simple_intersection) +add_subdirectory(tokamak) add_subdirectory(triangle_intersection) find_package(Boost COMPONENTS program_options) diff --git a/benchmarks/tokamak/CMakeLists.txt b/examples/tokamak/CMakeLists.txt similarity index 100% rename from benchmarks/tokamak/CMakeLists.txt rename to examples/tokamak/CMakeLists.txt diff --git a/benchmarks/tokamak/RK4Steps.txt b/examples/tokamak/RK4Steps.txt similarity index 100% rename from benchmarks/tokamak/RK4Steps.txt rename to examples/tokamak/RK4Steps.txt diff --git a/benchmarks/tokamak/RZGrid.stl b/examples/tokamak/RZGrid.stl similarity index 100% rename from benchmarks/tokamak/RZGrid.stl rename to examples/tokamak/RZGrid.stl diff --git a/benchmarks/tokamak/tokamak.cpp b/examples/tokamak/tokamak.cpp similarity index 100% rename from benchmarks/tokamak/tokamak.cpp rename to examples/tokamak/tokamak.cpp From 1cf69adfa440b3f3d9ed2c380a052d292ac6ec74 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Mon, 30 Oct 2023 16:58:15 -0400 Subject: [PATCH 21/23] Clean up KOKKOS_IMPL_DO_NOT_USE_PRINTF --- examples/tokamak/tokamak.cpp | 21 +++++++++++++-------- 1 file changed, 13 insertions(+), 8 deletions(-) diff --git a/examples/tokamak/tokamak.cpp b/examples/tokamak/tokamak.cpp index 2be4fda0b..3708ad979 100644 --- a/examples/tokamak/tokamak.cpp +++ b/examples/tokamak/tokamak.cpp @@ -248,13 +248,19 @@ int main() compute_barycentric_coordinates(triangle, point); bool intersects = test_coeffs[0] >= 0 && test_coeffs[1] >= 0 && test_coeffs[2] >= 0; + +#if KOKKOS_VERSION >= 40200 + using Kokkos::printf; +#elif defined(__SYCL_DEVICE_ONLY__) + using sycl::ext::oneapi::experimental::printf; +#endif + if (intersects) { coefficients = test_coeffs; - KOKKOS_IMPL_DO_NOT_USE_PRINTF( - "%d, %d: %f %f in %d (same), coefficients: %f, %f, %f\n", i, - j, point[0], point[1], triangle_index, coefficients[0], - coefficients[1], coefficients[2]); + printf("%d, %d: %f %f in %d (same), coefficients: %f, %f, %f\n", + i, j, point[0], point[1], triangle_index, coefficients[0], + coefficients[1], coefficients[2]); } else { @@ -263,10 +269,9 @@ int main() ArborX::attach(ArborX::intersects(point), Attachment{triangle_index, coefficients}), TriangleIntersectionCallback{triangles}); - KOKKOS_IMPL_DO_NOT_USE_PRINTF( - "%d, %d: %f %f in %d, coefficients: %f, %f, %f\n", i, j, - point[0], point[1], triangle_index, coefficients[0], - coefficients[1], coefficients[2]); + printf("%d, %d: %f %f in %d, coefficients: %f, %f, %f\n", i, j, + point[0], point[1], triangle_index, coefficients[0], + coefficients[1], coefficients[2]); } } }); From 164fafa35dd116628c78422965a6a79c47b2ed1e Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Mon, 30 Oct 2023 16:58:57 -0400 Subject: [PATCH 22/23] Don't check for intersecction in callback --- examples/tokamak/tokamak.cpp | 12 +++--------- 1 file changed, 3 insertions(+), 9 deletions(-) diff --git a/examples/tokamak/tokamak.cpp b/examples/tokamak/tokamak.cpp index 3708ad979..037d8d9a7 100644 --- a/examples/tokamak/tokamak.cpp +++ b/examples/tokamak/tokamak.cpp @@ -71,15 +71,9 @@ class TriangleIntersectionCallback ArborX::Point coeffs = compute_barycentric_coordinates(triangle, point); - bool intersects = coeffs[0] >= 0 && coeffs[1] >= 0 && coeffs[2] >= 0; - - if (intersects) - { - attachment.triangle_index = triangle_index; - attachment.coeffs = coeffs; - return ArborX::CallbackTreeTraversalControl::early_exit; - } - return ArborX::CallbackTreeTraversalControl::normal_continuation; + attachment.triangle_index = triangle_index; + attachment.coeffs = coeffs; + return ArborX::CallbackTreeTraversalControl::early_exit; } private: From 03348f1fd6f5c2b13950b080ce8d1af4ab3ce5e8 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Mon, 30 Oct 2023 17:02:10 -0400 Subject: [PATCH 23/23] Rename executable --- examples/tokamak/CMakeLists.txt | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/examples/tokamak/CMakeLists.txt b/examples/tokamak/CMakeLists.txt index 50d3635c0..53dea9a5a 100644 --- a/examples/tokamak/CMakeLists.txt +++ b/examples/tokamak/CMakeLists.txt @@ -1,8 +1,8 @@ -add_executable(ArborX_Benchmark_Tokamak.exe +add_executable(ArborX_Example_Tokamak.exe tokamak.cpp ) -target_link_libraries(ArborX_Benchmark_Tokamak.exe ArborX::ArborX) +target_link_libraries(ArborX_Example_Tokamak.exe ArborX::ArborX) file(COPY ${CMAKE_CURRENT_SOURCE_DIR}/RK4Steps.txt ${CMAKE_CURRENT_SOURCE_DIR}/RZGrid.stl DESTINATION ${CMAKE_CURRENT_BINARY_DIR}) -add_test(NAME ArborX_Benchmark_Tokamak.exe COMMAND ./ArborX_Benchmark_Tokamak.exe) +add_test(NAME ArborX_Example_Tokamak.exe COMMAND ./ArborX_Example_Tokamak.exe)