From 2ce4198ef105ef82fbf37dfd7453c26782b1e175 Mon Sep 17 00:00:00 2001 From: Joana Niermann Date: Fri, 7 Oct 2022 14:12:20 +0200 Subject: [PATCH] Grid implementation that addresses some issues found in the local navigation (grid navigation must be possible without reference surface). This PR provides: - 3D grids - Grid local coordinates for a global to local transform - Memory serialization: standalone grids and grids in grid collections - Implement populators, serializers and unittests, except for irregular attach populator. Due to the 'operator==' not being a __device__ function, the complete and attach populator unittests cannot pass the CI. The PR also splits the geometric shape information (cylinder, rectangle etc) from the navigation linking and boundary values, so that they can be reused for the grid geometry definition. With that comes the stringent use of local coordinates in the intersections and some refactoring (renaming etc). --- core/CMakeLists.txt | 32 +- .../detray/coordinates/cylindrical2.hpp | 4 +- core/include/detray/core/type_registry.hpp | 3 +- .../include/detray/definitions/containers.hpp | 100 +++- .../detray/definitions/detail/accessor.hpp | 41 ++ core/include/detray/definitions/indexing.hpp | 28 +- core/include/detray/grids/axis.hpp | 16 +- core/include/detray/grids/grid2.hpp | 10 +- core/include/detray/grids/populator.hpp | 8 +- .../concentric_cylinder_intersector.hpp | 11 +- .../intersection/cylinder_intersector.hpp | 19 +- .../detray/intersection/plane_intersector.hpp | 8 +- core/include/detray/masks/annulus2.hpp | 171 ------- core/include/detray/masks/annulus2D.hpp | 168 +++++++ core/include/detray/masks/cuboid3D.hpp | 108 +++++ core/include/detray/masks/cylinder2D.hpp | 111 +++++ core/include/detray/masks/cylinder3.hpp | 113 ----- core/include/detray/masks/cylinder3D.hpp | 106 +++++ core/include/detray/masks/line.hpp | 184 ++++---- core/include/detray/masks/mask_base.hpp | 120 ----- core/include/detray/masks/masks.hpp | 166 ++++++- core/include/detray/masks/rectangle2.hpp | 104 ----- core/include/detray/masks/rectangle2D.hpp | 99 ++++ core/include/detray/masks/ring2.hpp | 106 ----- core/include/detray/masks/ring2D.hpp | 101 +++++ core/include/detray/masks/single3.hpp | 93 ---- core/include/detray/masks/single3D.hpp | 97 ++++ core/include/detray/masks/trapezoid2.hpp | 113 ----- core/include/detray/masks/trapezoid2D.hpp | 109 +++++ core/include/detray/masks/unmasked.hpp | 94 ++-- .../detray/surface_finders/grid/axis.hpp | 339 ++++++++++++++ .../grid/detail/axis_binning.hpp | 298 ++++++++++++ .../grid/detail/axis_helpers.hpp | 169 +++++++ .../grid/detail/axis_shape.hpp | 272 +++++++++++ .../grid/detail/grid_helpers.hpp | 99 ++++ .../grid/detail/populator_impl.hpp | 342 ++++++++++++++ .../grid/detail/serializer_impl.hpp | 134 ++++++ .../detray/surface_finders/grid/grid.hpp | 303 +++++++++++++ .../surface_finders/grid/grid_builder.hpp | 130 ++++++ .../detray/surface_finders/grid/populator.hpp | 68 +++ .../surface_finders/grid/serializer.hpp | 11 + core/include/detray/tools/generators.hpp | 49 +- core/include/detray/utils/enumerate.hpp | 112 ++++- core/include/detray/utils/invalid_values.hpp | 41 +- core/include/detray/utils/type_traits.hpp | 4 +- .../tests/common/benchmark_find_volume.inl | 3 - .../tests/common/benchmark_intersect_all.inl | 3 - .../common/benchmark_intersect_surfaces.inl | 22 +- .../include/tests/common/benchmark_masks.inl | 226 ++++----- .../include/tests/common/coordinates.inl | 2 +- .../include/tests/common/core_detector.inl | 7 +- .../include/tests/common/core_mask_store.inl | 34 +- .../common/include/tests/common/grid_axis.inl | 420 +++++++++++++++++ .../common/include/tests/common/grid_grid.inl | 427 ++++++++++++++++++ .../include/tests/common/grid_populator.inl | 178 ++++++++ .../include/tests/common/grid_serializer.inl | 108 +++++ .../include/tests/common/masks_annulus2.inl | 82 ---- .../include/tests/common/masks_annulus2D.inl | 65 +++ .../include/tests/common/masks_cylinder.inl | 67 +++ .../include/tests/common/masks_cylinder3.inl | 52 --- .../include/tests/common/masks_line.inl | 56 ++- .../include/tests/common/masks_rectangle2.inl | 43 -- .../tests/common/masks_rectangle2D.inl | 38 ++ .../include/tests/common/masks_ring2.inl | 51 --- .../include/tests/common/masks_ring2D.inl | 35 ++ .../include/tests/common/masks_single3.inl | 93 ---- .../include/tests/common/masks_single3D.inl | 76 ++++ .../include/tests/common/masks_trapezoid2.inl | 45 -- .../tests/common/masks_trapezoid2D.inl | 40 ++ .../include/tests/common/masks_unmasked.inl | 21 +- .../common/include/tests/common/materials.inl | 6 +- .../common/include/tests/common/test_core.inl | 6 +- .../common/include/tests/common/test_defs.hpp | 6 - .../tests/common/test_toy_geometry.inl | 46 +- .../tools/create_telescope_detector.hpp | 2 +- .../common/tools/create_toy_geometry.hpp | 15 +- .../tests/common/tools/detector_metadata.hpp | 32 +- .../helix_cylinder_intersector.hpp | 21 +- .../helix_intersection_kernel.hpp | 10 +- .../intersectors/helix_plane_intersector.hpp | 7 +- .../tests/common/tools/test_surfaces.hpp | 7 +- .../common/tools_cylinder_intersection.inl | 49 +- .../common/tools_intersection_kernel.inl | 30 +- .../tests/common/tools_line_intersection.inl | 16 +- .../common/tools_planar_intersection.inl | 25 +- .../include/tests/common/tools_stepper.inl | 1 - .../include/tests/common/tools_track.inl | 4 +- tests/unit_tests/array/CMakeLists.txt | 55 ++- ...rray_cylinder3.cpp => array_annulus2D.cpp} | 2 +- .../{array_single3.cpp => array_cylinder.cpp} | 2 +- tests/unit_tests/array/array_grid.cpp | 9 + tests/unit_tests/array/array_grid_axis.cpp | 9 + .../unit_tests/array/array_grid_populator.cpp | 9 + .../array/array_grid_serializer.cpp | 9 + ...y_rectangle2.cpp => array_rectangle2D.cpp} | 2 +- .../{array_ring2.cpp => array_ring2D.cpp} | 2 +- ...{array_annulus2.cpp => array_single3D.cpp} | 2 +- tests/unit_tests/array/array_trapezoid2.cpp | 9 - tests/unit_tests/array/array_trapezoid2D.cpp | 9 + tests/unit_tests/cuda/CMakeLists.txt | 2 + tests/unit_tests/cuda/detector_cuda_kernel.cu | 2 +- .../unit_tests/cuda/detector_cuda_kernel.hpp | 2 +- tests/unit_tests/cuda/mask_store_cuda.cpp | 32 +- .../unit_tests/cuda/mask_store_cuda_kernel.cu | 18 +- .../cuda/mask_store_cuda_kernel.hpp | 16 +- .../unit_tests/cuda/sf_finders_grid_cuda.cpp | 325 +++++++++++++ .../cuda/sf_finders_grid_cuda_kernel.cu | 206 +++++++++ .../cuda/sf_finders_grid_cuda_kernel.hpp | 115 +++++ tests/unit_tests/eigen/CMakeLists.txt | 55 ++- ...igen_cylinder3.cpp => eigen_annulus2D.cpp} | 2 +- .../{eigen_single3.cpp => eigen_cylinder.cpp} | 2 +- tests/unit_tests/eigen/eigen_grid.cpp | 9 + tests/unit_tests/eigen/eigen_grid_axis.cpp | 9 + .../unit_tests/eigen/eigen_grid_populator.cpp | 9 + .../eigen/eigen_grid_serializer.cpp | 9 + ...n_rectangle2.cpp => eigen_rectangle2D.cpp} | 2 +- .../{eigen_ring2.cpp => eigen_ring2D.cpp} | 2 +- ...{eigen_annulus2.cpp => eigen_single3D.cpp} | 2 +- tests/unit_tests/eigen/eigen_trapezoid2.cpp | 9 - tests/unit_tests/eigen/eigen_trapezoid2D.cpp | 9 + tests/unit_tests/smatrix/CMakeLists.txt | 56 ++- ...ix_cylinder3.cpp => smatrix_annulus2D.cpp} | 2 +- ...trix_annulus2.cpp => smatrix_cylinder.cpp} | 2 +- tests/unit_tests/smatrix/smatrix_grid.cpp | 9 + .../unit_tests/smatrix/smatrix_grid_axis.cpp | 9 + .../smatrix/smatrix_grid_populator.cpp | 9 + .../smatrix/smatrix_grid_serializer.cpp | 9 + ...rectangle2.cpp => smatrix_rectangle2D.cpp} | 2 +- .../{smatrix_ring2.cpp => smatrix_ring2D.cpp} | 2 +- ...atrix_single3.cpp => smatrix_single3D.cpp} | 2 +- .../unit_tests/smatrix/smatrix_trapezoid2.cpp | 9 - .../smatrix/smatrix_trapezoid2D.cpp | 9 + tests/unit_tests/vc/CMakeLists.txt | 55 ++- ...y_cylinder3.cpp => vc_array_annulus2D.cpp} | 2 +- ...rray_single3.cpp => vc_array_cylinder.cpp} | 2 +- tests/unit_tests/vc/vc_array_grid.cpp | 9 + tests/unit_tests/vc/vc_array_grid_axis.cpp | 9 + .../unit_tests/vc/vc_array_grid_populator.cpp | 9 + .../vc/vc_array_grid_serializer.cpp | 9 + ...ectangle2.cpp => vc_array_rectangle2D.cpp} | 2 +- ...vc_array_ring2.cpp => vc_array_ring2D.cpp} | 2 +- ...ray_annulus2.cpp => vc_array_single3D.cpp} | 2 +- tests/unit_tests/vc/vc_array_trapezoid2.cpp | 9 - tests/unit_tests/vc/vc_array_trapezoid2D.cpp | 9 + 144 files changed, 6555 insertions(+), 1931 deletions(-) delete mode 100644 core/include/detray/masks/annulus2.hpp create mode 100644 core/include/detray/masks/annulus2D.hpp create mode 100644 core/include/detray/masks/cuboid3D.hpp create mode 100644 core/include/detray/masks/cylinder2D.hpp delete mode 100644 core/include/detray/masks/cylinder3.hpp create mode 100644 core/include/detray/masks/cylinder3D.hpp delete mode 100644 core/include/detray/masks/mask_base.hpp delete mode 100644 core/include/detray/masks/rectangle2.hpp create mode 100644 core/include/detray/masks/rectangle2D.hpp delete mode 100644 core/include/detray/masks/ring2.hpp create mode 100644 core/include/detray/masks/ring2D.hpp delete mode 100644 core/include/detray/masks/single3.hpp create mode 100644 core/include/detray/masks/single3D.hpp delete mode 100644 core/include/detray/masks/trapezoid2.hpp create mode 100644 core/include/detray/masks/trapezoid2D.hpp create mode 100644 core/include/detray/surface_finders/grid/axis.hpp create mode 100644 core/include/detray/surface_finders/grid/detail/axis_binning.hpp create mode 100644 core/include/detray/surface_finders/grid/detail/axis_helpers.hpp create mode 100644 core/include/detray/surface_finders/grid/detail/axis_shape.hpp create mode 100644 core/include/detray/surface_finders/grid/detail/grid_helpers.hpp create mode 100644 core/include/detray/surface_finders/grid/detail/populator_impl.hpp create mode 100644 core/include/detray/surface_finders/grid/detail/serializer_impl.hpp create mode 100644 core/include/detray/surface_finders/grid/grid.hpp create mode 100644 core/include/detray/surface_finders/grid/grid_builder.hpp create mode 100644 core/include/detray/surface_finders/grid/populator.hpp create mode 100644 core/include/detray/surface_finders/grid/serializer.hpp create mode 100644 tests/common/include/tests/common/grid_axis.inl create mode 100644 tests/common/include/tests/common/grid_grid.inl create mode 100644 tests/common/include/tests/common/grid_populator.inl create mode 100644 tests/common/include/tests/common/grid_serializer.inl delete mode 100644 tests/common/include/tests/common/masks_annulus2.inl create mode 100644 tests/common/include/tests/common/masks_annulus2D.inl create mode 100644 tests/common/include/tests/common/masks_cylinder.inl delete mode 100644 tests/common/include/tests/common/masks_cylinder3.inl delete mode 100644 tests/common/include/tests/common/masks_rectangle2.inl create mode 100644 tests/common/include/tests/common/masks_rectangle2D.inl delete mode 100644 tests/common/include/tests/common/masks_ring2.inl create mode 100644 tests/common/include/tests/common/masks_ring2D.inl delete mode 100644 tests/common/include/tests/common/masks_single3.inl create mode 100644 tests/common/include/tests/common/masks_single3D.inl delete mode 100644 tests/common/include/tests/common/masks_trapezoid2.inl create mode 100644 tests/common/include/tests/common/masks_trapezoid2D.inl rename tests/unit_tests/array/{array_cylinder3.cpp => array_annulus2D.cpp} (82%) rename tests/unit_tests/array/{array_single3.cpp => array_cylinder.cpp} (83%) create mode 100644 tests/unit_tests/array/array_grid.cpp create mode 100644 tests/unit_tests/array/array_grid_axis.cpp create mode 100644 tests/unit_tests/array/array_grid_populator.cpp create mode 100644 tests/unit_tests/array/array_grid_serializer.cpp rename tests/unit_tests/array/{array_rectangle2.cpp => array_rectangle2D.cpp} (82%) rename tests/unit_tests/array/{array_ring2.cpp => array_ring2D.cpp} (84%) rename tests/unit_tests/array/{array_annulus2.cpp => array_single3D.cpp} (83%) delete mode 100644 tests/unit_tests/array/array_trapezoid2.cpp create mode 100644 tests/unit_tests/array/array_trapezoid2D.cpp create mode 100644 tests/unit_tests/cuda/sf_finders_grid_cuda.cpp create mode 100644 tests/unit_tests/cuda/sf_finders_grid_cuda_kernel.cu create mode 100644 tests/unit_tests/cuda/sf_finders_grid_cuda_kernel.hpp rename tests/unit_tests/eigen/{eigen_cylinder3.cpp => eigen_annulus2D.cpp} (82%) rename tests/unit_tests/eigen/{eigen_single3.cpp => eigen_cylinder.cpp} (83%) create mode 100644 tests/unit_tests/eigen/eigen_grid.cpp create mode 100644 tests/unit_tests/eigen/eigen_grid_axis.cpp create mode 100644 tests/unit_tests/eigen/eigen_grid_populator.cpp create mode 100644 tests/unit_tests/eigen/eigen_grid_serializer.cpp rename tests/unit_tests/eigen/{eigen_rectangle2.cpp => eigen_rectangle2D.cpp} (82%) rename tests/unit_tests/eigen/{eigen_ring2.cpp => eigen_ring2D.cpp} (84%) rename tests/unit_tests/eigen/{eigen_annulus2.cpp => eigen_single3D.cpp} (83%) delete mode 100644 tests/unit_tests/eigen/eigen_trapezoid2.cpp create mode 100644 tests/unit_tests/eigen/eigen_trapezoid2D.cpp rename tests/unit_tests/smatrix/{smatrix_cylinder3.cpp => smatrix_annulus2D.cpp} (83%) rename tests/unit_tests/smatrix/{smatrix_annulus2.cpp => smatrix_cylinder.cpp} (83%) create mode 100644 tests/unit_tests/smatrix/smatrix_grid.cpp create mode 100644 tests/unit_tests/smatrix/smatrix_grid_axis.cpp create mode 100644 tests/unit_tests/smatrix/smatrix_grid_populator.cpp create mode 100644 tests/unit_tests/smatrix/smatrix_grid_serializer.cpp rename tests/unit_tests/smatrix/{smatrix_rectangle2.cpp => smatrix_rectangle2D.cpp} (82%) rename tests/unit_tests/smatrix/{smatrix_ring2.cpp => smatrix_ring2D.cpp} (84%) rename tests/unit_tests/smatrix/{smatrix_single3.cpp => smatrix_single3D.cpp} (83%) delete mode 100644 tests/unit_tests/smatrix/smatrix_trapezoid2.cpp create mode 100644 tests/unit_tests/smatrix/smatrix_trapezoid2D.cpp rename tests/unit_tests/vc/{vc_array_cylinder3.cpp => vc_array_annulus2D.cpp} (83%) rename tests/unit_tests/vc/{vc_array_single3.cpp => vc_array_cylinder.cpp} (83%) create mode 100644 tests/unit_tests/vc/vc_array_grid.cpp create mode 100644 tests/unit_tests/vc/vc_array_grid_axis.cpp create mode 100644 tests/unit_tests/vc/vc_array_grid_populator.cpp create mode 100644 tests/unit_tests/vc/vc_array_grid_serializer.cpp rename tests/unit_tests/vc/{vc_array_rectangle2.cpp => vc_array_rectangle2D.cpp} (82%) rename tests/unit_tests/vc/{vc_array_ring2.cpp => vc_array_ring2D.cpp} (84%) rename tests/unit_tests/vc/{vc_array_annulus2.cpp => vc_array_single3D.cpp} (83%) delete mode 100644 tests/unit_tests/vc/vc_array_trapezoid2.cpp create mode 100644 tests/unit_tests/vc/vc_array_trapezoid2D.cpp diff --git a/core/CMakeLists.txt b/core/CMakeLists.txt index ad0abc802..c4b6cd25c 100644 --- a/core/CMakeLists.txt +++ b/core/CMakeLists.txt @@ -56,21 +56,21 @@ detray_add_library( detray_core core "include/detray/intersection/line_intersector.hpp" "include/detray/intersection/plane_intersector.hpp" # masks include(s) - "include/detray/masks/annulus2.hpp" - "include/detray/masks/cylinder3.hpp" + "include/detray/masks/annulus2D.hpp" + "include/detray/masks/cuboid3D.hpp" + "include/detray/masks/cylinder3D.hpp" "include/detray/masks/line.hpp" "include/detray/masks/masks.hpp" - "include/detray/masks/masks_base.hpp" - "include/detray/masks/rectangle2.hpp" - "include/detray/masks/ring2.hpp" - "include/detray/masks/single3.hpp" - "include/detray/masks/trapezoid2.hpp" + "include/detray/masks/rectangle2D.hpp" + "include/detray/masks/ring2D.hpp" + "include/detray/masks/single3D.hpp" + "include/detray/masks/trapezoid2D.hpp" "include/detray/masks/unmasked.hpp" # materials include(s) - "include/detray/materials/mixture.hpp" + "include/detray/materials/mixture.hpp" "include/detray/materials/material.hpp" - "include/detray/materials/material_rod.hpp" - "include/detray/materials/material_slab.hpp" + "include/detray/materials/material_rod.hpp" + "include/detray/materials/material_slab.hpp" "include/detray/materials/predefined_materials.hpp" # propagator include(s) "include/detray/propagator/aborters.hpp" @@ -89,6 +89,18 @@ detray_add_library( detray_core core "include/detray/surface_finders/brute_force_finder.hpp" "include/detray/surface_finders/grid2_finder.hpp" "include/detray/surface_finders/neighborhood_kernel.hpp" + # surface finder include(s) + "include/detray/surface_finders/grid/detail/axis_binning.hpp" + "include/detray/surface_finders/grid/detail/axis_helpers.hpp" + "include/detray/surface_finders/grid/detail/axis_shape.hpp" + "include/detray/surface_finders/grid/detail/grid_helpers.hpp" + "include/detray/surface_finders/grid/detail/populator_impl.hpp" + "include/detray/surface_finders/grid/detail/serializer_impl.hpp" + "include/detray/surface_finders/grid/axis.hpp" + "include/detray/surface_finders/grid/grid_builder.hpp" + "include/detray/surface_finders/grid/grid.hpp" + "include/detray/surface_finders/grid/populator.hpp" + "include/detray/surface_finders/grid/serializer.hpp" # tools include(s) "include/detray/tools/associator.hpp" "include/detray/tools/bin_association.hpp" diff --git a/core/include/detray/coordinates/cylindrical2.hpp b/core/include/detray/coordinates/cylindrical2.hpp index 6ed134aa9..f08f9fde4 100644 --- a/core/include/detray/coordinates/cylindrical2.hpp +++ b/core/include/detray/coordinates/cylindrical2.hpp @@ -16,7 +16,7 @@ namespace detray { -/** Local frame projection into a polar coordinate frame +/** Local frame projection into a 2D cylindrical coordinate frame */ template struct cylindrical2 : public coordinate_base { @@ -60,7 +60,7 @@ struct cylindrical2 : public coordinate_base { DETRAY_HOST_DEVICE inline point3 local_to_global( const transform3_t &trf, const mask_t &mask, const point2 &p, const vector3 & /*d*/) const { - const scalar_type r = mask.radius(); + const scalar_type r = mask[0]; const scalar_type phi = p[0] / r; const scalar_type x = r * std::cos(phi); const scalar_type y = r * std::sin(phi); diff --git a/core/include/detray/core/type_registry.hpp b/core/include/detray/core/type_registry.hpp index a20f8c72f..9f0fac72f 100644 --- a/core/include/detray/core/type_registry.hpp +++ b/core/include/detray/core/type_registry.hpp @@ -79,6 +79,7 @@ class registry_base { template struct get_index { static constexpr ID value = get_id(); + DETRAY_HOST_DEVICE constexpr bool operator()() noexcept { return is_valid(value); } }; @@ -96,7 +97,7 @@ class registry_base { is_valid(ref_idx), "Index out of range: Please make sure that indices and type " "enums match the number of types in container."); - return static_cast(index); + return static_cast(ref_idx); } if constexpr (ref_idx < sizeof...(registered_types) - 1) { return to_id(index); diff --git a/core/include/detray/definitions/containers.hpp b/core/include/detray/definitions/containers.hpp index 1618ca758..57911f5db 100644 --- a/core/include/detray/definitions/containers.hpp +++ b/core/include/detray/definitions/containers.hpp @@ -1,20 +1,32 @@ /** Detray library, part of the ACTS project (R&D line) * - * (c) 2020 CERN for the benefit of the ACTS project + * (c) 2020-2022 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ #pragma once +// Project include(s) +#include "detray/definitions/qualifiers.hpp" +#include "detray/utils/type_traits.hpp" + +// Vecmem include(s) +#include +#include +#include +#include + +// Thrust include(s) +#include + +// System include(s) #include #include #include -#include +#include #include -#include "vecmem/containers/jagged_vector.hpp" - namespace detray { template using darray = std::array; @@ -31,4 +43,84 @@ using dmap = std::map; template using dtuple = std::tuple; +/// @brief Bundle container type definitions +template