From 14e8fab88d69be03a54dc08aefe9fd2eecb865f9 Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Mon, 19 May 2025 16:14:19 +0200 Subject: [PATCH 01/11] Introduced an SoA EDM for the track fit in traccc::core. --- core/CMakeLists.txt | 5 +- .../traccc/edm/impl/track_fit_collection.ipp | 21 ++ .../edm/impl/track_state_collection.ipp | 122 +++++++++ .../traccc/edm/track_fit_collection.hpp | 175 ++++++++++++ .../traccc/edm/track_fit_container.hpp | 75 +++++ core/include/traccc/edm/track_fit_outcome.hpp | 24 ++ core/include/traccc/edm/track_quality.hpp | 50 ---- core/include/traccc/edm/track_state.hpp | 244 ----------------- .../traccc/edm/track_state_collection.hpp | 259 ++++++++++++++++++ .../details/combinatorial_kalman_filter.hpp | 29 +- .../traccc/fitting/details/kalman_fitting.hpp | 60 ++-- .../kalman_filter/gain_matrix_updater.hpp | 40 +-- .../fitting/kalman_filter/kalman_actor.hpp | 76 ++--- .../fitting/kalman_filter/kalman_fitter.hpp | 107 ++++---- .../kalman_filter/statistics_updater.hpp | 27 +- .../kalman_filter/two_filters_smoother.hpp | 50 ++-- .../fitting/kalman_fitting_algorithm.hpp | 10 +- .../greedy_ambiguity_resolution_algorithm.cpp | 1 - ...man_fitting_algorithm_default_detector.cpp | 8 +- ...n_fitting_algorithm_telescope_detector.cpp | 8 +- performance/CMakeLists.txt | 5 +- .../efficiency/finding_performance_writer.hpp | 13 +- .../details/comparator_factory.hpp | 1 + .../performance/details/is_same_object.hpp | 3 +- .../traccc/performance/details/projector.hpp | 8 - .../impl/is_same_fitting_result.ipp | 35 --- .../performance/impl/is_same_track_fit.ipp | 105 +++++++ .../performance/impl/is_same_track_state.ipp | 82 ++++++ .../impl/track_fit_comparator_factory.ipp | 60 ++++ .../resolution/fitting_performance_writer.hpp | 44 ++- .../efficiency/finding_performance_writer.cpp | 39 ++- .../performance/details/is_same_object.cpp | 19 -- .../resolution/fitting_performance_writer.cpp | 14 +- performance/src/resolution/stat_plot_tool.cpp | 39 --- performance/src/resolution/stat_plot_tool.hpp | 12 +- performance/src/resolution/stat_plot_tool.ipp | 43 +++ 36 files changed, 1296 insertions(+), 617 deletions(-) create mode 100644 core/include/traccc/edm/impl/track_fit_collection.ipp create mode 100644 core/include/traccc/edm/impl/track_state_collection.ipp create mode 100644 core/include/traccc/edm/track_fit_collection.hpp create mode 100644 core/include/traccc/edm/track_fit_container.hpp create mode 100644 core/include/traccc/edm/track_fit_outcome.hpp delete mode 100644 core/include/traccc/edm/track_quality.hpp delete mode 100644 core/include/traccc/edm/track_state.hpp create mode 100644 core/include/traccc/edm/track_state_collection.hpp delete mode 100644 performance/include/traccc/performance/impl/is_same_fitting_result.ipp create mode 100644 performance/include/traccc/performance/impl/is_same_track_fit.ipp create mode 100644 performance/include/traccc/performance/impl/is_same_track_state.ipp create mode 100644 performance/include/traccc/performance/impl/track_fit_comparator_factory.ipp diff --git a/core/CMakeLists.txt b/core/CMakeLists.txt index 03648ec4b4..9aaf0e8f6b 100644 --- a/core/CMakeLists.txt +++ b/core/CMakeLists.txt @@ -24,7 +24,6 @@ traccc_add_library( traccc_core core TYPE SHARED "include/traccc/edm/particle.hpp" "include/traccc/edm/track_parameters.hpp" "include/traccc/edm/container.hpp" - "include/traccc/edm/track_state.hpp" "include/traccc/edm/silicon_cell_collection.hpp" "include/traccc/edm/impl/silicon_cell_collection.ipp" "include/traccc/edm/silicon_cluster_collection.hpp" @@ -35,6 +34,10 @@ traccc_add_library( traccc_core core TYPE SHARED "include/traccc/edm/track_candidate_collection.hpp" "include/traccc/edm/impl/track_candidate_collection.ipp" "include/traccc/edm/track_candidate_container.hpp" + "include/traccc/edm/track_state_collection.hpp" + "include/traccc/edm/impl/track_state_collection.ipp" + "include/traccc/edm/track_fit_outcome.hpp" + "include/traccc/edm/track_fit_collection.hpp" # Magnetic field description. "include/traccc/bfield/magnetic_field_types.hpp" "include/traccc/bfield/magnetic_field.hpp" diff --git a/core/include/traccc/edm/impl/track_fit_collection.ipp b/core/include/traccc/edm/impl/track_fit_collection.ipp new file mode 100644 index 0000000000..6acef48830 --- /dev/null +++ b/core/include/traccc/edm/impl/track_fit_collection.ipp @@ -0,0 +1,21 @@ +/** 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 + */ + +#pragma once + +namespace traccc::edm { + +template +TRACCC_HOST_DEVICE void track_fit::reset_quality() { + + ndf() = {}; + chi2() = {}; + pval() = {}; + nholes() = {}; +} + +} // namespace traccc::edm diff --git a/core/include/traccc/edm/impl/track_state_collection.ipp b/core/include/traccc/edm/impl/track_state_collection.ipp new file mode 100644 index 0000000000..ab386d7de3 --- /dev/null +++ b/core/include/traccc/edm/impl/track_state_collection.ipp @@ -0,0 +1,122 @@ +/** 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 + */ + +#pragma once + +namespace traccc::edm { + +template +TRACCC_HOST_DEVICE bool track_state::is_hole() const { + + return (state() & IS_HOLE_MASK); +} + +template +TRACCC_HOST_DEVICE void track_state::set_hole(bool value) { + + if (value) { + state() |= IS_HOLE_MASK; + } else { + state() &= ~IS_HOLE_MASK; + } +} + +template +TRACCC_HOST_DEVICE bool track_state::is_smoothed() const { + + return (state() & IS_SMOOTHED_MASK); +} + +template +TRACCC_HOST_DEVICE void track_state::set_smoothed(bool value) { + + if (value) { + state() |= IS_SMOOTHED_MASK; + } else { + state() &= ~IS_SMOOTHED_MASK; + } +} + +template +template +TRACCC_HOST_DEVICE void track_state::get_measurement_local( + const measurement_collection_types::const_device& measurements, + detray::dmatrix& pos) const { + + static_assert(((D == 1u) || (D == 2u)), + "The measurement dimension must be 1 or 2"); + + assert((measurements.at(measurement_index()).subs.get_indices()[0] == + e_bound_loc0) || + (measurements.at(measurement_index()).subs.get_indices()[0] == + e_bound_loc1)); + + const point2& local = measurements.at(measurement_index()).local; + + switch (measurements.at(measurement_index()).subs.get_indices()[0]) { + case e_bound_loc0: + getter::element(pos, 0, 0) = local[0]; + if constexpr (D == 2u) { + getter::element(pos, 1, 0) = local[1]; + } + break; + case e_bound_loc1: + getter::element(pos, 0, 0) = local[1]; + if constexpr (D == 2u) { + getter::element(pos, 1, 0) = local[0]; + } + break; + default: +#if defined(__GNUC__) + __builtin_unreachable(); +#endif + } +} + +template +template +TRACCC_HOST_DEVICE void track_state::get_measurement_covariance( + const measurement_collection_types::const_device& measurements, + detray::dmatrix& cov) const { + + static_assert(((D == 1u) || (D == 2u)), + "The measurement dimension must be 1 or 2"); + + assert((measurements.at(measurement_index()).subs.get_indices()[0] == + e_bound_loc0) || + (measurements.at(measurement_index()).subs.get_indices()[0] == + e_bound_loc1)); + + const variance2& variance = measurements.at(measurement_index()).variance; + + switch (measurements.at(measurement_index()).subs.get_indices()[0]) { + case e_bound_loc0: + getter::element(cov, 0, 0) = variance[0]; + if constexpr (D == 2u) { + getter::element(cov, 0, 1) = 0.f; + getter::element(cov, 1, 0) = 0.f; + getter::element(cov, 1, 1) = variance[1]; + } + break; + case e_bound_loc1: + getter::element(cov, 0, 0) = variance[1]; + if constexpr (D == 2u) { + getter::element(cov, 0, 1) = 0.f; + getter::element(cov, 1, 0) = 0.f; + getter::element(cov, 1, 1) = variance[0]; + } + break; + default: +#if defined(__GNUC__) + __builtin_unreachable(); +#endif + } +} + +} // namespace traccc::edm diff --git a/core/include/traccc/edm/track_fit_collection.hpp b/core/include/traccc/edm/track_fit_collection.hpp new file mode 100644 index 0000000000..a9f7cbc839 --- /dev/null +++ b/core/include/traccc/edm/track_fit_collection.hpp @@ -0,0 +1,175 @@ +/** 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 + */ + +#pragma once + +// Local include(s). +#include "traccc/definitions/qualifiers.hpp" +#include "traccc/edm/track_fit_outcome.hpp" +#include "traccc/edm/track_parameters.hpp" + +// Detray include(s). +#include + +// VecMem include(s). +#include + +namespace traccc::edm { + +/// Interface for the @c traccc::edm::track_fit_collection type. +/// +/// It provides the API that users would interact with, while using the +/// columns/arrays of the SoA containers, or the variables of the AoS proxies +/// created on top of the SoA containers. +/// +template +class track_fit : public BASE { + + public: + /// @name Functions inherited from the base class + /// @{ + + /// Inherit the base class's constructor(s) + using BASE::BASE; + /// Inherit the base class's assignment operator(s). + using BASE::operator=; + + /// @} + + /// @name Track Candidate Information + /// @{ + + /// The outcome of the track fit (non-const) + /// + /// @return A (non-const) vector of @c traccc::track_fit_outcome + /// + TRACCC_HOST_DEVICE + auto& fit_outcome() { return BASE::template get<0>(); } + /// The outcome of the track fit (non-const) + /// + /// @return A (const) vector of @c traccc::track_fit_outcome + /// + TRACCC_HOST_DEVICE + const auto& fit_outcome() const { return BASE::template get<0>(); } + + /// The parameters of the track (non-const) + /// + /// @return A (non-const) vector of bound track parameters + /// + TRACCC_HOST_DEVICE + auto& params() { return BASE::template get<1>(); } + /// The parameters of the track (const) + /// + /// @return A (const) vector of bound track parameters + /// + TRACCC_HOST_DEVICE + const auto& params() const { return BASE::template get<1>(); } + + /// The number of degrees of freedom of the track fit (non-const) + /// + /// @return A (non-const) vector of scalar values + /// + TRACCC_HOST_DEVICE + auto& ndf() { return BASE::template get<2>(); } + /// The number of degrees of freedom of the track fit (const) + /// + /// @return A (const) vector of scalar values + /// + TRACCC_HOST_DEVICE + const auto& ndf() const { return BASE::template get<2>(); } + + /// The chi square of the track fit (non-const) + /// + /// @return A (non-const) vector of scalar values + /// + TRACCC_HOST_DEVICE + auto& chi2() { return BASE::template get<3>(); } + /// The chi square of the track fit (const) + /// + /// @return A (const) vector of scalar values + /// + TRACCC_HOST_DEVICE + const auto& chi2() const { return BASE::template get<3>(); } + + /// The p-value of the track fit (non-const) + /// + /// @return A (non-const) vector of scalar values + /// + TRACCC_HOST_DEVICE + auto& pval() { return BASE::template get<4>(); } + /// The p-value of the track fit (const) + /// + /// @return A (const) vector of scalar values + /// + TRACCC_HOST_DEVICE + const auto& pval() const { return BASE::template get<4>(); } + + /// The number of holes in the track pattern (non-const) + /// + /// @return A (non-const) vector of unsigned integers + /// + TRACCC_HOST_DEVICE + auto& nholes() { return BASE::template get<5>(); } + /// The number of holes in the track pattern (const) + /// + /// @return A (const) vector of unsigned integers + /// + TRACCC_HOST_DEVICE + const auto& nholes() const { return BASE::template get<5>(); } + + /// The indices of the track states associated to the track fit (non-const) + /// + /// @return A (non-const) jagged vector of unsigned integers + /// + TRACCC_HOST_DEVICE + auto& state_indices() { return BASE::template get<6>(); } + /// The indices of the track states associated to the track fit (const) + /// + /// @return A (const) jagged vector of unsigned integers + /// + TRACCC_HOST_DEVICE + const auto& state_indices() const { return BASE::template get<6>(); } + + /// @} + + /// @name Utility functions + /// @{ + + /// Reset the fit quality variables + TRACCC_HOST_DEVICE + void reset_quality(); + + /// @} + +}; // class track_fit + +/// SoA container describing the fitted tracks +/// +/// @tparam ALGEBRA The algebra type used to describe the tracks +/// +template +using track_fit_collection = vecmem::edm::container< + track_fit, + // fit_outcome + vecmem::edm::type::vector, + // params + vecmem::edm::type::vector>, + // ndf + vecmem::edm::type::vector>, + // chi2 + vecmem::edm::type::vector>, + // pval + vecmem::edm::type::vector>, + // nholes + vecmem::edm::type::vector, + // state_indices + vecmem::edm::type::jagged_vector>; + +} // namespace traccc::edm + +// Include the implementation. +#include "traccc/edm/impl/track_fit_collection.ipp" diff --git a/core/include/traccc/edm/track_fit_container.hpp b/core/include/traccc/edm/track_fit_container.hpp new file mode 100644 index 0000000000..99d238b2ed --- /dev/null +++ b/core/include/traccc/edm/track_fit_container.hpp @@ -0,0 +1,75 @@ +/** 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 + */ + +#pragma once + +// Local include(s). +#include "traccc/edm/measurement.hpp" +#include "traccc/edm/track_fit_collection.hpp" +#include "traccc/edm/track_state_collection.hpp" + +// VecMem include(s). +#include + +namespace traccc::edm { + +/// Return type(s) for the track fitting algorithm(s) +template +struct track_fit_container { + + struct view { + /// The fitted tracks + track_fit_collection::view tracks; + /// The track states used for the fit + track_state_collection::view states; + /// The measurements used for the fit + measurement_collection_types::const_view measurements; + }; + + struct const_view { + /// The fitted tracks + track_fit_collection::const_view tracks; + /// The track states used for the fit + track_state_collection::const_view states; + /// The measurements used for the fit + measurement_collection_types::const_view measurements; + }; + + struct host { + /// Constructor using a memory resource + explicit host(vecmem::memory_resource& mr) : tracks{mr}, states{mr} {} + + /// The fitted tracks + track_fit_collection::host tracks; + /// The track states used for the fit + track_state_collection::host states; + }; + + struct buffer { + /// The fitted tracks + track_fit_collection::buffer tracks; + /// The track states used for the fit + track_state_collection::buffer states; + }; + + struct data { + /// The fitted tracks + track_fit_collection::data tracks; + /// The track states used for the fit + track_state_collection::data states; + }; + + struct const_data { + /// The fitted tracks + track_fit_collection::const_data tracks; + /// The track states used for the fit + track_state_collection::const_data states; + }; + +}; // struct track_fit_container + +} // namespace traccc::edm diff --git a/core/include/traccc/edm/track_fit_outcome.hpp b/core/include/traccc/edm/track_fit_outcome.hpp new file mode 100644 index 0000000000..509b540188 --- /dev/null +++ b/core/include/traccc/edm/track_fit_outcome.hpp @@ -0,0 +1,24 @@ +/** 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 + */ + +#pragma once + +// System include(s). +#include + +namespace traccc { + +/// Possible outcomes of a track fit +enum class track_fit_outcome : std::uint16_t { + UNKNOWN, + SUCCESS, + FAILURE_NON_POSITIVE_NDF, + FAILURE_NOT_ALL_SMOOTHED, + MAX_OUTCOME +}; + +} // namespace traccc diff --git a/core/include/traccc/edm/track_quality.hpp b/core/include/traccc/edm/track_quality.hpp deleted file mode 100644 index db2a7c19f5..0000000000 --- a/core/include/traccc/edm/track_quality.hpp +++ /dev/null @@ -1,50 +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 - */ - -#pragma once - -// Project include(s). -#include "traccc/definitions/primitives.hpp" -#include "traccc/definitions/qualifiers.hpp" - -namespace traccc { - -/// Track quality -struct track_quality { - - /// Number of degree of freedoms of fitted track - traccc::scalar ndf{0}; - - /// Chi square of fitted track - traccc::scalar chi2{0}; - - /// p-value - traccc::scalar pval{0}; - - // The number of holes (The number of sensitive surfaces which do not have a - // measurement for the track pattern) - unsigned int n_holes{0u}; - - /// Reset the summary - TRACCC_HOST_DEVICE - void reset_quality() { - ndf = 0.f; - chi2 = 0.f; - n_holes = 0u; - } -}; - -/// Equality operator for track quality -TRACCC_HOST_DEVICE -inline bool operator==(const track_quality& lhs, const track_quality& rhs) { - - return ((math::fabs(lhs.ndf - rhs.ndf) < float_epsilon) && - (math::fabs(lhs.chi2 - rhs.chi2) < float_epsilon) && - (lhs.n_holes == rhs.n_holes)); -} - -} // namespace traccc diff --git a/core/include/traccc/edm/track_state.hpp b/core/include/traccc/edm/track_state.hpp deleted file mode 100644 index 6ff40a01c3..0000000000 --- a/core/include/traccc/edm/track_state.hpp +++ /dev/null @@ -1,244 +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 - */ - -#pragma once - -// Project include(s). -#include "traccc/definitions/primitives.hpp" -#include "traccc/definitions/qualifiers.hpp" -#include "traccc/edm/container.hpp" -#include "traccc/edm/measurement.hpp" -#include "traccc/edm/track_parameters.hpp" -#include "traccc/edm/track_quality.hpp" - -namespace traccc { - -enum class fitter_outcome : uint32_t { - UNKNOWN, - SUCCESS, - FAILURE_NON_POSITIVE_NDF, - FAILURE_NOT_ALL_SMOOTHED, - MAX_OUTCOME -}; - -/// Fitting result per track -template -struct fitting_result { - using scalar_type = detray::dscalar; - - /// Fitting outcome - fitter_outcome fit_outcome = fitter_outcome::UNKNOWN; - - /// Fitted track parameter - traccc::bound_track_parameters fit_params; - - /// Track quality - traccc::track_quality trk_quality; -}; - -/// Fitting result per measurement -template -struct track_state { - - using scalar_type = detray::dscalar; - using size_type = detray::dsize_type; - - using bound_track_parameters_type = - traccc::bound_track_parameters; - using bound_matrix_type = traccc::bound_matrix; - template - using matrix_type = detray::dmatrix; - - track_state() = default; - - /// Construction with track candidate - TRACCC_HOST_DEVICE - track_state(const measurement& trk_cand) - : m_surface_link(trk_cand.surface_link), m_measurement(trk_cand) { - m_filtered.set_surface_link(m_surface_link); - m_smoothed.set_surface_link(m_surface_link); - } - - /// @return the surface link - TRACCC_HOST_DEVICE - inline detray::geometry::barcode surface_link() const { - return m_surface_link; - } - - /// @return the measurement - TRACCC_HOST_DEVICE - inline const measurement& get_measurement() const { return m_measurement; } - - /// @return the local position of measurement with 2 X 1 matrix - // FIXME: The conversion from vector to matrix is inefficient - template - TRACCC_HOST_DEVICE inline matrix_type measurement_local() const { - static_assert(((D == 1u) || (D == 2u)), - "The measurement dimension should be 1 or 2"); - assert((m_measurement.subs.get_indices()[0] == e_bound_loc0) || - (m_measurement.subs.get_indices()[0] == e_bound_loc1)); - - matrix_type ret; - switch (m_measurement.subs.get_indices()[0]) { - case e_bound_loc0: - getter::element(ret, 0, 0) = m_measurement.local[0]; - if constexpr (D == 2u) { - getter::element(ret, 1, 0) = m_measurement.local[1]; - } - break; - case e_bound_loc1: - getter::element(ret, 0, 0) = m_measurement.local[1]; - if constexpr (D == 2u) { - getter::element(ret, 1, 0) = m_measurement.local[0]; - } - break; - default: - __builtin_unreachable(); - } - return ret; - } - - /// @return the covariance of local position of measurement - template - TRACCC_HOST_DEVICE inline matrix_type measurement_covariance() const { - static_assert(((D == 1u) || (D == 2u)), - "The measurement dimension should be 1 or 2"); - assert((m_measurement.subs.get_indices()[0] == e_bound_loc0) || - (m_measurement.subs.get_indices()[0] == e_bound_loc1)); - - matrix_type ret; - switch (m_measurement.subs.get_indices()[0]) { - case e_bound_loc0: - getter::element(ret, 0, 0) = m_measurement.variance[0]; - if constexpr (D == 2u) { - getter::element(ret, 0, 1) = 0.f; - getter::element(ret, 1, 0) = 0.f; - getter::element(ret, 1, 1) = m_measurement.variance[1]; - } - break; - case e_bound_loc1: - getter::element(ret, 0, 0) = m_measurement.variance[1]; - if constexpr (D == 2u) { - getter::element(ret, 0, 1) = 0.f; - getter::element(ret, 1, 0) = 0.f; - getter::element(ret, 1, 1) = m_measurement.variance[0]; - } - break; - default: - __builtin_unreachable(); - } - return ret; - } - - /// @return the non-const chi square of filtered parameter - TRACCC_HOST_DEVICE - inline scalar_type& filtered_chi2() { return m_filtered_chi2; } - - /// @return the const chi square of filtered parameter - TRACCC_HOST_DEVICE - inline const scalar_type& filtered_chi2() const { return m_filtered_chi2; } - - /// @return the non-const chi square of backward filter - TRACCC_HOST_DEVICE - inline scalar_type& backward_chi2() { return m_backward_chi2; } - - /// @return the const chi square of backward filter - TRACCC_HOST_DEVICE - inline scalar_type backward_chi2() const { return m_backward_chi2; } - - /// @return the non-const filtered parameter - TRACCC_HOST_DEVICE - inline bound_track_parameters_type& filtered() { return m_filtered; } - - /// @return the const filtered parameter - TRACCC_HOST_DEVICE - inline const bound_track_parameters_type& filtered() const { - return m_filtered; - } - - /// @return the non-const chi square of smoothed parameter - TRACCC_HOST_DEVICE - inline scalar_type& smoothed_chi2() { return m_smoothed_chi2; } - - /// @return the const chi square of smoothed parameter - TRACCC_HOST_DEVICE - inline const scalar_type& smoothed_chi2() const { return m_smoothed_chi2; } - - /// @return the non-const smoothed parameter - TRACCC_HOST_DEVICE - inline bound_track_parameters_type& smoothed() { return m_smoothed; } - - /// @return the const smoothed parameter - TRACCC_HOST_DEVICE - inline const bound_track_parameters_type& smoothed() const { - return m_smoothed; - } - - public: - bool is_hole{true}; - bool is_smoothed{false}; - - private: - detray::geometry::barcode m_surface_link; - measurement m_measurement; - scalar_type m_filtered_chi2 = 0.f; - bound_track_parameters_type m_filtered; - scalar_type m_smoothed_chi2 = 0.f; - bound_track_parameters_type m_smoothed; - scalar_type m_backward_chi2 = 0.f; -}; - -/// Declare all track_state collection types -using track_state_collection_types = - collection_types>; - -/// Declare all track_state container types -using track_state_container_types = - container_types, - track_state>; - -inline void print_fitted_tracks_statistics( - const track_state_container_types::host& track_states) { - const std::size_t n_tracks = track_states.size(); - std::size_t success = 0; - std::size_t non_positive_ndf = 0; - std::size_t not_all_smoothed = 0; - - for (std::size_t i = 0; i < n_tracks; i++) { - if (track_states.at(i).header.fit_outcome == fitter_outcome::SUCCESS) { - success++; - } else if (track_states.at(i).header.fit_outcome == - fitter_outcome::FAILURE_NON_POSITIVE_NDF) { - non_positive_ndf++; - } else if (track_states.at(i).header.fit_outcome == - fitter_outcome::FAILURE_NOT_ALL_SMOOTHED) { - not_all_smoothed++; - } - } - - std::cout << "Success: " << success - << " Non positive NDF: " << non_positive_ndf - << " Not all smoothed: " << not_all_smoothed - << " Total: " << n_tracks << std::endl; -} - -inline std::size_t count_successfully_fitted_tracks( - const track_state_container_types::host& track_states) { - - const std::size_t n_tracks = track_states.size(); - std::size_t n_fitted_tracks = 0u; - - for (std::size_t i = 0; i < n_tracks; i++) { - if (track_states.at(i).header.fit_outcome == fitter_outcome::SUCCESS) { - n_fitted_tracks++; - } - } - - return n_fitted_tracks; -} - -} // namespace traccc diff --git a/core/include/traccc/edm/track_state_collection.hpp b/core/include/traccc/edm/track_state_collection.hpp new file mode 100644 index 0000000000..8019e90d82 --- /dev/null +++ b/core/include/traccc/edm/track_state_collection.hpp @@ -0,0 +1,259 @@ +/** 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 + */ + +#pragma once + +// Local include(s). +#include "traccc/definitions/qualifiers.hpp" +#include "traccc/edm/measurement.hpp" +#include "traccc/edm/track_parameters.hpp" + +// Detray include(s). +#include + +// VecMem include(s). +#include + +// System include(s). +#include +#include + +namespace traccc::edm { + +/// Interface for the @c traccc::edm::track_state_collection type. +/// +/// It provides the API that users would interact with, while using the +/// columns/arrays of the SoA containers, or the variables of the AoS proxies +/// created on top of the SoA containers. +/// +template +class track_state : public BASE { + + public: + /// @name Functions inherited from the base class + /// @{ + + /// Inherit the base class's constructor(s) + using BASE::BASE; + /// Inherit the base class's assignment operator(s). + using BASE::operator=; + + /// @} + + /// @name Constants + /// @{ + + /// @c is_hole bit in the state word + static constexpr std::uint8_t IS_HOLE_MASK = 0x01; + /// @c is_smoothed bit in the state word + static constexpr std::uint8_t IS_SMOOTHED_MASK = 0x02; + + /// @} + + /// @name Track State Information + /// @{ + + /// The "state word" of the track (non-const) + /// + /// @return A (non-const) vector of unsigned integers + /// + TRACCC_HOST_DEVICE + auto& state() { return BASE::template get<0>(); } + /// The "state word" of the track (const) + /// + /// @return A (const) vector of unsigned integers + /// + TRACCC_HOST_DEVICE + const auto& state() const { return BASE::template get<0>(); } + + /// Chi^2 of the fitered parameters (non-const) + /// + /// @return A (non-const) vector of scalar values + /// + TRACCC_HOST_DEVICE + auto& filtered_chi2() { return BASE::template get<1>(); } + /// Chi^2 of the filtered parameters (const) + /// + /// @return A (const) vector of scalar values + /// + TRACCC_HOST_DEVICE + const auto& filtered_chi2() const { return BASE::template get<1>(); } + + /// Chi^2 of the smoothed parameters (non-const) + /// + /// @return A (non-const) vector of scalar values + /// + TRACCC_HOST_DEVICE + auto& smoothed_chi2() { return BASE::template get<2>(); } + /// Chi^2 of the smoothed parameters (const) + /// + /// @return A (const) vector of scalar values + /// + TRACCC_HOST_DEVICE + const auto& smoothed_chi2() const { return BASE::template get<2>(); } + + /// Chi^2 of the backward parameters (non-const) + /// + /// @return A (non-const) vector of scalar values + /// + TRACCC_HOST_DEVICE + auto& backward_chi2() { return BASE::template get<3>(); } + /// Chi^2 of the backward parameters (const) + /// + /// @return A (const) vector of scalar values + /// + TRACCC_HOST_DEVICE + const auto& backward_chi2() const { return BASE::template get<3>(); } + + /// The filtered parameters of the track on the surface (non-const) + /// + /// @return A (non-const) vector of bound track parameters + /// + TRACCC_HOST_DEVICE + auto& filtered_params() { return BASE::template get<4>(); } + /// The filtered parameters of the track on the surface (const) + /// + /// @return A (const) vector of bound track parameters + /// + TRACCC_HOST_DEVICE + const auto& filtered_params() const { return BASE::template get<4>(); } + + /// The smoothed parameters of the track on the surface (non-const) + /// + /// @return A (non-const) vector of bound track parameters + /// + TRACCC_HOST_DEVICE + auto& smoothed_params() { return BASE::template get<5>(); } + /// The smoothed parameters of the track on the surface (const) + /// + /// @return A (const) vector of bound track parameters + /// + TRACCC_HOST_DEVICE + const auto& smoothed_params() const { return BASE::template get<5>(); } + + /// The index of the track's measurement on the current surface (non-const) + /// + /// @return A (non-const) vector of unsigned integers + /// + TRACCC_HOST_DEVICE + auto& measurement_index() { return BASE::template get<6>(); } + /// The index of the track's measurement on the current surface (const) + /// + /// @return A (const) vector of unsigned integers + /// + TRACCC_HOST_DEVICE + const auto& measurement_index() const { return BASE::template get<6>(); } + + /// @} + + /// @name Utility functions + /// @{ + + /// Check if the track state is on a hole + /// + /// @note This function must only be used on proxy objects, not on + /// containers! + /// + /// @return @c true if the track state is on a hole, @c false otherwise + /// + TRACCC_HOST_DEVICE + bool is_hole() const; + /// Set the track state to be on a hole + /// + /// @note This function must only be used on proxy objects, not on + /// containers! + /// + /// @param value The value to set + /// + TRACCC_HOST_DEVICE + void set_hole(bool value = true); + + /// Check if the track state is smoothed + /// + /// @note This function must only be used on proxy objects, not on + /// containers! + /// + /// @return @c true if the track state is smoothed, @c false otherwise + /// + TRACCC_HOST_DEVICE + bool is_smoothed() const; + /// Set the track state to be smoothed + /// + /// @note This function must only be used on proxy objects, not on + /// containers! + /// + /// @param value The value to set + /// + TRACCC_HOST_DEVICE + void set_smoothed(bool value = true); + + /// Get the local position of the measurement in a matrix + /// + /// @note This function must only be used on proxy objects, not on + /// containers! + /// + /// @tparam size_type The type of the matrix size variable + /// @tparam ALGEBRA The algebra type used to describe the tracks + /// @tparam D The dimension of the matrix + /// + /// @param measurements All measurements in the event + /// @param pos The matrix to fill with the local position of the measurement + /// + template + TRACCC_HOST_DEVICE void get_measurement_local( + const measurement_collection_types::const_device& measurements, + detray::dmatrix& pos) const; + + /// Get the covariance of the measurement in a matrix + /// + /// @note This function must only be used on proxy objects, not on + /// containers! + /// + /// @tparam size_type The type of the matrix size variable + /// @tparam ALGEBRA The algebra type used to describe the tracks + /// @tparam D The dimension of the matrix + /// + /// @param measurements All measurements in the event + /// @param cov The matrix to fill with the covariance of the measurement + /// + template + TRACCC_HOST_DEVICE void get_measurement_covariance( + const measurement_collection_types::const_device& measurements, + detray::dmatrix& cov) const; + + /// @} + +}; // class track_state + +/// SoA container describing the fit states of tracks on specific measurements +/// +/// @tparam ALGEBRA The algebra type used to describe the tracks +/// +template +using track_state_collection = vecmem::edm::container< + track_state, + // state + vecmem::edm::type::vector, + // filtered_chi2 + vecmem::edm::type::vector>, + // smoothed_chi2 + vecmem::edm::type::vector>, + // backward_chi2 + vecmem::edm::type::vector>, + // filtered_params + vecmem::edm::type::vector>, + // smoothed_params + vecmem::edm::type::vector>, + // measurement_index + vecmem::edm::type::vector>; + +} // namespace traccc::edm + +// Include the implementation. +#include "traccc/edm/impl/track_state_collection.ipp" diff --git a/core/include/traccc/finding/details/combinatorial_kalman_filter.hpp b/core/include/traccc/finding/details/combinatorial_kalman_filter.hpp index a125ba1072..7295beb87a 100644 --- a/core/include/traccc/finding/details/combinatorial_kalman_filter.hpp +++ b/core/include/traccc/finding/details/combinatorial_kalman_filter.hpp @@ -9,7 +9,6 @@ // Project include(s). #include "traccc/edm/measurement.hpp" #include "traccc/edm/track_candidate_collection.hpp" -#include "traccc/edm/track_state.hpp" #include "traccc/finding/actors/ckf_aborter.hpp" #include "traccc/finding/actors/interaction_register.hpp" #include "traccc/finding/candidate_link.hpp" @@ -253,23 +252,32 @@ combinatorial_kalman_filter( * Find tracks (CKF) *****************************************************************/ - std::vector>> + std::vector>> best_links; // Iterate over the measurements for (unsigned int item_id = range.first; item_id < range.second; item_id++) { - const auto& meas = measurements[item_id]; + // The measurement on surface to handle. + const measurement& meas = measurements.at(item_id); - track_state trk_state(meas); + // Create a standalone track state object. + typename edm::track_state_collection< + algebra_type>::host::object_type trk_state(0u, 0.f, 0.f, + 0.f, {}, {}, + item_id); + trk_state.set_hole(true); + trk_state.set_smoothed(false); + trk_state.filtered_params().set_surface_link(meas.surface_link); const bool is_line = sf.template visit_mask(); // Run the Kalman update on a copy of the track parameters const kalman_fitter_status res = - gain_matrix_updater{}(trk_state, in_param, - is_line); + gain_matrix_updater{}(trk_state, measurements, + in_param, is_line); const traccc::scalar chi2 = trk_state.filtered_chi2(); @@ -285,9 +293,8 @@ combinatorial_kalman_filter( .n_skipped = skip_counter, .chi2 = chi2, .chi2_sum = prev_chi2_sum + chi2, - .ndf_sum = prev_ndf_sum + - trk_state.get_measurement().meas_dim}, - trk_state}); + .ndf_sum = prev_ndf_sum + meas.meas_dim}, + trk_state.filtered_params()}); } } @@ -304,13 +311,13 @@ combinatorial_kalman_filter( << step << " and input parameter " << in_param_id); for (unsigned int i = 0; i < n_branches; ++i) { - const auto& [link, trk_state] = best_links[i]; + const auto& [link, filtered_params] = best_links[i]; // Add the link to the links container links[step].push_back(link); // Add the updated parameter to the updated parameters - updated_params.push_back(trk_state.filtered()); + updated_params.push_back(filtered_params); TRACCC_VERBOSE("updated_params[" << updated_params.size() - 1 << "] = " << updated_params.back()); diff --git a/core/include/traccc/fitting/details/kalman_fitting.hpp b/core/include/traccc/fitting/details/kalman_fitting.hpp index 201dafea57..eb025dd4b8 100644 --- a/core/include/traccc/fitting/details/kalman_fitting.hpp +++ b/core/include/traccc/fitting/details/kalman_fitting.hpp @@ -8,8 +8,11 @@ #pragma once // Project include(s). +#include "traccc/edm/measurement.hpp" #include "traccc/edm/track_candidate_container.hpp" -#include "traccc/edm/track_state.hpp" +#include "traccc/edm/track_fit_collection.hpp" +#include "traccc/edm/track_fit_container.hpp" +#include "traccc/edm/track_state_collection.hpp" #include "traccc/fitting/status_codes.hpp" // VecMem include(s). @@ -36,11 +39,11 @@ namespace traccc::host::details { /// /// @return A container of the fitted track states /// -template -track_state_container_types::host kalman_fitting( +template +typename edm::track_fit_container::host kalman_fitting( fitter_t& fitter, - const typename edm::track_candidate_container< - typename fitter_t::algebra_type>::const_view& track_container, + const typename edm::track_candidate_container::const_view& + track_container, vecmem::memory_resource& mr, vecmem::copy& copy) { // Create the input container(s). @@ -50,46 +53,61 @@ track_state_container_types::host kalman_fitting( typename fitter_t::algebra_type>::const_device track_candidates{ track_container.tracks}; - // Create the output container. - track_state_container_types::host result{&mr}; + // Create the output containers. + typename edm::track_fit_container::host result{mr}; // Iterate over the tracks, for (typename edm::track_candidate_collection< typename fitter_t::algebra_type>::const_device::size_type i = 0; i < track_candidates.size(); ++i) { - // Make a vector of track states for this track. - vecmem::vector > - input_states{&mr}; - input_states.reserve( - track_candidates.measurement_indices().at(i).size()); + // Create the objects that will describe this track fit. + result.tracks.push_back( + {track_fit_outcome::UNKNOWN, {}, 0.f, 0.f, 0.f, 0u, {}}); + auto fitted_track = result.tracks.at(result.tracks.size() - 1); for (unsigned int measurement_index : track_candidates.measurement_indices().at(i)) { - input_states.emplace_back(measurements.at(measurement_index)); + fitted_track.state_indices().push_back( + static_cast(result.states.size())); + result.states.push_back( + {0u, 0.f, 0.f, 0.f, {}, {}, measurement_index}); + auto state = result.states.at(result.states.size() - 1); + state.set_hole(true); + state.set_smoothed(false); + state.filtered_params().set_surface_link( + measurements.at(measurement_index).surface_link); + state.smoothed_params().set_surface_link( + measurements.at(measurement_index).surface_link); } vecmem::data::vector_buffer seqs_buffer{ static_cast::size_type>( - std::max(input_states.size() * + std::max(fitted_track.state_indices().size() * fitter.config().barcode_sequence_size_factor, fitter.config().min_barcode_sequence_capacity)), mr, vecmem::data::buffer_type::resizable}; copy.setup(seqs_buffer)->wait(); // Make a fitter state - typename fitter_t::state fitter_state(vecmem::get_data(input_states), - seqs_buffer); + auto result_tracks_view = vecmem::get_data(result.tracks); + typename edm::track_fit_collection::device + result_tracks_device{result_tracks_view}; + typename edm::track_fit_collection::device::proxy_type + fitted_track_device = + result_tracks_device.at(result_tracks_device.size() - 1); + auto result_states_view = vecmem::get_data(result.states); + typename fitter_t::state fitter_state( + fitted_track_device, + typename edm::track_state_collection::device{ + result_states_view}, + measurements, seqs_buffer); // Run the fitter. kalman_fitter_status fit_status = fitter.fit(track_candidates.params().at(i), fitter_state); - if (fit_status == kalman_fitter_status::SUCCESS) { - // Save the results into the output container. - result.push_back(std::move(fitter_state.m_fit_res), - std::move(input_states)); - } else { + if (fit_status != kalman_fitter_status::SUCCESS) { // TODO: Print a warning here. } } diff --git a/core/include/traccc/fitting/kalman_filter/gain_matrix_updater.hpp b/core/include/traccc/fitting/kalman_filter/gain_matrix_updater.hpp index 167a3bf2a8..87afb6f3e7 100644 --- a/core/include/traccc/fitting/kalman_filter/gain_matrix_updater.hpp +++ b/core/include/traccc/fitting/kalman_filter/gain_matrix_updater.hpp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2022-2024 CERN for the benefit of the ACTS project + * (c) 2022-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -10,7 +10,8 @@ // Project include(s). #include "traccc/definitions/qualifiers.hpp" #include "traccc/definitions/track_parametrization.hpp" -#include "traccc/edm/track_state.hpp" +#include "traccc/edm/measurement.hpp" +#include "traccc/edm/track_state_collection.hpp" #include "traccc/fitting/status_codes.hpp" namespace traccc { @@ -37,20 +38,24 @@ struct gain_matrix_updater { /// @param bound_params bound parameter /// /// @return true if the update succeeds + template [[nodiscard]] TRACCC_HOST_DEVICE inline kalman_fitter_status operator()( - track_state& trk_state, + typename edm::track_state& trk_state, + const measurement_collection_types::const_device& measurements, const bound_track_parameters& bound_params, const bool is_line) const { - const auto D = trk_state.get_measurement().meas_dim; + const auto D = measurements.at(trk_state.measurement_index()).meas_dim; assert(D == 1u || D == 2u); - return update(trk_state, bound_params, D, is_line); + return update(trk_state, measurements, bound_params, D, is_line); } + template [[nodiscard]] TRACCC_HOST_DEVICE inline kalman_fitter_status update( - track_state& trk_state, + typename edm::track_state& trk_state, + const measurement_collection_types::const_device& measurements, const bound_track_parameters& bound_params, const unsigned int dim, const bool is_line) const { @@ -61,16 +66,15 @@ struct gain_matrix_updater { assert(!bound_params.is_invalid()); assert(!bound_params.surface_link().is_invalid()); - const auto meas = trk_state.get_measurement(); - // Some identity matrices // @TODO: Make constexpr work const auto I66 = matrix::identity(); const auto I_m = matrix::identity>(); // Measurement data on surface - const matrix_type meas_local = - trk_state.template measurement_local(); + matrix_type meas_local; + trk_state.template get_measurement_local(measurements, + meas_local); assert((dim > 1) || (getter::element(meas_local, 1u, 0u) == 0.f)); @@ -80,7 +84,9 @@ struct gain_matrix_updater { // Predicted covaraince of bound track parameters const bound_matrix_type& predicted_cov = bound_params.covariance(); - matrix_type H = meas.subs.template projector(); + matrix_type H = + measurements.at(trk_state.measurement_index()) + .subs.template projector(); // Flip the sign of projector matrix element in case the first element // of line measurement is negative @@ -94,7 +100,9 @@ struct gain_matrix_updater { } // Spatial resolution (Measurement covariance) - matrix_type V = trk_state.template measurement_covariance(); + matrix_type V; + trk_state.template get_measurement_covariance(measurements, + V); if (dim == 1) { getter::element(V, 1u, 1u) = 1.f; @@ -148,14 +156,14 @@ struct gain_matrix_updater { } // Set the track state parameters - trk_state.filtered().set_vector(filtered_vec); - trk_state.filtered().set_covariance(filtered_cov); + trk_state.filtered_params().set_vector(filtered_vec); + trk_state.filtered_params().set_covariance(filtered_cov); trk_state.filtered_chi2() = getter::element(chi2, 0, 0); // Wrap the phi in the range of [-pi, pi] - wrap_phi(trk_state.filtered()); + wrap_phi(trk_state.filtered_params()); - assert(!trk_state.filtered().is_invalid()); + assert(!trk_state.filtered_params().is_invalid()); return kalman_fitter_status::SUCCESS; } diff --git a/core/include/traccc/fitting/kalman_filter/kalman_actor.hpp b/core/include/traccc/fitting/kalman_filter/kalman_actor.hpp index 7a54135f14..4ce9665a9d 100644 --- a/core/include/traccc/fitting/kalman_filter/kalman_actor.hpp +++ b/core/include/traccc/fitting/kalman_filter/kalman_actor.hpp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2022-2023 CERN for the benefit of the ACTS project + * (c) 2022-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -9,7 +9,9 @@ // Project include(s). #include "traccc/definitions/qualifiers.hpp" -#include "traccc/edm/track_state.hpp" +#include "traccc/edm/measurement.hpp" +#include "traccc/edm/track_fit_collection.hpp" +#include "traccc/edm/track_state_collection.hpp" #include "traccc/fitting/kalman_filter/gain_matrix_updater.hpp" #include "traccc/fitting/kalman_filter/is_line_visitor.hpp" #include "traccc/fitting/kalman_filter/two_filters_smoother.hpp" @@ -32,31 +34,38 @@ enum class kalman_actor_direction { template struct kalman_actor_state { - using track_state_coll = vecmem::device_vector>; /// Constructor with the vector of track states TRACCC_HOST_DEVICE - explicit kalman_actor_state(track_state_coll track_states) - : m_track_states(track_states) { - m_it = m_track_states.begin(); - m_it_rev = m_track_states.rbegin(); + kalman_actor_state( + const typename edm::track_fit_collection::device::proxy_type& + track, + const typename edm::track_state_collection::device& + track_states, + const measurement_collection_types::const_device& measurements) + : m_track{track}, + m_track_states{track_states}, + m_measurements{measurements} { + + reset(); } /// @return the reference of track state pointed by the iterator TRACCC_HOST_DEVICE - typename track_state_coll::value_type& operator()() { + typename edm::track_state_collection::device::proxy_type + operator()() { if (!backward_mode) { - return *m_it; + return m_track_states.at(*m_it); } else { - return *m_it_rev; + return m_track_states.at(*m_it_rev); } } /// Reset the iterator TRACCC_HOST_DEVICE void reset() { - m_it = m_track_states.begin(); - m_it_rev = m_track_states.rbegin(); + m_it = m_track.state_indices().begin(); + m_it_rev = m_track.state_indices().rbegin(); } /// Advance the iterator @@ -72,22 +81,27 @@ struct kalman_actor_state { /// @return true if the iterator reaches the end of vector TRACCC_HOST_DEVICE bool is_complete() { - if (!backward_mode && m_it == m_track_states.end()) { + if (!backward_mode && m_it == m_track.state_indices().end()) { return true; - } else if (backward_mode && m_it_rev == m_track_states.rend()) { + } else if (backward_mode && + m_it_rev == m_track.state_indices().rend()) { return true; } return false; } - // vector of track states - track_state_coll m_track_states; + /// Object describing the track fit + typename edm::track_fit_collection::device::proxy_type m_track; + /// All track states in the event + typename edm::track_state_collection::device m_track_states; + /// All measurements in the event + measurement_collection_types::const_device m_measurements; - // iterator for forward filtering - typename track_state_coll::iterator m_it; + /// Iterator for forward filtering over the track states + vecmem::device_vector::iterator m_it; - // iterator for backward filtering - typename track_state_coll::reverse_iterator m_it_rev; + /// Iterator for backward filtering over the track states + vecmem::device_vector::reverse_iterator m_it_rev; // The number of holes (The number of sensitive surfaces which do not // have a measurement for the track pattern) @@ -101,9 +115,6 @@ struct kalman_actor_state { template struct kalman_actor : detray::actor { - // Type declarations - using track_state_coll = vecmem::device_vector>; - // Actor state using state = kalman_actor_state; @@ -127,11 +138,14 @@ struct kalman_actor : detray::actor { // triggered only for sensitive surfaces if (navigation.is_on_sensitive()) { - auto& trk_state = actor_state(); + typename edm::track_state_collection::device::proxy_type + trk_state = actor_state(); // Increase the hole counts if the propagator fails to find the next // measurement - if (navigation.barcode() != trk_state.surface_link()) { + if (navigation.barcode() != + actor_state.m_measurements.at(trk_state.measurement_index()) + .surface_link) { if (!actor_state.backward_mode) { actor_state.n_holes++; } @@ -140,7 +154,7 @@ struct kalman_actor : detray::actor { // This track state is not a hole if (!actor_state.backward_mode) { - trk_state.is_hole = false; + trk_state.set_hole(false); } // Run Kalman Gain Updater @@ -157,11 +171,11 @@ struct kalman_actor : detray::actor { kalman_actor_direction::BIDIRECTIONAL) { // Forward filter res = gain_matrix_updater{}( - trk_state, propagation._stepping.bound_params(), - is_line); + trk_state, actor_state.m_measurements, + propagation._stepping.bound_params(), is_line); // Update the propagation flow - stepping.bound_params() = trk_state.filtered(); + stepping.bound_params() = trk_state.filtered_params(); } else { assert(false); } @@ -172,8 +186,8 @@ struct kalman_actor : detray::actor { kalman_actor_direction::BIDIRECTIONAL) { // Backward filter for smoothing res = two_filters_smoother{}( - trk_state, propagation._stepping.bound_params(), - is_line); + trk_state, actor_state.m_measurements, + propagation._stepping.bound_params(), is_line); } else { assert(false); } diff --git a/core/include/traccc/fitting/kalman_filter/kalman_fitter.hpp b/core/include/traccc/fitting/kalman_filter/kalman_fitter.hpp index 9cf771badc..147ded1048 100644 --- a/core/include/traccc/fitting/kalman_filter/kalman_fitter.hpp +++ b/core/include/traccc/fitting/kalman_filter/kalman_fitter.hpp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2022-2024 CERN for the benefit of the ACTS project + * (c) 2022-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -10,8 +10,10 @@ // Project include(s). #include "kalman_actor.hpp" #include "traccc/definitions/qualifiers.hpp" +#include "traccc/edm/measurement.hpp" +#include "traccc/edm/track_fit_collection.hpp" #include "traccc/edm/track_parameters.hpp" -#include "traccc/edm/track_state.hpp" +#include "traccc/edm/track_state_collection.hpp" #include "traccc/fitting/fitting_config.hpp" #include "traccc/fitting/kalman_filter/kalman_actor.hpp" #include "traccc/fitting/kalman_filter/kalman_step_aborter.hpp" @@ -102,29 +104,18 @@ class kalman_fitter { /// @param track_states the vector of track states TRACCC_HOST_DEVICE explicit state( - vecmem::data::vector_view> track_states, + const typename edm::track_fit_collection< + algebra_type>::device::proxy_type& track, + const typename edm::track_state_collection::device& + track_states, + const measurement_collection_types::const_device& measurements, vecmem::data::vector_view sequence_buffer) - : m_fit_actor_state( - vecmem::device_vector>( - track_states)), - m_sequencer_state( - vecmem::device_vector( - sequence_buffer)), - m_sequence_buffer(sequence_buffer) {} - - /// State constructor - /// - /// @param track_states the vector of track states - TRACCC_HOST_DEVICE - explicit state(const vecmem::device_vector>& - track_states, - vecmem::data::vector_view - sequence_buffer) - : m_fit_actor_state(track_states), + : m_fit_actor_state{track, track_states, measurements}, m_sequencer_state( vecmem::device_vector( sequence_buffer)), + m_fit_res{track}, m_sequence_buffer(sequence_buffer) {} /// @return the actor chain state @@ -151,7 +142,8 @@ class kalman_fitter { kalman_step_aborter::state m_step_aborter_state{}; /// Fitting result per track - fitting_result m_fit_res; + typename edm::track_fit_collection::device::proxy_type + m_fit_res; /// View object for barcode sequence vecmem::data::vector_view m_sequence_buffer; @@ -182,7 +174,10 @@ class kalman_fitter { // extrapolate from the filtered or smoothed state of next valid // track state. params = - fitter_state.m_fit_actor_state.m_track_states[0].smoothed(); + fitter_state.m_fit_actor_state.m_track_states + .at(fitter_state.m_fit_actor_state.m_track.state_indices() + .at(0)) + .filtered_params(); // Reset the iterator of kalman actor fitter_state.m_fit_actor_state.reset(); } @@ -243,7 +238,7 @@ class kalman_fitter { m_cfg.propagation.stepping.step_constraint); // Reset fitter statistics - fitter_state.m_fit_res.trk_quality.reset_quality(); + fitter_state.m_fit_res.reset_quality(); // Run forward filtering propagator.propagate(propagation, fitter_state()); @@ -264,32 +259,35 @@ class kalman_fitter { return kalman_fitter_status::ERROR_BARCODE_SEQUENCE_OVERFLOW; } + auto& track = fitter_state.m_fit_actor_state.m_track; auto& track_states = fitter_state.m_fit_actor_state.m_track_states; // Since the smoothed track parameter of the last surface can be // considered to be the filtered one, we can reversly iterate the // algorithm to obtain the smoothed parameter of other surfaces - for (auto it = track_states.rbegin(); it != track_states.rend(); ++it) { - if (!(*it).is_hole) { + for (auto it = track.state_indices().rbegin(); + it != track.state_indices().rend(); ++it) { + if (!track_states.at(*it).is_hole()) { fitter_state.m_fit_actor_state.m_it_rev = it; break; } // TODO: Return false because there is no valid track state // return false; } - auto& last = *fitter_state.m_fit_actor_state.m_it_rev; + auto last = track_states.at(*fitter_state.m_fit_actor_state.m_it_rev); - const scalar theta = last.filtered().theta(); + const scalar theta = last.filtered_params().theta(); if (theta <= 0.f || theta >= constant::pi) { return kalman_fitter_status::ERROR_THETA_ZERO; } - if (!std::isfinite(last.filtered().phi())) { + if (!std::isfinite(last.filtered_params().phi())) { return kalman_fitter_status::ERROR_INVERSION; } - last.smoothed().set_parameter_vector(last.filtered()); - last.smoothed().set_covariance(last.filtered().covariance()); + last.smoothed_params().set_parameter_vector(last.filtered_params()); + last.smoothed_params().set_covariance( + last.filtered_params().covariance()); last.smoothed_chi2() = last.filtered_chi2(); if (fitter_state.m_sequencer_state._sequence.empty()) { @@ -310,10 +308,10 @@ class kalman_fitter { m_cfg.propagation.stepping.path_limit); typename backward_propagator_type::state propagation( - last.smoothed(), m_field, m_detector, + last.smoothed_params(), m_field, m_detector, fitter_state.m_sequence_buffer, backward_cfg.context); propagation.set_particle(detail::correct_particle_hypothesis( - m_cfg.ptc_hypothesis, last.smoothed())); + m_cfg.ptc_hypothesis, last.smoothed_params())); assert(std::signbit( propagation._stepping.particle_hypothesis().charge()) == @@ -328,7 +326,7 @@ class kalman_fitter { // Synchronize the current barcode with the input track parameter while (propagation._navigation.get_target_barcode() != - last.smoothed().surface_link()) { + last.smoothed_params().surface_link()) { assert(!propagation._navigation.is_complete()); propagation._navigation.next(); } @@ -348,29 +346,31 @@ class kalman_fitter { // Fit parameter = smoothed track parameter of the first smoothed track // state - for (const auto& st : track_states) { - if (st.is_smoothed) { - fit_res.fit_params = st.smoothed(); + for (unsigned int i : fit_res.state_indices()) { + if (track_states.at(i).is_smoothed()) { + fit_res.params() = track_states.at(i).smoothed_params(); break; } } - for (const auto& trk_state : track_states) { + for (unsigned int i : fit_res.state_indices()) { - const detray::tracking_surface sf{m_detector, - trk_state.surface_link()}; - statistics_updater{}(fit_res, trk_state); + auto trk_state = track_states.at(i); + const detray::tracking_surface sf{ + m_detector, fitter_state.m_fit_actor_state.m_measurements + .at(trk_state.measurement_index()) + .surface_link}; + statistics_updater{}( + fit_res, trk_state, + fitter_state.m_fit_actor_state.m_measurements); } - // Track quality - auto& trk_quality = fit_res.trk_quality; - // Subtract the NDoF with the degree of freedom of the bound track (=5) - trk_quality.ndf = trk_quality.ndf - 5.f; - trk_quality.pval = prob(trk_quality.chi2, trk_quality.ndf); + fit_res.ndf() -= 5.f; + fit_res.pval() = prob(fit_res.chi2(), fit_res.ndf()); // The number of holes - trk_quality.n_holes = fitter_state.m_fit_actor_state.n_holes; + fit_res.nholes() = fitter_state.m_fit_actor_state.n_holes; } TRACCC_HOST_DEVICE @@ -380,22 +380,23 @@ class kalman_fitter { fitter_state.m_fit_actor_state.m_track_states; // NDF should always be positive for fitting - if (fit_res.trk_quality.ndf > 0) { - for (const auto& trk_state : track_states) { + if (fit_res.ndf() > 0) { + for (unsigned int i : fit_res.state_indices()) { + auto trk_state = track_states.at(i); // Fitting fails if any of non-hole track states is not smoothed - if (!trk_state.is_hole && !trk_state.is_smoothed) { - fit_res.fit_outcome = - fitter_outcome::FAILURE_NOT_ALL_SMOOTHED; + if (!trk_state.is_hole() && !trk_state.is_smoothed()) { + fit_res.fit_outcome() = + track_fit_outcome::FAILURE_NOT_ALL_SMOOTHED; return; } } // Fitting succeeds if any of non-hole track states is not smoothed - fit_res.fit_outcome = fitter_outcome::SUCCESS; + fit_res.fit_outcome() = track_fit_outcome::SUCCESS; return; } - fit_res.fit_outcome = fitter_outcome::FAILURE_NON_POSITIVE_NDF; + fit_res.fit_outcome() = track_fit_outcome::FAILURE_NON_POSITIVE_NDF; return; } diff --git a/core/include/traccc/fitting/kalman_filter/statistics_updater.hpp b/core/include/traccc/fitting/kalman_filter/statistics_updater.hpp index c4fce60960..733492f338 100644 --- a/core/include/traccc/fitting/kalman_filter/statistics_updater.hpp +++ b/core/include/traccc/fitting/kalman_filter/statistics_updater.hpp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2022-2024 CERN for the benefit of the ACTS project + * (c) 2022-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -9,7 +9,9 @@ // Project include(s). #include "traccc/definitions/qualifiers.hpp" -#include "traccc/edm/track_state.hpp" +#include "traccc/edm/measurement.hpp" +#include "traccc/edm/track_fit_collection.hpp" +#include "traccc/edm/track_state_collection.hpp" namespace traccc { @@ -27,21 +29,22 @@ struct statistics_updater { /// @param fit_res fitting information such as NDoF or Chi2 /// @param trk_state track state of the current surface TRACCC_HOST_DEVICE inline void operator()( - fitting_result& fit_res, - const track_state& trk_state) { + typename edm::track_fit_collection::device::proxy_type& + fit_res, + const typename edm::track_state_collection< + algebra_t>::const_device::const_proxy_type& trk_state, + const measurement_collection_types::const_device& measurements) { - if (!trk_state.is_hole) { + if (!trk_state.is_hole()) { // Measurement dimension - const unsigned int D = trk_state.get_measurement().meas_dim; + const unsigned int D = + measurements.at(trk_state.measurement_index()).meas_dim; - // Track quality - auto& trk_quality = fit_res.trk_quality; - - if (trk_state.is_smoothed) { + if (trk_state.is_smoothed()) { // NDoF = NDoF + number of coordinates per measurement - trk_quality.ndf += static_cast(D); - trk_quality.chi2 += trk_state.backward_chi2(); + fit_res.ndf() += static_cast(D); + fit_res.chi2() += trk_state.backward_chi2(); } } } diff --git a/core/include/traccc/fitting/kalman_filter/two_filters_smoother.hpp b/core/include/traccc/fitting/kalman_filter/two_filters_smoother.hpp index 57a6e4dcac..1c9c7b437f 100644 --- a/core/include/traccc/fitting/kalman_filter/two_filters_smoother.hpp +++ b/core/include/traccc/fitting/kalman_filter/two_filters_smoother.hpp @@ -10,7 +10,8 @@ // Project include(s). #include "traccc/definitions/qualifiers.hpp" #include "traccc/definitions/track_parametrization.hpp" -#include "traccc/edm/track_state.hpp" +#include "traccc/edm/measurement.hpp" +#include "traccc/edm/track_state_collection.hpp" #include "traccc/fitting/status_codes.hpp" namespace traccc { @@ -33,21 +34,24 @@ struct two_filters_smoother { /// /// @return true if the update succeeds [[nodiscard]] TRACCC_HOST_DEVICE inline kalman_fitter_status operator()( - track_state& trk_state, + typename edm::track_state_collection::device::proxy_type& + trk_state, + const measurement_collection_types::const_device& measurements, bound_track_parameters& bound_params, const bool is_line) const { - const auto D = trk_state.get_measurement().meas_dim; - + const auto D = measurements.at(trk_state.measurement_index()).meas_dim; assert(D == 1u || D == 2u); - return smoothe(trk_state, bound_params, D, is_line); + return smoothe(trk_state, measurements, bound_params, D, is_line); } // Reference: The Optimun Linear Smoother as a Combination of Two Optimum // Linear Filters [[nodiscard]] TRACCC_HOST_DEVICE inline kalman_fitter_status smoothe( - track_state& trk_state, + typename edm::track_state_collection::device::proxy_type& + trk_state, + const measurement_collection_types::const_device& measurements, bound_track_parameters& bound_params, const unsigned int dim, const bool is_line) const { @@ -57,19 +61,18 @@ struct two_filters_smoother { assert(!bound_params.is_invalid()); assert(!bound_params.surface_link().is_invalid()); - assert(trk_state.filtered().surface_link() == + assert(trk_state.filtered_params().surface_link() == bound_params.surface_link()); // Do not smoothe if the forward pass produced an error - if (trk_state.filtered().is_invalid()) { + if (trk_state.filtered_params().is_invalid()) { return kalman_fitter_status::ERROR_INVALID_TRACK_STATE; } - const auto meas = trk_state.get_measurement(); - // Measurement data on surface - const matrix_type meas_local = - trk_state.template measurement_local(); + matrix_type meas_local; + trk_state.template get_measurement_local(measurements, + meas_local); assert((dim > 1) || (getter::element(meas_local, 1u, 0u) == 0.f)); @@ -84,7 +87,7 @@ struct two_filters_smoother { const matrix_type predicted_cov_inv = matrix::inverse(predicted_cov); const matrix_type filtered_cov_inv = - matrix::inverse(trk_state.filtered().covariance()); + matrix::inverse(trk_state.filtered_params().covariance()); // Eq (3.38) of "Pattern Recognition, Tracking and Vertex // Reconstruction in Particle Detectors" @@ -98,14 +101,16 @@ struct two_filters_smoother { // Eq (3.38) of "Pattern Recognition, Tracking and Vertex // Reconstruction in Particle Detectors" const matrix_type smoothed_vec = - smoothed_cov * (filtered_cov_inv * trk_state.filtered().vector() + - predicted_cov_inv * predicted_vec); + smoothed_cov * + (filtered_cov_inv * trk_state.filtered_params().vector() + + predicted_cov_inv * predicted_vec); - trk_state.smoothed().set_vector(smoothed_vec); - trk_state.smoothed().set_covariance(smoothed_cov); - - matrix_type H = meas.subs.template projector(); + trk_state.smoothed_params().set_vector(smoothed_vec); + trk_state.smoothed_params().set_covariance(smoothed_cov); + matrix_type H = + measurements.at(trk_state.measurement_index()) + .subs.template projector(); if (dim == 1) { getter::element(H, 1u, 0u) = 0.f; getter::element(H, 1u, 1u) = 0.f; @@ -114,8 +119,9 @@ struct two_filters_smoother { const matrix_type residual_smt = meas_local - H * smoothed_vec; // Spatial resolution (Measurement covariance) - matrix_type V = trk_state.template measurement_covariance(); - + matrix_type V; + trk_state.template get_measurement_covariance(measurements, + V); if (dim == 1) { getter::element(V, 1u, 1u) = 1.f; } @@ -217,7 +223,7 @@ struct two_filters_smoother { // Wrap the phi in the range of [-pi, pi] wrap_phi(bound_params); - trk_state.is_smoothed = true; + trk_state.set_smoothed(); assert(!bound_params.is_invalid()); diff --git a/core/include/traccc/fitting/kalman_fitting_algorithm.hpp b/core/include/traccc/fitting/kalman_fitting_algorithm.hpp index a69016c696..11c1621f65 100644 --- a/core/include/traccc/fitting/kalman_fitting_algorithm.hpp +++ b/core/include/traccc/fitting/kalman_fitting_algorithm.hpp @@ -9,8 +9,9 @@ // Project include(s). #include "traccc/bfield/magnetic_field.hpp" +#include "traccc/edm/measurement.hpp" #include "traccc/edm/track_candidate_container.hpp" -#include "traccc/edm/track_state.hpp" +#include "traccc/edm/track_fit_container.hpp" #include "traccc/fitting/fitting_config.hpp" #include "traccc/geometry/detector.hpp" #include "traccc/utils/algorithm.hpp" @@ -24,15 +25,16 @@ // System include(s). #include +#include 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 edm::track_candidate_container::const_view&)>, public messaging { @@ -41,7 +43,7 @@ class kalman_fitting_algorithm /// Configuration type using config_type = fitting_config; /// Output type - using output_type = track_state_container_types::host; + using output_type = edm::track_fit_container::host; /// Constructor with the algorithm's configuration /// diff --git a/core/src/ambiguity_resolution/legacy/greedy_ambiguity_resolution_algorithm.cpp b/core/src/ambiguity_resolution/legacy/greedy_ambiguity_resolution_algorithm.cpp index f6e5839c14..703da51ea8 100644 --- a/core/src/ambiguity_resolution/legacy/greedy_ambiguity_resolution_algorithm.cpp +++ b/core/src/ambiguity_resolution/legacy/greedy_ambiguity_resolution_algorithm.cpp @@ -25,7 +25,6 @@ // Project include(s). #include "traccc/definitions/primitives.hpp" #include "traccc/definitions/qualifiers.hpp" -#include "traccc/edm/track_state.hpp" #include "traccc/utils/algorithm.hpp" // Greedy ambiguity resolution adapted from ACTS code diff --git a/core/src/fitting/kalman_fitting_algorithm_default_detector.cpp b/core/src/fitting/kalman_fitting_algorithm_default_detector.cpp index ba1da18e43..6d0f62d4d0 100644 --- a/core/src/fitting/kalman_fitting_algorithm_default_detector.cpp +++ b/core/src/fitting/kalman_fitting_algorithm_default_detector.cpp @@ -28,16 +28,16 @@ kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()( covfie::field>::view_t> fitter{det, bfield.as_view>(), m_config}; - return details::kalman_fitting(fitter, track_candidates, m_mr.get(), - m_copy.get()); + return details::kalman_fitting( + fitter, track_candidates, m_mr.get(), m_copy.get()); } else if (bfield.is>()) { traccc::details::kalman_fitter_t< default_detector::host, covfie::field>::view_t> fitter{det, bfield.as_view>(), m_config}; - return details::kalman_fitting(fitter, track_candidates, m_mr.get(), - m_copy.get()); + return details::kalman_fitting( + fitter, track_candidates, m_mr.get(), m_copy.get()); } else { throw std::invalid_argument( "Unsupported b-field type received in " diff --git a/core/src/fitting/kalman_fitting_algorithm_telescope_detector.cpp b/core/src/fitting/kalman_fitting_algorithm_telescope_detector.cpp index 122ade3a49..5e022ae2d1 100644 --- a/core/src/fitting/kalman_fitting_algorithm_telescope_detector.cpp +++ b/core/src/fitting/kalman_fitting_algorithm_telescope_detector.cpp @@ -28,16 +28,16 @@ kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()( covfie::field>::view_t> fitter{det, bfield.as_view>(), m_config}; - return details::kalman_fitting(fitter, track_candidates, m_mr.get(), - m_copy.get()); + return details::kalman_fitting( + fitter, track_candidates, m_mr.get(), m_copy.get()); } else if (bfield.is>()) { traccc::details::kalman_fitter_t< telescope_detector::host, covfie::field>::view_t> fitter{det, bfield.as_view>(), m_config}; - return details::kalman_fitting(fitter, track_candidates, m_mr.get(), - m_copy.get()); + return details::kalman_fitting( + fitter, track_candidates, m_mr.get(), m_copy.get()); } else { throw std::invalid_argument( "Unsupported b-field type received in " diff --git a/performance/CMakeLists.txt b/performance/CMakeLists.txt index 956e26fcce..7eb080cd6e 100644 --- a/performance/CMakeLists.txt +++ b/performance/CMakeLists.txt @@ -44,16 +44,19 @@ traccc_add_library( traccc_performance performance TYPE SHARED "src/performance/details/is_same_scalar.cpp" "include/traccc/performance/details/is_same_object.hpp" "include/traccc/performance/impl/is_same_object.ipp" - "include/traccc/performance/impl/is_same_fitting_result.ipp" "include/traccc/performance/impl/is_same_seed.ipp" "include/traccc/performance/impl/is_same_measurement.ipp" "include/traccc/performance/impl/is_same_spacepoint.ipp" "include/traccc/performance/impl/is_same_track_candidate.ipp" + "include/traccc/performance/impl/is_same_track_fit.ipp" "include/traccc/performance/impl/is_same_track_parameters.ipp" + "include/traccc/performance/impl/is_same_track_state.ipp" "src/performance/details/is_same_object.cpp" "include/traccc/performance/details/comparator_factory.hpp" "include/traccc/performance/impl/comparator_factory.ipp" "include/traccc/performance/impl/seed_comparator_factory.ipp" + "include/traccc/performance/impl/track_candidate_comparator_factory.ipp" + "include/traccc/performance/impl/track_fit_comparator_factory.ipp" # Collection/container comparison code. "include/traccc/performance/collection_comparator.hpp" "include/traccc/performance/impl/collection_comparator.ipp" diff --git a/performance/include/traccc/efficiency/finding_performance_writer.hpp b/performance/include/traccc/efficiency/finding_performance_writer.hpp index c24dda5cb2..27a193471b 100644 --- a/performance/include/traccc/efficiency/finding_performance_writer.hpp +++ b/performance/include/traccc/efficiency/finding_performance_writer.hpp @@ -15,8 +15,10 @@ #include "traccc/utils/truth_matching_config.hpp" // Project include(s). +#include "traccc/edm/measurement.hpp" #include "traccc/edm/track_candidate_collection.hpp" -#include "traccc/edm/track_state.hpp" +#include "traccc/edm/track_fit_collection.hpp" +#include "traccc/edm/track_state_collection.hpp" #include "traccc/utils/event_data.hpp" // System include(s). @@ -75,8 +77,13 @@ class finding_performance_writer : public messaging { const measurement_collection_types::const_view& measurements_view, const event_data& evt_data); - void write(const track_state_container_types::const_view& track_states_view, - const event_data& evt_data); + void write( + const edm::track_fit_collection::const_view& + track_fit_view, + const edm::track_state_collection::const_view& + track_states_view, + const measurement_collection_types::const_view& measurements_view, + const event_data& evt_data); void finalize(); diff --git a/performance/include/traccc/performance/details/comparator_factory.hpp b/performance/include/traccc/performance/details/comparator_factory.hpp index 09b185cba6..51410ed484 100644 --- a/performance/include/traccc/performance/details/comparator_factory.hpp +++ b/performance/include/traccc/performance/details/comparator_factory.hpp @@ -42,3 +42,4 @@ class comparator_factory { // Include the specialised implementation(s). #include "traccc/performance/impl/seed_comparator_factory.ipp" #include "traccc/performance/impl/track_candidate_comparator_factory.ipp" +#include "traccc/performance/impl/track_fit_comparator_factory.ipp" diff --git a/performance/include/traccc/performance/details/is_same_object.hpp b/performance/include/traccc/performance/details/is_same_object.hpp index 98954cdb79..675c950ae2 100644 --- a/performance/include/traccc/performance/details/is_same_object.hpp +++ b/performance/include/traccc/performance/details/is_same_object.hpp @@ -45,9 +45,10 @@ class is_same_object { #include "traccc/performance/impl/is_same_object.ipp" // Include specialisations for the core library types -#include "traccc/performance/impl/is_same_fitting_result.ipp" #include "traccc/performance/impl/is_same_measurement.ipp" #include "traccc/performance/impl/is_same_seed.ipp" #include "traccc/performance/impl/is_same_spacepoint.ipp" #include "traccc/performance/impl/is_same_track_candidate.ipp" +#include "traccc/performance/impl/is_same_track_fit.ipp" #include "traccc/performance/impl/is_same_track_parameters.ipp" +#include "traccc/performance/impl/is_same_track_state.ipp" diff --git a/performance/include/traccc/performance/details/projector.hpp b/performance/include/traccc/performance/details/projector.hpp index 93bfc64398..a6447385ee 100644 --- a/performance/include/traccc/performance/details/projector.hpp +++ b/performance/include/traccc/performance/details/projector.hpp @@ -47,12 +47,4 @@ struct projector> { } }; -template <> -struct projector> { - static constexpr bool exists = true; - - float operator()(const fitting_result& i) { - return static_cast(i.trk_quality.ndf); - } -}; } // namespace traccc::details diff --git a/performance/include/traccc/performance/impl/is_same_fitting_result.ipp b/performance/include/traccc/performance/impl/is_same_fitting_result.ipp deleted file mode 100644 index e270d31c86..0000000000 --- a/performance/include/traccc/performance/impl/is_same_fitting_result.ipp +++ /dev/null @@ -1,35 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2023 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -#pragma once - -// Project include(s). -#include "traccc/edm/track_state.hpp" - -namespace traccc::details { - -/// @c traccc::is_same_object specialisation for @c traccc::fitting_result -template <> -class is_same_object> { - - public: - /// Constructor with a reference object, and an allowed uncertainty - is_same_object(const fitting_result& ref, - scalar unc = float_epsilon); - - /// Specialised implementation for @c traccc::measurement - bool operator()(const fitting_result& obj) const; - - private: - /// The reference object - std::reference_wrapper> m_ref; - /// The uncertainty - scalar m_unc; - -}; // class is_same_object> - -} // namespace traccc::details diff --git a/performance/include/traccc/performance/impl/is_same_track_fit.ipp b/performance/include/traccc/performance/impl/is_same_track_fit.ipp new file mode 100644 index 0000000000..8d452f1060 --- /dev/null +++ b/performance/include/traccc/performance/impl/is_same_track_fit.ipp @@ -0,0 +1,105 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2023-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s). +#include "traccc/edm/measurement.hpp" +#include "traccc/edm/track_fit_collection.hpp" +#include "traccc/edm/track_state_collection.hpp" + +// Local include(s). +#include "traccc/performance/impl/is_same_track_state.ipp" + +namespace traccc::details { + +/// @c traccc::is_same_object specialisation for @c traccc::edm::track_fit +template +class is_same_object> { + + public: + /// Constructor with a reference object, and an allowed uncertainty + is_same_object( + const measurement_collection_types::const_view& ref_meas, + const measurement_collection_types::const_view& test_meas, + const edm::track_state_collection::const_view& + ref_states, + const edm::track_state_collection::const_view& + test_states, + const edm::track_fit& ref, scalar unc = float_epsilon) + : m_ref_meas(ref_meas), + m_test_meas(test_meas), + m_ref_states(ref_states), + m_test_states(test_states), + m_ref(ref), + m_unc(unc) {} + + /// Specialised implementation for @c traccc::edm::track_fit + bool operator()(const edm::track_fit& obj) const { + + // Compare the fit outcomes. + if (obj.fit_outcome() != m_ref.fit_outcome()) { + return false; + } + // Compare the parameters. + if (!is_same_object>(m_ref.params(), + m_unc)(obj.params())) { + return false; + } + // Compare the scalar values. + if (!is_same_scalar(obj.ndf(), m_ref.ndf(), m_unc) || + !is_same_scalar(obj.chi2(), m_ref.chi2(), m_unc) || + !is_same_scalar(obj.pval(), m_ref.pval(), m_unc)) { + return false; + } + // Compare the number of holes. + if (obj.nholes() != m_ref.nholes()) { + return false; + } + + // The two tracks need to have the same number of states. + if (obj.state_indices().size() != m_ref.state_indices().size()) { + return false; + } + + // Now compare the track states one by one. + const edm::track_state_collection::const_device + ref_states{m_ref_states}; + const edm::track_state_collection::const_device + test_states{m_test_states}; + for (unsigned int i = 0; i < obj.state_indices().size(); ++i) { + if (!is_same_object::const_device::const_proxy_type>( + m_ref_meas, m_test_meas, + ref_states.at(m_ref.state_indices()[i]), + m_unc)(test_states.at(obj.state_indices()[i]))) { + return false; + } + } + + // If we got here, the two track fits are the same. + return true; + } + + private: + /// Measurements for the reference object + const measurement_collection_types::const_view m_ref_meas; + /// Measurements for the test object + const measurement_collection_types::const_view m_test_meas; + /// States for the reference object + const edm::track_state_collection::const_view m_ref_states; + /// States for the test object + const edm::track_state_collection::const_view + m_test_states; + /// The reference object + const edm::track_fit m_ref; + /// The uncertainty + scalar m_unc; + +}; // class is_same_object> + +} // namespace traccc::details diff --git a/performance/include/traccc/performance/impl/is_same_track_state.ipp b/performance/include/traccc/performance/impl/is_same_track_state.ipp new file mode 100644 index 0000000000..453d0e7b2b --- /dev/null +++ b/performance/include/traccc/performance/impl/is_same_track_state.ipp @@ -0,0 +1,82 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2023-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s). +#include "traccc/edm/measurement.hpp" +#include "traccc/edm/track_state_collection.hpp" + +// Local include(s). +#include "traccc/performance/impl/is_same_measurement.ipp" + +namespace traccc::details { + +/// @c traccc::is_same_object specialisation for +/// @c traccc::edm::track_state_collection +template +class is_same_object> { + + public: + /// Constructor with a reference object, and an allowed uncertainty + is_same_object(const measurement_collection_types::const_view& ref_meas, + const measurement_collection_types::const_view& test_meas, + const edm::track_state& ref, scalar unc = float_epsilon) + : m_ref_meas(ref_meas), + m_test_meas(test_meas), + m_ref(ref), + m_unc(unc) {} + + /// Specialised implementation for @c traccc::edm::track_state + bool operator()(const edm::track_state& obj) const { + + // Compare the state words. + if (obj.state() != m_ref.state()) { + return false; + } + // Compare the chi2 values. + if (!is_same_scalar(obj.filtered_chi2(), m_ref.filtered_chi2(), + m_unc) || + !is_same_scalar(obj.smoothed_chi2(), m_ref.smoothed_chi2(), + m_unc) || + !is_same_scalar(obj.backward_chi2(), m_ref.backward_chi2(), + m_unc)) { + return false; + } + // Compare the fitted parameters. + if (!is_same_object>( + m_ref.filtered_params(), m_unc)(obj.filtered_params()) || + !is_same_object>( + m_ref.smoothed_params(), m_unc)(obj.smoothed_params())) { + return false; + } + // Compare the measurements that they point at. + const measurement_collection_types::const_device ref_meas{m_ref_meas}; + const measurement_collection_types::const_device test_meas{m_test_meas}; + if (!is_same_object( + ref_meas.at(m_ref.measurement_index()), + m_unc)(test_meas.at(obj.measurement_index()))) { + return false; + } + + // If we got here, the two track states are the same. + return true; + } + + private: + /// Measurements for the reference object + const measurement_collection_types::const_view m_ref_meas; + /// Measurements for the test object + const measurement_collection_types::const_view m_test_meas; + /// The reference object + const edm::track_state m_ref; + /// The uncertainty + scalar m_unc; + +}; // class is_same_object> + +} // namespace traccc::details diff --git a/performance/include/traccc/performance/impl/track_fit_comparator_factory.ipp b/performance/include/traccc/performance/impl/track_fit_comparator_factory.ipp new file mode 100644 index 0000000000..040ff55378 --- /dev/null +++ b/performance/include/traccc/performance/impl/track_fit_comparator_factory.ipp @@ -0,0 +1,60 @@ +/** 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 + +// Project include(s). +#include "traccc/edm/measurement.hpp" +#include "traccc/edm/track_fit_collection.hpp" +#include "traccc/edm/track_state_collection.hpp" + +// Local include(s). +#include "traccc/performance/details/is_same_object.hpp" + +namespace traccc::details { + +/// @c traccc::details::comparator_factory specialisation for +/// @c traccc::edm::track_fit +template +class comparator_factory> { + + public: + /// Constructor with all necessary arguments + comparator_factory( + const measurement_collection_types::const_view& ref_meas, + const measurement_collection_types::const_view& test_meas, + const edm::track_state_collection::const_view& + ref_states, + const edm::track_state_collection::const_view& + test_states) + : m_ref_meas(ref_meas), + m_test_meas(test_meas), + m_ref_states(ref_states), + m_test_states(test_states) {} + + /// Instantiate an instance of a comparator object + is_same_object> make_comparator( + const edm::track_fit& ref, scalar unc = float_epsilon) const { + + return is_same_object>( + m_ref_meas, m_test_meas, m_ref_states, m_test_states, ref, unc); + } + + private: + /// Measurement container for the reference track candidates + const measurement_collection_types::const_view m_ref_meas; + /// Measurement container for the test track candidates + const measurement_collection_types::const_view m_test_meas; + /// States for the reference object + const edm::track_state_collection::const_view m_ref_states; + /// States for the test object + const edm::track_state_collection::const_view + m_test_states; + +}; // class comparator_factory + +} // namespace traccc::details diff --git a/performance/include/traccc/resolution/fitting_performance_writer.hpp b/performance/include/traccc/resolution/fitting_performance_writer.hpp index ee18a367be..aba0d93df3 100644 --- a/performance/include/traccc/resolution/fitting_performance_writer.hpp +++ b/performance/include/traccc/resolution/fitting_performance_writer.hpp @@ -13,9 +13,11 @@ #include "traccc/utils/messaging.hpp" // Project include(s). +#include "traccc/edm/measurement.hpp" #include "traccc/edm/particle.hpp" +#include "traccc/edm/track_fit_collection.hpp" #include "traccc/edm/track_parameters.hpp" -#include "traccc/edm/track_state.hpp" +#include "traccc/edm/track_state_collection.hpp" #include "traccc/utils/event_data.hpp" // System include(s). @@ -52,28 +54,36 @@ class fitting_performance_writer : public messaging { /// Fill the tracking results into the histograms /// - /// @param track_states_per_track vector of track states of a track + /// @param track The fitted track to measure/write the performance of + /// @param track_states All reconstructed track states + /// @param measurements All reconstructed measurements /// @param det detector object /// @param evt_map event map to find the truth values template - void write(const track_state_collection_types::host& track_states_per_track, - const fitting_result& fit_res, + void write(const edm::track_fit_collection< + traccc::default_algebra>::host::proxy_type track, + const edm::track_state_collection::host& + track_states, + const measurement_collection_types::host& measurements, const detector_t& det, event_data& evt_data, const detector_t::geometry_context& ctx = {}) { static_assert(std::same_as); - if (fit_res.fit_outcome != fitter_outcome::SUCCESS) { + if (track.fit_outcome() != track_fit_outcome::SUCCESS) { return; } // Get the first smoothed track state - const auto& trk_state = *std::find_if( - track_states_per_track.begin(), track_states_per_track.end(), - [](const auto& st) { return st.is_smoothed; }); - assert(!trk_state.is_hole); - assert(trk_state.is_smoothed); + const unsigned int trk_state_idx = + *std::find_if(track.state_indices().begin(), + track.state_indices().end(), [&](unsigned int idx) { + return track_states.at(idx).is_smoothed(); + }); + const auto trk_state = track_states.at(trk_state_idx); + assert(!trk_state.is_hole()); + assert(trk_state.is_smoothed()); bool use_found = !evt_data.m_found_meas_to_ptc_map.empty(); @@ -84,7 +94,7 @@ class fitting_performance_writer : public messaging { meas_to_param_map = use_found ? evt_data.m_found_meas_to_param_map : evt_data.m_meas_to_param_map; - const measurement meas = trk_state.get_measurement(); + const measurement meas = measurements.at(trk_state.measurement_index()); // Find the contributing particle // @todo: Use identify_contributing_particles function @@ -111,8 +121,8 @@ class fitting_performance_writer : public messaging { truth_param.set_qop(ptc.charge / vector::norm(global_mom)); // For the moment, only fill with the first measurements - write_res(truth_param, trk_state.smoothed(), ptc); - write_stat(fit_res, track_states_per_track); + write_res(truth_param, trk_state.smoothed_params(), ptc); + write_stat(track, track_states, measurements); } /// Writing caches into the file @@ -125,8 +135,12 @@ class fitting_performance_writer : public messaging { const particle& ptc); /// Non-templated part of the @c write(...) function - void write_stat(const fitting_result& fit_res, - const track_state_collection_types::host& track_states); + void write_stat( + const edm::track_fit_collection< + traccc::default_algebra>::host::proxy_type track, + const edm::track_state_collection::host& + track_states, + const measurement_collection_types::host& measurements); /// Configuration for the tool config m_cfg; diff --git a/performance/src/efficiency/finding_performance_writer.cpp b/performance/src/efficiency/finding_performance_writer.cpp index 31f7b027bd..f593b62853 100644 --- a/performance/src/efficiency/finding_performance_writer.cpp +++ b/performance/src/efficiency/finding_performance_writer.cpp @@ -129,23 +129,33 @@ std::vector> prepare_data( * with its corresponding measurements. */ std::vector> prepare_data( - const track_state_container_types::const_view& track_states_view) { + const edm::track_fit_collection::const_view& + track_fit_view, + const edm::track_state_collection::const_view& + track_states_view, + const measurement_collection_types::const_view& measurements_view) { std::vector> result; - // Iterate over the tracks. - track_state_container_types::const_device track_states(track_states_view); + // Set up the input containers. + const edm::track_fit_collection::const_device track_fits( + track_fit_view); + const edm::track_state_collection::const_device + track_states(track_states_view); + const measurement_collection_types::const_device measurements{ + measurements_view}; - const unsigned int n_tracks = track_states.size(); + // Iterate over the tracks. + const unsigned int n_tracks = track_fits.size(); result.reserve(n_tracks); for (unsigned int i = 0; i < n_tracks; i++) { - auto const& [fit_res, states] = track_states.at(i); - std::vector measurements; - measurements.reserve(states.size()); - for (const auto& st : states) { - measurements.push_back(st.get_measurement()); + std::vector result_measurements; + result_measurements.reserve(track_fits.at(i).state_indices().size()); + for (unsigned int st_idx : track_fits.state_indices().at(i)) { + result_measurements.push_back( + measurements.at(track_states.at(st_idx).measurement_index())); } - result.push_back(std::move(measurements)); + result.push_back(std::move(result_measurements)); } return result; } @@ -337,10 +347,15 @@ void finding_performance_writer::write( /// For ambiguity resolution void finding_performance_writer::write( - const track_state_container_types::const_view& track_states_view, + const edm::track_fit_collection::const_view& + track_fit_view, + const edm::track_state_collection::const_view& + track_states_view, + const measurement_collection_types::const_view& measurements_view, const event_data& evt_data) { + std::vector> tracks = - prepare_data(track_states_view); + prepare_data(track_fit_view, track_states_view, measurements_view); write_common(tracks, evt_data); } diff --git a/performance/src/performance/details/is_same_object.cpp b/performance/src/performance/details/is_same_object.cpp index 4c62051b4d..7ff648ac44 100644 --- a/performance/src/performance/details/is_same_object.cpp +++ b/performance/src/performance/details/is_same_object.cpp @@ -52,23 +52,4 @@ bool is_same_object>::operator()( /// @} -/// @name Implementation for -/// @c traccc::details::is_same_object -/// @{ - -is_same_object>::is_same_object( - const fitting_result& ref, scalar unc) - : m_ref(ref), m_unc(unc) {} - -bool is_same_object>::operator()( - const fitting_result& obj) const { - - return (is_same_object>(m_ref.get().fit_params, - m_unc)(obj.fit_params) && - is_same_scalar(obj.trk_quality.ndf, m_ref.get().trk_quality.ndf, - m_unc)); -} - -/// @} - } // namespace traccc::details diff --git a/performance/src/resolution/fitting_performance_writer.cpp b/performance/src/resolution/fitting_performance_writer.cpp index d2483794d7..f4961eb874 100644 --- a/performance/src/resolution/fitting_performance_writer.cpp +++ b/performance/src/resolution/fitting_performance_writer.cpp @@ -83,13 +83,17 @@ void fitting_performance_writer::write_res( } void fitting_performance_writer::write_stat( - const fitting_result& fit_res, - const track_state_collection_types::host& track_states) { + const edm::track_fit_collection::host::proxy_type + track, + const edm::track_state_collection::host& + track_states, + const measurement_collection_types::host& measurements) { - m_data->m_stat_plot_tool.fill(m_data->m_stat_plot_cache, fit_res); + m_data->m_stat_plot_tool.fill(m_data->m_stat_plot_cache, track); - for (const auto& trk_state : track_states) { - m_data->m_stat_plot_tool.fill(m_data->m_stat_plot_cache, trk_state); + for (std::size_t i : track.state_indices()) { + m_data->m_stat_plot_tool.fill(m_data->m_stat_plot_cache, + track_states.at(i), measurements); } } diff --git a/performance/src/resolution/stat_plot_tool.cpp b/performance/src/resolution/stat_plot_tool.cpp index 2cc04387b0..adae416740 100644 --- a/performance/src/resolution/stat_plot_tool.cpp +++ b/performance/src/resolution/stat_plot_tool.cpp @@ -53,45 +53,6 @@ void stat_plot_tool::book(stat_plot_cache& cache) const { #endif // TRACCC_HAVE_ROOT } -void stat_plot_tool::fill( - stat_plot_cache& cache, - const fitting_result& fit_res) const { - - // Avoid unused variable warnings when building the code without ROOT. - (void)cache; - (void)fit_res; - -#ifdef TRACCC_HAVE_ROOT - const auto& ndf = fit_res.trk_quality.ndf; - const auto& chi2 = fit_res.trk_quality.chi2; - cache.ndf_hist->Fill(ndf); - cache.chi2_hist->Fill(chi2); - cache.reduced_chi2_hist->Fill(chi2 / ndf); - cache.pval_hist->Fill(fit_res.trk_quality.pval); -#endif // TRACCC_HAVE_ROOT -} - -void stat_plot_tool::fill( - stat_plot_cache& cache, - const track_state& trk_state) const { - - // Avoid unused variable warnings when building the code without ROOT. - (void)cache; - (void)trk_state; - -#ifdef TRACCC_HAVE_ROOT - const unsigned int D = trk_state.get_measurement().meas_dim; - const auto filtered_chi2 = trk_state.filtered_chi2(); - const auto smoothed_chi2 = trk_state.smoothed_chi2(); - cache.chi2_filtered_hist[D]->Fill(filtered_chi2); - cache.chi2_smoothed_hist[D]->Fill(smoothed_chi2); - cache.pval_filtered_hist[D]->Fill( - prob(filtered_chi2, static_cast(D))); - cache.pval_smoothed_hist[D]->Fill( - prob(smoothed_chi2, static_cast(D))); -#endif // TRACCC_HAVE_ROOT -} - void stat_plot_tool::fill(stat_plot_cache& cache, double purity, double completeness) const { diff --git a/performance/src/resolution/stat_plot_tool.hpp b/performance/src/resolution/stat_plot_tool.hpp index e49aeb7012..806b826cb1 100644 --- a/performance/src/resolution/stat_plot_tool.hpp +++ b/performance/src/resolution/stat_plot_tool.hpp @@ -13,7 +13,8 @@ // Project include(s). #include "traccc/edm/track_candidate_collection.hpp" -#include "traccc/edm/track_state.hpp" +#include "traccc/edm/track_fit_collection.hpp" +#include "traccc/edm/track_state_collection.hpp" namespace traccc { @@ -68,15 +69,16 @@ class stat_plot_tool { /// /// @param cache the cache for statistics plots /// @param fit_res fitting information that contains statistics - void fill(stat_plot_cache& cache, - const fitting_result& fit_res) const; + template + void fill(stat_plot_cache& cache, const edm::track_fit& fit_res) const; /// @brief fill the cache /// /// @param cache the cache for statistics plots /// @param trk_state track state at local measurements - void fill(stat_plot_cache& cache, - const track_state& trk_state) const; + template + void fill(stat_plot_cache& cache, const edm::track_state& trk_state, + const measurement_collection_types::host& measurements) const; /// @brief fill the cache /// @param cache the cache for statistics plots diff --git a/performance/src/resolution/stat_plot_tool.ipp b/performance/src/resolution/stat_plot_tool.ipp index b9f2bb9ed4..ad63a928cc 100644 --- a/performance/src/resolution/stat_plot_tool.ipp +++ b/performance/src/resolution/stat_plot_tool.ipp @@ -7,6 +7,7 @@ #pragma once +// Project include(s). #include "traccc/utils/prob.hpp" namespace traccc { @@ -29,4 +30,46 @@ void stat_plot_tool::fill(stat_plot_cache& cache, #endif // TRACCC_HAVE_ROOT } +template +void stat_plot_tool::fill(stat_plot_cache& cache, + const edm::track_fit& fit_res) const { + + // Avoid unused variable warnings when building the code without ROOT. + (void)cache; + (void)fit_res; + +#ifdef TRACCC_HAVE_ROOT + const auto& ndf = fit_res.ndf(); + const auto& chi2 = fit_res.chi2(); + cache.ndf_hist->Fill(ndf); + cache.chi2_hist->Fill(chi2); + cache.reduced_chi2_hist->Fill(chi2 / ndf); + cache.pval_hist->Fill(fit_res.pval()); +#endif // TRACCC_HAVE_ROOT +} + +template +void stat_plot_tool::fill( + stat_plot_cache& cache, const edm::track_state& trk_state, + const measurement_collection_types::host& measurements) const { + + // Avoid unused variable warnings when building the code without ROOT. + (void)cache; + (void)trk_state; + (void)measurements; + +#ifdef TRACCC_HAVE_ROOT + const unsigned int D = + measurements.at(trk_state.measurement_index()).meas_dim; + const auto filtered_chi2 = trk_state.filtered_chi2(); + const auto smoothed_chi2 = trk_state.smoothed_chi2(); + cache.chi2_filtered_hist[D]->Fill(filtered_chi2); + cache.chi2_smoothed_hist[D]->Fill(smoothed_chi2); + cache.pval_filtered_hist[D]->Fill( + prob(filtered_chi2, static_cast(D))); + cache.pval_smoothed_hist[D]->Fill( + prob(smoothed_chi2, static_cast(D))); +#endif // TRACCC_HAVE_ROOT +} + } // namespace traccc From 9ca8243a3f76a999d8bf1f52979ca27701e05e38 Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Mon, 4 Aug 2025 11:38:23 +0200 Subject: [PATCH 02/11] Adapted traccc::device_common and traccc::cuda to the SoA track fit EDM. --- .../finding/device/impl/find_tracks.ipp | 25 +++++++---- .../include/traccc/fitting/device/fit.hpp | 7 +-- .../traccc/fitting/device/fit_backward.hpp | 18 +++++--- .../traccc/fitting/device/fit_forward.hpp | 18 +++++--- .../traccc/fitting/device/fit_prelude.hpp | 45 ++++++++++++++----- .../cuda/fitting/kalman_fitting_algorithm.hpp | 8 ++-- device/cuda/src/fitting/kalman_fitting.cuh | 30 ++++++++----- .../cuda/src/fitting/kernels/fit_prelude.cu | 12 ++--- .../cuda/src/fitting/kernels/fit_prelude.hpp | 4 +- 9 files changed, 111 insertions(+), 56 deletions(-) diff --git a/device/common/include/traccc/finding/device/impl/find_tracks.ipp b/device/common/include/traccc/finding/device/impl/find_tracks.ipp index ab21f177b8..b7146191e2 100644 --- a/device/common/include/traccc/finding/device/impl/find_tracks.ipp +++ b/device/common/include/traccc/finding/device/impl/find_tracks.ipp @@ -178,8 +178,10 @@ TRACCC_HOST_DEVICE inline void find_tracks( barrier.blockBarrier(); - std::optional, - unsigned int, unsigned int>> + std::optional::device::object_type, + unsigned int, unsigned int>> result = std::nullopt; /* @@ -216,9 +218,13 @@ TRACCC_HOST_DEVICE inline void find_tracks( } if (use_measurement) { - const auto& meas = measurements.at(meas_idx); - track_state trk_state(meas); + typename edm::track_state_collection< + typename detector_t::algebra_type>::device::object_type + trk_state(0u, 0.f, 0.f, 0.f, {}, {}, meas_idx); + trk_state.filtered_params().set_surface_link( + measurements.at(meas_idx).surface_link); + const detray::tracking_surface sf{det, in_par.surface_link()}; const bool is_line = sf.template visit_mask(); @@ -226,7 +232,7 @@ TRACCC_HOST_DEVICE inline void find_tracks( // Run the Kalman update const kalman_fitter_status res = gain_matrix_updater{}( - trk_state, in_par, is_line); + trk_state, measurements, in_par, is_line); /* * The $\chi^2$ value from the Kalman update should be less than @@ -246,8 +252,7 @@ TRACCC_HOST_DEVICE inline void find_tracks( } if (use_measurement) { - result.emplace(std::move(trk_state), meas_idx, - owner_local_thread_id); + result.emplace(trk_state, meas_idx, owner_local_thread_id); } } } @@ -441,10 +446,12 @@ TRACCC_HOST_DEVICE inline void find_tracks( .chi2_sum = prev_chi2_sum + chi2, .ndf_sum = prev_ndf_sum + - std::get<0>(*result).get_measurement().meas_dim}; + measurements + .at(std::get<0>(*result).measurement_index()) + .meas_dim}; tmp_params.at(p_offset + l_pos) = - std::get<0>(*result).filtered(); + std::get<0>(*result).filtered_params(); /* * Reset the temporary state storage, as this is no longer diff --git a/device/common/include/traccc/fitting/device/fit.hpp b/device/common/include/traccc/fitting/device/fit.hpp index e4aa539fde..8e923eecbf 100644 --- a/device/common/include/traccc/fitting/device/fit.hpp +++ b/device/common/include/traccc/fitting/device/fit.hpp @@ -12,7 +12,7 @@ // Project include(s). #include "traccc/definitions/qualifiers.hpp" -#include "traccc/edm/track_state.hpp" +#include "traccc/edm/track_fit_container.hpp" // VecMem include(s). #include @@ -44,9 +44,10 @@ struct fit_payload { vecmem::data::vector_view param_liveness_view; /** - * @brief View object to the output track states + * @brief View object to the output tracks */ - track_state_container_types::view track_states_view; + typename edm::track_fit_container< + typename fitter_t::detector_type::algebra_type>::view tracks_view; /** * @brief View object to the output barcode sequence diff --git a/device/common/include/traccc/fitting/device/fit_backward.hpp b/device/common/include/traccc/fitting/device/fit_backward.hpp index 9cf3cc9c1a..5787a7cca9 100644 --- a/device/common/include/traccc/fitting/device/fit_backward.hpp +++ b/device/common/include/traccc/fitting/device/fit_backward.hpp @@ -22,20 +22,28 @@ TRACCC_HOST_DEVICE inline void fit_backward( vecmem::device_vector param_ids(payload.param_ids_view); vecmem::device_vector param_liveness( payload.param_liveness_view); - track_state_container_types::device track_states(payload.track_states_view); - - if (globalIndex >= track_states.size()) { + typename edm::track_fit_collection< + typename fitter_t::detector_type::algebra_type>::device + tracks(payload.tracks_view.tracks); + typename edm::track_state_collection< + typename fitter_t::detector_type::algebra_type>::device + track_states(payload.tracks_view.states); + measurement_collection_types::const_device measurements{ + payload.tracks_view.measurements}; + + if (globalIndex >= tracks.size()) { return; } const unsigned int param_id = param_ids.at(globalIndex); + auto track = tracks.at(param_id); // Run fitting fitter_t fitter(det, payload.field_data, cfg); if (param_liveness.at(param_id) > 0u) { typename fitter_t::state fitter_state( - track_states.at(param_id).items, + track, track_states, measurements, *(payload.barcodes_view.ptr() + param_id)); kalman_fitter_status fit_status = fitter.smooth(fitter_state); @@ -48,7 +56,7 @@ TRACCC_HOST_DEVICE inline void fit_backward( assert(fit_status == kalman_fitter_status::SUCCESS); - track_states.at(param_id).header = fitter_state.m_fit_res; + track = fitter_state.m_fit_res; } else { param_liveness.at(param_id) = 0u; } diff --git a/device/common/include/traccc/fitting/device/fit_forward.hpp b/device/common/include/traccc/fitting/device/fit_forward.hpp index 2d6f3221a0..9957cd1d98 100644 --- a/device/common/include/traccc/fitting/device/fit_forward.hpp +++ b/device/common/include/traccc/fitting/device/fit_forward.hpp @@ -22,9 +22,16 @@ TRACCC_HOST_DEVICE inline void fit_forward( vecmem::device_vector param_ids(payload.param_ids_view); vecmem::device_vector param_liveness( payload.param_liveness_view); - track_state_container_types::device track_states(payload.track_states_view); - - if (globalIndex >= track_states.size()) { + typename edm::track_fit_collection< + typename fitter_t::detector_type::algebra_type>::device + tracks(payload.tracks_view.tracks); + typename edm::track_state_collection< + typename fitter_t::detector_type::algebra_type>::device + track_states(payload.tracks_view.states); + measurement_collection_types::const_device measurements{ + payload.tracks_view.measurements}; + + if (globalIndex >= tracks.size()) { return; } @@ -32,13 +39,14 @@ TRACCC_HOST_DEVICE inline void fit_forward( fitter_t fitter(det, payload.field_data, cfg); - auto params = track_states.at(param_id).header.fit_params; + auto track = tracks.at(param_id); + auto params = track.params(); // TODO: Merge into filter? inflate_covariance(params, fitter.config().covariance_inflation_factor); typename fitter_t::state fitter_state( - track_states.at(param_id).items, + track, track_states, measurements, *(payload.barcodes_view.ptr() + param_id)); kalman_fitter_status fit_status = fitter.filter(params, fitter_state); diff --git a/device/common/include/traccc/fitting/device/fit_prelude.hpp b/device/common/include/traccc/fitting/device/fit_prelude.hpp index 3436e4bd6c..d8d9ef7cc0 100644 --- a/device/common/include/traccc/fitting/device/fit_prelude.hpp +++ b/device/common/include/traccc/fitting/device/fit_prelude.hpp @@ -7,45 +7,66 @@ #pragma once +// Local include(s). +#include "traccc/device/global_index.hpp" + +// Project include(s). #include "traccc/edm/track_candidate_container.hpp" -#include "traccc/fitting/device/fit.hpp" +#include "traccc/edm/track_fit_container.hpp" #include "traccc/fitting/status_codes.hpp" +// VecMem include(s). +#include + namespace traccc::device { +template TRACCC_HOST_DEVICE inline void fit_prelude( const global_index_t globalIndex, vecmem::data::vector_view param_ids_view, - edm::track_candidate_container::const_view + typename edm::track_candidate_container::const_view track_candidates_view, - track_state_container_types::view track_states_view, + typename edm::track_fit_container::view tracks_view, vecmem::data::vector_view param_liveness_view) { - track_state_container_types::device track_states(track_states_view); + typename edm::track_fit_collection::device tracks( + tracks_view.tracks); - if (globalIndex >= track_states.size()) { + if (globalIndex >= tracks.size()) { return; } + 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_states_per_track = track_states.at(param_id).items; + auto track = tracks.at(param_id); - const edm::track_candidate_collection::const_device + 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_measurement_indices = + track_candidate.measurement_indices(); const measurement_collection_types::const_device measurements{ track_candidates_view.measurements}; - for (unsigned int meas_idx : - track_candidates.measurement_indices().at(param_id)) { - track_states_per_track.emplace_back(measurements.at(meas_idx)); + for (unsigned int meas_idx : track_candidate_measurement_indices) { + const unsigned int track_state_index = + track_states.push_back({0, 0.f, 0.f, 0.f, {}, {}, meas_idx}); + auto state = track_states.at(track_state_index); + state.set_hole(true); + state.set_smoothed(false); + const auto surface_link = measurements.at(meas_idx).surface_link; + state.filtered_params().set_surface_link(surface_link); + state.smoothed_params().set_surface_link(surface_link); + track.state_indices().push_back(track_state_index); } // TODO: Set other stuff in the header? - track_states.at(param_id).header.fit_params = - track_candidates.at(param_id).params(); + track.params() = track_candidate.params(); param_liveness.at(param_id) = 1u; } 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 f76ebd58c6..00ecc465aa 100644 --- a/device/cuda/include/traccc/cuda/fitting/kalman_fitting_algorithm.hpp +++ b/device/cuda/include/traccc/cuda/fitting/kalman_fitting_algorithm.hpp @@ -13,7 +13,7 @@ // Project include(s). #include "traccc/bfield/magnetic_field.hpp" #include "traccc/edm/track_candidate_container.hpp" -#include "traccc/edm/track_state.hpp" +#include "traccc/edm/track_fit_container.hpp" #include "traccc/fitting/fitting_config.hpp" #include "traccc/geometry/detector.hpp" #include "traccc/utils/algorithm.hpp" @@ -30,10 +30,10 @@ namespace traccc::cuda { /// Kalman filter based track fitting algorithm class kalman_fitting_algorithm - : public algorithm::buffer( const default_detector::view&, 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&)>, public messaging { @@ -42,7 +42,7 @@ class kalman_fitting_algorithm /// Configuration type using config_type = fitting_config; /// Output type - using output_type = track_state_container_types::buffer; + using output_type = edm::track_fit_container::buffer; /// Constructor with the algorithm's configuration /// diff --git a/device/cuda/src/fitting/kalman_fitting.cuh b/device/cuda/src/fitting/kalman_fitting.cuh index 0c16490503..f1ebd61d8e 100644 --- a/device/cuda/src/fitting/kalman_fitting.cuh +++ b/device/cuda/src/fitting/kalman_fitting.cuh @@ -19,7 +19,7 @@ // Project include(s). #include "traccc/edm/device/sort_key.hpp" #include "traccc/edm/track_candidate_container.hpp" -#include "traccc/edm/track_state.hpp" +#include "traccc/edm/track_fit_container.hpp" #include "traccc/fitting/details/kalman_fitting_types.hpp" #include "traccc/fitting/device/fill_fitting_sort_keys.hpp" #include "traccc/fitting/fitting_config.hpp" @@ -32,6 +32,9 @@ #include #include +// System include(s). +#include + namespace traccc::cuda::details { /// Templated implementation of the CUDA track fitting algorithm. @@ -50,7 +53,8 @@ namespace traccc::cuda::details { /// @return A container of the fitted track states /// template -track_state_container_types::buffer kalman_fitting( +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< @@ -69,14 +73,17 @@ track_state_container_types::buffer kalman_fitting( // 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. - track_state_container_types::buffer track_states_buffer{ - {n_tracks, mr.main}, - {candidate_sizes, mr.main, mr.host, - vecmem::data::buffer_type::resizable}}; - copy.setup(track_states_buffer.headers)->ignore(); - copy.setup(track_states_buffer.items)->ignore(); + 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) { @@ -128,7 +135,9 @@ track_state_container_types::buffer kalman_fitting( // Run the fitting, using the sorted parameter IDs. fit_prelude(nBlocks, nThreads, 0, stream, param_ids_buffer, - track_candidates_view, track_states_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(); @@ -140,7 +149,8 @@ track_state_container_types::buffer kalman_fitting( .field_data = field_view, .param_ids_view = param_ids_buffer, .param_liveness_view = param_liveness_buffer, - .track_states_view = track_states_buffer, + .tracks_view = {track_states_buffer.tracks, track_states_buffer.states, + track_candidates_view.measurements}, .barcodes_view = seqs_buffer}; for (std::size_t i = 0; i < config.n_iterations; ++i) { diff --git a/device/cuda/src/fitting/kernels/fit_prelude.cu b/device/cuda/src/fitting/kernels/fit_prelude.cu index 319f2a29dc..ef8e63bcd9 100644 --- a/device/cuda/src/fitting/kernels/fit_prelude.cu +++ b/device/cuda/src/fitting/kernels/fit_prelude.cu @@ -15,11 +15,11 @@ __global__ void fit_prelude( vecmem::data::vector_view param_ids_view, edm::track_candidate_container::const_view track_candidates_view, - track_state_container_types::view track_states_view, + edm::track_fit_container::view tracks_view, vecmem::data::vector_view param_liveness_view) { - device::fit_prelude(details::global_index1(), param_ids_view, - track_candidates_view, track_states_view, - param_liveness_view); + device::fit_prelude(details::global_index1(), + param_ids_view, track_candidates_view, + tracks_view, param_liveness_view); } } // namespace kernels @@ -28,10 +28,10 @@ void fit_prelude(const dim3& grid_size, const dim3& block_size, vecmem::data::vector_view param_ids_view, edm::track_candidate_container::const_view track_candidates_view, - track_state_container_types::view track_states_view, + edm::track_fit_container::view tracks_view, vecmem::data::vector_view param_liveness_view) { kernels::fit_prelude<<>>( - param_ids_view, track_candidates_view, track_states_view, + param_ids_view, track_candidates_view, tracks_view, param_liveness_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 76d91be358..6f4098032c 100644 --- a/device/cuda/src/fitting/kernels/fit_prelude.hpp +++ b/device/cuda/src/fitting/kernels/fit_prelude.hpp @@ -10,7 +10,7 @@ #include #include "traccc/edm/track_candidate_container.hpp" -#include "traccc/edm/track_state.hpp" +#include "traccc/edm/track_fit_container.hpp" namespace traccc::cuda { void fit_prelude(const dim3& grid_size, const dim3& block_size, @@ -18,6 +18,6 @@ void fit_prelude(const dim3& grid_size, const dim3& block_size, vecmem::data::vector_view param_ids_view, edm::track_candidate_container::const_view track_candidates_view, - track_state_container_types::view track_states_view, + edm::track_fit_container::view tracks_view, vecmem::data::vector_view param_liveness_view); } From 4d01d5c08755707d205b8c4ee38c1d18cdbcfb81 Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Mon, 4 Aug 2025 13:29:01 +0200 Subject: [PATCH 03/11] Adapted traccc::alpaka and traccc::sycl to the SoA track fit EDM. --- .../fitting/kalman_fitting_algorithm.hpp | 8 +-- device/alpaka/src/fitting/kalman_fitting.hpp | 47 ++++++++++-------- .../sycl/fitting/kalman_fitting_algorithm.hpp | 8 +-- device/sycl/src/fitting/kalman_fitting.hpp | 49 ++++++++++++------- 4 files changed, 65 insertions(+), 47 deletions(-) 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 e0009a0090..6fa432b448 100644 --- a/device/alpaka/include/traccc/alpaka/fitting/kalman_fitting_algorithm.hpp +++ b/device/alpaka/include/traccc/alpaka/fitting/kalman_fitting_algorithm.hpp @@ -13,7 +13,7 @@ // Project include(s). #include "traccc/bfield/magnetic_field.hpp" #include "traccc/edm/track_candidate_container.hpp" -#include "traccc/edm/track_state.hpp" +#include "traccc/edm/track_fit_container.hpp" #include "traccc/fitting/fitting_config.hpp" #include "traccc/geometry/detector.hpp" #include "traccc/utils/algorithm.hpp" @@ -30,10 +30,10 @@ namespace traccc::alpaka { /// Kalman filter based track fitting algorithm class kalman_fitting_algorithm - : public algorithm::buffer( const default_detector::view&, 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&)>, public messaging { @@ -42,7 +42,7 @@ class kalman_fitting_algorithm /// Configuration type using config_type = fitting_config; /// Output type - using output_type = track_state_container_types::buffer; + using output_type = edm::track_fit_container::buffer; /// Constructor with the algorithm's configuration /// diff --git a/device/alpaka/src/fitting/kalman_fitting.hpp b/device/alpaka/src/fitting/kalman_fitting.hpp index 472f544fb7..be5ad7a0df 100644 --- a/device/alpaka/src/fitting/kalman_fitting.hpp +++ b/device/alpaka/src/fitting/kalman_fitting.hpp @@ -14,7 +14,7 @@ // Project include(s). #include "traccc/edm/device/sort_key.hpp" #include "traccc/edm/track_candidate_container.hpp" -#include "traccc/edm/track_state.hpp" +#include "traccc/edm/track_fit_container.hpp" #include "traccc/fitting/details/kalman_fitting_types.hpp" #include "traccc/fitting/device/fill_fitting_sort_keys.hpp" #include "traccc/fitting/device/fit.hpp" @@ -55,14 +55,14 @@ struct fit_prelude { vecmem::data::vector_view param_ids_view, edm::track_candidate_container::const_view track_candidates_view, - track_state_container_types::view track_states_view, + edm::track_fit_container::view track_states_view, vecmem::data::vector_view param_liveness_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); + device::fit_prelude( + globalThreadIdx, param_ids_view, track_candidates_view, + track_states_view, param_liveness_view); } }; @@ -112,7 +112,8 @@ struct fit_backward { /// @return A container of the fitted track states /// template -track_state_container_types::buffer kalman_fitting( +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< @@ -131,21 +132,24 @@ track_state_container_types::buffer kalman_fitting( // 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. - track_state_container_types::buffer track_states_buffer{ - {n_tracks, mr.main}, - {candidate_sizes, mr.main, mr.host, - vecmem::data::buffer_type::resizable}}; - vecmem::copy::event_type track_states_headers_setup_event = - copy.setup(track_states_buffer.headers); - vecmem::copy::event_type track_states_items_setup_event = - copy.setup(track_states_buffer.items); + 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) { - track_states_headers_setup_event->wait(); - track_states_items_setup_event->wait(); + tracks_setup_event->wait(); + track_states_setup_event->wait(); return track_states_buffer; } @@ -194,9 +198,12 @@ track_state_container_types::buffer kalman_fitting( param_ids_device.begin()); // Run the fitting, using the sorted parameter IDs. - track_state_container_types::view track_states_view = track_states_buffer; - track_states_headers_setup_event->wait(); - track_states_items_setup_event->wait(); + 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), @@ -211,7 +218,7 @@ track_state_container_types::buffer kalman_fitting( .field_data = field_view, .param_ids_view = param_ids_buffer, .param_liveness_view = param_liveness_buffer, - .track_states_view = track_states_view, + .tracks_view = track_states_view, .barcodes_view = seqs_buffer}; // Now copy it to device memory. vecmem::data::vector_buffer> device_payload( 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 94916e51bd..2612a9a274 100644 --- a/device/sycl/include/traccc/sycl/fitting/kalman_fitting_algorithm.hpp +++ b/device/sycl/include/traccc/sycl/fitting/kalman_fitting_algorithm.hpp @@ -13,7 +13,7 @@ // Project include(s). #include "traccc/bfield/magnetic_field.hpp" #include "traccc/edm/track_candidate_container.hpp" -#include "traccc/edm/track_state.hpp" +#include "traccc/edm/track_fit_container.hpp" #include "traccc/fitting/fitting_config.hpp" #include "traccc/geometry/detector.hpp" #include "traccc/utils/algorithm.hpp" @@ -30,10 +30,10 @@ namespace traccc::sycl { /// Kalman filter based track fitting algorithm class kalman_fitting_algorithm - : public algorithm::buffer( const default_detector::view&, 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&)>, public messaging { @@ -42,7 +42,7 @@ class kalman_fitting_algorithm /// Configuration type using config_type = fitting_config; /// Output type - using output_type = track_state_container_types::buffer; + using output_type = edm::track_fit_container::buffer; /// Constructor with the algorithm's configuration /// diff --git a/device/sycl/src/fitting/kalman_fitting.hpp b/device/sycl/src/fitting/kalman_fitting.hpp index e9c242a477..22c4a011b0 100644 --- a/device/sycl/src/fitting/kalman_fitting.hpp +++ b/device/sycl/src/fitting/kalman_fitting.hpp @@ -15,7 +15,7 @@ // Project include(s). #include "traccc/edm/device/sort_key.hpp" #include "traccc/edm/track_candidate_container.hpp" -#include "traccc/edm/track_state.hpp" +#include "traccc/edm/track_fit_container.hpp" #include "traccc/fitting/details/kalman_fitting_types.hpp" #include "traccc/fitting/device/fill_fitting_sort_keys.hpp" #include "traccc/fitting/device/fit.hpp" @@ -31,6 +31,9 @@ // SYCL include(s). #include +// System include(s). +#include + namespace traccc::sycl { namespace kernels { @@ -64,7 +67,8 @@ namespace details { /// @return A container of the fitted track states /// template -track_state_container_types::buffer kalman_fitting( +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< @@ -80,21 +84,24 @@ track_state_container_types::buffer kalman_fitting( // 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. - track_state_container_types::buffer track_states_buffer{ - {n_tracks, mr.main}, - {candidate_sizes, mr.main, mr.host, - vecmem::data::buffer_type::resizable}}; - vecmem::copy::event_type track_states_headers_setup_event = - copy.setup(track_states_buffer.headers); - vecmem::copy::event_type track_states_items_setup_event = - copy.setup(track_states_buffer.items); + 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) { - track_states_headers_setup_event->wait(); - track_states_items_setup_event->wait(); + tracks_setup_event->wait(); + track_states_setup_event->wait(); return track_states_buffer; } @@ -150,9 +157,12 @@ track_state_container_types::buffer kalman_fitting( param_ids_device.begin()); // Run the fitting, using the sorted parameter IDs. - track_state_container_types::view track_states_view = track_states_buffer; - track_states_headers_setup_event->wait(); - track_states_items_setup_event->wait(); + 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) { @@ -161,9 +171,10 @@ track_state_container_types::buffer kalman_fitting( 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); + device::fit_prelude( + details::global_index(item), param_ids_view, + track_candidates_view, track_states_view, + param_liveness_view); }); }) .wait_and_throw(); @@ -174,7 +185,7 @@ track_state_container_types::buffer kalman_fitting( .field_data = field_view, .param_ids_view = param_ids_buffer, .param_liveness_view = param_liveness_buffer, - .track_states_view = track_states_view, + .tracks_view = track_states_view, .barcodes_view = seqs_buffer}; for (std::size_t i = 0; i < config.n_iterations; ++i) { From 4dbf2252acdfe3b35e5b558786708dd6c8253b2e Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Mon, 4 Aug 2025 18:01:14 +0200 Subject: [PATCH 04/11] Adapted the tests and benchmarks to the SoA track fit EDM. --- benchmarks/cuda/toy_detector_cuda.cpp | 6 +-- ...alman_fitting_momentum_resolution_test.cpp | 6 ++- ...alman_fitting_momentum_resolution_test.hpp | 5 ++- .../tests/kalman_fitting_telescope_test.hpp | 8 ++-- tests/common/tests/kalman_fitting_test.cpp | 34 ++++++++++++---- tests/common/tests/kalman_fitting_test.hpp | 21 ++++++++-- .../kalman_fitting_wire_chamber_test.hpp | 8 ++-- .../cpu/test_ckf_sparse_tracks_telescope.cpp | 20 +++++----- tests/cpu/test_kalman_fitter_hole_count.cpp | 11 +++-- ...test_kalman_fitter_momentum_resolution.cpp | 21 +++++----- tests/cpu/test_kalman_fitter_telescope.cpp | 21 +++++----- tests/cpu/test_kalman_fitter_wire_chamber.cpp | 19 +++++---- .../cuda/test_ckf_combinatorics_telescope.cpp | 5 --- tests/cuda/test_ckf_toy_detector.cpp | 5 --- tests/cuda/test_kalman_fitter_telescope.cpp | 40 ++++++++++--------- .../sycl/test_ckf_combinatorics_telescope.cpp | 5 --- tests/sycl/test_ckf_toy_detector.cpp | 5 --- tests/sycl/test_kalman_fitter_telescope.cpp | 40 ++++++++++--------- 18 files changed, 150 insertions(+), 130 deletions(-) diff --git a/benchmarks/cuda/toy_detector_cuda.cpp b/benchmarks/cuda/toy_detector_cuda.cpp index 87c8523662..b2f14c542b 100644 --- a/benchmarks/cuda/toy_detector_cuda.cpp +++ b/benchmarks/cuda/toy_detector_cuda.cpp @@ -73,10 +73,6 @@ BENCHMARK_DEFINE_F(ToyDetectorBenchmark, CUDA)(benchmark::State& state) { // Detector view object auto det_view = detray::get_data(det_buffer); - // D2H copy object - traccc::device::container_d2h_copy_alg - track_state_d2h{{device_mr, &host_mr}, copy}; - for (auto _ : state) { state.PauseTiming(); @@ -129,7 +125,7 @@ BENCHMARK_DEFINE_F(ToyDetectorBenchmark, CUDA)(benchmark::State& state) { params_cuda_buffer); // Run track fitting - traccc::track_state_container_types::buffer + traccc::edm::track_fit_container::buffer track_states_cuda_buffer = device_fitting( det_view, field, {track_candidates_cuda_buffer, measurements_cuda_buffer}); diff --git a/tests/common/tests/kalman_fitting_momentum_resolution_test.cpp b/tests/common/tests/kalman_fitting_momentum_resolution_test.cpp index c757617caf..c74c776a6d 100644 --- a/tests/common/tests/kalman_fitting_momentum_resolution_test.cpp +++ b/tests/common/tests/kalman_fitting_momentum_resolution_test.cpp @@ -27,11 +27,13 @@ namespace traccc { void KalmanFittingMomentumResolutionTests::consistency_tests( - const track_state_collection_types::host& track_states_per_track) const { + const edm::track_fit_collection::host::const_proxy_type& + track, + const edm::track_state_collection::host&) const { // The nubmer of track states is supposed be equal to the number // of planes - ASSERT_EQ(track_states_per_track.size(), std::get<11>(GetParam())); + ASSERT_EQ(track.state_indices().size(), std::get<11>(GetParam())); } void KalmanFittingMomentumResolutionTests::momentum_resolution_tests( diff --git a/tests/common/tests/kalman_fitting_momentum_resolution_test.hpp b/tests/common/tests/kalman_fitting_momentum_resolution_test.hpp index e30126d0ee..7c7ea9450f 100644 --- a/tests/common/tests/kalman_fitting_momentum_resolution_test.hpp +++ b/tests/common/tests/kalman_fitting_momentum_resolution_test.hpp @@ -57,7 +57,10 @@ class KalmanFittingMomentumResolutionTests 1.f * traccc::unit::ns}; void consistency_tests( - const track_state_collection_types::host& track_states_per_track) const; + const edm::track_fit_collection< + default_algebra>::host::const_proxy_type& track, + const edm::track_state_collection::host& track_states) + const; void momentum_resolution_tests(std::string_view file_name) const; diff --git a/tests/common/tests/kalman_fitting_telescope_test.hpp b/tests/common/tests/kalman_fitting_telescope_test.hpp index 4d977f3545..2fcffaf74e 100644 --- a/tests/common/tests/kalman_fitting_telescope_test.hpp +++ b/tests/common/tests/kalman_fitting_telescope_test.hpp @@ -77,12 +77,14 @@ class KalmanFittingTelescopeTests std::get<11>(GetParam())); } - void consistency_tests(const track_state_collection_types::host& - track_states_per_track) const { + void consistency_tests( + const edm::track_fit_collection< + default_algebra>::host::const_proxy_type& track, + const edm::track_state_collection::host&) const { // The nubmer of track states is supposed be equal to the number // of planes - ASSERT_EQ(track_states_per_track.size(), std::get<11>(GetParam())); + ASSERT_EQ(track.state_indices().size(), std::get<11>(GetParam())); } protected: diff --git a/tests/common/tests/kalman_fitting_test.cpp b/tests/common/tests/kalman_fitting_test.cpp index e7f1d1e65d..fe6cde04de 100644 --- a/tests/common/tests/kalman_fitting_test.cpp +++ b/tests/common/tests/kalman_fitting_test.cpp @@ -151,31 +151,51 @@ void KalmanFittingTests::ndf_tests( } void KalmanFittingTests::ndf_tests( - const fitting_result& fit_res, - const track_state_collection_types::host& track_states_per_track) { + const edm::track_fit_collection::host::const_proxy_type& + track, + const edm::track_state_collection::host& track_states, + const measurement_collection_types::host& measurements) { scalar dim_sum = 0; std::size_t n_effective_states = 0; - for (const auto& state : track_states_per_track) { + for (unsigned int state_idx : track.state_indices()) { + + auto state = track_states.at(state_idx); - if (!state.is_hole && state.is_smoothed) { + if (!state.is_hole() && state.is_smoothed()) { - dim_sum += static_cast(state.get_measurement().meas_dim); + dim_sum += static_cast( + measurements.at(state.measurement_index()).meas_dim); n_effective_states++; } } // Check if the number of degree of freedoms is equal to (the sum of // measurement dimensions - 5) - ASSERT_FLOAT_EQ(static_cast(fit_res.trk_quality.ndf), + ASSERT_FLOAT_EQ(static_cast(track.ndf()), static_cast(dim_sum) - 5.f); // The number of track states is supposed to be eqaul to the number // of measurements unless KF failes in the middle of propagation - if (n_effective_states == track_states_per_track.size()) { + if (n_effective_states == track.state_indices().size()) { n_success++; } } +std::size_t KalmanFittingTests::count_successfully_fitted_tracks( + const edm::track_fit_collection::host& tracks) const { + + const std::size_t n_tracks = tracks.size(); + std::size_t n_fitted_tracks = 0u; + + for (std::size_t i = 0; i < n_tracks; ++i) { + if (tracks.at(i).fit_outcome() == track_fit_outcome::SUCCESS) { + n_fitted_tracks++; + } + } + + return n_fitted_tracks; +} + } // namespace traccc diff --git a/tests/common/tests/kalman_fitting_test.hpp b/tests/common/tests/kalman_fitting_test.hpp index d615eed1b1..9473b54655 100644 --- a/tests/common/tests/kalman_fitting_test.hpp +++ b/tests/common/tests/kalman_fitting_test.hpp @@ -12,6 +12,8 @@ #include "traccc/definitions/common.hpp" #include "traccc/edm/measurement.hpp" #include "traccc/edm/track_candidate_collection.hpp" +#include "traccc/edm/track_fit_collection.hpp" +#include "traccc/edm/track_state_collection.hpp" #include "traccc/fitting/kalman_filter/kalman_fitter.hpp" #include "traccc/geometry/detector.hpp" #include "traccc/simulation/event_generators.hpp" @@ -79,12 +81,23 @@ class KalmanFittingTests : public testing::Test { /// Validadte the NDF for track fitting output /// - /// @param fit_res Fitting statistics result of a track - /// @param track_states_per_track Track states of a track + /// @param track Fitting statistics result of a track + /// @param track_states All track states in the event + /// @param measurements All measurements in the event /// void ndf_tests( - const fitting_result& fit_res, - const track_state_collection_types::host& track_states_per_track); + const edm::track_fit_collection< + default_algebra>::host::const_proxy_type& track, + const edm::track_state_collection::host& track_states, + const measurement_collection_types::host& measurements); + + /// Count the number of tracks that were successfully fitted + /// + /// @param tracks The track fit collection to count on + /// @return The number of tracks that were successfully fitted + /// + std::size_t count_successfully_fitted_tracks( + const edm::track_fit_collection::host& tracks) const; // The number of tracks successful with KF std::size_t n_success{0u}; diff --git a/tests/common/tests/kalman_fitting_wire_chamber_test.hpp b/tests/common/tests/kalman_fitting_wire_chamber_test.hpp index 457ba3b1e0..c6634df2d6 100644 --- a/tests/common/tests/kalman_fitting_wire_chamber_test.hpp +++ b/tests/common/tests/kalman_fitting_wire_chamber_test.hpp @@ -70,12 +70,14 @@ class KalmanFittingWireChamberTests 0.05f / traccc::unit::GeV, 1.f * traccc::unit::ns}; - void consistency_tests(const track_state_collection_types::host& - track_states_per_track) const { + void consistency_tests( + const edm::track_fit_collection< + default_algebra>::host::const_proxy_type& track, + const edm::track_state_collection::host&) const { // The nubmer of track states is supposed be greater than or // equal to the number of layers - ASSERT_GE(track_states_per_track.size(), n_wire_layers); + ASSERT_GE(track.state_indices().size(), n_wire_layers); } protected: diff --git a/tests/cpu/test_ckf_sparse_tracks_telescope.cpp b/tests/cpu/test_ckf_sparse_tracks_telescope.cpp index 9800cfda14..25c9134747 100644 --- a/tests/cpu/test_ckf_sparse_tracks_telescope.cpp +++ b/tests/cpu/test_ckf_sparse_tracks_telescope.cpp @@ -181,22 +181,22 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { {vecmem::get_data(track_candidates), vecmem::get_data(measurements_per_event)}); const std::size_t n_fitted_tracks = - count_successfully_fitted_tracks(track_states); + count_successfully_fitted_tracks(track_states.tracks); - ASSERT_EQ(track_states.size(), n_truth_tracks); - ASSERT_EQ(track_states.size(), n_fitted_tracks); + ASSERT_EQ(track_states.tracks.size(), n_truth_tracks); + ASSERT_EQ(track_states.tracks.size(), n_fitted_tracks); for (unsigned int i_trk = 0; i_trk < n_truth_tracks; i_trk++) { - const auto& track_states_per_track = track_states[i_trk].items; - const auto& fit_res = track_states[i_trk].header; + consistency_tests(track_states.tracks.at(i_trk), + track_states.states); - consistency_tests(track_states_per_track); + ndf_tests(track_states.tracks.at(i_trk), track_states.states, + measurements_per_event); - ndf_tests(fit_res, track_states_per_track); - - fit_performance_writer.write(track_states_per_track, fit_res, - host_det, evt_data); + fit_performance_writer.write( + track_states.tracks.at(i_trk), track_states.states, + measurements_per_event, host_det, evt_data); } } diff --git a/tests/cpu/test_kalman_fitter_hole_count.cpp b/tests/cpu/test_kalman_fitter_hole_count.cpp index 0181e57794..b88f8dafa9 100644 --- a/tests/cpu/test_kalman_fitter_hole_count.cpp +++ b/tests/cpu/test_kalman_fitter_hole_count.cpp @@ -7,7 +7,6 @@ // Project include(s). #include "traccc/bfield/construct_const_bfield.hpp" -#include "traccc/edm/track_state.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" #include "traccc/io/utils.hpp" #include "traccc/resolution/fitting_performance_writer.hpp" @@ -161,19 +160,19 @@ TEST_P(KalmanFittingHoleCountTests, Run) { vecmem::get_data(track_candidates.measurements)}); // A sanity check - const std::size_t n_tracks = track_states.size(); + const std::size_t n_tracks = track_states.tracks.size(); ASSERT_EQ(n_tracks, n_truth_tracks); // Check the number of holes // The three holes at the end are not counted as KF aborts once it goes // through all track candidates - const auto& fit_res = track_states.at(0u).header; - ASSERT_EQ(fit_res.trk_quality.n_holes, 5u); + const auto track = track_states.tracks.at(0u); + ASSERT_EQ(track.nholes(), 5u); // Some sanity checks ASSERT_FLOAT_EQ( - static_cast(fit_res.trk_quality.ndf), - static_cast(track_states.at(0u).items.size()) * 2.f - 5.f); + static_cast(track.ndf()), + static_cast(track.state_indices().size()) * 2.f - 5.f); } INSTANTIATE_TEST_SUITE_P( diff --git a/tests/cpu/test_kalman_fitter_momentum_resolution.cpp b/tests/cpu/test_kalman_fitter_momentum_resolution.cpp index 725d5ce1e0..1a18b446af 100644 --- a/tests/cpu/test_kalman_fitter_momentum_resolution.cpp +++ b/tests/cpu/test_kalman_fitter_momentum_resolution.cpp @@ -7,7 +7,6 @@ // Project include(s). #include "traccc/bfield/construct_const_bfield.hpp" -#include "traccc/edm/track_state.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" #include "traccc/io/utils.hpp" #include "traccc/resolution/fitting_performance_writer.hpp" @@ -173,9 +172,9 @@ TEST_P(KalmanFittingMomentumResolutionTests, Run) { vecmem::get_data(track_candidates.measurements)}); // Iterator over tracks - const std::size_t n_tracks = track_states.size(); + const std::size_t n_tracks = track_states.tracks.size(); const std::size_t n_fitted_tracks = - count_successfully_fitted_tracks(track_states); + count_successfully_fitted_tracks(track_states.tracks); // n_trakcs = 100 ASSERT_GE(static_cast(n_tracks), @@ -185,17 +184,17 @@ TEST_P(KalmanFittingMomentumResolutionTests, Run) { for (std::size_t i_trk = 0; i_trk < n_tracks; i_trk++) { - const auto& track_states_per_track = track_states[i_trk].items; - const auto& fit_res = track_states[i_trk].header; + consistency_tests(track_states.tracks.at(i_trk), + track_states.states); - consistency_tests(track_states_per_track); + ndf_tests(track_states.tracks.at(i_trk), track_states.states, + track_candidates.measurements); - ndf_tests(fit_res, track_states_per_track); + ASSERT_EQ(track_states.tracks.at(i_trk).nholes(), 0u); - ASSERT_EQ(fit_res.trk_quality.n_holes, 0u); - - fit_performance_writer.write(track_states_per_track, fit_res, - host_det, evt_data); + fit_performance_writer.write( + track_states.tracks.at(i_trk), track_states.states, + track_candidates.measurements, host_det, evt_data); } } diff --git a/tests/cpu/test_kalman_fitter_telescope.cpp b/tests/cpu/test_kalman_fitter_telescope.cpp index aae6ac3072..c3a8fdb5c0 100644 --- a/tests/cpu/test_kalman_fitter_telescope.cpp +++ b/tests/cpu/test_kalman_fitter_telescope.cpp @@ -7,7 +7,6 @@ // Project include(s). #include "traccc/bfield/construct_const_bfield.hpp" -#include "traccc/edm/track_state.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" #include "traccc/io/utils.hpp" #include "traccc/resolution/fitting_performance_writer.hpp" @@ -144,9 +143,9 @@ TEST_P(KalmanFittingTelescopeTests, Run) { vecmem::get_data(track_candidates.measurements)}); // Iterator over tracks - const std::size_t n_tracks = track_states.size(); + const std::size_t n_tracks = track_states.tracks.size(); const std::size_t n_fitted_tracks = - count_successfully_fitted_tracks(track_states); + count_successfully_fitted_tracks(track_states.tracks); // n_trakcs = 100 ASSERT_EQ(n_tracks, n_truth_tracks); @@ -154,17 +153,17 @@ TEST_P(KalmanFittingTelescopeTests, Run) { for (std::size_t i_trk = 0; i_trk < n_tracks; i_trk++) { - const auto& track_states_per_track = track_states[i_trk].items; - const auto& fit_res = track_states[i_trk].header; + consistency_tests(track_states.tracks.at(i_trk), + track_states.states); - consistency_tests(track_states_per_track); + ndf_tests(track_states.tracks.at(i_trk), track_states.states, + track_candidates.measurements); - ndf_tests(fit_res, track_states_per_track); + ASSERT_EQ(track_states.tracks.at(i_trk).nholes(), 0u); - ASSERT_EQ(fit_res.trk_quality.n_holes, 0u); - - fit_performance_writer.write(track_states_per_track, fit_res, - host_det, evt_data); + fit_performance_writer.write( + track_states.tracks.at(i_trk), track_states.states, + track_candidates.measurements, host_det, evt_data); } } diff --git a/tests/cpu/test_kalman_fitter_wire_chamber.cpp b/tests/cpu/test_kalman_fitter_wire_chamber.cpp index 8386b284a3..612501dde2 100644 --- a/tests/cpu/test_kalman_fitter_wire_chamber.cpp +++ b/tests/cpu/test_kalman_fitter_wire_chamber.cpp @@ -7,7 +7,6 @@ // Project include(s). #include "traccc/bfield/construct_const_bfield.hpp" -#include "traccc/edm/track_state.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" #include "traccc/io/utils.hpp" #include "traccc/resolution/fitting_performance_writer.hpp" @@ -149,27 +148,27 @@ TEST_P(KalmanFittingWireChamberTests, Run) { vecmem::get_data(track_candidates.measurements)}); // Iterator over tracks - const std::size_t n_tracks = track_states.size(); + const std::size_t n_tracks = track_states.tracks.size(); ASSERT_GE(static_cast(n_tracks), 0.98 * static_cast(n_truth_tracks)); const std::size_t n_fitted_tracks = - count_successfully_fitted_tracks(track_states); + count_successfully_fitted_tracks(track_states.tracks); ASSERT_GE(static_cast(n_fitted_tracks), 0.95 * static_cast(n_truth_tracks)); for (std::size_t i_trk = 0; i_trk < n_tracks; i_trk++) { - const auto& fit_res = track_states[i_trk].header; - const auto& track_states_per_track = track_states[i_trk].items; + consistency_tests(track_states.tracks.at(i_trk), + track_states.states); - consistency_tests(track_states_per_track); + ndf_tests(track_states.tracks.at(i_trk), track_states.states, + track_candidates.measurements); - ndf_tests(fit_res, track_states_per_track); - - fit_performance_writer.write(track_states_per_track, fit_res, - host_det, evt_data); + fit_performance_writer.write( + track_states.tracks.at(i_trk), track_states.states, + track_candidates.measurements, host_det, evt_data); } } diff --git a/tests/cuda/test_ckf_combinatorics_telescope.cpp b/tests/cuda/test_ckf_combinatorics_telescope.cpp index e2dad94c43..2e0d0be4ab 100644 --- a/tests/cuda/test_ckf_combinatorics_telescope.cpp +++ b/tests/cuda/test_ckf_combinatorics_telescope.cpp @@ -9,8 +9,6 @@ #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/device/container_d2h_copy_alg.hpp" -#include "traccc/device/container_h2d_copy_alg.hpp" #include "traccc/io/read_measurements.hpp" #include "traccc/io/utils.hpp" #include "traccc/simulation/event_generators.hpp" @@ -126,9 +124,6 @@ TEST_P(CudaCkfCombinatoricsTelescopeTests, Run) { // Copy objects vecmem::cuda::async_copy copy{stream.cudaStream()}; - traccc::device::container_d2h_copy_alg - track_state_d2h{mr, copy}; - // Seed generator seed_generator sg(host_det, stddevs); diff --git a/tests/cuda/test_ckf_toy_detector.cpp b/tests/cuda/test_ckf_toy_detector.cpp index 5ae7a1023e..ee9c92b6b5 100644 --- a/tests/cuda/test_ckf_toy_detector.cpp +++ b/tests/cuda/test_ckf_toy_detector.cpp @@ -9,8 +9,6 @@ #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/device/container_d2h_copy_alg.hpp" -#include "traccc/device/container_h2d_copy_alg.hpp" #include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" #include "traccc/io/read_measurements.hpp" #include "traccc/io/utils.hpp" @@ -125,9 +123,6 @@ TEST_P(CkfToyDetectorTests, Run) { // Copy objects vecmem::cuda::async_copy copy{stream.cudaStream()}; - traccc::device::container_d2h_copy_alg - track_state_d2h{mr, copy}; - // Seed generator seed_generator sg(host_det, stddevs); diff --git a/tests/cuda/test_kalman_fitter_telescope.cpp b/tests/cuda/test_kalman_fitter_telescope.cpp index d5e720ee5b..062030aa81 100644 --- a/tests/cuda/test_kalman_fitter_telescope.cpp +++ b/tests/cuda/test_kalman_fitter_telescope.cpp @@ -9,9 +9,7 @@ #include "traccc/bfield/construct_const_bfield.hpp" #include "traccc/bfield/magnetic_field_types.hpp" #include "traccc/cuda/fitting/kalman_fitting_algorithm.hpp" -#include "traccc/device/container_d2h_copy_alg.hpp" -#include "traccc/device/container_h2d_copy_alg.hpp" -#include "traccc/edm/track_state.hpp" +#include "traccc/edm/track_fit_container.hpp" #include "traccc/io/utils.hpp" #include "traccc/performance/details/is_same_object.hpp" #include "traccc/resolution/fitting_performance_writer.hpp" @@ -135,9 +133,6 @@ TEST_P(KalmanFittingTelescopeTests, Run) { // Copy objects vecmem::cuda::async_copy copy{stream.cudaStream()}; - traccc::device::container_d2h_copy_alg - track_state_d2h{mr, copy}; - // Seed generator seed_generator sg(host_det, stddevs); @@ -169,30 +164,37 @@ TEST_P(KalmanFittingTelescopeTests, Run) { mr.main, vecmem::copy::type::host_to_device)}; // Run fitting - traccc::track_state_container_types::buffer track_states_cuda_buffer = + auto track_states_cuda_buffer = device_fitting(det_view, field, {track_candidates_buffer.tracks, track_candidates_buffer.measurements}); - traccc::track_state_container_types::host track_states_cuda = - track_state_d2h(track_states_cuda_buffer); + traccc::edm::track_fit_container::host + track_states_cuda{host_mr}; + copy(track_states_cuda_buffer.tracks, track_states_cuda.tracks, + vecmem::copy::type::device_to_host) + ->wait(); + copy(track_states_cuda_buffer.states, track_states_cuda.states, + vecmem::copy::type::device_to_host) + ->wait(); + const std::size_t n_fitted_tracks = - count_successfully_fitted_tracks(track_states_cuda); + count_successfully_fitted_tracks(track_states_cuda.tracks); - ASSERT_EQ(track_states_cuda.size(), n_truth_tracks); - ASSERT_EQ(track_states_cuda.size(), n_fitted_tracks); + ASSERT_EQ(track_states_cuda.tracks.size(), n_truth_tracks); + ASSERT_EQ(track_states_cuda.tracks.size(), n_fitted_tracks); for (std::size_t i_trk = 0; i_trk < n_truth_tracks; i_trk++) { - const auto& track_states_per_track = track_states_cuda[i_trk].items; - const auto& fit_res = track_states_cuda[i_trk].header; - - consistency_tests(track_states_per_track); + consistency_tests(track_states_cuda.tracks.at(i_trk), + track_states_cuda.states); - ndf_tests(fit_res, track_states_per_track); + ndf_tests(track_states_cuda.tracks.at(i_trk), + track_states_cuda.states, track_candidates.measurements); - fit_performance_writer.write(track_states_per_track, fit_res, - host_det, evt_data); + fit_performance_writer.write( + track_states_cuda.tracks.at(i_trk), track_states_cuda.states, + track_candidates.measurements, host_det, evt_data); } } diff --git a/tests/sycl/test_ckf_combinatorics_telescope.cpp b/tests/sycl/test_ckf_combinatorics_telescope.cpp index 675a256c85..38ebddc996 100644 --- a/tests/sycl/test_ckf_combinatorics_telescope.cpp +++ b/tests/sycl/test_ckf_combinatorics_telescope.cpp @@ -12,8 +12,6 @@ // Project include(s). #include "traccc/bfield/construct_const_bfield.hpp" #include "traccc/bfield/magnetic_field_types.hpp" -#include "traccc/device/container_d2h_copy_alg.hpp" -#include "traccc/device/container_h2d_copy_alg.hpp" #include "traccc/io/read_detector.hpp" #include "traccc/io/read_measurements.hpp" #include "traccc/io/utils.hpp" @@ -130,9 +128,6 @@ TEST_P(CkfCombinatoricsTelescopeTests, Run) { // Copy objects vecmem::sycl::async_copy copy{vecmem_queue}; - traccc::device::container_d2h_copy_alg - track_state_d2h{mr, copy}; - // Seed generator seed_generator sg(host_det, stddevs); diff --git a/tests/sycl/test_ckf_toy_detector.cpp b/tests/sycl/test_ckf_toy_detector.cpp index 013e47abe8..f67e928b8b 100644 --- a/tests/sycl/test_ckf_toy_detector.cpp +++ b/tests/sycl/test_ckf_toy_detector.cpp @@ -12,8 +12,6 @@ // Project include(s). #include "traccc/bfield/construct_const_bfield.hpp" #include "traccc/bfield/magnetic_field_types.hpp" -#include "traccc/device/container_d2h_copy_alg.hpp" -#include "traccc/device/container_h2d_copy_alg.hpp" #include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" #include "traccc/io/read_detector.hpp" #include "traccc/io/read_measurements.hpp" @@ -128,9 +126,6 @@ TEST_P(CkfToyDetectorTests, Run) { // Copy objects vecmem::sycl::async_copy copy{vecmem_queue}; - traccc::device::container_d2h_copy_alg - track_state_d2h{mr, copy}; - // Seed generator seed_generator sg(host_det, stddevs); diff --git a/tests/sycl/test_kalman_fitter_telescope.cpp b/tests/sycl/test_kalman_fitter_telescope.cpp index 4ff4840881..75849e883b 100644 --- a/tests/sycl/test_kalman_fitter_telescope.cpp +++ b/tests/sycl/test_kalman_fitter_telescope.cpp @@ -9,7 +9,7 @@ #include "traccc/bfield/construct_const_bfield.hpp" #include "traccc/bfield/magnetic_field_types.hpp" #include "traccc/device/container_d2h_copy_alg.hpp" -#include "traccc/edm/track_state.hpp" +#include "traccc/edm/track_fit_container.hpp" #include "traccc/io/utils.hpp" #include "traccc/performance/details/is_same_object.hpp" #include "traccc/resolution/fitting_performance_writer.hpp" @@ -141,9 +141,6 @@ TEST_P(KalmanFittingTelescopeTests, Run) { vecmem::sycl::copy copy{vecmem_queue}; - traccc::device::container_d2h_copy_alg - track_state_d2h{mr, copy}; - // Seed generator seed_generator sg(host_det, stddevs); @@ -176,32 +173,39 @@ TEST_P(KalmanFittingTelescopeTests, Run) { mr.main, vecmem::copy::type::host_to_device)}; // Run fitting - traccc::track_state_container_types::buffer track_states_sycl_buffer = + auto track_states_sycl_buffer = device_fitting(det_view, field, {track_candidates_buffer.tracks, track_candidates_buffer.measurements}); - traccc::track_state_container_types::host track_states_sycl = - track_state_d2h(track_states_sycl_buffer); + traccc::edm::track_fit_container::host + track_states_sycl{host_mr}; + copy(track_states_sycl_buffer.tracks, track_states_sycl.tracks, + vecmem::copy::type::device_to_host) + ->wait(); + copy(track_states_sycl_buffer.states, track_states_sycl.states, + vecmem::copy::type::device_to_host) + ->wait(); + const std::size_t n_fitted_tracks = - count_successfully_fitted_tracks(track_states_sycl); + count_successfully_fitted_tracks(track_states_sycl.tracks); - ASSERT_EQ(track_states_sycl.size(), n_truth_tracks); - ASSERT_EQ(track_states_sycl.size(), n_fitted_tracks); + ASSERT_EQ(track_states_sycl.tracks.size(), n_truth_tracks); + ASSERT_EQ(track_states_sycl.tracks.size(), n_fitted_tracks); for (std::size_t i_trk = 0; i_trk < n_truth_tracks; i_trk++) { - const auto& track_states_per_track = track_states_sycl[i_trk].items; - const auto& fit_res = track_states_sycl[i_trk].header; - - consistency_tests(track_states_per_track); + consistency_tests(track_states_sycl.tracks.at(i_trk), + track_states_sycl.states); - ndf_tests(fit_res, track_states_per_track); + ndf_tests(track_states_sycl.tracks.at(i_trk), + track_states_sycl.states, track_candidates.measurements); - ASSERT_EQ(fit_res.trk_quality.n_holes, 0u); + ASSERT_EQ(track_states_sycl.tracks.at(i_trk).nholes(), 0u); - fit_performance_writer.write(track_states_per_track, fit_res, - host_det, evt_data); + fit_performance_writer.write( + track_states_sycl.tracks.at(i_trk), track_states_sycl.states, + track_candidates.measurements, host_det, evt_data); } } From e35894e745da6be89236733f3cd4d3321ace9ca0 Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Tue, 5 Aug 2025 15:11:57 +0200 Subject: [PATCH 05/11] Adapted the examples to the SoA track fit EDM. --- examples/run/alpaka/full_chain_algorithm.cpp | 6 +- examples/run/alpaka/full_chain_algorithm.hpp | 4 +- .../run/alpaka/seeding_example_alpaka.cpp | 24 ++++---- examples/run/alpaka/seq_example_alpaka.cpp | 37 +++++++----- examples/run/cpu/full_chain_algorithm.cpp | 8 ++- examples/run/cpu/full_chain_algorithm.hpp | 11 ++-- .../cpu/misaligned_truth_fitting_example.cpp | 32 ++++------ examples/run/cpu/seeding_example.cpp | 16 +++-- examples/run/cpu/seq_example.cpp | 15 ++--- examples/run/cpu/truth_finding_example.cpp | 13 ++-- examples/run/cpu/truth_fitting_example.cpp | 14 ++--- examples/run/cuda/full_chain_algorithm.cpp | 6 +- examples/run/cuda/full_chain_algorithm.hpp | 4 +- examples/run/cuda/seeding_example_cuda.cpp | 34 +++++------ examples/run/cuda/seq_example_cuda.cpp | 42 ++++++++----- .../run/cuda/truth_finding_example_cuda.cpp | 36 +++++------ .../run/cuda/truth_fitting_example_cuda.cpp | 59 +++++++++++-------- examples/run/sycl/full_chain_algorithm.hpp | 2 +- examples/run/sycl/full_chain_algorithm.sycl | 6 +- 19 files changed, 192 insertions(+), 177 deletions(-) diff --git a/examples/run/alpaka/full_chain_algorithm.cpp b/examples/run/alpaka/full_chain_algorithm.cpp index db367ea8b5..d7a4744a4e 100644 --- a/examples/run/alpaka/full_chain_algorithm.cpp +++ b/examples/run/alpaka/full_chain_algorithm.cpp @@ -184,8 +184,8 @@ full_chain_algorithm::output_type full_chain_algorithm::operator()( m_device_detector_view, m_field, {track_candidates, measurements}); // Copy a limited amount of result data back to the host. - output_type result{&m_host_mr}; - m_vecmem_objects.async_copy()(track_states.headers, result)->wait(); + output_type result{m_host_mr}; + m_vecmem_objects.async_copy()(track_states.tracks, result)->wait(); return result; } @@ -198,7 +198,7 @@ full_chain_algorithm::output_type full_chain_algorithm::operator()( m_vecmem_objects.async_copy()(measurements, measurements_host)->wait(); // Return an empty object. - return {}; + return output_type{m_host_mr}; } } diff --git a/examples/run/alpaka/full_chain_algorithm.hpp b/examples/run/alpaka/full_chain_algorithm.hpp index 7e8074ee19..8cbd6fa7fe 100644 --- a/examples/run/alpaka/full_chain_algorithm.hpp +++ b/examples/run/alpaka/full_chain_algorithm.hpp @@ -20,8 +20,8 @@ #include "traccc/bfield/magnetic_field.hpp" #include "traccc/clusterization/clustering_config.hpp" #include "traccc/edm/silicon_cell_collection.hpp" +#include "traccc/edm/track_fit_collection.hpp" #include "traccc/edm/track_parameters.hpp" -#include "traccc/edm/track_state.hpp" #include "traccc/fitting/kalman_filter/kalman_fitter.hpp" #include "traccc/geometry/detector.hpp" #include "traccc/geometry/silicon_detector_description.hpp" @@ -44,7 +44,7 @@ namespace traccc::alpaka { /// At least as much as is implemented in the project at any given moment. /// class full_chain_algorithm - : public algorithm>( + : public algorithm::host( const edm::silicon_cell_collection::host&)>, public messaging { diff --git a/examples/run/alpaka/seeding_example_alpaka.cpp b/examples/run/alpaka/seeding_example_alpaka.cpp index 17c4e89aad..70114bdc8e 100644 --- a/examples/run/alpaka/seeding_example_alpaka.cpp +++ b/examples/run/alpaka/seeding_example_alpaka.cpp @@ -138,9 +138,6 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, vecmem::copy& copy = vo.copy(); vecmem::copy& async_copy = vo.async_copy(); - traccc::device::container_d2h_copy_alg - track_state_d2h{mr, copy, logger().clone("TrackStateD2HCopyAlg")}; - // Seeding algorithms const traccc::seedfinder_config seedfinder_config(seeding_opts); const traccc::seedfilter_config seedfilter_config(seeding_opts); @@ -199,7 +196,8 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, traccc::host::track_params_estimation::output_type params; traccc::edm::track_candidate_collection::host track_candidates{host_mr}; - traccc::track_state_container_types::host track_states; + traccc::edm::track_fit_container::host + track_states{host_mr}; traccc::edm::seed_collection::buffer seeds_alpaka_buffer; traccc::bound_track_parameters_collection_types::buffer @@ -208,8 +206,8 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, traccc::edm::track_candidate_collection::buffer track_candidates_alpaka_buffer; - traccc::track_state_container_types::buffer track_states_alpaka_buffer{ - {{}, *(mr.host)}, {{}, *(mr.host), mr.host}}; + traccc::edm::track_fit_container::buffer + track_states_alpaka_buffer; { // Start measuring wall time traccc::performance::timer wall_t("Wall time", elapsedTimes); @@ -355,8 +353,14 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, copy(track_candidates_alpaka_buffer, track_candidates_alpaka)->wait(); // Copy track states from device to host - traccc::track_state_container_types::host track_states_alpaka = - track_state_d2h(track_states_alpaka_buffer); + traccc::edm::track_fit_container::host + track_states_alpaka{host_mr}; + async_copy(track_states_alpaka_buffer.tracks, + track_states_alpaka.tracks) + ->wait(); + async_copy(track_states_alpaka_buffer.states, + track_states_alpaka.states) + ->wait(); if (accelerator_opts.compare_with_cpu) { // Show which event we are currently presenting the results for. @@ -403,8 +407,8 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, n_seeds += seeds.size(); n_found_tracks_alpaka += track_candidates_alpaka.size(); n_found_tracks += track_candidates.size(); - n_fitted_tracks_alpaka += track_states_alpaka.size(); - n_fitted_tracks += track_states.size(); + n_fitted_tracks_alpaka += track_states_alpaka.tracks.size(); + n_fitted_tracks += track_states.tracks.size(); /*------------ Writer diff --git a/examples/run/alpaka/seq_example_alpaka.cpp b/examples/run/alpaka/seq_example_alpaka.cpp index 705791a79d..0595e47ca8 100644 --- a/examples/run/alpaka/seq_example_alpaka.cpp +++ b/examples/run/alpaka/seq_example_alpaka.cpp @@ -185,9 +185,6 @@ int seq_run(const traccc::opts::detector& detector_opts, device_fitting_algorithm fitting_alg_alpaka( fitting_cfg, mr, copy, queue, logger().clone("AlpakaFittingAlg")); - traccc::device::container_d2h_copy_alg - copy_track_states(mr, copy, logger().clone("TrackStateD2HCopyAlg")); - // performance writer traccc::seeding_performance_writer sd_performance_writer( traccc::seeding_performance_writer::config{}, @@ -207,7 +204,7 @@ int seq_run(const traccc::opts::detector& detector_opts, traccc::host::seeding_algorithm::output_type seeds{host_mr}; traccc::host::track_params_estimation::output_type params{&host_mr}; host_finding_algorithm::output_type track_candidates{host_mr}; - host_fitting_algorithm::output_type track_states; + host_fitting_algorithm::output_type track_states{host_mr}; // Instantiate alpaka containers/collections traccc::measurement_collection_types::buffer measurements_alpaka_buffer( @@ -218,7 +215,8 @@ int seq_run(const traccc::opts::detector& detector_opts, params_alpaka_buffer(0, *mr.host); traccc::edm::track_candidate_collection::buffer track_candidates_buffer; - traccc::track_state_container_types::buffer track_states_buffer; + traccc::edm::track_fit_container::buffer + track_states_buffer; { traccc::performance::timer wall_t("Wall time", elapsedTimes); @@ -372,13 +370,16 @@ int seq_run(const traccc::opts::detector& detector_opts, &host_mr}; traccc::edm::track_candidate_collection::host track_candidates_alpaka{host_mr}; + traccc::edm::track_fit_container::host + track_states_alpaka{host_mr}; copy(measurements_alpaka_buffer, measurements_per_event_alpaka)->wait(); copy(spacepoints_alpaka_buffer, spacepoints_per_event_alpaka)->wait(); copy(seeds_alpaka_buffer, seeds_alpaka)->wait(); copy(params_alpaka_buffer, params_alpaka)->wait(); copy(track_candidates_buffer, track_candidates_alpaka)->wait(); - auto track_states_alpaka = copy_track_states(track_states_buffer); + copy(track_states_buffer.tracks, track_states_alpaka.tracks)->wait(); + copy(track_states_buffer.states, track_states_alpaka.states)->wait(); queue.synchronize(); if (accelerator_opts.compare_with_cpu) { @@ -430,12 +431,20 @@ int seq_run(const traccc::opts::detector& detector_opts, vecmem::get_data(track_candidates_alpaka)); // Compare tracks fitted on the host and on the device. - traccc::collection_comparator< - traccc::track_state_container_types::host::header_type> - compare_track_states{"track states"}; - compare_track_states( - vecmem::get_data(track_states.get_headers()), - vecmem::get_data(track_states_alpaka.get_headers())); + traccc::soa_comparator< + traccc::edm::track_fit_collection> + compare_track_fits{ + "track fits", + traccc::details::comparator_factory< + traccc::edm::track_fit_collection< + traccc::default_algebra>::const_device:: + const_proxy_type>{ + vecmem::get_data(measurements_per_event), + vecmem::get_data(measurements_per_event_alpaka), + vecmem::get_data(track_states.states), + vecmem::get_data(track_states_alpaka.states)}}; + compare_track_fits(vecmem::get_data(track_states.tracks), + vecmem::get_data(track_states_alpaka.tracks)); } /// Statistics n_measurements += measurements_per_event.size(); @@ -446,8 +455,8 @@ int seq_run(const traccc::opts::detector& detector_opts, n_seeds_alpaka += seeds_alpaka.size(); n_found_tracks += track_candidates.size(); n_found_tracks_alpaka += track_candidates_alpaka.size(); - n_fitted_tracks += track_states.size(); - n_fitted_tracks_alpaka += track_states_alpaka.size(); + n_fitted_tracks += track_states.tracks.size(); + n_fitted_tracks_alpaka += track_states_alpaka.tracks.size(); if (performance_opts.run) { diff --git a/examples/run/cpu/full_chain_algorithm.cpp b/examples/run/cpu/full_chain_algorithm.cpp index a9c30605e5..a07486a9fa 100644 --- a/examples/run/cpu/full_chain_algorithm.cpp +++ b/examples/run/cpu/full_chain_algorithm.cpp @@ -21,6 +21,7 @@ full_chain_algorithm::full_chain_algorithm( const magnetic_field& field, detector_type* detector, std::unique_ptr logger) : messaging(logger->clone()), + m_mr(mr), m_copy{std::make_unique()}, m_field_vec{0.f, 0.f, finder_config.bFieldInZ}, m_field(field), @@ -79,14 +80,15 @@ full_chain_algorithm::output_type full_chain_algorithm::operator()( // Run the track fitting, and return its results. return m_fitting( - *m_detector, m_field, - {vecmem::get_data(track_candidates), measurements_view}); + *m_detector, m_field, + {vecmem::get_data(track_candidates), measurements_view}) + .tracks; } // If not, just return an empty object. else { // Return an empty object. - return {}; + return output_type{m_mr.get()}; } } diff --git a/examples/run/cpu/full_chain_algorithm.hpp b/examples/run/cpu/full_chain_algorithm.hpp index 9555b8683c..0741796e57 100644 --- a/examples/run/cpu/full_chain_algorithm.hpp +++ b/examples/run/cpu/full_chain_algorithm.hpp @@ -11,8 +11,8 @@ #include "traccc/bfield/magnetic_field.hpp" #include "traccc/clusterization/clusterization_algorithm.hpp" #include "traccc/edm/silicon_cell_collection.hpp" +#include "traccc/edm/track_fit_collection.hpp" #include "traccc/edm/track_parameters.hpp" -#include "traccc/edm/track_state.hpp" #include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" #include "traccc/geometry/detector.hpp" @@ -38,9 +38,10 @@ namespace traccc { /// /// At least as much as is implemented in the project at any given moment. /// -class full_chain_algorithm : public algorithm, - public messaging { +class full_chain_algorithm + : public algorithm::host( + const edm::silicon_cell_collection::host&)>, + public messaging { public: /// @name Type declaration(s) @@ -97,6 +98,8 @@ class full_chain_algorithm : public algorithm m_mr; /// Vecmem copy object std::unique_ptr m_copy; /// Constant B field for the (seed) track parameter estimation diff --git a/examples/run/cpu/misaligned_truth_fitting_example.cpp b/examples/run/cpu/misaligned_truth_fitting_example.cpp index cd9e55e62a..637121be34 100644 --- a/examples/run/cpu/misaligned_truth_fitting_example.cpp +++ b/examples/run/cpu/misaligned_truth_fitting_example.cpp @@ -166,21 +166,17 @@ int main(int argc, char* argv[]) { {vecmem::get_data(truth_track_candidates.tracks), vecmem::get_data(truth_track_candidates.measurements)}); - print_fitted_tracks_statistics(track_states); + // print_fitted_tracks_statistics(track_states); - const decltype(track_states)::size_type n_fitted_tracks = - track_states.size(); + const std::size_t n_fitted_tracks = track_states.tracks.size(); if (performance_opts.run) { for (unsigned int i = 0; i < n_fitted_tracks; i++) { - const auto& trk_states_per_track = track_states.at(i).items; - - const auto& fit_res = track_states[i].header; - - fit_performance_writer.write(trk_states_per_track, fit_res, - host_det, evt_data, - fit_cfg0.propagation.context); + fit_performance_writer.write( + track_states.tracks.at(i), track_states.states, + truth_track_candidates.measurements, host_det, evt_data, + fit_cfg0.propagation.context); } } } else { @@ -195,21 +191,17 @@ int main(int argc, char* argv[]) { {vecmem::get_data(truth_track_candidates.tracks), vecmem::get_data(truth_track_candidates.measurements)}); - print_fitted_tracks_statistics(track_states); + // print_fitted_tracks_statistics(track_states); - const decltype(track_states)::size_type n_fitted_tracks = - track_states.size(); + const std::size_t n_fitted_tracks = track_states.tracks.size(); if (performance_opts.run) { for (unsigned int i = 0; i < n_fitted_tracks; i++) { - const auto& trk_states_per_track = track_states.at(i).items; - - const auto& fit_res = track_states[i].header; - - fit_performance_writer.write(trk_states_per_track, fit_res, - host_det, evt_data, - fit_cfg1.propagation.context); + fit_performance_writer.write( + track_states.tracks.at(i), track_states.states, + truth_track_candidates.measurements, host_det, evt_data, + fit_cfg1.propagation.context); } } } diff --git a/examples/run/cpu/seeding_example.cpp b/examples/run/cpu/seeding_example.cpp index fc5b77298d..c03516ea66 100644 --- a/examples/run/cpu/seeding_example.cpp +++ b/examples/run/cpu/seeding_example.cpp @@ -190,7 +190,8 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, track_candidates{host_mr}; traccc::edm::track_candidate_collection::host track_candidates_ar{host_mr}; - traccc::track_state_container_types::host track_states; + traccc::edm::track_fit_container::host + track_states{host_mr}; /*------------------------ Track Finding with CKF @@ -217,7 +218,7 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, track_states = host_fitting(detector, field, {vecmem::get_data(track_candidates_ar), vecmem::get_data(measurements_per_event)}); - n_fitted_tracks += track_states.size(); + n_fitted_tracks += track_states.tracks.size(); /*------------ Statistics @@ -249,13 +250,10 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, vecmem::get_data(track_candidates_ar), vecmem::get_data(measurements_per_event), evt_data); - for (unsigned int i = 0; i < track_states.size(); i++) { - const auto& trk_states_per_track = track_states.at(i).items; - - const auto& fit_res = track_states[i].header; - - fit_performance_writer.write(trk_states_per_track, fit_res, - detector, 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); } } } diff --git a/examples/run/cpu/seq_example.cpp b/examples/run/cpu/seq_example.cpp index 630588cb10..87673ef863 100644 --- a/examples/run/cpu/seq_example.cpp +++ b/examples/run/cpu/seq_example.cpp @@ -199,7 +199,7 @@ int seq_run(const traccc::opts::input_data& input_opts, finding_algorithm::output_type track_candidates{host_mr}; traccc::host::greedy_ambiguity_resolution_algorithm::output_type resolved_track_candidates{host_mr}; - fitting_algorithm::output_type track_states{&host_mr}; + fitting_algorithm::output_type track_states{host_mr}; { // Start measuring wall time. traccc::performance::timer timer_wall{"Wall time", elapsedTimes}; @@ -320,7 +320,7 @@ int seq_run(const traccc::opts::input_data& input_opts, n_seeds += seeds.size(); n_found_tracks += track_candidates.size(); n_ambiguity_free_tracks += resolved_track_candidates.size(); - n_fitted_tracks += track_states.size(); + n_fitted_tracks += track_states.tracks.size(); } // Stop measuring Wall time. @@ -347,13 +347,10 @@ int seq_run(const traccc::opts::input_data& input_opts, vecmem::get_data(resolved_track_candidates), vecmem::get_data(measurements_per_event), evt_data); - for (unsigned int i = 0; i < track_states.size(); i++) { - const auto& trk_states_per_track = track_states.at(i).items; - - const auto& fit_res = track_states[i].header; - - fit_performance_writer.write(trk_states_per_track, fit_res, - detector, 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); } } } diff --git a/examples/run/cpu/truth_finding_example.cpp b/examples/run/cpu/truth_finding_example.cpp index 96e8c9a705..ea795f153b 100644 --- a/examples/run/cpu/truth_finding_example.cpp +++ b/examples/run/cpu/truth_finding_example.cpp @@ -163,9 +163,9 @@ int seq_run(const traccc::opts::track_finding& finding_opts, {vecmem::get_data(track_candidates), vecmem::get_data(measurements_per_event)}); - print_fitted_tracks_statistics(track_states); + // print_fitted_tracks_statistics(track_states); - const std::size_t n_fitted_tracks = track_states.size(); + const std::size_t n_fitted_tracks = track_states.tracks.size(); if (performance_opts.run) { find_performance_writer.write( @@ -173,12 +173,9 @@ int seq_run(const traccc::opts::track_finding& finding_opts, vecmem::get_data(measurements_per_event), evt_data); for (std::size_t i = 0; i < n_fitted_tracks; i++) { - const auto& trk_states_per_track = track_states.at(i).items; - - const auto& fit_res = track_states[i].header; - - fit_performance_writer.write(trk_states_per_track, fit_res, - detector, evt_data); + fit_performance_writer.write( + track_states.tracks.at(i), track_states.states, + measurements_per_event, detector, evt_data); } } } diff --git a/examples/run/cpu/truth_fitting_example.cpp b/examples/run/cpu/truth_fitting_example.cpp index f4b3d798d4..0c626661ba 100644 --- a/examples/run/cpu/truth_fitting_example.cpp +++ b/examples/run/cpu/truth_fitting_example.cpp @@ -140,20 +140,16 @@ int main(int argc, char* argv[]) { {vecmem::get_data(truth_track_candidates.tracks), vecmem::get_data(truth_track_candidates.measurements)}); - print_fitted_tracks_statistics(track_states); + // print_fitted_tracks_statistics(track_states); - const decltype(track_states)::size_type n_fitted_tracks = - track_states.size(); + const std::size_t n_fitted_tracks = track_states.tracks.size(); if (performance_opts.run) { for (unsigned int i = 0; i < n_fitted_tracks; i++) { - const auto& trk_states_per_track = track_states.at(i).items; - - const auto& fit_res = track_states[i].header; - - fit_performance_writer.write(trk_states_per_track, fit_res, - host_det, evt_data); + fit_performance_writer.write( + track_states.tracks.at(i), track_states.states, + truth_track_candidates.measurements, host_det, evt_data); } } } diff --git a/examples/run/cuda/full_chain_algorithm.cpp b/examples/run/cuda/full_chain_algorithm.cpp index 78dfae2fc9..cb23475fa7 100644 --- a/examples/run/cuda/full_chain_algorithm.cpp +++ b/examples/run/cuda/full_chain_algorithm.cpp @@ -190,8 +190,8 @@ full_chain_algorithm::output_type full_chain_algorithm::operator()( m_device_detector_view, m_field, {track_candidates, measurements}); // Copy a limited amount of result data back to the host. - output_type result{&m_host_mr}; - m_copy(track_states.headers, result)->wait(); + output_type result{m_host_mr}; + m_copy(track_states.tracks, result)->wait(); return result; } @@ -204,7 +204,7 @@ full_chain_algorithm::output_type full_chain_algorithm::operator()( m_copy(measurements, measurements_host)->wait(); // Return an empty object. - return {}; + return output_type{m_host_mr}; } } diff --git a/examples/run/cuda/full_chain_algorithm.hpp b/examples/run/cuda/full_chain_algorithm.hpp index fd8c379d7c..fe9c3155de 100644 --- a/examples/run/cuda/full_chain_algorithm.hpp +++ b/examples/run/cuda/full_chain_algorithm.hpp @@ -18,8 +18,8 @@ #include "traccc/cuda/seeding/track_params_estimation.hpp" #include "traccc/cuda/utils/stream.hpp" #include "traccc/edm/silicon_cell_collection.hpp" +#include "traccc/edm/track_fit_collection.hpp" #include "traccc/edm/track_parameters.hpp" -#include "traccc/edm/track_state.hpp" #include "traccc/geometry/detector.hpp" #include "traccc/geometry/silicon_detector_description.hpp" #include "traccc/utils/algorithm.hpp" @@ -43,7 +43,7 @@ namespace traccc::cuda { /// At least as much as is implemented in the project at any given moment. /// class full_chain_algorithm - : public algorithm>( + : public algorithm::host( const edm::silicon_cell_collection::host&)>, public messaging { diff --git a/examples/run/cuda/seeding_example_cuda.cpp b/examples/run/cuda/seeding_example_cuda.cpp index 7de3b77fbd..76edcc69b7 100644 --- a/examples/run/cuda/seeding_example_cuda.cpp +++ b/examples/run/cuda/seeding_example_cuda.cpp @@ -146,9 +146,6 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, vecmem::copy host_copy; vecmem::cuda::copy copy; - traccc::device::container_d2h_copy_alg - track_state_d2h{mr, copy, logger().clone("TrackStateD2HCopyAlg")}; - // Seeding algorithm const traccc::seedfinder_config seedfinder_config(seeding_opts); const traccc::seedfilter_config seedfilter_config(seeding_opts); @@ -209,7 +206,8 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, traccc::host::track_params_estimation::output_type params; traccc::edm::track_candidate_collection::host track_candidates{host_mr}; - traccc::track_state_container_types::host track_states; + traccc::edm::track_fit_container::host + track_states{host_mr}; traccc::edm::seed_collection::buffer seeds_cuda_buffer; traccc::bound_track_parameters_collection_types::buffer @@ -218,8 +216,8 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, traccc::edm::track_candidate_collection::buffer track_candidates_cuda_buffer; - traccc::track_state_container_types::buffer track_states_cuda_buffer{ - {{}, *(mr.host)}, {{}, *(mr.host), mr.host}}; + traccc::edm::track_fit_container::buffer + track_states_cuda_buffer; { // Start measuring wall time traccc::performance::timer wall_t("Wall time", elapsedTimes); @@ -364,8 +362,12 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, async_copy(track_candidates_cuda_buffer, track_candidates_cuda)->wait(); // Copy track states from device to host - traccc::track_state_container_types::host track_states_cuda = - track_state_d2h(track_states_cuda_buffer); + traccc::edm::track_fit_container::host + track_states_cuda{host_mr}; + async_copy(track_states_cuda_buffer.tracks, track_states_cuda.tracks) + ->wait(); + async_copy(track_states_cuda_buffer.states, track_states_cuda.states) + ->wait(); if (accelerator_opts.compare_with_cpu) { // Show which event we are currently presenting the results for. @@ -412,8 +414,8 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, n_seeds += seeds.size(); n_found_tracks_cuda += track_candidates_cuda.size(); n_found_tracks += track_candidates.size(); - n_fitted_tracks_cuda += track_states_cuda.size(); - n_fitted_tracks += track_states.size(); + n_fitted_tracks_cuda += track_states_cuda.tracks.size(); + n_fitted_tracks += track_states.tracks.size(); /*------------ Writer @@ -434,14 +436,10 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, vecmem::get_data(track_candidates_cuda), vecmem::get_data(measurements_per_event), evt_data); - for (unsigned int i = 0; i < track_states_cuda.size(); i++) { - const auto& trk_states_per_track = - track_states_cuda.at(i).items; - - const auto& fit_res = track_states_cuda[i].header; - - fit_performance_writer.write(trk_states_per_track, fit_res, - host_det, 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); } } } diff --git a/examples/run/cuda/seq_example_cuda.cpp b/examples/run/cuda/seq_example_cuda.cpp index eed09614e2..8f1c64d2a4 100644 --- a/examples/run/cuda/seq_example_cuda.cpp +++ b/examples/run/cuda/seq_example_cuda.cpp @@ -211,9 +211,6 @@ int seq_run(const traccc::opts::detector& detector_opts, device_fitting_algorithm fitting_alg_cuda(fitting_cfg, mr, copy, stream, logger().clone("CudaFittingAlg")); - traccc::device::container_d2h_copy_alg - copy_track_states(mr, copy, logger().clone("TrackStateD2HCopyAlg")); - // performance writer traccc::seeding_performance_writer sd_performance_writer( traccc::seeding_performance_writer::config{}, @@ -235,7 +232,7 @@ int seq_run(const traccc::opts::detector& detector_opts, host_finding_algorithm::output_type track_candidates{host_mr}; traccc::host::greedy_ambiguity_resolution_algorithm::output_type res_track_candidates{host_mr}; - host_fitting_algorithm::output_type track_states; + host_fitting_algorithm::output_type track_states{host_mr}; // Instantiate cuda containers/collections traccc::measurement_collection_types::buffer measurements_cuda_buffer( @@ -248,7 +245,8 @@ int seq_run(const traccc::opts::detector& detector_opts, track_candidates_buffer; traccc::edm::track_candidate_collection::buffer res_track_candidates_buffer; - traccc::track_state_container_types::buffer track_states_buffer; + traccc::edm::track_fit_container::buffer + track_states_buffer; { traccc::performance::timer wall_t("Wall time", elapsedTimes); @@ -421,6 +419,8 @@ int seq_run(const traccc::opts::detector& detector_opts, track_candidates_cuda{host_mr}; traccc::edm::track_candidate_collection::host res_track_candidates_cuda{host_mr}; + traccc::edm::track_fit_container::host + track_states_cuda{host_mr}; copy(measurements_cuda_buffer, measurements_per_event_cuda)->wait(); copy(spacepoints_cuda_buffer, spacepoints_per_event_cuda)->wait(); @@ -432,8 +432,12 @@ int seq_run(const traccc::opts::detector& detector_opts, copy(res_track_candidates_buffer, res_track_candidates_cuda, vecmem::copy::type::device_to_host) ->wait(); - - auto track_states_cuda = copy_track_states(track_states_buffer); + copy(track_states_buffer.tracks, track_states_cuda.tracks, + vecmem::copy::type::device_to_host) + ->wait(); + copy(track_states_buffer.states, track_states_cuda.states, + vecmem::copy::type::device_to_host) + ->wait(); stream.synchronize(); if (accelerator_opts.compare_with_cpu) { @@ -499,12 +503,20 @@ int seq_run(const traccc::opts::detector& detector_opts, vecmem::get_data(res_track_candidates_cuda)); // Compare tracks fitted on the host and on the device. - traccc::collection_comparator< - traccc::track_state_container_types::host::header_type> - compare_track_states{"track states"}; - compare_track_states( - vecmem::get_data(track_states.get_headers()), - vecmem::get_data(track_states_cuda.get_headers())); + traccc::soa_comparator< + traccc::edm::track_fit_collection> + compare_track_fits{ + "track fits", + traccc::details::comparator_factory< + traccc::edm::track_fit_collection< + traccc::default_algebra>::const_device:: + const_proxy_type>{ + vecmem::get_data(measurements_per_event), + vecmem::get_data(measurements_per_event_cuda), + vecmem::get_data(track_states.states), + vecmem::get_data(track_states_cuda.states)}}; + compare_track_fits(vecmem::get_data(track_states.tracks), + vecmem::get_data(track_states_cuda.tracks)); } /// Statistics n_measurements += measurements_per_event.size(); @@ -517,8 +529,8 @@ int seq_run(const traccc::opts::detector& detector_opts, n_found_tracks_cuda += track_candidates_cuda.size(); n_ambiguity_free_tracks += res_track_candidates.size(); n_ambiguity_free_tracks_cuda += res_track_candidates_cuda.size(); - n_fitted_tracks += track_states.size(); - n_fitted_tracks_cuda += track_states_cuda.size(); + n_fitted_tracks += track_states.tracks.size(); + n_fitted_tracks_cuda += track_states_cuda.tracks.size(); if (performance_opts.run) { diff --git a/examples/run/cuda/truth_finding_example_cuda.cpp b/examples/run/cuda/truth_finding_example_cuda.cpp index a5a68ad0ef..e3a8a0c97b 100644 --- a/examples/run/cuda/truth_finding_example_cuda.cpp +++ b/examples/run/cuda/truth_finding_example_cuda.cpp @@ -122,9 +122,6 @@ int seq_run(const traccc::opts::track_finding& finding_opts, vecmem::copy host_copy; vecmem::cuda::async_copy async_copy{stream.cudaStream()}; - traccc::device::container_d2h_copy_alg - track_state_d2h{mr, async_copy, logger().clone("TrackStateD2HCopyAlg")}; - // Standard deviations for seed track parameters static constexpr std::array stddevs = {1e-4f * traccc::unit::mm, @@ -224,8 +221,8 @@ int seq_run(const traccc::opts::track_finding& finding_opts, ->wait(); // Instantiate cuda containers/collections - traccc::track_state_container_types::buffer track_states_cuda_buffer{ - {{}, *(mr.host)}, {{}, *(mr.host), mr.host}}; + traccc::edm::track_fit_container::buffer + track_states_cuda_buffer; { traccc::performance::timer t("Track fitting (cuda)", elapsedTimes); @@ -235,13 +232,20 @@ int seq_run(const traccc::opts::track_finding& finding_opts, det_view, device_field, {track_candidates_cuda_buffer, measurements_cuda_buffer}); } - traccc::track_state_container_types::host track_states_cuda = - track_state_d2h(track_states_cuda_buffer); + traccc::edm::track_fit_container::host + track_states_cuda{host_mr}; + async_copy(track_states_cuda_buffer.tracks, track_states_cuda.tracks, + vecmem::copy::type::device_to_host) + ->wait(); + async_copy(track_states_cuda_buffer.states, track_states_cuda.states, + vecmem::copy::type::device_to_host) + ->wait(); // CPU containers traccc::host::combinatorial_kalman_filter_algorithm::output_type track_candidates{host_mr}; - traccc::host::kalman_fitting_algorithm::output_type track_states; + traccc::host::kalman_fitting_algorithm::output_type track_states{ + host_mr}; if (accelerator_opts.compare_with_cpu) { @@ -290,23 +294,19 @@ int seq_run(const traccc::opts::track_finding& finding_opts, /// Statistics n_found_tracks += track_candidates.size(); - n_fitted_tracks += track_states.size(); + n_fitted_tracks += track_states.tracks.size(); n_found_tracks_cuda += track_candidates_cuda.size(); - n_fitted_tracks_cuda += track_states_cuda.size(); + n_fitted_tracks_cuda += track_states_cuda.tracks.size(); if (performance_opts.run) { find_performance_writer.write( vecmem::get_data(track_candidates_cuda), vecmem::get_data(measurements_per_event), evt_data); - for (unsigned int i = 0; i < track_states_cuda.size(); i++) { - const auto& trk_states_per_track = - track_states_cuda.at(i).items; - - const auto& fit_res = track_states_cuda[i].header; - - fit_performance_writer.write(trk_states_per_track, fit_res, - detector, 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); } } } diff --git a/examples/run/cuda/truth_fitting_example_cuda.cpp b/examples/run/cuda/truth_fitting_example_cuda.cpp index 810b1cbbaf..08c3d0b21a 100644 --- a/examples/run/cuda/truth_fitting_example_cuda.cpp +++ b/examples/run/cuda/truth_fitting_example_cuda.cpp @@ -27,8 +27,7 @@ #include "traccc/options/program_options.hpp" #include "traccc/options/track_fitting.hpp" #include "traccc/options/track_propagation.hpp" -#include "traccc/performance/collection_comparator.hpp" -#include "traccc/performance/container_comparator.hpp" +#include "traccc/performance/soa_comparator.hpp" #include "traccc/performance/timer.hpp" #include "traccc/resolution/fitting_performance_writer.hpp" #include "traccc/utils/propagation.hpp" @@ -131,9 +130,6 @@ int main(int argc, char* argv[]) { vecmem::copy host_copy; vecmem::cuda::async_copy async_copy{stream.cudaStream()}; - traccc::device::container_d2h_copy_alg - track_state_d2h{mr, async_copy, logger().clone("TrackStateD2HCopyAlg")}; - /// Standard deviations for seed track parameters static constexpr std::array stddevs = { 0.03f * traccc::unit::mm, @@ -181,8 +177,8 @@ int main(int argc, char* argv[]) { mr.main, vecmem::copy::type::host_to_device)}; // Instantiate cuda containers/collections - traccc::track_state_container_types::buffer track_states_cuda_buffer{ - {{}, *(mr.host)}, {{}, *(mr.host), mr.host}}; + traccc::edm::track_fit_container::buffer + track_states_cuda_buffer; { traccc::performance::timer t("Track fitting (cuda)", elapsedTimes); @@ -194,11 +190,18 @@ int main(int argc, char* argv[]) { truth_track_candidates_buffer.measurements}); } - traccc::track_state_container_types::host track_states_cuda = - track_state_d2h(track_states_cuda_buffer); + traccc::edm::track_fit_container::host + track_states_cuda{host_mr}; + async_copy(track_states_cuda_buffer.tracks, track_states_cuda.tracks, + vecmem::copy::type::device_to_host) + ->wait(); + async_copy(track_states_cuda_buffer.states, track_states_cuda.states, + vecmem::copy::type::device_to_host) + ->wait(); // CPU container(s) - traccc::host::kalman_fitting_algorithm::output_type track_states; + traccc::host::kalman_fitting_algorithm::output_type track_states{ + host_mr}; if (accelerator_opts.compare_with_cpu) { @@ -219,27 +222,31 @@ int main(int argc, char* argv[]) { std::cout << "===>>> Event " << event << " <<<===" << std::endl; // Compare the track parameters made on the host and on the device. - traccc::collection_comparator< - traccc::fitting_result> - compare_fitting_results{"fitted tracks"}; - compare_fitting_results( - vecmem::get_data(track_states.get_headers()), - vecmem::get_data(track_states_cuda.get_headers())); + traccc::soa_comparator< + traccc::edm::track_fit_collection> + compare_track_fits{ + "track fits", + traccc::details::comparator_factory< + traccc::edm::track_fit_collection< + traccc::default_algebra>::const_device:: + const_proxy_type>{ + vecmem::get_data(truth_track_candidates.measurements), + vecmem::get_data(truth_track_candidates.measurements), + vecmem::get_data(track_states.states), + vecmem::get_data(track_states_cuda.states)}}; + compare_track_fits(vecmem::get_data(track_states.tracks), + vecmem::get_data(track_states_cuda.tracks)); } // Statistics - n_fitted_tracks += track_states.size(); - n_fitted_tracks_cuda += track_states_cuda.size(); + n_fitted_tracks += track_states.tracks.size(); + n_fitted_tracks_cuda += track_states_cuda.tracks.size(); if (performance_opts.run) { - for (unsigned int i = 0; i < track_states_cuda.size(); i++) { - const auto& trk_states_per_track = - track_states_cuda.at(i).items; - - const auto& fit_res = track_states_cuda[i].header; - - fit_performance_writer.write(trk_states_per_track, fit_res, - host_det, 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, + truth_track_candidates.measurements, host_det, evt_data); } } } diff --git a/examples/run/sycl/full_chain_algorithm.hpp b/examples/run/sycl/full_chain_algorithm.hpp index e71a6fed17..a3d0f288ba 100644 --- a/examples/run/sycl/full_chain_algorithm.hpp +++ b/examples/run/sycl/full_chain_algorithm.hpp @@ -42,7 +42,7 @@ struct full_chain_algorithm_data; /// At least as much as is implemented in the project at any given moment. /// class full_chain_algorithm - : public algorithm>( + : public algorithm::host( const edm::silicon_cell_collection::host&)>, public messaging { diff --git a/examples/run/sycl/full_chain_algorithm.sycl b/examples/run/sycl/full_chain_algorithm.sycl index 6623fc3339..5ed191869e 100644 --- a/examples/run/sycl/full_chain_algorithm.sycl +++ b/examples/run/sycl/full_chain_algorithm.sycl @@ -222,8 +222,8 @@ full_chain_algorithm::output_type full_chain_algorithm::operator()( m_device_detector_view, m_field, {track_candidates, measurements}); // Copy a limited amount of result data back to the host. - output_type result{&(m_host_mr.get())}; - m_copy(track_states.headers, result)->wait(); + output_type result{m_host_mr.get()}; + m_copy(track_states.tracks, result)->wait(); return result; } // If not, copy the measurements back to the host, and return @@ -236,7 +236,7 @@ full_chain_algorithm::output_type full_chain_algorithm::operator()( m_copy(measurements, measurements_host)->wait(); // Return an empty object. - return {}; + return output_type{m_host_mr.get()}; } } From 5b6c3010b03ddfe54960adad79dc50b07cf5064b Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Mon, 11 Aug 2025 13:46:17 +0200 Subject: [PATCH 06/11] Introduced track_candidates_container::device and track_fit_container::device. To make it a bit easier to deal with the slightly complex EDMs of track finding and fitting in the code. --- .../traccc/edm/track_candidate_container.hpp | 23 ++++++++ .../traccc/edm/track_fit_container.hpp | 31 ++++++++++ .../traccc/fitting/device/fit_backward.hpp | 15 ++--- .../traccc/fitting/device/fit_forward.hpp | 15 ++--- examples/run/cpu/seeding_example.cpp | 10 ++-- examples/run/cpu/seq_example.cpp | 10 ++-- examples/run/cpu/truth_finding_example.cpp | 5 +- examples/run/cuda/seeding_example_cuda.cpp | 5 +- .../run/cuda/truth_finding_example_cuda.cpp | 5 +- .../efficiency/finding_performance_writer.hpp | 26 +++------ .../efficiency/finding_performance_writer.cpp | 57 +++++++------------ 11 files changed, 115 insertions(+), 87 deletions(-) diff --git a/core/include/traccc/edm/track_candidate_container.hpp b/core/include/traccc/edm/track_candidate_container.hpp index 742fc5cd4d..ef7c28c0bc 100644 --- a/core/include/traccc/edm/track_candidate_container.hpp +++ b/core/include/traccc/edm/track_candidate_container.hpp @@ -8,6 +8,7 @@ #pragma once // Local include(s). +#include "traccc/definitions/qualifiers.hpp" #include "traccc/edm/measurement.hpp" #include "traccc/edm/track_candidate_collection.hpp" @@ -71,6 +72,28 @@ struct track_candidate_container { measurement_collection_types::const_view measurements; }; + struct device { + /// Constructor from a view + TRACCC_HOST_DEVICE + explicit device(const view& v) + : tracks{v.tracks}, measurements{v.measurements} {} + /// The track candidates + track_candidate_collection::device tracks; + /// Measurements referenced by the tracks + measurement_collection_types::const_device measurements; + }; + + struct const_device { + /// Constructor from a view + TRACCC_HOST_DEVICE + explicit const_device(const const_view& v) + : tracks{v.tracks}, measurements{v.measurements} {} + /// The track candidates + track_candidate_collection::const_device tracks; + /// Measurements referenced by the tracks + measurement_collection_types::const_device measurements; + }; + }; // struct track_candidate_container } // namespace traccc::edm diff --git a/core/include/traccc/edm/track_fit_container.hpp b/core/include/traccc/edm/track_fit_container.hpp index 99d238b2ed..541f9f34bc 100644 --- a/core/include/traccc/edm/track_fit_container.hpp +++ b/core/include/traccc/edm/track_fit_container.hpp @@ -8,6 +8,7 @@ #pragma once // Local include(s). +#include "traccc/definitions/qualifiers.hpp" #include "traccc/edm/measurement.hpp" #include "traccc/edm/track_fit_collection.hpp" #include "traccc/edm/track_state_collection.hpp" @@ -70,6 +71,36 @@ struct track_fit_container { track_state_collection::const_data states; }; + struct device { + /// Constructor from a view + TRACCC_HOST_DEVICE + explicit device(const view& v) + : tracks{v.tracks}, + states{v.states}, + measurements{v.measurements} {} + /// The fitted tracks + track_fit_collection::device tracks; + /// The track states used for the fit + track_state_collection::device states; + /// The measurements used for the fit + measurement_collection_types::const_device measurements; + }; + + struct const_device { + /// Constructor from a view + TRACCC_HOST_DEVICE + explicit const_device(const const_view& v) + : tracks{v.tracks}, + states{v.states}, + measurements{v.measurements} {} + /// The fitted tracks + track_fit_collection::const_device tracks; + /// The track states used for the fit + track_state_collection::const_device states; + /// The measurements used for the fit + measurement_collection_types::const_device measurements; + }; + }; // struct track_fit_container } // namespace traccc::edm diff --git a/device/common/include/traccc/fitting/device/fit_backward.hpp b/device/common/include/traccc/fitting/device/fit_backward.hpp index 5787a7cca9..a090b1e48e 100644 --- a/device/common/include/traccc/fitting/device/fit_backward.hpp +++ b/device/common/include/traccc/fitting/device/fit_backward.hpp @@ -22,28 +22,23 @@ TRACCC_HOST_DEVICE inline void fit_backward( vecmem::device_vector param_ids(payload.param_ids_view); vecmem::device_vector param_liveness( payload.param_liveness_view); - typename edm::track_fit_collection< + typename edm::track_fit_container< typename fitter_t::detector_type::algebra_type>::device - tracks(payload.tracks_view.tracks); - typename edm::track_state_collection< - typename fitter_t::detector_type::algebra_type>::device - track_states(payload.tracks_view.states); - measurement_collection_types::const_device measurements{ - payload.tracks_view.measurements}; + tracks(payload.tracks_view); - if (globalIndex >= tracks.size()) { + if (globalIndex >= tracks.tracks.size()) { return; } const unsigned int param_id = param_ids.at(globalIndex); - auto track = tracks.at(param_id); + auto track = tracks.tracks.at(param_id); // Run fitting fitter_t fitter(det, payload.field_data, cfg); if (param_liveness.at(param_id) > 0u) { typename fitter_t::state fitter_state( - track, track_states, measurements, + track, tracks.states, tracks.measurements, *(payload.barcodes_view.ptr() + param_id)); kalman_fitter_status fit_status = fitter.smooth(fitter_state); diff --git a/device/common/include/traccc/fitting/device/fit_forward.hpp b/device/common/include/traccc/fitting/device/fit_forward.hpp index 9957cd1d98..e6eba1a1f6 100644 --- a/device/common/include/traccc/fitting/device/fit_forward.hpp +++ b/device/common/include/traccc/fitting/device/fit_forward.hpp @@ -22,16 +22,11 @@ TRACCC_HOST_DEVICE inline void fit_forward( vecmem::device_vector param_ids(payload.param_ids_view); vecmem::device_vector param_liveness( payload.param_liveness_view); - typename edm::track_fit_collection< + typename edm::track_fit_container< typename fitter_t::detector_type::algebra_type>::device - tracks(payload.tracks_view.tracks); - typename edm::track_state_collection< - typename fitter_t::detector_type::algebra_type>::device - track_states(payload.tracks_view.states); - measurement_collection_types::const_device measurements{ - payload.tracks_view.measurements}; + tracks(payload.tracks_view); - if (globalIndex >= tracks.size()) { + if (globalIndex >= tracks.tracks.size()) { return; } @@ -39,14 +34,14 @@ TRACCC_HOST_DEVICE inline void fit_forward( fitter_t fitter(det, payload.field_data, cfg); - auto track = tracks.at(param_id); + auto track = tracks.tracks.at(param_id); auto params = track.params(); // TODO: Merge into filter? inflate_covariance(params, fitter.config().covariance_inflation_factor); typename fitter_t::state fitter_state( - track, track_states, measurements, + track, tracks.states, tracks.measurements, *(payload.barcodes_view.ptr() + param_id)); kalman_fitter_status fit_status = fitter.filter(params, fitter_state); diff --git a/examples/run/cpu/seeding_example.cpp b/examples/run/cpu/seeding_example.cpp index c03516ea66..28c021a742 100644 --- a/examples/run/cpu/seeding_example.cpp +++ b/examples/run/cpu/seeding_example.cpp @@ -243,12 +243,14 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, vecmem::get_data(measurements_per_event), evt_data); find_performance_writer.write( - vecmem::get_data(track_candidates), - vecmem::get_data(measurements_per_event), evt_data); + {vecmem::get_data(track_candidates), + vecmem::get_data(measurements_per_event)}, + evt_data); ar_performance_writer.write( - vecmem::get_data(track_candidates_ar), - vecmem::get_data(measurements_per_event), evt_data); + {vecmem::get_data(track_candidates_ar), + vecmem::get_data(measurements_per_event)}, + evt_data); for (unsigned int i = 0; i < track_states.tracks.size(); i++) { fit_performance_writer.write( diff --git a/examples/run/cpu/seq_example.cpp b/examples/run/cpu/seq_example.cpp index 87673ef863..e59f31e339 100644 --- a/examples/run/cpu/seq_example.cpp +++ b/examples/run/cpu/seq_example.cpp @@ -341,11 +341,13 @@ int seq_run(const traccc::opts::input_data& input_opts, vecmem::get_data(spacepoints_per_event), vecmem::get_data(measurements_per_event), evt_data); find_performance_writer.write( - vecmem::get_data(track_candidates), - vecmem::get_data(measurements_per_event), evt_data); + {vecmem::get_data(track_candidates), + vecmem::get_data(measurements_per_event)}, + evt_data); ar_performance_writer.write( - vecmem::get_data(resolved_track_candidates), - vecmem::get_data(measurements_per_event), evt_data); + {vecmem::get_data(resolved_track_candidates), + vecmem::get_data(measurements_per_event)}, + evt_data); for (unsigned int i = 0; i < track_states.tracks.size(); i++) { fit_performance_writer.write( diff --git a/examples/run/cpu/truth_finding_example.cpp b/examples/run/cpu/truth_finding_example.cpp index ea795f153b..9002b0336d 100644 --- a/examples/run/cpu/truth_finding_example.cpp +++ b/examples/run/cpu/truth_finding_example.cpp @@ -169,8 +169,9 @@ int seq_run(const traccc::opts::track_finding& finding_opts, if (performance_opts.run) { find_performance_writer.write( - vecmem::get_data(track_candidates), - vecmem::get_data(measurements_per_event), evt_data); + {vecmem::get_data(track_candidates), + vecmem::get_data(measurements_per_event)}, + evt_data); for (std::size_t i = 0; i < n_fitted_tracks; i++) { fit_performance_writer.write( diff --git a/examples/run/cuda/seeding_example_cuda.cpp b/examples/run/cuda/seeding_example_cuda.cpp index 76edcc69b7..28785e1f52 100644 --- a/examples/run/cuda/seeding_example_cuda.cpp +++ b/examples/run/cuda/seeding_example_cuda.cpp @@ -433,8 +433,9 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts, vecmem::get_data(measurements_per_event), evt_data); find_performance_writer.write( - vecmem::get_data(track_candidates_cuda), - vecmem::get_data(measurements_per_event), evt_data); + {vecmem::get_data(track_candidates_cuda), + vecmem::get_data(measurements_per_event)}, + evt_data); for (unsigned int i = 0; i < track_states_cuda.tracks.size(); i++) { fit_performance_writer.write( diff --git a/examples/run/cuda/truth_finding_example_cuda.cpp b/examples/run/cuda/truth_finding_example_cuda.cpp index e3a8a0c97b..ebc40787d5 100644 --- a/examples/run/cuda/truth_finding_example_cuda.cpp +++ b/examples/run/cuda/truth_finding_example_cuda.cpp @@ -300,8 +300,9 @@ int seq_run(const traccc::opts::track_finding& finding_opts, if (performance_opts.run) { find_performance_writer.write( - vecmem::get_data(track_candidates_cuda), - vecmem::get_data(measurements_per_event), evt_data); + {vecmem::get_data(track_candidates_cuda), + vecmem::get_data(measurements_per_event)}, + evt_data); for (unsigned int i = 0; i < track_states_cuda.tracks.size(); i++) { fit_performance_writer.write( diff --git a/performance/include/traccc/efficiency/finding_performance_writer.hpp b/performance/include/traccc/efficiency/finding_performance_writer.hpp index 27a193471b..30abc444b6 100644 --- a/performance/include/traccc/efficiency/finding_performance_writer.hpp +++ b/performance/include/traccc/efficiency/finding_performance_writer.hpp @@ -15,10 +15,8 @@ #include "traccc/utils/truth_matching_config.hpp" // Project include(s). -#include "traccc/edm/measurement.hpp" -#include "traccc/edm/track_candidate_collection.hpp" -#include "traccc/edm/track_fit_collection.hpp" -#include "traccc/edm/track_state_collection.hpp" +#include "traccc/edm/track_candidate_container.hpp" +#include "traccc/edm/track_fit_container.hpp" #include "traccc/utils/event_data.hpp" // System include(s). @@ -71,19 +69,13 @@ class finding_performance_writer : public messaging { /// Destructor ~finding_performance_writer(); - void write( - const edm::track_candidate_collection::const_view& - track_candidates_view, - const measurement_collection_types::const_view& measurements_view, - const event_data& evt_data); - - void write( - const edm::track_fit_collection::const_view& - track_fit_view, - const edm::track_state_collection::const_view& - track_states_view, - const measurement_collection_types::const_view& measurements_view, - const event_data& evt_data); + void write(const edm::track_candidate_container< + default_algebra>::const_view& track_candidates_view, + const event_data& evt_data); + + void write(const edm::track_fit_container::const_view& + track_fit_view, + const event_data& evt_data); void finalize(); diff --git a/performance/src/efficiency/finding_performance_writer.cpp b/performance/src/efficiency/finding_performance_writer.cpp index f593b62853..6a8d24ddeb 100644 --- a/performance/src/efficiency/finding_performance_writer.cpp +++ b/performance/src/efficiency/finding_performance_writer.cpp @@ -91,28 +91,25 @@ namespace { * with its corresponding measurements. */ std::vector> prepare_data( - const edm::track_candidate_collection::const_view& - track_candidates_view, - const measurement_collection_types::const_view& measurements_view) { + const edm::track_candidate_container::const_view& + track_candidates_view) { std::vector> result; // Iterate over the tracks. - const edm::track_candidate_collection::const_device + const edm::track_candidate_container::const_device track_candidates(track_candidates_view); - const measurement_collection_types::const_device measurements{ - measurements_view}; - const unsigned int n_tracks = track_candidates.size(); + const unsigned int n_tracks = track_candidates.tracks.size(); result.reserve(n_tracks); for (unsigned int i = 0; i < n_tracks; i++) { std::vector m; - m.reserve(track_candidates.at(i).measurement_indices().size()); + m.reserve(track_candidates.tracks.at(i).measurement_indices().size()); const edm::track_candidate_collection< default_algebra>::const_device::const_proxy_type track = - track_candidates.at(i); + track_candidates.tracks.at(i); for (unsigned int midx : track.measurement_indices()) { - m.push_back(measurements.at(midx)); + m.push_back(track_candidates.measurements.at(midx)); } result.push_back(std::move(m)); } @@ -129,31 +126,25 @@ std::vector> prepare_data( * with its corresponding measurements. */ std::vector> prepare_data( - const edm::track_fit_collection::const_view& - track_fit_view, - const edm::track_state_collection::const_view& - track_states_view, - const measurement_collection_types::const_view& measurements_view) { + const edm::track_fit_container::const_view& + track_fit_view) { std::vector> result; // Set up the input containers. - const edm::track_fit_collection::const_device track_fits( + const edm::track_fit_container::const_device track_fits( track_fit_view); - const edm::track_state_collection::const_device - track_states(track_states_view); - const measurement_collection_types::const_device measurements{ - measurements_view}; // Iterate over the tracks. - const unsigned int n_tracks = track_fits.size(); + const unsigned int n_tracks = track_fits.tracks.size(); result.reserve(n_tracks); for (unsigned int i = 0; i < n_tracks; i++) { std::vector result_measurements; - result_measurements.reserve(track_fits.at(i).state_indices().size()); - for (unsigned int st_idx : track_fits.state_indices().at(i)) { - result_measurements.push_back( - measurements.at(track_states.at(st_idx).measurement_index())); + result_measurements.reserve( + track_fits.tracks.at(i).state_indices().size()); + for (unsigned int st_idx : track_fits.tracks.state_indices().at(i)) { + result_measurements.push_back(track_fits.measurements.at( + track_fits.states.at(st_idx).measurement_index())); } result.push_back(std::move(result_measurements)); } @@ -322,14 +313,13 @@ void finding_performance_writer::write_common( /// For track finding void finding_performance_writer::write( - const edm::track_candidate_collection::const_view& + const edm::track_candidate_container::const_view& track_candidates_view, - const measurement_collection_types::const_view& measurements_view, const event_data& evt_data) { // Set up the input containers. const edm::track_candidate_collection::const_device - track_candidates(track_candidates_view); + track_candidates(track_candidates_view.tracks); const unsigned int n_tracks = track_candidates.size(); @@ -341,21 +331,16 @@ void finding_performance_writer::write( } std::vector> tracks = - prepare_data(track_candidates_view, measurements_view); + prepare_data(track_candidates_view); write_common(tracks, evt_data); } /// For ambiguity resolution void finding_performance_writer::write( - const edm::track_fit_collection::const_view& - track_fit_view, - const edm::track_state_collection::const_view& - track_states_view, - const measurement_collection_types::const_view& measurements_view, + const edm::track_fit_container::const_view& track_fit_view, const event_data& evt_data) { - std::vector> tracks = - prepare_data(track_fit_view, track_states_view, measurements_view); + std::vector> tracks = prepare_data(track_fit_view); write_common(tracks, evt_data); } From 70d1cb3892d3d9d1fd94b9097fd3d8c76ccbb045 Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Mon, 11 Aug 2025 15:47:34 +0200 Subject: [PATCH 07/11] Moved "complex EDM operations" to helper functions. --- core/CMakeLists.txt | 4 + .../traccc/edm/impl/measurement_helpers.ipp | 80 +++++++++++++++++++ .../edm/impl/track_state_collection.ipp | 78 ------------------ .../traccc/edm/impl/track_state_helpers.ipp | 39 +++++++++ .../traccc/edm/measurement_helpers.hpp | 46 +++++++++++ .../traccc/edm/track_state_collection.hpp | 36 --------- .../traccc/edm/track_state_helpers.hpp | 32 ++++++++ .../details/combinatorial_kalman_filter.hpp | 10 +-- .../traccc/fitting/details/kalman_fitting.hpp | 12 +-- .../kalman_filter/gain_matrix_updater.hpp | 9 ++- .../kalman_filter/two_filters_smoother.hpp | 11 +-- .../finding/device/impl/find_tracks.ipp | 9 +-- .../traccc/fitting/device/fit_prelude.hpp | 11 +-- 13 files changed, 225 insertions(+), 152 deletions(-) create mode 100644 core/include/traccc/edm/impl/measurement_helpers.ipp create mode 100644 core/include/traccc/edm/impl/track_state_helpers.ipp create mode 100644 core/include/traccc/edm/measurement_helpers.hpp create mode 100644 core/include/traccc/edm/track_state_helpers.hpp diff --git a/core/CMakeLists.txt b/core/CMakeLists.txt index 9aaf0e8f6b..72093bbfec 100644 --- a/core/CMakeLists.txt +++ b/core/CMakeLists.txt @@ -21,6 +21,8 @@ traccc_add_library( traccc_core core TYPE SHARED "include/traccc/edm/details/device_container.hpp" "include/traccc/edm/details/host_container.hpp" "include/traccc/edm/measurement.hpp" + "include/traccc/edm/measurement_helpers.hpp" + "include/traccc/edm/impl/measurement_helpers.ipp" "include/traccc/edm/particle.hpp" "include/traccc/edm/track_parameters.hpp" "include/traccc/edm/container.hpp" @@ -36,6 +38,8 @@ traccc_add_library( traccc_core core TYPE SHARED "include/traccc/edm/track_candidate_container.hpp" "include/traccc/edm/track_state_collection.hpp" "include/traccc/edm/impl/track_state_collection.ipp" + "include/traccc/edm/track_state_helpers.hpp" + "include/traccc/edm/impl/track_state_helpers.ipp" "include/traccc/edm/track_fit_outcome.hpp" "include/traccc/edm/track_fit_collection.hpp" # Magnetic field description. diff --git a/core/include/traccc/edm/impl/measurement_helpers.ipp b/core/include/traccc/edm/impl/measurement_helpers.ipp new file mode 100644 index 0000000000..ca3d30ede6 --- /dev/null +++ b/core/include/traccc/edm/impl/measurement_helpers.ipp @@ -0,0 +1,80 @@ +/** 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 + +namespace traccc::edm { + +template +TRACCC_HOST_DEVICE void get_measurement_local( + const measurement& meas, detray::dmatrix& pos) { + + static_assert(((D == 1u) || (D == 2u)), + "The measurement dimension must be 1 or 2"); + + assert((meas.subs.get_indices()[0] == e_bound_loc0) || + (meas.subs.get_indices()[0] == e_bound_loc1)); + + const point2& local = meas.local; + + switch (meas.subs.get_indices()[0]) { + case e_bound_loc0: + getter::element(pos, 0, 0) = local[0]; + if constexpr (D == 2u) { + getter::element(pos, 1, 0) = local[1]; + } + break; + case e_bound_loc1: + getter::element(pos, 0, 0) = local[1]; + if constexpr (D == 2u) { + getter::element(pos, 1, 0) = local[0]; + } + break; + default: +#if defined(__GNUC__) + __builtin_unreachable(); +#endif + } +} + +template +TRACCC_HOST_DEVICE void get_measurement_covariance( + const measurement& meas, detray::dmatrix& cov) { + + static_assert(((D == 1u) || (D == 2u)), + "The measurement dimension must be 1 or 2"); + + assert((meas.subs.get_indices()[0] == e_bound_loc0) || + (meas.subs.get_indices()[0] == e_bound_loc1)); + + const variance2& variance = meas.variance; + + switch (meas.subs.get_indices()[0]) { + case e_bound_loc0: + getter::element(cov, 0, 0) = variance[0]; + if constexpr (D == 2u) { + getter::element(cov, 0, 1) = 0.f; + getter::element(cov, 1, 0) = 0.f; + getter::element(cov, 1, 1) = variance[1]; + } + break; + case e_bound_loc1: + getter::element(cov, 0, 0) = variance[1]; + if constexpr (D == 2u) { + getter::element(cov, 0, 1) = 0.f; + getter::element(cov, 1, 0) = 0.f; + getter::element(cov, 1, 1) = variance[0]; + } + break; + default: +#if defined(__GNUC__) + __builtin_unreachable(); +#endif + } +} + +} // namespace traccc::edm diff --git a/core/include/traccc/edm/impl/track_state_collection.ipp b/core/include/traccc/edm/impl/track_state_collection.ipp index ab386d7de3..7b66bdc59c 100644 --- a/core/include/traccc/edm/impl/track_state_collection.ipp +++ b/core/include/traccc/edm/impl/track_state_collection.ipp @@ -41,82 +41,4 @@ TRACCC_HOST_DEVICE void track_state::set_smoothed(bool value) { } } -template -template -TRACCC_HOST_DEVICE void track_state::get_measurement_local( - const measurement_collection_types::const_device& measurements, - detray::dmatrix& pos) const { - - static_assert(((D == 1u) || (D == 2u)), - "The measurement dimension must be 1 or 2"); - - assert((measurements.at(measurement_index()).subs.get_indices()[0] == - e_bound_loc0) || - (measurements.at(measurement_index()).subs.get_indices()[0] == - e_bound_loc1)); - - const point2& local = measurements.at(measurement_index()).local; - - switch (measurements.at(measurement_index()).subs.get_indices()[0]) { - case e_bound_loc0: - getter::element(pos, 0, 0) = local[0]; - if constexpr (D == 2u) { - getter::element(pos, 1, 0) = local[1]; - } - break; - case e_bound_loc1: - getter::element(pos, 0, 0) = local[1]; - if constexpr (D == 2u) { - getter::element(pos, 1, 0) = local[0]; - } - break; - default: -#if defined(__GNUC__) - __builtin_unreachable(); -#endif - } -} - -template -template -TRACCC_HOST_DEVICE void track_state::get_measurement_covariance( - const measurement_collection_types::const_device& measurements, - detray::dmatrix& cov) const { - - static_assert(((D == 1u) || (D == 2u)), - "The measurement dimension must be 1 or 2"); - - assert((measurements.at(measurement_index()).subs.get_indices()[0] == - e_bound_loc0) || - (measurements.at(measurement_index()).subs.get_indices()[0] == - e_bound_loc1)); - - const variance2& variance = measurements.at(measurement_index()).variance; - - switch (measurements.at(measurement_index()).subs.get_indices()[0]) { - case e_bound_loc0: - getter::element(cov, 0, 0) = variance[0]; - if constexpr (D == 2u) { - getter::element(cov, 0, 1) = 0.f; - getter::element(cov, 1, 0) = 0.f; - getter::element(cov, 1, 1) = variance[1]; - } - break; - case e_bound_loc1: - getter::element(cov, 0, 0) = variance[1]; - if constexpr (D == 2u) { - getter::element(cov, 0, 1) = 0.f; - getter::element(cov, 1, 0) = 0.f; - getter::element(cov, 1, 1) = variance[0]; - } - break; - default: -#if defined(__GNUC__) - __builtin_unreachable(); -#endif - } -} - } // namespace traccc::edm diff --git a/core/include/traccc/edm/impl/track_state_helpers.ipp b/core/include/traccc/edm/impl/track_state_helpers.ipp new file mode 100644 index 0000000000..7d72ca1ab6 --- /dev/null +++ b/core/include/traccc/edm/impl/track_state_helpers.ipp @@ -0,0 +1,39 @@ +/** 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 + +namespace traccc::edm { + +template +TRACCC_HOST_DEVICE + typename track_state_collection::device::object_type + make_track_state( + const measurement_collection_types::const_device& measurements, + unsigned int mindex) { + + // Create the result object. + typename track_state_collection::device::object_type state{ + track_state_collection::device::object_type::IS_HOLE_MASK, + 0.f, + 0.f, + 0.f, + {}, + {}, + mindex}; + + // Set the correct surface link for the track parameters. + state.filtered_params().set_surface_link( + measurements.at(mindex).surface_link); + state.smoothed_params().set_surface_link( + measurements.at(mindex).surface_link); + + // Return the initialized state. + return state; +} + +} // namespace traccc::edm diff --git a/core/include/traccc/edm/measurement_helpers.hpp b/core/include/traccc/edm/measurement_helpers.hpp new file mode 100644 index 0000000000..b038053a1e --- /dev/null +++ b/core/include/traccc/edm/measurement_helpers.hpp @@ -0,0 +1,46 @@ +/** 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 + +// Local include(s). +#include "traccc/definitions/primitives.hpp" +#include "traccc/definitions/qualifiers.hpp" +#include "traccc/edm/measurement.hpp" + +namespace traccc::edm { + +/// Get the local position of a measurement as a matrix +/// +/// @tparam algebra_t The algebra type used to describe the tracks +/// @tparam size_t The type of the matrix size variable +/// @tparam D The dimension of the matrix +/// +/// @param meas The measurement to extract the local position from +/// @param pos The matrix to fill with the local position of the measurement +/// +template +TRACCC_HOST_DEVICE void get_measurement_local( + const measurement& meas, detray::dmatrix& pos); + +/// Get the covariance of a measurement as a matrix +/// +/// @tparam algebra_t The algebra type used to describe the tracks +/// @tparam size_t The type of the matrix size variable +/// @tparam D The dimension of the matrix +/// +/// @param meas The measurement to extract the covariance from +/// @param cov The matrix to fill with the covariance of the measurement +/// +template +TRACCC_HOST_DEVICE void get_measurement_covariance( + const measurement& meas, detray::dmatrix& cov); + +} // namespace traccc::edm + +// Include the implementation. +#include "traccc/edm/impl/measurement_helpers.ipp" diff --git a/core/include/traccc/edm/track_state_collection.hpp b/core/include/traccc/edm/track_state_collection.hpp index 8019e90d82..fbf61db72c 100644 --- a/core/include/traccc/edm/track_state_collection.hpp +++ b/core/include/traccc/edm/track_state_collection.hpp @@ -191,42 +191,6 @@ class track_state : public BASE { TRACCC_HOST_DEVICE void set_smoothed(bool value = true); - /// Get the local position of the measurement in a matrix - /// - /// @note This function must only be used on proxy objects, not on - /// containers! - /// - /// @tparam size_type The type of the matrix size variable - /// @tparam ALGEBRA The algebra type used to describe the tracks - /// @tparam D The dimension of the matrix - /// - /// @param measurements All measurements in the event - /// @param pos The matrix to fill with the local position of the measurement - /// - template - TRACCC_HOST_DEVICE void get_measurement_local( - const measurement_collection_types::const_device& measurements, - detray::dmatrix& pos) const; - - /// Get the covariance of the measurement in a matrix - /// - /// @note This function must only be used on proxy objects, not on - /// containers! - /// - /// @tparam size_type The type of the matrix size variable - /// @tparam ALGEBRA The algebra type used to describe the tracks - /// @tparam D The dimension of the matrix - /// - /// @param measurements All measurements in the event - /// @param cov The matrix to fill with the covariance of the measurement - /// - template - TRACCC_HOST_DEVICE void get_measurement_covariance( - const measurement_collection_types::const_device& measurements, - detray::dmatrix& cov) const; - /// @} }; // class track_state diff --git a/core/include/traccc/edm/track_state_helpers.hpp b/core/include/traccc/edm/track_state_helpers.hpp new file mode 100644 index 0000000000..15feefcb3d --- /dev/null +++ b/core/include/traccc/edm/track_state_helpers.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 + +// Local include(s). +#include "traccc/edm/track_state_collection.hpp" + +namespace traccc::edm { + +/// Create a track state with default values. +/// +/// @param measurements The collection of measurements to use for initialization +/// @param mindex The index of the measurement to associate with the state +/// +/// @return A track state object initialized with default values +/// +template +TRACCC_HOST_DEVICE + typename track_state_collection::device::object_type + make_track_state( + const measurement_collection_types::const_device& measurements, + unsigned int mindex); + +} // namespace traccc::edm + +// Include the implementation. +#include "traccc/edm/impl/track_state_helpers.ipp" diff --git a/core/include/traccc/finding/details/combinatorial_kalman_filter.hpp b/core/include/traccc/finding/details/combinatorial_kalman_filter.hpp index 7295beb87a..9e4ef30d97 100644 --- a/core/include/traccc/finding/details/combinatorial_kalman_filter.hpp +++ b/core/include/traccc/finding/details/combinatorial_kalman_filter.hpp @@ -9,6 +9,7 @@ // Project include(s). #include "traccc/edm/measurement.hpp" #include "traccc/edm/track_candidate_collection.hpp" +#include "traccc/edm/track_state_helpers.hpp" #include "traccc/finding/actors/ckf_aborter.hpp" #include "traccc/finding/actors/interaction_register.hpp" #include "traccc/finding/candidate_link.hpp" @@ -264,13 +265,8 @@ combinatorial_kalman_filter( const measurement& meas = measurements.at(item_id); // Create a standalone track state object. - typename edm::track_state_collection< - algebra_type>::host::object_type trk_state(0u, 0.f, 0.f, - 0.f, {}, {}, - item_id); - trk_state.set_hole(true); - trk_state.set_smoothed(false); - trk_state.filtered_params().set_surface_link(meas.surface_link); + auto trk_state = + edm::make_track_state(measurements, item_id); const bool is_line = sf.template visit_mask(); diff --git a/core/include/traccc/fitting/details/kalman_fitting.hpp b/core/include/traccc/fitting/details/kalman_fitting.hpp index eb025dd4b8..237428d9c2 100644 --- a/core/include/traccc/fitting/details/kalman_fitting.hpp +++ b/core/include/traccc/fitting/details/kalman_fitting.hpp @@ -13,6 +13,7 @@ #include "traccc/edm/track_fit_collection.hpp" #include "traccc/edm/track_fit_container.hpp" #include "traccc/edm/track_state_collection.hpp" +#include "traccc/edm/track_state_helpers.hpp" #include "traccc/fitting/status_codes.hpp" // VecMem include(s). @@ -69,15 +70,8 @@ typename edm::track_fit_container::host kalman_fitting( track_candidates.measurement_indices().at(i)) { fitted_track.state_indices().push_back( static_cast(result.states.size())); - result.states.push_back( - {0u, 0.f, 0.f, 0.f, {}, {}, measurement_index}); - auto state = result.states.at(result.states.size() - 1); - state.set_hole(true); - state.set_smoothed(false); - state.filtered_params().set_surface_link( - measurements.at(measurement_index).surface_link); - state.smoothed_params().set_surface_link( - measurements.at(measurement_index).surface_link); + result.states.push_back(edm::make_track_state( + measurements, measurement_index)); } vecmem::data::vector_buffer seqs_buffer{ diff --git a/core/include/traccc/fitting/kalman_filter/gain_matrix_updater.hpp b/core/include/traccc/fitting/kalman_filter/gain_matrix_updater.hpp index 87afb6f3e7..c96b2d4f76 100644 --- a/core/include/traccc/fitting/kalman_filter/gain_matrix_updater.hpp +++ b/core/include/traccc/fitting/kalman_filter/gain_matrix_updater.hpp @@ -11,6 +11,7 @@ #include "traccc/definitions/qualifiers.hpp" #include "traccc/definitions/track_parametrization.hpp" #include "traccc/edm/measurement.hpp" +#include "traccc/edm/measurement_helpers.hpp" #include "traccc/edm/track_state_collection.hpp" #include "traccc/fitting/status_codes.hpp" @@ -73,8 +74,8 @@ struct gain_matrix_updater { // Measurement data on surface matrix_type meas_local; - trk_state.template get_measurement_local(measurements, - meas_local); + edm::get_measurement_local( + measurements.at(trk_state.measurement_index()), meas_local); assert((dim > 1) || (getter::element(meas_local, 1u, 0u) == 0.f)); @@ -101,8 +102,8 @@ struct gain_matrix_updater { // Spatial resolution (Measurement covariance) matrix_type V; - trk_state.template get_measurement_covariance(measurements, - V); + edm::get_measurement_covariance( + measurements.at(trk_state.measurement_index()), V); if (dim == 1) { getter::element(V, 1u, 1u) = 1.f; diff --git a/core/include/traccc/fitting/kalman_filter/two_filters_smoother.hpp b/core/include/traccc/fitting/kalman_filter/two_filters_smoother.hpp index 1c9c7b437f..81dcbe44d8 100644 --- a/core/include/traccc/fitting/kalman_filter/two_filters_smoother.hpp +++ b/core/include/traccc/fitting/kalman_filter/two_filters_smoother.hpp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2024 CERN for the benefit of the ACTS project + * (c) 2024-2025 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -11,6 +11,7 @@ #include "traccc/definitions/qualifiers.hpp" #include "traccc/definitions/track_parametrization.hpp" #include "traccc/edm/measurement.hpp" +#include "traccc/edm/measurement_helpers.hpp" #include "traccc/edm/track_state_collection.hpp" #include "traccc/fitting/status_codes.hpp" @@ -71,8 +72,8 @@ struct two_filters_smoother { // Measurement data on surface matrix_type meas_local; - trk_state.template get_measurement_local(measurements, - meas_local); + edm::get_measurement_local( + measurements.at(trk_state.measurement_index()), meas_local); assert((dim > 1) || (getter::element(meas_local, 1u, 0u) == 0.f)); @@ -120,8 +121,8 @@ struct two_filters_smoother { // Spatial resolution (Measurement covariance) matrix_type V; - trk_state.template get_measurement_covariance(measurements, - V); + edm::get_measurement_covariance( + measurements.at(trk_state.measurement_index()), V); if (dim == 1) { getter::element(V, 1u, 1u) = 1.f; } diff --git a/device/common/include/traccc/finding/device/impl/find_tracks.ipp b/device/common/include/traccc/finding/device/impl/find_tracks.ipp index b7146191e2..0700986883 100644 --- a/device/common/include/traccc/finding/device/impl/find_tracks.ipp +++ b/device/common/include/traccc/finding/device/impl/find_tracks.ipp @@ -22,6 +22,7 @@ // Project include(s). #include "traccc/device/array_insertion_mutex.hpp" +#include "traccc/edm/track_state_helpers.hpp" #include "traccc/fitting/kalman_filter/gain_matrix_updater.hpp" #include "traccc/fitting/kalman_filter/is_line_visitor.hpp" #include "traccc/fitting/status_codes.hpp" @@ -219,11 +220,9 @@ TRACCC_HOST_DEVICE inline void find_tracks( if (use_measurement) { - typename edm::track_state_collection< - typename detector_t::algebra_type>::device::object_type - trk_state(0u, 0.f, 0.f, 0.f, {}, {}, meas_idx); - trk_state.filtered_params().set_surface_link( - measurements.at(meas_idx).surface_link); + auto trk_state = + edm::make_track_state( + measurements, meas_idx); const detray::tracking_surface sf{det, in_par.surface_link()}; diff --git a/device/common/include/traccc/fitting/device/fit_prelude.hpp b/device/common/include/traccc/fitting/device/fit_prelude.hpp index d8d9ef7cc0..ba0dc1cd45 100644 --- a/device/common/include/traccc/fitting/device/fit_prelude.hpp +++ b/device/common/include/traccc/fitting/device/fit_prelude.hpp @@ -13,6 +13,7 @@ // Project include(s). #include "traccc/edm/track_candidate_container.hpp" #include "traccc/edm/track_fit_container.hpp" +#include "traccc/edm/track_state_helpers.hpp" #include "traccc/fitting/status_codes.hpp" // VecMem include(s). @@ -54,14 +55,8 @@ TRACCC_HOST_DEVICE inline void fit_prelude( const measurement_collection_types::const_device measurements{ track_candidates_view.measurements}; for (unsigned int meas_idx : track_candidate_measurement_indices) { - const unsigned int track_state_index = - track_states.push_back({0, 0.f, 0.f, 0.f, {}, {}, meas_idx}); - auto state = track_states.at(track_state_index); - state.set_hole(true); - state.set_smoothed(false); - const auto surface_link = measurements.at(meas_idx).surface_link; - state.filtered_params().set_surface_link(surface_link); - state.smoothed_params().set_surface_link(surface_link); + const unsigned int track_state_index = track_states.push_back( + edm::make_track_state(measurements, meas_idx)); track.state_indices().push_back(track_state_index); } From da0241c6fda247d772520e1ff65766be2ea968d9 Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Tue, 12 Aug 2025 08:57:49 +0200 Subject: [PATCH 08/11] Adjusted the host Kalman fit testing code a little. Taking into account that the updated code now returns tracks that have not been successfully fitted. Previously it would not return the tracks for which the fit failed. --- tests/cpu/test_kalman_fitter_momentum_resolution.cpp | 6 ++++++ tests/cpu/test_kalman_fitter_telescope.cpp | 3 +++ tests/cpu/test_kalman_fitter_wire_chamber.cpp | 6 ++++++ 3 files changed, 15 insertions(+) diff --git a/tests/cpu/test_kalman_fitter_momentum_resolution.cpp b/tests/cpu/test_kalman_fitter_momentum_resolution.cpp index 1a18b446af..ea1b2b8c57 100644 --- a/tests/cpu/test_kalman_fitter_momentum_resolution.cpp +++ b/tests/cpu/test_kalman_fitter_momentum_resolution.cpp @@ -184,6 +184,12 @@ TEST_P(KalmanFittingMomentumResolutionTests, Run) { for (std::size_t i_trk = 0; i_trk < n_tracks; i_trk++) { + // Some fits fail. The results of those cannot be reasonably tested. + if (track_states.tracks.at(i_trk).fit_outcome() != + traccc::track_fit_outcome::SUCCESS) { + continue; + } + consistency_tests(track_states.tracks.at(i_trk), track_states.states); diff --git a/tests/cpu/test_kalman_fitter_telescope.cpp b/tests/cpu/test_kalman_fitter_telescope.cpp index c3a8fdb5c0..bd29ae11f6 100644 --- a/tests/cpu/test_kalman_fitter_telescope.cpp +++ b/tests/cpu/test_kalman_fitter_telescope.cpp @@ -153,6 +153,9 @@ TEST_P(KalmanFittingTelescopeTests, Run) { for (std::size_t i_trk = 0; i_trk < n_tracks; i_trk++) { + EXPECT_EQ(track_states.tracks.at(i_trk).fit_outcome(), + traccc::track_fit_outcome::SUCCESS); + consistency_tests(track_states.tracks.at(i_trk), track_states.states); diff --git a/tests/cpu/test_kalman_fitter_wire_chamber.cpp b/tests/cpu/test_kalman_fitter_wire_chamber.cpp index 612501dde2..733b362522 100644 --- a/tests/cpu/test_kalman_fitter_wire_chamber.cpp +++ b/tests/cpu/test_kalman_fitter_wire_chamber.cpp @@ -160,6 +160,12 @@ TEST_P(KalmanFittingWireChamberTests, Run) { for (std::size_t i_trk = 0; i_trk < n_tracks; i_trk++) { + // Some fits fail. The results of those cannot be reasonably tested. + if (track_states.tracks.at(i_trk).fit_outcome() != + traccc::track_fit_outcome::SUCCESS) { + continue; + } + consistency_tests(track_states.tracks.at(i_trk), track_states.states); From ac0386e29b0cde818324966a2b64078a9bcfff7d Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Tue, 12 Aug 2025 09:30:17 +0200 Subject: [PATCH 09/11] Reintroduce print_fitted_tracks_statistics(...). But now as a function in the examples code. Not as part of traccc::core. --- examples/run/CMakeLists.txt | 4 +- .../common/print_fitted_tracks_statistics.cpp | 38 +++++++++++++++++++ .../common/print_fitted_tracks_statistics.hpp | 25 ++++++++++++ .../cpu/misaligned_truth_fitting_example.cpp | 5 ++- examples/run/cpu/truth_finding_example.cpp | 3 +- examples/run/cpu/truth_fitting_example.cpp | 3 +- 6 files changed, 73 insertions(+), 5 deletions(-) create mode 100644 examples/run/common/print_fitted_tracks_statistics.cpp create mode 100644 examples/run/common/print_fitted_tracks_statistics.hpp diff --git a/examples/run/CMakeLists.txt b/examples/run/CMakeLists.txt index 2cd1b613c5..9c92578070 100644 --- a/examples/run/CMakeLists.txt +++ b/examples/run/CMakeLists.txt @@ -1,6 +1,6 @@ # TRACCC library, part of the ACTS project (R&D line) # -# (c) 2021-2023 CERN for the benefit of the ACTS project +# (c) 2021-2025 CERN for the benefit of the ACTS project # # Mozilla Public License Version 2.0 @@ -11,6 +11,8 @@ include( traccc-compiler-options-cpp ) add_library(traccc_examples_common STATIC "common/make_magnetic_field.hpp" "common/make_magnetic_field.cpp" + "common/print_fitted_tracks_statistics.hpp" + "common/print_fitted_tracks_statistics.cpp" "common/throughput_mt.hpp" "common/throughput_mt.ipp" "common/throughput_st.hpp" diff --git a/examples/run/common/print_fitted_tracks_statistics.cpp b/examples/run/common/print_fitted_tracks_statistics.cpp new file mode 100644 index 0000000000..853dc774b0 --- /dev/null +++ b/examples/run/common/print_fitted_tracks_statistics.cpp @@ -0,0 +1,38 @@ +/** 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 "print_fitted_tracks_statistics.hpp" + +namespace traccc::details { + +void print_fitted_tracks_statistics( + const edm::track_fit_container::host& tracks, + const Logger& log) { + + std::size_t success = 0; + std::size_t non_positive_ndf = 0; + std::size_t not_all_smoothed = 0; + + for (track_fit_outcome outcome : tracks.tracks.fit_outcome()) { + if (outcome == track_fit_outcome::SUCCESS) { + ++success; + } else if (outcome == track_fit_outcome::FAILURE_NON_POSITIVE_NDF) { + ++non_positive_ndf; + } else if (outcome == track_fit_outcome::FAILURE_NOT_ALL_SMOOTHED) { + ++not_all_smoothed; + } + } + + auto logger = [&log]() -> const Logger& { return log; }; + TRACCC_INFO("Success: " << success + << " Non positive NDF: " << non_positive_ndf + << " Not all smoothed: " << not_all_smoothed + << " Total: " << tracks.tracks.size()); +} + +} // namespace traccc::details diff --git a/examples/run/common/print_fitted_tracks_statistics.hpp b/examples/run/common/print_fitted_tracks_statistics.hpp new file mode 100644 index 0000000000..6bc959d1a6 --- /dev/null +++ b/examples/run/common/print_fitted_tracks_statistics.hpp @@ -0,0 +1,25 @@ +/** 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 + +// Project include(s). +#include "traccc/edm/track_fit_container.hpp" +#include "traccc/utils/logging.hpp" + +namespace traccc::details { + +/// Print statistics for the fitted tracks. +/// +/// @param tracks The fitted tracks to print statistics for +/// @param log The logger to use for outputting the statistics +/// +void print_fitted_tracks_statistics( + const edm::track_fit_container::host& tracks, + const Logger& log); + +} // namespace traccc::details diff --git a/examples/run/cpu/misaligned_truth_fitting_example.cpp b/examples/run/cpu/misaligned_truth_fitting_example.cpp index 637121be34..475ed61ebb 100644 --- a/examples/run/cpu/misaligned_truth_fitting_example.cpp +++ b/examples/run/cpu/misaligned_truth_fitting_example.cpp @@ -7,6 +7,7 @@ // Project include(s). #include "../common/make_magnetic_field.hpp" +#include "../common/print_fitted_tracks_statistics.hpp" #include "traccc/definitions/common.hpp" #include "traccc/definitions/primitives.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" @@ -166,7 +167,7 @@ int main(int argc, char* argv[]) { {vecmem::get_data(truth_track_candidates.tracks), vecmem::get_data(truth_track_candidates.measurements)}); - // print_fitted_tracks_statistics(track_states); + details::print_fitted_tracks_statistics(track_states, logger()); const std::size_t n_fitted_tracks = track_states.tracks.size(); @@ -191,7 +192,7 @@ int main(int argc, char* argv[]) { {vecmem::get_data(truth_track_candidates.tracks), vecmem::get_data(truth_track_candidates.measurements)}); - // print_fitted_tracks_statistics(track_states); + details::print_fitted_tracks_statistics(track_states, logger()); const std::size_t n_fitted_tracks = track_states.tracks.size(); diff --git a/examples/run/cpu/truth_finding_example.cpp b/examples/run/cpu/truth_finding_example.cpp index 9002b0336d..c257abc600 100644 --- a/examples/run/cpu/truth_finding_example.cpp +++ b/examples/run/cpu/truth_finding_example.cpp @@ -7,6 +7,7 @@ // Project include(s). #include "../common/make_magnetic_field.hpp" +#include "../common/print_fitted_tracks_statistics.hpp" #include "traccc/definitions/common.hpp" #include "traccc/definitions/primitives.hpp" #include "traccc/efficiency/finding_performance_writer.hpp" @@ -163,7 +164,7 @@ int seq_run(const traccc::opts::track_finding& finding_opts, {vecmem::get_data(track_candidates), vecmem::get_data(measurements_per_event)}); - // print_fitted_tracks_statistics(track_states); + details::print_fitted_tracks_statistics(track_states, logger()); const std::size_t n_fitted_tracks = track_states.tracks.size(); diff --git a/examples/run/cpu/truth_fitting_example.cpp b/examples/run/cpu/truth_fitting_example.cpp index 0c626661ba..225862234b 100644 --- a/examples/run/cpu/truth_fitting_example.cpp +++ b/examples/run/cpu/truth_fitting_example.cpp @@ -7,6 +7,7 @@ // Project include(s). #include "../common/make_magnetic_field.hpp" +#include "../common/print_fitted_tracks_statistics.hpp" #include "traccc/definitions/common.hpp" #include "traccc/definitions/primitives.hpp" #include "traccc/fitting/kalman_fitting_algorithm.hpp" @@ -140,7 +141,7 @@ int main(int argc, char* argv[]) { {vecmem::get_data(truth_track_candidates.tracks), vecmem::get_data(truth_track_candidates.measurements)}); - // print_fitted_tracks_statistics(track_states); + details::print_fitted_tracks_statistics(track_states, logger()); const std::size_t n_fitted_tracks = track_states.tracks.size(); From c20b01baa3abc61ec43271f073de8908954a5a1b Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Fri, 15 Aug 2025 13:24:50 +0200 Subject: [PATCH 10/11] Small changes in the host fitting code. Based on the PR feedback. --- .../traccc/fitting/details/kalman_fitting.hpp | 19 ++++++------------- 1 file changed, 6 insertions(+), 13 deletions(-) diff --git a/core/include/traccc/fitting/details/kalman_fitting.hpp b/core/include/traccc/fitting/details/kalman_fitting.hpp index 237428d9c2..21d714289d 100644 --- a/core/include/traccc/fitting/details/kalman_fitting.hpp +++ b/core/include/traccc/fitting/details/kalman_fitting.hpp @@ -87,23 +87,16 @@ typename edm::track_fit_container::host kalman_fitting( auto result_tracks_view = vecmem::get_data(result.tracks); typename edm::track_fit_collection::device result_tracks_device{result_tracks_view}; - typename edm::track_fit_collection::device::proxy_type - fitted_track_device = - result_tracks_device.at(result_tracks_device.size() - 1); - auto result_states_view = vecmem::get_data(result.states); typename fitter_t::state fitter_state( - fitted_track_device, + result_tracks_device.at(result_tracks_device.size() - 1), typename edm::track_state_collection::device{ - result_states_view}, + vecmem::get_data(result.states)}, measurements, seqs_buffer); - // Run the fitter. - kalman_fitter_status fit_status = - fitter.fit(track_candidates.params().at(i), fitter_state); - - if (fit_status != kalman_fitter_status::SUCCESS) { - // TODO: Print a warning here. - } + // Run the fitter. The status that it returns is not used here. The main + // failure modes are saved onto the fitted track itself. Not sure what + // we may want to do with the more detailed status codes in the future. + (void)fitter.fit(track_candidates.params().at(i), fitter_state); } // Return the fitted track states. From ca7c1a01785dcd3ba845764f8344b28a107634a3 Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Fri, 15 Aug 2025 13:25:53 +0200 Subject: [PATCH 11/11] Made the plotting tool documentation more explicit. --- performance/src/resolution/stat_plot_tool.hpp | 23 +++++++++++++------ performance/src/resolution/stat_plot_tool.ipp | 19 ++++++++------- 2 files changed, 27 insertions(+), 15 deletions(-) diff --git a/performance/src/resolution/stat_plot_tool.hpp b/performance/src/resolution/stat_plot_tool.hpp index 806b826cb1..c75ad7c167 100644 --- a/performance/src/resolution/stat_plot_tool.hpp +++ b/performance/src/resolution/stat_plot_tool.hpp @@ -59,25 +59,34 @@ class stat_plot_tool { /// @brief fill the cache /// + /// @tparam track_candidate_backend_t the backend type for @c find_res + /// /// @param cache the cache for statistics plots /// @param find_res track finding result - template - void fill(stat_plot_cache& cache, - const edm::track_candidate& find_res) const; + template + void fill( + stat_plot_cache& cache, + const edm::track_candidate& find_res) const; /// @brief fill the cache /// + /// @tparam track_fit_backend_t the backend type for @c fit_res + /// /// @param cache the cache for statistics plots /// @param fit_res fitting information that contains statistics - template - void fill(stat_plot_cache& cache, const edm::track_fit& fit_res) const; + template + void fill(stat_plot_cache& cache, + const edm::track_fit& fit_res) const; /// @brief fill the cache /// + /// @tparam track_state_backend_t the backend type for @c trk_state + /// /// @param cache the cache for statistics plots /// @param trk_state track state at local measurements - template - void fill(stat_plot_cache& cache, const edm::track_state& trk_state, + template + void fill(stat_plot_cache& cache, + const edm::track_state& trk_state, const measurement_collection_types::host& measurements) const; /// @brief fill the cache diff --git a/performance/src/resolution/stat_plot_tool.ipp b/performance/src/resolution/stat_plot_tool.ipp index ad63a928cc..5471e72956 100644 --- a/performance/src/resolution/stat_plot_tool.ipp +++ b/performance/src/resolution/stat_plot_tool.ipp @@ -12,9 +12,10 @@ namespace traccc { -template -void stat_plot_tool::fill(stat_plot_cache& cache, - const edm::track_candidate& find_res) const { +template +void stat_plot_tool::fill( + stat_plot_cache& cache, + const edm::track_candidate& find_res) const { // Avoid unused variable warnings when building the code without ROOT. (void)cache; @@ -30,9 +31,10 @@ void stat_plot_tool::fill(stat_plot_cache& cache, #endif // TRACCC_HAVE_ROOT } -template -void stat_plot_tool::fill(stat_plot_cache& cache, - const edm::track_fit& fit_res) const { +template +void stat_plot_tool::fill( + stat_plot_cache& cache, + const edm::track_fit& fit_res) const { // Avoid unused variable warnings when building the code without ROOT. (void)cache; @@ -48,9 +50,10 @@ void stat_plot_tool::fill(stat_plot_cache& cache, #endif // TRACCC_HAVE_ROOT } -template +template void stat_plot_tool::fill( - stat_plot_cache& cache, const edm::track_state& trk_state, + stat_plot_cache& cache, + const edm::track_state& trk_state, const measurement_collection_types::host& measurements) const { // Avoid unused variable warnings when building the code without ROOT.