diff --git a/benchmarks/cpu/toy_detector_cpu.cpp b/benchmarks/cpu/toy_detector_cpu.cpp index e9415a3f24..b9203f4455 100644 --- a/benchmarks/cpu/toy_detector_cpu.cpp +++ b/benchmarks/cpu/toy_detector_cpu.cpp @@ -38,11 +38,8 @@ BENCHMARK_DEFINE_F(ToyDetectorBenchmark, CPU)(benchmark::State& state) { // VecMem copy object vecmem::copy copy; - // Type declarations - using host_detector_type = traccc::default_detector::host; - // Read back detector file - host_detector_type det{host_mr}; + traccc::host_detector det; traccc::io::read_detector( det, host_mr, sim_dir + "toy_detector_geometry.json", sim_dir + "toy_detector_homogeneous_material.json", diff --git a/benchmarks/cuda/toy_detector_cuda.cpp b/benchmarks/cuda/toy_detector_cuda.cpp index 4a047c563e..3f8ad3b808 100644 --- a/benchmarks/cuda/toy_detector_cuda.cpp +++ b/benchmarks/cuda/toy_detector_cuda.cpp @@ -49,7 +49,7 @@ BENCHMARK_DEFINE_F(ToyDetectorBenchmark, CUDA)(benchmark::State& state) { vecmem::cuda::async_copy async_copy{stream.cudaStream()}; // Read back detector file - traccc::default_detector::host det{cuda_host_mr}; + traccc::host_detector det; traccc::io::read_detector( det, cuda_host_mr, sim_dir + "toy_detector_geometry.json", sim_dir + "toy_detector_homogeneous_material.json", @@ -68,9 +68,8 @@ BENCHMARK_DEFINE_F(ToyDetectorBenchmark, CUDA)(benchmark::State& state) { async_copy, stream); // Copy detector to device - const auto det_buffer = detray::get_buffer(det, device_mr, copy); - // Detector view object - auto det_view = detray::get_data(det_buffer); + const traccc::detector_buffer det_buffer = + traccc::buffer_from_host_detector(det, device_mr, copy); for (auto _ : state) { @@ -120,13 +119,13 @@ BENCHMARK_DEFINE_F(ToyDetectorBenchmark, CUDA)(benchmark::State& state) { // Run CKF track finding traccc::edm::track_candidate_collection< traccc::default_algebra>::buffer track_candidates_cuda_buffer = - device_finding(det_view, field, measurements_cuda_buffer, + device_finding(det_buffer, field, measurements_cuda_buffer, params_cuda_buffer); // Run track fitting traccc::edm::track_fit_container::buffer track_states_cuda_buffer = device_fitting( - det_view, field, + det_buffer, field, {track_candidates_cuda_buffer, measurements_cuda_buffer}); // Create a temporary buffer that will receive the device memory. diff --git a/core/CMakeLists.txt b/core/CMakeLists.txt index 72093bbfec..ee78dcdbef 100644 --- a/core/CMakeLists.txt +++ b/core/CMakeLists.txt @@ -86,8 +86,6 @@ traccc_add_library( traccc_core core TYPE SHARED "include/traccc/finding/details/combinatorial_kalman_filter.hpp" "include/traccc/finding/combinatorial_kalman_filter_algorithm.hpp" "src/finding/combinatorial_kalman_filter_algorithm.cpp" - "src/finding/combinatorial_kalman_filter_algorithm_default_detector.cpp" - "src/finding/combinatorial_kalman_filter_algorithm_telescope_detector.cpp" # Fitting algorithmic code "include/traccc/fitting/kalman_filter/gain_matrix_updater.hpp" "include/traccc/fitting/kalman_filter/kalman_actor.hpp" @@ -99,8 +97,6 @@ traccc_add_library( traccc_core core TYPE SHARED "include/traccc/fitting/details/kalman_fitting.hpp" "include/traccc/fitting/kalman_fitting_algorithm.hpp" "src/fitting/kalman_fitting_algorithm.cpp" - "src/fitting/kalman_fitting_algorithm_default_detector.cpp" - "src/fitting/kalman_fitting_algorithm_telescope_detector.cpp" # Seed finding algorithmic code. "include/traccc/seeding/detail/lin_circle.hpp" "include/traccc/seeding/detail/doublet.hpp" @@ -130,8 +126,6 @@ traccc_add_library( traccc_core core TYPE SHARED "src/seeding/silicon_pixel_spacepoint_formation.hpp" "include/traccc/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" "src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp" - "src/seeding/silicon_pixel_spacepoint_formation_algorithm_defdet.cpp" - "src/seeding/silicon_pixel_spacepoint_formation_algorithm_teldet.cpp" # Ambiguity resolution "include/traccc/ambiguity_resolution/ambiguity_resolution_config.hpp" "include/traccc/ambiguity_resolution/greedy_ambiguity_resolution_algorithm.hpp" diff --git a/core/include/traccc/finding/combinatorial_kalman_filter_algorithm.hpp b/core/include/traccc/finding/combinatorial_kalman_filter_algorithm.hpp index 2e7d37bc40..b08cb608e7 100644 --- a/core/include/traccc/finding/combinatorial_kalman_filter_algorithm.hpp +++ b/core/include/traccc/finding/combinatorial_kalman_filter_algorithm.hpp @@ -14,6 +14,7 @@ #include "traccc/edm/track_parameters.hpp" #include "traccc/finding/finding_config.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" #include "traccc/utils/algorithm.hpp" #include "traccc/utils/messaging.hpp" @@ -32,11 +33,7 @@ namespace traccc::host { /// class combinatorial_kalman_filter_algorithm : public algorithm::host( - const default_detector::host&, const magnetic_field&, - const measurement_collection_types::const_view&, - const bound_track_parameters_collection_types::const_view&)>, - public algorithm::host( - const telescope_detector::host&, const magnetic_field&, + const host_detector&, const magnetic_field&, const measurement_collection_types::const_view&, const bound_track_parameters_collection_types::const_view&)>, public messaging { @@ -54,23 +51,7 @@ class combinatorial_kalman_filter_algorithm /// Execute the algorithm /// - /// @param det The (default) detector object - /// @param bfield The magnetic field object - /// @param measurements All measurements in an event - /// @param seeds All seeds in an event to start the track finding - /// with - /// - /// @return A container of the found track candidates - /// - output_type operator()( - const default_detector::host& det, const magnetic_field& bfield, - const measurement_collection_types::const_view& measurements, - const bound_track_parameters_collection_types::const_view& seeds) - const override; - - /// Execute the algorithm - /// - /// @param det The (telescope) detector object + /// @param det The detector object /// @param bfield The magnetic field object /// @param measurements All measurements in an event /// @param seeds All seeds in an event to start the track finding @@ -79,7 +60,7 @@ class combinatorial_kalman_filter_algorithm /// @return A container of the found track candidates /// output_type operator()( - const telescope_detector::host& det, const magnetic_field& bfield, + const host_detector& det, const magnetic_field& bfield, const measurement_collection_types::const_view& measurements, const bound_track_parameters_collection_types::const_view& seeds) const override; diff --git a/core/include/traccc/fitting/kalman_fitting_algorithm.hpp b/core/include/traccc/fitting/kalman_fitting_algorithm.hpp index 11c1621f65..9e7ef8ff63 100644 --- a/core/include/traccc/fitting/kalman_fitting_algorithm.hpp +++ b/core/include/traccc/fitting/kalman_fitting_algorithm.hpp @@ -14,6 +14,7 @@ #include "traccc/edm/track_fit_container.hpp" #include "traccc/fitting/fitting_config.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" #include "traccc/utils/algorithm.hpp" #include "traccc/utils/messaging.hpp" @@ -32,10 +33,7 @@ namespace traccc::host { /// Kalman filter based track fitting algorithm class kalman_fitting_algorithm : public algorithm::host( - const default_detector::host&, const magnetic_field&, - const edm::track_candidate_container::const_view&)>, - public algorithm::host( - const telescope_detector::host&, const magnetic_field&, + const host_detector&, const magnetic_field&, const edm::track_candidate_container::const_view&)>, public messaging { @@ -56,27 +54,14 @@ class kalman_fitting_algorithm /// Execute the algorithm /// - /// @param det The (default) detector object - /// @param bfield The magnetic field object - /// @param track_candidates All track candidates to fit - /// - /// @return A container of the fitted track states - /// - output_type operator()( - const default_detector::host& det, const magnetic_field& bfield, - const edm::track_candidate_container::const_view& - track_candidates) const override; - - /// Execute the algorithm - /// - /// @param det The (telescope) detector object + /// @param det The detector object /// @param bfield The magnetic field object /// @param track_candidates All track candidates to fit /// /// @return A container of the fitted track states /// output_type operator()( - const telescope_detector::host& det, const magnetic_field& bfield, + const host_detector& det, const magnetic_field& bfield, const edm::track_candidate_container::const_view& track_candidates) const override; diff --git a/core/include/traccc/geometry/detector.hpp b/core/include/traccc/geometry/detector.hpp index dc0b3de655..8464246d8c 100644 --- a/core/include/traccc/geometry/detector.hpp +++ b/core/include/traccc/geometry/detector.hpp @@ -11,6 +11,7 @@ #include "traccc/definitions/primitives.hpp" // Detray include(s). +#include #include #include #include @@ -47,7 +48,7 @@ struct device_detector_container_types { /// Base struct for the different detector types supported by the project. template -struct detector { +struct detector_traits { /// Metadata type of the detector. using metadata_type = metadata_t; @@ -66,15 +67,26 @@ struct detector { }; // struct default_detector +template +concept is_detector_traits = requires { + typename T::metadata_type; + typename T::host; + typename T::device; + typename T::view; + typename T::buffer; +}; + /// Default detector (also used for ODD) using default_detector = - detector>; + detector_traits>; /// Telescope detector -using telescope_detector = detector< +using telescope_detector = detector_traits< detray::telescope_metadata>; /// Toy detector -using toy_detector = detector>; +using toy_detector = + detector_traits>; +using detector_type_list = std::tuple; } // namespace traccc diff --git a/core/include/traccc/geometry/detector_buffer.hpp b/core/include/traccc/geometry/detector_buffer.hpp new file mode 100644 index 0000000000..60769e6de1 --- /dev/null +++ b/core/include/traccc/geometry/detector_buffer.hpp @@ -0,0 +1,117 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2024 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s). +#include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" +#include "traccc/geometry/move_only_any.hpp" + +// Detray include(s). +#include +#include +#include +#include +#include + +namespace traccc { + +class detector_buffer { + public: + detector_buffer() = default; + detector_buffer(const detector_buffer&) = delete; + detector_buffer(detector_buffer&&) = default; + detector_buffer& operator=(const detector_buffer&) = delete; + detector_buffer& operator=(detector_buffer&&) = default; + + template + void set(typename detector_traits_t::buffer&& obj) + requires(is_detector_traits) + { + m_obj.set(std::move(obj)); + } + + template + bool is() const + requires(is_detector_traits) + { + return (type() == typeid(typename detector_traits_t::buffer)); + } + + const std::type_info& type() const { return m_obj.type(); } + + template + const typename detector_traits_t::buffer& as() const + requires(is_detector_traits) + { + return m_obj.as(); + } + + template + typename detector_traits_t::view as_view() const + requires(is_detector_traits) + { + return detray::get_data(as()); + } + + private: + move_only_any m_obj; +}; // class bfield + +/// @brief Helper function for `detector_buffer_visitor` +template +auto detector_buffer_visitor_helper(const detector_buffer& detector_buffer, + callable_t&& callable, + std::tuple*) { + if (detector_buffer.is()) { + return callable.template operator()( + detector_buffer.as_view()); + } else { + if constexpr (sizeof...(detector_ts) > 0) { + return detector_buffer_visitor_helper( + detector_buffer, std::forward(callable), + static_cast*>(nullptr)); + } else { + std::stringstream exception_message; + + exception_message + << "Invalid detector type (" << detector_buffer.type().name() + << ") received, but this type is not supported" << std::endl; + + throw std::invalid_argument(exception_message.str()); + } + } +} + +/// @brief Visitor for polymorphic detector buffer types +/// +/// This function takes a list of supported detector trait types and checks +/// if the provided field is one of them. If it is, it will call the provided +/// callable on a view of it and otherwise it will throw an exception. +template +auto detector_buffer_visitor(const detector_buffer& detector_buffer, + callable_t&& callable) { + return detector_buffer_visitor_helper( + detector_buffer, std::forward(callable), + static_cast(nullptr)); +} + +// TODO: Docs +inline detector_buffer buffer_from_host_detector(const host_detector& det, + vecmem::memory_resource& mr, + vecmem::copy& copy) { + return host_detector_visitor( + det, [&mr, ©]( + const typename detector_traits_t::host& detector) { + traccc::detector_buffer rv; + rv.set(detray::get_buffer(detector, mr, copy)); + return rv; + }); +} + +} // namespace traccc diff --git a/core/include/traccc/geometry/host_detector.hpp b/core/include/traccc/geometry/host_detector.hpp new file mode 100644 index 0000000000..e381f1b507 --- /dev/null +++ b/core/include/traccc/geometry/host_detector.hpp @@ -0,0 +1,92 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2024 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s). +#include "traccc/geometry/detector.hpp" +#include "traccc/geometry/move_only_any.hpp" + +// Detray include(s). +#include +#include +#include +#include +#include + +namespace traccc { + +/// Typeless, owning, host detector object +class host_detector { + public: + host_detector() = default; + + template + void set(typename detector_traits_t::host&& obj) + requires(is_detector_traits) + { + m_obj.set(std::move(obj)); + } + + template + bool is() const + requires(is_detector_traits) + { + return (type() == typeid(typename detector_traits_t::host)); + } + + const std::type_info& type() const { return m_obj.type(); } + + template + const typename detector_traits_t::host& as() const + requires(is_detector_traits) + { + return m_obj.as(); + } + + private: + move_only_any m_obj; +}; + +/// @brief Helper function for `host_detector_visitor` +template +auto host_detector_visitor_helper(const host_detector& host_detector, + callable_t&& callable, + std::tuple*) { + if (host_detector.is()) { + return callable.template operator()( + host_detector.as()); + } else { + if constexpr (sizeof...(detector_ts) > 0) { + return host_detector_visitor_helper( + host_detector, std::forward(callable), + static_cast*>(nullptr)); + } else { + std::stringstream exception_message; + + exception_message + << "Invalid detector type (" << host_detector.type().name() + << ") received, but this type is not supported" << std::endl; + + throw std::invalid_argument(exception_message.str()); + } + } +} + +/// @brief Visitor for polymorphic host detector types +/// +/// This function takes a list of supported detector trait types and checks +/// if the provided field is one of them. If it is, it will call the provided +/// callable on a view of it and otherwise it will throw an exception. +template +auto host_detector_visitor(const host_detector& host_detector, + callable_t&& callable) { + return host_detector_visitor_helper(host_detector, + std::forward(callable), + static_cast(nullptr)); +} +} // namespace traccc diff --git a/core/include/traccc/geometry/move_only_any.hpp b/core/include/traccc/geometry/move_only_any.hpp new file mode 100644 index 0000000000..9e21b9a85d --- /dev/null +++ b/core/include/traccc/geometry/move_only_any.hpp @@ -0,0 +1,126 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2024 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +#include + +#include "traccc/definitions/primitives.hpp" + +namespace traccc { + +class move_only_any { + public: + move_only_any() = default; + move_only_any(const move_only_any &) = delete; + move_only_any &operator=(const move_only_any &) = delete; + + template + explicit move_only_any(obj_t &&obj) + requires(!std::same_as, move_only_any>) + : m_obj(std::malloc(sizeof(obj_t))), + m_type(&typeid(obj_t)), + m_destructor(get_destructor()) { + new (m_obj) obj_t(std::forward(obj)); + } + + move_only_any(move_only_any &&other) noexcept + : m_obj(other.m_obj), + m_type(other.m_type), + m_destructor(other.m_destructor) { + other.m_obj = nullptr; + other.m_type = nullptr; + other.m_destructor = nullptr; + } + + move_only_any &operator=(move_only_any &&other) noexcept { + if (m_obj != nullptr) { + assert(m_destructor != nullptr); + m_destructor(m_obj); + std::free(m_obj); + } + + m_obj = other.m_obj; + other.m_obj = nullptr; + m_type = other.m_type; + other.m_type = nullptr; + m_destructor = other.m_destructor; + other.m_destructor = nullptr; + + return *this; + } + + ~move_only_any() { + if (m_obj != nullptr) { + assert(m_destructor != nullptr); + m_destructor(m_obj); + std::free(m_obj); + } + } + + template + void set(obj_t &&obj) + requires(!std::same_as, move_only_any>) + { + if (m_obj != nullptr) { + assert(m_destructor != nullptr); + m_destructor(m_obj); + std::free(m_obj); + } + + m_obj = std::malloc(sizeof(obj_t)); + new (m_obj) obj_t(std::forward(obj)); + + m_type = &typeid(obj_t); + m_destructor = get_destructor(); + } + + template + bool is() const { + if (m_type == nullptr) { + return false; + } else { + return (*m_type == typeid(obj_t)); + } + } + + bool has_value() const { return m_type != nullptr; } + + const std::type_info &type() const { + if (!has_value()) { + throw std::logic_error( + "Type ID for `traccc::move_only_any` requested, but no value " + "exists."); + } + + return *m_type; + } + + template + obj_t &as() const { + if (!has_value()) { + throw std::logic_error( + "Value for `traccc::move_only_any` requested, but no value " + "exists."); + } + + return *static_cast(m_obj); + } + + private: + template + void (*get_destructor() const)(void *) { + return static_cast( + [](void *ptr) { static_cast(ptr)->~obj_t(); }); + } + + void *m_obj = nullptr; + const std::type_info *m_type = nullptr; + void (*m_destructor)(void *) = nullptr; +}; + +} // namespace traccc diff --git a/core/include/traccc/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp b/core/include/traccc/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp index 2050d6ff04..fa1b14cea8 100644 --- a/core/include/traccc/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp +++ b/core/include/traccc/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp @@ -11,6 +11,7 @@ #include "traccc/edm/measurement.hpp" #include "traccc/edm/spacepoint_collection.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" #include "traccc/utils/algorithm.hpp" #include "traccc/utils/messaging.hpp" @@ -29,10 +30,7 @@ namespace traccc::host { /// class silicon_pixel_spacepoint_formation_algorithm : public algorithm, - public algorithm, public messaging { @@ -56,18 +54,7 @@ class silicon_pixel_spacepoint_formation_algorithm /// silicon pixel measurement /// output_type operator()( - const default_detector::host& det, - const measurement_collection_types::const_view&) const override; - - /// Construct spacepoints from 2D silicon pixel measurements - /// - /// @param det Detector object - /// @param measurements A collection of measurements - /// @return A spacepoint container, with one spacepoint for every - /// silicon pixel measurement - /// - output_type operator()( - const telescope_detector::host& det, + const host_detector& det, const measurement_collection_types::const_view&) const override; private: diff --git a/core/include/traccc/utils/detector_buffer_bfield_visitor.hpp b/core/include/traccc/utils/detector_buffer_bfield_visitor.hpp new file mode 100644 index 0000000000..21edcdd0cd --- /dev/null +++ b/core/include/traccc/utils/detector_buffer_bfield_visitor.hpp @@ -0,0 +1,32 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +#include "traccc/bfield/magnetic_field.hpp" +#include "traccc/geometry/detector_buffer.hpp" + +namespace traccc { + +template +auto detector_buffer_magnetic_field_visitor( + const detector_buffer& detector_buffer, const magnetic_field& bfield, + callable_t&& callable) { + return magnetic_field_visitor( + bfield, [&detector_buffer, &callable]( + const bfield_t& concrete_bfield) { + return detector_buffer_visitor( + detector_buffer, + [&concrete_bfield, &callable]( + const detector_t::view& concrete_detector_view) { + return callable.template operator()( + concrete_detector_view, concrete_bfield); + }); + }); +} + +} // namespace traccc diff --git a/core/include/traccc/utils/host_detector_bfield_visitor.hpp b/core/include/traccc/utils/host_detector_bfield_visitor.hpp new file mode 100644 index 0000000000..2b8d6bc790 --- /dev/null +++ b/core/include/traccc/utils/host_detector_bfield_visitor.hpp @@ -0,0 +1,32 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +#include "traccc/bfield/magnetic_field.hpp" +#include "traccc/geometry/detector_buffer.hpp" + +namespace traccc { + +template +auto host_detector_magnetic_field_visitor(const host_detector& host_detector, + const magnetic_field& bfield, + callable_t&& callable) { + return magnetic_field_visitor( + bfield, [&host_detector, &callable]( + const bfield_t& concrete_bfield) { + return host_detector_visitor( + host_detector, + [&concrete_bfield, &callable]( + const detector_t::host& concrete_detector_host) { + return callable.template operator()( + concrete_detector_host, concrete_bfield); + }); + }); +} + +} // namespace traccc diff --git a/core/src/finding/combinatorial_kalman_filter_algorithm.cpp b/core/src/finding/combinatorial_kalman_filter_algorithm.cpp index 6b82dff0d7..8990a30c24 100644 --- a/core/src/finding/combinatorial_kalman_filter_algorithm.cpp +++ b/core/src/finding/combinatorial_kalman_filter_algorithm.cpp @@ -8,6 +8,10 @@ // Local include(s). #include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" +#include "traccc/bfield/magnetic_field_types.hpp" +#include "traccc/finding/details/combinatorial_kalman_filter.hpp" +#include "traccc/utils/host_detector_bfield_visitor.hpp" + // System include(s). #include @@ -26,4 +30,23 @@ combinatorial_kalman_filter_algorithm::combinatorial_kalman_filter_algorithm( } } +combinatorial_kalman_filter_algorithm::output_type +combinatorial_kalman_filter_algorithm::operator()( + const host_detector& det, const magnetic_field& bfield, + const measurement_collection_types::const_view& measurements, + const bound_track_parameters_collection_types::const_view& seeds) const { + + // Perform the track finding using the appropriate templated implementation. + return host_detector_magnetic_field_visitor>( + det, bfield, + [&]( + const typename detector_t::host& detector, + const bfield_view_t field) { + return details::combinatorial_kalman_filter( + detector, field, measurements, seeds, m_config, m_mr.get(), + logger()); + }); +} + } // namespace traccc::host diff --git a/core/src/finding/combinatorial_kalman_filter_algorithm_default_detector.cpp b/core/src/finding/combinatorial_kalman_filter_algorithm_default_detector.cpp deleted file mode 100644 index 57a80292de..0000000000 --- a/core/src/finding/combinatorial_kalman_filter_algorithm_default_detector.cpp +++ /dev/null @@ -1,33 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2024-2025 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// Local include(s). -#include "traccc/bfield/magnetic_field_types.hpp" -#include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" -#include "traccc/finding/details/combinatorial_kalman_filter.hpp" - -// System include(s). -#include - -namespace traccc::host { - -combinatorial_kalman_filter_algorithm::output_type -combinatorial_kalman_filter_algorithm::operator()( - const default_detector::host& det, const magnetic_field& bfield, - const measurement_collection_types::const_view& measurements, - const bound_track_parameters_collection_types::const_view& seeds) const { - - // Perform the track finding using the appropriate templated implementation. - return magnetic_field_visitor>( - bfield, [&](const bfield_view_t& bfield_view) { - return details::combinatorial_kalman_filter( - det, bfield_view, measurements, seeds, m_config, m_mr.get(), - logger()); - }); -} - -} // namespace traccc::host diff --git a/core/src/finding/combinatorial_kalman_filter_algorithm_telescope_detector.cpp b/core/src/finding/combinatorial_kalman_filter_algorithm_telescope_detector.cpp deleted file mode 100644 index 982d13386f..0000000000 --- a/core/src/finding/combinatorial_kalman_filter_algorithm_telescope_detector.cpp +++ /dev/null @@ -1,33 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2024-2025 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// Local include(s). -#include "traccc/bfield/magnetic_field_types.hpp" -#include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" -#include "traccc/finding/details/combinatorial_kalman_filter.hpp" - -// System include(s). -#include - -namespace traccc::host { - -combinatorial_kalman_filter_algorithm::output_type -combinatorial_kalman_filter_algorithm::operator()( - const telescope_detector::host& det, const magnetic_field& bfield, - const measurement_collection_types::const_view& measurements, - const bound_track_parameters_collection_types::const_view& seeds) const { - - // Perform the track finding using the appropriate templated implementation. - return magnetic_field_visitor>( - bfield, [&](const bfield_view_t& bfield_view) { - return details::combinatorial_kalman_filter( - det, bfield_view, measurements, seeds, m_config, m_mr.get(), - logger()); - }); -} - -} // namespace traccc::host diff --git a/core/src/fitting/kalman_fitting_algorithm.cpp b/core/src/fitting/kalman_fitting_algorithm.cpp index 17c49d7dbb..f79ae5b554 100644 --- a/core/src/fitting/kalman_fitting_algorithm.cpp +++ b/core/src/fitting/kalman_fitting_algorithm.cpp @@ -8,6 +8,11 @@ // Project include(s). #include "traccc/fitting/kalman_fitting_algorithm.hpp" +#include "traccc/bfield/magnetic_field_types.hpp" +#include "traccc/fitting/details/kalman_fitting.hpp" +#include "traccc/fitting/details/kalman_fitting_types.hpp" +#include "traccc/utils/host_detector_bfield_visitor.hpp" + namespace traccc::host { kalman_fitting_algorithm::kalman_fitting_algorithm( @@ -15,4 +20,24 @@ kalman_fitting_algorithm::kalman_fitting_algorithm( std::unique_ptr logger) : messaging(std::move(logger)), m_config{config}, m_mr{mr}, m_copy(copy) {} +kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()( + const host_detector& det, const magnetic_field& bfield, + const edm::track_candidate_container::const_view& + track_candidates) const { + + // Perform the track fitting using the appropriate templated implementation. + return host_detector_magnetic_field_visitor>( + det, bfield, + [&]( + const typename detector_t::host& detector, + const bfield_view_t field) { + traccc::details::kalman_fitter_t + fitter{detector, field, m_config}; + return details::kalman_fitting( + fitter, track_candidates, m_mr.get(), m_copy.get()); + }); +} + } // namespace traccc::host diff --git a/core/src/fitting/kalman_fitting_algorithm_default_detector.cpp b/core/src/fitting/kalman_fitting_algorithm_default_detector.cpp deleted file mode 100644 index 15e560746d..0000000000 --- a/core/src/fitting/kalman_fitting_algorithm_default_detector.cpp +++ /dev/null @@ -1,35 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2022-2025 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// Project include(s). -#include "traccc/bfield/magnetic_field_types.hpp" -#include "traccc/fitting/details/kalman_fitting.hpp" -#include "traccc/fitting/details/kalman_fitting_types.hpp" -#include "traccc/fitting/kalman_fitting_algorithm.hpp" - -// System include(s). -#include - -namespace traccc::host { - -kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()( - const default_detector::host& det, const magnetic_field& bfield, - const edm::track_candidate_container::const_view& - track_candidates) const { - - // Perform the track fitting using the appropriate templated implementation. - return magnetic_field_visitor>( - bfield, [&](const bfield_view_t& bfield_view) { - traccc::details::kalman_fitter_t - fitter{det, bfield_view, m_config}; - return details::kalman_fitting( - fitter, track_candidates, m_mr.get(), m_copy.get()); - }); -} - -} // namespace traccc::host diff --git a/core/src/fitting/kalman_fitting_algorithm_telescope_detector.cpp b/core/src/fitting/kalman_fitting_algorithm_telescope_detector.cpp deleted file mode 100644 index 65ac4f5ad9..0000000000 --- a/core/src/fitting/kalman_fitting_algorithm_telescope_detector.cpp +++ /dev/null @@ -1,35 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2022-2025 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// Project include(s). -#include "traccc/bfield/magnetic_field_types.hpp" -#include "traccc/fitting/details/kalman_fitting.hpp" -#include "traccc/fitting/details/kalman_fitting_types.hpp" -#include "traccc/fitting/kalman_fitting_algorithm.hpp" - -// System include(s). -#include - -namespace traccc::host { - -kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()( - const telescope_detector::host& det, const magnetic_field& bfield, - const edm::track_candidate_container::const_view& - track_candidates) const { - - // Perform the track fitting using the appropriate templated implementation. - return magnetic_field_visitor>( - bfield, [&](const bfield_view_t& bfield_view) { - traccc::details::kalman_fitter_t - fitter{det, bfield_view, m_config}; - return details::kalman_fitting( - fitter, track_candidates, m_mr.get(), m_copy.get()); - }); -} - -} // namespace traccc::host diff --git a/core/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp b/core/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp index 3442ae1ce4..937a500194 100644 --- a/core/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp +++ b/core/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp @@ -8,6 +8,8 @@ // Library include(s). #include "traccc/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" +#include "silicon_pixel_spacepoint_formation.hpp" + namespace traccc::host { silicon_pixel_spacepoint_formation_algorithm:: @@ -15,4 +17,17 @@ silicon_pixel_spacepoint_formation_algorithm:: vecmem::memory_resource& mr, std::unique_ptr logger) : messaging(std::move(logger)), m_mr(mr) {} +silicon_pixel_spacepoint_formation_algorithm::output_type +silicon_pixel_spacepoint_formation_algorithm::operator()( + const host_detector& det, + const measurement_collection_types::const_view& meas) const { + + return host_detector_visitor( + det, [&]( + const typename detector_traits_t::host& detector) { + return details::silicon_pixel_spacepoint_formation(detector, meas, + m_mr); + }); +} + } // namespace traccc::host diff --git a/core/src/seeding/silicon_pixel_spacepoint_formation_algorithm_defdet.cpp b/core/src/seeding/silicon_pixel_spacepoint_formation_algorithm_defdet.cpp deleted file mode 100644 index 70b6e3b7e1..0000000000 --- a/core/src/seeding/silicon_pixel_spacepoint_formation_algorithm_defdet.cpp +++ /dev/null @@ -1,22 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2022-2024 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// Library include(s). -#include "silicon_pixel_spacepoint_formation.hpp" -#include "traccc/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" - -namespace traccc::host { - -silicon_pixel_spacepoint_formation_algorithm::output_type -silicon_pixel_spacepoint_formation_algorithm::operator()( - const default_detector::host& det, - const measurement_collection_types::const_view& meas) const { - - return details::silicon_pixel_spacepoint_formation(det, meas, m_mr); -} - -} // namespace traccc::host diff --git a/core/src/seeding/silicon_pixel_spacepoint_formation_algorithm_teldet.cpp b/core/src/seeding/silicon_pixel_spacepoint_formation_algorithm_teldet.cpp deleted file mode 100644 index 8b5ff8f50b..0000000000 --- a/core/src/seeding/silicon_pixel_spacepoint_formation_algorithm_teldet.cpp +++ /dev/null @@ -1,22 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2022-2024 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// Library include(s). -#include "silicon_pixel_spacepoint_formation.hpp" -#include "traccc/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" - -namespace traccc::host { - -silicon_pixel_spacepoint_formation_algorithm::output_type -silicon_pixel_spacepoint_formation_algorithm::operator()( - const telescope_detector::host& det, - const measurement_collection_types::const_view& meas) const { - - return details::silicon_pixel_spacepoint_formation(det, meas, m_mr); -} - -} // namespace traccc::host diff --git a/device/alpaka/CMakeLists.txt b/device/alpaka/CMakeLists.txt index 86e0b95f12..7b3828d40e 100644 --- a/device/alpaka/CMakeLists.txt +++ b/device/alpaka/CMakeLists.txt @@ -56,14 +56,10 @@ traccc_add_alpaka_library( traccc_alpaka alpaka TYPE SHARED # Track finding algorithm(s). "include/traccc/alpaka/finding/combinatorial_kalman_filter_algorithm.hpp" "src/finding/combinatorial_kalman_filter_algorithm.cpp" - "src/finding/combinatorial_kalman_filter_algorithm_default_detector.cpp" - "src/finding/combinatorial_kalman_filter_algorithm_telescope_detector.cpp" "src/finding/combinatorial_kalman_filter.hpp" # Track fitting algorithm(s). "include/traccc/alpaka/fitting/kalman_fitting_algorithm.hpp" "src/fitting/kalman_fitting_algorithm.cpp" - "src/fitting/kalman_fitting_algorithm_default_detector.cpp" - "src/fitting/kalman_fitting_algorithm_telescope_detector.cpp" "src/fitting/kalman_fitting.hpp" ) diff --git a/device/alpaka/include/traccc/alpaka/finding/combinatorial_kalman_filter_algorithm.hpp b/device/alpaka/include/traccc/alpaka/finding/combinatorial_kalman_filter_algorithm.hpp index 22bb9fc4b2..bf7cea48cf 100644 --- a/device/alpaka/include/traccc/alpaka/finding/combinatorial_kalman_filter_algorithm.hpp +++ b/device/alpaka/include/traccc/alpaka/finding/combinatorial_kalman_filter_algorithm.hpp @@ -17,6 +17,7 @@ #include "traccc/edm/track_parameters.hpp" #include "traccc/finding/finding_config.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/detector_buffer.hpp" #include "traccc/utils/algorithm.hpp" #include "traccc/utils/memory_resource.hpp" #include "traccc/utils/messaging.hpp" @@ -32,11 +33,7 @@ namespace traccc::alpaka { /// CKF track finding algorithm class combinatorial_kalman_filter_algorithm : public algorithm::buffer( - const default_detector::view&, const magnetic_field&, - const measurement_collection_types::const_view&, - const bound_track_parameters_collection_types::const_view&)>, - public algorithm::buffer( - const telescope_detector::view&, const magnetic_field&, + const detector_buffer&, const magnetic_field&, const measurement_collection_types::const_view&, const bound_track_parameters_collection_types::const_view&)>, public messaging { @@ -56,23 +53,7 @@ class combinatorial_kalman_filter_algorithm /// Execute the algorithm /// - /// @param det The (default) detector object - /// @param bfield The magnetic field object - /// @param measurements All measurements in an event - /// @param seeds All seeds in an event to start the track finding - /// with - /// - /// @return A container of the found track candidates - /// - output_type operator()( - const default_detector::view& det, const magnetic_field& bfield, - const measurement_collection_types::const_view& measurements, - const bound_track_parameters_collection_types::const_view& seeds) - const override; - - /// Execute the algorithm - /// - /// @param det The (telescope) detector object + /// @param det The detector object /// @param bfield The magnetic field object /// @param measurements All measurements in an event /// @param seeds All seeds in an event to start the track finding @@ -81,7 +62,7 @@ class combinatorial_kalman_filter_algorithm /// @return A container of the found track candidates /// output_type operator()( - const telescope_detector::view& det, const magnetic_field& bfield, + const detector_buffer& det, const magnetic_field& bfield, const measurement_collection_types::const_view& measurements, const bound_track_parameters_collection_types::const_view& seeds) const override; diff --git a/device/alpaka/include/traccc/alpaka/fitting/kalman_fitting_algorithm.hpp b/device/alpaka/include/traccc/alpaka/fitting/kalman_fitting_algorithm.hpp index 6fa432b448..524b51f308 100644 --- a/device/alpaka/include/traccc/alpaka/fitting/kalman_fitting_algorithm.hpp +++ b/device/alpaka/include/traccc/alpaka/fitting/kalman_fitting_algorithm.hpp @@ -16,6 +16,7 @@ #include "traccc/edm/track_fit_container.hpp" #include "traccc/fitting/fitting_config.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/detector_buffer.hpp" #include "traccc/utils/algorithm.hpp" #include "traccc/utils/memory_resource.hpp" #include "traccc/utils/messaging.hpp" @@ -31,11 +32,12 @@ namespace traccc::alpaka { /// Kalman filter based track fitting algorithm class kalman_fitting_algorithm : public algorithm::buffer( - const default_detector::view&, const magnetic_field&, + const detector_buffer&, const magnetic_field&, const edm::track_candidate_container::const_view&)>, public algorithm::buffer( - const telescope_detector::view&, const magnetic_field&, - const edm::track_candidate_container::const_view&)>, + const detector_buffer&, const magnetic_field&, + edm::track_fit_container::buffer&&, + const measurement_collection_types::const_view&)>, public messaging { public: @@ -57,31 +59,32 @@ class kalman_fitting_algorithm vecmem::copy& copy, queue& q, std::unique_ptr logger = getDummyLogger().clone()); - /// Execute the algorithm + /// Execute the algorithm using unfitted tracks /// - /// @param det The (default) detector object + /// @param det The detector object /// @param bfield The magnetic field object /// @param track_candidates All track candidates to fit /// /// @return A container of the fitted track states /// output_type operator()( - const default_detector::view& det, const magnetic_field& bfield, + const detector_buffer& det, const magnetic_field& bfield, const edm::track_candidate_container::const_view& track_candidates) const override; - /// Execute the algorithm + /// Execute the algorithm using fitted tracks /// - /// @param det The (telescope) detector object + /// @param det The detector object /// @param bfield The magnetic field object /// @param track_candidates All track candidates to fit /// /// @return A container of the fitted track states /// output_type operator()( - const telescope_detector::view& det, const magnetic_field& bfield, - const edm::track_candidate_container::const_view& - track_candidates) const override; + const detector_buffer& det, const magnetic_field& bfield, + edm::track_fit_container::buffer&& track_states, + const measurement_collection_types::const_view& measurements) + const override; private: /// Algorithm configuration diff --git a/device/alpaka/include/traccc/alpaka/seeding/spacepoint_formation_algorithm.hpp b/device/alpaka/include/traccc/alpaka/seeding/spacepoint_formation_algorithm.hpp index 0cf606db3d..b61c92af63 100644 --- a/device/alpaka/include/traccc/alpaka/seeding/spacepoint_formation_algorithm.hpp +++ b/device/alpaka/include/traccc/alpaka/seeding/spacepoint_formation_algorithm.hpp @@ -13,6 +13,7 @@ // Project include(s). #include "traccc/edm/measurement.hpp" #include "traccc/edm/spacepoint_collection.hpp" +#include "traccc/geometry/detector_buffer.hpp" #include "traccc/utils/algorithm.hpp" #include "traccc/utils/memory_resource.hpp" #include "traccc/utils/messaging.hpp" @@ -30,10 +31,9 @@ namespace traccc::alpaka { /// This algorithm performs the local-to-global transformation of the 2D /// measurements made on every detector module, into 3D spacepoint coordinates. /// -template class spacepoint_formation_algorithm : public algorithm, public messaging { @@ -57,7 +57,7 @@ class spacepoint_formation_algorithm /// measurement /// edm::spacepoint_collection::buffer operator()( - const typename detector_t::const_view_type& det_view, + const detector_buffer& det, const measurement_collection_types::const_view& measurements_view) const override; diff --git a/device/alpaka/src/finding/combinatorial_kalman_filter_algorithm.cpp b/device/alpaka/src/finding/combinatorial_kalman_filter_algorithm.cpp index 5a141d850a..7224c8058a 100644 --- a/device/alpaka/src/finding/combinatorial_kalman_filter_algorithm.cpp +++ b/device/alpaka/src/finding/combinatorial_kalman_filter_algorithm.cpp @@ -8,6 +8,12 @@ // Local include(s). #include "traccc/alpaka/finding/combinatorial_kalman_filter_algorithm.hpp" +#include "../utils/get_queue.hpp" +#include "../utils/magnetic_field_types.hpp" +#include "combinatorial_kalman_filter.hpp" +#include "traccc/bfield/magnetic_field_types.hpp" +#include "traccc/utils/detector_buffer_bfield_visitor.hpp" + namespace traccc::alpaka { combinatorial_kalman_filter_algorithm::combinatorial_kalman_filter_algorithm( @@ -19,4 +25,24 @@ combinatorial_kalman_filter_algorithm::combinatorial_kalman_filter_algorithm( m_copy{copy}, m_queue{q} {} +combinatorial_kalman_filter_algorithm::output_type +combinatorial_kalman_filter_algorithm::operator()( + const detector_buffer& det, const magnetic_field& bfield, + const measurement_collection_types::const_view& measurements, + const bound_track_parameters_collection_types::const_view& seeds) const { + + // Perform the track finding using the templated implementation. + return detector_buffer_magnetic_field_visitor< + detector_type_list, alpaka::bfield_type_list>( + det, bfield, + [&]( + const typename detector_t::view& detector, + const bfield_view_t& field) { + return details::combinatorial_kalman_filter< + typename detector_t::device>( + detector, field, measurements, seeds, m_config, m_mr, m_copy, + logger(), details::get_queue(m_queue.get())); + }); +} + } // namespace traccc::alpaka diff --git a/device/alpaka/src/finding/combinatorial_kalman_filter_algorithm_default_detector.cpp b/device/alpaka/src/finding/combinatorial_kalman_filter_algorithm_default_detector.cpp deleted file mode 100644 index 515b241dec..0000000000 --- a/device/alpaka/src/finding/combinatorial_kalman_filter_algorithm_default_detector.cpp +++ /dev/null @@ -1,35 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2025 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// Local include(s). -#include "../utils/get_queue.hpp" -#include "../utils/magnetic_field_types.hpp" -#include "combinatorial_kalman_filter.hpp" -#include "traccc/alpaka/finding/combinatorial_kalman_filter_algorithm.hpp" - -// Project include(s). -#include "traccc/bfield/magnetic_field_types.hpp" - -namespace traccc::alpaka { - -combinatorial_kalman_filter_algorithm::output_type -combinatorial_kalman_filter_algorithm::operator()( - const default_detector::view& det, const magnetic_field& bfield, - const measurement_collection_types::const_view& measurements, - const bound_track_parameters_collection_types::const_view& seeds) const { - - // Perform the track finding using the templated implementation. - return magnetic_field_visitor>( - bfield, [&](const bfield_view_t& bfield_view) { - return details::combinatorial_kalman_filter< - default_detector::device>(det, bfield_view, measurements, seeds, - m_config, m_mr, m_copy, logger(), - details::get_queue(m_queue.get())); - }); -} - -} // namespace traccc::alpaka diff --git a/device/alpaka/src/finding/combinatorial_kalman_filter_algorithm_telescope_detector.cpp b/device/alpaka/src/finding/combinatorial_kalman_filter_algorithm_telescope_detector.cpp deleted file mode 100644 index 07f06cfda5..0000000000 --- a/device/alpaka/src/finding/combinatorial_kalman_filter_algorithm_telescope_detector.cpp +++ /dev/null @@ -1,35 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2025 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// Local include(s). -#include "../utils/get_queue.hpp" -#include "../utils/magnetic_field_types.hpp" -#include "combinatorial_kalman_filter.hpp" -#include "traccc/alpaka/finding/combinatorial_kalman_filter_algorithm.hpp" - -// Project include(s). -#include "traccc/bfield/magnetic_field_types.hpp" - -namespace traccc::alpaka { - -combinatorial_kalman_filter_algorithm::output_type -combinatorial_kalman_filter_algorithm::operator()( - const telescope_detector::view& det, const magnetic_field& bfield, - const measurement_collection_types::const_view& measurements, - const bound_track_parameters_collection_types::const_view& seeds) const { - - // Perform the track finding using the templated implementation. - return magnetic_field_visitor>( - bfield, [&](const bfield_view_t& bfield_view) { - return details::combinatorial_kalman_filter< - telescope_detector::device>( - det, bfield_view, measurements, seeds, m_config, m_mr, m_copy, - logger(), details::get_queue(m_queue.get())); - }); -} - -} // namespace traccc::alpaka diff --git a/device/alpaka/src/fitting/kalman_fitting.hpp b/device/alpaka/src/fitting/kalman_fitting.hpp index be5ad7a0df..821a43a95e 100644 --- a/device/alpaka/src/fitting/kalman_fitting.hpp +++ b/device/alpaka/src/fitting/kalman_fitting.hpp @@ -35,14 +35,13 @@ struct fill_fitting_sort_keys { template ALPAKA_FN_ACC void operator()( TAcc const& acc, - edm::track_candidate_collection::const_view - track_candidates_view, + edm::track_fit_collection::const_view track_fit_view, vecmem::data::vector_view keys_view, vecmem::data::vector_view ids_view) const { const device::global_index_t globalThreadIdx = ::alpaka::getIdx<::alpaka::Grid, ::alpaka::Threads>(acc)[0]; - device::fill_fitting_sort_keys(globalThreadIdx, track_candidates_view, + device::fill_fitting_sort_keys(globalThreadIdx, track_fit_view, keys_view, ids_view); } }; @@ -52,17 +51,15 @@ struct fit_prelude { template ALPAKA_FN_ACC void operator()( TAcc const& acc, - vecmem::data::vector_view param_ids_view, edm::track_candidate_container::const_view track_candidates_view, - edm::track_fit_container::view track_states_view, - vecmem::data::vector_view param_liveness_view) const { + edm::track_fit_container::view track_states_view) + const { const device::global_index_t globalThreadIdx = ::alpaka::getIdx<::alpaka::Grid, ::alpaka::Threads>(acc)[0]; device::fit_prelude( - globalThreadIdx, param_ids_view, track_candidates_view, - track_states_view, param_liveness_view); + globalThreadIdx, track_candidates_view, track_states_view); } }; @@ -96,14 +93,15 @@ struct fit_backward { } // namespace kernels -/// Templated implementation of the Alpaka track fitting algorithm. +/// Templated implementation of the Alpaka track fitting algorithm for +/// fitted tracks. /// /// @tparam detector_t The (device) detector type to use /// @tparam bfield_t The magnetic field type to use /// /// @param[in] det_view A view of the detector geometry /// @param[in] field_view A view of the magnetic field -/// @param[in] track_candidates_view All track candidates to fit +/// @param[in] track_fit_view All track candidates to fit /// @param[in] config The fitting configuration /// @param[in] mr Memory resource(s) to use /// @param[in] copy The copy object to use for memory transfers @@ -116,41 +114,32 @@ typename edm::track_fit_container::buffer kalman_fitting( const typename detector_t::const_view_type& det_view, const bfield_t& field_view, - const typename edm::track_candidate_container< - typename detector_t::algebra_type>::const_view& track_candidates_view, + typename edm::track_fit_container< + typename detector_t::algebra_type>::buffer&& track_fit_buffer, + const measurement_collection_types::const_view& measurements, const fitting_config& config, const memory_resource& mr, vecmem::copy& copy, - Queue& queue) { + Queue& queue, bool forward_on_first_iteration = false) { // Number of threads per block to use. const Idx threadsPerBlock = getWarpSize() * 2; + typename edm::track_fit_container< + typename detector_t::algebra_type>::const_view track_fit_view{ + vecmem::get_data(track_fit_buffer.tracks), + vecmem::get_data(track_fit_buffer.states), measurements}; + // Get the number of tracks. const edm::track_candidate_collection< default_algebra>::const_device::size_type n_tracks = - copy.get_size(track_candidates_view.tracks); + copy.get_size(track_fit_view.tracks); // Get the sizes of the track candidates in each track. const std::vector candidate_sizes = - copy.get_sizes(track_candidates_view.tracks); - const unsigned int n_states = - std::accumulate(candidate_sizes.begin(), candidate_sizes.end(), 0u); - - // Create the result buffer. - typename edm::track_fit_container::buffer - track_states_buffer{ - {candidate_sizes, mr.main, mr.host, - vecmem::data::buffer_type::resizable}, - {n_states, mr.main, vecmem::data::buffer_type::resizable}}; - vecmem::copy::event_type tracks_setup_event = - copy.setup(track_states_buffer.tracks); - vecmem::copy::event_type track_states_setup_event = - copy.setup(track_states_buffer.states); + copy.get_sizes(track_fit_view.tracks); // Return early, if there are no tracks. if (n_tracks == 0) { - tracks_setup_event->wait(); - track_states_setup_event->wait(); - return track_states_buffer; + return track_fit_buffer; } std::vector seqs_sizes(candidate_sizes.size()); @@ -178,6 +167,7 @@ kalman_fitting( keys_setup_event->wait(); param_ids_setup_event->wait(); param_liveness_setup_event->wait(); + copy.memset(param_liveness_buffer, 1)->ignore(); // The execution range for the two kernels of the function. const Idx blocksPerGrid = @@ -186,8 +176,7 @@ kalman_fitting( // Fill the keys and param_ids buffers. ::alpaka::exec(queue, workDiv, kernels::fill_fitting_sort_keys{}, - track_candidates_view.tracks, - vecmem::get_data(keys_buffer), + track_fit_view.tracks, vecmem::get_data(keys_buffer), vecmem::get_data(param_ids_buffer)); ::alpaka::wait(queue); @@ -197,20 +186,6 @@ kalman_fitting( details::sort_by_key(queue, mr, keys_device.begin(), keys_device.end(), param_ids_device.begin()); - // Run the fitting, using the sorted parameter IDs. - typename edm::track_fit_container::view - track_states_view{track_states_buffer.tracks, - track_states_buffer.states, - track_candidates_view.measurements}; - tracks_setup_event->wait(); - track_states_setup_event->wait(); - - ::alpaka::exec(queue, workDiv, kernels::fit_prelude{}, - vecmem::get_data(param_ids_buffer), - track_candidates_view, track_states_view, - vecmem::get_data(param_liveness_buffer)); - ::alpaka::wait(queue); - // Allocate the fitting kernels's payload in host memory. using fitter_t = traccc::details::kalman_fitter_t; device::fit_payload host_payload{ @@ -218,7 +193,8 @@ kalman_fitting( .field_data = field_view, .param_ids_view = param_ids_buffer, .param_liveness_view = param_liveness_buffer, - .tracks_view = track_states_view, + .tracks_view = {track_fit_buffer.tracks, track_fit_buffer.states, + measurements}, .barcodes_view = seqs_buffer}; // Now copy it to device memory. vecmem::data::vector_buffer> device_payload( @@ -231,16 +207,113 @@ kalman_fitting( for (std::size_t i = 0; i < config.n_iterations; ++i) { // Run the track fitting - ::alpaka::exec(queue, workDiv, kernels::fit_forward{}, - config, device_payload.ptr()); - ::alpaka::wait(queue); + if (i > 0 || forward_on_first_iteration) { + ::alpaka::exec(queue, workDiv, + kernels::fit_forward{}, config, + device_payload.ptr()); + ::alpaka::wait(queue); + } ::alpaka::exec(queue, workDiv, kernels::fit_backward{}, config, device_payload.ptr()); ::alpaka::wait(queue); } // Return the fitted tracks. - return track_states_buffer; + return track_fit_buffer; +} + +/// Templated implementation of the Alpaka track fitting algorithm for +/// unfitted tracks. +/// +/// @tparam detector_t The (device) detector type to use +/// @tparam bfield_t The magnetic field type to use +/// +/// @param[in] det_view A view of the detector geometry +/// @param[in] field_view A view of the magnetic field +/// @param[in] track_candidates_view All track candidates to fit +/// @param[in] config The fitting configuration +/// @param[in] mr Memory resource(s) to use +/// @param[in] copy The copy object to use for memory transfers +/// @param[in] queue The Alpaka queue to use for execution +/// +/// @return A container of the fitted track states +/// +template +typename edm::track_fit_container::buffer +kalman_fitting( + const typename detector_t::const_view_type& det_view, + const bfield_t& field_view, + const typename edm::track_candidate_container< + typename detector_t::algebra_type>::const_view& track_candidates_view, + const fitting_config& config, const memory_resource& mr, vecmem::copy& copy, + Queue& queue) { + + // Number of threads per block to use. + const Idx threadsPerBlock = getWarpSize() * 2; + + // Get the number of tracks. + const edm::track_candidate_collection< + default_algebra>::const_device::size_type n_tracks = + copy.get_size(track_candidates_view.tracks); + + // Get the sizes of the track candidates in each track. + const std::vector candidate_sizes = + copy.get_sizes(track_candidates_view.tracks); + const unsigned int n_states = + std::accumulate(candidate_sizes.begin(), candidate_sizes.end(), 0u); + + // Create the result buffer. + typename edm::track_fit_container::buffer + track_states_buffer{ + {candidate_sizes, mr.main, mr.host, + vecmem::data::buffer_type::resizable}, + {n_states, mr.main, vecmem::data::buffer_type::resizable}}; + vecmem::copy::event_type tracks_setup_event = + copy.setup(track_states_buffer.tracks); + vecmem::copy::event_type track_states_setup_event = + copy.setup(track_states_buffer.states); + + // Return early, if there are no tracks. + if (n_tracks == 0) { + tracks_setup_event->wait(); + track_states_setup_event->wait(); + return track_states_buffer; + } + + std::vector seqs_sizes(candidate_sizes.size()); + std::transform(candidate_sizes.begin(), candidate_sizes.end(), + seqs_sizes.begin(), [&config](const unsigned int sz) { + return std::max(sz * config.barcode_sequence_size_factor, + config.min_barcode_sequence_capacity); + }); + vecmem::data::jagged_vector_buffer seqs_buffer{ + seqs_sizes, mr.main, mr.host, vecmem::data::buffer_type::resizable}; + copy.setup(seqs_buffer)->wait(); + + // The execution range for the two kernels of the function. + const Idx blocksPerGrid = + (n_tracks + threadsPerBlock - 1) / threadsPerBlock; + const auto workDiv = makeWorkDiv(blocksPerGrid, threadsPerBlock); + + // Run the fitting, using the sorted parameter IDs. + typename edm::track_fit_container::view + track_states_view{track_states_buffer.tracks, + track_states_buffer.states, + track_candidates_view.measurements}; + tracks_setup_event->wait(); + track_states_setup_event->wait(); + + ::alpaka::exec(queue, workDiv, kernels::fit_prelude{}, + track_candidates_view, track_states_view); + ::alpaka::wait(queue); + + return kalman_fitting( + det_view, field_view, + typename edm::track_fit_container< + typename detector_t::algebra_type>::buffer{ + std::move(track_states_buffer.tracks), + std::move(track_states_buffer.states)}, + track_candidates_view.measurements, config, mr, copy, queue, true); } } // namespace traccc::alpaka::details diff --git a/device/alpaka/src/fitting/kalman_fitting_algorithm.cpp b/device/alpaka/src/fitting/kalman_fitting_algorithm.cpp index f7893f10bd..dc32a8b1da 100644 --- a/device/alpaka/src/fitting/kalman_fitting_algorithm.cpp +++ b/device/alpaka/src/fitting/kalman_fitting_algorithm.cpp @@ -8,6 +8,13 @@ // Local include(s). #include "traccc/alpaka/fitting/kalman_fitting_algorithm.hpp" +#include "../utils/get_queue.hpp" +#include "../utils/magnetic_field_types.hpp" +#include "kalman_fitting.hpp" +#include "traccc/alpaka/fitting/kalman_fitting_algorithm.hpp" +#include "traccc/bfield/magnetic_field_types.hpp" +#include "traccc/utils/detector_buffer_bfield_visitor.hpp" + namespace traccc::alpaka { kalman_fitting_algorithm::kalman_fitting_algorithm( @@ -19,4 +26,41 @@ kalman_fitting_algorithm::kalman_fitting_algorithm( m_copy{copy}, m_queue{q} {} +kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()( + const detector_buffer& det, const magnetic_field& bfield, + const edm::track_candidate_container::const_view& + track_candidates) const { + + // Run the track fitting. + return detector_buffer_magnetic_field_visitor< + detector_type_list, alpaka::bfield_type_list>( + det, bfield, + [&]( + const typename detector_t::view& detector, + const bfield_view_t& field) { + return details::kalman_fitting( + detector, field, track_candidates, m_config, m_mr, m_copy.get(), + details::get_queue(m_queue.get())); + }); +} + +kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()( + const detector_buffer& det, const magnetic_field& bfield, + edm::track_fit_container::buffer&& track_states, + const measurement_collection_types::const_view& measurements) const { + + // Run the track fitting. + return detector_buffer_magnetic_field_visitor< + detector_type_list, alpaka::bfield_type_list>( + det, bfield, + [&]( + const typename detector_t::view& detector, + const bfield_view_t& field) { + return details::kalman_fitting( + detector, field, std::move(track_states), measurements, + m_config, m_mr, m_copy.get(), + details::get_queue(m_queue.get())); + }); +} + } // namespace traccc::alpaka diff --git a/device/alpaka/src/fitting/kalman_fitting_algorithm_default_detector.cpp b/device/alpaka/src/fitting/kalman_fitting_algorithm_default_detector.cpp deleted file mode 100644 index 61e3321f7e..0000000000 --- a/device/alpaka/src/fitting/kalman_fitting_algorithm_default_detector.cpp +++ /dev/null @@ -1,33 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2022-2025 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// Local include(s). -#include "../utils/get_queue.hpp" -#include "../utils/magnetic_field_types.hpp" -#include "kalman_fitting.hpp" -#include "traccc/alpaka/fitting/kalman_fitting_algorithm.hpp" - -// Project include(s). -#include "traccc/bfield/magnetic_field_types.hpp" - -namespace traccc::alpaka { - -kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()( - const default_detector::view& det, const magnetic_field& bfield, - const edm::track_candidate_container::const_view& - track_candidates) const { - - // Run the track fitting. - return magnetic_field_visitor>( - bfield, [&](const bfield_view_t& bfield_view) { - return details::kalman_fitting( - det, bfield_view, track_candidates, m_config, m_mr, - m_copy.get(), details::get_queue(m_queue.get())); - }); -} - -} // namespace traccc::alpaka diff --git a/device/alpaka/src/fitting/kalman_fitting_algorithm_telescope_detector.cpp b/device/alpaka/src/fitting/kalman_fitting_algorithm_telescope_detector.cpp deleted file mode 100644 index 671646c69c..0000000000 --- a/device/alpaka/src/fitting/kalman_fitting_algorithm_telescope_detector.cpp +++ /dev/null @@ -1,33 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2022-2025 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// Local include(s). -#include "../utils/get_queue.hpp" -#include "../utils/magnetic_field_types.hpp" -#include "kalman_fitting.hpp" -#include "traccc/alpaka/fitting/kalman_fitting_algorithm.hpp" - -// Project include(s). -#include "traccc/bfield/magnetic_field_types.hpp" - -namespace traccc::alpaka { - -kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()( - const telescope_detector::view& det, const magnetic_field& bfield, - const edm::track_candidate_container::const_view& - track_candidates) const { - - // Run the track fitting. - return magnetic_field_visitor>( - bfield, [&](const bfield_view_t& bfield_view) { - return details::kalman_fitting( - det, bfield_view, track_candidates, m_config, m_mr, - m_copy.get(), details::get_queue(m_queue.get())); - }); -} - -} // namespace traccc::alpaka diff --git a/device/alpaka/src/seeding/spacepoint_formation_algorithm.cpp b/device/alpaka/src/seeding/spacepoint_formation_algorithm.cpp index b5e2fc33a6..af6f3858eb 100644 --- a/device/alpaka/src/seeding/spacepoint_formation_algorithm.cpp +++ b/device/alpaka/src/seeding/spacepoint_formation_algorithm.cpp @@ -22,7 +22,7 @@ template struct FormSpacepointsKernel { template ALPAKA_FN_ACC void operator()( - TAcc const& acc, typename detector_t::const_view_type det_view, + TAcc const& acc, typename detector_t::view det_view, measurement_collection_types::const_view measurements_view, edm::spacepoint_collection::view spacepoints_view) const { @@ -34,16 +34,13 @@ struct FormSpacepointsKernel { } }; -template -spacepoint_formation_algorithm::spacepoint_formation_algorithm( +spacepoint_formation_algorithm::spacepoint_formation_algorithm( const traccc::memory_resource& mr, vecmem::copy& copy, queue& q, std::unique_ptr logger) : messaging(std::move(logger)), m_mr(mr), m_copy(copy), m_queue(q) {} -template -edm::spacepoint_collection::buffer -spacepoint_formation_algorithm::operator()( - const typename detector_t::const_view_type& det_view, +edm::spacepoint_collection::buffer spacepoint_formation_algorithm::operator()( + const detector_buffer& det, const measurement_collection_types::const_view& measurements_view) const { // Get a convenience variable for the queue that we'll be using. @@ -69,14 +66,17 @@ spacepoint_formation_algorithm::operator()( const unsigned int nBlocks = (num_measurements + blockSize - 1) / blockSize; auto workDiv = makeWorkDiv(blockSize, nBlocks); - // Launch the spacepoint formation kernel. - ::alpaka::exec(queue, workDiv, FormSpacepointsKernel{}, - det_view, measurements_view, spacepoints_view); + detector_buffer_visitor( + det, [&]( + const typename detector_traits_t::view& det_view) { + // Launch the spacepoint formation kernel. + ::alpaka::exec(queue, workDiv, + FormSpacepointsKernel{}, + det_view, measurements_view, spacepoints_view); + }); // Return the reconstructed spacepoints. return spacepoints; } -// Explicit template instantiation -template class spacepoint_formation_algorithm; } // namespace traccc::alpaka diff --git a/device/common/include/traccc/fitting/device/fill_fitting_sort_keys.hpp b/device/common/include/traccc/fitting/device/fill_fitting_sort_keys.hpp index 66d39c3b43..75a76d4cc5 100644 --- a/device/common/include/traccc/fitting/device/fill_fitting_sort_keys.hpp +++ b/device/common/include/traccc/fitting/device/fill_fitting_sort_keys.hpp @@ -12,21 +12,21 @@ #include "traccc/edm/device/sort_key.hpp" // Project include(s). -#include "traccc/edm/track_candidate_collection.hpp" +#include "traccc/edm/track_fit_collection.hpp" namespace traccc::device { /// Function used to fill key container /// /// @param[in] globalIndex The index of the current thread -/// @param[in] track_candidates_view The input track candidates +/// @param[in] track_fit_view The input track states /// @param[out] keys_view The key values /// @param[out] ids_view The param ids /// TRACCC_HOST_DEVICE inline void fill_fitting_sort_keys( global_index_t globalIndex, - const edm::track_candidate_collection::const_view& - track_candidates_view, + const edm::track_fit_collection::const_view& + track_fit_view, vecmem::data::vector_view keys_view, vecmem::data::vector_view ids_view); diff --git a/device/common/include/traccc/fitting/device/fit_prelude.hpp b/device/common/include/traccc/fitting/device/fit_prelude.hpp index ba0dc1cd45..1305b723d1 100644 --- a/device/common/include/traccc/fitting/device/fit_prelude.hpp +++ b/device/common/include/traccc/fitting/device/fit_prelude.hpp @@ -24,11 +24,9 @@ namespace traccc::device { template TRACCC_HOST_DEVICE inline void fit_prelude( const global_index_t globalIndex, - vecmem::data::vector_view param_ids_view, typename edm::track_candidate_container::const_view track_candidates_view, - typename edm::track_fit_container::view tracks_view, - vecmem::data::vector_view param_liveness_view) { + typename edm::track_fit_container::view tracks_view) { typename edm::track_fit_collection::device tracks( tracks_view.tracks); @@ -40,16 +38,11 @@ TRACCC_HOST_DEVICE inline void fit_prelude( typename edm::track_state_collection::device track_states( tracks_view.states); - vecmem::device_vector param_ids(param_ids_view); - vecmem::device_vector param_liveness(param_liveness_view); - - const unsigned int param_id = param_ids.at(globalIndex); - - auto track = tracks.at(param_id); + auto track = tracks.at(globalIndex); const typename edm::track_candidate_collection::const_device track_candidates{track_candidates_view.tracks}; - const auto track_candidate = track_candidates.at(param_id); + const auto track_candidate = track_candidates.at(globalIndex); const auto track_candidate_measurement_indices = track_candidate.measurement_indices(); const measurement_collection_types::const_device measurements{ @@ -62,7 +55,6 @@ TRACCC_HOST_DEVICE inline void fit_prelude( // TODO: Set other stuff in the header? track.params() = track_candidate.params(); - param_liveness.at(param_id) = 1u; } } // namespace traccc::device diff --git a/device/common/include/traccc/fitting/device/impl/fill_fitting_sort_keys.ipp b/device/common/include/traccc/fitting/device/impl/fill_fitting_sort_keys.ipp index 9ddfb30713..85f5587084 100644 --- a/device/common/include/traccc/fitting/device/impl/fill_fitting_sort_keys.ipp +++ b/device/common/include/traccc/fitting/device/impl/fill_fitting_sort_keys.ipp @@ -11,13 +11,13 @@ namespace traccc::device { TRACCC_HOST_DEVICE inline void fill_fitting_sort_keys( const global_index_t globalIndex, - const edm::track_candidate_collection::const_view& - track_candidates_view, + const edm::track_fit_collection::const_view& + track_fit_view, vecmem::data::vector_view keys_view, vecmem::data::vector_view ids_view) { - const edm::track_candidate_collection::const_device - track_candidates(track_candidates_view); + const edm::track_fit_collection::const_device track_states( + track_fit_view); // Keys vecmem::device_vector keys_device(keys_view); @@ -31,7 +31,7 @@ TRACCC_HOST_DEVICE inline void fill_fitting_sort_keys( // Key = The number of measurements keys_device.at(globalIndex) = static_cast( - track_candidates.at(globalIndex).measurement_indices().size()); + track_states.at(globalIndex).state_indices().size()); ids_device.at(globalIndex) = globalIndex; } diff --git a/device/common/include/traccc/seeding/device/form_spacepoints.hpp b/device/common/include/traccc/seeding/device/form_spacepoints.hpp index 864d226d7c..c88dfeb454 100644 --- a/device/common/include/traccc/seeding/device/form_spacepoints.hpp +++ b/device/common/include/traccc/seeding/device/form_spacepoints.hpp @@ -27,7 +27,7 @@ namespace traccc::device { /// template TRACCC_HOST_DEVICE inline void form_spacepoints( - global_index_t globalIndex, typename detector_t::const_view_type det_view, + global_index_t globalIndex, typename detector_t::view det_view, const measurement_collection_types::const_view& measurements_view, edm::spacepoint_collection::view spacepoints_view); diff --git a/device/common/include/traccc/seeding/device/impl/form_spacepoints.ipp b/device/common/include/traccc/seeding/device/impl/form_spacepoints.ipp index eb8cd1b49c..8e61437943 100644 --- a/device/common/include/traccc/seeding/device/impl/form_spacepoints.ipp +++ b/device/common/include/traccc/seeding/device/impl/form_spacepoints.ipp @@ -17,8 +17,7 @@ namespace traccc::device { template TRACCC_HOST_DEVICE inline void form_spacepoints( - const global_index_t globalIndex, - typename detector_t::const_view_type det_view, + const global_index_t globalIndex, typename detector_t::view det_view, const measurement_collection_types::const_view& measurements_view, edm::spacepoint_collection::view spacepoints_view) { @@ -32,7 +31,7 @@ TRACCC_HOST_DEVICE inline void form_spacepoints( } // Create the tracking geometry - detector_t det(det_view); + typename detector_t::device det(det_view); // Set up the output container(s). edm::spacepoint_collection::device spacepoints(spacepoints_view); diff --git a/device/cuda/CMakeLists.txt b/device/cuda/CMakeLists.txt index 5fa19bbfd5..dfeb30ed1f 100644 --- a/device/cuda/CMakeLists.txt +++ b/device/cuda/CMakeLists.txt @@ -58,8 +58,7 @@ traccc_add_library( traccc_cuda cuda TYPE SHARED # Track finding algorithm(s). "include/traccc/cuda/finding/combinatorial_kalman_filter_algorithm.hpp" "src/finding/combinatorial_kalman_filter_algorithm.cpp" - "src/finding/combinatorial_kalman_filter_algorithm_default_detector.cu" - "src/finding/combinatorial_kalman_filter_algorithm_telescope_detector.cu" + "src/finding/combinatorial_kalman_filter_algorithm.cu" "src/finding/combinatorial_kalman_filter.cuh" "src/finding/kernels/make_barcode_sequence.cu" "src/finding/kernels/make_barcode_sequence.cuh" @@ -116,8 +115,7 @@ traccc_add_library( traccc_cuda cuda TYPE SHARED # Track fitting algorithm(s). "include/traccc/cuda/fitting/kalman_fitting_algorithm.hpp" "src/fitting/kalman_fitting_algorithm.cpp" - "src/fitting/kalman_fitting_algorithm_default_detector.cu" - "src/fitting/kalman_fitting_algorithm_telescope_detector.cu" + "src/fitting/kalman_fitting_algorithm.cu" "src/fitting/kalman_fitting.cuh" "src/fitting/kernels/fill_fitting_sort_keys.cu" "src/fitting/kernels/fit_prelude.cu" diff --git a/device/cuda/include/traccc/cuda/finding/combinatorial_kalman_filter_algorithm.hpp b/device/cuda/include/traccc/cuda/finding/combinatorial_kalman_filter_algorithm.hpp index f872e0f26a..8cf4bcefee 100644 --- a/device/cuda/include/traccc/cuda/finding/combinatorial_kalman_filter_algorithm.hpp +++ b/device/cuda/include/traccc/cuda/finding/combinatorial_kalman_filter_algorithm.hpp @@ -17,6 +17,7 @@ #include "traccc/edm/track_parameters.hpp" #include "traccc/finding/finding_config.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/detector_buffer.hpp" #include "traccc/utils/algorithm.hpp" #include "traccc/utils/memory_resource.hpp" #include "traccc/utils/messaging.hpp" @@ -32,11 +33,7 @@ namespace traccc::cuda { /// CKF track finding algorithm class combinatorial_kalman_filter_algorithm : public algorithm::buffer( - const default_detector::view&, const magnetic_field&, - const measurement_collection_types::const_view&, - const bound_track_parameters_collection_types::const_view&)>, - public algorithm::buffer( - const telescope_detector::view&, const magnetic_field&, + const detector_buffer&, const magnetic_field&, const measurement_collection_types::const_view&, const bound_track_parameters_collection_types::const_view&)>, public messaging { @@ -65,23 +62,7 @@ class combinatorial_kalman_filter_algorithm /// @return A container of the found track candidates /// output_type operator()( - const default_detector::view& det, const magnetic_field& bfield, - const measurement_collection_types::const_view& measurements, - const bound_track_parameters_collection_types::const_view& seeds) - const override; - - /// Execute the algorithm - /// - /// @param det The (telescope) detector object - /// @param bfield The magnetic field object - /// @param measurements All measurements in an event - /// @param seeds All seeds in an event to start the track finding - /// with - /// - /// @return A container of the found track candidates - /// - output_type operator()( - const telescope_detector::view& det, const magnetic_field& bfield, + const detector_buffer& det, const magnetic_field& bfield, const measurement_collection_types::const_view& measurements, const bound_track_parameters_collection_types::const_view& seeds) const override; diff --git a/device/cuda/include/traccc/cuda/fitting/kalman_fitting_algorithm.hpp b/device/cuda/include/traccc/cuda/fitting/kalman_fitting_algorithm.hpp index 00ecc465aa..99785f3a3d 100644 --- a/device/cuda/include/traccc/cuda/fitting/kalman_fitting_algorithm.hpp +++ b/device/cuda/include/traccc/cuda/fitting/kalman_fitting_algorithm.hpp @@ -16,6 +16,7 @@ #include "traccc/edm/track_fit_container.hpp" #include "traccc/fitting/fitting_config.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/detector_buffer.hpp" #include "traccc/utils/algorithm.hpp" #include "traccc/utils/memory_resource.hpp" #include "traccc/utils/messaging.hpp" @@ -31,11 +32,12 @@ namespace traccc::cuda { /// Kalman filter based track fitting algorithm class kalman_fitting_algorithm : public algorithm::buffer( - const default_detector::view&, const magnetic_field&, + const detector_buffer&, const magnetic_field&, const edm::track_candidate_container::const_view&)>, public algorithm::buffer( - const telescope_detector::view&, const magnetic_field&, - const edm::track_candidate_container::const_view&)>, + const detector_buffer&, const magnetic_field&, + edm::track_fit_container::buffer&&, + const measurement_collection_types::const_view&)>, public messaging { public: @@ -57,31 +59,32 @@ class kalman_fitting_algorithm vecmem::copy& copy, stream& str, std::unique_ptr logger = getDummyLogger().clone()); - /// Execute the algorithm + /// Execute the algorithm using unfitted tracks /// - /// @param det The (default) detector object - /// @param bfield The magnetic field object + /// @param det The detector object + /// @param field The magnetic field object /// @param track_candidates All track candidates to fit /// /// @return A container of the fitted track states /// output_type operator()( - const default_detector::view& det, const magnetic_field& bfield, + const detector_buffer& det, const magnetic_field& field, const edm::track_candidate_container::const_view& track_candidates) const override; - /// Execute the algorithm + /// Execute the algorithm using fitted tracks /// - /// @param det The (telescope) detector object - /// @param bfield The magnetic field object + /// @param det The detector object + /// @param field The magnetic field object /// @param track_candidates All track candidates to fit /// /// @return A container of the fitted track states /// output_type operator()( - const telescope_detector::view& det, const magnetic_field& bfield, - const edm::track_candidate_container::const_view& - track_candidates) const override; + const detector_buffer& det, const magnetic_field& field, + edm::track_fit_container::buffer&& track_candidates, + const measurement_collection_types::const_view& measurements) + const override; private: /// Algorithm configuration diff --git a/device/cuda/include/traccc/cuda/seeding/spacepoint_formation_algorithm.hpp b/device/cuda/include/traccc/cuda/seeding/spacepoint_formation_algorithm.hpp index 561ce62807..02039ee4f7 100644 --- a/device/cuda/include/traccc/cuda/seeding/spacepoint_formation_algorithm.hpp +++ b/device/cuda/include/traccc/cuda/seeding/spacepoint_formation_algorithm.hpp @@ -13,6 +13,7 @@ // Project include(s). #include "traccc/edm/measurement.hpp" #include "traccc/edm/spacepoint_collection.hpp" +#include "traccc/geometry/detector_buffer.hpp" #include "traccc/utils/algorithm.hpp" #include "traccc/utils/memory_resource.hpp" #include "traccc/utils/messaging.hpp" @@ -30,10 +31,9 @@ namespace traccc::cuda { /// This algorithm performs the local-to-global transformation of the 2D /// measurements made on every detector module, into 3D spacepoint coordinates. /// -template class spacepoint_formation_algorithm : public algorithm, public messaging { @@ -55,7 +55,7 @@ class spacepoint_formation_algorithm /// measurement /// edm::spacepoint_collection::buffer operator()( - const typename detector_t::const_view_type& det_view, + const detector_buffer& detector, const measurement_collection_types::const_view& measurements) const override; diff --git a/device/cuda/src/finding/combinatorial_kalman_filter_algorithm.cpp b/device/cuda/src/finding/combinatorial_kalman_filter_algorithm.cpp index ef635fb96e..445512a000 100644 --- a/device/cuda/src/finding/combinatorial_kalman_filter_algorithm.cpp +++ b/device/cuda/src/finding/combinatorial_kalman_filter_algorithm.cpp @@ -8,6 +8,7 @@ // Local include(s). #include "traccc/cuda/finding/combinatorial_kalman_filter_algorithm.hpp" +#include "../utils/magnetic_field_types.hpp" #include "../utils/utils.hpp" namespace traccc::cuda { diff --git a/device/cuda/src/finding/combinatorial_kalman_filter_algorithm_default_detector.cu b/device/cuda/src/finding/combinatorial_kalman_filter_algorithm.cu similarity index 58% rename from device/cuda/src/finding/combinatorial_kalman_filter_algorithm_default_detector.cu rename to device/cuda/src/finding/combinatorial_kalman_filter_algorithm.cu index ebef1c4ead..1d7a453bf8 100644 --- a/device/cuda/src/finding/combinatorial_kalman_filter_algorithm_default_detector.cu +++ b/device/cuda/src/finding/combinatorial_kalman_filter_algorithm.cu @@ -9,6 +9,7 @@ #include "../utils/magnetic_field_types.hpp" #include "combinatorial_kalman_filter.cuh" #include "traccc/cuda/finding/combinatorial_kalman_filter_algorithm.hpp" +#include "traccc/utils/detector_buffer_bfield_visitor.hpp" // Project include(s). #include "traccc/bfield/magnetic_field_types.hpp" @@ -20,17 +21,21 @@ namespace traccc::cuda { combinatorial_kalman_filter_algorithm::output_type combinatorial_kalman_filter_algorithm::operator()( - const default_detector::view& det, const magnetic_field& bfield, + const detector_buffer& det, const magnetic_field& field, const measurement_collection_types::const_view& measurements, const bound_track_parameters_collection_types::const_view& seeds) const { // Perform the track finding using the appropriate templated implementation. - return magnetic_field_visitor>( - bfield, [&](const bfield_view_t& bfield_view) { + return detector_buffer_magnetic_field_visitor< + detector_type_list, cuda::bfield_type_list>( + det, field, + [&]( + const typename detector_t::view& detector, + const bfield_view_t& bfield) { return details::combinatorial_kalman_filter< - default_detector::device>(det, bfield_view, measurements, seeds, - m_config, m_mr, m_copy, logger(), - m_stream, m_warp_size); + typename detector_t::device>(detector, bfield, measurements, + seeds, m_config, m_mr, m_copy, + logger(), m_stream, m_warp_size); }); } diff --git a/device/cuda/src/finding/combinatorial_kalman_filter_algorithm_telescope_detector.cu b/device/cuda/src/finding/combinatorial_kalman_filter_algorithm_telescope_detector.cu deleted file mode 100644 index 5413dc4465..0000000000 --- a/device/cuda/src/finding/combinatorial_kalman_filter_algorithm_telescope_detector.cu +++ /dev/null @@ -1,37 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2025 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// Local include(s). -#include "../utils/magnetic_field_types.hpp" -#include "combinatorial_kalman_filter.cuh" -#include "traccc/cuda/finding/combinatorial_kalman_filter_algorithm.hpp" - -// Project include(s). -#include "traccc/bfield/magnetic_field_types.hpp" - -// System include(s). -#include - -namespace traccc::cuda { - -combinatorial_kalman_filter_algorithm::output_type -combinatorial_kalman_filter_algorithm::operator()( - const telescope_detector::view& det, const magnetic_field& bfield, - const measurement_collection_types::const_view& measurements, - const bound_track_parameters_collection_types::const_view& seeds) const { - - // Perform the track finding using the appropriate templated implementation. - return magnetic_field_visitor>( - bfield, [&](const bfield_view_t& bfield_view) { - return details::combinatorial_kalman_filter< - telescope_detector::device>(det, bfield_view, measurements, - seeds, m_config, m_mr, m_copy, - logger(), m_stream, m_warp_size); - }); -} - -} // namespace traccc::cuda diff --git a/device/cuda/src/fitting/kalman_fitting.cuh b/device/cuda/src/fitting/kalman_fitting.cuh index f1ebd61d8e..e2b0aaec32 100644 --- a/device/cuda/src/fitting/kalman_fitting.cuh +++ b/device/cuda/src/fitting/kalman_fitting.cuh @@ -37,14 +37,15 @@ namespace traccc::cuda::details { -/// Templated implementation of the CUDA track fitting algorithm. +/// Templated implementation of the CUDA track fitting algorithm for fitted +/// tracks; reuses input memory /// /// @tparam detector_t The (device) detector type to use /// @tparam bfield_t The magnetic field type to use /// /// @param[in] det_view A view of the detector geometry /// @param[in] field_view A view of the magnetic field -/// @param[in] track_candidates_view All track candidates to fit +/// @param[in] track_fit_view All track candidates to fit /// @param[in] config The fitting configuration /// @param[in] mr Memory resource(s) to use /// @param[in] copy The copy object to use for memory transfers @@ -57,42 +58,39 @@ typename edm::track_fit_container::buffer kalman_fitting( const typename detector_t::const_view_type& det_view, const bfield_t& field_view, - const typename edm::track_candidate_container< - typename detector_t::algebra_type>::const_view& track_candidates_view, + typename edm::track_fit_container< + typename detector_t::algebra_type>::buffer&& track_fit_buffer, + const measurement_collection_types::const_view& measurements, const fitting_config& config, const memory_resource& mr, vecmem::copy& copy, - stream& str, unsigned int warp_size) { + stream& str, unsigned int warp_size, + bool forward_on_first_iteration = false) { // Get a convenience variable for the stream that we'll be using. cudaStream_t stream = details::get_stream(str); + typename edm::track_fit_container< + typename detector_t::algebra_type>::const_view track_fit_view{ + vecmem::get_data(track_fit_buffer.tracks), + vecmem::get_data(track_fit_buffer.states), measurements}; + // Get the number of tracks. - const edm::track_candidate_collection< - default_algebra>::const_device::size_type n_tracks = - copy.get_size(track_candidates_view.tracks); + const edm::track_fit_collection::const_device::size_type + n_tracks = copy.get_size(track_fit_view.tracks); // Get the sizes of the track candidates in each track. - const std::vector candidate_sizes = - copy.get_sizes(track_candidates_view.tracks); + const std::vector state_sizes = + copy.get_sizes(track_fit_view.tracks); const unsigned int n_states = - std::accumulate(candidate_sizes.begin(), candidate_sizes.end(), 0u); - - // Create the result buffer. - typename edm::track_fit_container::buffer - track_states_buffer{ - {candidate_sizes, mr.main, mr.host, - vecmem::data::buffer_type::resizable}, - {n_states, mr.main, vecmem::data::buffer_type::resizable}}; - copy.setup(track_states_buffer.tracks)->ignore(); - copy.setup(track_states_buffer.states)->ignore(); + std::accumulate(state_sizes.begin(), state_sizes.end(), 0u); // Return early, if there are no tracks. if (n_tracks == 0) { - return track_states_buffer; + return track_fit_buffer; } - std::vector seqs_sizes(candidate_sizes.size()); - std::transform(candidate_sizes.begin(), candidate_sizes.end(), - seqs_sizes.begin(), [&config](const unsigned int sz) { + std::vector seqs_sizes(state_sizes.size()); + std::transform(state_sizes.begin(), state_sizes.end(), seqs_sizes.begin(), + [&config](const unsigned int sz) { return std::max(sz * config.barcode_sequence_size_factor, config.min_barcode_sequence_capacity); }); @@ -115,15 +113,15 @@ kalman_fitting( keys_setup_event->ignore(); param_ids_setup_event->ignore(); param_liveness_setup_event->ignore(); + copy.memset(param_liveness_buffer, 1)->ignore(); // Launch parameters for all the kernels. const unsigned int nThreads = warp_size * 4; const unsigned int nBlocks = (n_tracks + nThreads - 1) / nThreads; // Fill the keys and param_ids buffers. - fill_fitting_sort_keys(nBlocks, nThreads, stream, - track_candidates_view.tracks, keys_buffer, - param_ids_buffer); + fill_fitting_sort_keys(nBlocks, nThreads, stream, track_fit_view.tracks, + keys_buffer, param_ids_buffer); // Sort the key to get the sorted parameter ids vecmem::device_vector keys_device(keys_buffer); @@ -133,15 +131,6 @@ kalman_fitting( .on(stream), keys_device.begin(), keys_device.end(), param_ids_device.begin()); - // Run the fitting, using the sorted parameter IDs. - fit_prelude(nBlocks, nThreads, 0, stream, param_ids_buffer, - track_candidates_view, - {track_states_buffer.tracks, track_states_buffer.states, - track_candidates_view.measurements}, - param_liveness_buffer); - TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); - str.synchronize(); - // Allocate the fitting kernels's payload in host memory. using fitter_t = traccc::details::kalman_fitter_t; device::fit_payload host_payload{ @@ -149,22 +138,101 @@ kalman_fitting( .field_data = field_view, .param_ids_view = param_ids_buffer, .param_liveness_view = param_liveness_buffer, - .tracks_view = {track_states_buffer.tracks, track_states_buffer.states, - track_candidates_view.measurements}, + .tracks_view = {track_fit_buffer.tracks, track_fit_buffer.states, + measurements}, .barcodes_view = seqs_buffer}; for (std::size_t i = 0; i < config.n_iterations; ++i) { // Run the track fitting - fit_forward(nBlocks, nThreads, 0, stream, config, - host_payload); - TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); + if (i > 0 || forward_on_first_iteration) { + // Don't run the forward step on the first iteration, as the + // input tracks are already fit. + fit_forward(nBlocks, nThreads, 0, stream, config, + host_payload); + TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); + } fit_backward(nBlocks, nThreads, 0, stream, config, host_payload); TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); } // Return the fitted tracks. - return track_states_buffer; + return track_fit_buffer; +} + +/// Templated implementation of the CUDA track fitting algorithm for unfitted +/// tracks. +/// +/// @tparam detector_t The (device) detector type to use +/// @tparam bfield_t The magnetic field type to use +/// +/// @param[in] det_view A view of the detector geometry +/// @param[in] field_view A view of the magnetic field +/// @param[in] track_candidates_view All track candidates to fit +/// @param[in] config The fitting configuration +/// @param[in] mr Memory resource(s) to use +/// @param[in] copy The copy object to use for memory transfers +/// @param[in] queue The Alpaka queue to use for execution +/// +/// @return A container of the fitted track states +/// +template +typename edm::track_fit_container::buffer +kalman_fitting( + const typename detector_t::const_view_type& det_view, + const bfield_t& field_view, + const typename edm::track_candidate_container< + typename detector_t::algebra_type>::const_view& track_candidates_view, + const fitting_config& config, const memory_resource& mr, vecmem::copy& copy, + stream& str, unsigned int warp_size) { + + // Get a convenience variable for the stream that we'll be using. + cudaStream_t stream = details::get_stream(str); + + // Get the number of tracks. + const edm::track_candidate_collection< + default_algebra>::const_device::size_type n_tracks = + copy.get_size(track_candidates_view.tracks); + + // Get the sizes of the track candidates in each track. + const std::vector candidate_sizes = + copy.get_sizes(track_candidates_view.tracks); + const unsigned int n_states = + std::accumulate(candidate_sizes.begin(), candidate_sizes.end(), 0u); + + // Create the result buffer. + typename edm::track_fit_container::buffer + track_states_buffer{ + {candidate_sizes, mr.main, mr.host, + vecmem::data::buffer_type::resizable}, + {n_states, mr.main, vecmem::data::buffer_type::resizable}}; + copy.setup(track_states_buffer.tracks)->ignore(); + copy.setup(track_states_buffer.states)->ignore(); + + // Return early, if there are no tracks. + if (n_tracks == 0) { + return track_states_buffer; + } + + // Launch parameters for all the kernels. + const unsigned int nThreads = warp_size * 4; + const unsigned int nBlocks = (n_tracks + nThreads - 1) / nThreads; + + // Run the fitting, using the sorted parameter IDs. + fit_prelude(nBlocks, nThreads, 0, stream, track_candidates_view, + {track_states_buffer.tracks, track_states_buffer.states, + track_candidates_view.measurements}); + TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); + str.synchronize(); + + return kalman_fitting( + det_view, field_view, + typename edm::track_fit_container< + typename detector_t::algebra_type>::buffer{ + std::move(track_states_buffer.tracks), + std::move(track_states_buffer.states)}, + track_candidates_view.measurements, config, mr, copy, str, warp_size, + true); } } // namespace traccc::cuda::details diff --git a/device/cuda/src/fitting/kalman_fitting_algorithm.cu b/device/cuda/src/fitting/kalman_fitting_algorithm.cu new file mode 100644 index 0000000000..2c9f4cdd8d --- /dev/null +++ b/device/cuda/src/fitting/kalman_fitting_algorithm.cu @@ -0,0 +1,53 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2022-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Local include(s). +#include "../utils/magnetic_field_types.hpp" +#include "kalman_fitting.cuh" +#include "traccc/cuda/fitting/kalman_fitting_algorithm.hpp" +#include "traccc/geometry/detector.hpp" +#include "traccc/utils/detector_buffer_bfield_visitor.hpp" + +namespace traccc::cuda { + +kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()( + const detector_buffer& det, const magnetic_field& field, + const edm::track_candidate_container::const_view& + track_candidates) const { + + // Run the track fitting. + return detector_buffer_magnetic_field_visitor< + detector_type_list, cuda::bfield_type_list>( + det, field, + [&]( + const typename detector_t::view& detector, + const bfield_view_t& bfield) { + return details::kalman_fitting( + detector, bfield, track_candidates, m_config, m_mr, + m_copy.get(), m_stream, m_warp_size); + }); +} + +kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()( + const detector_buffer& det, const magnetic_field& field, + edm::track_fit_container::buffer&& track_candidates, + const measurement_collection_types::const_view& measurements) const { + + // Run the track fitting. + return detector_buffer_magnetic_field_visitor< + detector_type_list, cuda::bfield_type_list>( + det, field, + [&]( + const typename detector_t::view& detector, + const bfield_view_t& bfield) { + return details::kalman_fitting( + detector, bfield, std::move(track_candidates), measurements, + m_config, m_mr, m_copy.get(), m_stream, m_warp_size); + }); +} + +} // namespace traccc::cuda diff --git a/device/cuda/src/fitting/kalman_fitting_algorithm_default_detector.cu b/device/cuda/src/fitting/kalman_fitting_algorithm_default_detector.cu deleted file mode 100644 index 2ce2e4dc16..0000000000 --- a/device/cuda/src/fitting/kalman_fitting_algorithm_default_detector.cu +++ /dev/null @@ -1,32 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2022-2025 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// Local include(s). -#include "../utils/magnetic_field_types.hpp" -#include "kalman_fitting.cuh" -#include "traccc/cuda/fitting/kalman_fitting_algorithm.hpp" - -// Project include(s). -#include "traccc/bfield/magnetic_field_types.hpp" - -namespace traccc::cuda { - -kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()( - const default_detector::view& det, const magnetic_field& bfield, - const edm::track_candidate_container::const_view& - track_candidates) const { - - // Run the track fitting. - return magnetic_field_visitor>( - bfield, [&](const bfield_view_t& bfield_view) { - return details::kalman_fitting( - det, bfield_view, track_candidates, m_config, m_mr, - m_copy.get(), m_stream, m_warp_size); - }); -} - -} // namespace traccc::cuda diff --git a/device/cuda/src/fitting/kalman_fitting_algorithm_telescope_detector.cu b/device/cuda/src/fitting/kalman_fitting_algorithm_telescope_detector.cu deleted file mode 100644 index 83433ef64b..0000000000 --- a/device/cuda/src/fitting/kalman_fitting_algorithm_telescope_detector.cu +++ /dev/null @@ -1,32 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2022-2025 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// Local include(s). -#include "../utils/magnetic_field_types.hpp" -#include "kalman_fitting.cuh" -#include "traccc/cuda/fitting/kalman_fitting_algorithm.hpp" - -// Project include(s). -#include "traccc/bfield/magnetic_field_types.hpp" - -namespace traccc::cuda { - -kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()( - const telescope_detector::view& det, const magnetic_field& bfield, - const edm::track_candidate_container::const_view& - track_candidates) const { - - // Run the track fitting. - return magnetic_field_visitor>( - bfield, [&](const bfield_view_t& bfield_view) { - return details::kalman_fitting( - det, bfield_view, track_candidates, m_config, m_mr, - m_copy.get(), m_stream, m_warp_size); - }); -} - -} // namespace traccc::cuda diff --git a/device/cuda/src/fitting/kernels/fill_fitting_sort_keys.cu b/device/cuda/src/fitting/kernels/fill_fitting_sort_keys.cu index 758b6bbab1..c8a0cf9cf6 100644 --- a/device/cuda/src/fitting/kernels/fill_fitting_sort_keys.cu +++ b/device/cuda/src/fitting/kernels/fill_fitting_sort_keys.cu @@ -17,26 +17,24 @@ namespace traccc::cuda { namespace kernels { __global__ void fill_fitting_sort_keys( - edm::track_candidate_collection::const_view - track_candidates_view, + edm::track_fit_collection::const_view track_fit_view, vecmem::data::vector_view keys_view, vecmem::data::vector_view ids_view) { - device::fill_fitting_sort_keys(details::global_index1(), - track_candidates_view, keys_view, ids_view); + device::fill_fitting_sort_keys(details::global_index1(), track_fit_view, + keys_view, ids_view); } } // namespace kernels void fill_fitting_sort_keys( const dim3& grid_size, const dim3& block_size, cudaStream_t stream, - edm::track_candidate_collection::const_view - track_candidates_view, + edm::track_fit_collection::const_view track_fit_view, vecmem::data::vector_view keys_view, vecmem::data::vector_view ids_view) { kernels::fill_fitting_sort_keys<<>>( - track_candidates_view, keys_view, ids_view); + track_fit_view, keys_view, ids_view); TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); } diff --git a/device/cuda/src/fitting/kernels/fill_fitting_sort_keys.hpp b/device/cuda/src/fitting/kernels/fill_fitting_sort_keys.hpp index 82d6a52d50..519cb1599a 100644 --- a/device/cuda/src/fitting/kernels/fill_fitting_sort_keys.hpp +++ b/device/cuda/src/fitting/kernels/fill_fitting_sort_keys.hpp @@ -9,7 +9,7 @@ // Project include(s). #include "traccc/edm/device/sort_key.hpp" -#include "traccc/edm/track_candidate_collection.hpp" +#include "traccc/edm/track_fit_collection.hpp" // CUDA include(s). #include @@ -22,8 +22,7 @@ namespace traccc::cuda { /// Function calling a kernel for @c traccc::device::fill_fitting_sort_keys void fill_fitting_sort_keys( const dim3& grid_size, const dim3& block_size, cudaStream_t stream, - edm::track_candidate_collection::const_view - track_candidates_view, + edm::track_fit_collection::const_view track_fit_view, vecmem::data::vector_view keys_view, vecmem::data::vector_view ids_view); diff --git a/device/cuda/src/fitting/kernels/fit_prelude.cu b/device/cuda/src/fitting/kernels/fit_prelude.cu index ef8e63bcd9..4209e3ea51 100644 --- a/device/cuda/src/fitting/kernels/fit_prelude.cu +++ b/device/cuda/src/fitting/kernels/fit_prelude.cu @@ -12,26 +12,20 @@ namespace traccc::cuda { namespace kernels { __global__ void fit_prelude( - vecmem::data::vector_view param_ids_view, edm::track_candidate_container::const_view track_candidates_view, - edm::track_fit_container::view tracks_view, - vecmem::data::vector_view param_liveness_view) { + edm::track_fit_container::view tracks_view) { device::fit_prelude(details::global_index1(), - param_ids_view, track_candidates_view, - tracks_view, param_liveness_view); + track_candidates_view, tracks_view); } } // namespace kernels void fit_prelude(const dim3& grid_size, const dim3& block_size, std::size_t shared_mem_size, const cudaStream_t& stream, - vecmem::data::vector_view param_ids_view, edm::track_candidate_container::const_view track_candidates_view, - edm::track_fit_container::view tracks_view, - vecmem::data::vector_view param_liveness_view) { + edm::track_fit_container::view tracks_view) { kernels::fit_prelude<<>>( - param_ids_view, track_candidates_view, tracks_view, - param_liveness_view); + track_candidates_view, tracks_view); } } // namespace traccc::cuda diff --git a/device/cuda/src/fitting/kernels/fit_prelude.hpp b/device/cuda/src/fitting/kernels/fit_prelude.hpp index 6f4098032c..0057bac72e 100644 --- a/device/cuda/src/fitting/kernels/fit_prelude.hpp +++ b/device/cuda/src/fitting/kernels/fit_prelude.hpp @@ -15,9 +15,7 @@ namespace traccc::cuda { void fit_prelude(const dim3& grid_size, const dim3& block_size, std::size_t shared_mem_size, const cudaStream_t& stream, - vecmem::data::vector_view param_ids_view, edm::track_candidate_container::const_view track_candidates_view, - edm::track_fit_container::view tracks_view, - vecmem::data::vector_view param_liveness_view); + edm::track_fit_container::view tracks_view); } diff --git a/device/cuda/src/seeding/spacepoint_formation_algorithm.cu b/device/cuda/src/seeding/spacepoint_formation_algorithm.cu index 582b8ff564..99edfe9a1d 100644 --- a/device/cuda/src/seeding/spacepoint_formation_algorithm.cu +++ b/device/cuda/src/seeding/spacepoint_formation_algorithm.cu @@ -20,9 +20,11 @@ namespace kernels { template __global__ void __launch_bounds__(1024, 1) - form_spacepoints(typename detector_t::const_view_type det_view, + form_spacepoints(typename detector_t::view det_view, measurement_collection_types::const_view measurements_view, - edm::spacepoint_collection::view spacepoints_view) { + edm::spacepoint_collection::view spacepoints_view) + requires(traccc::is_detector_traits) +{ device::form_spacepoints(details::global_index1(), det_view, measurements_view, spacepoints_view); @@ -30,16 +32,13 @@ __global__ void __launch_bounds__(1024, 1) } // namespace kernels -template -spacepoint_formation_algorithm::spacepoint_formation_algorithm( +spacepoint_formation_algorithm::spacepoint_formation_algorithm( const traccc::memory_resource& mr, vecmem::copy& copy, stream& str, std::unique_ptr logger) : messaging(std::move(logger)), m_mr(mr), m_copy(copy), m_stream(str) {} -template -edm::spacepoint_collection::buffer -spacepoint_formation_algorithm::operator()( - const typename detector_t::const_view_type& det_view, +edm::spacepoint_collection::buffer spacepoint_formation_algorithm::operator()( + const detector_buffer& detector, const measurement_collection_types::const_view& measurements_view) const { // Get the number of measurements. @@ -64,17 +63,17 @@ spacepoint_formation_algorithm::operator()( const unsigned int nBlocks = (num_measurements + blockSize - 1) / blockSize; // Launch the spacepoint formation kernel. - kernels::form_spacepoints<<>>( - det_view, measurements_view, spacepoints); + detector_buffer_visitor( + detector, [&]( + const typename detector_traits_t::view& det) { + kernels::form_spacepoints + <<>>(det, measurements_view, + spacepoints); + }); + TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); // Return the reconstructed spacepoints. return spacepoints; } - -// Explicit template instantiation -template class spacepoint_formation_algorithm; -template class spacepoint_formation_algorithm; -template class spacepoint_formation_algorithm; - } // namespace traccc::cuda diff --git a/device/sycl/CMakeLists.txt b/device/sycl/CMakeLists.txt index 17e554dc1a..d96f06110a 100644 --- a/device/sycl/CMakeLists.txt +++ b/device/sycl/CMakeLists.txt @@ -21,8 +21,7 @@ traccc_add_library( traccc_sycl sycl TYPE SHARED # Seeding algorithm(s). "include/traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" "src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp" - "src/seeding/silicon_pixel_spacepoint_formation_algorithm_default_detector.sycl" - "src/seeding/silicon_pixel_spacepoint_formation_algorithm_telescope_detector.sycl" + "src/seeding/silicon_pixel_spacepoint_formation_algorithm.sycl" "src/seeding/silicon_pixel_spacepoint_formation.hpp" "include/traccc/sycl/seeding/details/spacepoint_binning.hpp" "src/seeding/spacepoint_binning.sycl" @@ -35,14 +34,12 @@ traccc_add_library( traccc_sycl sycl TYPE SHARED # Track finding algorithm(s). "include/traccc/sycl/finding/combinatorial_kalman_filter_algorithm.hpp" "src/finding/combinatorial_kalman_filter_algorithm.cpp" - "src/finding/combinatorial_kalman_filter_algorithm_default_detector.sycl" - "src/finding/combinatorial_kalman_filter_algorithm_telescope_detector.sycl" + "src/finding/combinatorial_kalman_filter_algorithm.sycl" "src/finding/combinatorial_kalman_filter.hpp" # Track fitting algorithm(s). "include/traccc/sycl/fitting/kalman_fitting_algorithm.hpp" "src/fitting/kalman_fitting_algorithm.cpp" - "src/fitting/kalman_fitting_algorithm_default_detector.sycl" - "src/fitting/kalman_fitting_algorithm_telescope_detector.sycl" + "src/fitting/kalman_fitting_algorithm.sycl" "src/fitting/kalman_fitting.hpp" # Utilities. "include/traccc/sycl/utils/make_prefix_sum_buff.hpp" diff --git a/device/sycl/include/traccc/sycl/finding/combinatorial_kalman_filter_algorithm.hpp b/device/sycl/include/traccc/sycl/finding/combinatorial_kalman_filter_algorithm.hpp index b408552187..567a7728b7 100644 --- a/device/sycl/include/traccc/sycl/finding/combinatorial_kalman_filter_algorithm.hpp +++ b/device/sycl/include/traccc/sycl/finding/combinatorial_kalman_filter_algorithm.hpp @@ -17,6 +17,7 @@ #include "traccc/edm/track_parameters.hpp" #include "traccc/finding/finding_config.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/detector_buffer.hpp" #include "traccc/utils/algorithm.hpp" #include "traccc/utils/memory_resource.hpp" #include "traccc/utils/messaging.hpp" @@ -32,11 +33,7 @@ namespace traccc::sycl { /// CKF track finding algorithm class combinatorial_kalman_filter_algorithm : public algorithm::buffer( - const default_detector::view&, const magnetic_field&, - const measurement_collection_types::const_view&, - const bound_track_parameters_collection_types::const_view&)>, - public algorithm::buffer( - const telescope_detector::view&, const magnetic_field&, + const detector_buffer&, const magnetic_field&, const measurement_collection_types::const_view&, const bound_track_parameters_collection_types::const_view&)>, public messaging { @@ -56,23 +53,7 @@ class combinatorial_kalman_filter_algorithm /// Execute the algorithm /// - /// @param det The (default) detector object - /// @param bfield The magnetic field object - /// @param measurements All measurements in an event - /// @param seeds All seeds in an event to start the track finding - /// with - /// - /// @return A container of the found track candidates - /// - output_type operator()( - const default_detector::view& det, const magnetic_field& bfield, - const measurement_collection_types::const_view& measurements, - const bound_track_parameters_collection_types::const_view& seeds) - const override; - - /// Execute the algorithm - /// - /// @param det The (telescope) detector object + /// @param det The detector object /// @param bfield The magnetic field object /// @param measurements All measurements in an event /// @param seeds All seeds in an event to start the track finding @@ -81,7 +62,7 @@ class combinatorial_kalman_filter_algorithm /// @return A container of the found track candidates /// output_type operator()( - const telescope_detector::view& det, const magnetic_field& field, + const detector_buffer& det, const magnetic_field& bfield, const measurement_collection_types::const_view& measurements, const bound_track_parameters_collection_types::const_view& seeds) const override; diff --git a/device/sycl/include/traccc/sycl/fitting/kalman_fitting_algorithm.hpp b/device/sycl/include/traccc/sycl/fitting/kalman_fitting_algorithm.hpp index 2612a9a274..4932844ba1 100644 --- a/device/sycl/include/traccc/sycl/fitting/kalman_fitting_algorithm.hpp +++ b/device/sycl/include/traccc/sycl/fitting/kalman_fitting_algorithm.hpp @@ -16,6 +16,7 @@ #include "traccc/edm/track_fit_container.hpp" #include "traccc/fitting/fitting_config.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/detector_buffer.hpp" #include "traccc/utils/algorithm.hpp" #include "traccc/utils/memory_resource.hpp" #include "traccc/utils/messaging.hpp" @@ -31,11 +32,12 @@ namespace traccc::sycl { /// Kalman filter based track fitting algorithm class kalman_fitting_algorithm : public algorithm::buffer( - const default_detector::view&, const magnetic_field&, + const detector_buffer&, const magnetic_field&, const edm::track_candidate_container::const_view&)>, public algorithm::buffer( - const telescope_detector::view&, const magnetic_field&, - const edm::track_candidate_container::const_view&)>, + const detector_buffer&, const magnetic_field&, + edm::track_fit_container::buffer&&, + const measurement_collection_types::const_view&)>, public messaging { public: @@ -53,31 +55,31 @@ class kalman_fitting_algorithm vecmem::copy& copy, queue_wrapper queue, std::unique_ptr logger = getDummyLogger().clone()); - /// Execute the algorithm + /// Execute the algorithm with unfitted tracks /// - /// @param det The (default) detector object + /// @param det The detector object /// @param bfield The magnetic field object /// @param track_candidates All track candidates to fit /// /// @return A container of the fitted track states /// output_type operator()( - const default_detector::view& det, const magnetic_field& bfield, + const detector_buffer& det, const magnetic_field& bfield, const edm::track_candidate_container::const_view& track_candidates) const override; - /// Execute the algorithm + /// Execute the algorithm with fitted tracks /// - /// @param det The (telescope) detector object + /// @param det The detector object /// @param bfield The magnetic field object /// @param track_candidates All track candidates to fit /// /// @return A container of the fitted track states /// output_type operator()( - const telescope_detector::view& det, const magnetic_field& bfield, - const edm::track_candidate_container::const_view& - track_candidates) const override; + const detector_buffer& det, const magnetic_field& bfield, + edm::track_fit_container::buffer&& track_states, + const measurement_collection_types::const_view&) const override; private: /// Algorithm configuration diff --git a/device/sycl/include/traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp b/device/sycl/include/traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp index 6991153e2b..f085d93620 100644 --- a/device/sycl/include/traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp +++ b/device/sycl/include/traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp @@ -11,6 +11,7 @@ #include "traccc/edm/measurement.hpp" #include "traccc/edm/spacepoint_collection.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/detector_buffer.hpp" #include "traccc/sycl/utils/queue_wrapper.hpp" #include "traccc/utils/algorithm.hpp" #include "traccc/utils/memory_resource.hpp" @@ -31,10 +32,7 @@ namespace traccc::sycl { /// class silicon_pixel_spacepoint_formation_algorithm : public algorithm, - public algorithm, public messaging { @@ -58,18 +56,7 @@ class silicon_pixel_spacepoint_formation_algorithm /// @return A spacepoint buffer, with one spacepoint for every /// silicon pixel measurement /// - output_type operator()(const default_detector::view& det, - const measurement_collection_types::const_view& - measurements) const override; - - /// Construct spacepoints from 2D silicon pixel measurements - /// - /// @param det Detector object - /// @param measurements A collection of measurements - /// @return A spacepoint buffer, with one spacepoint for every - /// silicon pixel measurement - /// - output_type operator()(const telescope_detector::view& det, + output_type operator()(const detector_buffer& det, const measurement_collection_types::const_view& measurements) const override; diff --git a/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_default_detector.sycl b/device/sycl/src/finding/combinatorial_kalman_filter_algorithm.sycl similarity index 56% rename from device/sycl/src/finding/combinatorial_kalman_filter_algorithm_default_detector.sycl rename to device/sycl/src/finding/combinatorial_kalman_filter_algorithm.sycl index 92df2d8a54..38e9b440bd 100644 --- a/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_default_detector.sycl +++ b/device/sycl/src/finding/combinatorial_kalman_filter_algorithm.sycl @@ -6,6 +6,7 @@ */ // Local include(s). +#include "../utils/detector_types.hpp" #include "../utils/get_queue.hpp" #include "../utils/magnetic_field_types.hpp" #include "combinatorial_kalman_filter.hpp" @@ -13,6 +14,7 @@ // Project include(s). #include "traccc/bfield/magnetic_field_types.hpp" +#include "traccc/utils/detector_buffer_bfield_visitor.hpp" #include "traccc/utils/propagation.hpp" // System include(s). @@ -20,25 +22,29 @@ namespace traccc::sycl { namespace kernels { -template -struct ckf_default_detector; +template +struct combinatorial_kalman_filter; } // namespace kernels combinatorial_kalman_filter_algorithm::output_type combinatorial_kalman_filter_algorithm::operator()( - const default_detector::view& det, const magnetic_field& bfield, + const detector_buffer& detector, const magnetic_field& field, const measurement_collection_types::const_view& measurements, const bound_track_parameters_collection_types::const_view& seeds) const { // Perform the track finding using the templated implementation. - return magnetic_field_visitor>( - bfield, [&](const bfield_view_t& bfield_view) { + return detector_buffer_magnetic_field_visitor< + detector_type_list, sycl::bfield_type_list>( + detector, field, + [&]( + const typename detector_t::view& det, const bfield_view_t& bfield) { return details::combinatorial_kalman_filter< - kernels::ckf_default_detector< + kernels::combinatorial_kalman_filter< + detector_tag_selector_t, bfield_tag_selector_t>, - default_detector::device>(det, bfield_view, measurements, seeds, - m_config, m_mr, m_copy, - details::get_queue(m_queue)); + typename detector_t::device>(det, bfield, measurements, seeds, + m_config, m_mr, m_copy, + details::get_queue(m_queue)); }); } diff --git a/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_telescope_detector.sycl b/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_telescope_detector.sycl deleted file mode 100644 index 03cd109b58..0000000000 --- a/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_telescope_detector.sycl +++ /dev/null @@ -1,45 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2024-2025 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// Local include(s). -#include "../utils/get_queue.hpp" -#include "../utils/magnetic_field_types.hpp" -#include "combinatorial_kalman_filter.hpp" -#include "traccc/sycl/finding/combinatorial_kalman_filter_algorithm.hpp" - -// Project include(s). -#include "traccc/bfield/magnetic_field_types.hpp" -#include "traccc/utils/propagation.hpp" - -// System include(s). -#include - -namespace traccc::sycl { -namespace kernels { -template -struct ckf_telescope_detector; -} // namespace kernels - -combinatorial_kalman_filter_algorithm::output_type -combinatorial_kalman_filter_algorithm::operator()( - const telescope_detector::view& det, const magnetic_field& bfield, - const measurement_collection_types::const_view& measurements, - const bound_track_parameters_collection_types::const_view& seeds) const { - - // Perform the track finding using the templated implementation. - return magnetic_field_visitor>( - bfield, [&](const bfield_view_t& bfield_view) { - return details::combinatorial_kalman_filter< - kernels::ckf_telescope_detector< - bfield_tag_selector_t>, - telescope_detector::device>(det, bfield_view, measurements, - seeds, m_config, m_mr, m_copy, - details::get_queue(m_queue)); - }); -} - -} // namespace traccc::sycl diff --git a/device/sycl/src/fitting/kalman_fitting.hpp b/device/sycl/src/fitting/kalman_fitting.hpp index 22c4a011b0..d9ad1b27df 100644 --- a/device/sycl/src/fitting/kalman_fitting.hpp +++ b/device/sycl/src/fitting/kalman_fitting.hpp @@ -50,7 +50,8 @@ struct fit_backward; namespace details { -/// Templated implementation of the SYCL track fitting algorithm. +/// Templated implementation of the SYCL track fitting algorithm with fitted +/// tracks. /// /// @tparam kernel_t Structure to generate unique kernel names with /// @tparam detector_t The (device) detector type to use @@ -71,19 +72,25 @@ typename edm::track_fit_container::buffer kalman_fitting( const typename detector_t::const_view_type& det_view, const bfield_t& field_view, - const typename edm::track_candidate_container< - typename detector_t::algebra_type>::const_view& track_candidates_view, + typename edm::track_fit_container< + typename detector_t::algebra_type>::buffer&& track_fit_buffer, + const measurement_collection_types::const_view& measurements, const fitting_config& config, const memory_resource& mr, vecmem::copy& copy, - ::sycl::queue& queue) { + ::sycl::queue& queue, bool forward_on_first_iteration = false) { + + typename edm::track_fit_container< + typename detector_t::algebra_type>::const_view track_fit_view{ + vecmem::get_data(track_fit_buffer.tracks), + vecmem::get_data(track_fit_buffer.states), measurements}; // Get the number of tracks. const edm::track_candidate_collection< default_algebra>::const_device::size_type n_tracks = - copy.get_size(track_candidates_view.tracks); + copy.get_size(track_fit_view.tracks); // Get the sizes of the track candidates in each track. const std::vector candidate_sizes = - copy.get_sizes(track_candidates_view.tracks); + copy.get_sizes(track_fit_view.tracks); const unsigned int n_states = std::accumulate(candidate_sizes.begin(), candidate_sizes.end(), 0u); @@ -130,6 +137,7 @@ kalman_fitting( keys_setup_event->wait(); param_ids_setup_event->wait(); param_liveness_setup_event->wait(); + copy.memset(param_liveness_buffer, 1)->ignore(); // The execution range for the two kernels of the function. static constexpr unsigned int localSize = 64; @@ -138,13 +146,12 @@ kalman_fitting( // Fill the keys and param_ids buffers. ::sycl::event fill_keys_event = queue.submit([&](::sycl::handler& h) { h.parallel_for>( - range, - [track_candidates_view, keys_view = vecmem::get_data(keys_buffer), - param_ids_view = - vecmem::get_data(param_ids_buffer)](::sycl::nd_item<1> item) { + range, [track_fit_view, keys_view = vecmem::get_data(keys_buffer), + param_ids_view = vecmem::get_data(param_ids_buffer)]( + ::sycl::nd_item<1> item) { device::fill_fitting_sort_keys(details::global_index(item), - track_candidates_view.tracks, - keys_view, param_ids_view); + track_fit_view.tracks, keys_view, + param_ids_view); }); }); @@ -156,48 +163,28 @@ kalman_fitting( keys_device.begin(), keys_device.end(), param_ids_device.begin()); - // Run the fitting, using the sorted parameter IDs. - typename edm::track_fit_container::view - track_states_view{track_states_buffer.tracks, - track_states_buffer.states, - track_candidates_view.measurements}; - tracks_setup_event->wait(); - track_states_setup_event->wait(); - - queue - .submit([&](::sycl::handler& h) { - h.parallel_for>( - range, [param_ids_view = vecmem::get_data(param_ids_buffer), - track_candidates_view, track_states_view, - param_liveness_view = vecmem::get_data( - param_liveness_buffer)](::sycl::nd_item<1> item) { - device::fit_prelude( - details::global_index(item), param_ids_view, - track_candidates_view, track_states_view, - param_liveness_view); - }); - }) - .wait_and_throw(); - using fitter_t = traccc::details::kalman_fitter_t; device::fit_payload payload{ .det_data = det_view, .field_data = field_view, .param_ids_view = param_ids_buffer, .param_liveness_view = param_liveness_buffer, - .tracks_view = track_states_view, + .tracks_view = {track_fit_buffer.tracks, track_fit_buffer.states, + measurements}, .barcodes_view = seqs_buffer}; for (std::size_t i = 0; i < config.n_iterations; ++i) { - queue - .submit([&](::sycl::handler& h) { - h.parallel_for>( - range, [config, payload](::sycl::nd_item<1> item) { - device::fit_forward( - details::global_index(item), config, payload); - }); - }) - .wait_and_throw(); + if (i == 0 || forward_on_first_iteration) { + queue + .submit([&](::sycl::handler& h) { + h.parallel_for>( + range, [config, payload](::sycl::nd_item<1> item) { + device::fit_forward( + details::global_index(item), config, payload); + }); + }) + .wait_and_throw(); + } queue .submit([&](::sycl::handler& h) { @@ -214,5 +201,104 @@ kalman_fitting( return track_states_buffer; } +/// Templated implementation of the SYCL track fitting algorithm with unfitted +/// tracks. +/// +/// @tparam kernel_t Structure to generate unique kernel names with +/// @tparam detector_t The (device) detector type to use +/// @tparam bfield_t The magnetic field type to use +/// +/// @param[in] det_view A view of the detector geometry +/// @param[in] field_view A view of the magnetic field +/// @param[in] track_candidates_view All track candidates to fit +/// @param[in] config The fitting configuration +/// @param[in] mr Memory resource(s) to use +/// @param[in] copy The copy object to use for memory transfers +/// @param[in] queue The SYCL queue to use for execution +/// +/// @return A container of the fitted track states +/// +template +typename edm::track_fit_container::buffer +kalman_fitting( + const typename detector_t::const_view_type& det_view, + const bfield_t& field_view, + const typename edm::track_candidate_container< + typename detector_t::algebra_type>::const_view& track_candidates_view, + const fitting_config& config, const memory_resource& mr, vecmem::copy& copy, + ::sycl::queue& queue) { + + // Get the number of tracks. + const edm::track_candidate_collection< + default_algebra>::const_device::size_type n_tracks = + copy.get_size(track_candidates_view.tracks); + + // Get the sizes of the track candidates in each track. + const std::vector candidate_sizes = + copy.get_sizes(track_candidates_view.tracks); + const unsigned int n_states = + std::accumulate(candidate_sizes.begin(), candidate_sizes.end(), 0u); + + // Create the result buffer. + typename edm::track_fit_container::buffer + track_states_buffer{ + {candidate_sizes, mr.main, mr.host, + vecmem::data::buffer_type::resizable}, + {n_states, mr.main, vecmem::data::buffer_type::resizable}}; + vecmem::copy::event_type tracks_setup_event = + copy.setup(track_states_buffer.tracks); + vecmem::copy::event_type track_states_setup_event = + copy.setup(track_states_buffer.states); + + // Return early, if there are no tracks. + if (n_tracks == 0) { + tracks_setup_event->wait(); + track_states_setup_event->wait(); + return track_states_buffer; + } + + std::vector seqs_sizes(candidate_sizes.size()); + std::transform(candidate_sizes.begin(), candidate_sizes.end(), + seqs_sizes.begin(), [&config](const unsigned int sz) { + return std::max(sz * config.barcode_sequence_size_factor, + config.min_barcode_sequence_capacity); + }); + vecmem::data::jagged_vector_buffer seqs_buffer{ + seqs_sizes, mr.main, mr.host, vecmem::data::buffer_type::resizable}; + copy.setup(seqs_buffer)->wait(); + + // The execution range for the two kernels of the function. + static constexpr unsigned int localSize = 64; + ::sycl::nd_range<1> range = calculate1DimNdRange(n_tracks, localSize); + + // Run the fitting, using the sorted parameter IDs. + typename edm::track_fit_container::view + track_states_view{track_states_buffer.tracks, + track_states_buffer.states, + track_candidates_view.measurements}; + tracks_setup_event->wait(); + track_states_setup_event->wait(); + + queue + .submit([&](::sycl::handler& h) { + h.parallel_for>( + range, [track_candidates_view, + track_states_view](::sycl::nd_item<1> item) { + device::fit_prelude( + details::global_index(item), track_candidates_view, + track_states_view); + }); + }) + .wait_and_throw(); + + return kalman_fitting( + det_view, field_view, + typename edm::track_fit_container< + typename detector_t::algebra_type>::buffer{ + std::move(track_states_buffer.tracks), + std::move(track_states_buffer.states)}, + track_candidates_view.measurements, config, mr, copy, queue, true); +} + } // namespace details } // namespace traccc::sycl diff --git a/device/sycl/src/fitting/kalman_fitting_algorithm.sycl b/device/sycl/src/fitting/kalman_fitting_algorithm.sycl new file mode 100644 index 0000000000..ad59bb390c --- /dev/null +++ b/device/sycl/src/fitting/kalman_fitting_algorithm.sycl @@ -0,0 +1,67 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2022-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Local include(s). +#include "../utils/detector_types.hpp" +#include "../utils/get_queue.hpp" +#include "../utils/magnetic_field_types.hpp" +#include "kalman_fitting.hpp" +#include "traccc/sycl/fitting/kalman_fitting_algorithm.hpp" + +// Project include(s). +#include "traccc/bfield/magnetic_field_types.hpp" +#include "traccc/utils/detector_buffer_bfield_visitor.hpp" + +namespace traccc::sycl { +namespace kernels { +template +struct fit_tracks; +} // namespace kernels + +kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()( + const detector_buffer& detector, const magnetic_field& bfield, + const edm::track_candidate_container::const_view& + track_candidates) const { + + // Perform the track fitting using the templated implementation. + return detector_buffer_magnetic_field_visitor< + detector_type_list, sycl::bfield_type_list>( + detector, bfield, + [&]( + const typename detector_t::view& det, const bfield_view_t& field) { + return details::kalman_fitting< + kernels::fit_tracks< + detector_tag_selector_t, + bfield_tag_selector_t>, + typename detector_t::device>(det, field, track_candidates, + m_config, m_mr, m_copy.get(), + details::get_queue(m_queue)); + }); +} + +kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()( + const detector_buffer& detector, const magnetic_field& bfield, + edm::track_fit_container::buffer&& track_states, + const measurement_collection_types::const_view& measurements) const { + + // Perform the track fitting using the templated implementation. + return detector_buffer_magnetic_field_visitor< + detector_type_list, sycl::bfield_type_list>( + detector, bfield, + [&]( + const typename detector_t::view& det, const bfield_view_t& field) { + return details::kalman_fitting< + kernels::fit_tracks< + detector_tag_selector_t, + bfield_tag_selector_t>, + typename detector_t::device>( + det, field, std::move(track_states), measurements, m_config, + m_mr, m_copy.get(), details::get_queue(m_queue)); + }); +} + +} // namespace traccc::sycl diff --git a/device/sycl/src/fitting/kalman_fitting_algorithm_default_detector.sycl b/device/sycl/src/fitting/kalman_fitting_algorithm_default_detector.sycl deleted file mode 100644 index 50cdf67f2f..0000000000 --- a/device/sycl/src/fitting/kalman_fitting_algorithm_default_detector.sycl +++ /dev/null @@ -1,40 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2022-2025 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// Local include(s). -#include "../utils/get_queue.hpp" -#include "../utils/magnetic_field_types.hpp" -#include "kalman_fitting.hpp" -#include "traccc/sycl/fitting/kalman_fitting_algorithm.hpp" - -// Project include(s). -#include "traccc/bfield/magnetic_field_types.hpp" - -namespace traccc::sycl { -namespace kernels { -template -struct fit_tracks_default_detector; -} // namespace kernels - -kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()( - const default_detector::view& det, const magnetic_field& bfield, - const edm::track_candidate_container::const_view& - track_candidates) const { - - // Perform the track fitting using the templated implementation. - return magnetic_field_visitor>( - bfield, [&](const bfield_view_t& bfield_view) { - return details::kalman_fitting< - kernels::fit_tracks_default_detector< - bfield_tag_selector_t>, - default_detector::device>(det, bfield_view, track_candidates, - m_config, m_mr, m_copy.get(), - details::get_queue(m_queue)); - }); -} - -} // namespace traccc::sycl diff --git a/device/sycl/src/fitting/kalman_fitting_algorithm_telescope_detector.sycl b/device/sycl/src/fitting/kalman_fitting_algorithm_telescope_detector.sycl deleted file mode 100644 index f612e14443..0000000000 --- a/device/sycl/src/fitting/kalman_fitting_algorithm_telescope_detector.sycl +++ /dev/null @@ -1,40 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2022-2025 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// Local include(s). -#include "../utils/get_queue.hpp" -#include "../utils/magnetic_field_types.hpp" -#include "kalman_fitting.hpp" -#include "traccc/sycl/fitting/kalman_fitting_algorithm.hpp" - -// Project include(s). -#include "traccc/bfield/magnetic_field_types.hpp" - -namespace traccc::sycl { -namespace kernels { -template -struct fit_tracks_telescope_detector; -} // namespace kernels - -kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()( - const telescope_detector::view& det, const magnetic_field& bfield, - const edm::track_candidate_container::const_view& - track_candidates) const { - - // Perform the track fitting using the templated implementation. - return magnetic_field_visitor>( - bfield, [&](const bfield_view_t& bfield_view) { - return details::kalman_fitting< - kernels::fit_tracks_telescope_detector< - bfield_tag_selector_t>, - telescope_detector::device>(det, bfield_view, track_candidates, - m_config, m_mr, m_copy.get(), - details::get_queue(m_queue)); - }); -} - -} // namespace traccc::sycl diff --git a/device/sycl/src/seeding/silicon_pixel_spacepoint_formation.hpp b/device/sycl/src/seeding/silicon_pixel_spacepoint_formation.hpp index 82b2a57e50..5726dc0c20 100644 --- a/device/sycl/src/seeding/silicon_pixel_spacepoint_formation.hpp +++ b/device/sycl/src/seeding/silicon_pixel_spacepoint_formation.hpp @@ -39,7 +39,7 @@ namespace traccc::sycl::details { /// template edm::spacepoint_collection::buffer silicon_pixel_spacepoint_formation( - const typename detector_t::const_view_type& det_view, + const typename detector_t::view& det_view, const measurement_collection_types::const_view& measurements_view, vecmem::memory_resource& mr, vecmem::copy& copy, ::sycl::queue& queue) { diff --git a/device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm_default_detector.sycl b/device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm.sycl similarity index 56% rename from device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm_default_detector.sycl rename to device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm.sycl index c0ecd08c28..6defe43405 100644 --- a/device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm_default_detector.sycl +++ b/device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm.sycl @@ -8,18 +8,23 @@ // Local include(s). #include "../utils/get_queue.hpp" #include "silicon_pixel_spacepoint_formation.hpp" +#include "traccc/geometry/detector_buffer.hpp" #include "traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" namespace traccc::sycl { silicon_pixel_spacepoint_formation_algorithm::output_type silicon_pixel_spacepoint_formation_algorithm::operator()( - const default_detector::view& det, + const detector_buffer& det, const measurement_collection_types::const_view& meas) const { - return details::silicon_pixel_spacepoint_formation< - default_detector::device>(det, meas, m_mr.main, m_copy, - details::get_queue(m_queue)); + return detector_buffer_visitor( + det, [&]( + const typename detector_traits_t::view& det_view) { + return details::silicon_pixel_spacepoint_formation< + detector_traits_t>(det_view, meas, m_mr.main, m_copy, + details::get_queue(m_queue)); + }); } } // namespace traccc::sycl diff --git a/device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm_telescope_detector.sycl b/device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm_telescope_detector.sycl deleted file mode 100644 index 746d24cada..0000000000 --- a/device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm_telescope_detector.sycl +++ /dev/null @@ -1,25 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2023-2024 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// Local include(s). -#include "../utils/get_queue.hpp" -#include "silicon_pixel_spacepoint_formation.hpp" -#include "traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" - -namespace traccc::sycl { - -silicon_pixel_spacepoint_formation_algorithm::output_type -silicon_pixel_spacepoint_formation_algorithm::operator()( - const telescope_detector::view& det, - const measurement_collection_types::const_view& meas) const { - - return details::silicon_pixel_spacepoint_formation< - telescope_detector::device>(det, meas, m_mr.main, m_copy, - details::get_queue(m_queue)); -} - -} // namespace traccc::sycl diff --git a/device/sycl/src/utils/detector_types.hpp b/device/sycl/src/utils/detector_types.hpp new file mode 100644 index 0000000000..babc4fd073 --- /dev/null +++ b/device/sycl/src/utils/detector_types.hpp @@ -0,0 +1,54 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +#include "traccc/definitions/primitives.hpp" +#include "traccc/geometry/detector.hpp" + +namespace traccc::sycl { + +/* + * SYCL requires a little bit of extra massaging to make the kernel tags + * work... + */ +struct default_detector_kernel_tag {}; +struct telescope_detector_kernel_tag {}; + +template +struct detector_tag_selector {}; + +template <> +struct detector_tag_selector { + using type = default_detector_kernel_tag; +}; + +template <> +struct detector_tag_selector { + using type = telescope_detector_kernel_tag; +}; + +template +using detector_tag_selector_t = typename detector_tag_selector::type; + +template +concept detector_tag_exists_for_backend = + requires { typename detector_tag_selector_t; }; + +template +struct detector_tag_existance_validator {}; + +template +struct detector_tag_existance_validator> { + static constexpr bool value = (detector_tag_exists_for_backend && ...); +}; + +static_assert( + detector_tag_existance_validator::value, + "Not all detector types registered for SYCL have an accompanying tag"); + +} // namespace traccc::sycl diff --git a/examples/run/alpaka/full_chain_algorithm.cpp b/examples/run/alpaka/full_chain_algorithm.cpp index d08445d5e7..22c505daae 100644 --- a/examples/run/alpaka/full_chain_algorithm.cpp +++ b/examples/run/alpaka/full_chain_algorithm.cpp @@ -23,7 +23,7 @@ full_chain_algorithm::full_chain_algorithm( const finding_algorithm::config_type& finding_config, const fitting_algorithm::config_type& fitting_config, const silicon_detector_description::host& det_descr, - const magnetic_field& field, host_detector_type* detector, + const magnetic_field& field, host_detector* detector, std::unique_ptr logger) : messaging(logger->clone()), m_queue(), @@ -76,11 +76,9 @@ full_chain_algorithm::full_chain_algorithm( .async_copy()(::vecmem::get_data(m_det_descr.get()), m_device_det_descr) ->ignore(); if (m_detector != nullptr) { - m_device_detector = detray::get_buffer(detray::get_data(*m_detector), - m_vecmem_objects.device_mr(), - m_vecmem_objects.async_copy()); - const auto& const_device_detector = m_device_detector; - m_device_detector_view = detray::get_data(const_device_detector); + m_device_detector = traccc::buffer_from_host_detector( + *m_detector, m_vecmem_objects.device_mr(), + m_vecmem_objects.async_copy()); } } @@ -137,11 +135,9 @@ full_chain_algorithm::full_chain_algorithm(const full_chain_algorithm& parent) .async_copy()(::vecmem::get_data(m_det_descr.get()), m_device_det_descr) ->ignore(); if (m_detector != nullptr) { - m_device_detector = detray::get_buffer(detray::get_data(*m_detector), - m_vecmem_objects.device_mr(), - m_vecmem_objects.async_copy()); - const auto& const_device_detector = m_device_detector; - m_device_detector_view = detray::get_data(const_device_detector); + m_device_detector = traccc::buffer_from_host_detector( + *m_detector, m_vecmem_objects.device_mr(), + m_vecmem_objects.async_copy()); } } @@ -166,18 +162,18 @@ full_chain_algorithm::output_type full_chain_algorithm::operator()( // Run the seed-finding (asynchronously). const spacepoint_formation_algorithm::output_type spacepoints = - m_spacepoint_formation(m_device_detector_view, measurements); + m_spacepoint_formation(m_device_detector, measurements); const track_params_estimation::output_type track_params = m_track_parameter_estimation(measurements, spacepoints, m_seeding(spacepoints), m_field_vec); // Run the track finding (asynchronously). - const finding_algorithm::output_type track_candidates = m_finding( - m_device_detector_view, m_field, measurements, track_params); + const finding_algorithm::output_type track_candidates = + m_finding(m_device_detector, m_field, measurements, track_params); // Run the track fitting (asynchronously). const fitting_algorithm::output_type track_states = m_fitting( - m_device_detector_view, m_field, {track_candidates, measurements}); + m_device_detector, m_field, {track_candidates, measurements}); // Copy a limited amount of result data back to the host. const auto host_tracks = m_vecmem_objects.async_copy().to( @@ -221,7 +217,7 @@ bound_track_parameters_collection_types::host full_chain_algorithm::seeding( // Run the seed-finding (asynchronously). const spacepoint_formation_algorithm::output_type spacepoints = - m_spacepoint_formation(m_device_detector_view, measurements); + m_spacepoint_formation(m_device_detector, measurements); const track_params_estimation::output_type track_params = m_track_parameter_estimation(measurements, spacepoints, m_seeding(spacepoints), m_field_vec); diff --git a/examples/run/alpaka/full_chain_algorithm.hpp b/examples/run/alpaka/full_chain_algorithm.hpp index 564b138a37..51668913a2 100644 --- a/examples/run/alpaka/full_chain_algorithm.hpp +++ b/examples/run/alpaka/full_chain_algorithm.hpp @@ -24,6 +24,8 @@ #include "traccc/edm/track_parameters.hpp" #include "traccc/fitting/kalman_filter/kalman_fitter.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/detector_buffer.hpp" +#include "traccc/geometry/host_detector.hpp" #include "traccc/geometry/silicon_detector_description.hpp" #include "traccc/utils/algorithm.hpp" #include "traccc/utils/messaging.hpp" @@ -52,15 +54,9 @@ class full_chain_algorithm /// @name Type declaration(s) /// @{ - /// (Host) Detector type used during track finding and fitting - using host_detector_type = traccc::default_detector::host; - /// (Device) Detector type used during track finding and fitting - using device_detector_type = traccc::default_detector::device; - /// Spacepoint formation algorithm type using spacepoint_formation_algorithm = - traccc::alpaka::spacepoint_formation_algorithm< - traccc::default_detector::device>; + traccc::alpaka::spacepoint_formation_algorithm; /// Clustering algorithm type using clustering_algorithm = traccc::alpaka::clusterization_algorithm; /// Track finding algorithm type @@ -84,8 +80,7 @@ class full_chain_algorithm const finding_algorithm::config_type& finding_config, const fitting_algorithm::config_type& fitting_config, const silicon_detector_description::host& det_descr, - const magnetic_field& field, - host_detector_type* detector, + const magnetic_field& field, host_detector* detector, std::unique_ptr logger); /// Copy constructor @@ -140,12 +135,11 @@ class full_chain_algorithm m_det_descr; /// Detector description buffer silicon_detector_description::buffer m_device_det_descr; + /// Host detector - host_detector_type* m_detector; + host_detector* m_detector; /// Buffer holding the detector's payload on the device - host_detector_type::buffer_type m_device_detector; - /// View of the detector's payload on the device - host_detector_type::const_view_type m_device_detector_view; + detector_buffer m_device_detector; /// @name Sub-algorithms used by this full-chain algorithm /// @{ diff --git a/examples/run/alpaka/seeding_example_alpaka.cpp b/examples/run/alpaka/seeding_example_alpaka.cpp index 70114bdc8e..b017d51984 100644 --- a/examples/run/alpaka/seeding_example_alpaka.cpp +++ b/examples/run/alpaka/seeding_example_alpaka.cpp @@ -123,21 +123,20 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, const auto field = traccc::details::make_magnetic_field(bfield_opts); const traccc::vector3 field_vec(seeding_opts); - // Construct a Detray detector object, if supported by the configuration. - traccc::default_detector::host host_det{mng_mr}; + // Detector view object + traccc::host_detector host_det; traccc::io::read_detector(host_det, mng_mr, detector_opts.detector_file, detector_opts.material_file, detector_opts.grid_file); - // Detector view object - const traccc::default_detector::host& const_host_det = host_det; - traccc::default_detector::view det_view = detray::get_data(const_host_det); - // Copy objects vecmem::copy host_copy; vecmem::copy& copy = vo.copy(); vecmem::copy& async_copy = vo.async_copy(); + const traccc::detector_buffer detector_buffer = + traccc::buffer_from_host_detector(host_det, mng_mr, copy); + // Seeding algorithms const traccc::seedfinder_config seedfinder_config(seeding_opts); const traccc::seedfilter_config seedfilter_config(seeding_opts); @@ -298,9 +297,9 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, { traccc::performance::timer t("Track finding with CKF (alpaka)", elapsedTimes); - track_candidates_alpaka_buffer = - device_finding(det_view, field, measurements_alpaka_buffer, - params_alpaka_buffer); + track_candidates_alpaka_buffer = device_finding( + detector_buffer, field, measurements_alpaka_buffer, + params_alpaka_buffer); } if (accelerator_opts.compare_with_cpu) { @@ -320,7 +319,7 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, elapsedTimes); track_states_alpaka_buffer = - device_fitting(det_view, field, + device_fitting(detector_buffer, field, {track_candidates_alpaka_buffer, measurements_alpaka_buffer}); } diff --git a/examples/run/alpaka/seq_example_alpaka.cpp b/examples/run/alpaka/seq_example_alpaka.cpp index dcdad94575..1dc88ccec6 100644 --- a/examples/run/alpaka/seq_example_alpaka.cpp +++ b/examples/run/alpaka/seq_example_alpaka.cpp @@ -94,16 +94,13 @@ int seq_run(const traccc::opts::detector& detector_opts, copy(host_det_descr_data, device_det_descr)->wait(); // Construct a Detray detector object, if supported by the configuration. - traccc::default_detector::host host_detector{host_mr}; - traccc::default_detector::buffer device_detector; - traccc::default_detector::view device_detector_view; - traccc::io::read_detector( - host_detector, host_mr, detector_opts.detector_file, - detector_opts.material_file, detector_opts.grid_file); - device_detector = detray::get_buffer(host_detector, device_mr, copy); - queue.synchronize(); - const auto& const_device_detector = device_detector; - device_detector_view = detray::get_data(const_device_detector); + traccc::host_detector host_det; + traccc::io::read_detector(host_det, host_mr, detector_opts.detector_file, + detector_opts.material_file, + detector_opts.grid_file); + + const traccc::detector_buffer detector_buffer = + traccc::buffer_from_host_detector(host_det, device_mr, host_copy); // Output stats uint64_t n_cells = 0; @@ -122,8 +119,7 @@ int seq_run(const traccc::opts::detector& detector_opts, using host_spacepoint_formation_algorithm = traccc::host::silicon_pixel_spacepoint_formation_algorithm; using device_spacepoint_formation_algorithm = - traccc::alpaka::spacepoint_formation_algorithm< - traccc::default_detector::device>; + traccc::alpaka::spacepoint_formation_algorithm; using host_finding_algorithm = traccc::host::combinatorial_kalman_filter_algorithm; @@ -265,7 +261,7 @@ int seq_run(const traccc::opts::detector& detector_opts, traccc::performance::timer t("Spacepoint formation (alpaka)", elapsedTimes); spacepoints_alpaka_buffer = - sf_alpaka(device_detector_view, measurements_alpaka_buffer); + sf_alpaka(detector_buffer, measurements_alpaka_buffer); queue.synchronize(); } // stop measuring spacepoint formation alpaka timer @@ -274,7 +270,7 @@ int seq_run(const traccc::opts::detector& detector_opts, traccc::performance::timer t("Spacepoint formation (cpu)", elapsedTimes); spacepoints_per_event = - sf(host_detector, vecmem::get_data(measurements_per_event)); + sf(host_det, vecmem::get_data(measurements_per_event)); } // stop measuring spacepoint formation cpu timer // Alpaka @@ -314,7 +310,7 @@ int seq_run(const traccc::opts::detector& detector_opts, traccc::performance::timer timer{"Track finding (alpaka)", elapsedTimes}; track_candidates_buffer = finding_alg_alpaka( - device_detector_view, field, measurements_alpaka_buffer, + detector_buffer, field, measurements_alpaka_buffer, params_alpaka_buffer); } @@ -322,10 +318,9 @@ int seq_run(const traccc::opts::detector& detector_opts, if (accelerator_opts.compare_with_cpu) { traccc::performance::timer timer{"Track finding (cpu)", elapsedTimes}; - track_candidates = - finding_alg(host_detector, field, - vecmem::get_data(measurements_per_event), - vecmem::get_data(params)); + track_candidates = finding_alg( + host_det, field, vecmem::get_data(measurements_per_event), + vecmem::get_data(params)); } // Alpaka @@ -333,7 +328,7 @@ int seq_run(const traccc::opts::detector& detector_opts, traccc::performance::timer timer{"Track fitting (alpaka)", elapsedTimes}; track_states_buffer = fitting_alg_alpaka( - device_detector_view, field, + detector_buffer, field, {track_candidates_buffer, measurements_alpaka_buffer}); } @@ -342,7 +337,7 @@ int seq_run(const traccc::opts::detector& detector_opts, traccc::performance::timer timer{"Track fitting (cpu)", elapsedTimes}; track_states = - fitting_alg(host_detector, field, + fitting_alg(host_det, field, {vecmem::get_data(track_candidates), vecmem::get_data(measurements_per_event)}); } @@ -453,8 +448,7 @@ int seq_run(const traccc::opts::detector& detector_opts, traccc::event_data evt_data(input_opts.directory, event, host_mr, input_opts.use_acts_geom_source, - &host_detector, input_opts.format, - false); + &host_det, input_opts.format, false); sd_performance_writer.write( vecmem::get_data(seeds), diff --git a/examples/run/common/throughput_mt.ipp b/examples/run/common/throughput_mt.ipp index 94cef0dc98..19d3fa7829 100644 --- a/examples/run/common/throughput_mt.ipp +++ b/examples/run/common/throughput_mt.ipp @@ -12,6 +12,7 @@ // Project include(s) #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" // Command line option include(s). #include "traccc/options/clusterization.hpp" @@ -101,7 +102,7 @@ int throughput_mt(std::string_view description, int argc, char* argv[]) { traccc::data_format::json); // Construct a Detray detector object, if supported by the configuration. - traccc::default_detector::host detector{host_mr}; + traccc::host_detector detector; traccc::io::read_detector(detector, host_mr, detector_opts.detector_file, detector_opts.material_file, detector_opts.grid_file); diff --git a/examples/run/common/throughput_st.ipp b/examples/run/common/throughput_st.ipp index c60bb0949f..67ab2b7367 100644 --- a/examples/run/common/throughput_st.ipp +++ b/examples/run/common/throughput_st.ipp @@ -12,6 +12,7 @@ // Project include(s) #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" // Command line option include(s). #include "traccc/options/clusterization.hpp" @@ -89,7 +90,7 @@ int throughput_st(std::string_view description, int argc, char* argv[]) { traccc::data_format::json); // Construct a Detray detector object, if supported by the configuration. - traccc::default_detector::host detector{host_mr}; + traccc::host_detector detector; traccc::io::read_detector(detector, host_mr, detector_opts.detector_file, detector_opts.material_file, detector_opts.grid_file); diff --git a/examples/run/cpu/full_chain_algorithm.cpp b/examples/run/cpu/full_chain_algorithm.cpp index a07486a9fa..dcd99c05d5 100644 --- a/examples/run/cpu/full_chain_algorithm.cpp +++ b/examples/run/cpu/full_chain_algorithm.cpp @@ -18,7 +18,7 @@ full_chain_algorithm::full_chain_algorithm( const finding_algorithm::config_type& finding_config, const fitting_algorithm::config_type& fitting_config, const silicon_detector_description::host& det_descr, - const magnetic_field& field, detector_type* detector, + const magnetic_field& field, const host_detector* detector, std::unique_ptr logger) : messaging(logger->clone()), m_mr(mr), diff --git a/examples/run/cpu/full_chain_algorithm.hpp b/examples/run/cpu/full_chain_algorithm.hpp index 0741796e57..a34bac892b 100644 --- a/examples/run/cpu/full_chain_algorithm.hpp +++ b/examples/run/cpu/full_chain_algorithm.hpp @@ -16,6 +16,7 @@ #include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" #include "traccc/geometry/silicon_detector_description.hpp" #include "traccc/seeding/seeding_algorithm.hpp" #include "traccc/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" @@ -47,9 +48,6 @@ class full_chain_algorithm /// @name Type declaration(s) /// @{ - /// Detector type used during track finding and fitting - using detector_type = traccc::default_detector::host; - /// Clusterization algorithm type using clustering_algorithm = host::clusterization_algorithm; /// Spacepoint formation algorithm type @@ -78,7 +76,8 @@ class full_chain_algorithm const finding_algorithm::config_type& finding_config, const fitting_algorithm::config_type& fitting_config, const silicon_detector_description::host& det_descr, - const magnetic_field& field, detector_type* detector, + const magnetic_field& field, + const host_detector* detector, std::unique_ptr logger); /// Reconstruct track parameters in the entire detector @@ -111,7 +110,7 @@ class full_chain_algorithm std::reference_wrapper m_det_descr; /// Detector - detector_type* m_detector; + const host_detector* m_detector; /// @name Sub-algorithms used by this full-chain algorithm /// @{ diff --git a/examples/run/cpu/misaligned_truth_fitting_example.cpp b/examples/run/cpu/misaligned_truth_fitting_example.cpp index ce8b1c187e..553662743e 100644 --- a/examples/run/cpu/misaligned_truth_fitting_example.cpp +++ b/examples/run/cpu/misaligned_truth_fitting_example.cpp @@ -12,6 +12,7 @@ #include "traccc/definitions/primitives.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/io/read_detector.hpp" #include "traccc/io/utils.hpp" #include "traccc/options/detector.hpp" #include "traccc/options/input_data.hpp" @@ -84,19 +85,13 @@ int main(int argc, char* argv[]) { const auto field = details::make_magnetic_field(bfield_opts); // Read the detector - detray::io::detector_reader_config reader_cfg{}; - reader_cfg.add_file( - traccc::io::get_absolute_path(detector_opts.detector_file)); - if (!detector_opts.material_file.empty()) { - reader_cfg.add_file( - traccc::io::get_absolute_path(detector_opts.material_file)); - } - if (!detector_opts.grid_file.empty()) { - reader_cfg.add_file( - traccc::io::get_absolute_path(detector_opts.grid_file)); - } - const auto [host_det, names] = - detray::io::read_detector(host_mr, reader_cfg); + traccc::host_detector polymorphic_detector; + traccc::io::read_detector( + polymorphic_detector, host_mr, detector_opts.detector_file, + detector_opts.material_file, detector_opts.grid_file); + + const traccc::default_detector::host& host_det = + polymorphic_detector.as(); /// Create a "misaligned" context in the transform store using xf_container = host_detector_type::transform_container; @@ -150,8 +145,9 @@ int main(int argc, char* argv[]) { // Truth Track Candidates traccc::event_data evt_data(input_opts.directory, event, host_mr, - input_opts.use_acts_geom_source, &host_det, - input_opts.format, false); + input_opts.use_acts_geom_source, + &polymorphic_detector, input_opts.format, + false); // For the first half of events run Alg0 if ((event - input_opts.skip) / (input_opts.events / 2) == 0) { @@ -162,7 +158,7 @@ int main(int argc, char* argv[]) { // Run fitting auto track_states = host_fitting0( - host_det, field, + polymorphic_detector, field, {vecmem::get_data(truth_track_candidates.tracks), vecmem::get_data(truth_track_candidates.measurements)}); @@ -187,7 +183,7 @@ int main(int argc, char* argv[]) { // Run fitting auto track_states = host_fitting1( - host_det, field, + polymorphic_detector, field, {vecmem::get_data(truth_track_candidates.tracks), vecmem::get_data(truth_track_candidates.measurements)}); diff --git a/examples/run/cpu/seeding_example.cpp b/examples/run/cpu/seeding_example.cpp index 2e147c6f2a..811af8f85e 100644 --- a/examples/run/cpu/seeding_example.cpp +++ b/examples/run/cpu/seeding_example.cpp @@ -10,6 +10,7 @@ #include "traccc/definitions/common.hpp" #include "traccc/definitions/primitives.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" #include "traccc/utils/memory_resource.hpp" #include "traccc/utils/propagation.hpp" @@ -116,7 +117,7 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, const traccc::vector3 field_vec(seeding_opts); // Construct a Detray detector object, if supported by the configuration. - traccc::default_detector::host detector{host_mr}; + traccc::host_detector detector; traccc::io::read_detector(detector, host_mr, detector_opts.detector_file, detector_opts.material_file, detector_opts.grid_file); @@ -252,9 +253,13 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, evt_data); for (unsigned int i = 0; i < track_states.tracks.size(); i++) { - fit_performance_writer.write( - track_states.tracks.at(i), track_states.states, - measurements_per_event, detector, evt_data); + host_detector_visitor( + detector, [&]( + const typename detector_traits_t::host& det) { + fit_performance_writer.write( + track_states.tracks.at(i), track_states.states, + measurements_per_event, det, evt_data); + }); } } } diff --git a/examples/run/cpu/seq_example.cpp b/examples/run/cpu/seq_example.cpp index c769b76496..a5d38a415b 100644 --- a/examples/run/cpu/seq_example.cpp +++ b/examples/run/cpu/seq_example.cpp @@ -7,6 +7,7 @@ // core #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" #include "traccc/utils/memory_resource.hpp" #include "traccc/utils/propagation.hpp" @@ -96,10 +97,10 @@ int seq_run(const traccc::opts::input_data& input_opts, vecmem::get_data(det_descr)}; // Construct a Detray detector object, if supported by the configuration. - traccc::default_detector::host detector{host_mr}; - traccc::io::read_detector(detector, host_mr, detector_opts.detector_file, - detector_opts.material_file, - detector_opts.grid_file); + traccc::host_detector polymorphic_detector; + traccc::io::read_detector( + polymorphic_detector, host_mr, detector_opts.detector_file, + detector_opts.material_file, detector_opts.grid_file); // Output stats uint64_t n_cells = 0; @@ -235,7 +236,8 @@ int seq_run(const traccc::opts::input_data& input_opts, traccc::performance::timer timer{"Spacepoint formation", elapsedTimes}; spacepoints_per_event = - sf(detector, vecmem::get_data(measurements_per_event)); + sf(polymorphic_detector, + vecmem::get_data(measurements_per_event)); } if (output_opts.directory != "") { traccc::io::write(event, output_opts.directory, @@ -247,7 +249,6 @@ int seq_run(const traccc::opts::input_data& input_opts, /*----------------------- Seeding algorithm -----------------------*/ - { traccc::performance::timer timer{"Seeding", elapsedTimes}; seeds = sa(vecmem::get_data(spacepoints_per_event)); @@ -272,15 +273,17 @@ int seq_run(const traccc::opts::input_data& input_opts, { traccc::performance::timer timer{"Track finding", elapsedTimes}; - track_candidates = finding_alg( - detector, field, vecmem::get_data(measurements_per_event), - vecmem::get_data(params)); + track_candidates = + finding_alg(polymorphic_detector, field, + vecmem::get_data(measurements_per_event), + vecmem::get_data(params)); } if (output_opts.directory != "") { - traccc::io::write( - event, output_opts.directory, output_opts.format, - vecmem::get_data(track_candidates), - vecmem::get_data(measurements_per_event), detector); + traccc::io::write(event, output_opts.directory, + output_opts.format, + vecmem::get_data(track_candidates), + vecmem::get_data(measurements_per_event), + polymorphic_detector); } { @@ -295,7 +298,7 @@ int seq_run(const traccc::opts::input_data& input_opts, { traccc::performance::timer timer{"Track fitting", elapsedTimes}; track_states = - fitting_alg(detector, field, + fitting_alg(polymorphic_detector, field, {vecmem::get_data(resolved_track_candidates), vecmem::get_data(measurements_per_event)}); } @@ -322,7 +325,8 @@ int seq_run(const traccc::opts::input_data& input_opts, traccc::event_data evt_data(input_opts.directory, event, host_mr, input_opts.use_acts_geom_source, - &detector, input_opts.format, true); + &polymorphic_detector, + input_opts.format, true); evt_data.fill_cca_result(cells_per_event, clusters_per_event, measurements_per_event, det_descr); @@ -340,9 +344,14 @@ int seq_run(const traccc::opts::input_data& input_opts, evt_data); for (unsigned int i = 0; i < track_states.tracks.size(); i++) { - fit_performance_writer.write( - track_states.tracks.at(i), track_states.states, - measurements_per_event, detector, evt_data); + host_detector_visitor( + polymorphic_detector, + [&]( + const typename detector_traits_t::host& det) { + fit_performance_writer.write( + track_states.tracks.at(i), track_states.states, + measurements_per_event, det, evt_data); + }); } } } diff --git a/examples/run/cpu/truth_finding_example.cpp b/examples/run/cpu/truth_finding_example.cpp index 2462b6d6de..2d11a877ef 100644 --- a/examples/run/cpu/truth_finding_example.cpp +++ b/examples/run/cpu/truth_finding_example.cpp @@ -14,6 +14,7 @@ #include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" #include "traccc/io/read_detector.hpp" #include "traccc/io/read_detector_description.hpp" #include "traccc/io/read_measurements.hpp" @@ -80,10 +81,10 @@ int seq_run(const traccc::opts::track_finding& finding_opts, const auto field = traccc::details::make_magnetic_field(bfield_opts); // Construct a Detray detector object, if supported by the configuration. - traccc::default_detector::host detector{host_mr}; - traccc::io::read_detector(detector, host_mr, detector_opts.detector_file, - detector_opts.material_file, - detector_opts.grid_file); + traccc::host_detector polymorphic_detector; + traccc::io::read_detector( + polymorphic_detector, host_mr, detector_opts.detector_file, + detector_opts.material_file, detector_opts.grid_file); /***************************** * Do the reconstruction @@ -116,23 +117,30 @@ int seq_run(const traccc::opts::track_finding& finding_opts, traccc::host::kalman_fitting_algorithm host_fitting( fit_cfg, host_mr, copy, logger().clone("FittingAlg")); - // Seed generator - traccc::seed_generator sg(detector, - stddevs); - // Iterate over events for (std::size_t event = input_opts.skip; event < input_opts.events + input_opts.skip; ++event) { // Truth Track Candidates traccc::event_data evt_data(input_opts.directory, event, host_mr, - input_opts.use_acts_geom_source, &detector, - input_opts.format, false); + input_opts.use_acts_geom_source, + &polymorphic_detector, input_opts.format, + false); traccc::edm::track_candidate_container::host truth_track_candidates{host_mr}; - evt_data.generate_truth_candidates(truth_track_candidates, sg, host_mr, - truth_finding_opts.m_pT_min); + + // Seed generator + host_detector_visitor( + polymorphic_detector, + [&]( + const typename detector_traits_t::host& det) { + traccc::seed_generator sg( + det, stddevs); + evt_data.generate_truth_candidates(truth_track_candidates, sg, + host_mr, + truth_finding_opts.m_pT_min); + }); // Prepare truth seeds traccc::bound_track_parameters_collection_types::host seeds(&host_mr); @@ -146,20 +154,20 @@ int seq_run(const traccc::opts::track_finding& finding_opts, &host_mr}; traccc::io::read_measurements( measurements_per_event, event, input_opts.directory, - (input_opts.use_acts_geom_source ? &detector : nullptr), + (input_opts.use_acts_geom_source ? &polymorphic_detector : nullptr), input_opts.format); // Run finding auto track_candidates = host_finding( - detector, field, vecmem::get_data(measurements_per_event), - vecmem::get_data(seeds)); + polymorphic_detector, field, + vecmem::get_data(measurements_per_event), vecmem::get_data(seeds)); std::cout << "Number of found tracks: " << track_candidates.size() << std::endl; // Run fitting auto track_states = - host_fitting(detector, field, + host_fitting(polymorphic_detector, field, {vecmem::get_data(track_candidates), vecmem::get_data(measurements_per_event)}); @@ -174,9 +182,14 @@ int seq_run(const traccc::opts::track_finding& finding_opts, evt_data); for (std::size_t i = 0; i < n_fitted_tracks; i++) { - fit_performance_writer.write( - track_states.tracks.at(i), track_states.states, - measurements_per_event, detector, evt_data); + host_detector_visitor( + polymorphic_detector, + [&]( + const typename detector_traits_t::host& det) { + fit_performance_writer.write( + track_states.tracks.at(i), track_states.states, + measurements_per_event, det, evt_data); + }); } } } diff --git a/examples/run/cpu/truth_fitting_example.cpp b/examples/run/cpu/truth_fitting_example.cpp index 468784447f..4f6273eb4c 100644 --- a/examples/run/cpu/truth_fitting_example.cpp +++ b/examples/run/cpu/truth_fitting_example.cpp @@ -12,6 +12,7 @@ #include "traccc/definitions/primitives.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/io/read_detector.hpp" #include "traccc/io/utils.hpp" #include "traccc/options/detector.hpp" #include "traccc/options/input_data.hpp" @@ -63,9 +64,6 @@ int main(int argc, char* argv[]) { argv, logger().cloneWithSuffix("Options")}; - /// Type declarations - using host_detector_type = traccc::default_detector::host; - // Memory resources used by the application. vecmem::host_memory_resource host_mr; // Copy obejct @@ -84,19 +82,10 @@ int main(int argc, char* argv[]) { const auto field = traccc::details::make_magnetic_field(bfield_opts); // Read the detector - detray::io::detector_reader_config reader_cfg{}; - reader_cfg.add_file( - traccc::io::get_absolute_path(detector_opts.detector_file)); - if (!detector_opts.material_file.empty()) { - reader_cfg.add_file( - traccc::io::get_absolute_path(detector_opts.material_file)); - } - if (!detector_opts.grid_file.empty()) { - reader_cfg.add_file( - traccc::io::get_absolute_path(detector_opts.grid_file)); - } - const auto [host_det, names] = - detray::io::read_detector(host_mr, reader_cfg); + traccc::host_detector polymorphic_detector; + traccc::io::read_detector( + polymorphic_detector, host_mr, detector_opts.detector_file, + detector_opts.material_file, detector_opts.grid_file); /***************************** * Do the reconstruction @@ -118,25 +107,33 @@ int main(int argc, char* argv[]) { traccc::host::kalman_fitting_algorithm host_fitting( fit_cfg, host_mr, copy, logger().clone("FittingAlg")); - // Seed generator - traccc::seed_generator sg(host_det, stddevs); - // Iterate over events for (auto event = input_opts.skip; event < input_opts.events + input_opts.skip; ++event) { // Truth Track Candidates traccc::event_data evt_data(input_opts.directory, event, host_mr, - input_opts.use_acts_geom_source, &host_det, - input_opts.format, false); + input_opts.use_acts_geom_source, + &polymorphic_detector, input_opts.format, + false); traccc::edm::track_candidate_container::host truth_track_candidates{host_mr}; - evt_data.generate_truth_candidates(truth_track_candidates, sg, host_mr); + + host_detector_visitor( + polymorphic_detector, + [&]( + const typename detector_traits_t::host& det) { + // Seed generator + traccc::seed_generator sg( + det, stddevs); + evt_data.generate_truth_candidates(truth_track_candidates, sg, + host_mr); + }); // Run fitting auto track_states = host_fitting( - host_det, field, + polymorphic_detector, field, {vecmem::get_data(truth_track_candidates.tracks), vecmem::get_data(truth_track_candidates.measurements)}); @@ -147,9 +144,14 @@ int main(int argc, char* argv[]) { if (performance_opts.run) { for (unsigned int i = 0; i < n_fitted_tracks; i++) { - fit_performance_writer.write( - track_states.tracks.at(i), track_states.states, - truth_track_candidates.measurements, host_det, evt_data); + host_detector_visitor( + polymorphic_detector, + [&]( + const typename detector_traits_t::host& det) { + fit_performance_writer.write( + track_states.tracks.at(i), track_states.states, + truth_track_candidates.measurements, det, evt_data); + }); } } } diff --git a/examples/run/cuda/full_chain_algorithm.cpp b/examples/run/cuda/full_chain_algorithm.cpp index 367fbd2200..76b81bffa3 100644 --- a/examples/run/cuda/full_chain_algorithm.cpp +++ b/examples/run/cuda/full_chain_algorithm.cpp @@ -39,7 +39,7 @@ full_chain_algorithm::full_chain_algorithm( const finding_algorithm::config_type& finding_config, const fitting_algorithm::config_type& fitting_config, const silicon_detector_description::host& det_descr, - const magnetic_field& field, host_detector_type* detector, + const magnetic_field& field, host_detector* detector, std::unique_ptr logger) : messaging(logger->clone()), m_host_mr(host_mr), @@ -95,9 +95,7 @@ full_chain_algorithm::full_chain_algorithm( m_copy(vecmem::get_data(m_det_descr.get()), m_device_det_descr)->ignore(); if (m_detector != nullptr) { m_device_detector = - detray::get_buffer(*m_detector, m_device_mr, m_copy); - const auto& const_device_detector = m_device_detector; - m_device_detector_view = detray::get_data(const_device_detector); + traccc::buffer_from_host_detector(*m_detector, m_device_mr, m_copy); } } @@ -150,9 +148,7 @@ full_chain_algorithm::full_chain_algorithm(const full_chain_algorithm& parent) m_copy(vecmem::get_data(m_det_descr.get()), m_device_det_descr)->ignore(); if (m_detector != nullptr) { m_device_detector = - detray::get_buffer(*m_detector, m_device_mr, m_copy); - const auto& const_device_detector = m_device_detector; - m_device_detector_view = detray::get_data(const_device_detector); + traccc::buffer_from_host_detector(*m_detector, m_device_mr, m_copy); } } @@ -173,21 +169,20 @@ full_chain_algorithm::output_type full_chain_algorithm::operator()( // If we have a Detray detector, run the seeding, track finding and fitting. if (m_detector != nullptr) { - // Run the seed-finding (asynchronously). const spacepoint_formation_algorithm::output_type spacepoints = - m_spacepoint_formation(m_device_detector_view, measurements); + m_spacepoint_formation(m_device_detector, measurements); const track_params_estimation::output_type track_params = m_track_parameter_estimation(measurements, spacepoints, m_seeding(spacepoints), m_field_vec); // Run the track finding (asynchronously). - const finding_algorithm::output_type track_candidates = m_finding( - m_device_detector_view, m_field, measurements, track_params); + const finding_algorithm::output_type track_candidates = + m_finding(m_device_detector, m_field, measurements, track_params); // Run the track fitting (asynchronously). const fitting_algorithm::output_type track_states = m_fitting( - m_device_detector_view, m_field, {track_candidates, measurements}); + m_device_detector, m_field, {track_candidates, measurements}); // Copy a limited amount of result data back to the host. const auto host_tracks = @@ -230,7 +225,7 @@ bound_track_parameters_collection_types::host full_chain_algorithm::seeding( // Run the seed-finding (asynchronously). const spacepoint_formation_algorithm::output_type spacepoints = - m_spacepoint_formation(m_device_detector_view, measurements); + m_spacepoint_formation(m_device_detector, measurements); const track_params_estimation::output_type track_params = m_track_parameter_estimation(measurements, spacepoints, m_seeding(spacepoints), m_field_vec); diff --git a/examples/run/cuda/full_chain_algorithm.hpp b/examples/run/cuda/full_chain_algorithm.hpp index 1dd436cae4..6462895c66 100644 --- a/examples/run/cuda/full_chain_algorithm.hpp +++ b/examples/run/cuda/full_chain_algorithm.hpp @@ -21,6 +21,8 @@ #include "traccc/edm/track_fit_collection.hpp" #include "traccc/edm/track_parameters.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/detector_buffer.hpp" +#include "traccc/geometry/host_detector.hpp" #include "traccc/geometry/silicon_detector_description.hpp" #include "traccc/utils/algorithm.hpp" #include "traccc/utils/messaging.hpp" @@ -51,18 +53,9 @@ class full_chain_algorithm public: /// @name Type declaration(s) /// @{ - - /// (Host) Detector type used during track finding and fitting - using host_detector_type = traccc::default_detector::host; - /// (Device) Detector type used during track finding and fitting - using device_detector_type = traccc::default_detector::device; - - using scalar_type = device_detector_type::scalar_type; - /// Spacepoint formation algorithm type using spacepoint_formation_algorithm = - traccc::cuda::spacepoint_formation_algorithm< - traccc::default_detector::device>; + traccc::cuda::spacepoint_formation_algorithm; /// Clustering algorithm type using clustering_algorithm = traccc::cuda::clusterization_algorithm; /// Track finding algorithm type @@ -86,8 +79,7 @@ class full_chain_algorithm const finding_algorithm::config_type& finding_config, const fitting_algorithm::config_type& fitting_config, const silicon_detector_description::host& det_descr, - const magnetic_field& field, - host_detector_type* detector, + const magnetic_field& field, host_detector* detector, std::unique_ptr logger); /// Copy constructor @@ -146,11 +138,8 @@ class full_chain_algorithm /// Detector description buffer silicon_detector_description::buffer m_device_det_descr; /// Host detector - host_detector_type* m_detector; - /// Buffer holding the detector's payload on the device - host_detector_type::buffer_type m_device_detector; - /// View of the detector's payload on the device - host_detector_type::const_view_type m_device_detector_view; + host_detector* m_detector; + detector_buffer m_device_detector; /// @name Sub-algorithms used by this full-chain algorithm /// @{ diff --git a/examples/run/cuda/seeding_example_cuda.cpp b/examples/run/cuda/seeding_example_cuda.cpp index 28785e1f52..11d77240fb 100644 --- a/examples/run/cuda/seeding_example_cuda.cpp +++ b/examples/run/cuda/seeding_example_cuda.cpp @@ -22,6 +22,7 @@ #include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" #include "traccc/io/read_detector.hpp" #include "traccc/io/read_detector_description.hpp" #include "traccc/io/read_measurements.hpp" @@ -133,19 +134,18 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, : traccc::cuda::magnetic_field_storage::global_memory)); // Construct a Detray detector object, if supported by the configuration. - traccc::default_detector::host host_det{mng_mr}; + traccc::host_detector host_det; traccc::io::read_detector(host_det, mng_mr, detector_opts.detector_file, detector_opts.material_file, detector_opts.grid_file); - // Detector view object - const traccc::default_detector::host& const_host_det = host_det; - traccc::default_detector::view det_view = detray::get_data(const_host_det); - // Copy objects vecmem::copy host_copy; vecmem::cuda::copy copy; + const traccc::detector_buffer detector_buffer = + traccc::buffer_from_host_detector(host_det, mng_mr, host_copy); + // Seeding algorithm const traccc::seedfinder_config seedfinder_config(seeding_opts); const traccc::seedfilter_config seedfilter_config(seeding_opts); @@ -309,7 +309,7 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, traccc::performance::timer t("Track finding with CKF (cuda)", elapsedTimes); track_candidates_cuda_buffer = device_finding( - det_view, device_field, measurements_cuda_buffer, + detector_buffer, device_field, measurements_cuda_buffer, params_cuda_buffer); } @@ -331,7 +331,7 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, elapsedTimes); track_states_cuda_buffer = device_fitting( - det_view, device_field, + detector_buffer, device_field, {track_candidates_cuda_buffer, measurements_cuda_buffer}); } @@ -438,9 +438,14 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, evt_data); for (unsigned int i = 0; i < track_states_cuda.tracks.size(); i++) { - fit_performance_writer.write( - track_states_cuda.tracks.at(i), track_states_cuda.states, - measurements_per_event, host_det, evt_data); + host_detector_visitor( + host_det, [&]( + const typename detector_traits_t::host& det) { + fit_performance_writer.write( + track_states_cuda.tracks.at(i), + track_states_cuda.states, measurements_per_event, + det, evt_data); + }); } } } diff --git a/examples/run/cuda/seq_example_cuda.cpp b/examples/run/cuda/seq_example_cuda.cpp index 4f81ab8757..364923b1f5 100644 --- a/examples/run/cuda/seq_example_cuda.cpp +++ b/examples/run/cuda/seq_example_cuda.cpp @@ -24,6 +24,8 @@ #include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/detector_buffer.hpp" +#include "traccc/geometry/host_detector.hpp" #include "traccc/io/read_cells.hpp" #include "traccc/io/read_detector.hpp" #include "traccc/io/read_detector_description.hpp" @@ -103,17 +105,13 @@ int seq_run(const traccc::opts::detector& detector_opts, copy(host_det_descr_data, device_det_descr)->wait(); // Construct a Detray detector object, if supported by the configuration. - traccc::default_detector::host host_detector{host_mr}; - traccc::default_detector::buffer device_detector; - traccc::default_detector::view device_detector_view; + traccc::host_detector host_detector; traccc::io::read_detector( host_detector, host_mr, detector_opts.detector_file, detector_opts.material_file, detector_opts.grid_file); - device_detector = detray::get_buffer(host_detector, device_mr, copy); + const traccc::detector_buffer device_detector = + traccc::buffer_from_host_detector(host_detector, device_mr, copy); stream.synchronize(); - const traccc::default_detector::buffer& const_device_detector = - device_detector; - device_detector_view = detray::get_data(const_device_detector); // Output stats uint64_t n_cells = 0; @@ -134,8 +132,7 @@ int seq_run(const traccc::opts::detector& detector_opts, using host_spacepoint_formation_algorithm = traccc::host::silicon_pixel_spacepoint_formation_algorithm; using device_spacepoint_formation_algorithm = - traccc::cuda::spacepoint_formation_algorithm< - traccc::default_detector::device>; + traccc::cuda::spacepoint_formation_algorithm; using host_finding_algorithm = traccc::host::combinatorial_kalman_filter_algorithm; @@ -294,7 +291,7 @@ int seq_run(const traccc::opts::detector& detector_opts, traccc::performance::timer t("Spacepoint formation (cuda)", elapsedTimes); spacepoints_cuda_buffer = - sf_cuda(device_detector_view, measurements_cuda_buffer); + sf_cuda(device_detector, measurements_cuda_buffer); stream.synchronize(); } // stop measuring spacepoint formation cuda timer @@ -343,8 +340,8 @@ int seq_run(const traccc::opts::detector& detector_opts, traccc::performance::timer timer{"Track finding (cuda)", elapsedTimes}; track_candidates_buffer = finding_alg_cuda( - device_detector_view, device_field, - measurements_cuda_buffer, params_cuda_buffer); + device_detector, device_field, measurements_cuda_buffer, + params_cuda_buffer); } // CPU @@ -379,7 +376,7 @@ int seq_run(const traccc::opts::detector& detector_opts, traccc::performance::timer timer{"Track fitting (cuda)", elapsedTimes}; track_states_buffer = fitting_alg_cuda( - device_detector_view, device_field, + device_detector, device_field, {res_track_candidates_buffer, measurements_cuda_buffer}); } diff --git a/examples/run/cuda/truth_finding_example_cuda.cpp b/examples/run/cuda/truth_finding_example_cuda.cpp index 0ec2d9586c..4a6e17dcce 100644 --- a/examples/run/cuda/truth_finding_example_cuda.cpp +++ b/examples/run/cuda/truth_finding_example_cuda.cpp @@ -19,6 +19,7 @@ #include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" #include "traccc/io/read_detector.hpp" #include "traccc/io/read_detector_description.hpp" #include "traccc/io/read_measurements.hpp" @@ -101,14 +102,10 @@ int seq_run(const traccc::opts::track_finding& finding_opts, const auto device_field = traccc::cuda::make_magnetic_field(host_field); // Construct a Detray detector object, if supported by the configuration. - traccc::default_detector::host detector{mng_mr}; - traccc::io::read_detector(detector, mng_mr, detector_opts.detector_file, - detector_opts.material_file, - detector_opts.grid_file); - - // Detector view object - const traccc::default_detector::host& const_detector = detector; - traccc::default_detector::view det_view = detray::get_data(const_detector); + traccc::host_detector polymorphic_detector; + traccc::io::read_detector( + polymorphic_detector, mng_mr, detector_opts.detector_file, + detector_opts.material_file, detector_opts.grid_file); /***************************** * Do the reconstruction @@ -121,6 +118,10 @@ int seq_run(const traccc::opts::track_finding& finding_opts, vecmem::copy host_copy; vecmem::cuda::async_copy async_copy{stream.cudaStream()}; + const traccc::detector_buffer detector_buffer = + traccc::buffer_from_host_detector(polymorphic_detector, device_mr, + async_copy); + // Standard deviations for seed track parameters static constexpr std::array stddevs = {1e-4f * traccc::unit::mm, @@ -154,23 +155,30 @@ int seq_run(const traccc::opts::track_finding& finding_opts, traccc::performance::timing_info elapsedTimes; - // Seed generator - traccc::seed_generator sg(detector, - stddevs); - // Iterate over events for (std::size_t event = input_opts.skip; event < input_opts.events + input_opts.skip; ++event) { // Truth Track Candidates traccc::event_data evt_data(input_opts.directory, event, host_mr, - input_opts.use_acts_geom_source, &detector, - input_opts.format, false); + input_opts.use_acts_geom_source, + &polymorphic_detector, input_opts.format, + false); traccc::edm::track_candidate_container::host truth_track_candidates{host_mr}; - evt_data.generate_truth_candidates(truth_track_candidates, sg, host_mr, - truth_finding_opts.m_pT_min); + + host_detector_visitor( + polymorphic_detector, + [&]( + const typename detector_traits_t::host& det) { + // Seed generator + traccc::seed_generator sg( + det, stddevs); + evt_data.generate_truth_candidates(truth_track_candidates, sg, + host_mr, + truth_finding_opts.m_pT_min); + }); // Prepare truth seeds traccc::bound_track_parameters_collection_types::host seeds(mr.host); @@ -191,7 +199,7 @@ int seq_run(const traccc::opts::track_finding& finding_opts, mr.host}; traccc::io::read_measurements( measurements_per_event, event, input_opts.directory, - (input_opts.use_acts_geom_source ? &detector : nullptr), + (input_opts.use_acts_geom_source ? &polymorphic_detector : nullptr), input_opts.format); traccc::measurement_collection_types::buffer measurements_cuda_buffer( @@ -209,8 +217,9 @@ int seq_run(const traccc::opts::track_finding& finding_opts, traccc::performance::timer t("Track finding (cuda)", elapsedTimes); // Run finding - track_candidates_cuda_buffer = device_finding( - det_view, device_field, measurements_cuda_buffer, seeds_buffer); + track_candidates_cuda_buffer = + device_finding(detector_buffer, device_field, + measurements_cuda_buffer, seeds_buffer); } traccc::edm::track_candidate_collection::host @@ -228,7 +237,7 @@ int seq_run(const traccc::opts::track_finding& finding_opts, // Run fitting track_states_cuda_buffer = device_fitting( - det_view, device_field, + detector_buffer, device_field, {track_candidates_cuda_buffer, measurements_cuda_buffer}); } traccc::edm::track_fit_container::host @@ -254,7 +263,7 @@ int seq_run(const traccc::opts::track_finding& finding_opts, // Run finding track_candidates = - host_finding(detector, host_field, + host_finding(polymorphic_detector, host_field, vecmem::get_data(measurements_per_event), vecmem::get_data(seeds)); } @@ -265,7 +274,7 @@ int seq_run(const traccc::opts::track_finding& finding_opts, // Run fitting track_states = - host_fitting(detector, host_field, + host_fitting(polymorphic_detector, host_field, {vecmem::get_data(track_candidates), vecmem::get_data(measurements_per_event)}); } @@ -304,9 +313,15 @@ int seq_run(const traccc::opts::track_finding& finding_opts, evt_data); for (unsigned int i = 0; i < track_states_cuda.tracks.size(); i++) { - fit_performance_writer.write( - track_states_cuda.tracks.at(i), track_states_cuda.states, - measurements_per_event, detector, evt_data); + host_detector_visitor( + polymorphic_detector, + [&]( + const typename detector_traits_t::host& det) { + fit_performance_writer.write( + track_states_cuda.tracks.at(i), + track_states_cuda.states, measurements_per_event, + det, evt_data); + }); } } } diff --git a/examples/run/cuda/truth_fitting_example_cuda.cpp b/examples/run/cuda/truth_fitting_example_cuda.cpp index 53aeb10c7f..aec007858a 100644 --- a/examples/run/cuda/truth_fitting_example_cuda.cpp +++ b/examples/run/cuda/truth_fitting_example_cuda.cpp @@ -16,6 +16,8 @@ #include "traccc/device/container_h2d_copy_alg.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" +#include "traccc/io/read_detector.hpp" #include "traccc/io/read_measurements.hpp" #include "traccc/io/utils.hpp" #include "traccc/options/accelerator.hpp" @@ -73,9 +75,6 @@ int main(int argc, char* argv[]) { argv, logger().cloneWithSuffix("Options")}; - /// Type declarations - using host_detector_type = traccc::default_detector::host; - // Memory resources used by the application. vecmem::host_memory_resource host_mr; vecmem::cuda::host_memory_resource cuda_host_mr; @@ -101,22 +100,10 @@ int main(int argc, char* argv[]) { const auto device_field = traccc::cuda::make_magnetic_field(host_field); // Read the detector - detray::io::detector_reader_config reader_cfg{}; - reader_cfg.add_file( - traccc::io::get_absolute_path(detector_opts.detector_file)); - if (!detector_opts.material_file.empty()) { - reader_cfg.add_file( - traccc::io::get_absolute_path(detector_opts.material_file)); - } - if (!detector_opts.grid_file.empty()) { - reader_cfg.add_file( - traccc::io::get_absolute_path(detector_opts.grid_file)); - } - const auto [host_det, names] = - detray::io::read_detector(mng_mr, reader_cfg); - - // Detector view object - auto det_view = detray::get_data(host_det); + traccc::host_detector polymorphic_detector; + traccc::io::read_detector( + polymorphic_detector, mng_mr, detector_opts.detector_file, + detector_opts.material_file, detector_opts.grid_file); /***************************** * Do the reconstruction @@ -129,6 +116,10 @@ int main(int argc, char* argv[]) { vecmem::copy host_copy; vecmem::cuda::async_copy async_copy{stream.cudaStream()}; + const traccc::detector_buffer detector_buffer = + traccc::buffer_from_host_detector(polymorphic_detector, device_mr, + async_copy); + /// Standard deviations for seed track parameters static constexpr std::array stddevs = { 0.03f * traccc::unit::mm, @@ -147,9 +138,6 @@ int main(int argc, char* argv[]) { traccc::cuda::kalman_fitting_algorithm device_fitting( fit_cfg, mr, async_copy, stream, logger().clone("CudaFittingAlg")); - // Seed generator - traccc::seed_generator sg(host_det, stddevs); - traccc::performance::timing_info elapsedTimes; // Iterate over events @@ -158,12 +146,23 @@ int main(int argc, char* argv[]) { // Truth Track Candidates traccc::event_data evt_data(input_opts.directory, event, host_mr, - input_opts.use_acts_geom_source, &host_det, - input_opts.format, false); + input_opts.use_acts_geom_source, + &polymorphic_detector, input_opts.format, + false); traccc::edm::track_candidate_container::host truth_track_candidates{host_mr}; - evt_data.generate_truth_candidates(truth_track_candidates, sg, host_mr); + + host_detector_visitor( + polymorphic_detector, + [&]( + const typename detector_traits_t::host& det) { + // Seed generator + traccc::seed_generator sg( + det, stddevs); + evt_data.generate_truth_candidates(truth_track_candidates, sg, + host_mr); + }); // track candidates buffer traccc::edm::track_candidate_container::buffer @@ -184,7 +183,7 @@ int main(int argc, char* argv[]) { // Run fitting track_states_cuda_buffer = - device_fitting(det_view, device_field, + device_fitting(detector_buffer, device_field, {truth_track_candidates_buffer.tracks, truth_track_candidates_buffer.measurements}); } @@ -210,7 +209,7 @@ int main(int argc, char* argv[]) { // Run fitting track_states = host_fitting( - host_det, host_field, + polymorphic_detector, host_field, {vecmem::get_data(truth_track_candidates.tracks), vecmem::get_data(truth_track_candidates.measurements)}); } @@ -243,9 +242,15 @@ int main(int argc, char* argv[]) { if (performance_opts.run) { for (unsigned int i = 0; i < track_states_cuda.tracks.size(); i++) { - fit_performance_writer.write( - track_states_cuda.tracks.at(i), track_states_cuda.states, - truth_track_candidates.measurements, host_det, evt_data); + host_detector_visitor( + polymorphic_detector, + [&]( + const typename detector_traits_t::host& det) { + fit_performance_writer.write( + track_states_cuda.tracks.at(i), + track_states_cuda.states, + truth_track_candidates.measurements, det, evt_data); + }); } } } diff --git a/examples/run/kokkos/seeding_example_kokkos.cpp b/examples/run/kokkos/seeding_example_kokkos.cpp index 27fc7263f4..d1f3021a5c 100644 --- a/examples/run/kokkos/seeding_example_kokkos.cpp +++ b/examples/run/kokkos/seeding_example_kokkos.cpp @@ -53,7 +53,7 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, traccc::memory_resource mr{host_mr, &host_mr}; // Construct a Detray detector object, if supported by the configuration. - traccc::default_detector::host host_det{host_mr}; + traccc::host_detector host_det; traccc::io::read_detector(host_det, host_mr, detector_opts.detector_file, detector_opts.material_file, detector_opts.grid_file); diff --git a/examples/run/sycl/full_chain_algorithm.hpp b/examples/run/sycl/full_chain_algorithm.hpp index 4a906d4616..9d5a9b2b03 100644 --- a/examples/run/sycl/full_chain_algorithm.hpp +++ b/examples/run/sycl/full_chain_algorithm.hpp @@ -11,6 +11,8 @@ #include "traccc/edm/silicon_cell_collection.hpp" #include "traccc/edm/track_parameters.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/detector_buffer.hpp" +#include "traccc/geometry/host_detector.hpp" #include "traccc/geometry/silicon_detector_description.hpp" #include "traccc/sycl/clusterization/clusterization_algorithm.hpp" #include "traccc/sycl/clusterization/measurement_sorting_algorithm.hpp" @@ -51,11 +53,6 @@ class full_chain_algorithm /// @name (For now dummy...) Type declaration(s) /// @{ - /// (Host) Detector type used during track finding and fitting - using host_detector_type = traccc::default_detector::host; - /// (Device) Detector type used during track finding and fitting - using device_detector_type = traccc::default_detector::device; - /// Spacepoint formation algorithm type using spacepoint_formation_algorithm = traccc::sycl::silicon_pixel_spacepoint_formation_algorithm; @@ -82,8 +79,7 @@ class full_chain_algorithm const finding_algorithm::config_type& finding_config, const fitting_algorithm::config_type& fitting_config, const silicon_detector_description::host& det_descr, - const magnetic_field& field, - host_detector_type* detector, + const magnetic_field& field, host_detector* detector, std::unique_ptr logger); /// Copy constructor @@ -141,12 +137,11 @@ class full_chain_algorithm m_det_descr; /// Detector description buffer silicon_detector_description::buffer m_device_det_descr; + /// Host detector - host_detector_type* m_detector; + host_detector* m_detector; /// Buffer holding the detector's payload on the device - host_detector_type::buffer_type m_device_detector; - /// View of the detector's payload on the device - host_detector_type::const_view_type m_device_detector_view; + detector_buffer m_device_detector; /// @name Sub-algorithms used by this full-chain algorithm /// @{ diff --git a/examples/run/sycl/full_chain_algorithm.sycl b/examples/run/sycl/full_chain_algorithm.sycl index 772babcaf1..c8e1460ca3 100644 --- a/examples/run/sycl/full_chain_algorithm.sycl +++ b/examples/run/sycl/full_chain_algorithm.sycl @@ -60,7 +60,7 @@ full_chain_algorithm::full_chain_algorithm( const finding_algorithm::config_type& finding_config, const fitting_algorithm::config_type& fitting_config, const silicon_detector_description::host& det_descr, - const magnetic_field& field, host_detector_type* detector, + const magnetic_field& field, host_detector* detector, std::unique_ptr logger) : messaging(logger->clone()), m_data(std::make_unique( @@ -131,9 +131,7 @@ full_chain_algorithm::full_chain_algorithm( m_copy(vecmem::get_data(m_det_descr.get()), m_device_det_descr)->wait(); if (m_detector != nullptr) { m_device_detector = - detray::get_buffer(*m_detector, m_device_mr, m_copy); - const auto& const_device_detector = m_device_detector; - m_device_detector_view = detray::get_data(const_device_detector); + traccc::buffer_from_host_detector(*m_detector, m_device_mr, m_copy); } } @@ -201,9 +199,7 @@ full_chain_algorithm::full_chain_algorithm(const full_chain_algorithm& parent) m_copy(vecmem::get_data(m_det_descr.get()), m_device_det_descr)->wait(); if (m_detector != nullptr) { m_device_detector = - detray::get_buffer(*m_detector, m_device_mr, m_copy); - const auto& const_device_detector = m_device_detector; - m_device_detector_view = detray::get_data(const_device_detector); + traccc::buffer_from_host_detector(*m_detector, m_device_mr, m_copy); } } @@ -228,19 +224,19 @@ full_chain_algorithm::output_type full_chain_algorithm::operator()( // Run the seed-finding. const spacepoint_formation_algorithm::output_type spacepoints = - m_spacepoint_formation(m_device_detector_view, measurements); + m_spacepoint_formation(m_device_detector, measurements); const track_params_estimation::output_type track_params = m_track_parameter_estimation(measurements, spacepoints, m_seeding(spacepoints), {0.f, 0.f, m_finder_config.bFieldInZ}); // Run the track finding. - const finding_algorithm::output_type track_candidates = m_finding( - m_device_detector_view, m_field, measurements, track_params); + const finding_algorithm::output_type track_candidates = + m_finding(m_device_detector, m_field, measurements, track_params); // Run the track fitting (asynchronously). const fitting_algorithm::output_type track_states = m_fitting( - m_device_detector_view, m_field, {track_candidates, measurements}); + m_device_detector, m_field, {track_candidates, measurements}); // Copy a limited amount of result data back to the host. const auto host_tracks = @@ -284,7 +280,7 @@ bound_track_parameters_collection_types::host full_chain_algorithm::seeding( // Run the seed-finding. const spacepoint_formation_algorithm::output_type spacepoints = - m_spacepoint_formation(m_device_detector_view, measurements); + m_spacepoint_formation(m_device_detector, measurements); const track_params_estimation::output_type track_params = m_track_parameter_estimation(measurements, spacepoints, m_seeding(spacepoints), diff --git a/examples/run/sycl/seeding_example_sycl.sycl b/examples/run/sycl/seeding_example_sycl.sycl index 1e8fc2991b..c2fd622536 100644 --- a/examples/run/sycl/seeding_example_sycl.sycl +++ b/examples/run/sycl/seeding_example_sycl.sycl @@ -10,6 +10,8 @@ // core #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/detector_buffer.hpp" +#include "traccc/geometry/host_detector.hpp" #include "traccc/utils/propagation.hpp" // algorithms @@ -87,16 +89,14 @@ int seq_run(const traccc::opts::detector& detector_opts, *****************************/ // Construct a Detray detector object, if supported by the configuration. - traccc::default_detector::host host_det{host_mr}; - traccc::default_detector::buffer device_det; - traccc::default_detector::view device_det_view; + traccc::host_detector host_det; traccc::io::read_detector(host_det, host_mr, detector_opts.detector_file, detector_opts.material_file, detector_opts.grid_file); - device_det = detray::get_buffer(host_det, device_mr, copy); + + const traccc::detector_buffer detector_buffer = + traccc::buffer_from_host_detector(host_det, device_mr, copy); q.wait_and_throw(); - const auto& const_device_det = device_det; - device_det_view = detray::get_data(const_device_det); const traccc::vector3 field_vec(seeding_opts); diff --git a/examples/run/sycl/seq_example_sycl.sycl b/examples/run/sycl/seq_example_sycl.sycl index a7f9a0c07c..9296448516 100644 --- a/examples/run/sycl/seq_example_sycl.sycl +++ b/examples/run/sycl/seq_example_sycl.sycl @@ -131,16 +131,13 @@ int seq_run(const traccc::opts::detector& detector_opts, copy(host_det_descr_data, device_det_descr)->wait(); // Construct a Detray detector object, if supported by the configuration. - traccc::default_detector::host host_detector{host_mr}; - traccc::default_detector::buffer device_detector; - traccc::default_detector::view device_detector_view; - traccc::io::read_detector( - host_detector, host_mr, detector_opts.detector_file, - detector_opts.material_file, detector_opts.grid_file); - device_detector = detray::get_buffer(host_detector, device_mr, copy); - q.wait_and_throw(); - const auto& const_device_detector = device_detector; - device_detector_view = detray::get_data(const_device_detector); + traccc::host_detector host_det; + traccc::io::read_detector(host_det, host_mr, detector_opts.detector_file, + detector_opts.material_file, + detector_opts.grid_file); + + const traccc::detector_buffer detector_buffer = + traccc::buffer_from_host_detector(host_det, device_mr, copy); // Output stats uint64_t n_cells = 0; @@ -278,7 +275,7 @@ int seq_run(const traccc::opts::detector& detector_opts, elapsedTimes); // Reconstruct it into spacepoints on the device. spacepoints_sycl_buffer = - sf_sycl(device_detector_view, measurements_sycl_buffer); + sf_sycl(detector_buffer, measurements_sycl_buffer); q.wait_and_throw(); } // stop measuring clusterization sycl timer @@ -287,7 +284,7 @@ int seq_run(const traccc::opts::detector& detector_opts, traccc::performance::timer t("Spacepoint formation (cpu)", elapsedTimes); spacepoints_per_event = - sf(host_detector, vecmem::get_data(measurements_per_event)); + sf(host_det, vecmem::get_data(measurements_per_event)); } // SYCL @@ -327,8 +324,8 @@ int seq_run(const traccc::opts::detector& detector_opts, traccc::performance::timer timer{"Track finding (sycl)", elapsedTimes}; track_candidates_sycl_buffer = finding_alg_sycl( - device_detector_view, device_field, - measurements_sycl_buffer, params_sycl_buffer); + detector_buffer, device_field, measurements_sycl_buffer, + params_sycl_buffer); q.wait_and_throw(); } @@ -337,7 +334,7 @@ int seq_run(const traccc::opts::detector& detector_opts, traccc::performance::timer timer{"Track finding (cpu)", elapsedTimes}; track_candidates = - finding_alg(host_detector, host_field, + finding_alg(host_det, host_field, vecmem::get_data(measurements_per_event), vecmem::get_data(params)); } @@ -426,8 +423,7 @@ int seq_run(const traccc::opts::detector& detector_opts, traccc::event_data evt_data(input_opts.directory, event, host_mr, input_opts.use_acts_geom_source, - &host_detector, input_opts.format, - true); + &host_det, input_opts.format, true); sd_performance_writer.write( vecmem::get_data(seeds_sycl), diff --git a/io/include/traccc/io/read_detector.hpp b/io/include/traccc/io/read_detector.hpp index 40fe089901..215c83251e 100644 --- a/io/include/traccc/io/read_detector.hpp +++ b/io/include/traccc/io/read_detector.hpp @@ -8,7 +8,7 @@ #pragma once // Project include(s). -#include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" // VecMem include(s). #include @@ -29,8 +29,7 @@ namespace traccc::io { /// @param material_file The file containing the material description /// @param grid_file The file containing the detector grid description /// -void read_detector(default_detector::host& detector, - vecmem::memory_resource& mr, +void read_detector(host_detector& detector, vecmem::memory_resource& mr, const std::string_view& geometry_file, const std::string_view& material_file = "", const std::string_view& grid_file = ""); diff --git a/io/include/traccc/io/read_measurements.hpp b/io/include/traccc/io/read_measurements.hpp index ceab32f49a..ac233a28ce 100644 --- a/io/include/traccc/io/read_measurements.hpp +++ b/io/include/traccc/io/read_measurements.hpp @@ -13,6 +13,7 @@ // Project include(s). #include "traccc/edm/measurement.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" // System include(s). #include @@ -34,8 +35,7 @@ namespace traccc::io { /// std::vector read_measurements( measurement_collection_types::host& measurements, std::size_t event, - std::string_view directory, - const traccc::default_detector::host* detector = nullptr, + std::string_view directory, const traccc::host_detector* detector = nullptr, const bool sort_measurements = true, data_format format = data_format::csv); /// Read measurement data into memory @@ -49,7 +49,7 @@ std::vector read_measurements( /// std::vector read_measurements( measurement_collection_types::host& measurements, std::string_view filename, - const traccc::default_detector::host* detector = nullptr, + const traccc::host_detector* detector = nullptr, const bool sort_measurements = true, data_format format = data_format::csv); } // namespace traccc::io diff --git a/io/include/traccc/io/read_particles.hpp b/io/include/traccc/io/read_particles.hpp index f383d23150..27773244df 100644 --- a/io/include/traccc/io/read_particles.hpp +++ b/io/include/traccc/io/read_particles.hpp @@ -13,6 +13,7 @@ // Project include(s). #include "traccc/edm/particle.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" // System include(s). #include @@ -62,7 +63,7 @@ void read_particles(particle_collection_types::host &particles, /// void read_particles(particle_container_types::host &particles, std::size_t event, std::string_view directory, - const traccc::default_detector::host *detector = nullptr, + const traccc::host_detector *detector = nullptr, data_format format = data_format::csv, std::string_view filename_postfix = "-particles_initial"); @@ -82,7 +83,7 @@ void read_particles(particle_container_types::host &particles, std::string_view particles_file, std::string_view hits_file, std::string_view measurements_file, std::string_view hit_map_file, - const traccc::default_detector::host *detector = nullptr, + const traccc::host_detector *detector = nullptr, data_format format = data_format::csv); } // namespace traccc::io diff --git a/io/include/traccc/io/read_spacepoints.hpp b/io/include/traccc/io/read_spacepoints.hpp index 7889dc623e..015cf53e5e 100644 --- a/io/include/traccc/io/read_spacepoints.hpp +++ b/io/include/traccc/io/read_spacepoints.hpp @@ -14,6 +14,7 @@ #include "traccc/edm/measurement.hpp" #include "traccc/edm/spacepoint_collection.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" // System include(s). #include @@ -36,7 +37,7 @@ namespace traccc::io { void read_spacepoints(edm::spacepoint_collection::host& spacepoints, measurement_collection_types::host& measurements, std::size_t event, std::string_view directory, - const traccc::default_detector::host* detector = nullptr, + const traccc::host_detector* detector = nullptr, data_format format = data_format::csv); /// Read spacepoint data into memory @@ -57,7 +58,7 @@ void read_spacepoints(edm::spacepoint_collection::host& spacepoints, std::string_view hit_filename, std::string_view meas_filename, std::string_view meas_hit_map_filename, - const traccc::default_detector::host* detector = nullptr, + const traccc::host_detector* detector = nullptr, data_format format = data_format::csv); } // namespace traccc::io diff --git a/io/include/traccc/io/write.hpp b/io/include/traccc/io/write.hpp index 6fbc452ccc..ed977c7487 100644 --- a/io/include/traccc/io/write.hpp +++ b/io/include/traccc/io/write.hpp @@ -14,6 +14,7 @@ #include "traccc/edm/spacepoint_collection.hpp" #include "traccc/edm/track_candidate_collection.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" #include "traccc/geometry/silicon_detector_description.hpp" #include "traccc/io/data_format.hpp" #include "traccc/io/digitization_config.hpp" @@ -88,7 +89,7 @@ void write(std::size_t event, std::string_view directory, traccc::data_format format, edm::track_candidate_collection::const_view tracks, measurement_collection_types::const_view measurements, - const traccc::default_detector::host& detector); + const traccc::host_detector& detector); /// Write a digitization configuration to a file /// diff --git a/io/src/csv/read_measurements.cpp b/io/src/csv/read_measurements.cpp index db53325e1f..0d5b4d798f 100644 --- a/io/src/csv/read_measurements.cpp +++ b/io/src/csv/read_measurements.cpp @@ -18,7 +18,7 @@ namespace traccc::io::csv { std::vector read_measurements( measurement_collection_types::host& measurements, std::string_view filename, - const traccc::default_detector::host* detector, const bool do_sort) { + const traccc::host_detector* detector, const bool do_sort) { // Construct the measurement reader object. auto reader = make_measurement_reader(filename); @@ -27,10 +27,13 @@ std::vector read_measurements( std::map acts_to_detray_id; if (detector) { - for (const auto& surface_desc : detector->surfaces()) { - acts_to_detray_id[surface_desc.source] = - surface_desc.barcode().value(); - } + host_detector_visitor( + *detector, [&](const detector_t::host& det) { + for (const auto& surface_desc : det.surfaces()) { + acts_to_detray_id[surface_desc.source] = + surface_desc.barcode().value(); + } + }); } // Read the measurements from the input file. diff --git a/io/src/csv/read_measurements.hpp b/io/src/csv/read_measurements.hpp index bfab7326ed..76c5f56edc 100644 --- a/io/src/csv/read_measurements.hpp +++ b/io/src/csv/read_measurements.hpp @@ -11,6 +11,7 @@ #include "traccc/definitions/primitives.hpp" #include "traccc/edm/measurement.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" // System include(s). #include @@ -27,7 +28,6 @@ namespace traccc::io::csv { /// std::vector read_measurements( measurement_collection_types::host& measurements, std::string_view filename, - const traccc::default_detector::host* detector = nullptr, - const bool do_sort = true); + const traccc::host_detector* detector = nullptr, const bool do_sort = true); } // namespace traccc::io::csv diff --git a/io/src/csv/read_particles.cpp b/io/src/csv/read_particles.cpp index b8e9fb8085..5cc01f3f2a 100644 --- a/io/src/csv/read_particles.cpp +++ b/io/src/csv/read_particles.cpp @@ -44,7 +44,7 @@ void read_particles(particle_container_types::host& particles, std::string_view particles_file, std::string_view hits_file, std::string_view measurements_file, std::string_view hit_map_file, - const traccc::default_detector::host* detector, + const traccc::host_detector* detector, const bool sort_measurements) { // Memory resource used by the temporary collections. diff --git a/io/src/csv/read_particles.hpp b/io/src/csv/read_particles.hpp index 931ac7821d..8a4b9b3d4a 100644 --- a/io/src/csv/read_particles.hpp +++ b/io/src/csv/read_particles.hpp @@ -10,6 +10,7 @@ // Project include(s). #include "traccc/edm/particle.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" // System include(s). #include @@ -37,7 +38,7 @@ void read_particles(particle_container_types::host& particles, std::string_view particles_file, std::string_view hits_file, std::string_view measurements_file, std::string_view hit_map_file, - const traccc::default_detector::host* detector, + const traccc::host_detector* detector, const bool sort_measurements = true); } // namespace traccc::io::csv diff --git a/io/src/csv/read_spacepoints.cpp b/io/src/csv/read_spacepoints.cpp index 306081f7dd..185c4bd40a 100644 --- a/io/src/csv/read_spacepoints.cpp +++ b/io/src/csv/read_spacepoints.cpp @@ -27,7 +27,7 @@ void read_spacepoints(edm::spacepoint_collection::host& spacepoints, std::string_view hit_filename, std::string_view meas_filename, std::string_view meas_hit_map_filename, - const traccc::default_detector::host* detector, + const traccc::host_detector* detector, const bool sort_measurements) { // Read all measurements. diff --git a/io/src/csv/read_spacepoints.hpp b/io/src/csv/read_spacepoints.hpp index 3d673e7172..2bf5676e56 100644 --- a/io/src/csv/read_spacepoints.hpp +++ b/io/src/csv/read_spacepoints.hpp @@ -11,6 +11,7 @@ #include "traccc/edm/measurement.hpp" #include "traccc/edm/spacepoint_collection.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" // System include(s). #include @@ -32,7 +33,7 @@ void read_spacepoints(edm::spacepoint_collection::host& spacepoints, std::string_view hit_filename, std::string_view meas_filename, std::string_view meas_hit_map_filename, - const traccc::default_detector::host* detector = nullptr, + const traccc::host_detector* detector = nullptr, const bool sort_measurements = true); } // namespace traccc::io::csv diff --git a/io/src/obj/write_track_candidates.cpp b/io/src/obj/write_track_candidates.cpp index eb5108587a..5aaecbb2eb 100644 --- a/io/src/obj/write_track_candidates.cpp +++ b/io/src/obj/write_track_candidates.cpp @@ -21,7 +21,7 @@ void write_track_candidates( std::string_view filename, edm::track_candidate_collection::const_view tracks_view, measurement_collection_types::const_view measurements_view, - const traccc::default_detector::host& detector) { + const traccc::host_detector& detector) { // Open the output file. std::ofstream file{filename.data()}; @@ -58,10 +58,12 @@ void write_track_candidates( const measurement& m = measurements.at(midx); // Find the detector surface that this measurement sits on. - const detray::tracking_surface surface{detector, m.surface_link}; - - // Calculate a position for this measurement in global 3D space. - const auto global = surface.local_to_global({}, m.local, {}); + const auto global = host_detector_visitor( + detector, [&m]( + const typename detector_traits_t::host& d) { + detray::tracking_surface surface{d, m.surface_link}; + return surface.local_to_global({}, m.local, {}); + }); // Write the 3D coordinates of the measurement / spacepoint. assert(global.size() == 3); @@ -70,8 +72,8 @@ void write_track_candidates( } } - // Now loop over the track candidates again, and creates lines for each of - // them using the measurements / spacepoints written out earlier. + // Now loop over the track candidates again, and creates lines for each + // of them using the measurements / spacepoints written out earlier. file << "# Track candidates\n"; std::size_t vertex_counter = 1; for (size_type i = 0; i < tracks.size(); ++i) { diff --git a/io/src/obj/write_track_candidates.hpp b/io/src/obj/write_track_candidates.hpp index b316609257..b97b230eb7 100644 --- a/io/src/obj/write_track_candidates.hpp +++ b/io/src/obj/write_track_candidates.hpp @@ -11,6 +11,7 @@ #include "traccc/edm/measurement.hpp" #include "traccc/edm/track_candidate_collection.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" // System include(s). #include @@ -28,6 +29,6 @@ void write_track_candidates( std::string_view filename, edm::track_candidate_collection::const_view tracks, measurement_collection_types::const_view measurements, - const traccc::default_detector::host& detector); + const traccc::host_detector& detector); } // namespace traccc::io::obj diff --git a/io/src/read_detector.cpp b/io/src/read_detector.cpp index 880f86dcc1..9ef1075aa3 100644 --- a/io/src/read_detector.cpp +++ b/io/src/read_detector.cpp @@ -8,11 +8,14 @@ // Local include(s). #include "traccc/io/read_detector.hpp" +#include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" #include "traccc/io/utils.hpp" // Detray include(s). #include #include +#include // System include(s). #include @@ -21,7 +24,7 @@ namespace { /// Common implementation for constructing a detector from a set of input files template -void read_detector(detector_t& detector, vecmem::memory_resource& mr, +void read_detector(traccc::host_detector& detector, vecmem::memory_resource& mr, const std::string_view& geometry_file, const std::string_view& material_file, const std::string_view& grid_file) { @@ -37,21 +40,31 @@ void read_detector(detector_t& detector, vecmem::memory_resource& mr, } // Read the detector. - auto det = detray::io::read_detector(mr, cfg); - detector = std::move(det.first); + auto det = detray::io::read_detector(mr, cfg); + detector.set(std::move(det.first)); } } // namespace namespace traccc::io { -void read_detector(default_detector::host& detector, - vecmem::memory_resource& mr, +void read_detector(host_detector& detector, vecmem::memory_resource& mr, const std::string_view& geometry_file, const std::string_view& material_file, const std::string_view& grid_file) { - ::read_detector(detector, mr, geometry_file, material_file, grid_file); + // Peek at the header to determine the kind of reader that is needed + const auto header = detray::io::detail::deserialize_json_header( + traccc::io::get_absolute_path(geometry_file)); + + if (header.detector == "Cylindrical detector from DD4hep blueprint") { + ::read_detector(detector, mr, geometry_file, + material_file, grid_file); + } else { + // TODO: Warning here + ::read_detector(detector, mr, geometry_file, + material_file, grid_file); + } } } // namespace traccc::io diff --git a/io/src/read_detector_description.cpp b/io/src/read_detector_description.cpp index 6ca04ebfe0..fd7cfd2e81 100644 --- a/io/src/read_detector_description.cpp +++ b/io/src/read_detector_description.cpp @@ -8,6 +8,7 @@ // Library include(s). #include "traccc/io/read_detector_description.hpp" +#include "traccc/geometry/host_detector.hpp" #include "traccc/io/read_detector.hpp" #include "traccc/io/read_digitization_config.hpp" #include "traccc/io/utils.hpp" @@ -41,21 +42,20 @@ void fill_digi_info(traccc::silicon_detector_description::host& dd, dd.dimensions().back() = data.dimensions; } -void read_json_dd(traccc::silicon_detector_description::host& dd, - std::string_view geometry_file, - const traccc::digitization_config& digi) { - - // Construct a (temporary) Detray detector object from the geometry - // configuration file. - vecmem::host_memory_resource mr; - traccc::default_detector::host detector{mr}; - traccc::io::read_detector(detector, mr, geometry_file); +template +void read_json_dd_impl(traccc::silicon_detector_description::host& dd, + const traccc::host_detector& detector, + const traccc::digitization_config& digi) + requires(traccc::is_detector_traits) +{ + const traccc::default_detector::host& detector_host = + detector.as(); // Iterate over the surfaces of the detector. - const traccc::default_detector::host::surface_lookup_container& surfaces = - detector.surfaces(); + const typename detector_traits_t::host::surface_lookup_container& surfaces = + detector_host.surfaces(); dd.reserve(surfaces.size()); - for (const auto& surface_desc : detector.surfaces()) { + for (const auto& surface_desc : detector_host.surfaces()) { // Acts geometry identifier(s) for the surface. const traccc::geometry_id geom_id{surface_desc.source}; @@ -71,7 +71,7 @@ void read_json_dd(traccc::silicon_detector_description::host& dd, dd.resize(dd.size() + 1); // Construct a Detray surface object. - const detray::tracking_surface surface{detector, surface_desc}; + const detray::tracking_surface surface{detector_host, surface_desc}; // Fill the new element with the geometry ID and the transformation of // the surface in question. @@ -83,8 +83,7 @@ void read_json_dd(traccc::silicon_detector_description::host& dd, using annulus_t = detray::mask; if (surface_desc.mask().id() == - traccc::default_detector::host::masks::template get_id< - annulus_t>()) { + detector_traits_t::host::masks::template get_id()) { dd.subspace().back() = {1, 0}; } @@ -104,6 +103,20 @@ void read_json_dd(traccc::silicon_detector_description::host& dd, } } +void read_json_dd(traccc::silicon_detector_description::host& dd, + std::string_view geometry_file, + const traccc::digitization_config& digi) { + + // Construct a (temporary) Detray detector object from the geometry + // configuration file. + vecmem::host_memory_resource mr; + traccc::host_detector detector; + traccc::io::read_detector(detector, mr, geometry_file); + + read_json_dd_impl(dd, detector, digi); + // detector_buffer_visitor +} + } // namespace namespace traccc::io { diff --git a/io/src/read_measurements.cpp b/io/src/read_measurements.cpp index 4271116a96..e6ac4609bb 100644 --- a/io/src/read_measurements.cpp +++ b/io/src/read_measurements.cpp @@ -19,7 +19,7 @@ namespace traccc::io { std::vector read_measurements( measurement_collection_types::host& measurements, std::size_t event, - std::string_view directory, const traccc::default_detector::host* detector, + std::string_view directory, const traccc::host_detector* detector, const bool sort_measurements, data_format format) { switch (format) { @@ -49,8 +49,8 @@ std::vector read_measurements( std::vector read_measurements( measurement_collection_types::host& measurements, std::string_view filename, - const traccc::default_detector::host* detector, - const bool sort_measurements, data_format format) { + const traccc::host_detector* detector, const bool sort_measurements, + data_format format) { switch (format) { case data_format::csv: diff --git a/io/src/read_particles.cpp b/io/src/read_particles.cpp index a3f0503de4..15cf661dc7 100644 --- a/io/src/read_particles.cpp +++ b/io/src/read_particles.cpp @@ -50,8 +50,8 @@ void read_particles(particle_collection_types::host& particles, void read_particles(particle_container_types::host& particles, std::size_t event, std::string_view directory, - const traccc::default_detector::host* detector, - data_format format, std::string_view filename_postfix) { + const traccc::host_detector* detector, data_format format, + std::string_view filename_postfix) { switch (format) { case data_format::csv: @@ -85,8 +85,7 @@ void read_particles(particle_container_types::host& particles, std::string_view particles_file, std::string_view hits_file, std::string_view measurements_file, std::string_view hit_map_file, - const traccc::default_detector::host* detector, - data_format format) { + const traccc::host_detector* detector, data_format format) { switch (format) { case data_format::csv: diff --git a/io/src/read_spacepoints.cpp b/io/src/read_spacepoints.cpp index b8c12fda2d..d19161e8dd 100644 --- a/io/src/read_spacepoints.cpp +++ b/io/src/read_spacepoints.cpp @@ -20,7 +20,7 @@ namespace traccc::io { void read_spacepoints(edm::spacepoint_collection::host& spacepoints, measurement_collection_types::host& measurements, std::size_t event, std::string_view directory, - const traccc::default_detector::host* detector, + const traccc::host_detector* detector, data_format format) { switch (format) { @@ -67,7 +67,7 @@ void read_spacepoints(edm::spacepoint_collection::host& spacepoints, std::string_view hit_filename, std::string_view meas_filename, std::string_view meas_hit_map_filename, - const traccc::default_detector::host* detector, + const traccc::host_detector* detector, data_format format) { switch (format) { diff --git a/io/src/write.cpp b/io/src/write.cpp index f47667f83e..493ea5f833 100644 --- a/io/src/write.cpp +++ b/io/src/write.cpp @@ -13,6 +13,7 @@ #include "obj/write_seeds.hpp" #include "obj/write_spacepoints.hpp" #include "obj/write_track_candidates.hpp" +#include "traccc/geometry/host_detector.hpp" #include "traccc/io/utils.hpp" #include "write_binary.hpp" @@ -125,7 +126,7 @@ void write(std::size_t event, std::string_view directory, traccc::data_format format, edm::track_candidate_collection::const_view tracks, measurement_collection_types::const_view measurements, - const traccc::default_detector::host& detector) { + const traccc::host_detector& detector) { switch (format) { case data_format::obj: diff --git a/performance/include/traccc/utils/event_data.hpp b/performance/include/traccc/utils/event_data.hpp index f784942edc..033cb270c9 100644 --- a/performance/include/traccc/utils/event_data.hpp +++ b/performance/include/traccc/utils/event_data.hpp @@ -14,6 +14,7 @@ #include "traccc/edm/silicon_cluster_collection.hpp" #include "traccc/edm/track_candidate_container.hpp" #include "traccc/geometry/detector.hpp" +#include "traccc/geometry/host_detector.hpp" #include "traccc/geometry/silicon_detector_description.hpp" #include "traccc/io/csv/cell.hpp" #include "traccc/io/csv/hit.hpp" @@ -35,9 +36,6 @@ namespace traccc { struct event_data { public: - // Type definitions - using detector_type = traccc::default_detector::host; - event_data() = delete; /// Event data constructor @@ -54,7 +52,7 @@ struct event_data { event_data(const std::string& event_dir, const std::size_t event_id, vecmem::memory_resource& resource, bool use_acts_geom_source = false, - const detector_type* det = nullptr, + const host_detector* det = nullptr, data_format format = data_format::csv, bool include_silicon_cells = false); @@ -77,10 +75,46 @@ struct event_data { /// @param[in] sg Seed generator for fitting /// @param[in] resource vecmem memory resource /// + template void generate_truth_candidates( edm::track_candidate_container::host& truth_candidates, seed_generator& sg, vecmem::memory_resource& resource, - float pt_cut = 0.f); + float pt_cut = 0.f) { + for (auto const& [ptc, measurements] : m_ptc_to_meas_map) { + + const auto& param = m_meas_to_param_map.at(measurements[0]); + const free_track_parameters<> free_param(param.first, 0.f, + param.second, ptc.charge); + + auto ptc_particle = + detail::particle_from_pdg_number(ptc.particle_type); + + if (ptc_particle.pdg_num() == 0) { + // TODO: Add some debug logging here. + continue; + } else if (free_param.pT(ptc_particle.charge()) <= pt_cut) { + continue; + } + + auto seed_params = + sg(measurements[0].surface_link, free_param, ptc_particle); + + // Record the measurements, and remember their indices. + vecmem::vector meas_indices{&resource}; + truth_candidates.measurements.reserve( + truth_candidates.measurements.size() + measurements.size()); + meas_indices.reserve(measurements.size()); + for (const auto& meas : measurements) { + meas_indices.push_back(static_cast( + truth_candidates.measurements.size())); + truth_candidates.measurements.push_back(meas); + } + + // Record the truth track candidate. + truth_candidates.tracks.push_back( + {seed_params, 0.f, 0.f, 0.f, 0u, meas_indices}); + } + } // Measurement map std::map m_measurement_map; @@ -106,7 +140,7 @@ struct event_data { std::reference_wrapper m_mr; private: - void setup_csv(bool use_acts_geom_source, const detector_type* det, + void setup_csv(bool use_acts_geom_source, const host_detector* det, bool include_silicon_cells); }; diff --git a/performance/src/utils/event_data.cpp b/performance/src/utils/event_data.cpp index d046fc428c..1d803b209e 100644 --- a/performance/src/utils/event_data.cpp +++ b/performance/src/utils/event_data.cpp @@ -28,7 +28,7 @@ namespace traccc { event_data::event_data(const std::string& event_dir, const std::size_t event_id, vecmem::memory_resource& resource, - bool use_acts_geom_source, const detector_type* det, + bool use_acts_geom_source, const host_detector* det, data_format format, bool include_silicon_cells) : m_event_dir(event_dir), m_event_id(event_id), m_mr(resource) { @@ -39,7 +39,7 @@ event_data::event_data(const std::string& event_dir, const std::size_t event_id, } } -void event_data::setup_csv(bool use_acts_geom_source, const detector_type* det, +void event_data::setup_csv(bool use_acts_geom_source, const host_detector* det, bool include_silicon_cells) { /******************** @@ -136,10 +136,14 @@ void event_data::setup_csv(bool use_acts_geom_source, const detector_type* det, // For Acts data, build a map of acts->detray geometry IDs std::map acts_to_detray_id; if (use_acts_geom_source) { - for (const auto& surface_desc : det->surfaces()) { - acts_to_detray_id[surface_desc.source] = - surface_desc.barcode().value(); - } + host_detector_visitor( + *det, [&acts_to_detray_id]( + const typename detector_traits_t::host& d) { + for (const auto& surface_desc : d.surfaces()) { + acts_to_detray_id[surface_desc.source] = + surface_desc.barcode().value(); + } + }); } /******************** @@ -369,46 +373,4 @@ void event_data::fill_cca_result( } } } - -void event_data::generate_truth_candidates( - edm::track_candidate_container::host& truth_candidates, - seed_generator& sg, vecmem::memory_resource& resource, - float pt_cut) { - - for (auto const& [ptc, measurements] : m_ptc_to_meas_map) { - - const auto& param = m_meas_to_param_map.at(measurements[0]); - const free_track_parameters<> free_param(param.first, 0.f, param.second, - ptc.charge); - - auto ptc_particle = - detail::particle_from_pdg_number(ptc.particle_type); - - if (ptc_particle.pdg_num() == 0) { - // TODO: Add some debug logging here. - continue; - } else if (free_param.pT(ptc_particle.charge()) <= pt_cut) { - continue; - } - - auto seed_params = - sg(measurements[0].surface_link, free_param, ptc_particle); - - // Record the measurements, and remember their indices. - vecmem::vector meas_indices{&resource}; - truth_candidates.measurements.reserve( - truth_candidates.measurements.size() + measurements.size()); - meas_indices.reserve(measurements.size()); - for (const auto& meas : measurements) { - meas_indices.push_back(static_cast( - truth_candidates.measurements.size())); - truth_candidates.measurements.push_back(meas); - } - - // Record the truth track candidate. - truth_candidates.tracks.push_back( - {seed_params, 0.f, 0.f, 0.f, 0u, meas_indices}); - } -} - } // namespace traccc diff --git a/tests/common/tests/kalman_fitting_test.hpp b/tests/common/tests/kalman_fitting_test.hpp index 9473b54655..7b271d66dd 100644 --- a/tests/common/tests/kalman_fitting_test.hpp +++ b/tests/common/tests/kalman_fitting_test.hpp @@ -35,8 +35,9 @@ namespace traccc { class KalmanFittingTests : public testing::Test { public: /// Type declarations - using host_detector_type = traccc::default_detector::host; - using device_detector_type = traccc::default_detector::device; + using detector_traits = traccc::default_detector; + using host_detector_type = detector_traits::host; + using device_detector_type = detector_traits::device; using scalar_type = device_detector_type::scalar_type; using b_field_t = diff --git a/tests/cpu/test_ckf_combinatorics_telescope.cpp b/tests/cpu/test_ckf_combinatorics_telescope.cpp index c248f34927..3a13db1230 100644 --- a/tests/cpu/test_ckf_combinatorics_telescope.cpp +++ b/tests/cpu/test_ckf_combinatorics_telescope.cpp @@ -9,6 +9,7 @@ #include "traccc/bfield/construct_const_bfield.hpp" #include "traccc/bfield/magnetic_field.hpp" #include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" +#include "traccc/io/read_detector.hpp" #include "traccc/io/read_measurements.hpp" #include "traccc/io/utils.hpp" #include "traccc/resolution/fitting_performance_writer.hpp" @@ -58,12 +59,16 @@ TEST_P(CpuCkfCombinatoricsTelescopeTests, Run) { // Read back detector file const std::string path = name + "/"; - detray::io::detector_reader_config reader_cfg{}; - reader_cfg.add_file(path + "telescope_detector_geometry.json") - .add_file(path + "telescope_detector_homogeneous_material.json"); - - const auto [host_det, names] = - detray::io::read_detector(host_mr, reader_cfg); + traccc::host_detector detector; + traccc::io::read_detector( + detector, host_mr, + std::filesystem::absolute( + std::filesystem::path(path + "telescope_detector_geometry.json")) + .native(), + std::filesystem::absolute( + std::filesystem::path( + path + "telescope_detector_homogeneous_material.json")) + .native()); const auto field = traccc::construct_const_bfield(std::get<13>(GetParam())); @@ -99,7 +104,7 @@ TEST_P(CpuCkfCombinatoricsTelescopeTests, Run) { std::filesystem::create_directories(full_path); auto sim = traccc::simulator( - std::get<6>(GetParam()), n_events, host_det, + std::get<6>(GetParam()), n_events, detector.as(), field.as_field>(), std::move(generator), std::move(smearer_writer_cfg), full_path); sim.run(); @@ -109,7 +114,8 @@ TEST_P(CpuCkfCombinatoricsTelescopeTests, Run) { *****************************/ // Seed generator - seed_generator sg(host_det, stddevs); + seed_generator sg(detector.as(), + stddevs); // Finding algorithm configuration traccc::finding_config cfg_no_limit; @@ -161,10 +167,10 @@ TEST_P(CpuCkfCombinatoricsTelescopeTests, Run) { // Run finding auto track_candidates = - host_finding(host_det, field, measurements_view, seeds_view); + host_finding(detector, field, measurements_view, seeds_view); auto track_candidates_limit = - host_finding_limit(host_det, field, measurements_view, seeds_view); + host_finding_limit(detector, field, measurements_view, seeds_view); // Make sure that the number of found tracks = n_track ^ (n_planes + 1) ASSERT_TRUE(track_candidates.size() > track_candidates_limit.size()); diff --git a/tests/cpu/test_ckf_sparse_tracks_telescope.cpp b/tests/cpu/test_ckf_sparse_tracks_telescope.cpp index 019ff36392..5e6e402931 100644 --- a/tests/cpu/test_ckf_sparse_tracks_telescope.cpp +++ b/tests/cpu/test_ckf_sparse_tracks_telescope.cpp @@ -9,6 +9,7 @@ #include "traccc/bfield/construct_const_bfield.hpp" #include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" +#include "traccc/io/read_detector.hpp" #include "traccc/io/read_measurements.hpp" #include "traccc/io/utils.hpp" #include "traccc/resolution/fitting_performance_writer.hpp" @@ -69,12 +70,16 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { // Read back detector file const std::string path = name + "/"; - detray::io::detector_reader_config reader_cfg{}; - reader_cfg.add_file(path + "telescope_detector_geometry.json") - .add_file(path + "telescope_detector_homogeneous_material.json"); - - const auto [host_det, names] = - detray::io::read_detector(host_mr, reader_cfg); + traccc::host_detector detector; + traccc::io::read_detector( + detector, host_mr, + std::filesystem::absolute( + std::filesystem::path(path + "telescope_detector_geometry.json")) + .native(), + std::filesystem::absolute( + std::filesystem::path( + path + "telescope_detector_homogeneous_material.json")) + .native()); auto field = traccc::construct_const_bfield(std::get<13>(GetParam())); @@ -110,7 +115,7 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { std::filesystem::create_directories(full_path); auto sim = traccc::simulator( - ptc, n_events, host_det, + ptc, n_events, detector.as(), field.as_field>(), std::move(generator), std::move(smearer_writer_cfg), full_path); @@ -121,7 +126,8 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { *****************************/ // Seed generator - seed_generator sg(host_det, stddevs); + seed_generator sg(detector.as(), + stddevs); // Finding algorithm configuration typename traccc::finding_config cfg; @@ -165,7 +171,7 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { // Run finding auto track_candidates = host_finding( - host_det, field, vecmem::get_data(measurements_per_event), + detector, field, vecmem::get_data(measurements_per_event), vecmem::get_data(seeds)); ASSERT_EQ(track_candidates.size(), n_truth_tracks); @@ -179,7 +185,7 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { // Run fitting auto track_states = - host_fitting(host_det, field, + host_fitting(detector, field, {vecmem::get_data(track_candidates), vecmem::get_data(measurements_per_event)}); const std::size_t n_fitted_tracks = @@ -198,7 +204,8 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { fit_performance_writer.write( track_states.tracks.at(i_trk), track_states.states, - measurements_per_event, host_det, evt_data); + measurements_per_event, detector.as(), + evt_data); } } diff --git a/tests/cpu/test_clusterization_resolution.cpp b/tests/cpu/test_clusterization_resolution.cpp index ba0353fa5a..dca63f9bcb 100644 --- a/tests/cpu/test_clusterization_resolution.cpp +++ b/tests/cpu/test_clusterization_resolution.cpp @@ -45,7 +45,7 @@ TEST_P(SurfaceBinningTests, Run) { vecmem::get_data(dd); // Read the detector - traccc::default_detector::host detector{host_mr}; + traccc::host_detector detector; traccc::io::read_detector(detector, host_mr, detector_file); // Algorithms diff --git a/tests/cpu/test_kalman_fitter_hole_count.cpp b/tests/cpu/test_kalman_fitter_hole_count.cpp index b88f8dafa9..54fd41c861 100644 --- a/tests/cpu/test_kalman_fitter_hole_count.cpp +++ b/tests/cpu/test_kalman_fitter_hole_count.cpp @@ -8,6 +8,7 @@ // Project include(s). #include "traccc/bfield/construct_const_bfield.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" +#include "traccc/io/read_detector.hpp" #include "traccc/io/utils.hpp" #include "traccc/resolution/fitting_performance_writer.hpp" #include "traccc/simulation/event_generators.hpp" @@ -66,12 +67,16 @@ TEST_P(KalmanFittingHoleCountTests, Run) { // Read back detector file const std::string path = name + "/"; - detray::io::detector_reader_config reader_cfg{}; - reader_cfg.add_file(path + "telescope_detector_geometry.json") - .add_file(path + "telescope_detector_homogeneous_material.json"); - - const auto [host_det, names] = - detray::io::read_detector(host_mr, reader_cfg); + traccc::host_detector detector; + traccc::io::read_detector( + detector, host_mr, + std::filesystem::absolute( + std::filesystem::path(path + "telescope_detector_geometry.json")) + .native(), + std::filesystem::absolute( + std::filesystem::path( + path + "telescope_detector_homogeneous_material.json")) + .native()); auto field = traccc::construct_const_bfield(std::get<13>(GetParam())); /*************************** @@ -106,7 +111,7 @@ TEST_P(KalmanFittingHoleCountTests, Run) { std::filesystem::create_directories(full_path); auto sim = traccc::simulator( - ptc, n_events, host_det, + ptc, n_events, detector.as(), field.as_field>(), std::move(generator), std::move(smearer_writer_cfg), full_path); sim.run(); @@ -116,7 +121,8 @@ TEST_P(KalmanFittingHoleCountTests, Run) { ***************/ // Seed generator - seed_generator sg(host_det, stddevs); + seed_generator sg(detector.as(), + stddevs); // Fitting algorithm object traccc::fitting_config fit_cfg; @@ -155,7 +161,7 @@ TEST_P(KalmanFittingHoleCountTests, Run) { // Run fitting auto track_states = - fitting(host_det, field, + fitting(detector, field, {vecmem::get_data(track_candidates.tracks), vecmem::get_data(track_candidates.measurements)}); diff --git a/tests/cpu/test_kalman_fitter_momentum_resolution.cpp b/tests/cpu/test_kalman_fitter_momentum_resolution.cpp index ea1b2b8c57..310b7efbd1 100644 --- a/tests/cpu/test_kalman_fitter_momentum_resolution.cpp +++ b/tests/cpu/test_kalman_fitter_momentum_resolution.cpp @@ -8,6 +8,7 @@ // Project include(s). #include "traccc/bfield/construct_const_bfield.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" +#include "traccc/io/read_detector.hpp" #include "traccc/io/utils.hpp" #include "traccc/resolution/fitting_performance_writer.hpp" #include "traccc/simulation/event_generators.hpp" @@ -74,19 +75,23 @@ TEST_P(KalmanFittingMomentumResolutionTests, Run) { // Read back detector file const std::string path = name + "/"; - detray::io::detector_reader_config reader_cfg{}; - reader_cfg.add_file(path + "telescope_detector_geometry.json"); - // If the module material is not vacuum, read the material map - if (std::get<14>(GetParam()) != detray::vacuum()) { - reader_cfg.add_file(path + - "telescope_detector_homogeneous_material.json"); - } + traccc::host_detector detector; + traccc::io::read_detector( + detector, host_mr, + std::filesystem::absolute( + std::filesystem::path(path + "telescope_detector_geometry.json")) + .native(), + (std::get<14>(GetParam()) != detray::vacuum() + ? std::filesystem::absolute( + std::filesystem::path( + path + "telescope_detector_homogeneous_material.json")) + .native() + : "")); - const auto [host_det, names] = - detray::io::read_detector(host_mr, reader_cfg); auto field = traccc::construct_const_bfield(std::get<13>(GetParam())); - const auto vol0 = detray::tracking_volume{host_det, 0u}; + const auto vol0 = + detray::tracking_volume{detector.as(), 0u}; // The number of sensitive surfaces = # of total surfaces - # of portals // (=6) @@ -127,7 +132,7 @@ TEST_P(KalmanFittingMomentumResolutionTests, Run) { std::filesystem::create_directories(full_path); auto sim = traccc::simulator( - ptc, n_events, host_det, + ptc, n_events, detector.as(), field.as_field>(), std::move(generator), std::move(smearer_writer_cfg), full_path); sim.run(); @@ -137,7 +142,8 @@ TEST_P(KalmanFittingMomentumResolutionTests, Run) { ***************/ // Seed generator - seed_generator sg(host_det, stddevs); + seed_generator sg(detector.as(), + stddevs); // Fitting algorithm object traccc::fitting_config fit_cfg; @@ -167,7 +173,7 @@ TEST_P(KalmanFittingMomentumResolutionTests, Run) { // Run fitting auto track_states = - fitting(host_det, field, + fitting(detector, field, {vecmem::get_data(track_candidates.tracks), vecmem::get_data(track_candidates.measurements)}); @@ -200,7 +206,8 @@ TEST_P(KalmanFittingMomentumResolutionTests, Run) { fit_performance_writer.write( track_states.tracks.at(i_trk), track_states.states, - track_candidates.measurements, host_det, evt_data); + track_candidates.measurements, detector.as(), + evt_data); } } diff --git a/tests/cpu/test_kalman_fitter_telescope.cpp b/tests/cpu/test_kalman_fitter_telescope.cpp index bd29ae11f6..45b937314f 100644 --- a/tests/cpu/test_kalman_fitter_telescope.cpp +++ b/tests/cpu/test_kalman_fitter_telescope.cpp @@ -8,6 +8,7 @@ // Project include(s). #include "traccc/bfield/construct_const_bfield.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" +#include "traccc/io/read_detector.hpp" #include "traccc/io/utils.hpp" #include "traccc/resolution/fitting_performance_writer.hpp" #include "traccc/simulation/event_generators.hpp" @@ -66,12 +67,17 @@ TEST_P(KalmanFittingTelescopeTests, Run) { // Read back detector file const std::string path = name + "/"; - detray::io::detector_reader_config reader_cfg{}; - reader_cfg.add_file(path + "telescope_detector_geometry.json") - .add_file(path + "telescope_detector_homogeneous_material.json"); - const auto [host_det, names] = - detray::io::read_detector(host_mr, reader_cfg); + traccc::host_detector detector; + traccc::io::read_detector( + detector, host_mr, + std::filesystem::absolute( + std::filesystem::path(path + "telescope_detector_geometry.json")) + .native(), + std::filesystem::absolute( + std::filesystem::path( + path + "telescope_detector_homogeneous_material.json")) + .native()); auto field = traccc::construct_const_bfield(std::get<13>(GetParam())); /*************************** @@ -106,7 +112,7 @@ TEST_P(KalmanFittingTelescopeTests, Run) { std::filesystem::create_directories(full_path); auto sim = traccc::simulator( - ptc, n_events, host_det, + ptc, n_events, detector.as(), field.as_field>(), std::move(generator), std::move(smearer_writer_cfg), full_path); sim.run(); @@ -116,7 +122,8 @@ TEST_P(KalmanFittingTelescopeTests, Run) { ***************/ // Seed generator - seed_generator sg(host_det, stddevs); + seed_generator sg(detector.as(), + stddevs); // Fitting algorithm object traccc::fitting_config fit_cfg; @@ -138,7 +145,7 @@ TEST_P(KalmanFittingTelescopeTests, Run) { // Run fitting auto track_states = - fitting(host_det, field, + fitting(detector, field, {vecmem::get_data(track_candidates.tracks), vecmem::get_data(track_candidates.measurements)}); @@ -166,7 +173,8 @@ TEST_P(KalmanFittingTelescopeTests, Run) { fit_performance_writer.write( track_states.tracks.at(i_trk), track_states.states, - track_candidates.measurements, host_det, evt_data); + track_candidates.measurements, detector.as(), + evt_data); } } diff --git a/tests/cpu/test_kalman_fitter_wire_chamber.cpp b/tests/cpu/test_kalman_fitter_wire_chamber.cpp index 733b362522..e3fd187124 100644 --- a/tests/cpu/test_kalman_fitter_wire_chamber.cpp +++ b/tests/cpu/test_kalman_fitter_wire_chamber.cpp @@ -8,6 +8,7 @@ // Project include(s). #include "traccc/bfield/construct_const_bfield.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" +#include "traccc/io/read_detector.hpp" #include "traccc/io/utils.hpp" #include "traccc/resolution/fitting_performance_writer.hpp" #include "traccc/simulation/event_generators.hpp" @@ -62,14 +63,19 @@ TEST_P(KalmanFittingWireChamberTests, Run) { // Read back detector file const std::string path = name + "/"; - detray::io::detector_reader_config reader_cfg{}; - reader_cfg.add_file(path + "wire_chamber_geometry.json") - .add_file(path + "wire_chamber_homogeneous_material.json") - .add_file(path + "wire_chamber_surface_grids.json") - .do_check(true); - - const auto [host_det, names] = - detray::io::read_detector(host_mr, reader_cfg); + traccc::host_detector detector; + traccc::io::read_detector( + detector, host_mr, + std::filesystem::absolute( + std::filesystem::path(path + "wire_chamber_geometry.json")) + .native(), + std::filesystem::absolute( + std::filesystem::path(path + + "wire_chamber_homogeneous_material.json")) + .native(), + std::filesystem::absolute( + std::filesystem::path(path + "wire_chamber_surface_grids.json")) + .native()); const auto field = traccc::construct_const_bfield(B); /*************************** @@ -105,7 +111,7 @@ TEST_P(KalmanFittingWireChamberTests, Run) { std::filesystem::create_directories(full_path); auto sim = traccc::simulator( - ptc, n_events, host_det, + ptc, n_events, detector.as(), field.as_field>(), std::move(generator), std::move(smearer_writer_cfg), full_path); @@ -118,7 +124,8 @@ TEST_P(KalmanFittingWireChamberTests, Run) { ***************/ // Seed generator - seed_generator sg(host_det, stddevs); + seed_generator sg(detector.as(), + stddevs); // Fitting algorithm object traccc::fitting_config fit_cfg; @@ -143,7 +150,7 @@ TEST_P(KalmanFittingWireChamberTests, Run) { // Run fitting auto track_states = - fitting(host_det, field, + fitting(detector, field, {vecmem::get_data(track_candidates.tracks), vecmem::get_data(track_candidates.measurements)}); @@ -174,7 +181,8 @@ TEST_P(KalmanFittingWireChamberTests, Run) { fit_performance_writer.write( track_states.tracks.at(i_trk), track_states.states, - track_candidates.measurements, host_det, evt_data); + track_candidates.measurements, detector.as(), + evt_data); } } diff --git a/tests/cpu/test_spacepoint_formation.cpp b/tests/cpu/test_spacepoint_formation.cpp index e1b5eac32d..8ce2eab5c5 100644 --- a/tests/cpu/test_spacepoint_formation.cpp +++ b/tests/cpu/test_spacepoint_formation.cpp @@ -40,11 +40,15 @@ TEST(spacepoint_formation, cpu) { tel_cfg.pilot_track(traj); // Create telescope geometry - const auto [det, name_map] = build_telescope_detector(host_mr, tel_cfg); + auto [det, name_map] = build_telescope_detector(host_mr, tel_cfg); - // Surface lookup auto surfaces = det.surfaces(); + traccc::host_detector host_det; + host_det.set(std::move(det)); + + // Surface lookup + // Prepare measurement collection typename measurement_collection_types::host measurements{&host_mr}; @@ -56,7 +60,7 @@ TEST(spacepoint_formation, cpu) { // Run spacepoint formation host::silicon_pixel_spacepoint_formation_algorithm sp_formation(host_mr); - auto spacepoints = sp_formation(det, vecmem::get_data(measurements)); + auto spacepoints = sp_formation(host_det, vecmem::get_data(measurements)); // Check the results EXPECT_EQ(spacepoints.size(), 2u); diff --git a/tests/cuda/test_ckf_combinatorics_telescope.cpp b/tests/cuda/test_ckf_combinatorics_telescope.cpp index 2e0d0be4ab..f1ad9d3dfd 100644 --- a/tests/cuda/test_ckf_combinatorics_telescope.cpp +++ b/tests/cuda/test_ckf_combinatorics_telescope.cpp @@ -9,6 +9,7 @@ #include "traccc/bfield/construct_const_bfield.hpp" #include "traccc/bfield/magnetic_field_types.hpp" #include "traccc/cuda/finding/combinatorial_kalman_filter_algorithm.hpp" +#include "traccc/io/read_detector.hpp" #include "traccc/io/read_measurements.hpp" #include "traccc/io/utils.hpp" #include "traccc/simulation/event_generators.hpp" @@ -62,21 +63,26 @@ TEST_P(CudaCkfCombinatoricsTelescopeTests, Run) { vecmem::cuda::device_memory_resource device_mr; traccc::memory_resource mr{device_mr, &host_mr}; vecmem::cuda::managed_memory_resource mng_mr; + vecmem::copy host_copy; // Read back detector file const std::string path = name + "/"; - detray::io::detector_reader_config reader_cfg{}; - reader_cfg.add_file(path + "telescope_detector_geometry.json") - .add_file(path + "telescope_detector_homogeneous_material.json"); - - const auto [host_det, names] = - detray::io::read_detector(mng_mr, reader_cfg); + traccc::host_detector detector; + traccc::io::read_detector( + detector, mng_mr, + std::filesystem::absolute( + std::filesystem::path(path + "telescope_detector_geometry.json")) + .native(), + std::filesystem::absolute( + std::filesystem::path( + path + "telescope_detector_homogeneous_material.json")) + .native()); + + const traccc::detector_buffer detector_buffer = + traccc::buffer_from_host_detector(detector, mng_mr, host_copy); const auto field = traccc::construct_const_bfield(std::get<13>(GetParam())); - // Detector view object - auto det_view = detray::get_data(host_det); - /*************************** * Generate simulation data ***************************/ @@ -109,7 +115,7 @@ TEST_P(CudaCkfCombinatoricsTelescopeTests, Run) { std::filesystem::create_directories(full_path); auto sim = traccc::simulator( - ptc, n_events, host_det, + ptc, n_events, detector.as(), field.as_field>(), std::move(generator), std::move(smearer_writer_cfg), full_path); sim.run(); @@ -125,7 +131,8 @@ TEST_P(CudaCkfCombinatoricsTelescopeTests, Run) { vecmem::cuda::async_copy copy{stream.cudaStream()}; // Seed generator - seed_generator sg(host_det, stddevs); + seed_generator sg(detector.as(), + stddevs); // Finding algorithm configuration typename traccc::cuda::combinatorial_kalman_filter_algorithm::config_type @@ -189,12 +196,12 @@ TEST_P(CudaCkfCombinatoricsTelescopeTests, Run) { // Run device finding traccc::edm::track_candidate_collection::buffer track_candidates_cuda_buffer = device_finding( - det_view, field, measurements_buffer, seeds_buffer); + detector_buffer, field, measurements_buffer, seeds_buffer); // Run device finding (Limit) traccc::edm::track_candidate_collection::buffer track_candidates_limit_cuda_buffer = device_finding_limit( - det_view, field, measurements_buffer, seeds_buffer); + detector_buffer, field, measurements_buffer, seeds_buffer); traccc::edm::track_candidate_collection::host track_candidates_cuda{host_mr}, diff --git a/tests/cuda/test_ckf_toy_detector.cpp b/tests/cuda/test_ckf_toy_detector.cpp index ee9c92b6b5..2fc7f65fa9 100644 --- a/tests/cuda/test_ckf_toy_detector.cpp +++ b/tests/cuda/test_ckf_toy_detector.cpp @@ -10,6 +10,7 @@ #include "traccc/bfield/magnetic_field_types.hpp" #include "traccc/cuda/finding/combinatorial_kalman_filter_algorithm.hpp" #include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" +#include "traccc/io/read_detector.hpp" #include "traccc/io/read_measurements.hpp" #include "traccc/io/utils.hpp" #include "traccc/performance/container_comparator.hpp" @@ -58,22 +59,30 @@ TEST_P(CkfToyDetectorTests, Run) { vecmem::cuda::device_memory_resource device_mr; traccc::memory_resource mr{device_mr, &host_mr}; vecmem::cuda::managed_memory_resource mng_mr; + vecmem::copy host_copy; // Read back detector file const std::string path = name + "/"; - detray::io::detector_reader_config reader_cfg{}; - reader_cfg.add_file(path + "toy_detector_geometry.json") - .add_file(path + "toy_detector_homogeneous_material.json") - .add_file(path + "toy_detector_surface_grids.json"); - - const auto [host_det, names] = - detray::io::read_detector(mng_mr, reader_cfg); + traccc::host_detector detector; + traccc::io::read_detector( + detector, mng_mr, + std::filesystem::absolute( + std::filesystem::path(path + "toy_detector_geometry.json")) + .native(), + std::filesystem::absolute( + std::filesystem::path(path + + "toy_detector_homogeneous_material.json")) + .native(), + std::filesystem::absolute( + std::filesystem::path(path + "toy_detector_surface_grids.json")) + .native()); + + traccc::detector_buffer detector_buffer = + traccc::buffer_from_host_detector(detector, mng_mr, host_copy); + ; const auto field = traccc::construct_const_bfield(B); - // Detector view object - auto det_view = detray::get_data(host_det); - /*************************** * Generate simulation data ***************************/ @@ -107,7 +116,7 @@ TEST_P(CkfToyDetectorTests, Run) { std::filesystem::create_directories(full_path); auto sim = traccc::simulator( - ptc, n_events, host_det, + ptc, n_events, detector.as(), field.as_field>(), std::move(generator), std::move(smearer_writer_cfg), full_path); sim.get_config().propagation.navigation.search_window = search_window; @@ -124,7 +133,8 @@ TEST_P(CkfToyDetectorTests, Run) { vecmem::cuda::async_copy copy{stream.cudaStream()}; // Seed generator - seed_generator sg(host_det, stddevs); + seed_generator sg(detector.as(), + stddevs); // Finding algorithm configuration typename traccc::cuda::combinatorial_kalman_filter_algorithm::config_type @@ -182,13 +192,13 @@ TEST_P(CkfToyDetectorTests, Run) { // Run host finding auto track_candidates = host_finding( - host_det, field, vecmem::get_data(measurements_per_event), + detector, field, vecmem::get_data(measurements_per_event), vecmem::get_data(seeds)); // Run device finding traccc::edm::track_candidate_collection::buffer track_candidates_cuda_buffer = device_finding( - det_view, field, measurements_buffer, seeds_buffer); + detector_buffer, field, measurements_buffer, seeds_buffer); traccc::edm::track_candidate_collection::host track_candidates_cuda{host_mr}; diff --git a/tests/cuda/test_kalman_fitter_telescope.cpp b/tests/cuda/test_kalman_fitter_telescope.cpp index 062030aa81..d692ca57ef 100644 --- a/tests/cuda/test_kalman_fitter_telescope.cpp +++ b/tests/cuda/test_kalman_fitter_telescope.cpp @@ -10,6 +10,7 @@ #include "traccc/bfield/magnetic_field_types.hpp" #include "traccc/cuda/fitting/kalman_fitting_algorithm.hpp" #include "traccc/edm/track_fit_container.hpp" +#include "traccc/geometry/host_detector.hpp" #include "traccc/io/utils.hpp" #include "traccc/performance/details/is_same_object.hpp" #include "traccc/resolution/fitting_performance_writer.hpp" @@ -30,6 +31,7 @@ #include #include #include +#include // GTest include(s). #include @@ -78,11 +80,11 @@ TEST_P(KalmanFittingTelescopeTests, Run) { detray::io::detector_reader_config reader_cfg{}; reader_cfg.add_file(path + "telescope_detector_geometry.json") .add_file(path + "telescope_detector_homogeneous_material.json"); - const auto [host_det, names] = + auto [host_det, names] = detray::io::read_detector(mng_mr, reader_cfg); - // Detector view object - auto det_view = detray::get_data(host_det); + traccc::host_detector polymorphic_detector; + polymorphic_detector.set(std::move(host_det)); const auto field = traccc::construct_const_bfield(std::get<13>(GetParam())); @@ -118,7 +120,7 @@ TEST_P(KalmanFittingTelescopeTests, Run) { std::filesystem::create_directories(full_path); auto sim = traccc::simulator( - ptc, n_events, host_det, + ptc, n_events, polymorphic_detector.as(), field.as_field>(), std::move(generator), std::move(smearer_writer_cfg), full_path); sim.run(); @@ -133,8 +135,13 @@ TEST_P(KalmanFittingTelescopeTests, Run) { // Copy objects vecmem::cuda::async_copy copy{stream.cudaStream()}; + const traccc::detector_buffer detector_buffer = + traccc::buffer_from_host_detector(polymorphic_detector, device_mr, + copy); + // Seed generator - seed_generator sg(host_det, stddevs); + seed_generator sg( + polymorphic_detector.as(), stddevs); // Fitting algorithm object traccc::cuda::kalman_fitting_algorithm::config_type fit_cfg; @@ -165,7 +172,7 @@ TEST_P(KalmanFittingTelescopeTests, Run) { // Run fitting auto track_states_cuda_buffer = - device_fitting(det_view, field, + device_fitting(detector_buffer, field, {track_candidates_buffer.tracks, track_candidates_buffer.measurements}); @@ -194,7 +201,8 @@ TEST_P(KalmanFittingTelescopeTests, Run) { fit_performance_writer.write( track_states_cuda.tracks.at(i_trk), track_states_cuda.states, - track_candidates.measurements, host_det, evt_data); + track_candidates.measurements, + polymorphic_detector.as(), evt_data); } } diff --git a/tests/cuda/test_spacepoint_formation.cpp b/tests/cuda/test_spacepoint_formation.cpp index c59c867855..44b94df2f4 100644 --- a/tests/cuda/test_spacepoint_formation.cpp +++ b/tests/cuda/test_spacepoint_formation.cpp @@ -53,11 +53,16 @@ TEST(CUDASpacepointFormation, cuda) { tel_cfg.pilot_track(traj); // Create telescope geometry - const auto [det, name_map] = build_telescope_detector(mng_mr, tel_cfg); - using device_detector_type = traccc::telescope_detector::device; + auto [det, name_map] = build_telescope_detector(mng_mr, tel_cfg); + + traccc::host_detector host_det; + host_det.set(std::move(det)); // Surface lookup - auto surfaces = det.surfaces(); + auto surfaces = host_det.as().surfaces(); + + const traccc::detector_buffer device_det = + traccc::buffer_from_host_detector(host_det, mng_mr, copy); // Prepare measurement collection measurement_collection_types::host measurements{&mng_mr}; @@ -69,10 +74,9 @@ TEST(CUDASpacepointFormation, cuda) { measurements.push_back({{10.f, 15.f}, {0.f, 0.f}, surfaces[8u].barcode()}); // Run spacepoint formation - traccc::cuda::spacepoint_formation_algorithm - sp_formation(mr, copy, stream); + traccc::cuda::spacepoint_formation_algorithm sp_formation(mr, copy, stream); auto spacepoints_buffer = - sp_formation(detray::get_data(det), vecmem::get_data(measurements)); + sp_formation(device_det, vecmem::get_data(measurements)); edm::spacepoint_collection::device spacepoints(spacepoints_buffer); diff --git a/tests/io/test_csv.cpp b/tests/io/test_csv.cpp index a4f075d9ae..093bbe14cb 100644 --- a/tests/io/test_csv.cpp +++ b/tests/io/test_csv.cpp @@ -96,7 +96,7 @@ TEST_F(io, csv_read_odd_single_muon) { // Memory resource used by the test. vecmem::host_memory_resource mr; - traccc::default_detector::host detector{mr}; + traccc::host_detector detector; traccc::io::read_detector(detector, mr, "geometries/odd/odd-detray_geometry_detray.json"); diff --git a/tests/io/test_event_data.cpp b/tests/io/test_event_data.cpp index f8e11dfe22..aa243a944a 100644 --- a/tests/io/test_event_data.cpp +++ b/tests/io/test_event_data.cpp @@ -42,9 +42,13 @@ TEST(event_data, acts_odd) { auto [host_det, names] = detray::io::read_detector(resource, reader_cfg); + traccc::host_detector polymorphic_detector; + polymorphic_detector.set(std::move(host_det)); + { // without cell - traccc::event_data evt_data(path, 0u, resource, true, &host_det, + traccc::event_data evt_data(path, 0u, resource, true, + &polymorphic_detector, traccc::data_format::csv, false); EXPECT_EQ(evt_data.m_particle_map.size(), 4515u); EXPECT_EQ(evt_data.m_meas_to_ptc_map.size(), 58u); @@ -52,7 +56,8 @@ TEST(event_data, acts_odd) { } { // with cell - traccc::event_data evt_data(path, 0u, resource, true, &host_det, + traccc::event_data evt_data(path, 0u, resource, true, + &polymorphic_detector, traccc::data_format::csv, true); EXPECT_EQ(evt_data.m_particle_map.size(), 4515u); EXPECT_EQ(evt_data.m_meas_to_ptc_map.size(), 58u); @@ -107,7 +112,10 @@ TEST(event_data, mock_data) { auto [host_det, names] = detray::io::read_detector(resource, reader_cfg); - traccc::event_data evt_data(path, 0u, resource, true, &host_det, + traccc::host_detector polymorphic_detector; + polymorphic_detector.set(std::move(host_det)); + + traccc::event_data evt_data(path, 0u, resource, true, &polymorphic_detector, traccc::data_format::csv, true); // There are three measurements diff --git a/tests/sycl/test_ckf_combinatorics_telescope.cpp b/tests/sycl/test_ckf_combinatorics_telescope.cpp index 38ebddc996..4cbaa922bb 100644 --- a/tests/sycl/test_ckf_combinatorics_telescope.cpp +++ b/tests/sycl/test_ckf_combinatorics_telescope.cpp @@ -70,21 +70,27 @@ TEST_P(CkfCombinatoricsTelescopeTests, Run) { traccc::memory_resource mr{device_mr, &host_mr}; vecmem::sycl::shared_memory_resource shared_mr{vecmem_queue}; + // Copy objects + vecmem::sycl::async_copy copy{vecmem_queue}; + // Path to the working directory. const std::filesystem::path path = std::filesystem::current_path() / name; // Read in the detector geometry that was generated by the test fixture. - host_detector_type host_det{shared_mr}; + traccc::host_detector host_detector; traccc::io::read_detector( - host_det, shared_mr, - (path / "telescope_detector_geometry.json").native(), - (path / "telescope_detector_homogeneous_material.json").native(), ""); + host_detector, host_mr, + std::filesystem::path(path / "telescope_detector_geometry.json") + .native(), + std::filesystem::absolute( + path / "telescope_detector_homogeneous_material.json") + .native()); - auto field = traccc::construct_const_bfield(std::get<13>(GetParam())); + const traccc::detector_buffer detector_buffer = + traccc::buffer_from_host_detector(host_detector, device_mr, copy); + vecmem_queue.synchronize(); - // Detector view object - const host_detector_type& const_host_det = host_det; - auto det_view = detray::get_data(const_host_det); + auto field = traccc::construct_const_bfield(std::get<13>(GetParam())); /*************************** * Generate simulation data @@ -116,7 +122,7 @@ TEST_P(CkfCombinatoricsTelescopeTests, Run) { // Run simulator auto sim = traccc::simulator( - ptc, n_events, host_det, + ptc, n_events, host_detector.as(), field.as_field>(), std::move(generator), std::move(smearer_writer_cfg), path.native()); sim.run(); @@ -125,11 +131,9 @@ TEST_P(CkfCombinatoricsTelescopeTests, Run) { * Do the reconstruction *****************************/ - // Copy objects - vecmem::sycl::async_copy copy{vecmem_queue}; - // Seed generator - seed_generator sg(host_det, stddevs); + seed_generator sg(host_detector.as(), + stddevs); // Finding algorithm configuration traccc::sycl::combinatorial_kalman_filter_algorithm::config_type @@ -192,12 +196,12 @@ TEST_P(CkfCombinatoricsTelescopeTests, Run) { ->wait(); // Run device finding - auto track_candidates_buffer = - device_finding(det_view, field, measurements_buffer, seeds_buffer); + auto track_candidates_buffer = device_finding( + detector_buffer, field, measurements_buffer, seeds_buffer); // Run device finding (Limit) auto track_candidates_limit_buffer = device_finding_limit( - det_view, field, measurements_buffer, seeds_buffer); + detector_buffer, field, measurements_buffer, seeds_buffer); traccc::edm::track_candidate_collection::host track_candidates{host_mr}, diff --git a/tests/sycl/test_ckf_toy_detector.cpp b/tests/sycl/test_ckf_toy_detector.cpp index f67e928b8b..c03dc5e239 100644 --- a/tests/sycl/test_ckf_toy_detector.cpp +++ b/tests/sycl/test_ckf_toy_detector.cpp @@ -66,22 +66,29 @@ TEST_P(CkfToyDetectorTests, Run) { traccc::memory_resource mr{device_mr, &host_mr}; vecmem::sycl::shared_memory_resource shared_mr{vecmem_queue}; + // Copy objects + vecmem::sycl::async_copy copy{vecmem_queue}; + // Path to the working directory. const std::filesystem::path path = std::filesystem::current_path() / name; // Read in the detector geometry that was generated by the test fixture. - host_detector_type host_det{shared_mr}; + traccc::host_detector host_detector; traccc::io::read_detector( - host_det, shared_mr, (path / "toy_detector_geometry.json").native(), - (path / "toy_detector_homogeneous_material.json").native(), - (path / "toy_detector_surface_grids.json").native()); + host_detector, host_mr, + std::filesystem::path(path / "toy_detector_geometry.json").native(), + std::filesystem::absolute(path / + "toy_detector_homogeneous_material.json") + .native(), + std::filesystem::absolute(path / "toy_detector_surface_grids.json") + .native()); + + const traccc::detector_buffer detector_buffer = + traccc::buffer_from_host_detector(host_detector, device_mr, copy); + vecmem_queue.synchronize(); const auto field = traccc::construct_const_bfield(B); - // Detector view object - const host_detector_type& const_host_det = host_det; - auto det_view = detray::get_data(const_host_det); - /*************************** * Generate simulation data ***************************/ @@ -113,7 +120,7 @@ TEST_P(CkfToyDetectorTests, Run) { // Run simulator auto sim = traccc::simulator( - ptc, n_events, host_det, + ptc, n_events, host_detector.as(), field.as_field>(), std::move(generator), std::move(smearer_writer_cfg), path.native()); sim.get_config().propagation.navigation.search_window = search_window; @@ -123,11 +130,9 @@ TEST_P(CkfToyDetectorTests, Run) { * Do the reconstruction *****************************/ - // Copy objects - vecmem::sycl::async_copy copy{vecmem_queue}; - // Seed generator - seed_generator sg(host_det, stddevs); + seed_generator sg(host_detector.as(), + stddevs); // Finding algorithm configuration traccc::sycl::combinatorial_kalman_filter_algorithm::config_type cfg; @@ -185,13 +190,13 @@ TEST_P(CkfToyDetectorTests, Run) { // Run host finding auto track_candidates = host_finding( - host_det, field, vecmem::get_data(measurements_per_event), + host_detector, field, vecmem::get_data(measurements_per_event), vecmem::get_data(seeds)); // Run device finding traccc::edm::track_candidate_collection::buffer track_candidates_sycl_buffer = device_finding( - det_view, field, measurements_buffer, seeds_buffer); + detector_buffer, field, measurements_buffer, seeds_buffer); traccc::edm::track_candidate_collection::host track_candidates_sycl{host_mr}; diff --git a/tests/sycl/test_kalman_fitter_telescope.cpp b/tests/sycl/test_kalman_fitter_telescope.cpp index 75849e883b..023c8f0e79 100644 --- a/tests/sycl/test_kalman_fitter_telescope.cpp +++ b/tests/sycl/test_kalman_fitter_telescope.cpp @@ -93,9 +93,12 @@ TEST_P(KalmanFittingTelescopeTests, Run) { reader_cfg.add_file(path + "telescope_detector_geometry.json") .add_file(path + "telescope_detector_homogeneous_material.json"); - const auto [host_det, names] = + auto [host_det, names] = detray::io::read_detector(shared_mr, reader_cfg); - auto det_view = detray::get_data(host_det); + + traccc::host_detector polymorphic_detector; + polymorphic_detector.set(std::move(host_det)); + const auto field = traccc::construct_const_bfield(std::get<13>(GetParam())); /*************************** @@ -130,7 +133,7 @@ TEST_P(KalmanFittingTelescopeTests, Run) { std::filesystem::create_directories(full_path); auto sim = traccc::simulator( - ptc, n_events, host_det, + ptc, n_events, polymorphic_detector.as(), field.as_field>(), std::move(generator), std::move(smearer_writer_cfg), full_path); sim.run(); @@ -142,7 +145,8 @@ TEST_P(KalmanFittingTelescopeTests, Run) { vecmem::sycl::copy copy{vecmem_queue}; // Seed generator - seed_generator sg(host_det, stddevs); + seed_generator sg( + polymorphic_detector.as(), stddevs); // Fitting algorithm object typename traccc::sycl::kalman_fitting_algorithm::config_type fit_cfg; @@ -150,6 +154,10 @@ TEST_P(KalmanFittingTelescopeTests, Run) { traccc::sycl::kalman_fitting_algorithm device_fitting(fit_cfg, mr, copy, traccc_queue); + const traccc::detector_buffer detector_buffer = + traccc::buffer_from_host_detector(polymorphic_detector, device_mr, + copy); + // Iterate over events for (std::size_t i_evt = 0; i_evt < n_events; i_evt++) { @@ -174,7 +182,7 @@ TEST_P(KalmanFittingTelescopeTests, Run) { // Run fitting auto track_states_sycl_buffer = - device_fitting(det_view, field, + device_fitting(detector_buffer, field, {track_candidates_buffer.tracks, track_candidates_buffer.measurements}); @@ -205,7 +213,8 @@ TEST_P(KalmanFittingTelescopeTests, Run) { fit_performance_writer.write( track_states_sycl.tracks.at(i_trk), track_states_sycl.states, - track_candidates.measurements, host_det, evt_data); + track_candidates.measurements, + polymorphic_detector.as(), evt_data); } } diff --git a/tests/sycl/test_spacepoint_formation.cpp b/tests/sycl/test_spacepoint_formation.cpp index 70b473407b..88ebdde94b 100644 --- a/tests/sycl/test_spacepoint_formation.cpp +++ b/tests/sycl/test_spacepoint_formation.cpp @@ -52,11 +52,18 @@ TEST(SYCLSpacepointFormation, sycl) { tel_cfg.pilot_track(traj); // Create telescope geometry - const auto [det, name_map] = build_telescope_detector(shared_mr, tel_cfg); + auto [det, name_map] = build_telescope_detector(shared_mr, tel_cfg); // Surface lookup auto surfaces = det.surfaces(); + traccc::host_detector polymorphic_detector; + polymorphic_detector.set(std::move(det)); + + const traccc::detector_buffer detector_buffer = + traccc::buffer_from_host_detector(polymorphic_detector, shared_mr, + copy); + // Prepare measurement collection measurement_collection_types::host measurements{&shared_mr}; @@ -70,7 +77,7 @@ TEST(SYCLSpacepointFormation, sycl) { traccc::sycl::silicon_pixel_spacepoint_formation_algorithm sp_formation( mr, copy, queue.queue()); auto spacepoints_buffer = - sp_formation(detray::get_data(det), vecmem::get_data(measurements)); + sp_formation(detector_buffer, vecmem::get_data(measurements)); edm::spacepoint_collection::device spacepoints(spacepoints_buffer);