From d4921b7371e18a533fb9eafc757f771726b88823 Mon Sep 17 00:00:00 2001
From: Attila Krasznahorkay <Attila.Krasznahorkay@cern.ch>
Date: Fri, 10 Apr 2026 10:42:48 +0200
Subject: [PATCH] Concretized edm::spacepoint.

So that its memory layout would not depend on which algebra the code
is built with.

Adapted all client code to this change.
---
 core/CMakeLists.txt                           |  2 +
 .../traccc/edm/impl/spacepoint_collection.ipp |  6 +--
 .../traccc/edm/impl/spacepoint_helpers.ipp    | 22 ++++++++
 .../traccc/edm/spacepoint_collection.hpp      | 52 ++++++++++++-------
 .../include/traccc/edm/spacepoint_helpers.hpp | 33 ++++++++++++
 .../seeding/impl/spacepoint_formation.ipp     |  6 +--
 .../track_params_estimation_helper.hpp        |  7 ++-
 .../traccc/utils/detray_conversion.hpp        | 20 +++++++
 .../traccc/utils/impl/detray_conversion.ipp   | 23 ++++++++
 .../kernels/GbtsNodesMakingKernels.cuh        |  2 +-
 tests/cpu/test_track_params_estimation.cpp    | 25 ++++++---
 tests/cuda/test_spacepoint_formation.cpp      |  4 +-
 12 files changed, 165 insertions(+), 37 deletions(-)
 create mode 100644 core/include/traccc/edm/impl/spacepoint_helpers.ipp
 create mode 100644 core/include/traccc/edm/spacepoint_helpers.hpp

diff --git a/core/CMakeLists.txt b/core/CMakeLists.txt
index acd76cb359..95bc3ce0a0 100644
--- a/core/CMakeLists.txt
+++ b/core/CMakeLists.txt
@@ -32,6 +32,8 @@ traccc_add_library( traccc_core core TYPE SHARED
   "include/traccc/edm/silicon_cluster_collection.hpp"
   "include/traccc/edm/spacepoint_collection.hpp"
   "include/traccc/edm/impl/spacepoint_collection.ipp"
+  "include/traccc/edm/spacepoint_helpers.hpp"
+  "include/traccc/edm/impl/spacepoint_helpers.ipp"
   "include/traccc/edm/seed_collection.hpp"
   "include/traccc/edm/impl/seed_collection.ipp"
   "include/traccc/edm/track_state_collection.hpp"
diff --git a/core/include/traccc/edm/impl/spacepoint_collection.ipp b/core/include/traccc/edm/impl/spacepoint_collection.ipp
index 577929fdc9..ec6828aff7 100644
--- a/core/include/traccc/edm/impl/spacepoint_collection.ipp
+++ b/core/include/traccc/edm/impl/spacepoint_collection.ipp
@@ -1,6 +1,6 @@
 /** TRACCC library, part of the ACTS project (R&D line)
  *
- * (c) 2021-2025 CERN for the benefit of the ACTS project
+ * (c) 2021-2026 CERN for the benefit of the ACTS project
  *
  * Mozilla Public License Version 2.0
  */
@@ -51,8 +51,8 @@ TRACCC_HOST_DEVICE const auto& spacepoint<BASE>::z() const {
 template <typename BASE>
 TRACCC_HOST_DEVICE auto spacepoint<BASE>::radius() const {
 
-    const scalar xx = x();
-    const scalar yy = y();
+    const float xx = x();
+    const float yy = y();
     return math::sqrt(xx * xx + yy * yy);
 }
 
diff --git a/core/include/traccc/edm/impl/spacepoint_helpers.ipp b/core/include/traccc/edm/impl/spacepoint_helpers.ipp
new file mode 100644
index 0000000000..9ede0865ef
--- /dev/null
+++ b/core/include/traccc/edm/impl/spacepoint_helpers.ipp
@@ -0,0 +1,22 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2026 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+// Local include(s).
+#include "traccc/utils/detray_conversion.hpp"
+
+namespace traccc::edm {
+
+template <detray::concepts::algebra algebra_t, typename spacepoint_backend_t>
+TRACCC_HOST_DEVICE detray::dpoint3D<algebra_t> get_spacepoint_global(
+    const edm::spacepoint<spacepoint_backend_t>& sp) {
+
+    return utils::to_dpoint3D<algebra_t>(sp.global());
+}
+
+}  // namespace traccc::edm
diff --git a/core/include/traccc/edm/spacepoint_collection.hpp b/core/include/traccc/edm/spacepoint_collection.hpp
index caaa9812b7..587e844922 100644
--- a/core/include/traccc/edm/spacepoint_collection.hpp
+++ b/core/include/traccc/edm/spacepoint_collection.hpp
@@ -1,6 +1,6 @@
 /** TRACCC library, part of the ACTS project (R&D line)
  *
- * (c) 2021-2025 CERN for the benefit of the ACTS project
+ * (c) 2021-2026 CERN for the benefit of the ACTS project
  *
  * Mozilla Public License Version 2.0
  */
@@ -8,12 +8,16 @@
 #pragma once
 
 // Local include(s).
-#include "traccc/definitions/primitives.hpp"
 #include "traccc/definitions/qualifiers.hpp"
 
 // VecMem include(s).
 #include <vecmem/edm/container.hpp>
 
+// System include(s).
+#include <array>
+#include <compare>
+#include <limits>
+
 namespace traccc::edm {
 
 /// Interface for the @c traccc::edm::spacepoint_collection class.
@@ -71,12 +75,12 @@ class spacepoint : public BASE {
 
     /// Global / 3D position of the spacepoint
     ///
-    /// @return A (non-const) vector of @c traccc::point3 values
+    /// @return A (non-const) vector of @c std::array<float,3> values
     ///
     TRACCC_HOST_DEVICE auto& global() { return BASE::template get<2>(); }
     /// Global / 3D position of the spacepoint
     ///
-    /// @return A (const) vector of @c traccc::point3 values
+    /// @return A (const) vector of @c std::array<float,3> values
     ///
     TRACCC_HOST_DEVICE const auto& global() const {
         return BASE::template get<2>();
@@ -87,7 +91,7 @@ class spacepoint : public BASE {
     /// @note This function must only be used on proxy objects, not on
     ///       containers!
     ///
-    /// @return A (non-const) reference to a @c traccc::scalar value
+    /// @return A (non-const) reference to a @c float value
     ///
     TRACCC_HOST_DEVICE
     auto& x();
@@ -96,7 +100,7 @@ class spacepoint : public BASE {
     /// @note This function must only be used on proxy objects, not on
     ///       containers!
     ///
-    /// @return A (const) reference to a @c traccc::scalar value
+    /// @return A (const) reference to a @c float value
     ///
     TRACCC_HOST_DEVICE
     const auto& x() const;
@@ -106,13 +110,13 @@ class spacepoint : public BASE {
     /// @note This function must only be used on proxy objects, not on
     ///       containers!
     ///
-    /// @return A (non-const) reference to a @c traccc::scalar value
+    /// @return A (non-const) reference to a @c float value
     ///
     TRACCC_HOST_DEVICE
     auto& y();
     /// The Y position of the spacepoint (const)
     ///
-    /// @return A (const) reference to a @c traccc::scalar value
+    /// @return A (const) reference to a @c float value
     ///
     TRACCC_HOST_DEVICE
     const auto& y() const;
@@ -122,7 +126,7 @@ class spacepoint : public BASE {
     /// @note This function must only be used on proxy objects, not on
     ///       containers!
     ///
-    /// @return A (non-const) reference to a @c traccc::scalar value
+    /// @return A (non-const) reference to a @c float value
     ///
     TRACCC_HOST_DEVICE
     auto& z();
@@ -131,19 +135,19 @@ class spacepoint : public BASE {
     /// @note This function must only be used on proxy objects, not on
     ///       containers!
     ///
-    /// @return A (const) reference to a @c traccc::scalar value
+    /// @return A (const) reference to a @c float value
     ///
     TRACCC_HOST_DEVICE
     const auto& z() const;
 
     /// The variation on the spacepoint's Z coordinate (non-const)
     ///
-    /// @return A (non-const) vector of @c traccc::scalar values
+    /// @return A (non-const) vector of @c float values
     ///
     TRACCC_HOST_DEVICE auto& z_variance() { return BASE::template get<3>(); }
     /// The variation on the spacepoint's Z coordinate (const)
     ///
-    /// @return A (const) vector of @c traccc::scalar values
+    /// @return A (const) vector of @c float values
     ///
     TRACCC_HOST_DEVICE const auto& z_variance() const {
         return BASE::template get<3>();
@@ -154,20 +158,20 @@ class spacepoint : public BASE {
     /// @note This function must only be used on proxy objects, not on
     ///       containers!
     ///
-    /// @return A @c traccc::scalar value
+    /// @return A @c float value
     ///
     TRACCC_HOST_DEVICE auto radius() const;
 
     /// The variation on the spacepoint radious (non-const)
     ///
-    /// @return A (non-const) vector of @c traccc::scalar values
+    /// @return A (non-const) vector of @c float values
     ///
     TRACCC_HOST_DEVICE auto& radius_variance() {
         return BASE::template get<4>();
     }
     /// The variation on the spacepoint radious (const)
     ///
-    /// @return A (non-const) vector of @c traccc::scalar values
+    /// @return A (non-const) vector of @c float values
     ///
     TRACCC_HOST_DEVICE const auto& radius_variance() const {
         return BASE::template get<4>();
@@ -178,7 +182,7 @@ class spacepoint : public BASE {
     /// @note This function must only be used on proxy objects, not on
     ///       containers!
     ///
-    /// @return A @c traccc::scalar value
+    /// @return A @c float value
     ///
     TRACCC_HOST_DEVICE auto phi() const;
 
@@ -216,10 +220,18 @@ class spacepoint : public BASE {
 };  // class spacepoint
 
 /// SoA container describing reconstructed spacepoints
-using spacepoint_collection = vecmem::edm::container<
-    spacepoint, vecmem::edm::type::vector<unsigned int>,
-    vecmem::edm::type::vector<unsigned int>, vecmem::edm::type::vector<point3>,
-    vecmem::edm::type::vector<scalar>, vecmem::edm::type::vector<scalar> >;
+using spacepoint_collection =
+    vecmem::edm::container<spacepoint,
+                           // measurement_index_1
+                           vecmem::edm::type::vector<unsigned int>,
+                           // measurement_index_2
+                           vecmem::edm::type::vector<unsigned int>,
+                           // global
+                           vecmem::edm::type::vector<std::array<float, 3u>>,
+                           // z_variance
+                           vecmem::edm::type::vector<float>,
+                           // radius_variance
+                           vecmem::edm::type::vector<float>>;
 
 }  // namespace traccc::edm
 
diff --git a/core/include/traccc/edm/spacepoint_helpers.hpp b/core/include/traccc/edm/spacepoint_helpers.hpp
new file mode 100644
index 0000000000..1e3eafaefd
--- /dev/null
+++ b/core/include/traccc/edm/spacepoint_helpers.hpp
@@ -0,0 +1,33 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2026 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+// Local include(s).
+#include "traccc/definitions/qualifiers.hpp"
+#include "traccc/edm/spacepoint_collection.hpp"
+
+// Detray include(s).
+#include <detray/definitions/algebra.hpp>
+
+namespace traccc::edm {
+
+/// Get the global position of a spacepoint as a 3D point
+///
+/// @tparam algebra_t The algebra type used to describe the tracks
+///
+/// @param sp The spacepoint to extract the global position from
+/// @param pos The 3D point to fill with the global position of the spacepoint
+///
+template <detray::concepts::algebra algebra_t, typename spacepoint_backend_t>
+TRACCC_HOST_DEVICE detray::dpoint3D<algebra_t> get_spacepoint_global(
+    const edm::spacepoint<spacepoint_backend_t>& sp);
+
+}  // namespace traccc::edm
+
+// Implementation include(s).
+#include "traccc/edm/impl/spacepoint_helpers.ipp"
diff --git a/core/include/traccc/seeding/impl/spacepoint_formation.ipp b/core/include/traccc/seeding/impl/spacepoint_formation.ipp
index 0cd48b6411..dfbf32ec78 100644
--- a/core/include/traccc/seeding/impl/spacepoint_formation.ipp
+++ b/core/include/traccc/seeding/impl/spacepoint_formation.ipp
@@ -39,9 +39,9 @@ TRACCC_HOST_DEVICE inline void fill_pixel_spacepoint(
             {});
 
     // Fill the spacepoint with the global position and the measurement.
-    sp.x() = getter::element(global, 0u);
-    sp.y() = getter::element(global, 1u);
-    sp.z() = getter::element(global, 2u);
+    sp.x() = static_cast<float>(getter::element(global, 0u));
+    sp.y() = static_cast<float>(getter::element(global, 1u));
+    sp.z() = static_cast<float>(getter::element(global, 2u));
     sp.radius_variance() = 0.f;
     sp.z_variance() = 0.f;
 }
diff --git a/core/include/traccc/seeding/track_params_estimation_helper.hpp b/core/include/traccc/seeding/track_params_estimation_helper.hpp
index fdeaeae32a..0a9647bd01 100644
--- a/core/include/traccc/seeding/track_params_estimation_helper.hpp
+++ b/core/include/traccc/seeding/track_params_estimation_helper.hpp
@@ -12,6 +12,7 @@
 #include "traccc/edm/measurement_collection.hpp"
 #include "traccc/edm/seed_collection.hpp"
 #include "traccc/edm/spacepoint_collection.hpp"
+#include "traccc/edm/spacepoint_helpers.hpp"
 #include "traccc/edm/track_parameters.hpp"
 
 // System include(s).
@@ -57,8 +58,10 @@ inline TRACCC_HOST_DEVICE void seed_to_bound_param_vector(
     const edm::spacepoint_collection::const_device::const_proxy_type spT =
         spacepoints.at(seed.top_index());
 
-    std::array<vector3, 3> sp_global_positions{spB.global(), spM.global(),
-                                               spT.global()};
+    std::array<vector3, 3> sp_global_positions{
+        edm::get_spacepoint_global<default_algebra>(spB),
+        edm::get_spacepoint_global<default_algebra>(spM),
+        edm::get_spacepoint_global<default_algebra>(spT)};
 
     // Define a new coordinate frame with its origin at the bottom space
     // point, z axis long the magnetic field direction and y axis
diff --git a/core/include/traccc/utils/detray_conversion.hpp b/core/include/traccc/utils/detray_conversion.hpp
index 9884144891..0b18b2002b 100644
--- a/core/include/traccc/utils/detray_conversion.hpp
+++ b/core/include/traccc/utils/detray_conversion.hpp
@@ -28,6 +28,16 @@ template <detray::concepts::algebra ALGEBRA_TYPE>
 TRACCC_HOST_DEVICE std::array<float, 2u> to_float_array(
     const detray::dpoint2D<ALGEBRA_TYPE>& point);
 
+/// Convert an algebra specific 3D point to a float array of size 3
+///
+/// @tparam ALGEBRA_TYPE The algebra type of the input point
+/// @param point The input 3D point to be converted
+/// @return An @c std::array of size 3 containing the converted float values
+///
+template <detray::concepts::algebra ALGEBRA_TYPE>
+TRACCC_HOST_DEVICE std::array<float, 3u> to_float_array(
+    const detray::dpoint3D<ALGEBRA_TYPE>& point);
+
 /// Convert a float array of size 2 to an algebra specific 2D point
 ///
 /// @tparam ALGEBRA_TYPE The algebra type of the output point
@@ -38,6 +48,16 @@ template <detray::concepts::algebra ALGEBRA_TYPE>
 TRACCC_HOST_DEVICE detray::dpoint2D<ALGEBRA_TYPE> to_dpoint2D(
     const std::array<float, 2u>& arr);
 
+/// Convert a float array of size 3 to an algebra specific 3D point
+///
+/// @tparam ALGEBRA_TYPE The algebra type of the output point
+/// @param arr The input @c std::array of size 3 containing the float values
+/// @return A 3D point of type @c detray::dpoint3D with the converted values
+///
+template <detray::concepts::algebra ALGEBRA_TYPE>
+TRACCC_HOST_DEVICE detray::dpoint3D<ALGEBRA_TYPE> to_dpoint3D(
+    const std::array<float, 3u>& arr);
+
 }  // namespace traccc::utils
 
 // Implementation include(s).
diff --git a/core/include/traccc/utils/impl/detray_conversion.ipp b/core/include/traccc/utils/impl/detray_conversion.ipp
index 7a898c224f..9652419b55 100644
--- a/core/include/traccc/utils/impl/detray_conversion.ipp
+++ b/core/include/traccc/utils/impl/detray_conversion.ipp
@@ -17,6 +17,15 @@ TRACCC_HOST_DEVICE std::array<float, 2u> to_float_array(
             static_cast<float>(getter::element(point, 1u))};
 }
 
+template <detray::concepts::algebra ALGEBRA_TYPE>
+TRACCC_HOST_DEVICE std::array<float, 3u> to_float_array(
+    const detray::dpoint3D<ALGEBRA_TYPE>& point) {
+
+    return {static_cast<float>(getter::element(point, 0u)),
+            static_cast<float>(getter::element(point, 1u)),
+            static_cast<float>(getter::element(point, 2u))};
+}
+
 template <detray::concepts::algebra ALGEBRA_TYPE>
 TRACCC_HOST_DEVICE detray::dpoint2D<ALGEBRA_TYPE> to_dpoint2D(
     const std::array<float, 2u>& arr) {
@@ -29,4 +38,18 @@ TRACCC_HOST_DEVICE detray::dpoint2D<ALGEBRA_TYPE> to_dpoint2D(
     return point;
 }
 
+template <detray::concepts::algebra ALGEBRA_TYPE>
+TRACCC_HOST_DEVICE detray::dpoint3D<ALGEBRA_TYPE> to_dpoint3D(
+    const std::array<float, 3u>& arr) {
+
+    detray::dpoint3D<ALGEBRA_TYPE> point;
+    getter::element(point, 0u) =
+        static_cast<typename ALGEBRA_TYPE::value_type>(arr[0u]);
+    getter::element(point, 1u) =
+        static_cast<typename ALGEBRA_TYPE::value_type>(arr[1u]);
+    getter::element(point, 2u) =
+        static_cast<typename ALGEBRA_TYPE::value_type>(arr[2u]);
+    return point;
+}
+
 }  // namespace traccc::utils
diff --git a/device/cuda/src/gbts_seeding/kernels/GbtsNodesMakingKernels.cuh b/device/cuda/src/gbts_seeding/kernels/GbtsNodesMakingKernels.cuh
index 10533456a5..fd07e07c9d 100644
--- a/device/cuda/src/gbts_seeding/kernels/GbtsNodesMakingKernels.cuh
+++ b/device/cuda/src/gbts_seeding/kernels/GbtsNodesMakingKernels.cuh
@@ -86,7 +86,7 @@ __global__ void count_sp_by_layer(
         // count and store x,y,z,cw info
         atomicAdd(&d_layerCounts[layerIdx], 1);
         spacepointsLayer[spIdx] = layerIdx;
-        const traccc::point3 pos = spacepoint.global();
+        const std::array<float, 3u> pos = spacepoint.global();
         reducedSP[spIdx] =
             make_float4(pos[0], pos[1], pos[2], cluster_diameter);
     }
diff --git a/tests/cpu/test_track_params_estimation.cpp b/tests/cpu/test_track_params_estimation.cpp
index 7815c33343..1511f8d744 100644
--- a/tests/cpu/test_track_params_estimation.cpp
+++ b/tests/cpu/test_track_params_estimation.cpp
@@ -11,6 +11,7 @@
 #include "traccc/edm/spacepoint_collection.hpp"
 #include "traccc/seeding/seeding_algorithm.hpp"
 #include "traccc/seeding/track_params_estimation.hpp"
+#include "traccc/utils/detray_conversion.hpp"
 
 // Detray include(s).
 #include <detray/tracks/helix.hpp>
@@ -53,13 +54,19 @@ TEST(track_params_estimation, helix_negative_charge) {
     spacepoints.reserve(3);
     spacepoints.push_back(
         {0, traccc::edm::spacepoint_collection::host::INVALID_MEASUREMENT_INDEX,
-         hlx(50 * unit<scalar>::mm), 0.f, 0.f});
+         traccc::utils::to_float_array<traccc::default_algebra>(
+             hlx(50 * unit<scalar>::mm)),
+         0.f, 0.f});
     spacepoints.push_back(
         {1, traccc::edm::spacepoint_collection::host::INVALID_MEASUREMENT_INDEX,
-         hlx(100 * unit<scalar>::mm), 0.f, 0.f});
+         traccc::utils::to_float_array<traccc::default_algebra>(
+             hlx(100 * unit<scalar>::mm)),
+         0.f, 0.f});
     spacepoints.push_back(
         {2, traccc::edm::spacepoint_collection::host::INVALID_MEASUREMENT_INDEX,
-         hlx(150 * unit<scalar>::mm), 0.f, 0.f});
+         traccc::utils::to_float_array<traccc::default_algebra>(
+             hlx(150 * unit<scalar>::mm)),
+         0.f, 0.f});
 
     // Make a seed from the three spacepoints
     edm::seed_collection::host seeds{host_mr};
@@ -103,13 +110,19 @@ TEST(track_params_estimation, helix_positive_charge) {
     spacepoints.reserve(3);
     spacepoints.push_back(
         {0, traccc::edm::spacepoint_collection::host::INVALID_MEASUREMENT_INDEX,
-         hlx(50 * unit<scalar>::mm), 0.f, 0.f});
+         traccc::utils::to_float_array<traccc::default_algebra>(
+             hlx(50 * unit<scalar>::mm)),
+         0.f, 0.f});
     spacepoints.push_back(
         {1, traccc::edm::spacepoint_collection::host::INVALID_MEASUREMENT_INDEX,
-         hlx(100 * unit<scalar>::mm), 0.f, 0.f});
+         traccc::utils::to_float_array<traccc::default_algebra>(
+             hlx(100 * unit<scalar>::mm)),
+         0.f, 0.f});
     spacepoints.push_back(
         {2, traccc::edm::spacepoint_collection::host::INVALID_MEASUREMENT_INDEX,
-         hlx(150 * unit<scalar>::mm), 0.f, 0.f});
+         traccc::utils::to_float_array<traccc::default_algebra>(
+             hlx(150 * unit<scalar>::mm)),
+         0.f, 0.f});
 
     // Make a seed from the three spacepoints
     edm::seed_collection::host seeds{host_mr};
diff --git a/tests/cuda/test_spacepoint_formation.cpp b/tests/cuda/test_spacepoint_formation.cpp
index 1c5cb2f47a..d2a6fa6fce 100644
--- a/tests/cuda/test_spacepoint_formation.cpp
+++ b/tests/cuda/test_spacepoint_formation.cpp
@@ -101,11 +101,11 @@ TEST(CUDASpacepointFormation, cuda) {
 
     // Check the results
     EXPECT_EQ(copy.get_size(spacepoints_buffer), 2u);
-    std::set<point3> test;
+    std::set<std::array<float, 3u>> test;
     test.insert(spacepoints[0].global());
     test.insert(spacepoints[1].global());
 
-    std::set<point3> ref;
+    std::set<std::array<float, 3u>> ref;
     ref.insert({180.f, 10.f, 15.f});
     ref.insert({20.f, 7.f, 2.f});