From fad2304e978db2257e583a0328b510a497879147 Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Thu, 14 Nov 2024 15:21:54 +0100 Subject: [PATCH 01/12] Add a skeleton for the SYCL track finding algorithm. The code doesn't do anything, and is not used by anything yet. --- device/sycl/CMakeLists.txt | 6 + .../combinatorial_kalman_filter_algorithm.hpp | 103 ++++++++++++++++++ .../combinatorial_kalman_filter_algorithm.cpp | 18 +++ ...rithm_constant_field_default_detector.sycl | 38 +++++++ ...thm_constant_field_telescope_detector.sycl | 38 +++++++ device/sycl/src/finding/find_tracks.hpp | 65 +++++++++++ 6 files changed, 268 insertions(+) create mode 100644 device/sycl/include/traccc/sycl/finding/combinatorial_kalman_filter_algorithm.hpp create mode 100644 device/sycl/src/finding/combinatorial_kalman_filter_algorithm.cpp create mode 100644 device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_default_detector.sycl create mode 100644 device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_telescope_detector.sycl create mode 100644 device/sycl/src/finding/find_tracks.hpp diff --git a/device/sycl/CMakeLists.txt b/device/sycl/CMakeLists.txt index f3d08fdce..070369eff 100644 --- a/device/sycl/CMakeLists.txt +++ b/device/sycl/CMakeLists.txt @@ -30,6 +30,12 @@ traccc_add_library( traccc_sycl sycl TYPE SHARED "src/seeding/seeding_algorithm.cpp" "include/traccc/sycl/seeding/track_params_estimation.hpp" "src/seeding/track_params_estimation.sycl" + # Track finding algorithm(s). + "include/traccc/sycl/finding/combinatorial_kalman_filter_algorithm.hpp" + "src/finding/combinatorial_kalman_filter_algorithm.cpp" + "src/finding/combinatorial_kalman_filter_algorithm_constant_field_default_detector.sycl" + "src/finding/combinatorial_kalman_filter_algorithm_constant_field_telescope_detector.sycl" + "src/finding/find_tracks.hpp" # Track fitting algorithm(s). "include/traccc/sycl/fitting/kalman_fitting_algorithm.hpp" "src/fitting/kalman_fitting_algorithm.cpp" diff --git a/device/sycl/include/traccc/sycl/finding/combinatorial_kalman_filter_algorithm.hpp b/device/sycl/include/traccc/sycl/finding/combinatorial_kalman_filter_algorithm.hpp new file mode 100644 index 000000000..6552b7174 --- /dev/null +++ b/device/sycl/include/traccc/sycl/finding/combinatorial_kalman_filter_algorithm.hpp @@ -0,0 +1,103 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2024 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// SYCL library include(s). +#include "traccc/sycl/utils/queue_wrapper.hpp" + +// Project include(s). +#include "traccc/edm/measurement.hpp" +#include "traccc/edm/track_candidate.hpp" +#include "traccc/edm/track_parameters.hpp" +#include "traccc/finding/finding_config.hpp" +#include "traccc/geometry/detector.hpp" +#include "traccc/utils/algorithm.hpp" +#include "traccc/utils/memory_resource.hpp" + +// Detray include(s). +#include + +// VecMem include(s). +#include + +// System include(s). +#include + +namespace traccc::sycl { + +/// CKF track finding algorithm +class combinatorial_kalman_filter_algorithm + : public algorithm, + public algorithm { + + public: + /// Configuration type + using config_type = finding_config; + /// Output type + using output_type = track_candidate_container_types::buffer; + + /// Constructor with the algorithm's configuration + explicit combinatorial_kalman_filter_algorithm( + const config_type& config, const traccc::memory_resource& mr, + vecmem::copy& copy, queue_wrapper queue); + + /// Execute the algorithm + /// + /// @param det The (default) detector object + /// @param field The (constant) magnetic field object + /// @param measurements All measurements in an event + /// @param seeds All seeds in an event to start the track finding + /// with + /// + /// @return A container of the found track candidates + /// + output_type operator()( + const default_detector::view& det, + const detray::bfield::const_field_t::view_t& field, + const measurement_collection_types::const_view& measurements, + const bound_track_parameters_collection_types::const_view& seeds) + const override; + + /// Execute the algorithm + /// + /// @param det The (telescope) detector object + /// @param field The (constant) magnetic field object + /// @param measurements All measurements in an event + /// @param seeds All seeds in an event to start the track finding + /// with + /// + /// @return A container of the found track candidates + /// + output_type operator()( + const telescope_detector::view& det, + const detray::bfield::const_field_t::view_t& field, + const measurement_collection_types::const_view& measurements, + const bound_track_parameters_collection_types::const_view& seeds) + const override; + + private: + /// Algorithm configuration + config_type m_config; + /// Memory resource used by the algorithm + traccc::memory_resource m_mr; + /// Copy object used by the algorithm + std::reference_wrapper m_copy; + /// Queue wrapper + mutable queue_wrapper m_queue; + +}; // class combinatorial_kalman_filter_algorithm + +} // namespace traccc::sycl diff --git a/device/sycl/src/finding/combinatorial_kalman_filter_algorithm.cpp b/device/sycl/src/finding/combinatorial_kalman_filter_algorithm.cpp new file mode 100644 index 000000000..79747c2ae --- /dev/null +++ b/device/sycl/src/finding/combinatorial_kalman_filter_algorithm.cpp @@ -0,0 +1,18 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2024 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Local include(s). +#include "traccc/sycl/finding/combinatorial_kalman_filter_algorithm.hpp" + +namespace traccc::sycl { + +combinatorial_kalman_filter_algorithm::combinatorial_kalman_filter_algorithm( + const config_type& config, const traccc::memory_resource& mr, + vecmem::copy& copy, queue_wrapper queue) + : m_config{config}, m_mr{mr}, m_copy{copy}, m_queue{queue} {} + +} // namespace traccc::sycl diff --git a/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_default_detector.sycl b/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_default_detector.sycl new file mode 100644 index 000000000..fe2cb604e --- /dev/null +++ b/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_default_detector.sycl @@ -0,0 +1,38 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2024 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Local include(s). +#include "../utils/get_queue.hpp" +#include "find_tracks.hpp" +#include "traccc/sycl/finding/combinatorial_kalman_filter_algorithm.hpp" + +// Detray include(s). +#include +#include +#include +#include + +namespace traccc::sycl { + +combinatorial_kalman_filter_algorithm::output_type +combinatorial_kalman_filter_algorithm::operator()( + const default_detector::view& det, + const detray::bfield::const_field_t::view_t& field, + const measurement_collection_types::const_view& measurements, + const bound_track_parameters_collection_types::const_view& seeds) const { + + // Perform the track finding using the templated implementation. + return details::find_tracks< + detray::rk_stepper>, + detray::navigator>( + det, field, measurements, seeds, m_config, m_mr, m_copy, + details::get_queue(m_queue)); +} + +} // namespace traccc::sycl diff --git a/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_telescope_detector.sycl b/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_telescope_detector.sycl new file mode 100644 index 000000000..07d150861 --- /dev/null +++ b/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_telescope_detector.sycl @@ -0,0 +1,38 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2024 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Local include(s). +#include "../utils/get_queue.hpp" +#include "find_tracks.hpp" +#include "traccc/sycl/finding/combinatorial_kalman_filter_algorithm.hpp" + +// Detray include(s). +#include +#include +#include +#include + +namespace traccc::sycl { + +combinatorial_kalman_filter_algorithm::output_type +combinatorial_kalman_filter_algorithm::operator()( + const telescope_detector::view& det, + const detray::bfield::const_field_t::view_t& field, + const measurement_collection_types::const_view& measurements, + const bound_track_parameters_collection_types::const_view& seeds) const { + + // Perform the track finding using the templated implementation. + return details::find_tracks< + detray::rk_stepper>, + detray::navigator>( + det, field, measurements, seeds, m_config, m_mr, m_copy, + details::get_queue(m_queue)); +} + +} // namespace traccc::sycl diff --git a/device/sycl/src/finding/find_tracks.hpp b/device/sycl/src/finding/find_tracks.hpp new file mode 100644 index 000000000..d970b93ca --- /dev/null +++ b/device/sycl/src/finding/find_tracks.hpp @@ -0,0 +1,65 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2024 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s). +#include "traccc/edm/measurement.hpp" +#include "traccc/edm/track_candidate.hpp" +#include "traccc/finding/finding_config.hpp" +#include "traccc/utils/memory_resource.hpp" + +// Detray include(s). +#include "detray/propagator/actor_chain.hpp" +#include "detray/propagator/actors/aborters.hpp" +#include "detray/propagator/actors/parameter_resetter.hpp" +#include "detray/propagator/actors/parameter_transporter.hpp" +#include "detray/propagator/actors/pointwise_material_interactor.hpp" +#include "detray/propagator/propagator.hpp" + +// VecMem include(s). +#include + +// SYCL include(s). +#include + +namespace traccc::sycl::details { + +/// Templated implementation of the track finding algorithm. +/// +/// Concrete track finding algorithms can use this function with the appropriate +/// specializations, to find tracks on top of a specific detector type, magnetic +/// field type, and track finding configuration. +/// +/// @tparam stepper_t The stepper type used for the track propagation +/// @tparam navigator_t The navigator type used for the track navigation +/// +/// @param det A view of the detector object +/// @param field The magnetic field object +/// @param measurements_view All measurements in an event +/// @param seeds_view All seeds in an event to start the track finding +/// with +/// @param config The track finding configuration +/// @param mr The memory resource(s) to use +/// @param copy The copy object to use +/// @param queue The SYCL queue to use +/// +/// @return A buffer of the found track candidates +/// +template +track_candidate_container_types::buffer find_tracks( + const typename navigator_t::detector_type::view_type& /*det*/, + const typename stepper_t::magnetic_field_type& /*field*/, + const measurement_collection_types::const_view& /*measurements_view*/, + const bound_track_parameters_collection_types::const_view& /*seeds_view*/, + const finding_config& /*config*/, const memory_resource& /*mr*/, + vecmem::copy& /*copy*/, ::sycl::queue& /*queue*/) { + + return {}; +} + +} // namespace traccc::sycl::details From cc62f071fbd19bdfc343c3baab28c0f53dcfc652 Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Thu, 14 Nov 2024 17:25:31 +0100 Subject: [PATCH 02/12] Small track finding optimizations. Made the "track liveness" buffers use char instead of unsigned int. Since they store boolean information, char is enough. And it also makes memset(...) do a more expected thing on the buffers. Updated device::make_barcode_sequence not to narrow std::size_t into unsigned int. (Which oneAPI doesn't like.) --- .../traccc/finding/device/apply_interaction.hpp | 2 +- .../include/traccc/finding/device/find_tracks.hpp | 4 ++-- .../finding/device/impl/apply_interaction.ipp | 4 ++-- .../traccc/finding/device/impl/find_tracks.ipp | 14 +++++++------- .../finding/device/impl/make_barcode_sequence.ipp | 4 ++-- .../device/impl/propagate_to_next_surface.ipp | 13 ++++++------- .../finding/device/make_barcode_sequence.hpp | 4 ++-- .../finding/device/propagate_to_next_surface.hpp | 2 +- device/cuda/src/finding/finding_algorithm.cu | 6 +++--- .../src/finding/kernels/make_barcode_sequence.cu | 5 ++--- 10 files changed, 28 insertions(+), 30 deletions(-) diff --git a/device/common/include/traccc/finding/device/apply_interaction.hpp b/device/common/include/traccc/finding/device/apply_interaction.hpp index 5650d129e..cee6696ed 100644 --- a/device/common/include/traccc/finding/device/apply_interaction.hpp +++ b/device/common/include/traccc/finding/device/apply_interaction.hpp @@ -36,7 +36,7 @@ struct apply_interaction_payload { * @brief View object to the vector of boolean-like integers describing * whether each parameter is live. Has the same size as \ref params_view */ - vecmem::data::vector_view params_liveness_view; + vecmem::data::vector_view params_liveness_view; }; /// Function applying the Pre material interaction to tracks spawned by bound diff --git a/device/common/include/traccc/finding/device/find_tracks.hpp b/device/common/include/traccc/finding/device/find_tracks.hpp index 7d03ea3db..ecedc8a1e 100644 --- a/device/common/include/traccc/finding/device/find_tracks.hpp +++ b/device/common/include/traccc/finding/device/find_tracks.hpp @@ -43,7 +43,7 @@ struct find_tracks_payload { * @brief View object to the vector of boolean-like integers describing the * liveness of each parameter */ - vecmem::data::vector_view in_params_liveness_view; + vecmem::data::vector_view in_params_liveness_view; /** * @brief The total number of input parameters @@ -84,7 +84,7 @@ struct find_tracks_payload { /** * @brief View object to the output track parameter liveness vector */ - vecmem::data::vector_view out_params_liveness_view; + vecmem::data::vector_view out_params_liveness_view; /** * @brief View object to the output candidate links diff --git a/device/common/include/traccc/finding/device/impl/apply_interaction.ipp b/device/common/include/traccc/finding/device/impl/apply_interaction.ipp index bfbb587cc..d004f1e54 100644 --- a/device/common/include/traccc/finding/device/impl/apply_interaction.ipp +++ b/device/common/include/traccc/finding/device/impl/apply_interaction.ipp @@ -30,7 +30,7 @@ TRACCC_DEVICE inline void apply_interaction( // in param bound_track_parameters_collection_types::device params(payload.params_view); - vecmem::device_vector params_liveness( + vecmem::device_vector params_liveness( payload.params_liveness_view); if (globalIndex >= payload.n_params) { @@ -39,7 +39,7 @@ TRACCC_DEVICE inline void apply_interaction( auto& bound_param = params.at(globalIndex); - if (params_liveness.at(globalIndex) != 0u) { + if (params_liveness.at(globalIndex) != 0) { // Get surface corresponding to bound params const detray::tracking_surface sf{det, bound_param.surface_link()}; const typename detector_t::geometry_context ctx{}; diff --git a/device/common/include/traccc/finding/device/impl/find_tracks.ipp b/device/common/include/traccc/finding/device/impl/find_tracks.ipp index 0a860f77e..9de94f448 100644 --- a/device/common/include/traccc/finding/device/impl/find_tracks.ipp +++ b/device/common/include/traccc/finding/device/impl/find_tracks.ipp @@ -53,13 +53,13 @@ TRACCC_DEVICE inline void find_tracks( payload.measurements_view); bound_track_parameters_collection_types::const_device in_params( payload.in_params_view); - vecmem::device_vector in_params_liveness( + vecmem::device_vector in_params_liveness( payload.in_params_liveness_view); vecmem::device_vector prev_links( payload.prev_links_view); bound_track_parameters_collection_types::device out_params( payload.out_params_view); - vecmem::device_vector out_params_liveness( + vecmem::device_vector out_params_liveness( payload.out_params_liveness_view); vecmem::device_vector links(payload.links_view); vecmem::device_atomic_ref 0u) { + in_params_liveness.at(in_param_id) != 0) { /* * Get the barcode of this thread's parameters, then find the first * measurement that matches it. @@ -185,7 +185,7 @@ TRACCC_DEVICE inline void find_tracks( const unsigned int owner_global_thread_id = owner_local_thread_id + thread_id.getBlockDimX() * thread_id.getBlockIdX(); - assert(in_params_liveness.at(owner_global_thread_id) != 0u); + assert(in_params_liveness.at(owner_global_thread_id) != 0); const bound_track_parameters& in_par = in_params.at(owner_global_thread_id); const unsigned int meas_idx = @@ -236,7 +236,7 @@ TRACCC_DEVICE inline void find_tracks( .fetch_add(1u); out_params.at(l_pos) = trk_state.filtered(); - out_params_liveness.at(l_pos) = 1u; + out_params_liveness.at(l_pos) = 1; } } } @@ -268,7 +268,7 @@ TRACCC_DEVICE inline void find_tracks( * match any measurements. */ if (in_param_id < payload.n_in_params && - in_params_liveness.at(in_param_id) > 0u && + in_params_liveness.at(in_param_id) != 0 && shared_payload.shared_num_candidates[thread_id.getLocalThreadIdX()] == 0u) { // Add measurement candidates to link @@ -292,7 +292,7 @@ TRACCC_DEVICE inline void find_tracks( } out_params.at(l_pos) = in_params.at(in_param_id); - out_params_liveness.at(l_pos) = 1u; + out_params_liveness.at(l_pos) = 1; } } } diff --git a/device/common/include/traccc/finding/device/impl/make_barcode_sequence.ipp b/device/common/include/traccc/finding/device/impl/make_barcode_sequence.ipp index a70b6f684..bcc508a2a 100644 --- a/device/common/include/traccc/finding/device/impl/make_barcode_sequence.ipp +++ b/device/common/include/traccc/finding/device/impl/make_barcode_sequence.ipp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023 CERN for the benefit of the ACTS project + * (c) 2023-2024 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -15,7 +15,7 @@ namespace traccc::device { TRACCC_DEVICE inline void make_barcode_sequence( - std::size_t globalIndex, const make_barcode_sequence_payload& payload) { + unsigned int globalIndex, const make_barcode_sequence_payload& payload) { measurement_collection_types::const_device uniques(payload.uniques_view); vecmem::device_vector barcodes(payload.barcodes_view); diff --git a/device/common/include/traccc/finding/device/impl/propagate_to_next_surface.ipp b/device/common/include/traccc/finding/device/impl/propagate_to_next_surface.ipp index 06ea73949..baec59f0c 100644 --- a/device/common/include/traccc/finding/device/impl/propagate_to_next_surface.ipp +++ b/device/common/include/traccc/finding/device/impl/propagate_to_next_surface.ipp @@ -49,8 +49,7 @@ TRACCC_DEVICE inline void propagate_to_next_surface( n_tracks_per_seed.at(orig_param_id)); const unsigned int s_pos = num_tracks_per_seed.fetch_add(1); - vecmem::device_vector params_liveness( - payload.params_liveness_view); + vecmem::device_vector params_liveness(payload.params_liveness_view); if (s_pos >= cfg.max_num_branches_per_seed) { params_liveness[param_id] = 0u; @@ -62,7 +61,7 @@ TRACCC_DEVICE inline void propagate_to_next_surface( payload.tips_view); if (links.at(param_id).n_skipped > cfg.max_num_skipping_per_cand) { - params_liveness[param_id] = 0u; + params_liveness[param_id] = 0; tips.push_back({payload.step, param_id}); return; } @@ -73,7 +72,7 @@ TRACCC_DEVICE inline void propagate_to_next_surface( // Parameters bound_track_parameters_collection_types::device params(payload.params_view); - if (params_liveness.at(param_id) == 0u) { + if (params_liveness.at(param_id) == 0) { return; } @@ -121,12 +120,12 @@ TRACCC_DEVICE inline void propagate_to_next_surface( if (payload.step == cfg.max_track_candidates_per_track - 1) { tips.push_back({payload.step, param_id}); - params_liveness[param_id] = 0u; + params_liveness[param_id] = 0; } else { - params_liveness[param_id] = 1u; + params_liveness[param_id] = 1; } } else { - params_liveness[param_id] = 0u; + params_liveness[param_id] = 0; if (payload.step >= cfg.min_track_candidates_per_track - 1) { tips.push_back({payload.step, param_id}); diff --git a/device/common/include/traccc/finding/device/make_barcode_sequence.hpp b/device/common/include/traccc/finding/device/make_barcode_sequence.hpp index f7d17c6ed..057d7d92a 100644 --- a/device/common/include/traccc/finding/device/make_barcode_sequence.hpp +++ b/device/common/include/traccc/finding/device/make_barcode_sequence.hpp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023 CERN for the benefit of the ACTS project + * (c) 2023-2024 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -30,7 +30,7 @@ struct make_barcode_sequence_payload { /// @param[in] globalIndex The index of the current thread /// @param[inout] payload The function call payload TRACCC_DEVICE inline void make_barcode_sequence( - std::size_t globalIndex, const make_barcode_sequence_payload& payload); + unsigned int globalIndex, const make_barcode_sequence_payload& payload); } // namespace traccc::device #include "./impl/make_barcode_sequence.ipp" diff --git a/device/common/include/traccc/finding/device/propagate_to_next_surface.hpp b/device/common/include/traccc/finding/device/propagate_to_next_surface.hpp index e88639571..1a03da9b8 100644 --- a/device/common/include/traccc/finding/device/propagate_to_next_surface.hpp +++ b/device/common/include/traccc/finding/device/propagate_to_next_surface.hpp @@ -36,7 +36,7 @@ struct propagate_to_next_surface_payload { /** * @brief View object to the vector of track parameter liveness values */ - vecmem::data::vector_view params_liveness_view; + vecmem::data::vector_view params_liveness_view; /** * @brief View object to the access order of parameters so they are sorted diff --git a/device/cuda/src/finding/finding_algorithm.cu b/device/cuda/src/finding/finding_algorithm.cu index 4fe8f760b..ef6f82b7a 100644 --- a/device/cuda/src/finding/finding_algorithm.cu +++ b/device/cuda/src/finding/finding_algorithm.cu @@ -148,8 +148,7 @@ finding_algorithm::operator()( m_copy.setup(in_params_buffer)->ignore(); m_copy(vecmem::get_data(seeds_buffer), vecmem::get_data(in_params_buffer)) ->ignore(); - vecmem::data::vector_buffer param_liveness_buffer(n_seeds, - m_mr.main); + vecmem::data::vector_buffer param_liveness_buffer(n_seeds, m_mr.main); m_copy.setup(param_liveness_buffer)->ignore(); m_copy.memset(param_liveness_buffer, 1)->ignore(); @@ -215,9 +214,10 @@ finding_algorithm::operator()( m_mr.main); m_copy.setup(updated_params_buffer)->ignore(); - vecmem::data::vector_buffer updated_liveness_buffer( + vecmem::data::vector_buffer updated_liveness_buffer( n_in_params * m_cfg.max_num_branches_per_surface, m_mr.main); m_copy.setup(updated_liveness_buffer)->ignore(); + m_copy.memset(updated_liveness_buffer, 0)->ignore(); // Create the link map link_map[step] = {n_in_params * m_cfg.max_num_branches_per_surface, diff --git a/device/cuda/src/finding/kernels/make_barcode_sequence.cu b/device/cuda/src/finding/kernels/make_barcode_sequence.cu index e6587b553..461bc74bc 100644 --- a/device/cuda/src/finding/kernels/make_barcode_sequence.cu +++ b/device/cuda/src/finding/kernels/make_barcode_sequence.cu @@ -14,8 +14,7 @@ namespace traccc::cuda::kernels { __global__ void make_barcode_sequence( device::make_barcode_sequence_payload payload) { - int gid = threadIdx.x + blockIdx.x * blockDim.x; - - device::make_barcode_sequence(gid, payload); + device::make_barcode_sequence(threadIdx.x + blockIdx.x * blockDim.x, + payload); } } // namespace traccc::cuda::kernels From 1290c0f5ebdae8451765d1203712d850fdd6a2c0 Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Fri, 15 Nov 2024 11:26:49 +0100 Subject: [PATCH 03/12] Further tweaks necessitated by oneAPI. Mainly to avoid type conversion / narrowing in the code, but also to simplify it slightly in some places. --- .../finding/device/apply_interaction.hpp | 6 +++--- .../traccc/finding/device/build_tracks.hpp | 4 ++-- .../traccc/finding/device/find_tracks.hpp | 6 +++--- .../finding/device/impl/apply_interaction.ipp | 4 ++-- .../finding/device/impl/build_tracks.ipp | 7 ++----- .../traccc/finding/device/impl/find_tracks.ipp | 11 +++++++---- .../device/impl/propagate_to_next_surface.ipp | 4 ++-- .../finding/device/impl/prune_tracks.ipp | 2 +- .../device/propagate_to_next_surface.hpp | 7 ++++--- .../traccc/finding/device/prune_tracks.hpp | 2 +- .../include/traccc/cuda/utils/thread_id.hpp | 18 +++++++++--------- device/cuda/src/finding/finding_algorithm.cu | 4 ++-- .../cuda/src/finding/kernels/build_tracks.cu | 2 +- .../cuda/src/finding/kernels/prune_tracks.cu | 2 +- .../specializations/apply_interaction_src.cuh | 2 +- .../specializations/find_tracks_src.cuh | 3 +-- .../propagate_to_next_surface_src.cuh | 6 +++--- 17 files changed, 45 insertions(+), 45 deletions(-) diff --git a/device/common/include/traccc/finding/device/apply_interaction.hpp b/device/common/include/traccc/finding/device/apply_interaction.hpp index cee6696ed..204f4f041 100644 --- a/device/common/include/traccc/finding/device/apply_interaction.hpp +++ b/device/common/include/traccc/finding/device/apply_interaction.hpp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023 CERN for the benefit of the ACTS project + * (c) 2023-2024 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -25,7 +25,7 @@ struct apply_interaction_payload { /** * @brief Total number of input parameters (including non-live ones) */ - const int n_params; + const unsigned int n_params; /** * @brief View object to the vector of bound track parameters @@ -47,7 +47,7 @@ struct apply_interaction_payload { /// @param[inout] payload The function call payload template TRACCC_DEVICE inline void apply_interaction( - std::size_t globalIndex, const finding_config& cfg, + unsigned int globalIndex, const finding_config& cfg, const apply_interaction_payload& payload); } // namespace traccc::device diff --git a/device/common/include/traccc/finding/device/build_tracks.hpp b/device/common/include/traccc/finding/device/build_tracks.hpp index b364d3130..42add735b 100644 --- a/device/common/include/traccc/finding/device/build_tracks.hpp +++ b/device/common/include/traccc/finding/device/build_tracks.hpp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023 CERN for the benefit of the ACTS project + * (c) 2023-2024 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -65,7 +65,7 @@ struct build_tracks_payload { /// @param[in] cfg Track finding config object /// @param[inout] payload The function call payload template -TRACCC_DEVICE inline void build_tracks(std::size_t globalIndex, +TRACCC_DEVICE inline void build_tracks(unsigned int globalIndex, const config_t cfg, const build_tracks_payload& payload); diff --git a/device/common/include/traccc/finding/device/find_tracks.hpp b/device/common/include/traccc/finding/device/find_tracks.hpp index ecedc8a1e..8a9cec972 100644 --- a/device/common/include/traccc/finding/device/find_tracks.hpp +++ b/device/common/include/traccc/finding/device/find_tracks.hpp @@ -129,10 +129,10 @@ struct find_tracks_shared_payload { /// @param[in] cfg Track finding config object /// @param[inout] payload The global memory payload /// @param[inout] shared_payload The shared memory payload -template +template TRACCC_DEVICE inline void find_tracks( - thread_id_t& thread_id, barrier_t& barrier, const config_t cfg, + thread_id_t& thread_id, barrier_t& barrier, const finding_config& cfg, const find_tracks_payload& payload, const find_tracks_shared_payload& shared_payload); } // namespace traccc::device diff --git a/device/common/include/traccc/finding/device/impl/apply_interaction.ipp b/device/common/include/traccc/finding/device/impl/apply_interaction.ipp index d004f1e54..26a276746 100644 --- a/device/common/include/traccc/finding/device/impl/apply_interaction.ipp +++ b/device/common/include/traccc/finding/device/impl/apply_interaction.ipp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023 CERN for the benefit of the ACTS project + * (c) 2023-2024 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -18,7 +18,7 @@ namespace traccc::device { template TRACCC_DEVICE inline void apply_interaction( - std::size_t globalIndex, const finding_config& cfg, + unsigned int globalIndex, const finding_config& cfg, const apply_interaction_payload& payload) { // Type definitions diff --git a/device/common/include/traccc/finding/device/impl/build_tracks.ipp b/device/common/include/traccc/finding/device/impl/build_tracks.ipp index 149d9aa55..d038fa668 100644 --- a/device/common/include/traccc/finding/device/impl/build_tracks.ipp +++ b/device/common/include/traccc/finding/device/impl/build_tracks.ipp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023 CERN for the benefit of the ACTS project + * (c) 2023-2024 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -17,7 +17,7 @@ namespace traccc::device { template -TRACCC_DEVICE inline void build_tracks(std::size_t globalIndex, +TRACCC_DEVICE inline void build_tracks(unsigned int globalIndex, const config_t cfg, const build_tracks_payload& payload) { @@ -73,12 +73,9 @@ TRACCC_DEVICE inline void build_tracks(std::size_t globalIndex, // Resize the candidates with the exact size cands_per_track.resize(n_cands); - unsigned int i = 0; - // Reversely iterate to fill the track candidates for (auto it = cands_per_track.rbegin(); it != cands_per_track.rend(); it++) { - i++; while (L.meas_idx > n_meas && L.previous.first != diff --git a/device/common/include/traccc/finding/device/impl/find_tracks.ipp b/device/common/include/traccc/finding/device/impl/find_tracks.ipp index 9de94f448..e4a9124bb 100644 --- a/device/common/include/traccc/finding/device/impl/find_tracks.ipp +++ b/device/common/include/traccc/finding/device/impl/find_tracks.ipp @@ -25,10 +25,10 @@ namespace traccc::device { -template +template TRACCC_DEVICE inline void find_tracks( - thread_id_t& thread_id, barrier_t& barrier, const config_t cfg, + thread_id_t& thread_id, barrier_t& barrier, const finding_config& cfg, const find_tracks_payload& payload, const find_tracks_shared_payload& shared_payload) { @@ -116,7 +116,10 @@ TRACCC_DEVICE inline void find_tracks( * this thread. */ else { - const auto bcd_id = std::distance(barcodes.begin(), lo); + const vecmem::device_vector::size_type bcd_id = + static_cast< + vecmem::device_vector::size_type>( + std::distance(barcodes.begin(), lo)); init_meas = lo == barcodes.begin() ? 0u : upper_bounds[bcd_id - 1]; num_meas = upper_bounds[bcd_id] - init_meas; diff --git a/device/common/include/traccc/finding/device/impl/propagate_to_next_surface.ipp b/device/common/include/traccc/finding/device/impl/propagate_to_next_surface.ipp index baec59f0c..da4f41c93 100644 --- a/device/common/include/traccc/finding/device/impl/propagate_to_next_surface.ipp +++ b/device/common/include/traccc/finding/device/impl/propagate_to_next_surface.ipp @@ -20,9 +20,9 @@ namespace traccc::device { -template +template TRACCC_DEVICE inline void propagate_to_next_surface( - std::size_t globalIndex, const config_t cfg, + unsigned int globalIndex, const finding_config& cfg, const propagate_to_next_surface_payload& payload) { if (globalIndex >= payload.n_in_params) { diff --git a/device/common/include/traccc/finding/device/impl/prune_tracks.ipp b/device/common/include/traccc/finding/device/impl/prune_tracks.ipp index d9979241f..6f97d0778 100644 --- a/device/common/include/traccc/finding/device/impl/prune_tracks.ipp +++ b/device/common/include/traccc/finding/device/impl/prune_tracks.ipp @@ -14,7 +14,7 @@ namespace traccc::device { -TRACCC_DEVICE inline void prune_tracks(std::size_t globalIndex, +TRACCC_DEVICE inline void prune_tracks(unsigned int globalIndex, const prune_tracks_payload& payload) { track_candidate_container_types::const_device track_candidates( diff --git a/device/common/include/traccc/finding/device/propagate_to_next_surface.hpp b/device/common/include/traccc/finding/device/propagate_to_next_surface.hpp index 1a03da9b8..c17cf3ea4 100644 --- a/device/common/include/traccc/finding/device/propagate_to_next_surface.hpp +++ b/device/common/include/traccc/finding/device/propagate_to_next_surface.hpp @@ -1,6 +1,6 @@ /** TRACCC library, part of the ACTS project (R&D line) * - * (c) 2023 CERN for the benefit of the ACTS project + * (c) 2023-2024 CERN for the benefit of the ACTS project * * Mozilla Public License Version 2.0 */ @@ -13,6 +13,7 @@ #include "traccc/edm/measurement.hpp" #include "traccc/edm/track_parameters.hpp" #include "traccc/finding/candidate_link.hpp" +#include "traccc/finding/finding_config.hpp" #include "traccc/utils/particle.hpp" namespace traccc::device { @@ -81,9 +82,9 @@ struct propagate_to_next_surface_payload { /// @param[in] globalIndex The index of the current thread /// @param[in] cfg Track finding config object /// @param[inout] payload The function call payload -template +template TRACCC_DEVICE inline void propagate_to_next_surface( - std::size_t globalIndex, const config_t cfg, + unsigned int globalIndex, const finding_config& cfg, const propagate_to_next_surface_payload& payload); } // namespace traccc::device diff --git a/device/common/include/traccc/finding/device/prune_tracks.hpp b/device/common/include/traccc/finding/device/prune_tracks.hpp index bc3692565..ba258ff2c 100644 --- a/device/common/include/traccc/finding/device/prune_tracks.hpp +++ b/device/common/include/traccc/finding/device/prune_tracks.hpp @@ -35,7 +35,7 @@ struct prune_tracks_payload { /// /// @param[in] globalIndex The index of the current thread /// @param[inout] payload The function call payload -TRACCC_DEVICE inline void prune_tracks(std::size_t globalIndex, +TRACCC_DEVICE inline void prune_tracks(unsigned int globalIndex, const prune_tracks_payload& payload); } // namespace traccc::device diff --git a/device/cuda/include/traccc/cuda/utils/thread_id.hpp b/device/cuda/include/traccc/cuda/utils/thread_id.hpp index ede321665..dfd124dc7 100644 --- a/device/cuda/include/traccc/cuda/utils/thread_id.hpp +++ b/device/cuda/include/traccc/cuda/utils/thread_id.hpp @@ -8,34 +8,34 @@ #pragma once -#include - #include "traccc/definitions/qualifiers.hpp" namespace traccc::cuda { struct thread_id1 { TRACCC_DEVICE thread_id1() {} - std::size_t inline TRACCC_DEVICE getLocalThreadId() const { + unsigned int inline TRACCC_DEVICE getLocalThreadId() const { return threadIdx.x; } - std::size_t inline TRACCC_DEVICE getLocalThreadIdX() const { + unsigned int inline TRACCC_DEVICE getLocalThreadIdX() const { return threadIdx.x; } - std::size_t inline TRACCC_DEVICE getGlobalThreadId() const { + unsigned int inline TRACCC_DEVICE getGlobalThreadId() const { return threadIdx.x + blockIdx.x * blockDim.x; } - std::size_t inline TRACCC_DEVICE getGlobalThreadIdX() const { + unsigned int inline TRACCC_DEVICE getGlobalThreadIdX() const { return threadIdx.x + blockIdx.x * blockDim.x; } - std::size_t inline TRACCC_DEVICE getBlockIdX() const { return blockIdx.x; } + unsigned int inline TRACCC_DEVICE getBlockIdX() const { return blockIdx.x; } - std::size_t inline TRACCC_DEVICE getBlockDimX() const { return blockDim.x; } + unsigned int inline TRACCC_DEVICE getBlockDimX() const { + return blockDim.x; + } - std::size_t inline TRACCC_DEVICE getGridDimX() const { return gridDim.x; } + unsigned int inline TRACCC_DEVICE getGridDimX() const { return gridDim.x; } }; } // namespace traccc::cuda diff --git a/device/cuda/src/finding/finding_algorithm.cu b/device/cuda/src/finding/finding_algorithm.cu index ef6f82b7a..e38df16d9 100644 --- a/device/cuda/src/finding/finding_algorithm.cu +++ b/device/cuda/src/finding/finding_algorithm.cu @@ -188,8 +188,8 @@ finding_algorithm::operator()( kernels::apply_interaction> <<>>( - m_cfg, {det_view, static_cast(n_in_params), - in_params_buffer, param_liveness_buffer}); + m_cfg, {det_view, n_in_params, in_params_buffer, + param_liveness_buffer}); TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); } diff --git a/device/cuda/src/finding/kernels/build_tracks.cu b/device/cuda/src/finding/kernels/build_tracks.cu index 801bf118d..9b80fd8d6 100644 --- a/device/cuda/src/finding/kernels/build_tracks.cu +++ b/device/cuda/src/finding/kernels/build_tracks.cu @@ -18,7 +18,7 @@ namespace traccc::cuda::kernels { __global__ void build_tracks(const finding_config cfg, device::build_tracks_payload payload) { - int gid = threadIdx.x + blockIdx.x * blockDim.x; + const unsigned int gid = threadIdx.x + blockIdx.x * blockDim.x; device::build_tracks(gid, cfg, payload); } diff --git a/device/cuda/src/finding/kernels/prune_tracks.cu b/device/cuda/src/finding/kernels/prune_tracks.cu index f431676a4..525a3d6a9 100644 --- a/device/cuda/src/finding/kernels/prune_tracks.cu +++ b/device/cuda/src/finding/kernels/prune_tracks.cu @@ -12,7 +12,7 @@ namespace traccc::cuda::kernels { __global__ void prune_tracks(device::prune_tracks_payload payload) { - int gid = threadIdx.x + blockIdx.x * blockDim.x; + const unsigned int gid = threadIdx.x + blockIdx.x * blockDim.x; device::prune_tracks(gid, payload); } diff --git a/device/cuda/src/finding/kernels/specializations/apply_interaction_src.cuh b/device/cuda/src/finding/kernels/specializations/apply_interaction_src.cuh index 53b6b2e2a..da8f98395 100644 --- a/device/cuda/src/finding/kernels/specializations/apply_interaction_src.cuh +++ b/device/cuda/src/finding/kernels/specializations/apply_interaction_src.cuh @@ -18,7 +18,7 @@ __global__ void apply_interaction( const finding_config cfg, device::apply_interaction_payload payload) { - int gid = threadIdx.x + blockIdx.x * blockDim.x; + const unsigned int gid = threadIdx.x + blockIdx.x * blockDim.x; device::apply_interaction(gid, cfg, payload); } diff --git a/device/cuda/src/finding/kernels/specializations/find_tracks_src.cuh b/device/cuda/src/finding/kernels/specializations/find_tracks_src.cuh index 1e9fbafc7..ab25795bb 100644 --- a/device/cuda/src/finding/kernels/specializations/find_tracks_src.cuh +++ b/device/cuda/src/finding/kernels/specializations/find_tracks_src.cuh @@ -29,8 +29,7 @@ __global__ void find_tracks(const finding_config cfg, cuda::barrier barrier; cuda::thread_id1 thread_id; - device::find_tracks( + device::find_tracks( thread_id, barrier, cfg, payload, {shared_num_candidates, shared_candidates, shared_candidates_size}); } diff --git a/device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_src.cuh b/device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_src.cuh index d35724a75..20419c799 100644 --- a/device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_src.cuh +++ b/device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_src.cuh @@ -18,10 +18,10 @@ __global__ void propagate_to_next_surface( const finding_config cfg, device::propagate_to_next_surface_payload payload) { - int gid = threadIdx.x + blockIdx.x * blockDim.x; + const unsigned int gid = threadIdx.x + blockIdx.x * blockDim.x; - device::propagate_to_next_surface( - gid, cfg, payload); + device::propagate_to_next_surface(gid, cfg, + payload); } } // namespace traccc::cuda::kernels From 7965296a74a4db93e6d058c8cd5c965b5a012698 Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Fri, 15 Nov 2024 14:52:22 +0100 Subject: [PATCH 04/12] Add a first version of SYCL track finding. --- core/include/traccc/edm/measurement.hpp | 4 +- device/sycl/src/finding/find_tracks.hpp | 449 +++++++++++++++++++++++- device/sycl/src/utils/thread_id.hpp | 28 +- 3 files changed, 462 insertions(+), 19 deletions(-) diff --git a/core/include/traccc/edm/measurement.hpp b/core/include/traccc/edm/measurement.hpp index 015345df6..68b77fbe2 100644 --- a/core/include/traccc/edm/measurement.hpp +++ b/core/include/traccc/edm/measurement.hpp @@ -88,14 +88,14 @@ inline bool operator==(const measurement& lhs, const measurement& rhs) { /// Comparator based on detray barcode value struct measurement_sort_comp { TRACCC_HOST_DEVICE - bool operator()(const measurement& lhs, const measurement& rhs) { + bool operator()(const measurement& lhs, const measurement& rhs) const { return lhs.surface_link < rhs.surface_link; } }; struct measurement_equal_comp { TRACCC_HOST_DEVICE - bool operator()(const measurement& lhs, const measurement& rhs) { + bool operator()(const measurement& lhs, const measurement& rhs) const { return lhs.surface_link == rhs.surface_link; } }; diff --git a/device/sycl/src/finding/find_tracks.hpp b/device/sycl/src/finding/find_tracks.hpp index d970b93ca..c95111f44 100644 --- a/device/sycl/src/finding/find_tracks.hpp +++ b/device/sycl/src/finding/find_tracks.hpp @@ -7,11 +7,28 @@ #pragma once +// Local include(s). +#include "../utils/barrier.hpp" +#include "../utils/calculate1DimNdRange.hpp" +#include "../utils/thread_id.hpp" + // Project include(s). +#include "../sanity/contiguous_on.hpp" #include "traccc/edm/measurement.hpp" #include "traccc/edm/track_candidate.hpp" +#include "traccc/finding/actors/ckf_aborter.hpp" +#include "traccc/finding/actors/interaction_register.hpp" +#include "traccc/finding/candidate_link.hpp" +#include "traccc/finding/device/apply_interaction.hpp" +#include "traccc/finding/device/build_tracks.hpp" +#include "traccc/finding/device/fill_sort_keys.hpp" +#include "traccc/finding/device/find_tracks.hpp" +#include "traccc/finding/device/make_barcode_sequence.hpp" +#include "traccc/finding/device/propagate_to_next_surface.hpp" +#include "traccc/finding/device/prune_tracks.hpp" #include "traccc/finding/finding_config.hpp" #include "traccc/utils/memory_resource.hpp" +#include "traccc/utils/projections.hpp" // Detray include(s). #include "detray/propagator/actor_chain.hpp" @@ -23,6 +40,11 @@ // VecMem include(s). #include +#include + +// oneDPL include(s). +#include +#include // SYCL include(s). #include @@ -52,14 +74,425 @@ namespace traccc::sycl::details { /// template track_candidate_container_types::buffer find_tracks( - const typename navigator_t::detector_type::view_type& /*det*/, - const typename stepper_t::magnetic_field_type& /*field*/, - const measurement_collection_types::const_view& /*measurements_view*/, - const bound_track_parameters_collection_types::const_view& /*seeds_view*/, - const finding_config& /*config*/, const memory_resource& /*mr*/, - vecmem::copy& /*copy*/, ::sycl::queue& /*queue*/) { - - return {}; + const typename navigator_t::detector_type::view_type& det, + const typename stepper_t::magnetic_field_type& field, + const measurement_collection_types::const_view& measurements, + const bound_track_parameters_collection_types::const_view& seeds, + const finding_config& config, const memory_resource& mr, vecmem::copy& copy, + ::sycl::queue& queue) { + + assert(is_contiguous_on( + measurement_module_projection(), mr.main, copy, queue, measurements)); + + // oneDPL policy to use, forcing execution onto the same device that the + // hand-written kernels would run on. + auto policy = oneapi::dpl::execution::device_policy{queue}; + + /***************************************************************** + * Measurement Operations + *****************************************************************/ + + const measurement_collection_types::const_view::size_type n_measurements = + copy.get_size(measurements); + + // Get copy of barcode uniques + measurement_collection_types::buffer uniques_buffer{n_measurements, + mr.main}; + copy.setup(uniques_buffer)->wait(); + measurement_collection_types::device uniques(uniques_buffer); + + measurement_collection_types::device::iterator uniques_end = + oneapi::dpl::unique_copy(policy, measurements.ptr(), + measurements.ptr() + n_measurements, + uniques.begin(), measurement_equal_comp()); + const unsigned int n_modules = + static_cast(uniques_end - uniques.begin()); + + // Get upper bounds of unique elements + vecmem::data::vector_buffer upper_bounds_buffer{n_modules, + mr.main}; + copy.setup(upper_bounds_buffer)->wait(); + vecmem::device_vector upper_bounds(upper_bounds_buffer); + + oneapi::dpl::upper_bound(policy, measurements.ptr(), + measurements.ptr() + n_measurements, + uniques.begin(), uniques.begin() + n_modules, + upper_bounds.begin(), measurement_sort_comp()); + + /***************************************************************** + * Kernel1: Create barcode sequence + *****************************************************************/ + + vecmem::data::vector_buffer barcodes_buffer{ + n_modules, mr.main}; + copy.setup(barcodes_buffer)->wait(); + + queue + .submit([&](::sycl::handler& h) { + h.parallel_for( + calculate1DimNdRange(n_modules, 64), + [uniques_view = vecmem::get_data(uniques_buffer), + barcodes_view = vecmem::get_data(barcodes_buffer)]( + ::sycl::nd_item<1> item) { + device::make_barcode_sequence( + static_cast(item.get_global_id(0)), + {uniques_view, barcodes_view}); + }); + }) + .wait_and_throw(); + + const unsigned int n_seeds = copy.get_size(seeds); + + // Prepare input parameters with seeds + bound_track_parameters_collection_types::buffer in_params_buffer(n_seeds, + mr.main); + copy.setup(in_params_buffer)->wait(); + copy(seeds, in_params_buffer, vecmem::copy::type::device_to_device)->wait(); + vecmem::data::vector_buffer param_liveness_buffer(n_seeds, mr.main); + copy.setup(param_liveness_buffer)->wait(); + copy.memset(param_liveness_buffer, 1)->wait(); + + // Number of tracks per seed + vecmem::data::vector_buffer n_tracks_per_seed_buffer(n_seeds, + mr.main); + copy.setup(n_tracks_per_seed_buffer)->wait(); + + // Create a map for links + std::map> + link_map; + + // Create a buffer of tip links + vecmem::data::vector_buffer + tips_buffer{config.max_num_branches_per_seed * n_seeds, mr.main, + vecmem::data::buffer_type::resizable}; + copy.setup(tips_buffer)->wait(); + + // Link size + std::vector n_candidates_per_step; + n_candidates_per_step.reserve(config.max_track_candidates_per_track); + + unsigned int n_in_params = n_seeds; + for (unsigned int step = 0; + step < config.max_track_candidates_per_track && n_in_params > 0; + step++) { + + /***************************************************************** + * Kernel2: Apply material interaction + ****************************************************************/ + + queue + .submit([&](::sycl::handler& h) { + h.parallel_for( + calculate1DimNdRange(n_in_params, 64), + [config, det, n_in_params, + in_params = vecmem::get_data(in_params_buffer), + param_liveness = vecmem::get_data(param_liveness_buffer)]( + ::sycl::nd_item<1> item) { + device::apply_interaction< + typename navigator_t::detector_type>( + static_cast(item.get_global_id(0)), + config, + {det, n_in_params, in_params, param_liveness}); + }); + }) + .wait_and_throw(); + + /***************************************************************** + * Kernel3: Find valid tracks + *****************************************************************/ + + // Previous step + const unsigned int prev_step = (step == 0 ? 0 : step - 1); + + // Buffer for kalman-updated parameters spawned by the + // measurement candidates + const unsigned int n_max_candidates = + n_in_params * config.max_num_branches_per_surface; + + bound_track_parameters_collection_types::buffer updated_params_buffer( + n_in_params * config.max_num_branches_per_surface, mr.main); + copy.setup(updated_params_buffer)->wait(); + + vecmem::data::vector_buffer updated_liveness_buffer( + n_in_params * config.max_num_branches_per_surface, mr.main); + copy.setup(updated_liveness_buffer)->wait(); + copy.memset(updated_liveness_buffer, 0)->wait(); + + // Create the link map + link_map[step] = {n_in_params * config.max_num_branches_per_surface, + mr.main}; + copy.setup(link_map[step])->wait(); + + vecmem::unique_alloc_ptr n_candidates_device = + vecmem::make_unique_alloc(mr.main); + queue.memset(n_candidates_device.get(), 0, sizeof(unsigned int)) + .wait_and_throw(); + + // The number of threads to use per block in the track finding. + static const unsigned int nFindTracksThreads = 64; + + // Submit the kernel to the queue. + queue + .submit([&](::sycl::handler& h) { + // Allocate shared memory for the kernel. + vecmem::sycl::local_accessor + shared_num_candidates(nFindTracksThreads, h); + vecmem::sycl::local_accessor< + std::pair> + shared_candidates(nFindTracksThreads, h); + vecmem::sycl::local_accessor + shared_candidates_size(1, h); + + // Launch the kernel. + h.parallel_for( + calculate1DimNdRange(n_in_params, nFindTracksThreads), + [config, det, measurements, + in_params = vecmem::get_data(in_params_buffer), + param_liveness = vecmem::get_data(param_liveness_buffer), + n_in_params, barcodes = vecmem::get_data(barcodes_buffer), + upper_bounds = vecmem::get_data(upper_bounds_buffer), + previous_candidate_links = + vecmem::get_data(link_map.at(prev_step)), + step, n_max_candidates, + updated_params = vecmem::get_data(updated_params_buffer), + updated_liveness = + vecmem::get_data(updated_liveness_buffer), + current_candidate_links = + vecmem::get_data(link_map.at(step)), + n_candidates = n_candidates_device.get(), + shared_candidates_size, shared_num_candidates, + shared_candidates](::sycl::nd_item<1> item) { + // SYCL wrappers used in the algorithm. + const details::barrier barrier{item}; + const details::thread_id thread_id{item}; + + // Call the device function to find tracks. + device::find_tracks< + std::decay_t>( + thread_id, barrier, config, + {det, measurements, in_params, param_liveness, + n_in_params, barcodes, upper_bounds, + previous_candidate_links, step, n_max_candidates, + updated_params, updated_liveness, + current_candidate_links, n_candidates}, + {&(shared_num_candidates[0]), + &(shared_candidates[0]), + shared_candidates_size[0]}); + }); + }) + .wait_and_throw(); + + std::swap(in_params_buffer, updated_params_buffer); + std::swap(param_liveness_buffer, updated_liveness_buffer); + + // Get the number of candidates back to the host. + unsigned int n_candidates = 0; + queue + .memcpy(&n_candidates, n_candidates_device.get(), + sizeof(unsigned int)) + .wait_and_throw(); + + if (n_candidates > 0) { + /***************************************************************** + * Kernel4: Get key and value for parameter sorting + *****************************************************************/ + + vecmem::data::vector_buffer param_ids_buffer( + n_candidates, mr.main); + copy.setup(param_ids_buffer)->wait(); + + { + vecmem::data::vector_buffer keys_buffer( + n_candidates, mr.main); + copy.setup(keys_buffer)->wait(); + + queue + .submit([&](::sycl::handler& h) { + h.parallel_for( + calculate1DimNdRange(n_candidates, 256), + [in_params = vecmem::get_data(in_params_buffer), + keys = vecmem::get_data(keys_buffer), + param_ids = vecmem::get_data(param_ids_buffer)]( + ::sycl::nd_item<1> item) { + device::fill_sort_keys( + static_cast( + item.get_global_id(0)), + {in_params, keys, param_ids}); + }); + }) + .wait_and_throw(); + + // Sort the keys and values. + vecmem::device_vector keys_device( + keys_buffer); + vecmem::device_vector param_ids_device( + param_ids_buffer); + oneapi::dpl::sort_by_key(policy, keys_device.begin(), + keys_device.end(), + param_ids_device.begin()); + } + + /***************************************************************** + * Kernel5: Propagate to the next surface + *****************************************************************/ + + // Reset the number of tracks per seed + copy.memset(n_tracks_per_seed_buffer, 0)->wait(); + + /// Actor types + using algebra_type = + typename navigator_t::detector_type::algebra_type; + using interactor_type = + detray::pointwise_material_interactor; + using actor_type = + detray::actor_chain, + interaction_register, + interactor_type, ckf_aborter>; + using propagator_type = + detray::propagator; + + // Launch the kernel to propagate all active tracks to the next + // surface. + queue + .submit([&](::sycl::handler& h) { + h.parallel_for( + calculate1DimNdRange(n_candidates, 64), + [config, det, field, + in_params = vecmem::get_data(in_params_buffer), + param_liveness = + vecmem::get_data(param_liveness_buffer), + param_ids = vecmem::get_data(param_ids_buffer), + current_candidate_links = + vecmem::get_data(link_map.at(step)), + step, n_candidates, + tips = vecmem::get_data(tips_buffer), + n_tracks_per_seed = + vecmem::get_data(n_tracks_per_seed_buffer)]( + ::sycl::nd_item<1> item) { + device::propagate_to_next_surface< + propagator_type, + typename stepper_t::magnetic_field_type>( + static_cast( + item.get_global_id(0)), + config, + {det, field, in_params, param_liveness, + param_ids, current_candidate_links, step, + n_candidates, tips, n_tracks_per_seed}); + }); + }) + .wait_and_throw(); + } + + // Fill the candidate size vector + n_candidates_per_step.push_back(n_candidates); + + n_in_params = n_candidates; + } + + // Create link buffer + vecmem::data::jagged_vector_buffer links_buffer( + n_candidates_per_step, mr.main, mr.host); + copy.setup(links_buffer)->wait(); + + // Copy link map to link buffer + for (unsigned int it = 0; + it < static_cast(n_candidates_per_step.size()); it++) { + + vecmem::device_vector in(link_map.at(it)); + vecmem::device_vector out( + *(links_buffer.host_ptr() + it)); + + oneapi::dpl::copy(policy, in.begin(), + in.begin() + n_candidates_per_step[it], out.begin()); + } + + /***************************************************************** + * Kernel6: Build tracks + *****************************************************************/ + + // Get the number of tips + auto n_tips_total = copy.get_size(tips_buffer); + + // Create track candidate buffer + track_candidate_container_types::buffer track_candidates_buffer{ + {n_tips_total, mr.main}, + {std::vector(n_tips_total, + config.max_track_candidates_per_track), + mr.main, mr.host, vecmem::data::buffer_type::resizable}}; + copy.setup(track_candidates_buffer.headers)->wait(); + copy.setup(track_candidates_buffer.items)->wait(); + track_candidate_container_types::view track_candidates = + track_candidates_buffer; + + // Create buffer for valid indices + vecmem::data::vector_buffer valid_indices_buffer(n_tips_total, + mr.main); + copy.setup(valid_indices_buffer)->wait(); + + unsigned int n_valid_tracks = 0u; + + if (n_tips_total > 0) { + vecmem::unique_alloc_ptr n_valid_tracks_device = + vecmem::make_unique_alloc(mr.main); + queue.memset(n_valid_tracks_device.get(), 0, sizeof(unsigned int)) + .wait_and_throw(); + + queue + .submit([&](::sycl::handler& h) { + h.parallel_for( + calculate1DimNdRange(n_tips_total, 64), + [config, measurements, seeds, + links = vecmem::get_data(links_buffer), + tips = vecmem::get_data(tips_buffer), track_candidates, + valid_indices = vecmem::get_data(valid_indices_buffer), + n_valid_tracks = + n_valid_tracks_device.get()](::sycl::nd_item<1> item) { + device::build_tracks( + static_cast(item.get_global_id(0)), + config, + {measurements, seeds, links, tips, track_candidates, + valid_indices, n_valid_tracks}); + }); + }) + .wait_and_throw(); + + queue + .memcpy(&n_valid_tracks, n_valid_tracks_device.get(), + sizeof(unsigned int)) + .wait_and_throw(); + } + + // Create pruned candidate buffer + track_candidate_container_types::buffer prune_candidates_buffer{ + {n_valid_tracks, mr.main}, + {std::vector(n_valid_tracks, + config.max_track_candidates_per_track), + mr.main, mr.host, vecmem::data::buffer_type::resizable}}; + copy.setup(prune_candidates_buffer.headers)->wait(); + copy.setup(prune_candidates_buffer.items)->wait(); + track_candidate_container_types::view prune_candidates = + prune_candidates_buffer; + + if (n_valid_tracks > 0) { + + queue + .submit([&](::sycl::handler& h) { + h.parallel_for( + calculate1DimNdRange(n_valid_tracks, 64), + [track_candidates, + valid_indices = vecmem::get_data(valid_indices_buffer), + prune_candidates](::sycl::nd_item<1> item) { + device::prune_tracks( + static_cast(item.get_global_id(0)), + {track_candidates, valid_indices, + prune_candidates}); + }); + }) + .wait_and_throw(); + } + + return prune_candidates_buffer; } } // namespace traccc::sycl::details diff --git a/device/sycl/src/utils/thread_id.hpp b/device/sycl/src/utils/thread_id.hpp index 8e62ae135..637249006 100644 --- a/device/sycl/src/utils/thread_id.hpp +++ b/device/sycl/src/utils/thread_id.hpp @@ -33,23 +33,33 @@ requires(DIMENSIONS >= 1 && DIMENSIONS <= 3) struct thread_id { /// @name Function(s) implementing @c traccc::device::concepts::thread_id1 /// @{ - inline auto getLocalThreadId() const { - return m_item.get_local_linear_id(); + inline unsigned int getLocalThreadId() const { + return static_cast(m_item.get_local_linear_id()); } - inline auto getLocalThreadIdX() const { return m_item.get_local_id(0); } + inline unsigned int getLocalThreadIdX() const { + return static_cast(m_item.get_local_id(0)); + } - inline auto getGlobalThreadId() const { - return m_item.get_global_linear_id(); + inline unsigned int getGlobalThreadId() const { + return static_cast(m_item.get_global_linear_id()); } - inline auto getGlobalThreadIdX() const { return m_item.get_global_id(0); } + inline unsigned int getGlobalThreadIdX() const { + return static_cast(m_item.get_global_id(0)); + } - inline auto getBlockIdX() const { return m_item.get_group(0); } + inline unsigned int getBlockIdX() const { + return static_cast(m_item.get_group(0)); + } - inline auto getBlockDimX() const { return m_item.get_local_range(0); } + inline unsigned int getBlockDimX() const { + return static_cast(m_item.get_local_range(0)); + } - inline auto getGridDimX() const { return m_item.get_global_range(0); } + inline unsigned int getGridDimX() const { + return static_cast(m_item.get_global_range(0)); + } /// @} From a6a23fe7f26006255e15676dde3b82a616bc7ef0 Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Sat, 16 Nov 2024 16:00:33 +0100 Subject: [PATCH 05/12] Added a not yet functional test for the SYCL CKF. --- .../traccc/simulation/smearing_writer.hpp | 29 +- tests/sycl/CMakeLists.txt | 3 +- tests/sycl/test_ckf_toy_detector.cpp | 253 ++++++++++++++++++ 3 files changed, 275 insertions(+), 10 deletions(-) create mode 100644 tests/sycl/test_ckf_toy_detector.cpp diff --git a/simulation/include/traccc/simulation/smearing_writer.hpp b/simulation/include/traccc/simulation/smearing_writer.hpp index 523bcef2e..b12f3975c 100644 --- a/simulation/include/traccc/simulation/smearing_writer.hpp +++ b/simulation/include/traccc/simulation/smearing_writer.hpp @@ -26,6 +26,10 @@ #include #include +// System include(s). +#include +#include + namespace traccc { template @@ -47,16 +51,23 @@ struct smearing_writer : detray::actor { struct state { state(std::size_t event_id, config&& writer_cfg, const std::string directory) - : m_particle_writer(directory + - traccc::io::get_event_filename( - event_id, "-particles_initial.csv")), - m_hit_writer(directory + traccc::io::get_event_filename( - event_id, "-hits.csv")), - m_meas_writer(directory + traccc::io::get_event_filename( - event_id, "-measurements.csv")), + : m_particle_writer((std::filesystem::path{directory} / + traccc::io::get_event_filename( + event_id, "-particles_initial.csv")) + .native()), + m_hit_writer( + (std::filesystem::path{directory} / + traccc::io::get_event_filename(event_id, "-hits.csv")) + .native()), + m_meas_writer((std::filesystem::path{directory} / + traccc::io::get_event_filename( + event_id, "-measurements.csv")) + .native()), m_measurement_hit_id_writer( - directory + traccc::io::get_event_filename( - event_id, "-measurement-simhit-map.csv")), + (std::filesystem::path{directory} / + traccc::io::get_event_filename( + event_id, "-measurement-simhit-map.csv")) + .native()), m_meas_smearer(writer_cfg.smearer) {} uint64_t particle_id = 0u; diff --git a/tests/sycl/CMakeLists.txt b/tests/sycl/CMakeLists.txt index 1b01c21c5..8c62f0c92 100644 --- a/tests/sycl/CMakeLists.txt +++ b/tests/sycl/CMakeLists.txt @@ -1,6 +1,6 @@ # TRACCC library, part of the ACTS project (R&D line) # -# (c) 2022 CERN for the benefit of the ACTS project +# (c) 2022-2024 CERN for the benefit of the ACTS project # # Mozilla Public License Version 2.0 @@ -14,6 +14,7 @@ traccc_add_test( # Define the sources for the test. # TODO: Reactivate this once #655 is fixed. # test_kalman_fitter_telescope.sycl + test_ckf_toy_detector.cpp test_clusterization.sycl test_dpl.sycl test_spacepoint_formation.sycl diff --git a/tests/sycl/test_ckf_toy_detector.cpp b/tests/sycl/test_ckf_toy_detector.cpp new file mode 100644 index 000000000..3660a53fa --- /dev/null +++ b/tests/sycl/test_ckf_toy_detector.cpp @@ -0,0 +1,253 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2023-2024 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Project include(s). +#include "traccc/device/container_d2h_copy_alg.hpp" +#include "traccc/device/container_h2d_copy_alg.hpp" +#include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp" +#include "traccc/io/read_detector.hpp" +#include "traccc/io/read_measurements.hpp" +#include "traccc/io/utils.hpp" +#include "traccc/performance/container_comparator.hpp" +#include "traccc/simulation/simulator.hpp" +#include "traccc/sycl/finding/combinatorial_kalman_filter_algorithm.hpp" +#include "traccc/utils/event_data.hpp" +#include "traccc/utils/ranges.hpp" + +// Test include(s). +#include "tests/ckf_toy_detector_test.hpp" +#include "traccc/utils/seed_generator.hpp" + +// detray include(s). +#include "detray/propagator/propagator.hpp" +#include "detray/test/utils/simulation/event_generator/track_generators.hpp" + +// VecMem include(s). +#include +#include +#include +#include + +// GTest include(s). +#include + +// System include(s). +#include +#include + +namespace traccc { + +TEST_P(CkfToyDetectorTests, Run) { + + // Get the parameters + const std::string name = std::get<0>(GetParam()); + const detray::pdg_particle ptc = std::get<6>(GetParam()); + const unsigned int n_truth_tracks = std::get<7>(GetParam()); + const unsigned int n_events = std::get<8>(GetParam()); + const bool random_charge = std::get<9>(GetParam()); + + /***************************** + * Build a toy detector + *****************************/ + + // Memory resources used by the application. + vecmem::host_memory_resource host_mr; + vecmem::sycl::queue_wrapper queue; + vecmem::sycl::device_memory_resource device_mr{queue}; + traccc::memory_resource mr{device_mr, &host_mr}; + vecmem::sycl::shared_memory_resource shared_mr{queue}; + + // Path to the working directory. + const std::filesystem::path path = std::filesystem::current_path() / name; + + // Read in the detector geometry that was generated by the test fixture. + host_detector_type host_det{shared_mr}; + traccc::io::read_detector( + host_det, shared_mr, (path / "toy_detector_geometry.json").native(), + (path / "toy_detector_homogeneous_material.json").native(), + (path / "toy_detector_surface_grids.json").native()); + + auto field = detray::bfield::create_const_field(B); + + // Detector view object + auto det_view = detray::get_data(host_det); + + /*************************** + * Generate simulation data + ***************************/ + + // Track generator + using generator_type = + detray::random_track_generator; + generator_type::configuration gen_cfg{}; + gen_cfg.n_tracks(n_truth_tracks); + gen_cfg.origin(std::get<1>(GetParam())); + gen_cfg.origin_stddev(std::get<2>(GetParam())); + gen_cfg.phi_range(std::get<5>(GetParam())); + gen_cfg.eta_range(std::get<4>(GetParam())); + gen_cfg.mom_range(std::get<3>(GetParam())); + gen_cfg.randomize_charge(random_charge); + gen_cfg.seed(42); + generator_type generator(gen_cfg); + + // Smearing value for measurements + traccc::measurement_smearer meas_smearer( + smearing[0], smearing[1]); + + using writer_type = traccc::smearing_writer< + traccc::measurement_smearer>; + + typename writer_type::config smearer_writer_cfg{meas_smearer}; + + // Run simulator + auto sim = traccc::simulator( + ptc, n_events, host_det, field, std::move(generator), + std::move(smearer_writer_cfg), path.native()); + sim.get_config().propagation.stepping.step_constraint = step_constraint; + sim.get_config().propagation.navigation.search_window = search_window; + sim.run(); + + /***************************** + * Do the reconstruction + *****************************/ + + // Copy objects + vecmem::sycl::async_copy copy{queue}; + + traccc::device::container_d2h_copy_alg< + traccc::track_candidate_container_types> + track_candidate_d2h{mr, copy}; + + traccc::device::container_d2h_copy_alg + track_state_d2h{mr, copy}; + + // Seed generator + seed_generator sg(host_det, stddevs); + + // Finding algorithm configuration + traccc::sycl::combinatorial_kalman_filter_algorithm::config_type cfg; + cfg.ptc_hypothesis = ptc; + cfg.max_num_branches_per_seed = 500; + cfg.chi2_max = 30.f; + cfg.propagation.navigation.search_window = search_window; + + // Finding algorithm object + traccc::host::combinatorial_kalman_filter_algorithm host_finding(cfg); + + // Finding algorithm object + traccc::sycl::combinatorial_kalman_filter_algorithm device_finding{ + cfg, mr, copy, queue.queue()}; + + // Iterate over events + for (std::size_t i_evt = 0; i_evt < n_events; i_evt++) { + + // Truth Track Candidates + traccc::event_data evt_data(path.native(), i_evt, host_mr); + + traccc::track_candidate_container_types::host truth_track_candidates = + evt_data.generate_truth_candidates(sg, host_mr); + + ASSERT_EQ(truth_track_candidates.size(), n_truth_tracks); + + // Prepare truth seeds + traccc::bound_track_parameters_collection_types::host seeds(&host_mr); + for (unsigned int i_trk = 0; i_trk < n_truth_tracks; i_trk++) { + seeds.push_back(truth_track_candidates.at(i_trk).header); + } + ASSERT_EQ(seeds.size(), n_truth_tracks); + + traccc::bound_track_parameters_collection_types::buffer seeds_buffer{ + static_cast(seeds.size()), mr.main}; + copy.setup(seeds_buffer)->wait(); + copy(vecmem::get_data(seeds), seeds_buffer, + vecmem::copy::type::host_to_device) + ->wait(); + + // Read measurements + traccc::measurement_collection_types::host measurements_per_event{ + &host_mr}; + traccc::io::read_measurements(measurements_per_event, i_evt, + path.native()); + + traccc::measurement_collection_types::buffer measurements_buffer( + static_cast(measurements_per_event.size()), mr.main); + copy.setup(measurements_buffer)->wait(); + copy(vecmem::get_data(measurements_per_event), measurements_buffer) + ->wait(); + + // Run host finding + auto track_candidates = host_finding( + host_det, field, vecmem::get_data(measurements_per_event), + vecmem::get_data(seeds)); + + // Run device finding + traccc::track_candidate_container_types::buffer + track_candidates_sycl_buffer = device_finding( + det_view, field, measurements_buffer, seeds_buffer); + + traccc::track_candidate_container_types::host track_candidates_sycl = + track_candidate_d2h(track_candidates_sycl_buffer); + + // Simple check + ASSERT_TRUE( + std::llabs(static_cast(track_candidates.size()) - + static_cast(track_candidates_sycl.size())) <= 1u); + ASSERT_GE(track_candidates.size(), n_truth_tracks); + + // Make sure that the outputs from cpu and cuda CKF are equivalent + unsigned int n_matches = 0u; + for (unsigned int i = 0u; i < track_candidates.size(); i++) { + auto iso = + traccc::details::is_same_object(track_candidates.at(i).items); + + for (unsigned int j = 0u; j < track_candidates_sycl.size(); j++) { + if (iso(track_candidates_sycl.at(j).items)) { + n_matches++; + break; + } + } + } + + float matching_rate = + float(n_matches) / + static_cast(std::max(track_candidates.size(), + track_candidates_sycl.size())); + EXPECT_GE(matching_rate, 0.999f); + } +} + +INSTANTIATE_TEST_SUITE_P( + SYCLCkfToyDetectorValidation, CkfToyDetectorTests, + ::testing::Values( + std::make_tuple("toy_n_particles_1", + std::array{0.f, 0.f, 0.f}, + std::array{0.f, 0.f, 0.f}, + std::array{1.f, 100.f}, + std::array{-4.f, 4.f}, + std::array{-detray::constant::pi, + detray::constant::pi}, + detray::muon(), 1, 1, false), + std::make_tuple("toy_n_particles_10000", + std::array{0.f, 0.f, 0.f}, + std::array{0.f, 0.f, 0.f}, + std::array{1.f, 100.f}, + std::array{-4.f, 4.f}, + std::array{-detray::constant::pi, + detray::constant::pi}, + detray::muon(), 10000, 1, false), + std::make_tuple("toy_n_particles_10000_random_charge", + std::array{0.f, 0.f, 0.f}, + std::array{0.f, 0.f, 0.f}, + std::array{1.f, 100.f}, + std::array{-4.f, 4.f}, + std::array{-detray::constant::pi, + detray::constant::pi}, + detray::muon(), 10000, 1, true))); + +} // namespace traccc From d277b9926be70c8d71602d526a9351d08adadfcc Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Mon, 18 Nov 2024 16:07:41 +0100 Subject: [PATCH 06/12] Make sure that all SYCL kernels would have a unique kernel class. Trying to avoid confusion at runtime about which kernel is which. --- ...rithm_constant_field_default_detector.sycl | 29 ++++++++++++++++-- ...thm_constant_field_telescope_detector.sycl | 30 +++++++++++++++++-- device/sycl/src/finding/find_tracks.hpp | 21 ++++++++----- .../silicon_pixel_spacepoint_formation.hpp | 5 ++-- ..._formation_algorithm_default_detector.sycl | 9 ++++-- ...ormation_algorithm_telescope_detector.sycl | 10 +++++-- 6 files changed, 84 insertions(+), 20 deletions(-) diff --git a/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_default_detector.sycl b/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_default_detector.sycl index fe2cb604e..3ce363f4e 100644 --- a/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_default_detector.sycl +++ b/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_default_detector.sycl @@ -17,6 +17,28 @@ #include namespace traccc::sycl { +namespace kernels::combinatorial_kalman_filter_constant_field_default_detector { + +struct make_barcode_sequence; +struct apply_interaction; +struct find_tracks; +struct fill_sort_keys; +struct propagate_to_next_surface; +struct build_tracks; +struct prune_tracks; + +struct kernels { + using make_barcode_sequence_kernel_type = make_barcode_sequence; + using apply_interaction_kernel_type = apply_interaction; + using find_tracks_kernel_type = find_tracks; + using fill_sort_keys_kernel_type = fill_sort_keys; + using propagate_to_next_surface_kernel_type = propagate_to_next_surface; + using build_tracks_kernel_type = build_tracks; + using prune_tracks_kernel_type = prune_tracks; +}; // namespace kernels + +} // namespace + // kernels::combinatorial_kalman_filter_constant_field_default_detector combinatorial_kalman_filter_algorithm::output_type combinatorial_kalman_filter_algorithm::operator()( @@ -30,9 +52,10 @@ combinatorial_kalman_filter_algorithm::operator()( detray::rk_stepper>, - detray::navigator>( - det, field, measurements, seeds, m_config, m_mr, m_copy, - details::get_queue(m_queue)); + detray::navigator, + kernels::combinatorial_kalman_filter_constant_field_default_detector:: + kernels>(det, field, measurements, seeds, m_config, m_mr, m_copy, + details::get_queue(m_queue)); } } // namespace traccc::sycl diff --git a/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_telescope_detector.sycl b/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_telescope_detector.sycl index 07d150861..424583910 100644 --- a/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_telescope_detector.sycl +++ b/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_telescope_detector.sycl @@ -17,6 +17,29 @@ #include namespace traccc::sycl { +namespace kernels:: + combinatorial_kalman_filter_constant_field_telescope_detector { + +struct make_barcode_sequence; +struct apply_interaction; +struct find_tracks; +struct fill_sort_keys; +struct propagate_to_next_surface; +struct build_tracks; +struct prune_tracks; + +struct kernels { + using make_barcode_sequence_kernel_type = make_barcode_sequence; + using apply_interaction_kernel_type = apply_interaction; + using find_tracks_kernel_type = find_tracks; + using fill_sort_keys_kernel_type = fill_sort_keys; + using propagate_to_next_surface_kernel_type = propagate_to_next_surface; + using build_tracks_kernel_type = build_tracks; + using prune_tracks_kernel_type = prune_tracks; +}; // namespace kernels + +} // namespace + // kernels::combinatorial_kalman_filter_constant_field_telescope_detector combinatorial_kalman_filter_algorithm::output_type combinatorial_kalman_filter_algorithm::operator()( @@ -30,9 +53,10 @@ combinatorial_kalman_filter_algorithm::operator()( detray::rk_stepper>, - detray::navigator>( - det, field, measurements, seeds, m_config, m_mr, m_copy, - details::get_queue(m_queue)); + detray::navigator, + kernels::combinatorial_kalman_filter_constant_field_telescope_detector:: + kernels>(det, field, measurements, seeds, m_config, m_mr, m_copy, + details::get_queue(m_queue)); } } // namespace traccc::sycl diff --git a/device/sycl/src/finding/find_tracks.hpp b/device/sycl/src/finding/find_tracks.hpp index c95111f44..53ee6d48e 100644 --- a/device/sycl/src/finding/find_tracks.hpp +++ b/device/sycl/src/finding/find_tracks.hpp @@ -59,6 +59,7 @@ namespace traccc::sycl::details { /// /// @tparam stepper_t The stepper type used for the track propagation /// @tparam navigator_t The navigator type used for the track navigation +/// @tparam kernels_t Structure with unique "kernel structures" /// /// @param det A view of the detector object /// @param field The magnetic field object @@ -72,7 +73,7 @@ namespace traccc::sycl::details { /// /// @return A buffer of the found track candidates /// -template +template track_candidate_container_types::buffer find_tracks( const typename navigator_t::detector_type::view_type& det, const typename stepper_t::magnetic_field_type& field, @@ -129,7 +130,8 @@ track_candidate_container_types::buffer find_tracks( queue .submit([&](::sycl::handler& h) { - h.parallel_for( + h.parallel_for< + typename kernels_t::make_barcode_sequence_kernel_type>( calculate1DimNdRange(n_modules, 64), [uniques_view = vecmem::get_data(uniques_buffer), barcodes_view = vecmem::get_data(barcodes_buffer)]( @@ -182,7 +184,8 @@ track_candidate_container_types::buffer find_tracks( queue .submit([&](::sycl::handler& h) { - h.parallel_for( + h.parallel_for< + typename kernels_t::apply_interaction_kernel_type>( calculate1DimNdRange(n_in_params, 64), [config, det, n_in_params, in_params = vecmem::get_data(in_params_buffer), @@ -244,7 +247,7 @@ track_candidate_container_types::buffer find_tracks( shared_candidates_size(1, h); // Launch the kernel. - h.parallel_for( + h.parallel_for( calculate1DimNdRange(n_in_params, nFindTracksThreads), [config, det, measurements, in_params = vecmem::get_data(in_params_buffer), @@ -308,7 +311,8 @@ track_candidate_container_types::buffer find_tracks( queue .submit([&](::sycl::handler& h) { - h.parallel_for( + h.parallel_for< + typename kernels_t::fill_sort_keys_kernel_type>( calculate1DimNdRange(n_candidates, 256), [in_params = vecmem::get_data(in_params_buffer), keys = vecmem::get_data(keys_buffer), @@ -356,7 +360,8 @@ track_candidate_container_types::buffer find_tracks( // surface. queue .submit([&](::sycl::handler& h) { - h.parallel_for( + h.parallel_for( calculate1DimNdRange(n_candidates, 64), [config, det, field, in_params = vecmem::get_data(in_params_buffer), @@ -440,7 +445,7 @@ track_candidate_container_types::buffer find_tracks( queue .submit([&](::sycl::handler& h) { - h.parallel_for( + h.parallel_for( calculate1DimNdRange(n_tips_total, 64), [config, measurements, seeds, links = vecmem::get_data(links_buffer), @@ -478,7 +483,7 @@ track_candidate_container_types::buffer find_tracks( queue .submit([&](::sycl::handler& h) { - h.parallel_for( + h.parallel_for( calculate1DimNdRange(n_valid_tracks, 64), [track_candidates, valid_indices = vecmem::get_data(valid_indices_buffer), diff --git a/device/sycl/src/seeding/silicon_pixel_spacepoint_formation.hpp b/device/sycl/src/seeding/silicon_pixel_spacepoint_formation.hpp index ecd76b04a..1b1d52fce 100644 --- a/device/sycl/src/seeding/silicon_pixel_spacepoint_formation.hpp +++ b/device/sycl/src/seeding/silicon_pixel_spacepoint_formation.hpp @@ -28,6 +28,7 @@ namespace traccc::sycl::details { /// functions /// /// @tparam detector_t The detector type to use +/// @tparam kernel_t The kernel type to use /// /// @param det_view The view of the detector to use /// @param measurements_view The view of the measurements to process @@ -36,7 +37,7 @@ namespace traccc::sycl::details { /// @param queue The queue to use for the computation /// @return A buffer of the created spacepoints /// -template +template spacepoint_collection_types::buffer silicon_pixel_spacepoint_formation( const typename detector_t::view_type& det_view, const measurement_collection_types::const_view& measurements_view, @@ -64,7 +65,7 @@ spacepoint_collection_types::buffer silicon_pixel_spacepoint_formation( // Run the spacepoint formation on the device. queue .submit([&](::sycl::handler& h) { - h.parallel_for( + h.parallel_for( countRange, [det_view, measurements_view, n_measurements, spacepoints_view = vecmem::get_data(result)]( ::sycl::nd_item<1> item) { diff --git a/device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm_default_detector.sycl b/device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm_default_detector.sycl index c0ecd08c2..7dde18502 100644 --- a/device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm_default_detector.sycl +++ b/device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm_default_detector.sycl @@ -11,6 +11,11 @@ #include "traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" namespace traccc::sycl { +namespace kernels { + +struct form_spacepoints_default_detector; + +} // namespace kernels silicon_pixel_spacepoint_formation_algorithm::output_type silicon_pixel_spacepoint_formation_algorithm::operator()( @@ -18,8 +23,8 @@ silicon_pixel_spacepoint_formation_algorithm::operator()( const measurement_collection_types::const_view& meas) const { return details::silicon_pixel_spacepoint_formation< - default_detector::device>(det, meas, m_mr.main, m_copy, - details::get_queue(m_queue)); + default_detector::device, kernels::form_spacepoints_default_detector>( + det, meas, m_mr.main, m_copy, details::get_queue(m_queue)); } } // namespace traccc::sycl diff --git a/device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm_telescope_detector.sycl b/device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm_telescope_detector.sycl index 746d24cad..80b0478f7 100644 --- a/device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm_telescope_detector.sycl +++ b/device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm_telescope_detector.sycl @@ -11,6 +11,11 @@ #include "traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" namespace traccc::sycl { +namespace kernels { + +struct form_spacepoints_telescope_detector; + +} // namespace kernels silicon_pixel_spacepoint_formation_algorithm::output_type silicon_pixel_spacepoint_formation_algorithm::operator()( @@ -18,8 +23,9 @@ silicon_pixel_spacepoint_formation_algorithm::operator()( const measurement_collection_types::const_view& meas) const { return details::silicon_pixel_spacepoint_formation< - telescope_detector::device>(det, meas, m_mr.main, m_copy, - details::get_queue(m_queue)); + telescope_detector::device, + kernels::form_spacepoints_telescope_detector>( + det, meas, m_mr.main, m_copy, details::get_queue(m_queue)); } } // namespace traccc::sycl From 27643443dd0e1126a0b2ae3a101254d35888bf78 Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Wed, 20 Nov 2024 21:32:45 +0100 Subject: [PATCH 07/12] Make sorting work with oneDPL-2022.6.0. --- device/sycl/src/finding/find_tracks.hpp | 10 +++++++--- 1 file changed, 7 insertions(+), 3 deletions(-) diff --git a/device/sycl/src/finding/find_tracks.hpp b/device/sycl/src/finding/find_tracks.hpp index 53ee6d48e..53b429fff 100644 --- a/device/sycl/src/finding/find_tracks.hpp +++ b/device/sycl/src/finding/find_tracks.hpp @@ -331,9 +331,13 @@ track_candidate_container_types::buffer find_tracks( keys_buffer); vecmem::device_vector param_ids_device( param_ids_buffer); - oneapi::dpl::sort_by_key(policy, keys_device.begin(), - keys_device.end(), - param_ids_device.begin()); + auto zipped_first = oneapi::dpl::make_zip_iterator( + keys_device.begin(), param_ids_device.begin()); + oneapi::dpl::sort( + policy, zipped_first, zipped_first + keys_device.size(), + [](auto lhs, auto rhs) { + return std::get<0>(lhs) < std::get<0>(rhs); + }); } /***************************************************************** From 948f49a95546a3be095defe54e24614ff3ff2bc8 Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Wed, 20 Nov 2024 22:00:19 +0100 Subject: [PATCH 08/12] Shorten the namespace name, to stop freaking out clang-format. --- ...er_algorithm_constant_field_default_detector.sycl | 11 +++++------ ..._algorithm_constant_field_telescope_detector.sycl | 12 +++++------- 2 files changed, 10 insertions(+), 13 deletions(-) diff --git a/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_default_detector.sycl b/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_default_detector.sycl index 3ce363f4e..467ca244a 100644 --- a/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_default_detector.sycl +++ b/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_default_detector.sycl @@ -17,7 +17,7 @@ #include namespace traccc::sycl { -namespace kernels::combinatorial_kalman_filter_constant_field_default_detector { +namespace kernels::ckf_cfield_defdet { struct make_barcode_sequence; struct apply_interaction; @@ -37,8 +37,7 @@ struct kernels { using prune_tracks_kernel_type = prune_tracks; }; // namespace kernels -} // namespace - // kernels::combinatorial_kalman_filter_constant_field_default_detector +} // namespace kernels::ckf_cfield_defdet combinatorial_kalman_filter_algorithm::output_type combinatorial_kalman_filter_algorithm::operator()( @@ -53,9 +52,9 @@ combinatorial_kalman_filter_algorithm::operator()( default_detector::device::algebra_type, detray::constrained_step<>>, detray::navigator, - kernels::combinatorial_kalman_filter_constant_field_default_detector:: - kernels>(det, field, measurements, seeds, m_config, m_mr, m_copy, - details::get_queue(m_queue)); + kernels::ckf_cfield_defdet::kernels>(det, field, measurements, seeds, + m_config, m_mr, m_copy, + details::get_queue(m_queue)); } } // namespace traccc::sycl diff --git a/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_telescope_detector.sycl b/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_telescope_detector.sycl index 424583910..2591e9df0 100644 --- a/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_telescope_detector.sycl +++ b/device/sycl/src/finding/combinatorial_kalman_filter_algorithm_constant_field_telescope_detector.sycl @@ -17,8 +17,7 @@ #include namespace traccc::sycl { -namespace kernels:: - combinatorial_kalman_filter_constant_field_telescope_detector { +namespace kernels::ckf_cfield_teldet { struct make_barcode_sequence; struct apply_interaction; @@ -38,8 +37,7 @@ struct kernels { using prune_tracks_kernel_type = prune_tracks; }; // namespace kernels -} // namespace - // kernels::combinatorial_kalman_filter_constant_field_telescope_detector +} // namespace kernels::ckf_cfield_teldet combinatorial_kalman_filter_algorithm::output_type combinatorial_kalman_filter_algorithm::operator()( @@ -54,9 +52,9 @@ combinatorial_kalman_filter_algorithm::operator()( telescope_detector::device::algebra_type, detray::constrained_step<>>, detray::navigator, - kernels::combinatorial_kalman_filter_constant_field_telescope_detector:: - kernels>(det, field, measurements, seeds, m_config, m_mr, m_copy, - details::get_queue(m_queue)); + kernels::ckf_cfield_teldet::kernels>(det, field, measurements, seeds, + m_config, m_mr, m_copy, + details::get_queue(m_queue)); } } // namespace traccc::sycl From bd1ff94d1b242690efc048ad6622c98b10a97eea Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Wed, 11 Dec 2024 11:13:23 +0100 Subject: [PATCH 09/12] Update to oneDPL 2022.7.1. --- extern/dpl/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/extern/dpl/CMakeLists.txt b/extern/dpl/CMakeLists.txt index dc88cf5ae..e42245b5e 100644 --- a/extern/dpl/CMakeLists.txt +++ b/extern/dpl/CMakeLists.txt @@ -13,7 +13,7 @@ message( STATUS "Building oneDPL as part of the TRACCC project" ) # Declare where to get DPL from. set( TRACCC_DPL_SOURCE - "URL;https://github.com/oneapi-src/oneDPL/archive/refs/tags/oneDPL-2022.6.0-rc1.tar.gz;URL_MD5;f52a2ed5c9e4cdb3c65c2465b50abecf" + "URL;https://github.com/oneapi-src/oneDPL/archive/refs/tags/oneDPL-2022.7.1-release.tar.gz;URL_MD5;21d45dc27ba3133bfb282ec7383177f4" CACHE STRING "Source for DPL, when built as part of this project" ) mark_as_advanced( TRACCC_DPL_SOURCE ) FetchContent_Declare( DPL SYSTEM ${TRACCC_DPL_SOURCE} ) From 0fe1acb072cfdfdcef08e3f56d414322b35d5209 Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Wed, 11 Dec 2024 11:14:06 +0100 Subject: [PATCH 10/12] Made track finding work with the SYCL OpenCL backend! --- device/sycl/src/finding/find_tracks.hpp | 15 +++++++-------- 1 file changed, 7 insertions(+), 8 deletions(-) diff --git a/device/sycl/src/finding/find_tracks.hpp b/device/sycl/src/finding/find_tracks.hpp index 53b429fff..1e14ab5aa 100644 --- a/device/sycl/src/finding/find_tracks.hpp +++ b/device/sycl/src/finding/find_tracks.hpp @@ -119,6 +119,7 @@ track_candidate_container_types::buffer find_tracks( measurements.ptr() + n_measurements, uniques.begin(), uniques.begin() + n_modules, upper_bounds.begin(), measurement_sort_comp()); + queue.wait_and_throw(); /***************************************************************** * Kernel1: Create barcode sequence @@ -242,7 +243,7 @@ track_candidate_container_types::buffer find_tracks( shared_num_candidates(nFindTracksThreads, h); vecmem::sycl::local_accessor< std::pair> - shared_candidates(nFindTracksThreads, h); + shared_candidates(2 * nFindTracksThreads, h); vecmem::sycl::local_accessor shared_candidates_size(1, h); @@ -331,13 +332,10 @@ track_candidate_container_types::buffer find_tracks( keys_buffer); vecmem::device_vector param_ids_device( param_ids_buffer); - auto zipped_first = oneapi::dpl::make_zip_iterator( - keys_device.begin(), param_ids_device.begin()); - oneapi::dpl::sort( - policy, zipped_first, zipped_first + keys_device.size(), - [](auto lhs, auto rhs) { - return std::get<0>(lhs) < std::get<0>(rhs); - }); + oneapi::dpl::sort_by_key(policy, keys_device.begin(), + keys_device.end(), + param_ids_device.begin()); + queue.wait_and_throw(); } /***************************************************************** @@ -415,6 +413,7 @@ track_candidate_container_types::buffer find_tracks( oneapi::dpl::copy(policy, in.begin(), in.begin() + n_candidates_per_step[it], out.begin()); } + queue.wait_and_throw(); /***************************************************************** * Kernel6: Build tracks From b174baa26181fa2693e969b50b56aa9c2da09272 Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Wed, 11 Dec 2024 14:47:24 +0100 Subject: [PATCH 11/12] Introduced traccc::sycl::test_queue. Just as a convenience method for seeing in the SYCL unit tests which exact device is used. --- tests/sycl/CMakeLists.txt | 2 ++ tests/sycl/test_ckf_toy_detector.cpp | 11 +++++--- tests/sycl/test_queue.hpp | 40 ++++++++++++++++++++++++++++ tests/sycl/test_queue.sycl | 39 +++++++++++++++++++++++++++ 4 files changed, 88 insertions(+), 4 deletions(-) create mode 100644 tests/sycl/test_queue.hpp create mode 100644 tests/sycl/test_queue.sycl diff --git a/tests/sycl/CMakeLists.txt b/tests/sycl/CMakeLists.txt index 8c62f0c92..bcec6dd61 100644 --- a/tests/sycl/CMakeLists.txt +++ b/tests/sycl/CMakeLists.txt @@ -25,6 +25,8 @@ traccc_add_test( test_sanity_contiguous_on.sycl test_sanity_ordered_on.sycl test_sort.sycl + test_queue.hpp + test_queue.sycl LINK_LIBRARIES GTest::gtest_main diff --git a/tests/sycl/test_ckf_toy_detector.cpp b/tests/sycl/test_ckf_toy_detector.cpp index 3660a53fa..3aeace95f 100644 --- a/tests/sycl/test_ckf_toy_detector.cpp +++ b/tests/sycl/test_ckf_toy_detector.cpp @@ -19,6 +19,7 @@ #include "traccc/utils/ranges.hpp" // Test include(s). +#include "test_queue.hpp" #include "tests/ckf_toy_detector_test.hpp" #include "traccc/utils/seed_generator.hpp" @@ -54,12 +55,14 @@ TEST_P(CkfToyDetectorTests, Run) { * Build a toy detector *****************************/ + // SYCL queue. + sycl::test_queue queue; + // Memory resources used by the application. vecmem::host_memory_resource host_mr; - vecmem::sycl::queue_wrapper queue; - vecmem::sycl::device_memory_resource device_mr{queue}; + vecmem::sycl::device_memory_resource device_mr{queue.queue().queue()}; traccc::memory_resource mr{device_mr, &host_mr}; - vecmem::sycl::shared_memory_resource shared_mr{queue}; + vecmem::sycl::shared_memory_resource shared_mr{queue.queue().queue()}; // Path to the working directory. const std::filesystem::path path = std::filesystem::current_path() / name; @@ -118,7 +121,7 @@ TEST_P(CkfToyDetectorTests, Run) { *****************************/ // Copy objects - vecmem::sycl::async_copy copy{queue}; + vecmem::sycl::async_copy copy{queue.queue().queue()}; traccc::device::container_d2h_copy_alg< traccc::track_candidate_container_types> diff --git a/tests/sycl/test_queue.hpp b/tests/sycl/test_queue.hpp new file mode 100644 index 000000000..2a4483f26 --- /dev/null +++ b/tests/sycl/test_queue.hpp @@ -0,0 +1,40 @@ +/** + * traccc library, part of the ACTS project (R&D line) + * + * (c) 2024 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s). +#include "traccc/sycl/utils/queue_wrapper.hpp" + +// System include(s). +#include + +namespace traccc::sycl { + +/// Queue to use in the SYCL tests +class test_queue { + + public: + /// Default constructor + test_queue(); + /// Destructor + ~test_queue(); + + /// Get the SYCL queue + queue_wrapper queue(); + + private: + /// Internal data type + struct impl; + + /// Pointer to the internal data + std::unique_ptr m_impl; + +}; // struct test_queue + +} // namespace traccc::sycl diff --git a/tests/sycl/test_queue.sycl b/tests/sycl/test_queue.sycl new file mode 100644 index 000000000..39696db33 --- /dev/null +++ b/tests/sycl/test_queue.sycl @@ -0,0 +1,39 @@ +/** + * traccc library, part of the ACTS project (R&D line) + * + * (c) 2024 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Local include(s). +#include "test_queue.hpp" + +// SYCL include(s). +#include + +// System include(s). +#include + +namespace traccc::sycl { + +struct test_queue::impl { + ::sycl::queue m_queue; +}; + +test_queue::test_queue() : m_impl(std::make_unique()) { + + // Print the device name + std::cout + << "Created test queue on device: " + << m_impl->m_queue.get_device().get_info<::sycl::info::device::name>() + << std::endl; +} + +test_queue::~test_queue() = default; + +queue_wrapper test_queue::queue() { + return {&(m_impl->m_queue)}; +} + +} // namespace traccc::sycl From 65fd176953ff36788a5deea357d6f97ec0a791ac Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Tue, 7 Jan 2025 11:34:29 +0100 Subject: [PATCH 12/12] Only running the CKF test on NVIDIA and AMD backends. Taught traccc::sycl::test_queue how to figure out what sort of a queue it is. So that the CKF test could be skipped on OpenCL and Level-0 backends as long as those are still not working. --- tests/sycl/test_ckf_toy_detector.cpp | 5 ++++ tests/sycl/test_queue.hpp | 9 +++++++ tests/sycl/test_queue.sycl | 37 ++++++++++++++++++++++++++++ 3 files changed, 51 insertions(+) diff --git a/tests/sycl/test_ckf_toy_detector.cpp b/tests/sycl/test_ckf_toy_detector.cpp index 3aeace95f..f3dea30b9 100644 --- a/tests/sycl/test_ckf_toy_detector.cpp +++ b/tests/sycl/test_ckf_toy_detector.cpp @@ -58,6 +58,11 @@ TEST_P(CkfToyDetectorTests, Run) { // SYCL queue. sycl::test_queue queue; + // Only run this test on NVIDIA and AMD backends. + if (!(queue.is_cuda() || queue.is_hip())) { + GTEST_SKIP(); + } + // Memory resources used by the application. vecmem::host_memory_resource host_mr; vecmem::sycl::device_memory_resource device_mr{queue.queue().queue()}; diff --git a/tests/sycl/test_queue.hpp b/tests/sycl/test_queue.hpp index 2a4483f26..5b19824a3 100644 --- a/tests/sycl/test_queue.hpp +++ b/tests/sycl/test_queue.hpp @@ -28,6 +28,15 @@ class test_queue { /// Get the SYCL queue queue_wrapper queue(); + /// Check if it's an OpenCL queue + bool is_opencl() const; + /// Check if it's a Level-0 queue + bool is_level0() const; + /// Check if it's a CUDA queue + bool is_cuda() const; + /// Check if it's a HIP queue + bool is_hip() const; + private: /// Internal data type struct impl; diff --git a/tests/sycl/test_queue.sycl b/tests/sycl/test_queue.sycl index 39696db33..7499b0005 100644 --- a/tests/sycl/test_queue.sycl +++ b/tests/sycl/test_queue.sycl @@ -36,4 +36,41 @@ queue_wrapper test_queue::queue() { return {&(m_impl->m_queue)}; } +bool test_queue::is_opencl() const { + +#if SYCL_BACKEND_OPENCL + return (m_impl->m_queue.get_backend() == ::sycl::backend::opencl); +#else + return false; +#endif +} + +bool test_queue::is_level0() const { + +#if SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO + return (m_impl->m_queue.get_backend() == + ::sycl::backend::ext_oneapi_level_zero); +#else + return false; +#endif +} + +bool test_queue::is_cuda() const { + +#if SYCL_EXT_ONEAPI_BACKEND_CUDA + return (m_impl->m_queue.get_backend() == ::sycl::backend::ext_oneapi_cuda); +#else + return false; +#endif +} + +bool test_queue::is_hip() const { + +#if SYCL_EXT_ONEAPI_BACKEND_HIP + return (m_impl->m_queue.get_backend() == ::sycl::backend::ext_oneapi_hip); +#else + return false; +#endif +} + } // namespace traccc::sycl