diff --git a/examples/options/include/traccc/options/track_seeding.hpp b/examples/options/include/traccc/options/track_seeding.hpp index fe383ff6b3..2d6fcdcd9b 100644 --- a/examples/options/include/traccc/options/track_seeding.hpp +++ b/examples/options/include/traccc/options/track_seeding.hpp @@ -8,6 +8,7 @@ #pragma once // Project include(s). +#include "traccc/definitions/primitives.hpp" #include "traccc/options/details/config_provider.hpp" #include "traccc/options/details/interface.hpp" #include "traccc/options/details/value_array.hpp" diff --git a/examples/run/CMakeLists.txt b/examples/run/CMakeLists.txt index 2a33464969..c7537c1388 100644 --- a/examples/run/CMakeLists.txt +++ b/examples/run/CMakeLists.txt @@ -9,6 +9,9 @@ include( traccc-compiler-options-cpp ) # Create the common library. add_library(traccc_examples_common STATIC + "common/device_backend.hpp" + "common/device_track_finding_validation.hpp" + "common/device_track_finding_validation.ipp" "common/make_magnetic_field.hpp" "common/make_magnetic_field.cpp" "common/print_fitted_tracks_statistics.hpp" @@ -18,8 +21,8 @@ add_library(traccc_examples_common STATIC "common/throughput_st.hpp" "common/throughput_st.ipp") target_link_libraries(traccc_examples_common - PUBLIC traccc::core traccc::options - PRIVATE traccc::io) + PUBLIC vecmem::core traccc::core traccc::options traccc::io + traccc::performance ) # Add all the subdirectories that can be built. add_subdirectory(cpu) diff --git a/examples/run/alpaka/CMakeLists.txt b/examples/run/alpaka/CMakeLists.txt index 21b62cfb98..d2e8a9755f 100644 --- a/examples/run/alpaka/CMakeLists.txt +++ b/examples/run/alpaka/CMakeLists.txt @@ -21,25 +21,35 @@ endif() set(LIBRARIES vecmem::core traccc::io traccc::performance traccc::core traccc::device_common traccc::alpaka traccc::options traccc_examples_common ${EXTRA_LIBS}) -set(DETRAY detray::io detray::detectors) - -traccc_add_executable( seq_example_alpaka "seq_example_alpaka.cpp" - LINK_LIBRARIES ${LIBRARIES} ${DETRAY} ) -traccc_add_executable( seeding_example_alpaka "seeding_example_alpaka.cpp" - LINK_LIBRARIES ${LIBRARIES} ) # -# Set up the "throughput applications". +# Set up the library used by the applications. # add_library( traccc_examples_alpaka STATIC + "device_backend.hpp" + "device_backend.cpp" "full_chain_algorithm.hpp" "full_chain_algorithm.cpp" ) target_link_libraries( traccc_examples_alpaka PUBLIC vecmem::core detray::core detray::detectors - traccc::core traccc::device_common traccc::alpaka traccc_examples_common ${EXTRA_LIBS}) + traccc::core traccc::device_common traccc::alpaka traccc_examples_common + ${EXTRA_LIBS} ) +# +# Set up the applications. +# traccc_add_executable( throughput_st_alpaka "throughput_st.cpp" - LINK_LIBRARIES indicators::indicators ${LIBRARIES} ${DETRAY} traccc_examples_alpaka ) + LINK_LIBRARIES indicators::indicators traccc_examples_common + traccc_examples_alpaka ) traccc_add_executable( throughput_mt_alpaka "throughput_mt.cpp" - LINK_LIBRARIES TBB::tbb indicators::indicators ${LIBRARIES} ${DETRAY} traccc_examples_alpaka ) + LINK_LIBRARIES TBB::tbb indicators::indicators traccc_examples_common + traccc_examples_alpaka ) + +traccc_add_executable( track_finding_validation_alpaka + "track_finding_validation.cpp" + LINK_LIBRARIES traccc_examples_common traccc_examples_alpaka ) + +traccc_add_executable( reconstruction_validation_alpaka + "reconstruction_validation.cpp" + LINK_LIBRARIES traccc_examples_common traccc_examples_alpaka ) diff --git a/examples/run/alpaka/device_backend.cpp b/examples/run/alpaka/device_backend.cpp new file mode 100644 index 0000000000..be50d315b7 --- /dev/null +++ b/examples/run/alpaka/device_backend.cpp @@ -0,0 +1,156 @@ +/** 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 "device_backend.hpp" + +// Project include(s). +#include "traccc/alpaka/clusterization/clusterization_algorithm.hpp" +#include "traccc/alpaka/clusterization/measurement_sorting_algorithm.hpp" +#include "traccc/alpaka/finding/combinatorial_kalman_filter_algorithm.hpp" +#include "traccc/alpaka/fitting/kalman_fitting_algorithm.hpp" +#include "traccc/alpaka/seeding/seeding_algorithm.hpp" +#include "traccc/alpaka/seeding/spacepoint_formation_algorithm.hpp" +#include "traccc/alpaka/seeding/track_params_estimation.hpp" +#include "traccc/alpaka/utils/queue.hpp" +#include "traccc/alpaka/utils/vecmem_objects.hpp" + +namespace traccc::alpaka { + +struct device_backend::impl { + + /// Alpaka queue to use + queue m_queue; + /// VecMem objects to use + vecmem_objects m_vo{m_queue}; + + /// Traccc memory resource + memory_resource m_mr{m_vo.device_mr(), &(m_vo.host_mr())}; + +}; // struct device_backend::impl + +device_backend::device_backend(std::unique_ptr logger) + : messaging(std::move(logger)), m_impl{std::make_unique()} {} + +device_backend::~device_backend() = default; + +vecmem::copy& device_backend::copy() const { + + return m_impl->m_vo.async_copy(); +} + +memory_resource& device_backend::mr() const { + + return m_impl->m_mr; +} + +void device_backend::synchronize() const { + + m_impl->m_queue.synchronize(); +} + +magnetic_field device_backend::make_magnetic_field(const magnetic_field& bfield, + bool) const { + + return bfield; +} + +std::unique_ptr::buffer( + const edm::silicon_cell_collection::const_view&, + const silicon_detector_description::const_view&)>> +device_backend::make_clusterization_algorithm( + const clustering_config& config) const { + + TRACCC_VERBOSE("Constructing alpaka::clusterization_algorithm"); + return std::make_unique( + m_impl->m_mr, m_impl->m_vo.async_copy(), m_impl->m_queue, config, + logger().clone("alpaka::clusterization_algorithm")); +} + +std::unique_ptr::buffer( + const edm::measurement_collection::const_view&)>> +device_backend::make_measurement_sorting_algorithm() const { + + TRACCC_VERBOSE("Constructing alpaka::measurement_sorting_algorithm"); + return std::make_unique( + m_impl->m_mr, m_impl->m_vo.async_copy(), m_impl->m_queue, + logger().clone("alpaka::measurement_sorting_algorithm")); +} + +std::unique_ptr::const_view&)>> +device_backend::make_spacepoint_formation_algorithm() const { + + TRACCC_VERBOSE("Constructing alpaka::spacepoint_formation_algorithm"); + return std::make_unique( + m_impl->m_mr, m_impl->m_vo.async_copy(), m_impl->m_queue, + logger().clone("alpaka::spacepoint_formation_algorithm")); +} + +std::unique_ptr> +device_backend::make_seeding_algorithm( + const seedfinder_config& finder_config, + const spacepoint_grid_config& grid_config, + const seedfilter_config& filter_config) const { + + TRACCC_VERBOSE("Constructing alpaka::seeding_algorithm"); + return std::make_unique( + finder_config, grid_config, filter_config, m_impl->m_mr, + m_impl->m_vo.async_copy(), m_impl->m_queue, + logger().clone("alpaka::seeding_algorithm")); +} + +std::unique_ptr::const_view&, + const edm::spacepoint_collection::const_view&, + const edm::seed_collection::const_view&, const vector3&)>> +device_backend::make_track_params_estimation_algorithm( + const track_params_estimation_config& config) const { + + TRACCC_VERBOSE("Constructing alpaka::track_params_estimation"); + return std::make_unique( + config, m_impl->m_mr, m_impl->m_vo.async_copy(), m_impl->m_queue, + logger().clone("alpaka::track_params_estimation")); +} + +std::unique_ptr::buffer( + const detector_buffer&, const magnetic_field&, + const edm::measurement_collection::const_view&, + const bound_track_parameters_collection_types::const_view&)>> +device_backend::make_finding_algorithm(const finding_config& config) const { + + TRACCC_VERBOSE( + "Constructing alpaka::combinatorial_kalman_filter_algorithm"); + return std::make_unique( + config, m_impl->m_mr, m_impl->m_vo.async_copy(), m_impl->m_queue, + logger().clone("alpaka::combinatorial_kalman_filter_algorithm")); +} + +std::unique_ptr::buffer( + const edm::track_container::const_view&)>> +device_backend::make_ambiguity_resolution_algorithm( + const ambiguity_resolution_config&) const { + + TRACCC_DEBUG( + "No ambiguity resolution algorithm implemented for the Alpaka backend"); + return {}; +} + +std::unique_ptr::buffer( + const detector_buffer&, const magnetic_field&, + const edm::track_container::const_view&)>> +device_backend::make_fitting_algorithm(const fitting_config& config) const { + + TRACCC_VERBOSE("Constructing alpaka::kalman_fitting_algorithm"); + return std::make_unique( + config, m_impl->m_mr, m_impl->m_vo.async_copy(), m_impl->m_queue, + logger().clone("alpaka::kalman_fitting_algorithm")); +} + +} // namespace traccc::alpaka diff --git a/examples/run/alpaka/device_backend.hpp b/examples/run/alpaka/device_backend.hpp new file mode 100644 index 0000000000..09e5719ba9 --- /dev/null +++ b/examples/run/alpaka/device_backend.hpp @@ -0,0 +1,114 @@ +/** 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 "../common/device_backend.hpp" +#include "traccc/utils/messaging.hpp" + +// System include(s). +#include + +namespace traccc::alpaka { + +/// Alpaka Device Backend +class device_backend : public traccc::device_backend, public messaging { + + public: + /// Constructor + /// + /// @param logger The logger to use + /// + device_backend( + std::unique_ptr logger = getDummyLogger().clone()); + /// Destructor + ~device_backend(); + + /// @name Function(s) implemented from @c traccc::device_backend + /// @{ + + /// Access a copy object for the used device + vecmem::copy& copy() const override; + + /// Get the memory resource(s) used by the algorithms + memory_resource& mr() const override; + + /// Wait for the used device to finish all scheduled operations + void synchronize() const override; + + /// Set up the magnetic field for the device + magnetic_field make_magnetic_field( + const magnetic_field& bfield, + bool texture_memory = false) const override; + + /// Construct a clusterization algorithm instance + std::unique_ptr< + algorithm::buffer( + const edm::silicon_cell_collection::const_view&, + const silicon_detector_description::const_view&)>> + make_clusterization_algorithm( + const clustering_config& config) const override; + + /// Construct a measurement sorting algorithm instance + std::unique_ptr< + algorithm::buffer( + const edm::measurement_collection::const_view&)>> + make_measurement_sorting_algorithm() const override; + + /// Construct a spacepoint formation algorithm instance + std::unique_ptr::const_view&)>> + make_spacepoint_formation_algorithm() const override; + + /// Construct a seeding algorithm instance + std::unique_ptr> + make_seeding_algorithm( + const seedfinder_config& finder_config, + const spacepoint_grid_config& grid_config, + const seedfilter_config& filter_config) const override; + + /// Construct a track parameter estimation algorithm instance + std::unique_ptr::const_view&, + const edm::spacepoint_collection::const_view&, + const edm::seed_collection::const_view&, const vector3&)>> + make_track_params_estimation_algorithm( + const track_params_estimation_config& config) const override; + + /// Construct a track finding algorithm instance + std::unique_ptr::buffer( + const detector_buffer&, const magnetic_field&, + const edm::measurement_collection::const_view&, + const bound_track_parameters_collection_types::const_view&)>> + make_finding_algorithm(const finding_config& config) const override; + + /// Construct an ambiguity resolution algorithm instance + std::unique_ptr::buffer( + const edm::track_container::const_view&)>> + make_ambiguity_resolution_algorithm( + const ambiguity_resolution_config& config) const override; + + /// Construct a track fitting algorithm instance + std::unique_ptr::buffer( + const detector_buffer&, const magnetic_field&, + const edm::track_container::const_view&)>> + make_fitting_algorithm(const fitting_config& config) const override; + + /// @} + + private: + /// Implementation class + struct impl; + /// PIMPL data object + std::unique_ptr m_impl; + +}; // class device_backend + +} // namespace traccc::alpaka diff --git a/examples/run/alpaka/reconstruction_validation.cpp b/examples/run/alpaka/reconstruction_validation.cpp new file mode 100644 index 0000000000..6591f1e621 --- /dev/null +++ b/examples/run/alpaka/reconstruction_validation.cpp @@ -0,0 +1,20 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2021-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Project include(s). +#include "../common/device_reconstruction_validation.hpp" + +// Local include(s). +#include "device_backend.hpp" + +int main(int argc, char* argv[]) { + + return traccc::device_reconstruction_validation< + traccc::alpaka::device_backend>("reconstruction_validation_alpaka", + "Alpaka Reconstruction Validation", + argc, argv); +} diff --git a/examples/run/alpaka/seeding_example_alpaka.cpp b/examples/run/alpaka/seeding_example_alpaka.cpp deleted file mode 100644 index 8b1b0ddbca..0000000000 --- a/examples/run/alpaka/seeding_example_alpaka.cpp +++ /dev/null @@ -1,492 +0,0 @@ -/** 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 - */ - -// Project include(s). -#include "../common/make_magnetic_field.hpp" -#include "traccc/alpaka/finding/combinatorial_kalman_filter_algorithm.hpp" -#include "traccc/alpaka/fitting/kalman_fitting_algorithm.hpp" -#include "traccc/alpaka/seeding/seeding_algorithm.hpp" -#include "traccc/alpaka/seeding/track_params_estimation.hpp" -#include "traccc/alpaka/utils/queue.hpp" -#include "traccc/alpaka/utils/vecmem_objects.hpp" -#include "traccc/definitions/common.hpp" -#include "traccc/device/container_d2h_copy_alg.hpp" -#include "traccc/device/container_h2d_copy_alg.hpp" -#include "traccc/efficiency/finding_performance_writer.hpp" -#include "traccc/efficiency/nseed_performance_writer.hpp" -#include "traccc/efficiency/seeding_performance_writer.hpp" -#include "traccc/efficiency/track_filter.hpp" -#include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" -#include "traccc/fitting/kalman_filter/kalman_fitter.hpp" -#include "traccc/fitting/kalman_fitting_algorithm.hpp" -#include "traccc/geometry/detector.hpp" -#include "traccc/io/read_detector.hpp" -#include "traccc/io/read_detector_description.hpp" -#include "traccc/io/read_measurements.hpp" -#include "traccc/io/read_spacepoints.hpp" -#include "traccc/io/utils.hpp" -#include "traccc/options/accelerator.hpp" -#include "traccc/options/detector.hpp" -#include "traccc/options/input_data.hpp" -#include "traccc/options/performance.hpp" -#include "traccc/options/program_options.hpp" -#include "traccc/options/seed_matching.hpp" -#include "traccc/options/track_finding.hpp" -#include "traccc/options/track_fitting.hpp" -#include "traccc/options/track_matching.hpp" -#include "traccc/options/track_propagation.hpp" -#include "traccc/options/track_seeding.hpp" -#include "traccc/options/truth_finding.hpp" -#include "traccc/performance/collection_comparator.hpp" -#include "traccc/performance/soa_comparator.hpp" -#include "traccc/performance/timer.hpp" -#include "traccc/resolution/fitting_performance_writer.hpp" -#include "traccc/seeding/detail/track_params_estimation_config.hpp" -#include "traccc/seeding/seeding_algorithm.hpp" -#include "traccc/seeding/track_params_estimation.hpp" -#include "traccc/utils/propagation.hpp" - -// System include(s). -#include -#include -#include -#include - -using namespace traccc; - -int seq_run(const traccc::opts::track_seeding& seeding_opts, - const traccc::opts::track_finding& finding_opts, - const traccc::opts::track_propagation& propagation_opts, - const traccc::opts::track_fitting& fitting_opts, - const traccc::opts::input_data& input_opts, - const traccc::opts::detector& detector_opts, - const traccc::opts::magnetic_field& bfield_opts, - const traccc::opts::performance& performance_opts, - const traccc::opts::accelerator& accelerator_opts, - const traccc::opts::truth_finding& truth_finding_opts, - const traccc::opts::seed_matching& seed_matching_opts, - const traccc::opts::track_matching& track_matching_opts, - [[maybe_unused]] std::unique_ptr ilogger) { - TRACCC_LOCAL_LOGGER(std::move(ilogger)); - - // Memory resources used by the application. - traccc::alpaka::queue queue; - traccc::alpaka::vecmem_objects vo(queue); - - vecmem::memory_resource& host_mr = vo.host_mr(); - vecmem::memory_resource& device_mr = vo.device_mr(); - vecmem::memory_resource& mng_mr = vo.shared_mr(); - traccc::memory_resource mr{device_mr, &host_mr}; - - // Performance writer - traccc::seeding_performance_writer sd_performance_writer( - traccc::seeding_performance_writer::config{ - .truth_config = truth_finding_opts, - .seed_truth_config = seed_matching_opts}, - logger().clone("SeedingPerformanceWriter")); - traccc::finding_performance_writer find_performance_writer( - traccc::finding_performance_writer::config{ - .truth_config = truth_finding_opts, - .track_truth_config = track_matching_opts}, - logger().clone("FindingPerformanceWriter")); - traccc::fitting_performance_writer fit_performance_writer( - traccc::fitting_performance_writer::config{}, - logger().clone("FittingPerformanceWriter")); - - traccc::nseed_performance_writer nsd_performance_writer( - "nseed_performance_", - std::make_unique( - 2.7f, 1.f * traccc::unit::GeV), - std::make_unique(0.6f)); - - if (performance_opts.run) { - nsd_performance_writer.initialize(); - } - - // Output stats - uint64_t n_spacepoints = 0; - uint64_t n_seeds = 0; - uint64_t n_seeds_alpaka = 0; - uint64_t n_found_tracks = 0; - uint64_t n_found_tracks_alpaka = 0; - uint64_t n_fitted_tracks = 0; - uint64_t n_fitted_tracks_alpaka = 0; - - /***************************** - * Build a geometry - *****************************/ - - // B field value and its type - const auto field = traccc::details::make_magnetic_field(bfield_opts); - const traccc::vector3 field_vec(seeding_opts); - - // Detector view object - traccc::host_detector host_det; - traccc::io::read_detector(host_det, mng_mr, detector_opts.detector_file, - detector_opts.material_file, - detector_opts.grid_file); - - // Copy objects - vecmem::copy host_copy; - vecmem::copy& copy = vo.copy(); - vecmem::copy& async_copy = vo.async_copy(); - - const traccc::detector_buffer detector_buffer = - traccc::buffer_from_host_detector(host_det, mng_mr, copy); - - // Seeding algorithms - const traccc::seedfinder_config seedfinder_config(seeding_opts); - const traccc::seedfilter_config seedfilter_config(seeding_opts); - const traccc::spacepoint_grid_config spacepoint_grid_config(seeding_opts); - traccc::host::seeding_algorithm sa( - seedfinder_config, spacepoint_grid_config, seedfilter_config, host_mr, - logger().clone("HostSeedingAlg")); - const traccc::track_params_estimation_config track_params_estimation_config; - traccc::host::track_params_estimation tp( - track_params_estimation_config, host_mr, - logger().clone("HostTrackParEstAlg")); - - // Alpaka Algorithms - traccc::alpaka::seeding_algorithm sa_alpaka{ - seedfinder_config, - spacepoint_grid_config, - seedfilter_config, - mr, - async_copy, - queue, - logger().clone("AlpakaSeedingAlg")}; - traccc::alpaka::track_params_estimation tp_alpaka{ - track_params_estimation_config, mr, async_copy, queue, - logger().clone("AlpakaTrackParEstAlg")}; - - // Propagation configuration - detray::propagation::config propagation_config(propagation_opts); - - // Finding algorithm configuration - traccc::finding_config cfg(finding_opts); - cfg.propagation = propagation_config; - - // Finding algorithm object - traccc::host::combinatorial_kalman_filter_algorithm host_finding( - cfg, host_mr, logger().clone("HostFindingAlg")); - traccc::alpaka::combinatorial_kalman_filter_algorithm device_finding( - cfg, mr, copy, queue, logger().clone("AlpakaFindingAlg")); - - // Fitting algorithm object - traccc::fitting_config fit_cfg(fitting_opts); - fit_cfg.propagation = propagation_config; - - traccc::host::kalman_fitting_algorithm host_fitting( - fit_cfg, host_mr, host_copy, logger().clone("HostFittingAlg")); - traccc::alpaka::kalman_fitting_algorithm device_fitting( - fit_cfg, mr, copy, queue, logger().clone("AlpakaFittingAlg")); - - traccc::performance::timing_info elapsedTimes; - - // Loop over events - for (std::size_t event = input_opts.skip; - event < input_opts.events + input_opts.skip; ++event) { - - // Instantiate host containers/collections - traccc::edm::spacepoint_collection::host spacepoints_per_event{host_mr}; - traccc::edm::measurement_collection::host - measurements_per_event{host_mr}; - traccc::host::seeding_algorithm::output_type seeds{host_mr}; - traccc::host::track_params_estimation::output_type params; - traccc::edm::track_container::host - track_candidates{host_mr}; - traccc::edm::track_container::host - track_states{host_mr}; - - traccc::edm::seed_collection::buffer seeds_alpaka_buffer; - traccc::bound_track_parameters_collection_types::buffer - params_alpaka_buffer(0, *mr.host); - - traccc::edm::track_container::buffer - track_candidates_alpaka_buffer; - - traccc::edm::track_container::buffer - track_states_alpaka_buffer; - - { // Start measuring wall time - traccc::performance::timer wall_t("Wall time", elapsedTimes); - - /*----------------- - hit file reading - -----------------*/ - { - traccc::performance::timer t("Hit reading (cpu)", - elapsedTimes); - // Read the hits from the relevant event file - traccc::io::read_spacepoints( - spacepoints_per_event, measurements_per_event, event, - input_opts.directory, - (input_opts.use_acts_geom_source ? &host_det : nullptr), - nullptr, input_opts.format); - - } // stop measuring hit reading timer - - /*---------------------------- - Seeding algorithm - ----------------------------*/ - - // Alpaka - - // Copy the spacepoint data to the device. - traccc::edm::spacepoint_collection::buffer - spacepoints_alpaka_buffer( - static_cast(spacepoints_per_event.size()), - mr.main); - async_copy.setup(spacepoints_alpaka_buffer)->wait(); - async_copy(vecmem::get_data(spacepoints_per_event), - spacepoints_alpaka_buffer) - ->wait(); - - traccc::edm::measurement_collection::buffer - measurements_alpaka_buffer( - static_cast(measurements_per_event.size()), - mr.main); - async_copy.setup(measurements_alpaka_buffer)->wait(); - async_copy(vecmem::get_data(measurements_per_event), - measurements_alpaka_buffer) - ->wait(); - - { - traccc::performance::timer t("Seeding (alpaka)", elapsedTimes); - // Reconstruct the spacepoints into seeds. - seeds_alpaka_buffer = - sa_alpaka(vecmem::get_data(spacepoints_alpaka_buffer)); - queue.synchronize(); - } - - // CPU - - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Seeding (cpu)", elapsedTimes); - seeds = sa(vecmem::get_data(spacepoints_per_event)); - } // stop measuring seeding cpu timer - - /*---------------------------- - Track params estimation - ----------------------------*/ - - // Alpaka - - { - traccc::performance::timer t("Track params (alpaka)", - elapsedTimes); - params_alpaka_buffer = tp_alpaka( - measurements_alpaka_buffer, spacepoints_alpaka_buffer, - seeds_alpaka_buffer, field_vec); - queue.synchronize(); - } // stop measuring track params alpaka timer - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Track params (cpu)", - elapsedTimes); - params = tp(vecmem::get_data(measurements_per_event), - vecmem::get_data(spacepoints_per_event), - vecmem::get_data(seeds), field_vec); - } // stop measuring track params cpu timer - - /*------------------------ - Track Finding with CKF - ------------------------*/ - - { - traccc::performance::timer t("Track finding with CKF (alpaka)", - elapsedTimes); - track_candidates_alpaka_buffer = device_finding( - detector_buffer, field, measurements_alpaka_buffer, - params_alpaka_buffer); - } - - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Track finding with CKF (cpu)", - elapsedTimes); - track_candidates = host_finding( - host_det, field, vecmem::get_data(measurements_per_event), - vecmem::get_data(params)); - } - - /*------------------------ - Track Fitting with KF - ------------------------*/ - - { - traccc::performance::timer t("Track fitting with KF (alpaka)", - elapsedTimes); - - track_states_alpaka_buffer = device_fitting( - detector_buffer, field, track_candidates_alpaka_buffer); - } - - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Track fitting with KF (cpu)", - elapsedTimes); - track_states = host_fitting( - host_det, field, - traccc::edm::track_container< - traccc::default_algebra>::const_data(track_candidates)); - } - - } // Stop measuring wall time - - /*---------------------------------- - compare seeds from cpu and alpaka - ----------------------------------*/ - - // Copy the seeds to the host for comparisons - traccc::edm::seed_collection::host seeds_alpaka{host_mr}; - traccc::bound_track_parameters_collection_types::host params_alpaka{ - &host_mr}; - async_copy(seeds_alpaka_buffer, seeds_alpaka)->wait(); - async_copy(params_alpaka_buffer, params_alpaka)->wait(); - - // Copy track candidates from device to host - traccc::edm::track_collection::host - track_candidates_alpaka{host_mr}; - copy(track_candidates_alpaka_buffer.tracks, track_candidates_alpaka) - ->wait(); - - // Copy track states from device to host - traccc::edm::track_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. - std::cout << "===>>> Event " << event << " <<<===" << std::endl; - - // Compare the seeds made on the host and on the device - traccc::soa_comparator compare_seeds{ - "seeds", traccc::details::comparator_factory< - traccc::edm::seed_collection::const_device:: - const_proxy_type>{ - vecmem::get_data(spacepoints_per_event), - vecmem::get_data(spacepoints_per_event)}}; - compare_seeds(vecmem::get_data(seeds), - vecmem::get_data(seeds_alpaka)); - - // Compare the track parameters made on the host and on the device. - traccc::collection_comparator> - compare_track_parameters{"track parameters"}; - compare_track_parameters(vecmem::get_data(params), - vecmem::get_data(params_alpaka)); - - // Compare the track candidates made on the host and on the - // device - traccc::soa_comparator< - traccc::edm::track_collection> - compare_track_candidates{ - "track candidates", - traccc::details::comparator_factory< - traccc::edm::track_collection:: - const_device::const_proxy_type>{ - vecmem::get_data(measurements_per_event), - vecmem::get_data(measurements_per_event), - {}, - {}}}; - compare_track_candidates(vecmem::get_data(track_candidates.tracks), - vecmem::get_data(track_candidates_alpaka)); - } - - /*---------------- - Statistics - ---------------*/ - - n_spacepoints += spacepoints_per_event.size(); - n_seeds_alpaka += seeds_alpaka.size(); - n_seeds += seeds.size(); - n_found_tracks_alpaka += track_candidates_alpaka.size(); - n_found_tracks += track_candidates.tracks.size(); - n_fitted_tracks_alpaka += track_states_alpaka.tracks.size(); - n_fitted_tracks += track_states.tracks.size(); - - /*------------ - Writer - ------------*/ - - if (performance_opts.run) { - - traccc::event_data evt_data(input_opts.directory, event, host_mr, - input_opts.use_acts_geom_source, - &host_det, input_opts.format, false); - - sd_performance_writer.write( - vecmem::get_data(seeds), - vecmem::get_data(spacepoints_per_event), - vecmem::get_data(measurements_per_event), evt_data); - } - } - - if (performance_opts.run) { - sd_performance_writer.finalize(); - nsd_performance_writer.finalize(); - find_performance_writer.finalize(); - fit_performance_writer.finalize(); - std::cout << nsd_performance_writer.generate_report_str(); - } - - std::cout << "==> Statistics ... " << std::endl; - std::cout << "- read " << n_spacepoints << " spacepoints" << std::endl; - std::cout << "- created (cpu) " << n_seeds << " seeds" << std::endl; - std::cout << "- created (alpaka) " << n_seeds_alpaka << " seeds" - << std::endl; - std::cout << "- created (cpu) " << n_found_tracks << " found tracks" - << std::endl; - std::cout << "- created (alpaka) " << n_found_tracks_alpaka - << " found tracks" << std::endl; - std::cout << "- created (cpu) " << n_fitted_tracks << " fitted tracks" - << std::endl; - std::cout << "- created (alpaka) " << n_fitted_tracks_alpaka - << " fitted tracks" << std::endl; - std::cout << "==>Elapsed times...\n" << elapsedTimes << std::endl; - - return 0; -} - -// The main routine -// -int main(int argc, char* argv[]) { - std::unique_ptr logger = traccc::getDefaultLogger( - "TracccExampleSeedingAlpaka", traccc::Logging::Level::INFO); - - // Program options. - traccc::opts::detector detector_opts; - traccc::opts::magnetic_field bfield_opts; - traccc::opts::input_data input_opts; - traccc::opts::track_seeding seeding_opts; - traccc::opts::track_finding finding_opts; - traccc::opts::track_propagation propagation_opts; - traccc::opts::track_fitting fitting_opts; - traccc::opts::performance performance_opts; - traccc::opts::accelerator accelerator_opts; - traccc::opts::truth_finding truth_finding_opts; - traccc::opts::seed_matching seed_matching_opts; - traccc::opts::track_matching track_matching_opts; - traccc::opts::program_options program_opts{ - "Full Tracking Chain Using Alpaka (without clusterization)", - {detector_opts, bfield_opts, input_opts, seeding_opts, finding_opts, - propagation_opts, fitting_opts, performance_opts, accelerator_opts, - truth_finding_opts, seed_matching_opts, track_matching_opts}, - argc, - argv, - logger->cloneWithSuffix("Options")}; - - // Run the application. - return seq_run(seeding_opts, finding_opts, propagation_opts, fitting_opts, - input_opts, detector_opts, bfield_opts, performance_opts, - accelerator_opts, truth_finding_opts, seed_matching_opts, - track_matching_opts, logger->clone()); -} diff --git a/examples/run/alpaka/seq_example_alpaka.cpp b/examples/run/alpaka/seq_example_alpaka.cpp deleted file mode 100644 index 9c3abad83a..0000000000 --- a/examples/run/alpaka/seq_example_alpaka.cpp +++ /dev/null @@ -1,517 +0,0 @@ -/** 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 - */ - -// Project include(s). -#include "../common/make_magnetic_field.hpp" -#include "traccc/alpaka/clusterization/clusterization_algorithm.hpp" -#include "traccc/alpaka/clusterization/measurement_sorting_algorithm.hpp" -#include "traccc/alpaka/finding/combinatorial_kalman_filter_algorithm.hpp" -#include "traccc/alpaka/fitting/kalman_fitting_algorithm.hpp" -#include "traccc/alpaka/seeding/seeding_algorithm.hpp" -#include "traccc/alpaka/seeding/spacepoint_formation_algorithm.hpp" -#include "traccc/alpaka/seeding/track_params_estimation.hpp" -#include "traccc/alpaka/utils/queue.hpp" -#include "traccc/alpaka/utils/vecmem_objects.hpp" -#include "traccc/clusterization/clusterization_algorithm.hpp" -#include "traccc/device/container_d2h_copy_alg.hpp" -#include "traccc/efficiency/seeding_performance_writer.hpp" -#include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" -#include "traccc/fitting/kalman_filter/kalman_fitter.hpp" -#include "traccc/fitting/kalman_fitting_algorithm.hpp" -#include "traccc/geometry/detector.hpp" -#include "traccc/io/read_cells.hpp" -#include "traccc/io/read_detector.hpp" -#include "traccc/io/read_detector_description.hpp" -#include "traccc/io/utils.hpp" -#include "traccc/options/accelerator.hpp" -#include "traccc/options/clusterization.hpp" -#include "traccc/options/detector.hpp" -#include "traccc/options/input_data.hpp" -#include "traccc/options/performance.hpp" -#include "traccc/options/program_options.hpp" -#include "traccc/options/track_finding.hpp" -#include "traccc/options/track_fitting.hpp" -#include "traccc/options/track_propagation.hpp" -#include "traccc/options/track_seeding.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/seeding/seeding_algorithm.hpp" -#include "traccc/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" -#include "traccc/seeding/track_params_estimation.hpp" -#include "traccc/utils/propagation.hpp" - -// System include(s). -#include -#include -#include -#include - -int seq_run(const traccc::opts::detector& detector_opts, - const traccc::opts::magnetic_field& bfield_opts, - const traccc::opts::input_data& input_opts, - const traccc::opts::clusterization& clusterization_opts, - const traccc::opts::track_seeding& seeding_opts, - const traccc::opts::track_finding& finding_opts, - const traccc::opts::track_propagation& propagation_opts, - const traccc::opts::track_fitting& fitting_opts, - const traccc::opts::performance& performance_opts, - const traccc::opts::accelerator& accelerator_opts, - std::unique_ptr ilogger) { - TRACCC_LOCAL_LOGGER(std::move(ilogger)); - - // Memory resources used by the application. - traccc::alpaka::queue queue; - traccc::alpaka::vecmem_objects vo(queue); - - vecmem::memory_resource& host_mr = vo.host_mr(); - vecmem::memory_resource& device_mr = vo.device_mr(); - traccc::memory_resource mr{device_mr, &host_mr}; - - // Host copy object - vecmem::copy host_copy; - - // Device types used. - vecmem::copy& copy = vo.async_copy(); - - // Construct the detector description object. - traccc::silicon_detector_description::host host_det_descr{host_mr}; - traccc::io::read_detector_description( - host_det_descr, detector_opts.detector_file, - detector_opts.digitization_file, traccc::data_format::json); - traccc::silicon_detector_description::data host_det_descr_data{ - vecmem::get_data(host_det_descr)}; - traccc::silicon_detector_description::buffer device_det_descr{ - static_cast( - host_det_descr.size()), - device_mr}; - copy.setup(device_det_descr)->wait(); - copy(host_det_descr_data, device_det_descr)->wait(); - - // Construct a Detray detector object, if supported by the configuration. - traccc::host_detector host_det; - traccc::io::read_detector(host_det, host_mr, detector_opts.detector_file, - detector_opts.material_file, - detector_opts.grid_file); - - const traccc::detector_buffer detector_buffer = - traccc::buffer_from_host_detector(host_det, device_mr, copy); - - // Output stats - uint64_t n_cells = 0; - uint64_t n_measurements = 0; - uint64_t n_measurements_alpaka = 0; - uint64_t n_spacepoints = 0; - uint64_t n_spacepoints_alpaka = 0; - uint64_t n_seeds = 0; - uint64_t n_seeds_alpaka = 0; - uint64_t n_found_tracks = 0; - uint64_t n_found_tracks_alpaka = 0; - uint64_t n_fitted_tracks = 0; - uint64_t n_fitted_tracks_alpaka = 0; - - // Type definitions - using host_spacepoint_formation_algorithm = - traccc::host::silicon_pixel_spacepoint_formation_algorithm; - using device_spacepoint_formation_algorithm = - traccc::alpaka::spacepoint_formation_algorithm; - - using host_finding_algorithm = - traccc::host::combinatorial_kalman_filter_algorithm; - using device_finding_algorithm = - traccc::alpaka::combinatorial_kalman_filter_algorithm; - - using host_fitting_algorithm = traccc::host::kalman_fitting_algorithm; - using device_fitting_algorithm = traccc::alpaka::kalman_fitting_algorithm; - - // Algorithm configuration(s). - const traccc::seedfinder_config seedfinder_config(seeding_opts); - const traccc::seedfilter_config seedfilter_config(seeding_opts); - const traccc::spacepoint_grid_config spacepoint_grid_config(seeding_opts); - - detray::propagation::config propagation_config(propagation_opts); - - traccc::finding_config finding_cfg(finding_opts); - finding_cfg.propagation = propagation_config; - - traccc::fitting_config fitting_cfg(fitting_opts); - fitting_cfg.propagation = propagation_config; - - // Constant B field for the track finding and fitting - const traccc::vector3 field_vec(seeding_opts); - const auto field = traccc::details::make_magnetic_field(bfield_opts); - - traccc::host::clusterization_algorithm ca( - host_mr, logger().clone("HostClusteringAlg")); - host_spacepoint_formation_algorithm sf( - host_mr, logger().clone("HostSpFormationAlg")); - traccc::host::seeding_algorithm sa( - seedfinder_config, spacepoint_grid_config, seedfilter_config, host_mr, - logger().clone("HostSeedingAlg")); - traccc::track_params_estimation_config track_params_estimation_config; - traccc::host::track_params_estimation tp( - track_params_estimation_config, host_mr, - logger().clone("HostTrackParEstAlg")); - host_finding_algorithm finding_alg(finding_cfg, host_mr, - logger().clone("HostFindingAlg")); - host_fitting_algorithm fitting_alg(fitting_cfg, host_mr, host_copy, - logger().clone("HostFittingAlg")); - - traccc::alpaka::clusterization_algorithm ca_alpaka( - mr, copy, queue, clusterization_opts, - logger().clone("AlpakaClusteringAlg")); - traccc::alpaka::measurement_sorting_algorithm ms_alpaka( - mr, copy, queue, logger().clone("AlpakaMeasSortingAlg")); - device_spacepoint_formation_algorithm sf_alpaka( - mr, copy, queue, logger().clone("AlpakaSpFormationAlg")); - traccc::alpaka::seeding_algorithm sa_alpaka( - seedfinder_config, spacepoint_grid_config, seedfilter_config, mr, copy, - queue, logger().clone("AlpakaSeedingAlg")); - traccc::alpaka::track_params_estimation tp_alpaka( - track_params_estimation_config, mr, copy, queue, - logger().clone("AlpakaTrackParEstAlg")); - device_finding_algorithm finding_alg_alpaka( - finding_cfg, mr, copy, queue, logger().clone("AlpakaFindingAlg")); - device_fitting_algorithm fitting_alg_alpaka( - fitting_cfg, mr, copy, queue, logger().clone("AlpakaFittingAlg")); - - // performance writer - traccc::seeding_performance_writer sd_performance_writer( - traccc::seeding_performance_writer::config{}, - logger().clone("SeedingPerformanceWriter")); - - traccc::performance::timing_info elapsedTimes; - - // Loop over events - for (std::size_t event = input_opts.skip; - event < input_opts.events + input_opts.skip; ++event) { - - // Instantiate host containers/collections - traccc::host::clusterization_algorithm::output_type - measurements_per_event{host_mr}; - host_spacepoint_formation_algorithm::output_type spacepoints_per_event{ - host_mr}; - 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_mr}; - - // Instantiate alpaka containers/collections - traccc::edm::measurement_collection::buffer - measurements_alpaka_buffer; - traccc::edm::spacepoint_collection::buffer spacepoints_alpaka_buffer; - traccc::edm::seed_collection::buffer seeds_alpaka_buffer; - traccc::bound_track_parameters_collection_types::buffer - params_alpaka_buffer(0, *mr.host); - traccc::edm::track_container::buffer - track_candidates_buffer; - traccc::edm::track_container::buffer - track_states_buffer; - - { - traccc::performance::timer wall_t("Wall time", elapsedTimes); - - traccc::edm::silicon_cell_collection::host cells_per_event{host_mr}; - - { - traccc::performance::timer t("File reading (cpu)", - elapsedTimes); - // Read the cells from the relevant event file into host memory. - static constexpr bool DEDUPLICATE = true; - traccc::io::read_cells( - cells_per_event, event, input_opts.directory, - logger().clone(), &host_det_descr, input_opts.format, - DEDUPLICATE, input_opts.use_acts_geom_source); - } // stop measuring file reading timer - - n_cells += cells_per_event.size(); - - // Create device copy of input collections - traccc::edm::silicon_cell_collection::buffer cells_buffer( - static_cast(cells_per_event.size()), mr.main); - copy.setup(cells_buffer)->wait(); - copy(vecmem::get_data(cells_per_event), cells_buffer)->wait(); - - // Alpaka - { - traccc::performance::timer t("Clusterization (alpaka)", - elapsedTimes); - // Reconstruct it into spacepoints on the device. - auto unsorted_measurements = - ca_alpaka(cells_buffer, device_det_descr); - measurements_alpaka_buffer = ms_alpaka(unsorted_measurements); - queue.synchronize(); - } // stop measuring clusterization alpaka timer - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Clusterization (cpu)", - elapsedTimes); - measurements_per_event = - ca(vecmem::get_data(cells_per_event), host_det_descr_data); - } // stop measuring clusterization cpu timer - - // Perform seeding, track finding and fitting only when using a - // Detray geometry. - - // Alpaka - { - traccc::performance::timer t("Spacepoint formation (alpaka)", - elapsedTimes); - spacepoints_alpaka_buffer = - sf_alpaka(detector_buffer, measurements_alpaka_buffer); - queue.synchronize(); - } // stop measuring spacepoint formation alpaka timer - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Spacepoint formation (cpu)", - elapsedTimes); - spacepoints_per_event = - sf(host_det, vecmem::get_data(measurements_per_event)); - } // stop measuring spacepoint formation cpu timer - - // Alpaka - { - traccc::performance::timer t("Seeding (alpaka)", elapsedTimes); - seeds_alpaka_buffer = sa_alpaka(spacepoints_alpaka_buffer); - queue.synchronize(); - } // stop measuring seeding alpaka timer - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Seeding (cpu)", elapsedTimes); - seeds = sa(vecmem::get_data(spacepoints_per_event)); - } // stop measuring seeding cpu timer - - // Alpaka - { - traccc::performance::timer t("Track params (alpaka)", - elapsedTimes); - params_alpaka_buffer = tp_alpaka( - measurements_alpaka_buffer, spacepoints_alpaka_buffer, - seeds_alpaka_buffer, field_vec); - queue.synchronize(); - } // stop measuring track params timer - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Track params (cpu)", - elapsedTimes); - params = tp(vecmem::get_data(measurements_per_event), - vecmem::get_data(spacepoints_per_event), - vecmem::get_data(seeds), field_vec); - } // stop measuring track params cpu timer - - // Alpaka - { - traccc::performance::timer timer{"Track finding (alpaka)", - elapsedTimes}; - track_candidates_buffer = finding_alg_alpaka( - detector_buffer, field, measurements_alpaka_buffer, - params_alpaka_buffer); - } - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer timer{"Track finding (cpu)", - elapsedTimes}; - track_candidates = finding_alg( - host_det, field, vecmem::get_data(measurements_per_event), - vecmem::get_data(params)); - } - - // Alpaka - { - traccc::performance::timer timer{"Track fitting (alpaka)", - elapsedTimes}; - track_states_buffer = fitting_alg_alpaka( - detector_buffer, field, track_candidates_buffer); - } - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer timer{"Track fitting (cpu)", - elapsedTimes}; - track_states = fitting_alg( - host_det, field, - traccc::edm::track_container< - traccc::default_algebra>::const_data(track_candidates)); - } - } // Stop measuring wall time - - /*---------------------------------- - compare cpu and alpaka result - ----------------------------------*/ - - traccc::edm::measurement_collection::host - measurements_per_event_alpaka{host_mr}; - traccc::edm::spacepoint_collection::host spacepoints_per_event_alpaka{ - host_mr}; - traccc::edm::seed_collection::host seeds_alpaka{host_mr}; - traccc::bound_track_parameters_collection_types::host params_alpaka{ - &host_mr}; - traccc::edm::track_collection::host - track_candidates_alpaka{host_mr}; - traccc::edm::track_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.tracks, track_candidates_alpaka)->wait(); - 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) { - - // Show which event we are currently presenting the results for. - TRACCC_INFO("===>>> Event " << event << " <<<==="); - - // Compare the measurements made on the host and on the device. - traccc::soa_comparator< - traccc::edm::measurement_collection> - compare_measurements{"measurements"}; - compare_measurements( - vecmem::get_data(measurements_per_event), - vecmem::get_data(measurements_per_event_alpaka)); - - // Compare the spacepoints made on the host and on the device. - traccc::soa_comparator - compare_spacepoints{"spacepoints"}; - compare_spacepoints(vecmem::get_data(spacepoints_per_event), - vecmem::get_data(spacepoints_per_event_alpaka)); - - // Compare the seeds made on the host and on the device - traccc::soa_comparator compare_seeds{ - "seeds", traccc::details::comparator_factory< - traccc::edm::seed_collection::const_device:: - const_proxy_type>{ - vecmem::get_data(spacepoints_per_event), - vecmem::get_data(spacepoints_per_event_alpaka)}}; - compare_seeds(vecmem::get_data(seeds), - vecmem::get_data(seeds_alpaka)); - - // Compare the track parameters made on the host and on the device. - traccc::collection_comparator> - compare_track_parameters{"track parameters"}; - compare_track_parameters(vecmem::get_data(params), - vecmem::get_data(params_alpaka)); - - // Compare tracks found on the host and on the device. - traccc::soa_comparator< - traccc::edm::track_collection> - compare_track_candidates{ - "track candidates", - traccc::details::comparator_factory< - traccc::edm::track_collection:: - const_device::const_proxy_type>{ - vecmem::get_data(measurements_per_event), - vecmem::get_data(measurements_per_event_alpaka), - {}, - {}}}; - compare_track_candidates(vecmem::get_data(track_candidates.tracks), - vecmem::get_data(track_candidates_alpaka)); - - // Compare tracks fitted on the host and on the device. - traccc::soa_comparator< - traccc::edm::track_collection> - compare_track_fits{ - "track fits", - traccc::details::comparator_factory< - traccc::edm::track_collection:: - 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(); - n_spacepoints += spacepoints_per_event.size(); - n_seeds += seeds.size(); - n_measurements_alpaka += measurements_per_event_alpaka.size(); - n_spacepoints_alpaka += spacepoints_per_event_alpaka.size(); - n_seeds_alpaka += seeds_alpaka.size(); - n_found_tracks += track_candidates.tracks.size(); - n_found_tracks_alpaka += track_candidates_alpaka.size(); - n_fitted_tracks += track_states.tracks.size(); - n_fitted_tracks_alpaka += track_states_alpaka.tracks.size(); - - if (performance_opts.run) { - - traccc::event_data evt_data(input_opts.directory, event, host_mr, - input_opts.use_acts_geom_source, - &host_det, input_opts.format, false); - - sd_performance_writer.write( - vecmem::get_data(seeds), - vecmem::get_data(spacepoints_per_event), - vecmem::get_data(measurements_per_event), evt_data); - } - } - - if (performance_opts.run) { - sd_performance_writer.finalize(); - } - - TRACCC_INFO("==> Statistics ... "); - TRACCC_INFO("- read " << n_cells << " cells"); - TRACCC_INFO("- created (cpu) " << n_measurements << " measurements "); - TRACCC_INFO("- created (alpaka) " << n_measurements_alpaka - << " measurements "); - TRACCC_INFO("- created (cpu) " << n_spacepoints << " spacepoints "); - TRACCC_INFO("- created (alpaka) " << n_spacepoints_alpaka - << " spacepoints "); - - TRACCC_INFO("- created (cpu) " << n_seeds << " seeds"); - TRACCC_INFO("- created (alpaka) " << n_seeds_alpaka << " seeds"); - TRACCC_INFO("- found (cpu) " << n_found_tracks << " tracks"); - TRACCC_INFO("- found (alpaka) " << n_found_tracks_alpaka << " tracks"); - TRACCC_INFO("- fitted (cpu) " << n_fitted_tracks << " tracks"); - TRACCC_INFO("- fitted (alpaka) " << n_fitted_tracks_alpaka << " tracks"); - TRACCC_INFO("==>Elapsed times... " << elapsedTimes); - - return 0; -} - -// The main routine -// -int main(int argc, char* argv[]) { - std::unique_ptr logger = traccc::getDefaultLogger( - "TracccExampleSeqAlpaka", traccc::Logging::Level::INFO); - - // Program options. - traccc::opts::detector detector_opts; - traccc::opts::magnetic_field bfield_opts; - traccc::opts::input_data input_opts; - traccc::opts::clusterization clusterization_opts; - traccc::opts::track_seeding seeding_opts; - traccc::opts::track_finding finding_opts; - traccc::opts::track_propagation propagation_opts; - traccc::opts::track_fitting fitting_opts; - traccc::opts::performance performance_opts; - traccc::opts::accelerator accelerator_opts; - traccc::opts::program_options program_opts{ - "Full Tracking Chain Using Alpaka", - {detector_opts, bfield_opts, input_opts, clusterization_opts, - seeding_opts, finding_opts, propagation_opts, performance_opts, - fitting_opts, accelerator_opts}, - argc, - argv, - logger->cloneWithSuffix("Options")}; - - // Run the application. - return seq_run(detector_opts, bfield_opts, input_opts, clusterization_opts, - seeding_opts, finding_opts, propagation_opts, fitting_opts, - performance_opts, accelerator_opts, logger->clone()); -} diff --git a/examples/run/alpaka/track_finding_validation.cpp b/examples/run/alpaka/track_finding_validation.cpp new file mode 100644 index 0000000000..c7814a30bc --- /dev/null +++ b/examples/run/alpaka/track_finding_validation.cpp @@ -0,0 +1,20 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2021-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Project include(s). +#include "../common/device_track_finding_validation.hpp" + +// Local include(s). +#include "device_backend.hpp" + +int main(int argc, char* argv[]) { + + return traccc::device_track_finding_validation< + traccc::alpaka::device_backend>("track_finding_validation_alpaka", + "Alpaka Track Finding Validation", argc, + argv); +} diff --git a/examples/run/common/device_backend.hpp b/examples/run/common/device_backend.hpp new file mode 100644 index 0000000000..451c61e2cf --- /dev/null +++ b/examples/run/common/device_backend.hpp @@ -0,0 +1,124 @@ +/** 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/ambiguity_resolution/ambiguity_resolution_config.hpp" +#include "traccc/bfield/magnetic_field.hpp" +#include "traccc/clusterization/clustering_config.hpp" +#include "traccc/definitions/primitives.hpp" +#include "traccc/edm/measurement_collection.hpp" +#include "traccc/edm/seed_collection.hpp" +#include "traccc/edm/silicon_cell_collection.hpp" +#include "traccc/edm/spacepoint_collection.hpp" +#include "traccc/edm/track_container.hpp" +#include "traccc/edm/track_parameters.hpp" +#include "traccc/finding/finding_config.hpp" +#include "traccc/fitting/fitting_config.hpp" +#include "traccc/geometry/detector_buffer.hpp" +#include "traccc/geometry/silicon_detector_description.hpp" +#include "traccc/seeding/detail/seeding_config.hpp" +#include "traccc/seeding/detail/track_params_estimation_config.hpp" +#include "traccc/utils/algorithm.hpp" +#include "traccc/utils/memory_resource.hpp" + +// VecMem include(s). +#include + +// System include(s). +#include +#include + +namespace traccc { + +/// Interface for a "device backend" +struct device_backend { + + /// Virtual destructor + virtual ~device_backend() {} + + /// Access a copy object for the used device + virtual vecmem::copy& copy() const = 0; + + /// Get the memory resource(s) used by the algorithms + virtual memory_resource& mr() const = 0; + + /// Wait for the used device to finish all scheduled operations + virtual void synchronize() const = 0; + + /// Set up the magnetic field for the device + virtual magnetic_field make_magnetic_field( + const magnetic_field& bfield, bool texture_memory = false) const = 0; + + /// Construct a clusterization algorithm instance + virtual std::unique_ptr< + algorithm::buffer( + const edm::silicon_cell_collection::const_view&, + const silicon_detector_description::const_view&)>> + make_clusterization_algorithm(const clustering_config& config) const = 0; + + /// Construct a measurement sorting algorithm instance + virtual std::unique_ptr< + algorithm::buffer( + const edm::measurement_collection::const_view&)>> + make_measurement_sorting_algorithm() const = 0; + + /// Construct a spacepoint formation algorithm instance + virtual std::unique_ptr::const_view&)>> + make_spacepoint_formation_algorithm() const = 0; + + /// Construct a seeding algorithm instance + virtual std::unique_ptr> + make_seeding_algorithm(const seedfinder_config& finder_config, + const spacepoint_grid_config& grid_config, + const seedfilter_config& filter_config) const = 0; + + /// Construct a track parameter estimation algorithm instance + virtual std::unique_ptr< + algorithm::const_view&, + const edm::spacepoint_collection::const_view&, + const edm::seed_collection::const_view&, const vector3&)>> + make_track_params_estimation_algorithm( + const track_params_estimation_config& config) const = 0; + + /// Construct a track finding algorithm instance + virtual std::unique_ptr< + algorithm::buffer( + const detector_buffer&, const magnetic_field&, + const edm::measurement_collection::const_view&, + const bound_track_parameters_collection_types::const_view&)>> + make_finding_algorithm(const finding_config& config) const = 0; + + /// Construct an ambiguity resolution algorithm instance + virtual std::unique_ptr< + algorithm::buffer( + const edm::track_container::const_view&)>> + make_ambiguity_resolution_algorithm( + const ambiguity_resolution_config& config) const = 0; + + /// Construct a track fitting algorithm instance + virtual std::unique_ptr< + algorithm::buffer( + const detector_buffer&, const magnetic_field&, + const edm::track_container::const_view&)>> + make_fitting_algorithm(const fitting_config& config) const = 0; + +}; // struct algorithm_maker + +namespace concepts { + +/// Concept specifying a device backend +template +concept device_backend = std::derived_from; + +} // namespace concepts +} // namespace traccc diff --git a/examples/run/common/device_reconstruction_validation.hpp b/examples/run/common/device_reconstruction_validation.hpp new file mode 100644 index 0000000000..721de8ead5 --- /dev/null +++ b/examples/run/common/device_reconstruction_validation.hpp @@ -0,0 +1,40 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2021-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Local include(s). +#include "device_backend.hpp" + +// System include(s). +#include + +namespace traccc { + +/// Helper function implementing a device reconstruction validation application +/// +/// Applications that used to be called "traccc_seq_example_" in the +/// past. +/// +/// @tparam backend_t The device backend type to use +/// +/// @param logger_name The name to use for the logger +/// @param description A description for the application +/// @param argc The @c argc argument coming from @c main(...) +/// @param argv The @c argc argument coming from @c main(...) +/// +/// @return The value to be returned from @c main(...) +/// +template +int device_reconstruction_validation(std::string_view logger_name, + std::string_view description, int argc, + char* argv[]); + +} // namespace traccc + +// Include the implementation. +#include "device_reconstruction_validation.ipp" diff --git a/examples/run/common/device_reconstruction_validation.ipp b/examples/run/common/device_reconstruction_validation.ipp new file mode 100644 index 0000000000..76348a6b2e --- /dev/null +++ b/examples/run/common/device_reconstruction_validation.ipp @@ -0,0 +1,581 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2021-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Local include(s). +#include "make_magnetic_field.hpp" +#include "print_fitted_tracks_statistics.hpp" + +// Core include(s). +#include "traccc/geometry/detector_buffer.hpp" +#include "traccc/geometry/host_detector.hpp" +#include "traccc/utils/logging.hpp" + +// Host algorithm(s). +#include "traccc/ambiguity_resolution/greedy_ambiguity_resolution_algorithm.hpp" +#include "traccc/clusterization/clusterization_algorithm.hpp" +#include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" +#include "traccc/fitting/kalman_fitting_algorithm.hpp" +#include "traccc/seeding/seeding_algorithm.hpp" +#include "traccc/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" +#include "traccc/seeding/track_params_estimation.hpp" + +// Command line option include(s). +#include "traccc/options/accelerator.hpp" +#include "traccc/options/clusterization.hpp" +#include "traccc/options/detector.hpp" +#include "traccc/options/input_data.hpp" +#include "traccc/options/logging.hpp" +#include "traccc/options/magnetic_field.hpp" +#include "traccc/options/program_options.hpp" +#include "traccc/options/seed_matching.hpp" +#include "traccc/options/track_finding.hpp" +#include "traccc/options/track_fitting.hpp" +#include "traccc/options/track_matching.hpp" +#include "traccc/options/track_propagation.hpp" +#include "traccc/options/track_resolution.hpp" +#include "traccc/options/track_seeding.hpp" +#include "traccc/options/truth_finding.hpp" + +// Performance include(s). +#include "traccc/performance/collection_comparator.hpp" +#include "traccc/performance/soa_comparator.hpp" +#include "traccc/performance/timer.hpp" + +// I/O include(s). +#include "traccc/io/read_cells.hpp" +#include "traccc/io/read_detector.hpp" +#include "traccc/io/read_detector_description.hpp" + +// VecMem include(s). +#include +#include + +// System include(s). +#include + +namespace traccc { + +template +int device_reconstruction_validation(std::string_view logger_name, + std::string_view description, int argc, + char* argv[]) { + + // Logger object to use during the command line option reading. + std::unique_ptr prelogger = + getDefaultLogger(std::string{logger_name}, Logging::Level::INFO); + + // Program options. + opts::detector detector_opts; + opts::magnetic_field bfield_opts; + opts::input_data input_opts; + opts::clusterization clusterization_opts; + opts::track_seeding seeding_opts; + opts::track_finding finding_opts; + opts::track_resolution resolution_opts; + opts::track_propagation propagation_opts; + opts::track_fitting fitting_opts; + opts::accelerator accelerator_opts; + opts::truth_finding truth_finding_opts; + opts::seed_matching seed_matching_opts; + opts::track_matching track_matching_opts; + opts::logging logging_opts; + opts::program_options program_opts{ + description, + {detector_opts, bfield_opts, input_opts, clusterization_opts, + seeding_opts, finding_opts, propagation_opts, resolution_opts, + fitting_opts, accelerator_opts, truth_finding_opts, seed_matching_opts, + track_matching_opts, logging_opts}, + argc, + argv, + prelogger->clone()}; + + // The logger to use for the rest of the application. + TRACCC_LOCAL_LOGGER( + prelogger->clone(std::nullopt, Logging::Level(logging_opts))); + + // Create the device backend. + const backend_t backend{logger().clone("device_backend")}; + + // Memory resource for the host algorithm(s). + vecmem::host_memory_resource host_mr; + // Copy object for the host algorithm(s). + vecmem::copy host_copy; + + // Set up the detector description. + silicon_detector_description::host host_det_descr{host_mr}; + io::read_detector_description(host_det_descr, detector_opts.detector_file, + detector_opts.digitization_file, + traccc::data_format::json); + const silicon_detector_description::buffer device_det_descr = + backend.copy().to(vecmem::get_data(host_det_descr), backend.mr().main, + backend.mr().host, + vecmem::copy::type::host_to_device); + + // Set up the magnetic field. + const vector3 bfield_vec(seeding_opts); + const auto host_field = details::make_magnetic_field(bfield_opts); + const auto device_field = backend.make_magnetic_field( + host_field, accelerator_opts.use_gpu_texture_memory); + + // Set up the tracking geometry. + host_detector host_det; + io::read_detector(host_det, host_mr, detector_opts.detector_file, + detector_opts.material_file, detector_opts.grid_file); + const detector_buffer device_det = + buffer_from_host_detector(host_det, backend.mr().main, backend.copy()); + + // Set up the clusterization algorithm(s). + traccc::host::clusterization_algorithm host_clusterization{ + host_mr, logger().clone("host::clusterization_algorithm")}; + auto device_clusterization = + backend.make_clusterization_algorithm(clusterization_opts); + + // Set up the measurement sorting algorithm(s). + auto device_measurement_sorting = + backend.make_measurement_sorting_algorithm(); + + // Set up the spacepoint formation algorithm(s). + traccc::host::silicon_pixel_spacepoint_formation_algorithm + host_spacepoint_formation{ + host_mr, logger().clone( + "host::silicon_pixel_spacepoint_formation_algorithm")}; + auto device_spacepoint_formation = + backend.make_spacepoint_formation_algorithm(); + + // Set up the seeding algorithm(s). + const seedfinder_config sfinder_config(seeding_opts); + const seedfilter_config sfilter_config(seeding_opts); + const spacepoint_grid_config sg_config(seeding_opts); + + host::seeding_algorithm host_seeding( + sfinder_config, sg_config, sfilter_config, host_mr, + logger().clone("host::seeding_algorithm")); + auto device_seeding = backend.make_seeding_algorithm( + sfinder_config, sg_config, sfilter_config); + + // Set up the track parameter estimation algorithm(s). + const track_params_estimation_config tp_config; + + host::track_params_estimation host_tp_estimation( + tp_config, host_mr, logger().clone("host::track_params_estimation")); + auto device_tp_estimation = + backend.make_track_params_estimation_algorithm(tp_config); + + // Set up the track finding algorithm(s). + const detray::propagation::config prop_config(propagation_opts); + finding_config find_config(finding_opts); + find_config.propagation = prop_config; + + host::combinatorial_kalman_filter_algorithm host_finding( + find_config, host_mr, + logger().clone("host::combinatorial_kalman_filter_algorithm")); + auto device_finding = backend.make_finding_algorithm(find_config); + + // Set up the ambiguity resolution algorithm(s). + host::greedy_ambiguity_resolution_algorithm host_ambiguity_resolution( + resolution_opts, host_mr, + logger().clone("host::greedy_ambiguity_resolution_algorithm")); + auto device_ambiguity_resolution = + backend.make_ambiguity_resolution_algorithm(resolution_opts); + + // Set up the track fitting algorithm(s). + fitting_config fit_config(fitting_opts); + fit_config.propagation = prop_config; + + host::kalman_fitting_algorithm host_fitting( + fit_config, host_mr, host_copy, + logger().clone("host::kalman_fitting_algorithm")); + auto device_fitting = backend.make_fitting_algorithm(fit_config); + + // Counters for various reconstructed objects. + std::size_t n_cells = 0; + std::size_t n_host_measurements = 0; + std::size_t n_device_measurements = 0; + std::size_t n_host_spacepoints = 0; + std::size_t n_device_spacepoints = 0; + std::size_t n_host_seeds = 0; + std::size_t n_device_seeds = 0; + std::size_t n_host_found_tracks = 0; + std::size_t n_device_found_tracks = 0; + std::size_t n_host_resolved_tracks = 0; + std::size_t n_device_resolved_tracks = 0; + std::size_t n_host_fitted_tracks = 0; + std::size_t n_device_fitted_tracks = 0; + + // Times elapsed in the various reconstruction steps. + performance::timing_info times; + + // Process the requested number of events. + for (std::size_t event = input_opts.skip; + event < input_opts.events + input_opts.skip; ++event) { + + // Instantiate host containers/collections. + edm::silicon_cell_collection::host host_cells{host_mr}; + edm::measurement_collection::host host_measurements{ + host_mr}; + edm::spacepoint_collection::host host_spacepoints{host_mr}; + edm::seed_collection::host host_seeds{host_mr}; + bound_track_parameters_collection_types::host host_track_params{ + &host_mr}; + edm::track_container::host host_found_tracks{host_mr}; + edm::track_container::host host_resolved_tracks{ + host_mr}; + edm::track_container::host host_fitted_tracks{host_mr}; + + // Instantiate device containers/collections. + edm::silicon_cell_collection::buffer device_cells; + edm::measurement_collection::buffer + device_measurements; + edm::spacepoint_collection::buffer device_spacepoints; + edm::seed_collection::buffer device_seeds; + bound_track_parameters_collection_types::buffer device_track_params; + edm::track_container::buffer device_found_tracks; + edm::track_container::buffer device_resolved_tracks; + edm::track_container::buffer device_fitted_tracks; + + { + // Measure the total wall time. + performance::timer wall_t("Wall time", times); + + { + // Read the spacepoints and measurements from the relevant event + // files. + performance::timer t("Host data reading", times); + static constexpr bool DEDUPLICATE = true; + io::read_cells(host_cells, event, input_opts.directory, + logger().clone("io::read_cells"), + &host_det_descr, input_opts.format, DEDUPLICATE, + input_opts.use_acts_geom_source); + } + + { + // Copy the cell data to the device. + performance::timer t{"Host->Device data transfers", times}; + device_cells = backend.copy().to( + vecmem::get_data(host_cells), backend.mr().main, + backend.mr().host, vecmem::copy::type::host_to_device); + } + + { + // Reconstruct the measurements on the device. + performance::timer t("Device clusterization", times); + auto unsorted_device_measurements = + (*device_clusterization)(device_cells, device_det_descr); + device_measurements = + (*device_measurement_sorting)(unsorted_device_measurements); + backend.synchronize(); + } + + if (accelerator_opts.compare_with_cpu) { + // Reconstruct the measurements on the host. + performance::timer t("Host clusterization", times); + host_measurements = + host_clusterization(vecmem::get_data(host_cells), + vecmem::get_data(host_det_descr)); + } + + { + // Reconstruct the spacepoints on the device. + performance::timer t("Device spacepoint formation", times); + device_spacepoints = (*device_spacepoint_formation)( + device_det, device_measurements); + backend.synchronize(); + } + + if (accelerator_opts.compare_with_cpu) { + // Reconstruct the spacepoints on the host. + performance::timer t("Host spacepoint formation", times); + host_spacepoints = host_spacepoint_formation( + host_det, vecmem::get_data(host_measurements)); + } + + { + // Reconstruct the spacepoints into seeds on the device. + performance::timer t("Device seeding", times); + device_seeds = (*device_seeding)(device_spacepoints); + backend.synchronize(); + } + + if (accelerator_opts.compare_with_cpu) { + // Reconstruct the spacepoints into seeds on the host. + performance::timer t("Host seeding", times); + host_seeds = host_seeding(vecmem::get_data(host_spacepoints)); + } + + { + // Run track parameter estimation on the device. + performance::timer t("Device T/P estimation", times); + device_track_params = (*device_tp_estimation)( + device_measurements, device_spacepoints, device_seeds, + bfield_vec); + backend.synchronize(); + } + + if (accelerator_opts.compare_with_cpu) { + // Run track parameter estimation on the host. + performance::timer t("Host T/P esimation", times); + host_track_params = host_tp_estimation( + vecmem::get_data(host_measurements), + vecmem::get_data(host_spacepoints), + vecmem::get_data(host_seeds), bfield_vec); + } + + { + // Run track finding on the device. + performance::timer t("Device track finding", times); + device_found_tracks = + (*device_finding)(device_det, device_field, + device_measurements, device_track_params); + } + + if (accelerator_opts.compare_with_cpu) { + // Run track finding on the host. + traccc::performance::timer t("Host track finding", times); + host_found_tracks = host_finding( + host_det, host_field, vecmem::get_data(host_measurements), + vecmem::get_data(host_track_params)); + } + + if (device_ambiguity_resolution) { + { + // Run ambiguity resolution on the device. + performance::timer t("Device ambiguity resolution", times); + device_resolved_tracks = + (*device_ambiguity_resolution)(device_found_tracks); + } + + if (accelerator_opts.compare_with_cpu) { + // Run ambiguity resolution on the host. + traccc::performance::timer t("Host ambiguity resolution", + times); + host_resolved_tracks = host_ambiguity_resolution( + edm::track_container::const_data( + host_found_tracks)); + } + + { + // Run track fitting on the device. + performance::timer t("Device track fitting", times); + device_fitted_tracks = (*device_fitting)( + device_det, device_field, device_resolved_tracks); + } + + if (accelerator_opts.compare_with_cpu) { + // Run track fitting on the host. + traccc::performance::timer t("Host track fitting", times); + host_fitted_tracks = host_fitting( + host_det, host_field, + edm::track_container::const_data( + host_resolved_tracks)); + } + } else { + { + // Run track fitting on the device. + performance::timer t("Device track fitting", times); + device_fitted_tracks = (*device_fitting)( + device_det, device_field, device_found_tracks); + } + + if (accelerator_opts.compare_with_cpu) { + // Run track fitting on the host. + traccc::performance::timer t("Host track fitting", times); + host_fitted_tracks = host_fitting( + host_det, host_field, + edm::track_container::const_data( + host_found_tracks)); + } + } + } + + // Copy device containers/collections back to the host for validation. + edm::measurement_collection::host + device_host_measurements{host_mr}; + backend.copy()(device_measurements, device_host_measurements)->wait(); + + edm::spacepoint_collection::host device_host_spacepoints{host_mr}; + backend.copy()(device_spacepoints, device_host_spacepoints)->wait(); + + edm::seed_collection::host device_host_seeds{host_mr}; + backend.copy()(device_seeds, device_host_seeds)->wait(); + + bound_track_parameters_collection_types::host device_host_track_params{ + &host_mr}; + backend.copy()(device_track_params, device_host_track_params)->wait(); + + edm::track_container::host + device_host_found_tracks{ + host_mr, vecmem::get_data(device_host_measurements)}; + backend + .copy()(device_found_tracks.tracks, device_host_found_tracks.tracks) + ->wait(); + backend + .copy()(device_found_tracks.states, device_host_found_tracks.states) + ->wait(); + + edm::track_container::host + device_host_resolved_tracks{ + host_mr, vecmem::get_data(device_host_measurements)}; + if (device_ambiguity_resolution) { + backend + .copy()(device_resolved_tracks.tracks, + device_host_resolved_tracks.tracks) + ->wait(); + backend + .copy()(device_resolved_tracks.states, + device_host_resolved_tracks.states) + ->wait(); + } + + edm::track_container::host + device_host_fitted_tracks{ + host_mr, vecmem::get_data(device_host_measurements)}; + backend + .copy()(device_fitted_tracks.tracks, + device_host_fitted_tracks.tracks) + ->wait(); + backend + .copy()(device_fitted_tracks.states, + device_host_fitted_tracks.states) + ->wait(); + + if (accelerator_opts.compare_with_cpu) { + // Show which event we are currently presenting the results for. + TRACCC_INFO("===>>> Event " << event << " <<<==="); + + // Compare the measurements made on the host and on the device. + soa_comparator> + compare_measurements{"measurements"}; + compare_measurements(vecmem::get_data(host_measurements), + vecmem::get_data(device_host_measurements)); + + // Compare the spacepoints made on the host and on the device. + soa_comparator compare_spacepoints{ + "spacepoints"}; + compare_spacepoints(vecmem::get_data(host_spacepoints), + vecmem::get_data(device_host_spacepoints)); + + // Compare the seeds made on the host and on the device + soa_comparator compare_seeds{ + "seeds", + details::comparator_factory< + edm::seed_collection::const_device::const_proxy_type>{ + vecmem::get_data(host_spacepoints), + vecmem::get_data(device_host_spacepoints)}}; + compare_seeds(vecmem::get_data(host_seeds), + vecmem::get_data(device_host_seeds)); + + // Compare the track parameters made on the host and on the device. + collection_comparator> + compare_track_parameters{"track parameters"}; + compare_track_parameters( + vecmem::get_data(host_track_params), + vecmem::get_data(device_host_track_params)); + + // Compare the found tracks made on the host and on the device. + soa_comparator> + compare_found_tracks{ + "found tracks", + details::comparator_factory::const_device::const_proxy_type>{ + host_found_tracks.measurements, + device_host_found_tracks.measurements, + vecmem::get_data(host_found_tracks.states), + vecmem::get_data(device_host_found_tracks.states)}}; + compare_found_tracks( + vecmem::get_data(host_found_tracks.tracks), + vecmem::get_data(device_host_found_tracks.tracks)); + + // Compare the resolved tracks made on the host and on the device. + soa_comparator> + compare_resolved_tracks{ + "recolved tracks", + details::comparator_factory::const_device::const_proxy_type>{ + host_resolved_tracks.measurements, + device_host_resolved_tracks.measurements, + vecmem::get_data(host_resolved_tracks.states), + vecmem::get_data(device_host_resolved_tracks.states)}}; + compare_resolved_tracks( + vecmem::get_data(host_resolved_tracks.tracks), + vecmem::get_data(device_host_resolved_tracks.tracks)); + + // Compare the fitted tracks made on the host and on the device. + soa_comparator> + compare_fitted_tracks{ + "fitted tracks", + details::comparator_factory::const_device::const_proxy_type>{ + host_found_tracks.measurements, + device_host_found_tracks.measurements, + vecmem::get_data(host_fitted_tracks.states), + vecmem::get_data(device_host_fitted_tracks.states)}}; + compare_fitted_tracks( + vecmem::get_data(host_fitted_tracks.tracks), + vecmem::get_data(device_host_fitted_tracks.tracks)); + } + + // Print information about the fitted tracks. + details::print_fitted_tracks_statistics(device_host_fitted_tracks, + logger()); + + // Collect overall statistics. + n_cells += host_cells.size(); + n_host_measurements += host_measurements.size(); + n_device_measurements += device_host_measurements.size(); + n_host_spacepoints += host_spacepoints.size(); + n_device_spacepoints += device_host_spacepoints.size(); + n_host_seeds += host_seeds.size(); + n_device_seeds += device_host_seeds.size(); + n_host_found_tracks += host_found_tracks.tracks.size(); + n_device_found_tracks += device_host_found_tracks.tracks.size(); + n_host_resolved_tracks += host_resolved_tracks.tracks.size(); + n_device_resolved_tracks += device_host_resolved_tracks.tracks.size(); + n_host_fitted_tracks += host_fitted_tracks.tracks.size(); + n_device_fitted_tracks += device_host_fitted_tracks.tracks.size(); + } + + // Print some final statistics about the job. + TRACCC_INFO("===>>> Statistics <<<==="); + TRACCC_INFO(" Procssed cells: " << n_cells); + if (accelerator_opts.compare_with_cpu) { + TRACCC_INFO( + " Reconstructed measurements (host): " << n_host_measurements); + } + TRACCC_INFO( + " Reconstructed measurements (device): " << n_device_measurements); + if (accelerator_opts.compare_with_cpu) { + TRACCC_INFO( + " Reconstructed spacepoints (host): " << n_host_spacepoints); + } + TRACCC_INFO( + " Reconstructed spacepoints (device): " << n_device_spacepoints); + if (accelerator_opts.compare_with_cpu) { + TRACCC_INFO(" Found seeds (host): " << n_host_seeds); + } + TRACCC_INFO(" Found seeds (device): " << n_device_seeds); + if (accelerator_opts.compare_with_cpu) { + TRACCC_INFO(" Found tracks (host): " << n_host_found_tracks); + } + TRACCC_INFO(" Found tracks (device): " << n_device_found_tracks); + if (accelerator_opts.compare_with_cpu) { + TRACCC_INFO(" Resolved tracks (host): " << n_host_resolved_tracks); + } + TRACCC_INFO(" Resolved tracks (device): " << n_device_resolved_tracks); + if (accelerator_opts.compare_with_cpu) { + TRACCC_INFO(" Fitted tracks (host): " << n_host_fitted_tracks); + } + TRACCC_INFO(" Fitted tracks (device): " << n_device_fitted_tracks); + TRACCC_INFO("===>>> Timing information <<<==="); + TRACCC_INFO(times); + + // Return gracefully. + return EXIT_SUCCESS; +} + +} // namespace traccc diff --git a/examples/run/common/device_track_finding_validation.hpp b/examples/run/common/device_track_finding_validation.hpp new file mode 100644 index 0000000000..cd1d155609 --- /dev/null +++ b/examples/run/common/device_track_finding_validation.hpp @@ -0,0 +1,40 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2021-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Local include(s). +#include "device_backend.hpp" + +// System include(s). +#include + +namespace traccc { + +/// Helper function implementing a device track finding validation application +/// +/// Applications that used to be called "traccc_seeding_example_" in +/// the past. +/// +/// @tparam backend_t The device backend type to use +/// +/// @param logger_name The name to use for the logger +/// @param description A description for the application +/// @param argc The @c argc argument coming from @c main(...) +/// @param argv The @c argc argument coming from @c main(...) +/// +/// @return The value to be returned from @c main(...) +/// +template +int device_track_finding_validation(std::string_view logger_name, + std::string_view description, int argc, + char* argv[]); + +} // namespace traccc + +// Include the implementation. +#include "device_track_finding_validation.ipp" diff --git a/examples/run/common/device_track_finding_validation.ipp b/examples/run/common/device_track_finding_validation.ipp new file mode 100644 index 0000000000..755dc3eb28 --- /dev/null +++ b/examples/run/common/device_track_finding_validation.ipp @@ -0,0 +1,464 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2021-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Local include(s). +#include "make_magnetic_field.hpp" +#include "print_fitted_tracks_statistics.hpp" + +// Core include(s). +#include "traccc/geometry/detector_buffer.hpp" +#include "traccc/geometry/host_detector.hpp" +#include "traccc/utils/logging.hpp" + +// Host algorithm(s). +#include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" +#include "traccc/fitting/kalman_fitting_algorithm.hpp" +#include "traccc/seeding/seeding_algorithm.hpp" +#include "traccc/seeding/track_params_estimation.hpp" + +// Command line option include(s). +#include "traccc/options/accelerator.hpp" +#include "traccc/options/detector.hpp" +#include "traccc/options/input_data.hpp" +#include "traccc/options/logging.hpp" +#include "traccc/options/magnetic_field.hpp" +#include "traccc/options/performance.hpp" +#include "traccc/options/program_options.hpp" +#include "traccc/options/seed_matching.hpp" +#include "traccc/options/track_finding.hpp" +#include "traccc/options/track_fitting.hpp" +#include "traccc/options/track_matching.hpp" +#include "traccc/options/track_propagation.hpp" +#include "traccc/options/track_seeding.hpp" +#include "traccc/options/truth_finding.hpp" + +// Performance include(s). +#include "traccc/efficiency/finding_performance_writer.hpp" +#include "traccc/efficiency/seeding_performance_writer.hpp" +#include "traccc/performance/collection_comparator.hpp" +#include "traccc/performance/soa_comparator.hpp" +#include "traccc/performance/timer.hpp" +#include "traccc/resolution/fitting_performance_writer.hpp" + +// I/O include(s). +#include "traccc/io/read_detector.hpp" +#include "traccc/io/read_detector_description.hpp" +#include "traccc/io/read_spacepoints.hpp" + +// VecMem include(s). +#include +#include + +// System include(s). +#include + +namespace traccc { + +template +int device_track_finding_validation(std::string_view logger_name, + std::string_view description, int argc, + char* argv[]) { + + // Logger object to use during the command line option reading. + std::unique_ptr prelogger = + getDefaultLogger(std::string{logger_name}, Logging::Level::INFO); + + // Program options. + opts::detector detector_opts; + opts::magnetic_field bfield_opts; + opts::input_data input_opts; + opts::track_seeding seeding_opts; + opts::track_finding finding_opts; + opts::track_propagation propagation_opts; + opts::track_fitting fitting_opts; + opts::performance performance_opts; + opts::accelerator accelerator_opts; + opts::truth_finding truth_finding_opts; + opts::seed_matching seed_matching_opts; + opts::track_matching track_matching_opts; + opts::logging logging_opts; + opts::program_options program_opts{ + description, + {detector_opts, bfield_opts, input_opts, seeding_opts, finding_opts, + propagation_opts, fitting_opts, performance_opts, accelerator_opts, + truth_finding_opts, seed_matching_opts, track_matching_opts, + logging_opts}, + argc, + argv, + prelogger->clone()}; + + // The logger to use for the rest of the application. + TRACCC_LOCAL_LOGGER( + prelogger->clone(std::nullopt, Logging::Level(logging_opts))); + + // Create the device backend. + const backend_t backend{logger().clone("device_backend")}; + + // Performance writer + seeding_performance_writer seeding_pw( + {.truth_config = truth_finding_opts, + .seed_truth_config = seed_matching_opts}, + logger().clone("seeding_performance_writer")); + finding_performance_writer finding_pw( + {.truth_config = truth_finding_opts, + .track_truth_config = track_matching_opts}, + logger().clone("finding_performance_writer")); + finding_performance_writer postfit_finding_pw( + {.file_path = "performance_track_postfit_finding.root", + .truth_config = truth_finding_opts, + .track_truth_config = track_matching_opts, + .require_fit = true}, + logger().clone("post_fit_finding_performance_writer")); + fitting_performance_writer fitting_pw( + {}, logger().clone("fitting_performance_writer")); + + // Memory resource for the host algorithm(s). + vecmem::host_memory_resource host_mr; + // Copy object for the host algorithm(s). + vecmem::copy host_copy; + + // Set up the detector description. + silicon_detector_description::host det_descr{host_mr}; + io::read_detector_description(det_descr, detector_opts.detector_file, + detector_opts.digitization_file, + traccc::data_format::json); + + // Set up the magnetic field. + const vector3 bfield_vec(seeding_opts); + const auto host_field = details::make_magnetic_field(bfield_opts); + const auto device_field = backend.make_magnetic_field( + host_field, accelerator_opts.use_gpu_texture_memory); + + // Set up the tracking geometry. + host_detector host_det; + io::read_detector(host_det, host_mr, detector_opts.detector_file, + detector_opts.material_file, detector_opts.grid_file); + const detector_buffer device_det = + buffer_from_host_detector(host_det, backend.mr().main, backend.copy()); + + // Set up the seeding algorithm(s). + const seedfinder_config sfinder_config(seeding_opts); + const seedfilter_config sfilter_config(seeding_opts); + const spacepoint_grid_config sg_config(seeding_opts); + + host::seeding_algorithm host_seeding( + sfinder_config, sg_config, sfilter_config, host_mr, + logger().clone("host::seeding_algorithm")); + auto device_seeding = backend.make_seeding_algorithm( + sfinder_config, sg_config, sfilter_config); + + // Set up the track parameter estimation algorithm(s). + const track_params_estimation_config tp_config; + + host::track_params_estimation host_tp_estimation( + tp_config, host_mr, logger().clone("host::track_params_estimation")); + auto device_tp_estimation = + backend.make_track_params_estimation_algorithm(tp_config); + + // Set up the track finding algorithm(s). + const detray::propagation::config prop_config(propagation_opts); + finding_config find_config(finding_opts); + find_config.propagation = prop_config; + + host::combinatorial_kalman_filter_algorithm host_finding( + find_config, host_mr, + logger().clone("host::combinatorial_kalman_filter_algorithm")); + auto device_finding = backend.make_finding_algorithm(find_config); + + // Set up the track fitting algorithm(s). + fitting_config fit_config(fitting_opts); + fit_config.propagation = prop_config; + + host::kalman_fitting_algorithm host_fitting( + fit_config, host_mr, host_copy, + logger().clone("host::kalman_fitting_algorithm")); + auto device_fitting = backend.make_fitting_algorithm(fit_config); + + // Counters for various reconstructed objects. + std::size_t n_spacepoints = 0; + std::size_t n_host_seeds = 0; + std::size_t n_device_seeds = 0; + std::size_t n_host_found_tracks = 0; + std::size_t n_device_found_tracks = 0; + std::size_t n_host_fitted_tracks = 0; + std::size_t n_device_fitted_tracks = 0; + + // Times elapsed in the various reconstruction steps. + performance::timing_info times; + + // Process the requested number of events. + for (std::size_t event = input_opts.skip; + event < input_opts.events + input_opts.skip; ++event) { + + // Instantiate host containers/collections. + edm::spacepoint_collection::host host_spacepoints{host_mr}; + edm::measurement_collection::host host_measurements{ + host_mr}; + edm::seed_collection::host host_seeds{host_mr}; + bound_track_parameters_collection_types::host host_track_params{ + &host_mr}; + edm::track_container::host host_found_tracks{host_mr}; + edm::track_container::host host_fitted_tracks{host_mr}; + + edm::spacepoint_collection::buffer device_spacepoints; + edm::measurement_collection::buffer + device_measurements; + edm::seed_collection::buffer device_seeds; + bound_track_parameters_collection_types::buffer device_track_params; + edm::track_container::buffer device_found_tracks; + edm::track_container::buffer device_fitted_tracks; + + { + // Measure the total wall time. + performance::timer wall_t("Wall time", times); + + { + // Read the spacepoints and measurements from the relevant event + // files. + performance::timer t("Host data reading", times); + io::read_spacepoints( + host_spacepoints, host_measurements, event, + input_opts.directory, + (input_opts.use_acts_geom_source ? &host_det : nullptr), + &det_descr, input_opts.format); + } + + { + // Copy the spacepoint and measurement data to the device. + performance::timer t{"Host->Device data transfers", times}; + device_spacepoints = backend.copy().to( + vecmem::get_data(host_spacepoints), backend.mr().main, + backend.mr().host, vecmem::copy::type::host_to_device); + device_measurements = backend.copy().to( + vecmem::get_data(host_measurements), backend.mr().main, + backend.mr().host, vecmem::copy::type::host_to_device); + } + + { + // Reconstruct the spacepoints into seeds on the device. + performance::timer t("Device seeding", times); + device_seeds = (*device_seeding)(device_spacepoints); + backend.synchronize(); + } + + if (accelerator_opts.compare_with_cpu) { + // Reconstruct the spacepoints into seeds on the host. + performance::timer t("Host seeding", times); + host_seeds = host_seeding(vecmem::get_data(host_spacepoints)); + } + + { + // Run track parameter estimation on the device. + performance::timer t("Device T/P estimation", times); + device_track_params = (*device_tp_estimation)( + device_measurements, device_spacepoints, device_seeds, + bfield_vec); + backend.synchronize(); + } + + if (accelerator_opts.compare_with_cpu) { + // Run track parameter estimation on the host. + performance::timer t("Host T/P esimation", times); + host_track_params = host_tp_estimation( + vecmem::get_data(host_measurements), + vecmem::get_data(host_spacepoints), + vecmem::get_data(host_seeds), bfield_vec); + } + + { + // Run track finding on the device. + performance::timer t("Device track finding", times); + device_found_tracks = + (*device_finding)(device_det, device_field, + device_measurements, device_track_params); + } + + if (accelerator_opts.compare_with_cpu) { + // Run track finding on the host. + traccc::performance::timer t("Host track finding", times); + host_found_tracks = host_finding( + host_det, host_field, vecmem::get_data(host_measurements), + vecmem::get_data(host_track_params)); + } + + { + // Run track fitting on the device. + performance::timer t("Device track fitting", times); + device_fitted_tracks = (*device_fitting)( + device_det, device_field, device_found_tracks); + } + + if (accelerator_opts.compare_with_cpu) { + // Run track fitting on the host. + traccc::performance::timer t("Host track fitting", times); + host_fitted_tracks = host_fitting( + host_det, host_field, + traccc::edm::track_container:: + const_data(host_found_tracks)); + } + } + + // Copy device containers/collections back to the host for validation. + edm::seed_collection::host device_host_seeds{host_mr}; + backend.copy()(device_seeds, device_host_seeds)->wait(); + + bound_track_parameters_collection_types::host device_host_track_params{ + &host_mr}; + backend.copy()(device_track_params, device_host_track_params)->wait(); + + edm::track_container::host + device_host_found_tracks{host_mr, + vecmem::get_data(host_measurements)}; + backend + .copy()(device_found_tracks.tracks, device_host_found_tracks.tracks) + ->wait(); + backend + .copy()(device_found_tracks.states, device_host_found_tracks.states) + ->wait(); + + edm::track_container::host + device_host_fitted_tracks{host_mr, + vecmem::get_data(host_measurements)}; + backend + .copy()(device_fitted_tracks.tracks, + device_host_fitted_tracks.tracks) + ->wait(); + backend + .copy()(device_fitted_tracks.states, + device_host_fitted_tracks.states) + ->wait(); + + if (accelerator_opts.compare_with_cpu) { + // Show which event we are currently presenting the results for. + TRACCC_INFO("===>>> Event " << event << " <<<==="); + + // Compare the seeds made on the host and on the device + soa_comparator compare_seeds{ + "seeds", + details::comparator_factory< + edm::seed_collection::const_device::const_proxy_type>{ + vecmem::get_data(host_spacepoints), + vecmem::get_data(host_spacepoints)}}; + compare_seeds(vecmem::get_data(host_seeds), + vecmem::get_data(device_host_seeds)); + + // Compare the track parameters made on the host and on the device. + collection_comparator> + compare_track_parameters{"track parameters"}; + compare_track_parameters( + vecmem::get_data(host_track_params), + vecmem::get_data(device_host_track_params)); + + // Compare the found tracks made on the host and on the device. + soa_comparator> + compare_found_tracks{ + "found tracks", + details::comparator_factory::const_device::const_proxy_type>{ + vecmem::get_data(host_measurements), + vecmem::get_data(host_measurements), + vecmem::get_data(host_found_tracks.states), + vecmem::get_data(device_host_found_tracks.states)}}; + compare_found_tracks( + vecmem::get_data(host_found_tracks.tracks), + vecmem::get_data(device_host_found_tracks.tracks)); + + // Compare the fitted tracks made on the host and on the device. + soa_comparator> + compare_fitted_tracks{ + "fitted tracks", + details::comparator_factory::const_device::const_proxy_type>{ + vecmem::get_data(host_measurements), + vecmem::get_data(host_measurements), + vecmem::get_data(host_fitted_tracks.states), + vecmem::get_data(device_host_fitted_tracks.states)}}; + compare_fitted_tracks( + vecmem::get_data(host_fitted_tracks.tracks), + vecmem::get_data(device_host_fitted_tracks.tracks)); + } + + // Print information about the fitted tracks. + details::print_fitted_tracks_statistics(device_host_fitted_tracks, + logger()); + + // Collect overall statistics. + n_spacepoints += host_spacepoints.size(); + n_host_seeds += host_seeds.size(); + n_device_seeds += device_host_seeds.size(); + n_host_found_tracks += host_found_tracks.tracks.size(); + n_device_found_tracks += device_host_found_tracks.tracks.size(); + n_host_fitted_tracks += host_fitted_tracks.tracks.size(); + n_device_fitted_tracks += device_host_fitted_tracks.tracks.size(); + + // Write detailed performance data if requested. + if (performance_opts.run) { + + static constexpr bool USE_SILICON_CELLS = false; + event_data evt_data(input_opts.directory, event, host_mr, + input_opts.use_acts_geom_source, &host_det, + input_opts.format, USE_SILICON_CELLS); + + seeding_pw.write(vecmem::get_data(device_host_seeds), + vecmem::get_data(host_spacepoints), + vecmem::get_data(host_measurements), evt_data); + + finding_pw.write(edm::track_container::const_data( + device_host_found_tracks), + evt_data); + + postfit_finding_pw.write( + edm::track_container::const_data( + device_host_fitted_tracks), + evt_data); + + for (unsigned int i = 0; + i < device_host_fitted_tracks.tracks.size(); ++i) { + host_detector_visitor( + host_det, [&]( + const typename detector_traits_t::host& det) { + fitting_pw.write(device_host_fitted_tracks.tracks.at(i), + device_host_fitted_tracks.states, + host_measurements, det, evt_data); + }); + } + } + } + + // Finalize the performance writers if necessary. + if (performance_opts.run) { + seeding_pw.finalize(); + finding_pw.finalize(); + postfit_finding_pw.finalize(); + fitting_pw.finalize(); + } + + // Print some final statistics about the job. + TRACCC_INFO("===>>> Statistics <<<==="); + TRACCC_INFO(" Procssed measurements/spacepoints: " << n_spacepoints); + if (accelerator_opts.compare_with_cpu) { + TRACCC_INFO(" Found seeds (host): " << n_host_seeds); + } + TRACCC_INFO(" Found seeds (device): " << n_device_seeds); + if (accelerator_opts.compare_with_cpu) { + TRACCC_INFO(" Found tracks (host): " << n_host_found_tracks); + } + TRACCC_INFO(" Found tracks (device): " << n_device_found_tracks); + if (accelerator_opts.compare_with_cpu) { + TRACCC_INFO(" Fitted tracks (host): " << n_host_fitted_tracks); + } + TRACCC_INFO(" Fitted tracks (device): " << n_device_fitted_tracks); + TRACCC_INFO("===>>> Timing information <<<==="); + TRACCC_INFO(times); + + // Return gracefully. + return EXIT_SUCCESS; +} + +} // namespace traccc diff --git a/examples/run/cuda/CMakeLists.txt b/examples/run/cuda/CMakeLists.txt index eeeb648bd6..c4a01fb431 100644 --- a/examples/run/cuda/CMakeLists.txt +++ b/examples/run/cuda/CMakeLists.txt @@ -10,15 +10,6 @@ include( traccc-compiler-options-cuda ) # External(s). find_package( CUDAToolkit REQUIRED ) -traccc_add_executable( seq_example_cuda "seq_example_cuda.cpp" - LINK_LIBRARIES vecmem::core vecmem::cuda traccc::io traccc::performance - traccc::core traccc::device_common traccc::cuda - traccc::options detray::detectors detray::io - traccc_examples_common ) -traccc_add_executable( seeding_example_cuda "seeding_example_cuda.cpp" - LINK_LIBRARIES vecmem::core vecmem::cuda traccc::io traccc::performance - traccc::core traccc::device_common traccc::cuda - traccc::options traccc_examples_common ) traccc_add_executable( truth_finding_example_cuda "truth_finding_example_cuda.cpp" LINK_LIBRARIES vecmem::core vecmem::cuda traccc::io traccc::performance traccc::core traccc::device_common traccc::cuda @@ -28,9 +19,11 @@ traccc_add_executable( truth_fitting_example_cuda "truth_fitting_example_cuda.cp traccc::core traccc::device_common traccc::cuda traccc::options traccc_examples_common ) # -# Set up the "throughput applications". +# Set up the library used by the applications. # add_library( traccc_examples_cuda STATIC + "device_backend.hpp" + "device_backend.cpp" "full_chain_algorithm.hpp" "full_chain_algorithm.cpp" ) target_link_libraries( traccc_examples_cuda @@ -38,14 +31,21 @@ target_link_libraries( traccc_examples_cuda traccc::core traccc::device_common traccc::cuda traccc_examples_common ) +# +# Set up the applications. +# traccc_add_executable( throughput_st_cuda "throughput_st.cpp" - LINK_LIBRARIES indicators::indicators vecmem::core vecmem::cuda - detray::detectors detray::io traccc::io traccc::performance - traccc::core traccc::device_common traccc::cuda - traccc::options traccc_examples_cuda ) + LINK_LIBRARIES indicators::indicators traccc_examples_common + traccc_examples_cuda ) traccc_add_executable( throughput_mt_cuda "throughput_mt.cpp" - LINK_LIBRARIES indicators::indicators TBB::tbb vecmem::core vecmem::cuda - detray::detectors detray::io traccc::io traccc::performance - traccc::core traccc::device_common traccc::cuda - traccc::options traccc_examples_cuda ) + LINK_LIBRARIES TBB::tbb indicators::indicators traccc_examples_common + traccc_examples_cuda ) + +traccc_add_executable( track_finding_validation_cuda + "track_finding_validation.cpp" + LINK_LIBRARIES traccc_examples_common traccc_examples_cuda ) + +traccc_add_executable( reconstruction_validation_cuda + "reconstruction_validation.cpp" + LINK_LIBRARIES traccc_examples_common traccc_examples_cuda ) diff --git a/examples/run/cuda/device_backend.cpp b/examples/run/cuda/device_backend.cpp new file mode 100644 index 0000000000..0d44515860 --- /dev/null +++ b/examples/run/cuda/device_backend.cpp @@ -0,0 +1,168 @@ +/** 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 "device_backend.hpp" + +// Project include(s). +#include "traccc/cuda/ambiguity_resolution/greedy_ambiguity_resolution_algorithm.hpp" +#include "traccc/cuda/clusterization/clusterization_algorithm.hpp" +#include "traccc/cuda/clusterization/measurement_sorting_algorithm.hpp" +#include "traccc/cuda/finding/combinatorial_kalman_filter_algorithm.hpp" +#include "traccc/cuda/fitting/kalman_fitting_algorithm.hpp" +#include "traccc/cuda/seeding/seeding_algorithm.hpp" +#include "traccc/cuda/seeding/spacepoint_formation_algorithm.hpp" +#include "traccc/cuda/seeding/track_params_estimation.hpp" +#include "traccc/cuda/utils/make_magnetic_field.hpp" +#include "traccc/cuda/utils/stream.hpp" + +// VecMem include(s). +#include +#include +#include + +namespace traccc::cuda { + +struct device_backend::impl { + + /// CUDA stream to use + stream m_stream; + + /// Host memory resource + vecmem::cuda::host_memory_resource m_host_mr; + /// Device memory resource + vecmem::cuda::device_memory_resource m_device_mr{m_stream.device()}; + /// Traccc memory resource + memory_resource m_mr{m_device_mr, &m_host_mr}; + + /// (Asynchronous) Memory copy object + vecmem::cuda::async_copy m_copy{m_stream.cudaStream()}; + +}; // struct device_backend::impl + +device_backend::device_backend(std::unique_ptr logger) + : messaging(std::move(logger)), m_impl{std::make_unique()} {} + +device_backend::~device_backend() = default; + +vecmem::copy& device_backend::copy() const { + + return m_impl->m_copy; +} + +memory_resource& device_backend::mr() const { + + return m_impl->m_mr; +} + +void device_backend::synchronize() const { + + m_impl->m_stream.synchronize(); +} + +magnetic_field device_backend::make_magnetic_field( + const magnetic_field& bfield, const bool texture_memory) const { + + return cuda::make_magnetic_field( + bfield, (texture_memory ? magnetic_field_storage::texture_memory + : magnetic_field_storage::global_memory)); +} + +std::unique_ptr::buffer( + const edm::silicon_cell_collection::const_view&, + const silicon_detector_description::const_view&)>> +device_backend::make_clusterization_algorithm( + const clustering_config& config) const { + + TRACCC_VERBOSE("Constructing cuda::clusterization_algorithm"); + return std::make_unique( + m_impl->m_mr, m_impl->m_copy, m_impl->m_stream, config, + logger().clone("cuda::clusterization_algorithm")); +} + +std::unique_ptr::buffer( + const edm::measurement_collection::const_view&)>> +device_backend::make_measurement_sorting_algorithm() const { + + TRACCC_VERBOSE("Constructing cuda::measurement_sorting_algorithm"); + return std::make_unique( + m_impl->m_mr, m_impl->m_copy, m_impl->m_stream, + logger().clone("cuda::measurement_sorting_algorithm")); +} + +std::unique_ptr::const_view&)>> +device_backend::make_spacepoint_formation_algorithm() const { + + TRACCC_VERBOSE("Constructing cuda::spacepoint_formation_algorithm"); + return std::make_unique( + m_impl->m_mr, m_impl->m_copy, m_impl->m_stream, + logger().clone("cuda::spacepoint_formation_algorithm")); +} + +std::unique_ptr> +device_backend::make_seeding_algorithm( + const seedfinder_config& finder_config, + const spacepoint_grid_config& grid_config, + const seedfilter_config& filter_config) const { + + TRACCC_VERBOSE("Constructing cuda::seeding_algorithm"); + return std::make_unique( + finder_config, grid_config, filter_config, m_impl->m_mr, m_impl->m_copy, + m_impl->m_stream, logger().clone("cuda::seeding_algorithm")); +} + +std::unique_ptr::const_view&, + const edm::spacepoint_collection::const_view&, + const edm::seed_collection::const_view&, const vector3&)>> +device_backend::make_track_params_estimation_algorithm( + const track_params_estimation_config& config) const { + + TRACCC_VERBOSE("Constructing cuda::track_params_estimation"); + return std::make_unique( + config, m_impl->m_mr, m_impl->m_copy, m_impl->m_stream, + logger().clone("cuda::track_params_estimation")); +} + +std::unique_ptr::buffer( + const detector_buffer&, const magnetic_field&, + const edm::measurement_collection::const_view&, + const bound_track_parameters_collection_types::const_view&)>> +device_backend::make_finding_algorithm(const finding_config& config) const { + + TRACCC_VERBOSE("Constructing cuda::combinatorial_kalman_filter_algorithm"); + return std::make_unique( + config, m_impl->m_mr, m_impl->m_copy, m_impl->m_stream, + logger().clone("cuda::combinatorial_kalman_filter_algorithm")); +} + +std::unique_ptr::buffer( + const edm::track_container::const_view&)>> +device_backend::make_ambiguity_resolution_algorithm( + const ambiguity_resolution_config& config) const { + + TRACCC_VERBOSE("Constructing cuda::greedy_ambiguity_resolution_algorithm"); + return std::make_unique( + config, m_impl->m_mr, m_impl->m_copy, m_impl->m_stream, + logger().clone("cuda::greedy_ambiguity_resolution_algorithm")); +} + +std::unique_ptr::buffer( + const detector_buffer&, const magnetic_field&, + const edm::track_container::const_view&)>> +device_backend::make_fitting_algorithm(const fitting_config& config) const { + + TRACCC_VERBOSE("Constructing cuda::kalman_fitting_algorithm"); + return std::make_unique( + config, m_impl->m_mr, m_impl->m_copy, m_impl->m_stream, + logger().clone("cuda::kalman_fitting_algorithm")); +} + +} // namespace traccc::cuda diff --git a/examples/run/cuda/device_backend.hpp b/examples/run/cuda/device_backend.hpp new file mode 100644 index 0000000000..b91ba4d454 --- /dev/null +++ b/examples/run/cuda/device_backend.hpp @@ -0,0 +1,114 @@ +/** 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 "../common/device_backend.hpp" +#include "traccc/utils/messaging.hpp" + +// System include(s). +#include + +namespace traccc::cuda { + +/// CUDA Device Backend +class device_backend : public traccc::device_backend, public messaging { + + public: + /// Constructor + /// + /// @param logger The logger to use + /// + device_backend( + std::unique_ptr logger = getDummyLogger().clone()); + /// Destructor + ~device_backend(); + + /// @name Function(s) implemented from @c traccc::device_backend + /// @{ + + /// Access a copy object for the used device + vecmem::copy& copy() const override; + + /// Get the memory resource(s) used by the algorithms + memory_resource& mr() const override; + + /// Wait for the used device to finish all scheduled operations + void synchronize() const override; + + /// Set up the magnetic field for the device + magnetic_field make_magnetic_field( + const magnetic_field& bfield, + bool texture_memory = false) const override; + + /// Construct a clusterization algorithm instance + std::unique_ptr< + algorithm::buffer( + const edm::silicon_cell_collection::const_view&, + const silicon_detector_description::const_view&)>> + make_clusterization_algorithm( + const clustering_config& config) const override; + + /// Construct a measurement sorting algorithm instance + std::unique_ptr< + algorithm::buffer( + const edm::measurement_collection::const_view&)>> + make_measurement_sorting_algorithm() const override; + + /// Construct a spacepoint formation algorithm instance + std::unique_ptr::const_view&)>> + make_spacepoint_formation_algorithm() const override; + + /// Construct a seeding algorithm instance + std::unique_ptr> + make_seeding_algorithm( + const seedfinder_config& finder_config, + const spacepoint_grid_config& grid_config, + const seedfilter_config& filter_config) const override; + + /// Construct a track parameter estimation algorithm instance + std::unique_ptr::const_view&, + const edm::spacepoint_collection::const_view&, + const edm::seed_collection::const_view&, const vector3&)>> + make_track_params_estimation_algorithm( + const track_params_estimation_config& config) const override; + + /// Construct a track finding algorithm instance + std::unique_ptr::buffer( + const detector_buffer&, const magnetic_field&, + const edm::measurement_collection::const_view&, + const bound_track_parameters_collection_types::const_view&)>> + make_finding_algorithm(const finding_config& config) const override; + + /// Construct an ambiguity resolution algorithm instance + std::unique_ptr::buffer( + const edm::track_container::const_view&)>> + make_ambiguity_resolution_algorithm( + const ambiguity_resolution_config& config) const override; + + /// Construct a track fitting algorithm instance + std::unique_ptr::buffer( + const detector_buffer&, const magnetic_field&, + const edm::track_container::const_view&)>> + make_fitting_algorithm(const fitting_config& config) const override; + + /// @} + + private: + /// Implementation class + struct impl; + /// PIMPL data object + std::unique_ptr m_impl; + +}; // class device_backend + +} // namespace traccc::cuda diff --git a/examples/run/cuda/reconstruction_validation.cpp b/examples/run/cuda/reconstruction_validation.cpp new file mode 100644 index 0000000000..c2bafd7f62 --- /dev/null +++ b/examples/run/cuda/reconstruction_validation.cpp @@ -0,0 +1,20 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2021-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Project include(s). +#include "../common/device_reconstruction_validation.hpp" + +// Local include(s). +#include "device_backend.hpp" + +int main(int argc, char* argv[]) { + + return traccc::device_reconstruction_validation< + traccc::cuda::device_backend>("reconstruction_validation_cuda", + "CUDA Reconstruction Validation", argc, + argv); +} diff --git a/examples/run/cuda/seeding_example_cuda.cpp b/examples/run/cuda/seeding_example_cuda.cpp deleted file mode 100644 index 182078db69..0000000000 --- a/examples/run/cuda/seeding_example_cuda.cpp +++ /dev/null @@ -1,542 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2021-2025 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// Project include(s). -#include "../common/make_magnetic_field.hpp" -#include "../common/print_fitted_tracks_statistics.hpp" -#include "traccc/cuda/finding/combinatorial_kalman_filter_algorithm.hpp" -#include "traccc/cuda/fitting/kalman_fitting_algorithm.hpp" -#include "traccc/cuda/seeding/seeding_algorithm.hpp" -#include "traccc/cuda/seeding/track_params_estimation.hpp" -#include "traccc/cuda/utils/make_magnetic_field.hpp" -#include "traccc/definitions/common.hpp" -#include "traccc/device/container_d2h_copy_alg.hpp" -#include "traccc/device/container_h2d_copy_alg.hpp" -#include "traccc/efficiency/finding_performance_writer.hpp" -#include "traccc/efficiency/nseed_performance_writer.hpp" -#include "traccc/efficiency/seeding_performance_writer.hpp" -#include "traccc/efficiency/track_filter.hpp" -#include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" -#include "traccc/fitting/kalman_fitting_algorithm.hpp" -#include "traccc/geometry/detector.hpp" -#include "traccc/geometry/host_detector.hpp" -#include "traccc/io/read_detector.hpp" -#include "traccc/io/read_detector_description.hpp" -#include "traccc/io/read_measurements.hpp" -#include "traccc/io/read_spacepoints.hpp" -#include "traccc/io/utils.hpp" -#include "traccc/options/accelerator.hpp" -#include "traccc/options/detector.hpp" -#include "traccc/options/input_data.hpp" -#include "traccc/options/magnetic_field.hpp" -#include "traccc/options/performance.hpp" -#include "traccc/options/program_options.hpp" -#include "traccc/options/seed_matching.hpp" -#include "traccc/options/track_finding.hpp" -#include "traccc/options/track_fitting.hpp" -#include "traccc/options/track_matching.hpp" -#include "traccc/options/track_propagation.hpp" -#include "traccc/options/track_seeding.hpp" -#include "traccc/options/truth_finding.hpp" -#include "traccc/performance/collection_comparator.hpp" -#include "traccc/performance/soa_comparator.hpp" -#include "traccc/performance/timer.hpp" -#include "traccc/resolution/fitting_performance_writer.hpp" -#include "traccc/seeding/detail/track_params_estimation_config.hpp" -#include "traccc/seeding/seeding_algorithm.hpp" -#include "traccc/seeding/track_params_estimation.hpp" -#include "traccc/utils/propagation.hpp" - -// VecMem include(s). -#include -#include -#include -#include -#include -#include - -// System include(s). -#include -#include -#include - -using namespace traccc; - -int seq_run(const traccc::opts::track_seeding& seeding_opts, - const traccc::opts::track_finding& finding_opts, - const traccc::opts::track_propagation& propagation_opts, - const traccc::opts::track_fitting& fitting_opts, - const traccc::opts::input_data& input_opts, - const traccc::opts::detector& detector_opts, - const traccc::opts::magnetic_field& bfield_opts, - const traccc::opts::performance& performance_opts, - const traccc::opts::accelerator& accelerator_opts, - const traccc::opts::truth_finding& truth_finding_opts, - const traccc::opts::seed_matching& seed_matching_opts, - const traccc::opts::track_matching& track_matching_opts, - std::unique_ptr ilogger) { - TRACCC_LOCAL_LOGGER(std::move(ilogger)); - - // Memory resources used by the application. - vecmem::host_memory_resource host_mr; - vecmem::cuda::host_memory_resource cuda_host_mr; - vecmem::cuda::managed_memory_resource mng_mr; - vecmem::cuda::device_memory_resource device_mr; - traccc::memory_resource mr{device_mr, &cuda_host_mr}; - - // Performance writer - traccc::seeding_performance_writer sd_performance_writer( - traccc::seeding_performance_writer::config{ - .truth_config = truth_finding_opts, - .seed_truth_config = seed_matching_opts}, - logger().clone("SeedingPerformanceWriter")); - traccc::finding_performance_writer find_performance_writer( - traccc::finding_performance_writer::config{ - .truth_config = truth_finding_opts, - .track_truth_config = track_matching_opts}, - logger().clone("FindingPerformanceWriter")); - traccc::finding_performance_writer postfit_find_performance_writer( - traccc::finding_performance_writer::config{ - .file_path = "performance_track_postfit_finding.root", - .truth_config = truth_finding_opts, - .track_truth_config = track_matching_opts, - .require_fit = true}, - logger().clone("PostFitFindingPerformanceWriter")); - traccc::fitting_performance_writer fit_performance_writer( - traccc::fitting_performance_writer::config{}, - logger().clone("FittingPerformanceWriter")); - - traccc::nseed_performance_writer nsd_performance_writer( - "nseed_performance_", - std::make_unique( - 2.7f, 1.f * traccc::unit::GeV), - std::make_unique(0.6f)); - - if (performance_opts.run) { - nsd_performance_writer.initialize(); - } - - // Output stats - uint64_t n_spacepoints = 0; - uint64_t n_seeds = 0; - uint64_t n_seeds_cuda = 0; - uint64_t n_found_tracks = 0; - uint64_t n_found_tracks_cuda = 0; - uint64_t n_fitted_tracks = 0; - uint64_t n_fitted_tracks_cuda = 0; - - /***************************** - * Build a geometry - *****************************/ - - traccc::silicon_detector_description::host host_det_descr{host_mr}; - traccc::io::read_detector_description( - host_det_descr, detector_opts.detector_file, - detector_opts.digitization_file, traccc::data_format::json); - - // B field value - const traccc::vector3 field_vec(seeding_opts); - const auto host_field = traccc::details::make_magnetic_field(bfield_opts); - const auto device_field = traccc::cuda::make_magnetic_field( - host_field, - (accelerator_opts.use_gpu_texture_memory - ? traccc::cuda::magnetic_field_storage::texture_memory - : traccc::cuda::magnetic_field_storage::global_memory)); - - // Construct a Detray detector object, if supported by the configuration. - traccc::host_detector host_det; - traccc::io::read_detector(host_det, mng_mr, detector_opts.detector_file, - detector_opts.material_file, - detector_opts.grid_file); - - // Copy objects - vecmem::copy host_copy; - vecmem::cuda::copy copy; - - const traccc::detector_buffer detector_buffer = - traccc::buffer_from_host_detector(host_det, mng_mr, host_copy); - - // Seeding algorithm - const traccc::seedfinder_config seedfinder_config(seeding_opts); - const traccc::seedfilter_config seedfilter_config(seeding_opts); - const traccc::spacepoint_grid_config spacepoint_grid_config(seeding_opts); - traccc::host::seeding_algorithm sa( - seedfinder_config, spacepoint_grid_config, seedfilter_config, host_mr, - logger().clone("HostSeedingAlg")); - const traccc::track_params_estimation_config track_params_estimation_config; - traccc::host::track_params_estimation tp( - track_params_estimation_config, host_mr, - logger().clone("HostTrackParEstAlg")); - - traccc::cuda::stream stream; - - vecmem::cuda::async_copy async_copy{stream.cudaStream()}; - - traccc::cuda::seeding_algorithm sa_cuda{seedfinder_config, - spacepoint_grid_config, - seedfilter_config, - mr, - async_copy, - stream, - logger().clone("CudaSeedingAlg")}; - traccc::cuda::track_params_estimation tp_cuda{ - track_params_estimation_config, mr, async_copy, stream, - logger().clone("CudaTrackParEstAlg")}; - - // Propagation configuration - detray::propagation::config propagation_config(propagation_opts); - - // Finding algorithm configuration - traccc::finding_config cfg(finding_opts); - cfg.propagation = propagation_config; - - // Finding algorithm object - traccc::host::combinatorial_kalman_filter_algorithm host_finding( - cfg, host_mr, logger().clone("HostFindingAlg")); - traccc::cuda::combinatorial_kalman_filter_algorithm device_finding( - cfg, mr, async_copy, stream, logger().clone("CudaFindingAlg")); - - // Fitting algorithm object - traccc::fitting_config fit_cfg(fitting_opts); - fit_cfg.propagation = propagation_config; - - traccc::host::kalman_fitting_algorithm host_fitting( - fit_cfg, host_mr, host_copy, logger().clone("HostFittingAlg")); - traccc::cuda::kalman_fitting_algorithm device_fitting( - fit_cfg, mr, async_copy, stream, logger().clone("CudaFittingAlg")); - - traccc::performance::timing_info elapsedTimes; - - // Loop over events - for (std::size_t event = input_opts.skip; - event < input_opts.events + input_opts.skip; ++event) { - - // Instantiate host containers/collections - traccc::edm::spacepoint_collection::host spacepoints_per_event{host_mr}; - traccc::edm::measurement_collection::host - measurements_per_event{host_mr}; - traccc::host::seeding_algorithm::output_type seeds{host_mr}; - traccc::host::track_params_estimation::output_type params; - traccc::edm::track_container::host - track_candidates{host_mr}; - traccc::edm::track_container::host - track_states{host_mr}; - - traccc::edm::seed_collection::buffer seeds_cuda_buffer; - traccc::bound_track_parameters_collection_types::buffer - params_cuda_buffer(0, *mr.host); - - traccc::edm::track_container::buffer - track_candidates_cuda_buffer; - - traccc::edm::track_container::buffer - track_states_cuda_buffer; - - { // Start measuring wall time - traccc::performance::timer wall_t("Wall time", elapsedTimes); - - /*----------------- - hit file reading - -----------------*/ - { - traccc::performance::timer t("Hit reading (cpu)", - elapsedTimes); - // Read the hits and measurements from the relevant event files - traccc::io::read_spacepoints( - spacepoints_per_event, measurements_per_event, event, - input_opts.directory, - (input_opts.use_acts_geom_source ? &host_det : nullptr), - &host_det_descr, input_opts.format); - - } // stop measuring hit reading timer - - /*---------------------------- - Seeding algorithm - ----------------------------*/ - - /// CUDA - - // Copy the spacepoint and module data to the device. - traccc::edm::spacepoint_collection::buffer spacepoints_cuda_buffer( - static_cast(spacepoints_per_event.size()), - mr.main); - async_copy.setup(spacepoints_cuda_buffer)->wait(); - async_copy(vecmem::get_data(spacepoints_per_event), - spacepoints_cuda_buffer) - ->wait(); - - traccc::edm::measurement_collection::buffer - measurements_cuda_buffer( - static_cast(measurements_per_event.size()), - mr.main); - async_copy.setup(measurements_cuda_buffer)->wait(); - async_copy(vecmem::get_data(measurements_per_event), - measurements_cuda_buffer) - ->wait(); - - { - traccc::performance::timer t("Seeding (cuda)", elapsedTimes); - // Reconstruct the spacepoints into seeds. - seeds_cuda_buffer = sa_cuda(spacepoints_cuda_buffer); - stream.synchronize(); - } // stop measuring seeding cuda timer - - // CPU - - if (accelerator_opts.compare_with_cpu) { - { - traccc::performance::timer t("Seeding (cpu)", - elapsedTimes); - seeds = sa(vecmem::get_data(spacepoints_per_event)); - } - } // stop measuring seeding cpu timer - - /*---------------------------- - Track params estimation - ----------------------------*/ - - // CUDA - { - traccc::performance::timer t("Track params (cuda)", - elapsedTimes); - params_cuda_buffer = - tp_cuda(measurements_cuda_buffer, spacepoints_cuda_buffer, - seeds_cuda_buffer, field_vec); - stream.synchronize(); - } // stop measuring track params cuda timer - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Track params (cpu)", - elapsedTimes); - params = tp(vecmem::get_data(measurements_per_event), - vecmem::get_data(spacepoints_per_event), - vecmem::get_data(seeds), field_vec); - } // stop measuring track params cpu timer - - /*------------------------ - Track Finding with CKF - ------------------------*/ - - { - traccc::performance::timer t("Track finding with CKF (cuda)", - elapsedTimes); - track_candidates_cuda_buffer = device_finding( - detector_buffer, device_field, measurements_cuda_buffer, - params_cuda_buffer); - } - - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Track finding with CKF (cpu)", - elapsedTimes); - track_candidates = - host_finding(host_det, host_field, - vecmem::get_data(measurements_per_event), - vecmem::get_data(params)); - } - - /*------------------------ - Track Fitting with KF - ------------------------*/ - - { - traccc::performance::timer t("Track fitting with KF (cuda)", - elapsedTimes); - - track_states_cuda_buffer = - device_fitting(detector_buffer, device_field, - track_candidates_cuda_buffer); - } - - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Track fitting with KF (cpu)", - elapsedTimes); - track_states = host_fitting( - host_det, host_field, - traccc::edm::track_container< - traccc::default_algebra>::const_data(track_candidates)); - } - - } // Stop measuring wall time - - /*---------------------------------- - compare seeds from cpu and cuda - ----------------------------------*/ - - // Copy the seeds to the host for comparisons - traccc::edm::seed_collection::host seeds_cuda{host_mr}; - traccc::bound_track_parameters_collection_types::host params_cuda; - async_copy(seeds_cuda_buffer, seeds_cuda)->wait(); - async_copy(params_cuda_buffer, params_cuda)->wait(); - - // Copy track candidates from device to host - traccc::edm::track_container::host - track_candidates_cuda{host_mr, - vecmem::get_data(measurements_per_event)}; - async_copy(track_candidates_cuda_buffer.tracks, - track_candidates_cuda.tracks) - ->wait(); - - // Copy track states from device to host - traccc::edm::track_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(); - track_states_cuda.measurements = - vecmem::get_data(measurements_per_event); - - if (accelerator_opts.compare_with_cpu) { - // Show which event we are currently presenting the results for. - std::cout << "===>>> Event " << event << " <<<===" << std::endl; - - // Compare the seeds made on the host and on the device - traccc::soa_comparator compare_seeds{ - "seeds", traccc::details::comparator_factory< - traccc::edm::seed_collection::const_device:: - const_proxy_type>{ - vecmem::get_data(spacepoints_per_event), - vecmem::get_data(spacepoints_per_event)}}; - compare_seeds(vecmem::get_data(seeds), - vecmem::get_data(seeds_cuda)); - - // Compare the track parameters made on the host and on the device. - traccc::collection_comparator> - compare_track_parameters{"track parameters"}; - compare_track_parameters(vecmem::get_data(params), - vecmem::get_data(params_cuda)); - - // Compare the track candidates made on the host and on the - // device - traccc::soa_comparator< - traccc::edm::track_collection> - compare_track_candidates{ - "track candidates", - traccc::details::comparator_factory< - traccc::edm::track_collection:: - const_device::const_proxy_type>{ - vecmem::get_data(measurements_per_event), - vecmem::get_data(measurements_per_event), - {}, - {}}}; - compare_track_candidates( - vecmem::get_data(track_candidates.tracks), - vecmem::get_data(track_candidates_cuda.tracks)); - } - - /*---------------- - Statistics - ---------------*/ - - details::print_fitted_tracks_statistics(track_states_cuda, logger()); - n_spacepoints += spacepoints_per_event.size(); - n_seeds_cuda += seeds_cuda.size(); - n_seeds += seeds.size(); - n_found_tracks_cuda += track_candidates_cuda.tracks.size(); - n_found_tracks += track_candidates.tracks.size(); - n_fitted_tracks_cuda += track_states_cuda.tracks.size(); - n_fitted_tracks += track_states.tracks.size(); - - /*------------ - Writer - ------------*/ - - if (performance_opts.run) { - - traccc::event_data evt_data(input_opts.directory, event, host_mr, - input_opts.use_acts_geom_source, - &host_det, input_opts.format, false); - - sd_performance_writer.write( - vecmem::get_data(seeds_cuda), - vecmem::get_data(spacepoints_per_event), - vecmem::get_data(measurements_per_event), evt_data); - - find_performance_writer.write( - traccc::edm::track_container< - traccc::default_algebra>::const_data(track_candidates_cuda), - evt_data); - - postfit_find_performance_writer.write( - traccc::edm::track_container< - traccc::default_algebra>::const_data(track_states_cuda), - evt_data); - - for (unsigned int i = 0; i < track_states_cuda.tracks.size(); i++) { - host_detector_visitor( - host_det, [&]( - const typename detector_traits_t::host& det) { - fit_performance_writer.write( - track_states_cuda.tracks.at(i), - track_states_cuda.states, measurements_per_event, - det, evt_data); - }); - } - } - } - - if (performance_opts.run) { - sd_performance_writer.finalize(); - nsd_performance_writer.finalize(); - find_performance_writer.finalize(); - postfit_find_performance_writer.finalize(); - fit_performance_writer.finalize(); - std::cout << nsd_performance_writer.generate_report_str(); - } - - std::cout << "==> Statistics ... " << std::endl; - std::cout << "- read " << n_spacepoints << " spacepoints" << std::endl; - std::cout << "- created (cpu) " << n_seeds << " seeds" << std::endl; - std::cout << "- created (cuda) " << n_seeds_cuda << " seeds" << std::endl; - std::cout << "- created (cpu) " << n_found_tracks << " found tracks" - << std::endl; - std::cout << "- created (cuda) " << n_found_tracks_cuda << " found tracks" - << std::endl; - std::cout << "- created (cpu) " << n_fitted_tracks << " fitted tracks" - << std::endl; - std::cout << "- created (cuda) " << n_fitted_tracks_cuda << " fitted tracks" - << std::endl; - std::cout << "==>Elapsed times...\n" << elapsedTimes << std::endl; - - return 0; -} - -// The main routine -// -int main(int argc, char* argv[]) { - std::unique_ptr logger = traccc::getDefaultLogger( - "TracccExampleSeedingCuda", traccc::Logging::Level::INFO); - - // Program options. - traccc::opts::detector detector_opts; - traccc::opts::magnetic_field bfield_opts; - traccc::opts::input_data input_opts; - traccc::opts::track_seeding seeding_opts; - traccc::opts::track_finding finding_opts; - traccc::opts::track_propagation propagation_opts; - traccc::opts::track_fitting fitting_opts; - traccc::opts::performance performance_opts; - traccc::opts::accelerator accelerator_opts; - traccc::opts::truth_finding truth_finding_opts; - traccc::opts::seed_matching seed_matching_opts; - traccc::opts::track_matching track_matching_opts; - traccc::opts::program_options program_opts{ - "Full Tracking Chain Using CUDA (without clusterization)", - {detector_opts, bfield_opts, input_opts, seeding_opts, finding_opts, - propagation_opts, fitting_opts, performance_opts, accelerator_opts, - truth_finding_opts, seed_matching_opts, track_matching_opts}, - argc, - argv, - logger->cloneWithSuffix("Options")}; - - // Run the application. - return seq_run(seeding_opts, finding_opts, propagation_opts, fitting_opts, - input_opts, detector_opts, bfield_opts, performance_opts, - accelerator_opts, truth_finding_opts, seed_matching_opts, - track_matching_opts, logger->clone()); -} diff --git a/examples/run/cuda/seq_example_cuda.cpp b/examples/run/cuda/seq_example_cuda.cpp deleted file mode 100644 index 64a0d158a3..0000000000 --- a/examples/run/cuda/seq_example_cuda.cpp +++ /dev/null @@ -1,594 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2021-2025 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// Project include(s). -#include "../common/make_magnetic_field.hpp" -#include "traccc/ambiguity_resolution/greedy_ambiguity_resolution_algorithm.hpp" -#include "traccc/clusterization/clusterization_algorithm.hpp" -#include "traccc/cuda/ambiguity_resolution/greedy_ambiguity_resolution_algorithm.hpp" -#include "traccc/cuda/clusterization/clusterization_algorithm.hpp" -#include "traccc/cuda/clusterization/measurement_sorting_algorithm.hpp" -#include "traccc/cuda/finding/combinatorial_kalman_filter_algorithm.hpp" -#include "traccc/cuda/fitting/kalman_fitting_algorithm.hpp" -#include "traccc/cuda/seeding/seeding_algorithm.hpp" -#include "traccc/cuda/seeding/spacepoint_formation_algorithm.hpp" -#include "traccc/cuda/seeding/track_params_estimation.hpp" -#include "traccc/cuda/utils/make_magnetic_field.hpp" -#include "traccc/cuda/utils/stream.hpp" -#include "traccc/device/container_d2h_copy_alg.hpp" -#include "traccc/efficiency/seeding_performance_writer.hpp" -#include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" -#include "traccc/fitting/kalman_fitting_algorithm.hpp" -#include "traccc/geometry/detector.hpp" -#include "traccc/geometry/detector_buffer.hpp" -#include "traccc/geometry/host_detector.hpp" -#include "traccc/io/read_cells.hpp" -#include "traccc/io/read_detector.hpp" -#include "traccc/io/read_detector_description.hpp" -#include "traccc/io/utils.hpp" -#include "traccc/options/accelerator.hpp" -#include "traccc/options/clusterization.hpp" -#include "traccc/options/detector.hpp" -#include "traccc/options/input_data.hpp" -#include "traccc/options/magnetic_field.hpp" -#include "traccc/options/performance.hpp" -#include "traccc/options/program_options.hpp" -#include "traccc/options/track_finding.hpp" -#include "traccc/options/track_fitting.hpp" -#include "traccc/options/track_propagation.hpp" -#include "traccc/options/track_resolution.hpp" -#include "traccc/options/track_seeding.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/seeding/seeding_algorithm.hpp" -#include "traccc/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" -#include "traccc/seeding/track_params_estimation.hpp" -#include "traccc/utils/propagation.hpp" - -// VecMem include(s). -#include -#include -#include -#include - -// System include(s). -#include -#include -#include -#include - -int seq_run(const traccc::opts::detector& detector_opts, - const traccc::opts::magnetic_field& bfield_opts, - const traccc::opts::input_data& input_opts, - const traccc::opts::clusterization& clusterization_opts, - const traccc::opts::track_seeding& seeding_opts, - const traccc::opts::track_finding& finding_opts, - const traccc::opts::track_propagation& propagation_opts, - const traccc::opts::track_resolution& resolution_opts, - const traccc::opts::track_fitting& fitting_opts, - const traccc::opts::performance& performance_opts, - const traccc::opts::accelerator& accelerator_opts, - std::unique_ptr ilogger) { - TRACCC_LOCAL_LOGGER(std::move(ilogger)); - - // Memory resources used by the application. - vecmem::host_memory_resource host_mr; - vecmem::cuda::host_memory_resource cuda_host_mr; - vecmem::cuda::device_memory_resource device_mr; - traccc::memory_resource mr{device_mr, &cuda_host_mr}; - - // Host copy object - vecmem::copy host_copy; - - // CUDA types used. - traccc::cuda::stream stream; - vecmem::cuda::async_copy copy{stream.cudaStream()}; - - // Construct the detector description object. - traccc::silicon_detector_description::host host_det_descr{host_mr}; - traccc::io::read_detector_description( - host_det_descr, detector_opts.detector_file, - detector_opts.digitization_file, traccc::data_format::json); - traccc::silicon_detector_description::data host_det_descr_data{ - vecmem::get_data(host_det_descr)}; - traccc::silicon_detector_description::buffer device_det_descr{ - static_cast( - host_det_descr.size()), - device_mr}; - copy.setup(device_det_descr)->wait(); - copy(host_det_descr_data, device_det_descr)->wait(); - - // Construct a Detray detector object, if supported by the configuration. - traccc::host_detector host_detector; - traccc::io::read_detector( - host_detector, host_mr, detector_opts.detector_file, - detector_opts.material_file, detector_opts.grid_file); - const traccc::detector_buffer device_detector = - traccc::buffer_from_host_detector(host_detector, device_mr, copy); - stream.synchronize(); - - // Output stats - uint64_t n_cells = 0; - uint64_t n_measurements = 0; - uint64_t n_measurements_cuda = 0; - uint64_t n_spacepoints = 0; - uint64_t n_spacepoints_cuda = 0; - uint64_t n_seeds = 0; - uint64_t n_seeds_cuda = 0; - uint64_t n_found_tracks = 0; - uint64_t n_found_tracks_cuda = 0; - uint64_t n_ambiguity_free_tracks = 0; - uint64_t n_ambiguity_free_tracks_cuda = 0; - uint64_t n_fitted_tracks = 0; - uint64_t n_fitted_tracks_cuda = 0; - - // Type definitions - using host_spacepoint_formation_algorithm = - traccc::host::silicon_pixel_spacepoint_formation_algorithm; - using device_spacepoint_formation_algorithm = - traccc::cuda::spacepoint_formation_algorithm; - - using host_finding_algorithm = - traccc::host::combinatorial_kalman_filter_algorithm; - using device_finding_algorithm = - traccc::cuda::combinatorial_kalman_filter_algorithm; - - using host_fitting_algorithm = traccc::host::kalman_fitting_algorithm; - using device_fitting_algorithm = traccc::cuda::kalman_fitting_algorithm; - - // Algorithm configuration(s). - detray::propagation::config propagation_config(propagation_opts); - - const traccc::seedfinder_config seedfinder_config(seeding_opts); - const traccc::seedfilter_config seedfilter_config(seeding_opts); - const traccc::spacepoint_grid_config spacepoint_grid_config(seeding_opts); - - traccc::finding_config finding_cfg(finding_opts); - finding_cfg.propagation = propagation_config; - - traccc::host::greedy_ambiguity_resolution_algorithm::config_type - resolution_config(resolution_opts); - - traccc::fitting_config fitting_cfg(fitting_opts); - fitting_cfg.propagation = propagation_config; - - // Constant B field for the track finding and fitting - const traccc::vector3 field_vec(seeding_opts); - const auto host_field = traccc::details::make_magnetic_field(bfield_opts); - const auto device_field = traccc::cuda::make_magnetic_field( - host_field, - (accelerator_opts.use_gpu_texture_memory - ? traccc::cuda::magnetic_field_storage::texture_memory - : traccc::cuda::magnetic_field_storage::global_memory)); - - traccc::host::clusterization_algorithm ca( - host_mr, logger().clone("HostClusteringAlg")); - host_spacepoint_formation_algorithm sf( - host_mr, logger().clone("HostSpFormationAlg")); - traccc::host::seeding_algorithm sa( - seedfinder_config, spacepoint_grid_config, seedfilter_config, host_mr, - logger().clone("HostSeedingAlg")); - traccc::track_params_estimation_config track_params_estimation_config; - traccc::host::track_params_estimation tp( - track_params_estimation_config, host_mr, - logger().clone("HostTrackParEstAlg")); - host_finding_algorithm finding_alg(finding_cfg, host_mr, - logger().clone("HostFindingAlg")); - traccc::host::greedy_ambiguity_resolution_algorithm resolution_alg_cpu( - resolution_config, host_mr, - logger().clone("HostAmbiguityResolutionAlg")); - host_fitting_algorithm fitting_alg(fitting_cfg, host_mr, host_copy, - logger().clone("HostFittingAlg")); - - traccc::cuda::clusterization_algorithm ca_cuda( - mr, copy, stream, clusterization_opts, - logger().clone("CudaClusteringAlg")); - traccc::cuda::measurement_sorting_algorithm ms_cuda( - mr, copy, stream, logger().clone("CudaMeasSortingAlg")); - device_spacepoint_formation_algorithm sf_cuda( - mr, copy, stream, logger().clone("CudaSpFormationAlg")); - traccc::cuda::seeding_algorithm sa_cuda( - seedfinder_config, spacepoint_grid_config, seedfilter_config, mr, copy, - stream, logger().clone("CudaSeedingAlg")); - traccc::cuda::track_params_estimation tp_cuda( - track_params_estimation_config, mr, copy, stream, - logger().clone("CudaTrackParEstAlg")); - device_finding_algorithm finding_alg_cuda(finding_cfg, mr, copy, stream, - logger().clone("CudaFindingAlg")); - traccc::cuda::greedy_ambiguity_resolution_algorithm resolution_alg_cuda( - resolution_config, mr, copy, stream, - logger().clone("CudaAmbiguityResolutionAlg")); - device_fitting_algorithm fitting_alg_cuda(fitting_cfg, mr, copy, stream, - logger().clone("CudaFittingAlg")); - - // performance writer - traccc::seeding_performance_writer sd_performance_writer( - traccc::seeding_performance_writer::config{}, - logger().clone("SeedingPerformanceWriter")); - - traccc::performance::timing_info elapsedTimes; - - // Loop over events - for (std::size_t event = input_opts.skip; - event < input_opts.events + input_opts.skip; ++event) { - - // Instantiate host containers/collections - traccc::host::clusterization_algorithm::output_type - measurements_per_event{host_mr}; - host_spacepoint_formation_algorithm::output_type spacepoints_per_event{ - host_mr}; - traccc::host::seeding_algorithm::output_type seeds{host_mr}; - traccc::host::track_params_estimation::output_type params; - 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_mr}; - - // Instantiate cuda containers/collections - traccc::edm::measurement_collection::buffer - measurements_cuda_buffer; - traccc::edm::spacepoint_collection::buffer spacepoints_cuda_buffer; - traccc::edm::seed_collection::buffer seeds_cuda_buffer; - traccc::bound_track_parameters_collection_types::buffer - params_cuda_buffer(0, *mr.host); - traccc::edm::track_container::buffer - track_candidates_buffer; - traccc::edm::track_container::buffer - res_track_candidates_buffer; - traccc::edm::track_container::buffer - track_states_buffer; - - { - traccc::performance::timer wall_t("Wall time", elapsedTimes); - - traccc::edm::silicon_cell_collection::host cells_per_event{host_mr}; - - { - traccc::performance::timer t("File reading (cpu)", - elapsedTimes); - // Read the cells from the relevant event file into host memory. - static constexpr bool DEDUPLICATE = true; - traccc::io::read_cells( - cells_per_event, event, input_opts.directory, - logger().clone(), &host_det_descr, input_opts.format, - DEDUPLICATE, input_opts.use_acts_geom_source); - } // stop measuring file reading timer - - n_cells += cells_per_event.size(); - - // Create device copy of input collections - traccc::edm::silicon_cell_collection::buffer cells_buffer( - static_cast(cells_per_event.size()), mr.main); - copy.setup(cells_buffer)->wait(); - copy(vecmem::get_data(cells_per_event), cells_buffer)->wait(); - - // CUDA - { - traccc::performance::timer t("Clusterization (cuda)", - elapsedTimes); - // Reconstruct it into spacepoints on the device. - auto unsorted_measurements = - ca_cuda(cells_buffer, device_det_descr); - measurements_cuda_buffer = ms_cuda(unsorted_measurements); - stream.synchronize(); - } // stop measuring clusterization cuda timer - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Clusterization (cpu)", - elapsedTimes); - measurements_per_event = - ca(vecmem::get_data(cells_per_event), host_det_descr_data); - } // stop measuring clusterization cpu timer - - // Perform seeding, track finding and fitting only when using a - // Detray geometry. - // CUDA - { - traccc::performance::timer t("Spacepoint formation (cuda)", - elapsedTimes); - spacepoints_cuda_buffer = - sf_cuda(device_detector, measurements_cuda_buffer); - stream.synchronize(); - } // stop measuring spacepoint formation cuda timer - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Spacepoint formation (cpu)", - elapsedTimes); - spacepoints_per_event = - sf(host_detector, vecmem::get_data(measurements_per_event)); - } // stop measuring spacepoint formation cpu timer - - // CUDA - { - traccc::performance::timer t("Seeding (cuda)", elapsedTimes); - seeds_cuda_buffer = sa_cuda(spacepoints_cuda_buffer); - stream.synchronize(); - } // stop measuring seeding cuda timer - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Seeding (cpu)", elapsedTimes); - seeds = sa(vecmem::get_data(spacepoints_per_event)); - } // stop measuring seeding cpu timer - - // CUDA - { - traccc::performance::timer t("Track params (cuda)", - elapsedTimes); - params_cuda_buffer = - tp_cuda(measurements_cuda_buffer, spacepoints_cuda_buffer, - seeds_cuda_buffer, field_vec); - stream.synchronize(); - } // stop measuring track params timer - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Track params (cpu)", - elapsedTimes); - params = tp(vecmem::get_data(measurements_per_event), - vecmem::get_data(spacepoints_per_event), - vecmem::get_data(seeds), field_vec); - } // stop measuring track params cpu timer - - // CUDA - { - traccc::performance::timer timer{"Track finding (cuda)", - elapsedTimes}; - track_candidates_buffer = finding_alg_cuda( - device_detector, device_field, measurements_cuda_buffer, - params_cuda_buffer); - } - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer timer{"Track finding (cpu)", - elapsedTimes}; - track_candidates = - finding_alg(host_detector, host_field, - vecmem::get_data(measurements_per_event), - vecmem::get_data(params)); - } - - // CUDA - { - traccc::performance::timer timer{"Ambiguity resolution (cuda)", - elapsedTimes}; - res_track_candidates_buffer = - resolution_alg_cuda(track_candidates_buffer); - } - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer timer{"Ambiguity resolution (cpu)", - elapsedTimes}; - res_track_candidates = resolution_alg_cpu( - traccc::edm::track_container< - traccc::default_algebra>::const_data(track_candidates)); - } - - // CUDA - { - traccc::performance::timer timer{"Track fitting (cuda)", - elapsedTimes}; - track_states_buffer = fitting_alg_cuda( - device_detector, device_field, res_track_candidates_buffer); - } - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer timer{"Track fitting (cpu)", - elapsedTimes}; - track_states = fitting_alg( - host_detector, host_field, - traccc::edm::track_container:: - const_data(res_track_candidates)); - } - - } // Stop measuring wall time - - /*---------------------------------- - compare cpu and cuda result - ----------------------------------*/ - - traccc::edm::measurement_collection::host - measurements_per_event_cuda{host_mr}; - traccc::edm::spacepoint_collection::host spacepoints_per_event_cuda{ - host_mr}; - traccc::edm::seed_collection::host seeds_cuda{host_mr}; - traccc::bound_track_parameters_collection_types::host params_cuda; - traccc::edm::track_collection::host - track_candidates_cuda{host_mr}; - traccc::edm::track_collection::host - res_track_candidates_cuda{host_mr}; - traccc::edm::track_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(); - copy(seeds_cuda_buffer, seeds_cuda)->wait(); - copy(params_cuda_buffer, params_cuda)->wait(); - copy(track_candidates_buffer.tracks, track_candidates_cuda, - vecmem::copy::type::device_to_host) - ->wait(); - copy(res_track_candidates_buffer.tracks, res_track_candidates_cuda, - vecmem::copy::type::device_to_host) - ->wait(); - 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(); - track_states_cuda.measurements = - vecmem::get_data(measurements_per_event_cuda); - stream.synchronize(); - - if (accelerator_opts.compare_with_cpu) { - - // Show which event we are currently presenting the results for. - TRACCC_INFO("===>>> Event " << event << " <<<==="); - - // Compare the measurements made on the host and on the device. - traccc::soa_comparator< - traccc::edm::measurement_collection> - compare_measurements{"measurements"}; - compare_measurements(vecmem::get_data(measurements_per_event), - vecmem::get_data(measurements_per_event_cuda)); - - // Compare the spacepoints made on the host and on the device. - traccc::soa_comparator - compare_spacepoints{"spacepoints"}; - compare_spacepoints(vecmem::get_data(spacepoints_per_event), - vecmem::get_data(spacepoints_per_event_cuda)); - - // Compare the seeds made on the host and on the device - traccc::soa_comparator compare_seeds{ - "seeds", traccc::details::comparator_factory< - traccc::edm::seed_collection::const_device:: - const_proxy_type>{ - vecmem::get_data(spacepoints_per_event), - vecmem::get_data(spacepoints_per_event_cuda)}}; - compare_seeds(vecmem::get_data(seeds), - vecmem::get_data(seeds_cuda)); - - // Compare the track parameters made on the host and on the device. - traccc::collection_comparator> - compare_track_parameters{"track parameters"}; - compare_track_parameters(vecmem::get_data(params), - vecmem::get_data(params_cuda)); - - // Compare tracks found on the host and on the device. - traccc::soa_comparator< - traccc::edm::track_collection> - compare_track_candidates{ - "track candidates", - traccc::details::comparator_factory< - traccc::edm::track_collection:: - const_device::const_proxy_type>{ - vecmem::get_data(measurements_per_event), - vecmem::get_data(measurements_per_event_cuda), - {}, - {}}}; - compare_track_candidates(vecmem::get_data(track_candidates.tracks), - vecmem::get_data(track_candidates_cuda)); - - // Compare tracks resolved on the host and on the device. - traccc::soa_comparator< - traccc::edm::track_collection> - compare_resolved_track_candidates{ - "resolved track candidates", - traccc::details::comparator_factory< - traccc::edm::track_collection:: - const_device::const_proxy_type>{ - vecmem::get_data(measurements_per_event), - vecmem::get_data(measurements_per_event_cuda), - {}, - {}}}; - compare_resolved_track_candidates( - vecmem::get_data(res_track_candidates.tracks), - vecmem::get_data(res_track_candidates_cuda)); - - // Compare tracks fitted on the host and on the device. - traccc::soa_comparator< - traccc::edm::track_collection> - compare_track_fits{ - "track fits", - traccc::details::comparator_factory< - traccc::edm::track_collection:: - 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(); - n_spacepoints += spacepoints_per_event.size(); - n_seeds += seeds.size(); - n_measurements_cuda += measurements_per_event_cuda.size(); - n_spacepoints_cuda += spacepoints_per_event_cuda.size(); - n_seeds_cuda += seeds_cuda.size(); - n_found_tracks += track_candidates.tracks.size(); - n_found_tracks_cuda += track_candidates_cuda.size(); - n_ambiguity_free_tracks += res_track_candidates.tracks.size(); - n_ambiguity_free_tracks_cuda += res_track_candidates_cuda.size(); - n_fitted_tracks += track_states.tracks.size(); - n_fitted_tracks_cuda += track_states_cuda.tracks.size(); - - if (performance_opts.run) { - - // TODO: Do evt_data.fill_cca_result(...) with cuda clusters and - // measurements - } - } - - if (performance_opts.run) { - sd_performance_writer.finalize(); - } - - TRACCC_INFO("==> Statistics ... "); - TRACCC_INFO("- read " << n_cells << " cells"); - TRACCC_INFO("- created (cpu) " << n_measurements << " measurements "); - TRACCC_INFO("- created (cuda) " << n_measurements_cuda - << " measurements "); - TRACCC_INFO("- created (cpu) " << n_spacepoints << " spacepoints "); - TRACCC_INFO("- created (cuda) " << n_spacepoints_cuda - << " spacepoints "); - - TRACCC_INFO("- created (cpu) " << n_seeds << " seeds"); - TRACCC_INFO("- created (cuda) " << n_seeds_cuda << " seeds"); - TRACCC_INFO("- found (cpu) " << n_found_tracks << " tracks"); - TRACCC_INFO("- found (cuda) " << n_found_tracks_cuda << " tracks"); - TRACCC_INFO("- resolved (cpu) " << n_ambiguity_free_tracks << " tracks"); - TRACCC_INFO("- resolved (cuda) " << n_ambiguity_free_tracks_cuda - << " tracks"); - TRACCC_INFO("- fitted (cpu) " << n_fitted_tracks << " tracks"); - TRACCC_INFO("- fitted (cuda) " << n_fitted_tracks_cuda << " tracks"); - TRACCC_INFO("==>Elapsed times... " << elapsedTimes); - - return 0; -} - -// The main routine -// -int main(int argc, char* argv[]) { - std::unique_ptr logger = traccc::getDefaultLogger( - "CudaSeqExample", traccc::Logging::Level::INFO); - - // Program options. - traccc::opts::detector detector_opts; - traccc::opts::magnetic_field bfield_opts; - traccc::opts::input_data input_opts; - traccc::opts::clusterization clusterization_opts; - traccc::opts::track_seeding seeding_opts; - traccc::opts::track_finding finding_opts; - traccc::opts::track_propagation propagation_opts; - traccc::opts::track_resolution resolution_opts; - traccc::opts::track_fitting fitting_opts; - traccc::opts::performance performance_opts; - traccc::opts::accelerator accelerator_opts; - traccc::opts::program_options program_opts{ - "Full Tracking Chain Using CUDA", - {detector_opts, bfield_opts, input_opts, clusterization_opts, - seeding_opts, finding_opts, propagation_opts, resolution_opts, - performance_opts, fitting_opts, accelerator_opts}, - argc, - argv, - logger->cloneWithSuffix("Options")}; - - // Run the application. - return seq_run(detector_opts, bfield_opts, input_opts, clusterization_opts, - seeding_opts, finding_opts, propagation_opts, - resolution_opts, fitting_opts, performance_opts, - accelerator_opts, logger->clone()); -} diff --git a/examples/run/cuda/track_finding_validation.cpp b/examples/run/cuda/track_finding_validation.cpp new file mode 100644 index 0000000000..b41744973d --- /dev/null +++ b/examples/run/cuda/track_finding_validation.cpp @@ -0,0 +1,20 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2021-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Project include(s). +#include "../common/device_track_finding_validation.hpp" + +// Local include(s). +#include "device_backend.hpp" + +int main(int argc, char* argv[]) { + + return traccc::device_track_finding_validation< + traccc::cuda::device_backend>("track_finding_validation_cuda", + "CUDA Track Finding Validation", argc, + argv); +} diff --git a/examples/run/sycl/CMakeLists.txt b/examples/run/sycl/CMakeLists.txt index de5816e84d..752b88274a 100644 --- a/examples/run/sycl/CMakeLists.txt +++ b/examples/run/sycl/CMakeLists.txt @@ -13,21 +13,12 @@ include( traccc-compiler-options-sycl ) # We need the SYCL language in this directory. enable_language( SYCL ) -# SYCL seeding executable(s). -traccc_add_executable( seeding_example_sycl "seeding_example_sycl.sycl" - LINK_LIBRARIES traccc::options vecmem::core vecmem::sycl traccc::io - traccc::core traccc::device_common traccc::sycl - traccc::performance traccc_examples_common ) - -traccc_add_executable( seq_example_sycl "seq_example_sycl.sycl" - LINK_LIBRARIES traccc::options vecmem::core vecmem::sycl traccc::io - traccc::core traccc::device_common traccc::sycl - traccc::performance traccc_examples_common ) - # -# Set up the "throughput applications". +# Set up the library used by the applications. # add_library( traccc_examples_sycl OBJECT + "device_backend.hpp" + "device_backend.cpp" "full_chain_algorithm.hpp" "full_chain_algorithm.sycl" ) target_link_libraries( traccc_examples_sycl @@ -35,14 +26,21 @@ target_link_libraries( traccc_examples_sycl traccc::core traccc::device_common traccc::sycl traccc_examples_common ) +# +# Set up the applications. +# traccc_add_executable( throughput_st_sycl "throughput_st.cpp" - LINK_LIBRARIES indicators::indicators vecmem::core vecmem::sycl - detray::detectors detray::io traccc::io traccc::performance - traccc::core traccc::device_common traccc::sycl - traccc::options traccc_examples_sycl ) + LINK_LIBRARIES indicators::indicators traccc_examples_common + traccc_examples_sycl ) traccc_add_executable( throughput_mt_sycl "throughput_mt.cpp" - LINK_LIBRARIES indicators::indicators TBB::tbb vecmem::core vecmem::sycl - detray::detectors detray::io traccc::io traccc::performance - traccc::core traccc::device_common traccc::sycl - traccc::options traccc_examples_sycl ) + LINK_LIBRARIES TBB::tbb indicators::indicators traccc_examples_common + traccc_examples_sycl ) + +traccc_add_executable( track_finding_validation_sycl + "track_finding_validation.cpp" + LINK_LIBRARIES traccc_examples_common traccc_examples_sycl ) + +traccc_add_executable( reconstruction_validation_sycl + "reconstruction_validation.cpp" + LINK_LIBRARIES traccc_examples_common traccc_examples_sycl ) diff --git a/examples/run/sycl/device_backend.cpp b/examples/run/sycl/device_backend.cpp new file mode 100644 index 0000000000..31a7e4cf03 --- /dev/null +++ b/examples/run/sycl/device_backend.cpp @@ -0,0 +1,168 @@ +/** 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 "device_backend.hpp" + +// Project include(s). +#include "traccc/sycl/clusterization/clusterization_algorithm.hpp" +#include "traccc/sycl/clusterization/measurement_sorting_algorithm.hpp" +#include "traccc/sycl/finding/combinatorial_kalman_filter_algorithm.hpp" +#include "traccc/sycl/fitting/kalman_fitting_algorithm.hpp" +#include "traccc/sycl/seeding/seeding_algorithm.hpp" +#include "traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" +#include "traccc/sycl/seeding/track_params_estimation.hpp" +#include "traccc/sycl/utils/make_magnetic_field.hpp" +#include "traccc/sycl/utils/queue_wrapper.hpp" + +// VecMem include(s). +#include +#include +#include +#include + +namespace traccc::sycl { + +struct device_backend::impl { + + /// (VecMem) SYCL queue to use + vecmem::sycl::queue_wrapper m_vecmem_queue; + /// (Traccc) SYCL queue wrapper + traccc::sycl::queue_wrapper m_traccc_queue{m_vecmem_queue.queue()}; + + /// Host memory resource + vecmem::sycl::host_memory_resource m_host_mr{m_vecmem_queue}; + /// Device memory resource + vecmem::sycl::device_memory_resource m_device_mr{m_vecmem_queue}; + /// Traccc memory resource + memory_resource m_mr{m_device_mr, &m_host_mr}; + + /// (Asynchronous) Memory copy object + vecmem::sycl::async_copy m_copy{m_vecmem_queue}; + +}; // struct device_backend::impl + +device_backend::device_backend(std::unique_ptr logger) + : messaging(std::move(logger)), m_impl{std::make_unique()} {} + +device_backend::~device_backend() = default; + +vecmem::copy& device_backend::copy() const { + + return m_impl->m_copy; +} + +memory_resource& device_backend::mr() const { + + return m_impl->m_mr; +} + +void device_backend::synchronize() const { + + m_impl->m_vecmem_queue.synchronize(); +} + +magnetic_field device_backend::make_magnetic_field(const magnetic_field& bfield, + bool) const { + + return sycl::make_magnetic_field(bfield, m_impl->m_traccc_queue); +} + +std::unique_ptr::buffer( + const edm::silicon_cell_collection::const_view&, + const silicon_detector_description::const_view&)>> +device_backend::make_clusterization_algorithm( + const clustering_config& config) const { + + TRACCC_VERBOSE("Constructing sycl::clusterization_algorithm"); + return std::make_unique( + m_impl->m_mr, m_impl->m_copy, m_impl->m_traccc_queue, config, + logger().clone("sycl::clusterization_algorithm")); +} + +std::unique_ptr::buffer( + const edm::measurement_collection::const_view&)>> +device_backend::make_measurement_sorting_algorithm() const { + + TRACCC_VERBOSE("Constructing sycl::measurement_sorting_algorithm"); + return std::make_unique( + m_impl->m_mr, m_impl->m_copy, m_impl->m_traccc_queue, + logger().clone("sycl::measurement_sorting_algorithm")); +} + +std::unique_ptr::const_view&)>> +device_backend::make_spacepoint_formation_algorithm() const { + + TRACCC_VERBOSE( + "Constructing sycl::silicon_pixel_spacepoint_formation_algorithm"); + return std::make_unique( + m_impl->m_mr, m_impl->m_copy, m_impl->m_traccc_queue, + logger().clone("sycl::silicon_pixel_spacepoint_formation_algorithm")); +} + +std::unique_ptr> +device_backend::make_seeding_algorithm( + const seedfinder_config& finder_config, + const spacepoint_grid_config& grid_config, + const seedfilter_config& filter_config) const { + + TRACCC_VERBOSE("Constructing sycl::seeding_algorithm"); + return std::make_unique( + finder_config, grid_config, filter_config, m_impl->m_mr, m_impl->m_copy, + m_impl->m_traccc_queue, logger().clone("sycl::seeding_algorithm")); +} + +std::unique_ptr::const_view&, + const edm::spacepoint_collection::const_view&, + const edm::seed_collection::const_view&, const vector3&)>> +device_backend::make_track_params_estimation_algorithm( + const track_params_estimation_config& config) const { + + TRACCC_VERBOSE("Constructing sycl::track_params_estimation"); + return std::make_unique( + config, m_impl->m_mr, m_impl->m_copy, m_impl->m_traccc_queue, + logger().clone("sycl::track_params_estimation")); +} + +std::unique_ptr::buffer( + const detector_buffer&, const magnetic_field&, + const edm::measurement_collection::const_view&, + const bound_track_parameters_collection_types::const_view&)>> +device_backend::make_finding_algorithm(const finding_config& config) const { + + TRACCC_VERBOSE("Constructing sycl::combinatorial_kalman_filter_algorithm"); + return std::make_unique( + config, m_impl->m_mr, m_impl->m_copy, m_impl->m_traccc_queue, + logger().clone("sycl::combinatorial_kalman_filter_algorithm")); +} + +std::unique_ptr::buffer( + const edm::track_container::const_view&)>> +device_backend::make_ambiguity_resolution_algorithm( + const ambiguity_resolution_config&) const { + + TRACCC_DEBUG( + "No ambiguity resolution algorithm implemented for the SYCL backend"); + return {}; +} + +std::unique_ptr::buffer( + const detector_buffer&, const magnetic_field&, + const edm::track_container::const_view&)>> +device_backend::make_fitting_algorithm(const fitting_config& config) const { + + TRACCC_VERBOSE("Constructing sycl::kalman_fitting_algorithm"); + return std::make_unique( + config, m_impl->m_mr, m_impl->m_copy, m_impl->m_traccc_queue, + logger().clone("sycl::kalman_fitting_algorithm")); +} + +} // namespace traccc::sycl diff --git a/examples/run/sycl/device_backend.hpp b/examples/run/sycl/device_backend.hpp new file mode 100644 index 0000000000..09d107510b --- /dev/null +++ b/examples/run/sycl/device_backend.hpp @@ -0,0 +1,114 @@ +/** 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 "../common/device_backend.hpp" +#include "traccc/utils/messaging.hpp" + +// System include(s). +#include + +namespace traccc::sycl { + +/// SYCL Device Backend +class device_backend : public traccc::device_backend, public messaging { + + public: + /// Constructor + /// + /// @param logger The logger to use + /// + device_backend( + std::unique_ptr logger = getDummyLogger().clone()); + /// Destructor + ~device_backend(); + + /// @name Function(s) implemented from @c traccc::device_backend + /// @{ + + /// Access a copy object for the used device + vecmem::copy& copy() const override; + + /// Get the memory resource(s) used by the algorithms + memory_resource& mr() const override; + + /// Wait for the used device to finish all scheduled operations + void synchronize() const override; + + /// Set up the magnetic field for the device + magnetic_field make_magnetic_field( + const magnetic_field& bfield, + bool texture_memory = false) const override; + + /// Construct a clusterization algorithm instance + std::unique_ptr< + algorithm::buffer( + const edm::silicon_cell_collection::const_view&, + const silicon_detector_description::const_view&)>> + make_clusterization_algorithm( + const clustering_config& config) const override; + + /// Construct a measurement sorting algorithm instance + std::unique_ptr< + algorithm::buffer( + const edm::measurement_collection::const_view&)>> + make_measurement_sorting_algorithm() const override; + + /// Construct a spacepoint formation algorithm instance + std::unique_ptr::const_view&)>> + make_spacepoint_formation_algorithm() const override; + + /// Construct a seeding algorithm instance + std::unique_ptr> + make_seeding_algorithm( + const seedfinder_config& finder_config, + const spacepoint_grid_config& grid_config, + const seedfilter_config& filter_config) const override; + + /// Construct a track parameter estimation algorithm instance + std::unique_ptr::const_view&, + const edm::spacepoint_collection::const_view&, + const edm::seed_collection::const_view&, const vector3&)>> + make_track_params_estimation_algorithm( + const track_params_estimation_config& config) const override; + + /// Construct a track finding algorithm instance + std::unique_ptr::buffer( + const detector_buffer&, const magnetic_field&, + const edm::measurement_collection::const_view&, + const bound_track_parameters_collection_types::const_view&)>> + make_finding_algorithm(const finding_config& config) const override; + + /// Construct an ambiguity resolution algorithm instance + std::unique_ptr::buffer( + const edm::track_container::const_view&)>> + make_ambiguity_resolution_algorithm( + const ambiguity_resolution_config& config) const override; + + /// Construct a track fitting algorithm instance + std::unique_ptr::buffer( + const detector_buffer&, const magnetic_field&, + const edm::track_container::const_view&)>> + make_fitting_algorithm(const fitting_config& config) const override; + + /// @} + + private: + /// Implementation class + struct impl; + /// PIMPL data object + std::unique_ptr m_impl; + +}; // class device_backend + +} // namespace traccc::sycl diff --git a/examples/run/sycl/reconstruction_validation.cpp b/examples/run/sycl/reconstruction_validation.cpp new file mode 100644 index 0000000000..93257359b4 --- /dev/null +++ b/examples/run/sycl/reconstruction_validation.cpp @@ -0,0 +1,20 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2021-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Project include(s). +#include "../common/device_reconstruction_validation.hpp" + +// Local include(s). +#include "device_backend.hpp" + +int main(int argc, char* argv[]) { + + return traccc::device_reconstruction_validation< + traccc::sycl::device_backend>("reconstruction_validation_sycl", + "SYCL Reconstruction Validation", argc, + argv); +} diff --git a/examples/run/sycl/seeding_example_sycl.sycl b/examples/run/sycl/seeding_example_sycl.sycl deleted file mode 100644 index f0a9687cb3..0000000000 --- a/examples/run/sycl/seeding_example_sycl.sycl +++ /dev/null @@ -1,316 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2021-2025 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// SYCL include(s) -#include - -// core -#include "traccc/geometry/detector.hpp" -#include "traccc/geometry/detector_buffer.hpp" -#include "traccc/geometry/host_detector.hpp" -#include "traccc/utils/propagation.hpp" - -// algorithms -#include "traccc/seeding/seeding_algorithm.hpp" -#include "traccc/seeding/track_params_estimation.hpp" -#include "traccc/sycl/seeding/seeding_algorithm.hpp" -#include "traccc/sycl/seeding/track_params_estimation.hpp" - -// io -#include "traccc/io/read_detector.hpp" -#include "traccc/io/read_spacepoints.hpp" -#include "traccc/io/utils.hpp" - -// performance -#include "traccc/efficiency/seeding_performance_writer.hpp" -#include "traccc/performance/collection_comparator.hpp" -#include "traccc/performance/soa_comparator.hpp" -#include "traccc/performance/timer.hpp" - -// options -#include "traccc/options/accelerator.hpp" -#include "traccc/options/detector.hpp" -#include "traccc/options/input_data.hpp" -#include "traccc/options/performance.hpp" -#include "traccc/options/program_options.hpp" -#include "traccc/options/track_seeding.hpp" - -// Vecmem include(s) -#include -#include -#include -#include -#include - -// System include(s). -#include -#include -#include - -int seq_run(const traccc::opts::detector& detector_opts, - const traccc::opts::track_seeding& seeding_opts, - const traccc::opts::input_data& input_opts, - const traccc::opts::performance& performance_opts, - const traccc::opts::accelerator& accelerator_opts, - std::unique_ptr ilogger) { - TRACCC_LOCAL_LOGGER(std::move(ilogger)); - - // Creating sycl queue object - ::sycl::queue q; - TRACCC_INFO("Running on device: " - << q.get_device().get_info<::sycl::info::device::name>()); - - // Memory resources used by the application. - vecmem::host_memory_resource host_mr; - vecmem::sycl::host_memory_resource sycl_host_mr{&q}; - vecmem::sycl::shared_memory_resource shared_mr{&q}; - vecmem::sycl::device_memory_resource device_mr{&q}; - traccc::memory_resource mr{device_mr, &sycl_host_mr}; - - // Copy object for asynchronous data transfers. - vecmem::sycl::async_copy copy{&q}; - - // Performance writer - traccc::seeding_performance_writer sd_performance_writer( - traccc::seeding_performance_writer::config{}, - logger().clone("SeedingPerformanceWriter")); - - // Output stats - uint64_t n_spacepoints = 0; - uint64_t n_seeds = 0; - uint64_t n_seeds_sycl = 0; - - /***************************** - * Build a geometry - *****************************/ - - // Construct a Detray detector object, if supported by the configuration. - traccc::host_detector host_det; - traccc::io::read_detector(host_det, host_mr, detector_opts.detector_file, - detector_opts.material_file, - detector_opts.grid_file); - - const traccc::detector_buffer detector_buffer = - traccc::buffer_from_host_detector(host_det, device_mr, copy); - q.wait_and_throw(); - - const traccc::vector3 field_vec(seeding_opts); - - // Seeding algorithm - const traccc::seedfinder_config seedfinder_config(seeding_opts); - const traccc::seedfilter_config seedfilter_config(seeding_opts); - const traccc::spacepoint_grid_config spacepoint_grid_config(seeding_opts); - traccc::host::seeding_algorithm sa( - seedfinder_config, spacepoint_grid_config, seedfilter_config, host_mr, - logger().clone("HostSeedingAlg")); - const traccc::track_params_estimation_config track_params_estimation_config; - traccc::host::track_params_estimation tp( - track_params_estimation_config, host_mr, - logger().clone("HostTrackParEstAlg")); - - traccc::sycl::seeding_algorithm sa_sycl{seedfinder_config, - spacepoint_grid_config, - seedfilter_config, - mr, - copy, - &q, - logger().clone("SyclSeedingAlg")}; - traccc::sycl::track_params_estimation tp_sycl{ - track_params_estimation_config, mr, copy, &q, - logger().clone("SyclTrackParEstAlg")}; - - traccc::performance::timing_info elapsedTimes; - - // Loop over events - for (std::size_t event = input_opts.skip; - event < input_opts.events + input_opts.skip; ++event) { - - // Instantiate host containers/collections - traccc::edm::measurement_collection::host - measurements_per_event{host_mr}; - traccc::edm::spacepoint_collection::host spacepoints_per_event{host_mr}; - traccc::host::seeding_algorithm::output_type seeds{host_mr}; - traccc::host::track_params_estimation::output_type params{&host_mr}; - - // Instantiate sycl containers/collections - traccc::edm::seed_collection::buffer seeds_sycl_buffer; - traccc::bound_track_parameters_collection_types::buffer - params_sycl_buffer(0, *mr.host); - - { // Start measuring wall time - traccc::performance::timer wall_t("Wall time", elapsedTimes); - - /*----------------- - hit file reading - -----------------*/ - - { - traccc::performance::timer t("Hit reading (cpu)", - elapsedTimes); - // Read the hits from the relevant event file - traccc::io::read_spacepoints( - spacepoints_per_event, measurements_per_event, event, - input_opts.directory, - (input_opts.use_acts_geom_source ? &host_det : nullptr), - nullptr, input_opts.format); - - } // stop measuring hit reading timer - - /*---------------------------- - Seeding algorithm - ----------------------------*/ - - /// SYCL - - // Copy the measurements and spacepoint and module data to the - // device. - traccc::edm::measurement_collection::buffer - measurements_sycl_buffer( - static_cast(measurements_per_event.size()), - mr.main); - copy(vecmem::get_data(measurements_per_event), - measurements_sycl_buffer) - ->wait(); - traccc::edm::spacepoint_collection::buffer spacepoints_sycl_buffer( - static_cast(spacepoints_per_event.size()), - mr.main); - copy(vecmem::get_data(spacepoints_per_event), - spacepoints_sycl_buffer) - ->wait(); - - { - traccc::performance::timer t("Seeding (sycl)", elapsedTimes); - // Reconstruct the spacepoints into seeds. - seeds_sycl_buffer = sa_sycl(spacepoints_sycl_buffer); - } // stop measuring seeding sycl timer - - // CPU - - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Seeding (cpu)", elapsedTimes); - seeds = sa(vecmem::get_data(spacepoints_per_event)); - } // stop measuring seeding cpu timer - - /*---------------------------- - Track params estimation - ----------------------------*/ - - // SYCL - - { - traccc::performance::timer t("Track params (sycl)", - elapsedTimes); - params_sycl_buffer = - tp_sycl(measurements_sycl_buffer, spacepoints_sycl_buffer, - seeds_sycl_buffer, field_vec); - } // stop measuring track params sycl timer - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Track params (cpu)", - elapsedTimes); - params = tp(vecmem::get_data(measurements_per_event), - vecmem::get_data(spacepoints_per_event), - vecmem::get_data(seeds), field_vec); - } // stop measuring track params cpu timer - - } // Stop measuring wall time - - /*---------------------------------- - compare seeds from cpu and sycl - ----------------------------------*/ - - // Copy the seeds to the host for comparison. - traccc::edm::seed_collection::host seeds_sycl{host_mr}; - traccc::bound_track_parameters_collection_types::host params_sycl{ - &host_mr}; - copy(seeds_sycl_buffer, seeds_sycl)->wait(); - copy(params_sycl_buffer, params_sycl)->wait(); - - if (accelerator_opts.compare_with_cpu) { - // Show which event we are currently presenting the results for. - TRACCC_INFO("===>>> Event " << event << " <<<==="); - - // Compare the seeds made on the host and on the device - traccc::soa_comparator compare_seeds{ - "seeds", traccc::details::comparator_factory< - traccc::edm::seed_collection::const_device:: - const_proxy_type>{ - vecmem::get_data(spacepoints_per_event), - vecmem::get_data(spacepoints_per_event)}}; - compare_seeds(vecmem::get_data(seeds), - vecmem::get_data(seeds_sycl)); - - // Compare the track parameters made on the host and on the device. - traccc::collection_comparator> - compare_track_parameters{"track parameters"}; - compare_track_parameters(vecmem::get_data(params), - vecmem::get_data(params_sycl)); - } - - /*---------------- - Statistics - ---------------*/ - - n_spacepoints += spacepoints_per_event.size(); - n_seeds_sycl += seeds_sycl.size(); - n_seeds += seeds.size(); - - /*------------ - Writer - ------------*/ - - if (performance_opts.run) { - - traccc::event_data evt_data(input_opts.directory, event, host_mr, - input_opts.use_acts_geom_source, - &host_det, input_opts.format, false); - - sd_performance_writer.write( - vecmem::get_data(seeds_sycl), - vecmem::get_data(spacepoints_per_event), - vecmem::get_data(measurements_per_event), evt_data); - } - } - - if (performance_opts.run) { - sd_performance_writer.finalize(); - } - - TRACCC_INFO("==> Statistics ... "); - TRACCC_INFO("- read " << n_spacepoints << " spacepoints"); - TRACCC_INFO("- created (cpu) " << n_seeds << " seeds"); - TRACCC_INFO("- created (sycl) " << n_seeds_sycl << " seeds"); - TRACCC_INFO("==>Elapsed times... " << elapsedTimes); - - return 0; -} - -// The main routine -// -int main(int argc, char* argv[]) { - std::unique_ptr logger = traccc::getDefaultLogger( - "TracccExampleSeedingSycl", traccc::Logging::Level::INFO); - - // Program options. - traccc::opts::detector detector_opts; - traccc::opts::input_data input_opts; - traccc::opts::track_seeding seeding_opts; - traccc::opts::performance performance_opts; - traccc::opts::accelerator accelerator_opts; - traccc::opts::program_options program_opts{ - "Full Tracking Chain Using SYCL (without clusterization)", - {detector_opts, input_opts, seeding_opts, performance_opts, - accelerator_opts}, - argc, - argv, - logger->cloneWithSuffix("Options")}; - - // Run the application. - return seq_run(detector_opts, seeding_opts, input_opts, performance_opts, - accelerator_opts, logger->clone()); -} diff --git a/examples/run/sycl/seq_example_sycl.sycl b/examples/run/sycl/seq_example_sycl.sycl deleted file mode 100644 index 6d1f6e7fdb..0000000000 --- a/examples/run/sycl/seq_example_sycl.sycl +++ /dev/null @@ -1,490 +0,0 @@ -/* TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2021-2025 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// SYCL include(s) -#include - -// core -#include "traccc/geometry/detector.hpp" - -// io -#include "traccc/io/read_cells.hpp" -#include "traccc/io/read_detector.hpp" -#include "traccc/io/read_detector_description.hpp" -#include "traccc/io/utils.hpp" - -// algorithms -#include "traccc/clusterization/clusterization_algorithm.hpp" -#include "traccc/device/container_d2h_copy_alg.hpp" -#include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" -#include "traccc/seeding/seeding_algorithm.hpp" -#include "traccc/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" -#include "traccc/seeding/track_params_estimation.hpp" -#include "traccc/sycl/clusterization/clusterization_algorithm.hpp" -#include "traccc/sycl/clusterization/measurement_sorting_algorithm.hpp" -#include "traccc/sycl/finding/combinatorial_kalman_filter_algorithm.hpp" -#include "traccc/sycl/seeding/seeding_algorithm.hpp" -#include "traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" -#include "traccc/sycl/seeding/track_params_estimation.hpp" -#include "traccc/sycl/utils/make_magnetic_field.hpp" - -// performance -#include "traccc/efficiency/seeding_performance_writer.hpp" -#include "traccc/performance/collection_comparator.hpp" -#include "traccc/performance/soa_comparator.hpp" -#include "traccc/performance/timer.hpp" - -// options -#include "traccc/options/accelerator.hpp" -#include "traccc/options/clusterization.hpp" -#include "traccc/options/detector.hpp" -#include "traccc/options/input_data.hpp" -#include "traccc/options/magnetic_field.hpp" -#include "traccc/options/performance.hpp" -#include "traccc/options/program_options.hpp" -#include "traccc/options/track_finding.hpp" -#include "traccc/options/track_fitting.hpp" -#include "traccc/options/track_propagation.hpp" -#include "traccc/options/track_seeding.hpp" - -// examples -#include "../common/make_magnetic_field.hpp" - -// Vecmem include(s) -#include -#include -#include -#include -#include - -// Project include(s). -#include "traccc/utils/memory_resource.hpp" - -// System include(s). -#include -#include -#include -#include - -// Simple asynchronous handler function -class handle_async_error { - public: - handle_async_error(const traccc::Logger& l) : logger(l) {} - - auto operator()(::sycl::exception_list elist) { - for (auto& e : elist) { - try { - std::rethrow_exception(e); - } catch (::sycl::exception& e) { - TRACCC_ERROR("Asynchronous exception: " << e.what()); - } - } - } - - private: - const traccc::Logger& logger; -}; - -int seq_run(const traccc::opts::detector& detector_opts, - const traccc::opts::magnetic_field& bfield_opts, - const traccc::opts::input_data& input_opts, - const traccc::opts::clusterization& clusterization_opts, - const traccc::opts::track_seeding& seeding_opts, - const traccc::opts::track_finding& finding_opts, - const traccc::opts::track_propagation& propagation_opts, - const traccc::opts::track_fitting& /*fitting_opts*/, - const traccc::opts::performance& performance_opts, - const traccc::opts::accelerator& accelerator_opts, - std::unique_ptr ilogger) { - TRACCC_LOCAL_LOGGER(std::move(ilogger)); - - // Creating SYCL queue object - ::sycl::queue q(handle_async_error{logger()}); - traccc::sycl::queue_wrapper queue{&q}; - TRACCC_INFO("Running on device: " - << q.get_device().get_info<::sycl::info::device::name>()); - - // Memory resources used by the application. - vecmem::host_memory_resource host_mr; - vecmem::sycl::host_memory_resource sycl_host_mr{&q}; - vecmem::sycl::device_memory_resource device_mr{&q}; - traccc::memory_resource mr{device_mr, &sycl_host_mr}; - - // Copy object for asynchronous data transfers. - vecmem::sycl::async_copy copy{&q}; - - // Construct the detector description object. - traccc::silicon_detector_description::host host_det_descr{host_mr}; - traccc::io::read_detector_description( - host_det_descr, detector_opts.detector_file, - detector_opts.digitization_file, traccc::data_format::json); - traccc::silicon_detector_description::data host_det_descr_data{ - vecmem::get_data(host_det_descr)}; - traccc::silicon_detector_description::buffer device_det_descr{ - static_cast( - host_det_descr.size()), - device_mr}; - copy(host_det_descr_data, device_det_descr)->wait(); - - // Construct a Detray detector object, if supported by the configuration. - traccc::host_detector host_det; - traccc::io::read_detector(host_det, host_mr, detector_opts.detector_file, - detector_opts.material_file, - detector_opts.grid_file); - - const traccc::detector_buffer detector_buffer = - traccc::buffer_from_host_detector(host_det, device_mr, copy); - - // Output stats - uint64_t n_cells = 0; - // uint64_t n_clusters = 0; - uint64_t n_measurements = 0; - uint64_t n_spacepoints = 0; - uint64_t n_spacepoints_sycl = 0; - uint64_t n_seeds = 0; - uint64_t n_seeds_sycl = 0; - uint64_t n_found_tracks = 0; - uint64_t n_found_tracks_sycl = 0; - - // Constant B field for the track finding and fitting - const traccc::vector3 field_vec(seeding_opts); - const auto host_field = traccc::details::make_magnetic_field(bfield_opts); - const auto device_field = - traccc::sycl::make_magnetic_field(host_field, queue); - - // Algorithm configuration(s). - const traccc::seedfinder_config seedfinder_config(seeding_opts); - const traccc::seedfilter_config seedfilter_config(seeding_opts); - const traccc::spacepoint_grid_config spacepoint_grid_config(seeding_opts); - - detray::propagation::config propagation_config(propagation_opts); - - traccc::finding_config finding_cfg(finding_opts); - finding_cfg.propagation = propagation_config; - - // Algorithms. - traccc::host::clusterization_algorithm ca( - host_mr, logger().clone("HostClusteringAlg")); - traccc::host::silicon_pixel_spacepoint_formation_algorithm sf( - host_mr, logger().clone("HostSpFormationAlg")); - traccc::host::seeding_algorithm sa( - seedfinder_config, spacepoint_grid_config, seedfilter_config, host_mr, - logger().clone("HostSeedingAlg")); - traccc::track_params_estimation_config track_params_estimation_config; - traccc::host::track_params_estimation tp( - track_params_estimation_config, host_mr, - logger().clone("HostTrackParEstAlg")); - traccc::host::combinatorial_kalman_filter_algorithm finding_alg{ - finding_cfg, host_mr, logger().clone("HostFindingAlg")}; - - traccc::sycl::clusterization_algorithm ca_sycl( - mr, copy, queue, clusterization_opts, - logger().clone("SyclClusteringAlg")); - traccc::sycl::measurement_sorting_algorithm ms_sycl( - mr, copy, queue, logger().clone("SyclMeasSortingAlg")); - traccc::sycl::silicon_pixel_spacepoint_formation_algorithm sf_sycl( - mr, copy, queue, logger().clone("SyclSpFormationAlg")); - traccc::sycl::seeding_algorithm sa_sycl( - seedfinder_config, spacepoint_grid_config, seedfilter_config, mr, copy, - &q, logger().clone("SyclSeedingAlg")); - traccc::sycl::track_params_estimation tp_sycl( - track_params_estimation_config, mr, copy, &q, - logger().clone("SyclTrackParEstAlg")); - traccc::sycl::combinatorial_kalman_filter_algorithm finding_alg_sycl{ - finding_cfg, mr, copy, queue, logger().clone("SyclFindingAlg")}; - - // performance writer - traccc::seeding_performance_writer sd_performance_writer( - traccc::seeding_performance_writer::config{}, - logger().clone("SeedingPerformanceWriter")); - - traccc::performance::timing_info elapsedTimes; - - // Loop over events - for (std::size_t event = input_opts.skip; - event < input_opts.events + input_opts.skip; ++event) { - // Instantiate host containers/collections - traccc::host::clusterization_algorithm::output_type - measurements_per_event{host_mr}; - traccc::host::silicon_pixel_spacepoint_formation_algorithm::output_type - spacepoints_per_event{host_mr}; - traccc::host::seeding_algorithm::output_type seeds{host_mr}; - traccc::host::track_params_estimation::output_type params{&host_mr}; - traccc::host::combinatorial_kalman_filter_algorithm::output_type - track_candidates{host_mr}; - - // Instantiate SYCL containers/collections - traccc::sycl::clusterization_algorithm::output_type - measurements_sycl_buffer; - traccc::sycl::silicon_pixel_spacepoint_formation_algorithm::output_type - spacepoints_sycl_buffer; - traccc::sycl::seeding_algorithm::output_type seeds_sycl_buffer; - traccc::sycl::track_params_estimation::output_type params_sycl_buffer( - 0, *mr.host); - traccc::sycl::combinatorial_kalman_filter_algorithm::output_type - track_candidates_sycl_buffer; - - { - traccc::performance::timer wall_t("Wall time", elapsedTimes); - - traccc::edm::silicon_cell_collection::host cells_per_event{host_mr}; - - { - traccc::performance::timer t("File reading (cpu)", - elapsedTimes); - // Read the cells from the relevant event file into host memory. - static constexpr bool DEDUPLICATE = true; - traccc::io::read_cells( - cells_per_event, event, input_opts.directory, - logger().clone(), &host_det_descr, input_opts.format, - DEDUPLICATE, input_opts.use_acts_geom_source); - } // stop measuring file reading timer - - n_cells += cells_per_event.size(); - - // Create device copy of input collections - traccc::edm::silicon_cell_collection::buffer cells_buffer( - static_cast(cells_per_event.size()), mr.main); - copy(vecmem::get_data(cells_per_event), cells_buffer)->wait(); - - // SYCL - { - traccc::performance::timer t("Clusterization (sycl)", - elapsedTimes); - // Reconstruct it into spacepoints on the device. - auto unsorted_measurements = - ca_sycl(cells_buffer, device_det_descr); - measurements_sycl_buffer = ms_sycl(unsorted_measurements); - q.wait_and_throw(); - } // stop measuring clusterization sycl timer - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Clusterization (cpu)", - elapsedTimes); - measurements_per_event = - ca(vecmem::get_data(cells_per_event), host_det_descr_data); - } - - // Perform seeding, track finding and fitting only when using a - // Detray geometry. - // SYCL - { - traccc::performance::timer t("Spacepoint formation (sycl)", - elapsedTimes); - // Reconstruct it into spacepoints on the device. - spacepoints_sycl_buffer = - sf_sycl(detector_buffer, measurements_sycl_buffer); - q.wait_and_throw(); - } // stop measuring clusterization sycl timer - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Spacepoint formation (cpu)", - elapsedTimes); - spacepoints_per_event = - sf(host_det, vecmem::get_data(measurements_per_event)); - } - - // SYCL - { - traccc::performance::timer t("Seeding (sycl)", elapsedTimes); - seeds_sycl_buffer = sa_sycl(spacepoints_sycl_buffer); - q.wait_and_throw(); - } // stop measuring seeding sycl timer - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Seeding (cpu)", elapsedTimes); - seeds = sa(vecmem::get_data(spacepoints_per_event)); - } // stop measuring seeding cpu timer - - // SYCL - { - traccc::performance::timer t("Track params (sycl)", - elapsedTimes); - params_sycl_buffer = - tp_sycl(measurements_sycl_buffer, spacepoints_sycl_buffer, - seeds_sycl_buffer, field_vec); - q.wait_and_throw(); - } // stop measuring track params timer - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer t("Track params (cpu)", - elapsedTimes); - params = tp(vecmem::get_data(measurements_per_event), - vecmem::get_data(spacepoints_per_event), - vecmem::get_data(seeds), field_vec); - } // stop measuring track params cpu timer - - // SYCL - { - traccc::performance::timer timer{"Track finding (sycl)", - elapsedTimes}; - track_candidates_sycl_buffer = finding_alg_sycl( - detector_buffer, device_field, measurements_sycl_buffer, - params_sycl_buffer); - q.wait_and_throw(); - } - - // CPU - if (accelerator_opts.compare_with_cpu) { - traccc::performance::timer timer{"Track finding (cpu)", - elapsedTimes}; - track_candidates = - finding_alg(host_det, host_field, - vecmem::get_data(measurements_per_event), - vecmem::get_data(params)); - } - } // stop measuring wall time - - /*---------------------------------- - compare cpu and sycl result - ----------------------------------*/ - - traccc::edm::measurement_collection::host - measurements_per_event_sycl{host_mr}; - traccc::edm::spacepoint_collection::host spacepoints_per_event_sycl{ - host_mr}; - traccc::edm::seed_collection::host seeds_sycl{host_mr}; - traccc::bound_track_parameters_collection_types::host params_sycl{ - &host_mr}; - traccc::edm::track_collection::host - track_candidates_sycl{host_mr}; - - copy(measurements_sycl_buffer, measurements_per_event_sycl)->wait(); - copy(spacepoints_sycl_buffer, spacepoints_per_event_sycl)->wait(); - copy(seeds_sycl_buffer, seeds_sycl)->wait(); - copy(params_sycl_buffer, params_sycl)->wait(); - copy(track_candidates_sycl_buffer.tracks, track_candidates_sycl, - vecmem::copy::type::device_to_host) - ->wait(); - - if (accelerator_opts.compare_with_cpu) { - - // Show which event we are currently presenting the results for. - TRACCC_INFO("===>>> Event " << event << " <<<==="); - - // Compare the measurements made on the host and on the device. - traccc::soa_comparator< - traccc::edm::measurement_collection> - compare_measurements{"measurements"}; - compare_measurements(vecmem::get_data(measurements_per_event), - vecmem::get_data(measurements_per_event_sycl)); - - // Compare the spacepoints made on the host and on the device. - traccc::soa_comparator - compare_spacepoints{"spacepoints"}; - compare_spacepoints(vecmem::get_data(spacepoints_per_event), - vecmem::get_data(spacepoints_per_event_sycl)); - - // Compare the seeds made on the host and on the device - traccc::soa_comparator compare_seeds{ - "seeds", traccc::details::comparator_factory< - traccc::edm::seed_collection::const_device:: - const_proxy_type>{ - vecmem::get_data(spacepoints_per_event), - vecmem::get_data(spacepoints_per_event_sycl)}}; - compare_seeds(vecmem::get_data(seeds), - vecmem::get_data(seeds_sycl)); - - // Compare the track parameters made on the host and on the device. - traccc::collection_comparator> - compare_track_parameters{"track parameters"}; - compare_track_parameters(vecmem::get_data(params), - vecmem::get_data(params_sycl)); - - // Compare tracks found on the host and on the device. - traccc::soa_comparator< - traccc::edm::track_collection> - compare_track_candidates{ - "track candidates", - traccc::details::comparator_factory< - traccc::edm::track_collection:: - const_device::const_proxy_type>{ - vecmem::get_data(measurements_per_event), - vecmem::get_data(measurements_per_event_sycl), - {}, - {}}}; - compare_track_candidates(vecmem::get_data(track_candidates.tracks), - vecmem::get_data(track_candidates_sycl)); - } - - /// Statistics - n_measurements += measurements_per_event.size(); - n_spacepoints += spacepoints_per_event.size(); - n_spacepoints_sycl += spacepoints_per_event_sycl.size(); - n_seeds_sycl += seeds_sycl.size(); - n_seeds += seeds.size(); - n_found_tracks += track_candidates.tracks.size(); - n_found_tracks_sycl += track_candidates_sycl.size(); - - if (performance_opts.run) { - - traccc::event_data evt_data(input_opts.directory, event, host_mr, - input_opts.use_acts_geom_source, - &host_det, input_opts.format, true); - - sd_performance_writer.write( - vecmem::get_data(seeds_sycl), - vecmem::get_data(spacepoints_per_event), - vecmem::get_data(measurements_per_event), evt_data); - } - } - - if (performance_opts.run) { - sd_performance_writer.finalize(); - } - - TRACCC_INFO("==> Statistics ... "); - TRACCC_INFO("- read " << n_cells << " cells"); - TRACCC_INFO("- created (cpu) " << n_measurements << " measurements"); - TRACCC_INFO("- created (cpu) " << n_spacepoints << " spacepoints"); - TRACCC_INFO("- created (sycl) " << n_spacepoints_sycl - << " spacepoints "); - - TRACCC_INFO("- created (cpu) " << n_seeds << " seeds"); - TRACCC_INFO("- created (sycl) " << n_seeds_sycl << " seeds"); - TRACCC_INFO("- found (cpu) " << n_found_tracks << " tracks"); - TRACCC_INFO("- found (sycl) " << n_found_tracks_sycl << " tracks"); - TRACCC_INFO("==>Elapsed times...\n" << elapsedTimes); - - return 0; -} -// -// The main routine -// -int main(int argc, char* argv[]) { - std::unique_ptr logger = traccc::getDefaultLogger( - "TracccExampleSeqSycl", traccc::Logging::Level::INFO); - - // Program options. - traccc::opts::detector detector_opts; - traccc::opts::magnetic_field bfield_opts; - traccc::opts::input_data input_opts; - traccc::opts::clusterization clusterization_opts; - traccc::opts::track_seeding seeding_opts; - traccc::opts::track_finding finding_opts; - traccc::opts::track_propagation propagation_opts; - traccc::opts::track_fitting fitting_opts; - traccc::opts::performance performance_opts; - traccc::opts::accelerator accelerator_opts; - traccc::opts::program_options program_opts{ - "Full Tracking Chain Using SYCL", - {detector_opts, bfield_opts, input_opts, clusterization_opts, - seeding_opts, finding_opts, propagation_opts, fitting_opts, - performance_opts, accelerator_opts}, - argc, - argv, - logger->cloneWithSuffix("Options")}; - - // Run the application. - return seq_run(detector_opts, bfield_opts, input_opts, clusterization_opts, - seeding_opts, finding_opts, propagation_opts, fitting_opts, - performance_opts, accelerator_opts, logger->clone()); -} diff --git a/examples/run/sycl/track_finding_validation.cpp b/examples/run/sycl/track_finding_validation.cpp new file mode 100644 index 0000000000..b623f6c650 --- /dev/null +++ b/examples/run/sycl/track_finding_validation.cpp @@ -0,0 +1,20 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2021-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Project include(s). +#include "../common/device_track_finding_validation.hpp" + +// Local include(s). +#include "device_backend.hpp" + +int main(int argc, char* argv[]) { + + return traccc::device_track_finding_validation< + traccc::sycl::device_backend>("track_finding_validation_sycl", + "SYCL Track Finding Validation", argc, + argv); +}