From b5137563ceaa93bd859eb48445d4427341551e6c Mon Sep 17 00:00:00 2001 From: Attila Krasznahorkay Date: Wed, 30 Oct 2024 17:51:14 +0100 Subject: [PATCH] De-templated the SYCL version of spacepoint formation. In the same way as I did for traccc::host::silicon_pixel_spacepoint_formation_algorithm earlier. --- device/sycl/CMakeLists.txt | 8 +- ...n_pixel_spacepoint_formation_algorithm.hpp | 82 +++++++++++++++++++ .../spacepoint_formation_algorithm.hpp | 66 --------------- .../silicon_pixel_spacepoint_formation.hpp | 82 +++++++++++++++++++ ...n_pixel_spacepoint_formation_algorithm.cpp | 19 +++++ ..._formation_algorithm_default_detector.sycl | 25 ++++++ ...ormation_algorithm_telescope_detector.sycl | 25 ++++++ .../spacepoint_formation_algorithm.sycl | 77 ----------------- examples/run/sycl/full_chain_algorithm.hpp | 5 +- examples/run/sycl/seq_example_sycl.sycl | 5 +- tests/sycl/test_spacepoint_formation.sycl | 9 +- 11 files changed, 246 insertions(+), 157 deletions(-) create mode 100644 device/sycl/include/traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp delete mode 100644 device/sycl/include/traccc/sycl/seeding/spacepoint_formation_algorithm.hpp create mode 100644 device/sycl/src/seeding/silicon_pixel_spacepoint_formation.hpp create mode 100644 device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp create mode 100644 device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm_default_detector.sycl create mode 100644 device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm_telescope_detector.sycl delete mode 100644 device/sycl/src/seeding/spacepoint_formation_algorithm.sycl diff --git a/device/sycl/CMakeLists.txt b/device/sycl/CMakeLists.txt index a3424f790e..bff349cb6b 100644 --- a/device/sycl/CMakeLists.txt +++ b/device/sycl/CMakeLists.txt @@ -13,9 +13,14 @@ enable_language( SYCL ) # Set up the build of the traccc::sycl library. traccc_add_library( traccc_sycl sycl TYPE SHARED + # Spacepoint formation algorithm. + "include/traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" + "src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp" + "src/seeding/silicon_pixel_spacepoint_formation_algorithm_default_detector.sycl" + "src/seeding/silicon_pixel_spacepoint_formation_algorithm_telescope_detector.sycl" + "src/seeding/silicon_pixel_spacepoint_formation.hpp" # header files "include/traccc/sycl/fitting/fitting_algorithm.hpp" - "include/traccc/sycl/seeding/spacepoint_formation_algorithm.hpp" "include/traccc/sycl/seeding/seeding_algorithm.hpp" "include/traccc/sycl/seeding/seed_finding.hpp" "include/traccc/sycl/seeding/spacepoint_binning.hpp" @@ -26,7 +31,6 @@ traccc_add_library( traccc_sycl sycl TYPE SHARED # implementation files "src/clusterization/clusterization_algorithm.sycl" "src/fitting/fitting_algorithm.sycl" - "src/seeding/spacepoint_formation_algorithm.sycl" "src/seeding/seed_finding.sycl" "src/seeding/seeding_algorithm.cpp" "src/seeding/spacepoint_binning.sycl" diff --git a/device/sycl/include/traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp b/device/sycl/include/traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp new file mode 100644 index 0000000000..27773340c6 --- /dev/null +++ b/device/sycl/include/traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp @@ -0,0 +1,82 @@ +/** 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 + */ + +#pragma once + +// Library include(s). +#include "traccc/edm/measurement.hpp" +#include "traccc/edm/spacepoint.hpp" +#include "traccc/geometry/detector.hpp" +#include "traccc/sycl/utils/queue_wrapper.hpp" +#include "traccc/utils/algorithm.hpp" +#include "traccc/utils/memory_resource.hpp" + +// VecMem include(s). +#include + +// System include(s). +#include + +namespace traccc::sycl { + +/// Algorithm forming space points out of measurements +/// +/// This algorithm performs the local-to-global transformation of the 2D +/// measurements made on every detector module, into 3D spacepoint coordinates. +/// +class silicon_pixel_spacepoint_formation_algorithm + : public algorithm, + public algorithm { + + public: + /// Output type + using output_type = spacepoint_collection_types::buffer; + + /// Constructor for spacepoint_formation + /// + /// @param mr is the memory resource + /// + silicon_pixel_spacepoint_formation_algorithm( + const traccc::memory_resource& mr, vecmem::copy& copy, + queue_wrapper queue); + + /// Construct spacepoints from 2D silicon pixel measurements + /// + /// @param det Detector object + /// @param measurements A collection of measurements + /// @return A spacepoint buffer, with one spacepoint for every + /// silicon pixel measurement + /// + output_type operator()(const default_detector::view& det, + const measurement_collection_types::const_view& + measurements) const override; + + /// Construct spacepoints from 2D silicon pixel measurements + /// + /// @param det Detector object + /// @param measurements A collection of measurements + /// @return A spacepoint buffer, with one spacepoint for every + /// silicon pixel measurement + /// + output_type operator()(const telescope_detector::view& det, + const measurement_collection_types::const_view& + measurements) const override; + + private: + /// Memory resource used by the algorithm + traccc::memory_resource m_mr; + /// The copy object to use + std::reference_wrapper m_copy; + /// SYCL queue object + mutable queue_wrapper m_queue; +}; + +} // namespace traccc::sycl diff --git a/device/sycl/include/traccc/sycl/seeding/spacepoint_formation_algorithm.hpp b/device/sycl/include/traccc/sycl/seeding/spacepoint_formation_algorithm.hpp deleted file mode 100644 index b1466d5f38..0000000000 --- a/device/sycl/include/traccc/sycl/seeding/spacepoint_formation_algorithm.hpp +++ /dev/null @@ -1,66 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2023 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -#pragma once - -// Library include(s). -#include "traccc/edm/measurement.hpp" -#include "traccc/edm/spacepoint.hpp" -#include "traccc/sycl/utils/queue_wrapper.hpp" -#include "traccc/utils/algorithm.hpp" -#include "traccc/utils/memory_resource.hpp" - -// VecMem include(s). -#include -#include - -// System include(s). -#include - -namespace traccc::sycl { - -/// Algorithm forming space points out of measurements -/// -/// This algorithm performs the local-to-global transformation of the 2D -/// measurements made on every detector module, into 3D spacepoint coordinates. -/// -template -class spacepoint_formation_algorithm - : public algorithm { - - public: - /// Constructor for spacepoint_formation - /// - /// @param mr the memory resource - /// @param copy vecmem copy object - /// @param queue is a wrapper for the sycl queue for kernel - /// - spacepoint_formation_algorithm(const traccc::memory_resource& mr, - vecmem::copy& copy, queue_wrapper queue); - - /// Callable operator for spacepoint formation - /// - /// @param det_view a detector view object - /// @param measurements a collection of measurements - /// @return a spacepoint collection (buffer) - spacepoint_collection_types::buffer operator()( - const typename detector_t::view_type& det_view, - const measurement_collection_types::const_view& measurements_view) - const override; - - private: - /// Memory resource used by the algorithm - traccc::memory_resource m_mr; - /// The copy object to use - std::reference_wrapper m_copy; - /// SYCL queue object - mutable queue_wrapper m_queue; -}; - -} // namespace traccc::sycl \ No newline at end of file diff --git a/device/sycl/src/seeding/silicon_pixel_spacepoint_formation.hpp b/device/sycl/src/seeding/silicon_pixel_spacepoint_formation.hpp new file mode 100644 index 0000000000..2c3a7d5967 --- /dev/null +++ b/device/sycl/src/seeding/silicon_pixel_spacepoint_formation.hpp @@ -0,0 +1,82 @@ +/** 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 + */ + +#pragma once + +// Local include(s). +#include "../utils/get_queue.hpp" +#include "traccc/sycl/utils/calculate1DimNdRange.hpp" + +// Project include(s). +#include "traccc/edm/measurement.hpp" +#include "traccc/edm/spacepoint.hpp" +#include "traccc/seeding/device/form_spacepoints.hpp" + +// VecMem include(s). +#include + +// SYCL include(s). +#include + +namespace traccc::sycl::details { + +/// Common implementation for the spacepoint formation algorithm's execute +/// functions +/// +/// @tparam detector_t The detector type to use +/// +/// @param det_view The view of the detector to use +/// @param measurements_view The view of the measurements to process +/// @param mr The memory resource to create the output with +/// @param copy The copy object to use for the output buffer +/// @param queue The queue to use for the computation +/// @return A buffer of the created spacepoints +/// +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, + vecmem::memory_resource& mr, vecmem::copy& copy, cl::sycl::queue& queue) { + + // Get the number of measurements. + const measurement_collection_types::const_view::size_type n_measurements = + copy.get_size(measurements_view); + if (n_measurements == 0) { + return {}; + } + + // Create the result buffer. + spacepoint_collection_types::buffer result( + n_measurements, mr, vecmem::data::buffer_type::resizable); + vecmem::copy::event_type spacepoints_setup_event = copy.setup(result); + + // Calculate the range to run the spacepoint formation for. + static constexpr unsigned int localSize = 32 * 2; + auto countRange = calculate1DimNdRange(n_measurements, localSize); + + // Wait for the output buffer to be ready. + spacepoints_setup_event->wait(); + + // Run the spacepoint formation on the device. + queue + .submit([&](cl::sycl::handler& h) { + h.parallel_for( + countRange, [det_view, measurements_view, n_measurements, + spacepoints_view = vecmem::get_data(result)]( + cl::sycl::nd_item<1> item) { + device::form_spacepoints( + item.get_global_linear_id(), det_view, + measurements_view, n_measurements, spacepoints_view); + }); + }) + .wait_and_throw(); + + // Return the spacepoint buffer. + return result; +} + +} // namespace traccc::sycl::details diff --git a/device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp b/device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp new file mode 100644 index 0000000000..0de21fcaa7 --- /dev/null +++ b/device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm.cpp @@ -0,0 +1,19 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2023-2024 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Local include(s). +#include "traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" + +namespace traccc::sycl { + +silicon_pixel_spacepoint_formation_algorithm:: + silicon_pixel_spacepoint_formation_algorithm( + const traccc::memory_resource& mr, vecmem::copy& copy, + queue_wrapper queue) + : m_mr(mr), m_copy(copy), m_queue(queue) {} + +} // namespace traccc::sycl 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 new file mode 100644 index 0000000000..c0ecd08c28 --- /dev/null +++ b/device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm_default_detector.sycl @@ -0,0 +1,25 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2023-2024 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Local include(s). +#include "../utils/get_queue.hpp" +#include "silicon_pixel_spacepoint_formation.hpp" +#include "traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" + +namespace traccc::sycl { + +silicon_pixel_spacepoint_formation_algorithm::output_type +silicon_pixel_spacepoint_formation_algorithm::operator()( + const default_detector::view& det, + const measurement_collection_types::const_view& meas) const { + + return details::silicon_pixel_spacepoint_formation< + default_detector::device>(det, meas, m_mr.main, m_copy, + details::get_queue(m_queue)); +} + +} // 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 new file mode 100644 index 0000000000..746d24cada --- /dev/null +++ b/device/sycl/src/seeding/silicon_pixel_spacepoint_formation_algorithm_telescope_detector.sycl @@ -0,0 +1,25 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2023-2024 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Local include(s). +#include "../utils/get_queue.hpp" +#include "silicon_pixel_spacepoint_formation.hpp" +#include "traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" + +namespace traccc::sycl { + +silicon_pixel_spacepoint_formation_algorithm::output_type +silicon_pixel_spacepoint_formation_algorithm::operator()( + const telescope_detector::view& det, + const measurement_collection_types::const_view& meas) const { + + return details::silicon_pixel_spacepoint_formation< + telescope_detector::device>(det, meas, m_mr.main, m_copy, + details::get_queue(m_queue)); +} + +} // namespace traccc::sycl diff --git a/device/sycl/src/seeding/spacepoint_formation_algorithm.sycl b/device/sycl/src/seeding/spacepoint_formation_algorithm.sycl deleted file mode 100644 index c69547658a..0000000000 --- a/device/sycl/src/seeding/spacepoint_formation_algorithm.sycl +++ /dev/null @@ -1,77 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2023-2024 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -// Local include(s). -#include "../utils/get_queue.hpp" -#include "traccc/seeding/device/form_spacepoints.hpp" -#include "traccc/sycl/seeding/spacepoint_formation_algorithm.hpp" -#include "traccc/sycl/utils/calculate1DimNdRange.hpp" - -// Project include(s). -#include "traccc/geometry/detector.hpp" - -// detray include(s). -#include "detray/core/detector.hpp" -#include "detray/detectors/telescope_metadata.hpp" -#include "detray/geometry/shapes/rectangle2D.hpp" - -namespace traccc::sycl { - -template -spacepoint_formation_algorithm::spacepoint_formation_algorithm( - const traccc::memory_resource& mr, vecmem::copy& copy, queue_wrapper queue) - : m_mr(mr), m_copy(copy), m_queue(queue) {} - -template -spacepoint_collection_types::buffer -spacepoint_formation_algorithm::operator()( - const typename detector_t::view_type& det_view, - const measurement_collection_types::const_view& measurements_view) const { - - // Get the number of measurements. - const measurement_collection_types::const_view::size_type num_measurements = - m_copy.get().get_size(measurements_view); - - // Create the result buffer. - spacepoint_collection_types::buffer spacepoints_buffer( - num_measurements, m_mr.main, vecmem::data::buffer_type::resizable); - vecmem::copy::event_type spacepoints_setup_event = - m_copy.get().setup(spacepoints_buffer); - - // If there are no measurements, we can conclude here. - if (num_measurements == 0) { - return spacepoints_buffer; - } - - spacepoint_collection_types::view spacepoints_view = spacepoints_buffer; - - // Calculate the range to run the doublet counting for. - static constexpr unsigned int measLocalSize = 32 * 2; - auto measCountRange = - traccc::sycl::calculate1DimNdRange(num_measurements, measLocalSize); - - spacepoints_setup_event->wait(); - details::get_queue(m_queue) - .submit([&](::sycl::handler& h) { - h.parallel_for( - measCountRange, [det_view, measurements_view, num_measurements, - spacepoints_view](::sycl::nd_item<1> item) { - device::form_spacepoints( - item.get_global_linear_id(), det_view, - measurements_view, num_measurements, spacepoints_view); - }); - }) - .wait_and_throw(); - - return spacepoints_buffer; -} - -// Explicit template instantiation -template class spacepoint_formation_algorithm; -template class spacepoint_formation_algorithm; - -} // namespace traccc::sycl diff --git a/examples/run/sycl/full_chain_algorithm.hpp b/examples/run/sycl/full_chain_algorithm.hpp index e2f9f30d0a..ee3a99df17 100644 --- a/examples/run/sycl/full_chain_algorithm.hpp +++ b/examples/run/sycl/full_chain_algorithm.hpp @@ -16,7 +16,7 @@ #include "traccc/geometry/silicon_detector_description.hpp" #include "traccc/sycl/clusterization/clusterization_algorithm.hpp" #include "traccc/sycl/seeding/seeding_algorithm.hpp" -#include "traccc/sycl/seeding/spacepoint_formation_algorithm.hpp" +#include "traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" #include "traccc/sycl/seeding/track_params_estimation.hpp" #include "traccc/utils/algorithm.hpp" @@ -68,8 +68,7 @@ class full_chain_algorithm using navigator_type = detray::navigator; /// Spacepoint formation algorithm type using spacepoint_formation_algorithm = - traccc::sycl::spacepoint_formation_algorithm< - traccc::default_detector::device>; + traccc::sycl::silicon_pixel_spacepoint_formation_algorithm; /// Clustering algorithm type using clustering_algorithm = clusterization_algorithm; /// Track finding algorithm type diff --git a/examples/run/sycl/seq_example_sycl.sycl b/examples/run/sycl/seq_example_sycl.sycl index 57c5311d2c..68d730b08b 100644 --- a/examples/run/sycl/seq_example_sycl.sycl +++ b/examples/run/sycl/seq_example_sycl.sycl @@ -21,7 +21,7 @@ #include "traccc/seeding/track_params_estimation.hpp" #include "traccc/sycl/clusterization/clusterization_algorithm.hpp" #include "traccc/sycl/seeding/seeding_algorithm.hpp" -#include "traccc/sycl/seeding/spacepoint_formation_algorithm.hpp" +#include "traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" #include "traccc/sycl/seeding/track_params_estimation.hpp" // performance @@ -134,8 +134,7 @@ int seq_run(const traccc::opts::detector& detector_opts, using host_spacepoint_formation_algorithm = traccc::host::silicon_pixel_spacepoint_formation_algorithm; using device_spacepoint_formation_algorithm = - traccc::sycl::spacepoint_formation_algorithm< - traccc::default_detector::device>; + traccc::sycl::silicon_pixel_spacepoint_formation_algorithm; // Constant B field for the track finding and fitting const traccc::vector3 field_vec = {0.f, 0.f, diff --git a/tests/sycl/test_spacepoint_formation.sycl b/tests/sycl/test_spacepoint_formation.sycl index 757432dd93..76ac5cece8 100644 --- a/tests/sycl/test_spacepoint_formation.sycl +++ b/tests/sycl/test_spacepoint_formation.sycl @@ -8,7 +8,7 @@ // Project include(s). #include "traccc/definitions/common.hpp" #include "traccc/edm/spacepoint.hpp" -#include "traccc/sycl/seeding/spacepoint_formation_algorithm.hpp" +#include "traccc/sycl/seeding/silicon_pixel_spacepoint_formation_algorithm.hpp" // Detray include(s). #include "detray/geometry/mask.hpp" @@ -71,9 +71,6 @@ TEST(SYCLSpacepointFormation, sycl) { // Create telescope geometry auto [det, name_map] = build_telescope_detector(shared_mr, tel_cfg); - using device_detector_type = - detray::detector, - detray::device_container_types>; // Surface lookup auto surfaces = det.surfaces(); @@ -88,8 +85,8 @@ TEST(SYCLSpacepointFormation, sycl) { measurements.push_back({{10.f, 15.f}, {0.f, 0.f}, surfaces[8u].barcode()}); // Run spacepoint formation - traccc::sycl::spacepoint_formation_algorithm - sp_formation(mr, copy, &q); + traccc::sycl::silicon_pixel_spacepoint_formation_algorithm sp_formation( + mr, copy, &q); auto spacepoints_buffer = sp_formation(detray::get_data(det), vecmem::get_data(measurements));