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

@@ -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"

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 <vecmem/utils/copy.hpp>
+
+// System include(s).
+#include <functional>
+
+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<spacepoint_collection_types::buffer(
+          const default_detector::view&,
+          const measurement_collection_types::const_view&)>,
+      public algorithm<spacepoint_collection_types::buffer(
+          const telescope_detector::view&,
+          const measurement_collection_types::const_view&)> {
+
+    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<vecmem::copy> m_copy;
+    /// SYCL queue object
+    mutable queue_wrapper m_queue;
+};
+
+}  // namespace traccc::sycl

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 <vecmem/memory/memory_resource.hpp>
+
+// SYCL include(s).
+#include <CL/sycl.hpp>
+
+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 <typename detector_t>
+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<detector_t>(
+                        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

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

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

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