diff --git a/device/sycl/CMakeLists.txt b/device/sycl/CMakeLists.txt
index a3424f790e..ae0cb0fdf2 100644
--- a/device/sycl/CMakeLists.txt
+++ b/device/sycl/CMakeLists.txt
@@ -13,9 +13,19 @@ 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"
+  # Track fitting algorithm.
+  "include/traccc/sycl/fitting/kalman_fitting_algorithm.hpp"
+  "src/fitting/kalman_fitting_algorithm.cpp"
+  "src/fitting/kalman_fitting_algorithm_constant_field_default_detector.sycl"
+  "src/fitting/kalman_fitting_algorithm_constant_field_telescope_detector.sycl"
+  "src/fitting/fit_tracks.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"
@@ -25,8 +35,6 @@ traccc_add_library( traccc_sycl sycl TYPE SHARED
   "include/traccc/sycl/utils/make_prefix_sum_buff.hpp"
   # 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/fitting/fitting_algorithm.hpp b/device/sycl/include/traccc/sycl/fitting/fitting_algorithm.hpp
deleted file mode 100644
index 2395f990bc..0000000000
--- a/device/sycl/include/traccc/sycl/fitting/fitting_algorithm.hpp
+++ /dev/null
@@ -1,66 +0,0 @@
-/** TRACCC library, part of the ACTS project (R&D line)
- *
- * (c) 2022 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/track_candidate.hpp"
-#include "traccc/edm/track_state.hpp"
-#include "traccc/utils/algorithm.hpp"
-#include "traccc/utils/memory_resource.hpp"
-
-// VecMem include(s).
-#include <vecmem/utils/copy.hpp>
-#include <vecmem/utils/sycl/copy.hpp>
-
-// System include(s).
-#include <functional>
-
-namespace traccc::sycl {
-
-/// Fitting algorithm for a set of tracks
-template <typename fitter_t>
-class fitting_algorithm
-    : public algorithm<track_state_container_types::buffer(
-          const typename fitter_t::detector_type::view_type&,
-          const typename fitter_t::bfield_type&,
-          const typename track_candidate_container_types::const_view&)> {
-
-    public:
-    using algebra_type = typename fitter_t::algebra_type;
-    /// Configuration type
-    using config_type = typename fitter_t::config_type;
-
-    /// Constructor for the fitting algorithm
-    ///
-    /// @param mr The memory resource to use
-    /// @param queue is a wrapper for the sycl queue for kernel invocation
-    fitting_algorithm(const config_type& cfg, const traccc::memory_resource& mr,
-                      queue_wrapper queue);
-
-    /// Run the algorithm
-    track_state_container_types::buffer operator()(
-        const typename fitter_t::detector_type::view_type& det_view,
-        const typename fitter_t::bfield_type& field_view,
-        const typename track_candidate_container_types::const_view&
-            track_candidates_view) const override;
-
-    private:
-    /// Config object
-    config_type m_cfg;
-    /// Memory resource used by the algorithm
-    traccc::memory_resource m_mr;
-    /// Queue wrapper
-    mutable queue_wrapper m_queue;
-    /// Copy object used by the algorithm
-    std::unique_ptr<vecmem::copy> m_copy;
-};
-
-}  // namespace traccc::sycl
diff --git a/device/sycl/include/traccc/sycl/fitting/kalman_fitting_algorithm.hpp b/device/sycl/include/traccc/sycl/fitting/kalman_fitting_algorithm.hpp
new file mode 100644
index 0000000000..0ea43ec697
--- /dev/null
+++ b/device/sycl/include/traccc/sycl/fitting/kalman_fitting_algorithm.hpp
@@ -0,0 +1,95 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2022-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/track_candidate.hpp"
+#include "traccc/edm/track_state.hpp"
+#include "traccc/fitting/fitting_config.hpp"
+#include "traccc/geometry/detector.hpp"
+#include "traccc/utils/algorithm.hpp"
+#include "traccc/utils/memory_resource.hpp"
+
+// Detray include(s).
+#include <detray/detectors/bfield.hpp>
+
+// VecMem include(s).
+#include <vecmem/utils/copy.hpp>
+
+// System include(s).
+#include <functional>
+
+namespace traccc::sycl {
+
+/// Kalman filter based track fitting algorithm
+class kalman_fitting_algorithm
+    : public algorithm<track_state_container_types::buffer(
+          const default_detector::view&,
+          const detray::bfield::const_field_t::view_t&,
+          const track_candidate_container_types::const_view&)>,
+      public algorithm<track_state_container_types::buffer(
+          const telescope_detector::view&,
+          const detray::bfield::const_field_t::view_t&,
+          const track_candidate_container_types::const_view&)> {
+
+    public:
+    /// Configuration type
+    using config_type = fitting_config;
+    /// Output type
+    using output_type = track_state_container_types::buffer;
+
+    /// Constructor with the algorithm's configuration
+    ///
+    /// @param config The configuration object
+    ///
+    kalman_fitting_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 track_candidates All track candidates to fit
+    ///
+    /// @return A container of the fitted track states
+    ///
+    output_type operator()(const default_detector::view& det,
+                           const detray::bfield::const_field_t::view_t& field,
+                           const track_candidate_container_types::const_view&
+                               track_candidates) const override;
+
+    /// Execute the algorithm
+    ///
+    /// @param det             The (telescope) detector object
+    /// @param field           The (constant) magnetic field object
+    /// @param track_candidates All track candidates to fit
+    ///
+    /// @return A container of the fitted track states
+    ///
+    output_type operator()(const telescope_detector::view& det,
+                           const detray::bfield::const_field_t::view_t& field,
+                           const track_candidate_container_types::const_view&
+                               track_candidates) 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<vecmem::copy> m_copy;
+    /// Queue wrapper
+    mutable queue_wrapper m_queue;
+
+};  // class kalman_fitting_algorithm
+
+}  // namespace traccc::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 <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
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 <vecmem/memory/memory_resource.hpp>
-#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.
-///
-template <typename detector_t>
-class spacepoint_formation_algorithm
-    : public algorithm<spacepoint_collection_types::buffer(
-          const typename detector_t::view_type&,
-          const measurement_collection_types::const_view&)> {
-
-    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<vecmem::copy> 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/fitting/fit_tracks.hpp b/device/sycl/src/fitting/fit_tracks.hpp
new file mode 100644
index 0000000000..1c5e063c26
--- /dev/null
+++ b/device/sycl/src/fitting/fit_tracks.hpp
@@ -0,0 +1,136 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2022-2024 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+// Local include(s).
+#include "traccc/sycl/utils/calculate1DimNdRange.hpp"
+
+// Project include(s).
+#include "traccc/edm/device/sort_key.hpp"
+#include "traccc/edm/track_candidate.hpp"
+#include "traccc/edm/track_state.hpp"
+#include "traccc/fitting/device/fill_sort_keys.hpp"
+#include "traccc/fitting/device/fit.hpp"
+#include "traccc/fitting/fitting_config.hpp"
+#include "traccc/utils/memory_resource.hpp"
+
+// VecMem include(s).
+#include <vecmem/utils/copy.hpp>
+
+// oneDPL include(s).
+#include <oneapi/dpl/algorithm>
+#include <oneapi/dpl/execution>
+
+// SYCL include(s).
+#include <CL/sycl.hpp>
+
+namespace traccc::sycl {
+namespace kernels {
+
+/// Identifier for the kernel that fills the sorting keys.
+struct fill_sort_keys;
+
+}  // namespace kernels
+
+namespace details {
+
+template <typename fitter_t, typename fit_kernel_t>
+track_state_container_types::buffer fit_tracks(
+    const typename fitter_t::detector_type::view_type& det_view,
+    const typename fitter_t::bfield_type& field_view,
+    const typename track_candidate_container_types::const_view&
+        track_candidates_view,
+    const fitting_config& config, const memory_resource& mr, vecmem::copy& copy,
+    cl::sycl::queue& queue) {
+
+    // Get the number of tracks.
+    const track_candidate_container_types::const_device::header_vector::
+        size_type n_tracks = copy.get_size(track_candidates_view.headers);
+
+    // Get the number of the track candidates (measurements) in each track.
+    const std::vector<track_candidate_container_types::const_device::
+                          item_vector::value_type::size_type>
+        candidate_sizes = copy.get_sizes(track_candidates_view.items);
+
+    // Create the result buffer.
+    track_state_container_types::buffer track_states_buffer{
+        {n_tracks, mr.main},
+        {candidate_sizes, mr.main, mr.host,
+         vecmem::data::buffer_type::resizable}};
+    vecmem::copy::event_type track_states_headers_setup_event =
+        copy.setup(track_states_buffer.headers);
+    vecmem::copy::event_type track_states_items_setup_event =
+        copy.setup(track_states_buffer.items);
+
+    // Return early, if there are no tracks.
+    if (n_tracks == 0) {
+        track_states_headers_setup_event->wait();
+        track_states_items_setup_event->wait();
+        return track_states_buffer;
+    }
+
+    // Create the buffers for sorting the parameter IDs.
+    vecmem::data::vector_buffer<device::sort_key> keys_buffer(n_tracks,
+                                                              mr.main);
+    vecmem::data::vector_buffer<unsigned int> param_ids_buffer(n_tracks,
+                                                               mr.main);
+    vecmem::copy::event_type keys_setup_event = copy.setup(keys_buffer);
+    vecmem::copy::event_type param_ids_setup_event =
+        copy.setup(param_ids_buffer);
+    keys_setup_event->wait();
+    param_ids_setup_event->wait();
+
+    // The execution range for the two kernels of the function.
+    static constexpr unsigned int localSize = 64;
+    cl::sycl::nd_range<1> range = calculate1DimNdRange(n_tracks, localSize);
+
+    // Fill the keys and param_ids buffers.
+    cl::sycl::event fill_keys_event = queue.submit([&](cl::sycl::handler& h) {
+        h.parallel_for<kernels::fill_sort_keys>(
+            range,
+            [track_candidates_view, keys_view = vecmem::get_data(keys_buffer),
+             param_ids_view = vecmem::get_data(param_ids_buffer)](
+                cl::sycl::nd_item<1> item) {
+                device::fill_sort_keys(item.get_global_linear_id(),
+                                       track_candidates_view, keys_view,
+                                       param_ids_view);
+            });
+    });
+
+    // Sort the key to get the sorted parameter ids
+    vecmem::device_vector<device::sort_key> keys_device(keys_buffer);
+    vecmem::device_vector<unsigned int> param_ids_device(param_ids_buffer);
+    fill_keys_event.wait_and_throw();
+    oneapi::dpl::sort_by_key(oneapi::dpl::execution::dpcpp_default,
+                             keys_device.begin(), keys_device.end(),
+                             param_ids_device.begin());
+
+    // Run the fitting, using the sorted parameter IDs.
+    track_state_container_types::view track_states_view = track_states_buffer;
+    track_states_headers_setup_event->wait();
+    track_states_items_setup_event->wait();
+    queue
+        .submit([&](cl::sycl::handler& h) {
+            h.parallel_for<fit_kernel_t>(
+                range, [det_view, field_view, config, track_candidates_view,
+                        param_ids_view = vecmem::get_data(param_ids_buffer),
+                        track_states_view](cl::sycl::nd_item<1> item) {
+                    device::fit<fitter_t>(item.get_global_linear_id(), det_view,
+                                          field_view, config,
+                                          track_candidates_view, param_ids_view,
+                                          track_states_view);
+                });
+        })
+        .wait_and_throw();
+
+    // Return the fitted tracks.
+    return track_states_buffer;
+}
+
+}  // namespace details
+}  // namespace traccc::sycl
diff --git a/device/sycl/src/fitting/fitting_algorithm.sycl b/device/sycl/src/fitting/fitting_algorithm.sycl
deleted file mode 100644
index fa6945e473..0000000000
--- a/device/sycl/src/fitting/fitting_algorithm.sycl
+++ /dev/null
@@ -1,148 +0,0 @@
-/** TRACCC library, part of the ACTS project (R&D line)
- *
- * (c) 2022-2024 CERN for the benefit of the ACTS project
- *
- * Mozilla Public License Version 2.0
- */
-
-// Project include(s).
-#include "../utils/get_queue.hpp"
-#include "traccc/fitting/device/fill_sort_keys.hpp"
-#include "traccc/fitting/device/fit.hpp"
-#include "traccc/fitting/kalman_filter/kalman_fitter.hpp"
-#include "traccc/sycl/fitting/fitting_algorithm.hpp"
-#include "traccc/sycl/utils/calculate1DimNdRange.hpp"
-
-// detray include(s).
-#include "detray/core/detector_metadata.hpp"
-#include "detray/detectors/bfield.hpp"
-#include "detray/navigation/navigator.hpp"
-#include "detray/propagator/rk_stepper.hpp"
-
-// DPL include(s).
-#include <oneapi/dpl/algorithm>
-#include <oneapi/dpl/execution>
-
-// System include(s).
-#include <vector>
-
-namespace traccc::sycl {
-
-namespace kernels {
-/// Class identifying the kernel running @c
-/// traccc::device::fit
-class fit;
-/// Class identifying the kernel running @c
-/// traccc::device::fill_sort_keys
-class fill_sort_keys;
-}  // namespace kernels
-
-template <typename fitter_t>
-fitting_algorithm<fitter_t>::fitting_algorithm(
-    const config_type& cfg, const traccc::memory_resource& mr,
-    queue_wrapper queue)
-    : m_cfg(cfg), m_mr(mr), m_queue(queue) {
-
-    // Initialize m_copy ptr based on memory resources that were given
-    if (mr.host) {
-        m_copy = std::make_unique<vecmem::sycl::copy>(queue.queue());
-    } else {
-        m_copy = std::make_unique<vecmem::copy>();
-    }
-}
-
-template <typename fitter_t>
-track_state_container_types::buffer fitting_algorithm<fitter_t>::operator()(
-    const typename fitter_t::detector_type::view_type& det_view,
-    const typename fitter_t::bfield_type& field_view,
-    const typename track_candidate_container_types::const_view&
-        track_candidates_view) const {
-
-    // Number of tracks
-    const track_candidate_container_types::const_device::header_vector::
-        size_type n_tracks = m_copy->get_size(track_candidates_view.headers);
-
-    // Get the sizes of the track candidates in each track
-    const std::vector<track_candidate_container_types::const_device::
-                          item_vector::value_type::size_type>
-        candidate_sizes = m_copy->get_sizes(track_candidates_view.items);
-
-    track_state_container_types::buffer track_states_buffer{
-        {n_tracks, m_mr.main},
-        {candidate_sizes, m_mr.main, m_mr.host,
-         vecmem::data::buffer_type::resizable}};
-
-    vecmem::copy::event_type track_states_headers_setup_event =
-        m_copy->setup(track_states_buffer.headers);
-    vecmem::copy::event_type track_states_items_setup_event =
-        m_copy->setup(track_states_buffer.items);
-
-    track_state_container_types::view track_states_view(track_states_buffer);
-
-    // -- localSize
-    // The dimension of workGroup (block) is the integer multiple of WARP_SIZE
-    // (=32)
-    unsigned int localSize = 64;
-
-    vecmem::data::vector_buffer<device::sort_key> keys_buffer(n_tracks,
-                                                              m_mr.main);
-    vecmem::data::vector_buffer<unsigned int> param_ids_buffer(n_tracks,
-                                                               m_mr.main);
-    vecmem::data::vector_view<device::sort_key> keys_view(keys_buffer);
-    vecmem::data::vector_view<unsigned int> param_ids_view(param_ids_buffer);
-
-    // Sort the key to get the sorted parameter ids
-    vecmem::device_vector<device::sort_key> keys_device(keys_buffer);
-    vecmem::device_vector<unsigned int> param_ids_device(param_ids_buffer);
-
-    // 1 dim ND Range for the kernel
-    auto trackParamsNdRange =
-        traccc::sycl::calculate1DimNdRange(n_tracks, localSize);
-
-    details::get_queue(m_queue)
-        .submit([&](::sycl::handler& h) {
-            h.parallel_for<kernels::fill_sort_keys>(
-                trackParamsNdRange, [track_candidates_view, keys_view,
-                                     param_ids_view](::sycl::nd_item<1> item) {
-                    device::fill_sort_keys(item.get_global_linear_id(),
-                                           track_candidates_view, keys_view,
-                                           param_ids_view);
-                });
-        })
-        .wait_and_throw();
-
-    oneapi::dpl::sort_by_key(oneapi::dpl::execution::dpcpp_default,
-                             keys_device.begin(), keys_device.end(),
-                             param_ids_device.begin());
-
-    track_states_headers_setup_event->wait();
-    track_states_items_setup_event->wait();
-    details::get_queue(m_queue)
-        .submit([&](::sycl::handler& h) {
-            h.parallel_for<kernels::fit>(
-                trackParamsNdRange,
-                [det_view, field_view, config = m_cfg, track_candidates_view,
-                 param_ids_view, track_states_view](::sycl::nd_item<1> item) {
-                    device::fit<fitter_t>(item.get_global_linear_id(), det_view,
-                                          field_view, config,
-                                          track_candidates_view, param_ids_view,
-                                          track_states_view);
-                });
-        })
-        .wait_and_throw();
-
-    return track_states_buffer;
-}
-
-// Explicit template instantiation
-using default_detector_type =
-    detray::detector<detray::default_metadata, detray::device_container_types>;
-using default_stepper_type =
-    detray::rk_stepper<covfie::field<detray::bfield::const_bknd_t>::view_t,
-                       default_algebra, detray::constrained_step<>>;
-using default_navigator_type = detray::navigator<const default_detector_type>;
-using default_fitter_type =
-    kalman_fitter<default_stepper_type, default_navigator_type>;
-template class fitting_algorithm<default_fitter_type>;
-
-}  // namespace traccc::sycl
diff --git a/device/sycl/src/fitting/kalman_fitting_algorithm.cpp b/device/sycl/src/fitting/kalman_fitting_algorithm.cpp
new file mode 100644
index 0000000000..715ddbaca0
--- /dev/null
+++ b/device/sycl/src/fitting/kalman_fitting_algorithm.cpp
@@ -0,0 +1,18 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2022-2024 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+// Local include(s).
+#include "traccc/sycl/fitting/kalman_fitting_algorithm.hpp"
+
+namespace traccc::sycl {
+
+kalman_fitting_algorithm::kalman_fitting_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/fitting/kalman_fitting_algorithm_constant_field_default_detector.sycl b/device/sycl/src/fitting/kalman_fitting_algorithm_constant_field_default_detector.sycl
new file mode 100644
index 0000000000..eca547e4ce
--- /dev/null
+++ b/device/sycl/src/fitting/kalman_fitting_algorithm_constant_field_default_detector.sycl
@@ -0,0 +1,48 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2022-2024 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+// Local include(s).
+#include "../utils/get_queue.hpp"
+#include "fit_tracks.hpp"
+#include "traccc/sycl/fitting/kalman_fitting_algorithm.hpp"
+
+// Project include(s).
+#include "traccc/fitting/kalman_filter/kalman_fitter.hpp"
+
+// Detray include(s).
+#include <detray/navigation/navigator.hpp>
+#include <detray/propagator/rk_stepper.hpp>
+
+namespace traccc::sycl {
+namespace kernels {
+
+/// Identifier for the track fitting kernel.
+struct fit_tracks_constant_field_default_detector;
+
+}  // namespace kernels
+
+kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()(
+    const default_detector::view& det,
+    const detray::bfield::const_field_t::view_t& field,
+    const track_candidate_container_types::const_view& track_candidates) const {
+
+    // Construct the fitter type.
+    using stepper_type =
+        detray::rk_stepper<detray::bfield::const_field_t::view_t,
+                           default_detector::device::algebra_type,
+                           detray::constrained_step<>>;
+    using navigator_type = detray::navigator<const default_detector::device>;
+    using fitter_type = kalman_fitter<stepper_type, navigator_type>;
+
+    // Run the track fitting.
+    return details::fit_tracks<
+        fitter_type, kernels::fit_tracks_constant_field_default_detector>(
+        det, field, track_candidates, m_config, m_mr, m_copy.get(),
+        details::get_queue(m_queue));
+}
+
+}  // namespace traccc::sycl
diff --git a/device/sycl/src/fitting/kalman_fitting_algorithm_constant_field_telescope_detector.sycl b/device/sycl/src/fitting/kalman_fitting_algorithm_constant_field_telescope_detector.sycl
new file mode 100644
index 0000000000..073209ef1c
--- /dev/null
+++ b/device/sycl/src/fitting/kalman_fitting_algorithm_constant_field_telescope_detector.sycl
@@ -0,0 +1,48 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2022-2024 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+// Local include(s).
+#include "../utils/get_queue.hpp"
+#include "fit_tracks.hpp"
+#include "traccc/sycl/fitting/kalman_fitting_algorithm.hpp"
+
+// Project include(s).
+#include "traccc/fitting/kalman_filter/kalman_fitter.hpp"
+
+// Detray include(s).
+#include <detray/navigation/navigator.hpp>
+#include <detray/propagator/rk_stepper.hpp>
+
+namespace traccc::sycl {
+namespace kernels {
+
+/// Identifier for the track fitting kernel.
+struct fit_tracks_constant_field_telescope_detector;
+
+}  // namespace kernels
+
+kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()(
+    const telescope_detector::view& det,
+    const detray::bfield::const_field_t::view_t& field,
+    const track_candidate_container_types::const_view& track_candidates) const {
+
+    // Construct the fitter type.
+    using stepper_type =
+        detray::rk_stepper<detray::bfield::const_field_t::view_t,
+                           telescope_detector::device::algebra_type,
+                           detray::constrained_step<>>;
+    using navigator_type = detray::navigator<const telescope_detector::device>;
+    using fitter_type = kalman_fitter<stepper_type, navigator_type>;
+
+    // Run the track fitting.
+    return details::fit_tracks<
+        fitter_type, kernels::fit_tracks_constant_field_telescope_detector>(
+        det, field, track_candidates, m_config, m_mr, m_copy.get(),
+        details::get_queue(m_queue));
+}
+
+}  // namespace traccc::sycl
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 <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
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 <typename detector_t>
-spacepoint_formation_algorithm<detector_t>::spacepoint_formation_algorithm(
-    const traccc::memory_resource& mr, vecmem::copy& copy, queue_wrapper queue)
-    : m_mr(mr), m_copy(copy), m_queue(queue) {}
-
-template <typename detector_t>
-spacepoint_collection_types::buffer
-spacepoint_formation_algorithm<detector_t>::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<detector_t>(
-                        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<default_detector::device>;
-template class spacepoint_formation_algorithm<telescope_detector::device>;
-
-}  // namespace traccc::sycl
diff --git a/examples/run/sycl/full_chain_algorithm.hpp b/examples/run/sycl/full_chain_algorithm.hpp
index 5a280baad0..27215a8588 100644
--- a/examples/run/sycl/full_chain_algorithm.hpp
+++ b/examples/run/sycl/full_chain_algorithm.hpp
@@ -15,7 +15,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"
 
@@ -67,8 +67,7 @@ class full_chain_algorithm
     using navigator_type = detray::navigator<const device_detector_type>;
     /// 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::telescope_metadata<detray::rectangle2D>,
-                         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<device_detector_type>
-        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));