diff --git a/device/common/include/traccc/finding/device/build_fitted_tracks.hpp b/device/common/include/traccc/finding/device/build_fitted_tracks.hpp
new file mode 100644
index 0000000000..5620aa7389
--- /dev/null
+++ b/device/common/include/traccc/finding/device/build_fitted_tracks.hpp
@@ -0,0 +1,75 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2023-2025 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+// Local include(s).
+#include "traccc/device/global_index.hpp"
+
+// Project include(s).
+#include "traccc/definitions/qualifiers.hpp"
+#include "traccc/edm/track_candidate_container.hpp"
+#include "traccc/edm/track_parameters.hpp"
+#include "traccc/edm/track_state.hpp"
+#include "traccc/finding/candidate_link.hpp"
+#include "traccc/finding/finding_config.hpp"
+
+// VecMem include(s).
+#include <vecmem/containers/data/jagged_vector_view.hpp>
+#include <vecmem/containers/data/vector_view.hpp>
+
+namespace traccc::device {
+
+/// (Event Data) Payload for the @c traccc::device::build_tracks function
+struct build_fitted_tracks_payload {
+    /**
+     * @brief View objects to the vector of measurements
+     */
+    const measurement_collection_types::const_view measurements_view;
+
+    /**
+     * @brief View object to the vector of seeds
+     */
+    bound_track_parameters_collection_types::const_view seeds_view;
+
+    /**
+     * @brief View object to the track parameters
+     */
+    bound_track_parameters_collection_types::const_view track_param_view;
+
+    /**
+     * @brief View object to the vector of candidate links
+     */
+    vecmem::data::vector_view<const candidate_link> links_view;
+
+    /**
+     * @brief View object to the vector of tips
+     */
+    vecmem::data::vector_view<const unsigned int> tips_view;
+
+    /**
+     * @brief View object to the vector of track candidates
+     */
+    track_state_container_types::view track_states_view;
+};
+
+/// Function for building full tracks from the link container:
+/// The full tracks are built using the link container and tip link container.
+/// Since every link holds an information of the link from the previous step,
+/// we can build a full track by iterating from a tip link backwardly.
+///
+/// @param[in] globalIndex         The index of the current thread
+/// @param[in] cfg                    Track finding config object
+/// @param[inout] payload      The function call payload
+///
+TRACCC_HOST_DEVICE inline void build_fitted_tracks(
+    global_index_t globalIndex, const build_fitted_tracks_payload& payload);
+
+}  // namespace traccc::device
+
+// Include the implementation.
+#include "./impl/build_fitted_tracks.ipp"
diff --git a/device/common/include/traccc/finding/device/build_tracks.hpp b/device/common/include/traccc/finding/device/build_unfitted_tracks.hpp
similarity index 85%
rename from device/common/include/traccc/finding/device/build_tracks.hpp
rename to device/common/include/traccc/finding/device/build_unfitted_tracks.hpp
index f82b4bf7f1..f27b4b4f16 100644
--- a/device/common/include/traccc/finding/device/build_tracks.hpp
+++ b/device/common/include/traccc/finding/device/build_unfitted_tracks.hpp
@@ -23,8 +23,9 @@
 
 namespace traccc::device {
 
-/// (Event Data) Payload for the @c traccc::device::build_tracks function
-struct build_tracks_payload {
+/// (Event Data) Payload for the @c traccc::device::build_unfitted_tracks
+/// function
+struct build_unfitted_tracks_payload {
     /**
      * @brief View object to the vector of measurements
      */
@@ -55,10 +56,10 @@ struct build_tracks_payload {
 /// @param[in] cfg                    Track finding config object
 /// @param[inout] payload      The function call payload
 ///
-TRACCC_HOST_DEVICE inline void build_tracks(
-    global_index_t globalIndex, const build_tracks_payload& payload);
+TRACCC_HOST_DEVICE inline void build_unfitted_tracks(
+    global_index_t globalIndex, const build_unfitted_tracks_payload& payload);
 
 }  // namespace traccc::device
 
 // Include the implementation.
-#include "./impl/build_tracks.ipp"
+#include "./impl/build_unfitted_tracks.ipp"
diff --git a/device/common/include/traccc/finding/device/find_tracks.hpp b/device/common/include/traccc/finding/device/find_tracks.hpp
index 3c12f558c1..ebc49ccaf5 100644
--- a/device/common/include/traccc/finding/device/find_tracks.hpp
+++ b/device/common/include/traccc/finding/device/find_tracks.hpp
@@ -125,6 +125,18 @@ struct find_tracks_payload {
      * @brief View object to the temporary link vector
      */
     vecmem::data::vector_view<candidate_link> tmp_links_view;
+
+    /**
+     * @brief View object to the persistent track parameters, estabilishing
+     * a direct map between links and parameters at the same index.
+     */
+    bound_track_parameters_collection_types::view persistent_parameters_view;
+
+    /**
+     * @brief Flag that, if true, indicates whether holes should be counted
+     * in the length of the tip in the tip length output.
+     */
+    bool count_holes = false;
 };
 
 /// (Shared Event Data) Payload for the @c traccc::device::find_tracks function
diff --git a/device/common/include/traccc/finding/device/impl/build_fitted_tracks.ipp b/device/common/include/traccc/finding/device/impl/build_fitted_tracks.ipp
new file mode 100644
index 0000000000..3b1e24e375
--- /dev/null
+++ b/device/common/include/traccc/finding/device/impl/build_fitted_tracks.ipp
@@ -0,0 +1,87 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2023-2025 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+// Project include(s).
+#include "traccc/utils/prob.hpp"
+
+namespace traccc::device {
+
+TRACCC_HOST_DEVICE inline void build_fitted_tracks(
+    const global_index_t globalIndex,
+    const build_fitted_tracks_payload& payload) {
+
+    const measurement_collection_types::const_device measurements(
+        payload.measurements_view);
+
+    const bound_track_parameters_collection_types::const_device track_params(
+        payload.track_param_view);
+    const bound_track_parameters_collection_types::const_device seeds(
+        payload.seeds_view);
+
+    const vecmem::device_vector<const candidate_link> links(payload.links_view);
+
+    const vecmem::device_vector<const unsigned int> tips(payload.tips_view);
+
+    track_state_container_types::device track_states(payload.track_states_view);
+
+    if (globalIndex >= tips.size()) {
+        return;
+    }
+
+    const auto tip = tips.at(globalIndex);
+
+    auto track = track_states.at(globalIndex).items;
+    auto header = track_states.at(globalIndex).header;
+
+    // Get the link corresponding to tip
+    unsigned int link_idx = tip;
+    auto L = links.at(link_idx);
+    const unsigned int n_meas = measurements.size();
+
+    // Track summary variables
+    scalar ndf_sum = 0.f;
+    scalar chi2_sum = 0.f;
+
+    // Reversely iterate to fill the track candidates
+    for (auto it = track.rbegin(); it != track.rend(); it++) {
+        if (L.meas_idx >= n_meas) {
+            it->is_hole = true;
+        } else {
+            *it = track_state<default_algebra>(measurements.at(L.meas_idx));
+            it->is_hole = false;
+            it->filtered_chi2() = L.chi2;
+            it->filtered() = track_params.at(link_idx);
+
+            // Sanity check on chi2
+            assert(L.chi2 < std::numeric_limits<traccc::scalar>::max());
+            assert(L.chi2 >= 0.f);
+
+            ndf_sum +=
+                static_cast<scalar>(measurements.at(L.meas_idx).meas_dim);
+            chi2_sum += L.chi2;
+        }
+
+        // Break the loop if the iterator is at the first candidate and fill the
+        // seed and track quality
+        if (it != track.rend() - 1) {
+            link_idx = L.previous_candidate_idx;
+            L = links.at(link_idx);
+        }
+    }
+
+    header.fit_outcome = fitter_outcome::SUCCESS;
+    header.fit_params = seeds.at(L.seed_idx);
+    header.trk_quality.ndf = ndf_sum - 5.f;
+    header.trk_quality.chi2 = chi2_sum;
+    header.trk_quality.pval =
+        prob(header.trk_quality.chi2, header.trk_quality.ndf);
+    header.trk_quality.n_holes = L.n_skipped;
+}
+
+}  // namespace traccc::device
diff --git a/device/common/include/traccc/finding/device/impl/build_tracks.ipp b/device/common/include/traccc/finding/device/impl/build_unfitted_tracks.ipp
similarity index 95%
rename from device/common/include/traccc/finding/device/impl/build_tracks.ipp
rename to device/common/include/traccc/finding/device/impl/build_unfitted_tracks.ipp
index 69e8a8db11..3dd298f507 100644
--- a/device/common/include/traccc/finding/device/impl/build_tracks.ipp
+++ b/device/common/include/traccc/finding/device/impl/build_unfitted_tracks.ipp
@@ -12,8 +12,9 @@
 
 namespace traccc::device {
 
-TRACCC_HOST_DEVICE inline void build_tracks(
-    const global_index_t globalIndex, const build_tracks_payload& payload) {
+TRACCC_HOST_DEVICE inline void build_unfitted_tracks(
+    const global_index_t globalIndex,
+    const build_unfitted_tracks_payload& payload) {
 
     const measurement_collection_types::const_device measurements(
         payload.track_candidates_view.measurements);
diff --git a/device/common/include/traccc/finding/device/impl/find_tracks.ipp b/device/common/include/traccc/finding/device/impl/find_tracks.ipp
index 801ecd8527..8e3e1d5977 100644
--- a/device/common/include/traccc/finding/device/impl/find_tracks.ipp
+++ b/device/common/include/traccc/finding/device/impl/find_tracks.ipp
@@ -88,6 +88,8 @@ TRACCC_HOST_DEVICE inline void find_tracks(
     vecmem::device_vector<candidate_link> tmp_links(payload.tmp_links_view);
     bound_track_parameters_collection_types::device tmp_params(
         payload.tmp_params_view);
+    bound_track_parameters_collection_types::device persistent_params(
+        payload.persistent_parameters_view);
     vecmem::device_vector<const detray::geometry::barcode> barcodes(
         payload.barcodes_view);
     vecmem::device_vector<const unsigned int> upper_bounds(
@@ -630,6 +632,10 @@ TRACCC_HOST_DEVICE inline void find_tracks(
                     .chi2 = std::numeric_limits<traccc::scalar>::max(),
                     .chi2_sum = prev_chi2_sum,
                     .ndf_sum = prev_ndf_sum};
+                if (persistent_params.capacity() > 0) {
+                    persistent_params.at(out_offset) =
+                        in_params.at(in_param_id);
+                }
 
                 unsigned int param_pos = out_offset - payload.curr_links_idx;
 
@@ -641,7 +647,12 @@ TRACCC_HOST_DEVICE inline void find_tracks(
 
                 if (n_cands >= cfg.min_track_candidates_per_track) {
                     auto tip_pos = tips.push_back(prev_link_idx);
-                    tip_lengths.at(tip_pos) = n_cands;
+
+                    if (payload.count_holes) {
+                        tip_lengths.at(tip_pos) = payload.step + 1u;
+                    } else {
+                        tip_lengths.at(tip_pos) = n_cands;
+                    }
                 }
             }
         } else {
@@ -659,13 +670,21 @@ TRACCC_HOST_DEVICE inline void find_tracks(
                 out_params_liveness.at(param_pos) =
                     static_cast<unsigned int>(!last_step);
                 links.at(out_offset) = tmp_links.at(in_offset);
+                if (persistent_params.capacity() > 0) {
+                    persistent_params.at(out_offset) = tmp_params.at(in_offset);
+                }
 
                 const unsigned int n_cands = payload.step + 1 - n_skipped;
 
                 if (last_step &&
                     n_cands >= cfg.min_track_candidates_per_track) {
                     auto tip_pos = tips.push_back(param_pos);
-                    tip_lengths.at(tip_pos) = n_cands;
+
+                    if (payload.count_holes) {
+                        tip_lengths.at(tip_pos) = payload.step + 1u;
+                    } else {
+                        tip_lengths.at(tip_pos) = n_cands;
+                    }
                 }
             }
         }
diff --git a/device/common/include/traccc/finding/device/impl/propagate_to_next_surface.ipp b/device/common/include/traccc/finding/device/impl/propagate_to_next_surface.ipp
index a23e555886..a87951a666 100644
--- a/device/common/include/traccc/finding/device/impl/propagate_to_next_surface.ipp
+++ b/device/common/include/traccc/finding/device/impl/propagate_to_next_surface.ipp
@@ -109,7 +109,12 @@ TRACCC_HOST_DEVICE inline void propagate_to_next_surface(
 
         if (n_cands >= cfg.min_track_candidates_per_track) {
             auto tip_pos = tips.push_back(link_idx);
-            tip_lengths.at(tip_pos) = n_cands;
+
+            if (payload.count_holes) {
+                tip_lengths.at(tip_pos) = link.step + 1;
+            } else {
+                tip_lengths.at(tip_pos) = n_cands;
+            }
         }
     }
 }
diff --git a/device/common/include/traccc/finding/device/propagate_to_next_surface.hpp b/device/common/include/traccc/finding/device/propagate_to_next_surface.hpp
index fe827101d3..943bd93d45 100644
--- a/device/common/include/traccc/finding/device/propagate_to_next_surface.hpp
+++ b/device/common/include/traccc/finding/device/propagate_to_next_surface.hpp
@@ -79,6 +79,12 @@ struct propagate_to_next_surface_payload {
      * @brief Vector to hold the number of track states per tip
      */
     vecmem::data::vector_view<unsigned int> tip_lengths_view;
+
+    /**
+     * @brief Flag that, if true, indicates whether holes should be counted
+     * in the length of the tip in the tip length output.
+     */
+    bool count_holes = false;
 };
 
 /// Function for propagating the kalman-updated tracks to the next surface
diff --git a/device/common/include/traccc/finding/device/tags.hpp b/device/common/include/traccc/finding/device/tags.hpp
new file mode 100644
index 0000000000..158afced4b
--- /dev/null
+++ b/device/common/include/traccc/finding/device/tags.hpp
@@ -0,0 +1,28 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2025 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+namespace traccc::device {
+/**
+ * @defgroup Track finding return type specifiers
+ *
+ * Optional parameters to track finding algorithms which instruct the
+ * algorithm to return either unfitted tracks or fitted tracks directly.
+ *
+ * @{
+ * @brief Return tracks with fitted track states.
+ */
+struct finding_return_fitted {};
+/*
+ * @brief Return tracks with unfitted track states.
+ */
+struct finding_return_unfitted {};
+/*
+ * @}
+ */
+}  // namespace traccc::device
diff --git a/device/common/include/traccc/fitting/device/fill_fitting_state_sort_keys.hpp b/device/common/include/traccc/fitting/device/fill_fitting_state_sort_keys.hpp
new file mode 100644
index 0000000000..d443aed6c0
--- /dev/null
+++ b/device/common/include/traccc/fitting/device/fill_fitting_state_sort_keys.hpp
@@ -0,0 +1,50 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2024-2025 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+// Local include(s).
+#include "traccc/device/global_index.hpp"
+#include "traccc/edm/device/sort_key.hpp"
+
+// Project include(s).
+#include "traccc/edm/track_state.hpp"
+
+namespace traccc::device {
+
+/// Function used to fill key container
+///
+/// @param[in] globalIndex   The index of the current thread
+/// @param[in] track_candidates_view The input track states
+/// @param[out] keys_view    The key values
+/// @param[out] ids_view     The param ids
+///
+TRACCC_HOST_DEVICE inline void fill_fitting_state_sort_keys(
+    global_index_t globalIndex,
+    track_state_container_types::const_view& track_states_view,
+    vecmem::data::vector_view<device::sort_key> keys_view,
+    vecmem::data::vector_view<unsigned int> ids_view) {
+    const track_state_container_types::const_device track_states(
+        track_states_view);
+
+    // Keys
+    vecmem::device_vector<device::sort_key> keys_device(keys_view);
+
+    // Param id
+    vecmem::device_vector<unsigned int> ids_device(ids_view);
+
+    if (globalIndex >= keys_device.size()) {
+        return;
+    }
+
+    // Key = The number of measurements
+    keys_device.at(globalIndex) =
+        static_cast<traccc::scalar>(track_states.at(globalIndex).items.size());
+    ids_device.at(globalIndex) = globalIndex;
+}
+
+}  // namespace traccc::device
diff --git a/device/cuda/CMakeLists.txt b/device/cuda/CMakeLists.txt
index a55bbc7785..fd2e0eaf64 100644
--- a/device/cuda/CMakeLists.txt
+++ b/device/cuda/CMakeLists.txt
@@ -66,8 +66,10 @@ traccc_add_library( traccc_cuda cuda TYPE SHARED
   "src/finding/kernels/apply_interaction.hpp"
   "src/finding/kernels/fill_finding_propagation_sort_keys.cu"
   "src/finding/kernels/fill_finding_propagation_sort_keys.cuh"
-  "src/finding/kernels/build_tracks.cu"
-  "src/finding/kernels/build_tracks.cuh"
+  "src/finding/kernels/build_unfitted_tracks.cu"
+  "src/finding/kernels/build_unfitted_tracks.cuh"
+  "src/finding/kernels/build_fitted_tracks.cu"
+  "src/finding/kernels/build_fitted_tracks.cuh"
   "src/finding/kernels/find_tracks.cuh"
   "src/finding/kernels/propagate_to_next_surface.hpp"
   "src/finding/kernels/specializations/find_tracks_default_detector.cu"
@@ -122,6 +124,7 @@ traccc_add_library( traccc_cuda cuda TYPE SHARED
   "src/fitting/kalman_fitting_algorithm_telescope_detector.cu"
   "src/fitting/kalman_fitting.cuh"
   "src/fitting/kernels/fill_fitting_sort_keys.cu"
+  "src/fitting/kernels/fill_fitting_state_sort_keys.cu"
   "src/fitting/kernels/fit_prelude.cu"
   "src/fitting/kernels/specializations/fit_forward_constant_field_default_detector.cu"
   "src/fitting/kernels/specializations/fit_forward_constant_field_telescope_detector.cu"
diff --git a/device/cuda/include/traccc/cuda/finding/combinatorial_kalman_filter_algorithm.hpp b/device/cuda/include/traccc/cuda/finding/combinatorial_kalman_filter_algorithm.hpp
index fc916464e4..781a176ed2 100644
--- a/device/cuda/include/traccc/cuda/finding/combinatorial_kalman_filter_algorithm.hpp
+++ b/device/cuda/include/traccc/cuda/finding/combinatorial_kalman_filter_algorithm.hpp
@@ -14,6 +14,8 @@
 #include "traccc/edm/measurement.hpp"
 #include "traccc/edm/track_candidate_collection.hpp"
 #include "traccc/edm/track_parameters.hpp"
+#include "traccc/edm/track_state.hpp"
+#include "traccc/finding/device/tags.hpp"
 #include "traccc/finding/finding_config.hpp"
 #include "traccc/geometry/detector.hpp"
 #include "traccc/utils/algorithm.hpp"
@@ -39,14 +41,37 @@ class combinatorial_kalman_filter_algorithm
           const telescope_detector::view&, const bfield&,
           const measurement_collection_types::const_view&,
           const bound_track_parameters_collection_types::const_view&)>,
+
+      public algorithm<track_state_container_types::buffer(
+          const default_detector::view&, const bfield&,
+          const measurement_collection_types::const_view&,
+          const bound_track_parameters_collection_types::const_view&,
+          device::finding_return_fitted&&)>,
+      public algorithm<track_state_container_types::buffer(
+          const telescope_detector::view&, const bfield&,
+          const measurement_collection_types::const_view&,
+          const bound_track_parameters_collection_types::const_view&,
+          device::finding_return_fitted&&)>,
+
+      public algorithm<edm::track_candidate_collection<default_algebra>::buffer(
+          const default_detector::view&, const bfield&,
+          const measurement_collection_types::const_view&,
+          const bound_track_parameters_collection_types::const_view&,
+          device::finding_return_unfitted&&)>,
+      public algorithm<edm::track_candidate_collection<default_algebra>::buffer(
+          const telescope_detector::view&, const bfield&,
+          const measurement_collection_types::const_view&,
+          const bound_track_parameters_collection_types::const_view&,
+          device::finding_return_unfitted&&)>,
       public messaging {
 
     public:
     /// Configuration type
     using config_type = finding_config;
-    /// Output type
-    using output_type =
+    /// Output types
+    using unfitted_output_type =
         edm::track_candidate_collection<default_algebra>::buffer;
+    using fitted_output_type = track_state_container_types::buffer;
 
     /// Constructor with the algorithm's configuration
     combinatorial_kalman_filter_algorithm(
@@ -56,35 +81,50 @@ class combinatorial_kalman_filter_algorithm
 
     /// Execute the algorithm
     ///
-    /// @param det          The (default) detector object
+    /// @param det          The detector object
     /// @param field        The magnetic field object
     /// @param measurements All measurements in an event
     /// @param seeds        All seeds in an event to start the track finding
     ///                     with
     ///
-    /// @return A container of the found track candidates
-    ///
-    output_type operator()(
+    /// @{
+    unfitted_output_type operator()(
         const default_detector::view& det, const bfield& field,
         const measurement_collection_types::const_view& measurements,
         const bound_track_parameters_collection_types::const_view& seeds)
-        const override;
-
-    /// Execute the algorithm
-    ///
-    /// @param det          The (telescope) detector object
-    /// @param field        The magnetic field object
-    /// @param measurements All measurements in an event
-    /// @param seeds        All seeds in an event to start the track finding
-    ///                     with
-    ///
-    /// @return A container of the found track candidates
-    ///
-    output_type operator()(
+        const override {
+        return this->operator()(det, field, measurements, seeds,
+                                device::finding_return_unfitted{});
+    }
+    unfitted_output_type operator()(
         const telescope_detector::view& det, const bfield& field,
         const measurement_collection_types::const_view& measurements,
         const bound_track_parameters_collection_types::const_view& seeds)
-        const override;
+        const override {
+        return this->operator()(det, field, measurements, seeds,
+                                device::finding_return_unfitted{});
+    }
+    fitted_output_type operator()(
+        const default_detector::view& det, const bfield& field,
+        const measurement_collection_types::const_view& measurements,
+        const bound_track_parameters_collection_types::const_view& seeds,
+        device::finding_return_fitted&&) const override;
+    fitted_output_type operator()(
+        const telescope_detector::view& det, const bfield& field,
+        const measurement_collection_types::const_view& measurements,
+        const bound_track_parameters_collection_types::const_view& seeds,
+        device::finding_return_fitted&&) const override;
+    unfitted_output_type operator()(
+        const default_detector::view& det, const bfield& field,
+        const measurement_collection_types::const_view& measurements,
+        const bound_track_parameters_collection_types::const_view& seeds,
+        device::finding_return_unfitted&&) const override;
+    unfitted_output_type operator()(
+        const telescope_detector::view& det, const bfield& field,
+        const measurement_collection_types::const_view& measurements,
+        const bound_track_parameters_collection_types::const_view& seeds,
+        device::finding_return_unfitted&&) const override;
+    /// @}
 
     private:
     /// Algorithm configuration
diff --git a/device/cuda/include/traccc/cuda/fitting/kalman_fitting_algorithm.hpp b/device/cuda/include/traccc/cuda/fitting/kalman_fitting_algorithm.hpp
index 0f8151cc78..03e4704151 100644
--- a/device/cuda/include/traccc/cuda/fitting/kalman_fitting_algorithm.hpp
+++ b/device/cuda/include/traccc/cuda/fitting/kalman_fitting_algorithm.hpp
@@ -36,6 +36,12 @@ class kalman_fitting_algorithm
       public algorithm<track_state_container_types::buffer(
           const telescope_detector::view&, const bfield&,
           const edm::track_candidate_container<default_algebra>::const_view&)>,
+      public algorithm<track_state_container_types::buffer(
+          const default_detector::view&, const bfield&,
+          track_state_container_types::buffer&&)>,
+      public algorithm<track_state_container_types::buffer(
+          const telescope_detector::view&, const bfield&,
+          track_state_container_types::buffer&&)>,
       public messaging {
 
     public:
@@ -57,31 +63,41 @@ class kalman_fitting_algorithm
         vecmem::copy& copy, stream& str,
         std::unique_ptr<const Logger> logger = getDummyLogger().clone());
 
-    /// Execute the algorithm
+    /// Execute the algorithm from track candidates
     ///
-    /// @param det             The (default) detector object
+    /// @param det             The detector object
     /// @param field           The 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 bfield& field,
         const edm::track_candidate_container<default_algebra>::const_view&
             track_candidates) const override;
+    output_type operator()(
+        const telescope_detector::view& det, const bfield& field,
+        const edm::track_candidate_container<default_algebra>::const_view&
+            track_candidates) const override;
+    /// @}
 
-    /// Execute the algorithm
+    /// Execute the algorithm from track states
     ///
-    /// @param det             The (telescope) detector object
+    /// @param det             The detector object
     /// @param field           The magnetic field object
-    /// @param track_candidates All track candidates to fit
+    /// @param track_candidates All track states to fit
     ///
     /// @return A container of the fitted track states
     ///
+    /// @{
+    output_type operator()(
+        const default_detector::view& det, const bfield& field,
+        track_state_container_types::buffer&& track_states) const override;
     output_type operator()(
         const telescope_detector::view& det, const bfield& field,
-        const edm::track_candidate_container<default_algebra>::const_view&
-            track_candidates) const override;
+        track_state_container_types::buffer&& track_states) const override;
+    /// @}
 
     private:
     /// Algorithm configuration
diff --git a/device/cuda/src/finding/combinatorial_kalman_filter.cuh b/device/cuda/src/finding/combinatorial_kalman_filter.cuh
index 43c65cc1ce..1eeb660f76 100644
--- a/device/cuda/src/finding/combinatorial_kalman_filter.cuh
+++ b/device/cuda/src/finding/combinatorial_kalman_filter.cuh
@@ -14,7 +14,8 @@
 #include "../utils/thread_id.hpp"
 #include "../utils/utils.hpp"
 #include "./kernels/apply_interaction.hpp"
-#include "./kernels/build_tracks.cuh"
+#include "./kernels/build_fitted_tracks.cuh"
+#include "./kernels/build_unfitted_tracks.cuh"
 #include "./kernels/fill_finding_propagation_sort_keys.cuh"
 #include "./kernels/find_tracks.cuh"
 #include "./kernels/make_barcode_sequence.cuh"
@@ -25,6 +26,7 @@
 #include "traccc/edm/track_candidate_collection.hpp"
 #include "traccc/finding/candidate_link.hpp"
 #include "traccc/finding/details/combinatorial_kalman_filter_types.hpp"
+#include "traccc/finding/device/tags.hpp"
 #include "traccc/finding/finding_config.hpp"
 #include "traccc/utils/logging.hpp"
 #include "traccc/utils/memory_resource.hpp"
@@ -32,6 +34,7 @@
 #include "traccc/utils/propagation.hpp"
 
 // VecMem include(s).
+#include <type_traits>
 #include <vecmem/utils/copy.hpp>
 
 // Thrust include(s).
@@ -67,14 +70,17 @@ namespace traccc::cuda::details {
 ///
 /// @return A buffer of the found track candidates
 ///
-template <typename detector_t, typename bfield_t>
-edm::track_candidate_collection<default_algebra>::buffer
+template <typename detector_t, typename bfield_t, typename return_type_tag>
+std::conditional_t<
+    std::is_same_v<return_type_tag, device::finding_return_fitted>,
+    track_state_container_types::buffer,
+    edm::track_candidate_collection<default_algebra>::buffer>
 combinatorial_kalman_filter(
     const typename detector_t::view_type& det, const bfield_t& field,
     const measurement_collection_types::const_view& measurements,
     const bound_track_parameters_collection_types::const_view& seeds,
     const finding_config& config, const memory_resource& mr, vecmem::copy& copy,
-    const Logger& log, stream& str, unsigned int warp_size) {
+    const Logger& log, stream& str, unsigned int warp_size, return_type_tag&&) {
 
     assert(config.min_step_length_for_next_surface >
                math::fabs(config.propagation.navigation.overstep_tolerance) &&
@@ -165,12 +171,28 @@ combinatorial_kalman_filter(
                                                                        mr.main);
     copy.setup(n_tracks_per_seed_buffer)->ignore();
 
+    constexpr bool making_fitted_params =
+        std::is_same_v<return_type_tag, device::finding_return_fitted>;
+
     // Create a buffer for links
     unsigned int link_buffer_capacity = config.initial_links_per_seed * n_seeds;
     vecmem::data::vector_buffer<candidate_link> links_buffer(
         link_buffer_capacity, mr.main, vecmem::data::buffer_type::resizable);
     copy.setup(links_buffer)->ignore();
 
+    // Create a buffer for track parameters, if we want to output making fitted
+    // tracks.
+    bound_track_parameters_collection_types::buffer persistent_track_parameters;
+    if (making_fitted_params) {
+        persistent_track_parameters =
+            bound_track_parameters_collection_types::buffer(
+                link_buffer_capacity, mr.main);
+    } else {
+        persistent_track_parameters =
+            bound_track_parameters_collection_types::buffer(0u, mr.main);
+    }
+    copy.setup(persistent_track_parameters)->ignore();
+
     // Create a buffer of tip links
     vecmem::data::vector_buffer<unsigned int> tips_buffer{
         config.max_num_branches_per_seed * n_seeds, mr.main,
@@ -251,6 +273,19 @@ combinatorial_kalman_filter(
             copy(links_buffer, new_links_buffer)->wait();
 
             links_buffer = std::move(new_links_buffer);
+
+            if (making_fitted_params) {
+                bound_track_parameters_collection_types::buffer
+                    new_persistent_track_parameters(link_buffer_capacity,
+                                                    mr.main);
+                copy.setup(new_persistent_track_parameters)->ignore();
+                copy(persistent_track_parameters,
+                     new_persistent_track_parameters)
+                    ->wait();
+
+                persistent_track_parameters =
+                    std::move(new_persistent_track_parameters);
+            }
         }
 
         {
@@ -282,7 +317,9 @@ combinatorial_kalman_filter(
                 .tip_lengths_view = tip_length_buffer,
                 .n_tracks_per_seed_view = n_tracks_per_seed_buffer,
                 .tmp_params_view = tmp_params_buffer,
-                .tmp_links_view = tmp_links_buffer};
+                .tmp_links_view = tmp_links_buffer,
+                .persistent_parameters_view = persistent_track_parameters,
+                .count_holes = making_fitted_params};
 
             // The number of threads, blocks and shared memory to use.
             const unsigned int nThreads = warp_size * 2;
@@ -368,7 +405,8 @@ combinatorial_kalman_filter(
                     .step = step,
                     .n_in_params = n_candidates,
                     .tips_view = tips_buffer,
-                    .tip_lengths_view = tip_length_buffer};
+                    .tip_lengths_view = tip_length_buffer,
+                    .count_holes = making_fitted_params};
 
                 const unsigned int nThreads = warp_size * 4;
                 const unsigned int nBlocks =
@@ -407,28 +445,57 @@ combinatorial_kalman_filter(
         tips_length_host.resize(n_tips_total);
     }
 
-    // Create track candidate buffer
-    edm::track_candidate_collection<default_algebra>::buffer
-        track_candidates_buffer{tips_length_host, mr.main, mr.host};
-    copy.setup(track_candidates_buffer)->ignore();
+    if constexpr (std::is_same_v<return_type_tag,
+                                 device::finding_return_unfitted>) {
+        // Create track candidate buffer
+        edm::track_candidate_collection<default_algebra>::buffer
+            track_candidates_buffer{tips_length_host, mr.main, mr.host};
+        copy.setup(track_candidates_buffer)->ignore();
 
-    // @Note: nBlocks can be zero in case there is no tip. This happens when
-    // chi2_max config is set tightly and no tips are found
-    if (n_tips_total > 0) {
-        const unsigned int nThreads = warp_size * 2;
-        const unsigned int nBlocks = (n_tips_total + nThreads - 1) / nThreads;
+        // @Note: nBlocks can be zero in case there is no tip. This happens when
+        // chi2_max config is set tightly and no tips are found
+        if (n_tips_total > 0) {
+            const unsigned int nThreads = warp_size * 2;
+            const unsigned int nBlocks =
+                (n_tips_total + nThreads - 1) / nThreads;
+
+            kernels::build_unfitted_tracks<<<nBlocks, nThreads, 0, stream>>>(
+                {.seeds_view = seeds,
+                 .links_view = links_buffer,
+                 .tips_view = tips_buffer,
+                 .track_candidates_view = {track_candidates_buffer,
+                                           measurements}});
+            TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
 
-        kernels::build_tracks<<<nBlocks, nThreads, 0, stream>>>(
-            {.seeds_view = seeds,
-             .links_view = links_buffer,
-             .tips_view = tips_buffer,
-             .track_candidates_view = {track_candidates_buffer, measurements}});
-        TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
+            str.synchronize();
+        }
 
-        str.synchronize();
-    }
+        return track_candidates_buffer;
+    } else {
+        track_state_container_types::buffer track_states_buffer{
+            {n_tips_total, mr.main}, {tips_length_host, mr.main, mr.host}};
+        copy.setup(track_states_buffer.headers)->ignore();
+        copy.setup(track_states_buffer.items)->ignore();
 
-    return track_candidates_buffer;
+        if (n_tips_total > 0) {
+            const unsigned int nThreads = warp_size * 2;
+            const unsigned int nBlocks =
+                (n_tips_total + nThreads - 1) / nThreads;
+
+            kernels::build_fitted_tracks<<<nBlocks, nThreads, 0, stream>>>(
+                {.measurements_view = measurements,
+                 .seeds_view = seeds,
+                 .track_param_view = persistent_track_parameters,
+                 .links_view = links_buffer,
+                 .tips_view = tips_buffer,
+                 .track_states_view = track_states_buffer});
+            TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
+
+            str.synchronize();
+        }
+
+        return track_states_buffer;
+    }
 }
 
 }  // namespace traccc::cuda::details
diff --git a/device/cuda/src/finding/combinatorial_kalman_filter_algorithm_default_detector.cu b/device/cuda/src/finding/combinatorial_kalman_filter_algorithm_default_detector.cu
index 54840b9e3c..bfef3c1322 100644
--- a/device/cuda/src/finding/combinatorial_kalman_filter_algorithm_default_detector.cu
+++ b/device/cuda/src/finding/combinatorial_kalman_filter_algorithm_default_detector.cu
@@ -9,27 +9,56 @@
 #include "../utils/bfield.cuh"
 #include "combinatorial_kalman_filter.cuh"
 #include "traccc/cuda/finding/combinatorial_kalman_filter_algorithm.hpp"
+#include "traccc/finding/device/tags.hpp"
 
 // System include(s).
 #include <stdexcept>
 
 namespace traccc::cuda {
 
-combinatorial_kalman_filter_algorithm::output_type
+combinatorial_kalman_filter_algorithm::unfitted_output_type
 combinatorial_kalman_filter_algorithm::operator()(
     const default_detector::view& det, const bfield& field,
     const measurement_collection_types::const_view& measurements,
-    const bound_track_parameters_collection_types::const_view& seeds) const {
+    const bound_track_parameters_collection_types::const_view& seeds,
+    device::finding_return_unfitted&&) const {
 
     // Perform the track finding using the templated implementation.
     if (field.is<const_bfield_backend_t<scalar>>()) {
         return details::combinatorial_kalman_filter<default_detector::device>(
             det, field.as<const_bfield_backend_t<scalar>>(), measurements,
-            seeds, m_config, m_mr, m_copy, logger(), m_stream, m_warp_size);
+            seeds, m_config, m_mr, m_copy, logger(), m_stream, m_warp_size,
+            device::finding_return_unfitted{});
     } else if (field.is<cuda::inhom_bfield_backend_t<scalar>>()) {
         return details::combinatorial_kalman_filter<default_detector::device>(
             det, field.as<cuda::inhom_bfield_backend_t<scalar>>(), measurements,
-            seeds, m_config, m_mr, m_copy, logger(), m_stream, m_warp_size);
+            seeds, m_config, m_mr, m_copy, logger(), m_stream, m_warp_size,
+            device::finding_return_unfitted{});
+    } else {
+        throw std::invalid_argument(
+            "Unsupported b-field type received in "
+            "traccc::cuda::combinatorial_kalman_filter_algorithm");
+    }
+}
+
+combinatorial_kalman_filter_algorithm::fitted_output_type
+combinatorial_kalman_filter_algorithm::operator()(
+    const default_detector::view& det, const bfield& field,
+    const measurement_collection_types::const_view& measurements,
+    const bound_track_parameters_collection_types::const_view& seeds,
+    device::finding_return_fitted&&) const {
+
+    // Perform the track finding using the templated implementation.
+    if (field.is<const_bfield_backend_t<scalar>>()) {
+        return details::combinatorial_kalman_filter<default_detector::device>(
+            det, field.as<const_bfield_backend_t<scalar>>(), measurements,
+            seeds, m_config, m_mr, m_copy, logger(), m_stream, m_warp_size,
+            device::finding_return_fitted{});
+    } else if (field.is<cuda::inhom_bfield_backend_t<scalar>>()) {
+        return details::combinatorial_kalman_filter<default_detector::device>(
+            det, field.as<cuda::inhom_bfield_backend_t<scalar>>(), measurements,
+            seeds, m_config, m_mr, m_copy, logger(), m_stream, m_warp_size,
+            device::finding_return_fitted{});
     } else {
         throw std::invalid_argument(
             "Unsupported b-field type received in "
diff --git a/device/cuda/src/finding/combinatorial_kalman_filter_algorithm_telescope_detector.cu b/device/cuda/src/finding/combinatorial_kalman_filter_algorithm_telescope_detector.cu
index 784ae98635..b65071348b 100644
--- a/device/cuda/src/finding/combinatorial_kalman_filter_algorithm_telescope_detector.cu
+++ b/device/cuda/src/finding/combinatorial_kalman_filter_algorithm_telescope_detector.cu
@@ -9,27 +9,56 @@
 #include "../utils/bfield.cuh"
 #include "combinatorial_kalman_filter.cuh"
 #include "traccc/cuda/finding/combinatorial_kalman_filter_algorithm.hpp"
+#include "traccc/finding/device/tags.hpp"
 
 // System include(s).
 #include <stdexcept>
 
 namespace traccc::cuda {
 
-combinatorial_kalman_filter_algorithm::output_type
+combinatorial_kalman_filter_algorithm::unfitted_output_type
 combinatorial_kalman_filter_algorithm::operator()(
     const telescope_detector::view& det, const bfield& field,
     const measurement_collection_types::const_view& measurements,
-    const bound_track_parameters_collection_types::const_view& seeds) const {
+    const bound_track_parameters_collection_types::const_view& seeds,
+    device::finding_return_unfitted&&) const {
 
     // Perform the track finding using the templated implementation.
     if (field.is<const_bfield_backend_t<scalar>>()) {
         return details::combinatorial_kalman_filter<telescope_detector::device>(
             det, field.as<const_bfield_backend_t<scalar>>(), measurements,
-            seeds, m_config, m_mr, m_copy, logger(), m_stream, m_warp_size);
+            seeds, m_config, m_mr, m_copy, logger(), m_stream, m_warp_size,
+            device::finding_return_unfitted{});
     } else if (field.is<cuda::inhom_bfield_backend_t<scalar>>()) {
         return details::combinatorial_kalman_filter<telescope_detector::device>(
             det, field.as<cuda::inhom_bfield_backend_t<scalar>>(), measurements,
-            seeds, m_config, m_mr, m_copy, logger(), m_stream, m_warp_size);
+            seeds, m_config, m_mr, m_copy, logger(), m_stream, m_warp_size,
+            device::finding_return_unfitted{});
+    } else {
+        throw std::invalid_argument(
+            "Unsupported b-field type received in "
+            "traccc::cuda::combinatorial_kalman_filter_algorithm");
+    }
+}
+
+combinatorial_kalman_filter_algorithm::fitted_output_type
+combinatorial_kalman_filter_algorithm::operator()(
+    const telescope_detector::view& det, const bfield& field,
+    const measurement_collection_types::const_view& measurements,
+    const bound_track_parameters_collection_types::const_view& seeds,
+    device::finding_return_fitted&&) const {
+
+    // Perform the track finding using the templated implementation.
+    if (field.is<const_bfield_backend_t<scalar>>()) {
+        return details::combinatorial_kalman_filter<telescope_detector::device>(
+            det, field.as<const_bfield_backend_t<scalar>>(), measurements,
+            seeds, m_config, m_mr, m_copy, logger(), m_stream, m_warp_size,
+            device::finding_return_fitted{});
+    } else if (field.is<cuda::inhom_bfield_backend_t<scalar>>()) {
+        return details::combinatorial_kalman_filter<telescope_detector::device>(
+            det, field.as<cuda::inhom_bfield_backend_t<scalar>>(), measurements,
+            seeds, m_config, m_mr, m_copy, logger(), m_stream, m_warp_size,
+            device::finding_return_fitted{});
     } else {
         throw std::invalid_argument(
             "Unsupported b-field type received in "
diff --git a/device/cuda/src/finding/kernels/build_tracks.cu b/device/cuda/src/finding/kernels/build_fitted_tracks.cu
similarity index 66%
rename from device/cuda/src/finding/kernels/build_tracks.cu
rename to device/cuda/src/finding/kernels/build_fitted_tracks.cu
index ce0d54cda0..af15734f13 100644
--- a/device/cuda/src/finding/kernels/build_tracks.cu
+++ b/device/cuda/src/finding/kernels/build_fitted_tracks.cu
@@ -7,19 +7,20 @@
 
 // Local include(s).
 #include "../../utils/global_index.hpp"
-#include "build_tracks.cuh"
+#include "build_fitted_tracks.cuh"
 
 // Project include(s).
 #include "traccc/edm/measurement.hpp"
 #include "traccc/edm/track_parameters.hpp"
 #include "traccc/finding/candidate_link.hpp"
-#include "traccc/finding/device/build_tracks.hpp"
+#include "traccc/finding/device/build_fitted_tracks.hpp"
 #include "traccc/finding/finding_config.hpp"
 
 namespace traccc::cuda::kernels {
 
-__global__ void build_tracks(device::build_tracks_payload payload) {
+__global__ void build_fitted_tracks(
+    device::build_fitted_tracks_payload payload) {
 
-    device::build_tracks(details::global_index1(), payload);
+    device::build_fitted_tracks(details::global_index1(), payload);
 }
 }  // namespace traccc::cuda::kernels
diff --git a/device/cuda/src/finding/kernels/build_tracks.cuh b/device/cuda/src/finding/kernels/build_fitted_tracks.cuh
similarity index 66%
rename from device/cuda/src/finding/kernels/build_tracks.cuh
rename to device/cuda/src/finding/kernels/build_fitted_tracks.cuh
index b203c3147f..2b66b75562 100644
--- a/device/cuda/src/finding/kernels/build_tracks.cuh
+++ b/device/cuda/src/finding/kernels/build_fitted_tracks.cuh
@@ -8,10 +8,11 @@
 #pragma once
 
 // Project include(s).
-#include "traccc/finding/device/build_tracks.hpp"
+#include "traccc/finding/device/build_fitted_tracks.hpp"
 #include "traccc/finding/finding_config.hpp"
 
 namespace traccc::cuda::kernels {
 
-__global__ void build_tracks(device::build_tracks_payload payload);
+__global__ void build_fitted_tracks(
+    device::build_fitted_tracks_payload payload);
 }
diff --git a/device/cuda/src/finding/kernels/build_unfitted_tracks.cu b/device/cuda/src/finding/kernels/build_unfitted_tracks.cu
new file mode 100644
index 0000000000..cb6902bf26
--- /dev/null
+++ b/device/cuda/src/finding/kernels/build_unfitted_tracks.cu
@@ -0,0 +1,26 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2023-2025 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+// Local include(s).
+#include "../../utils/global_index.hpp"
+#include "build_unfitted_tracks.cuh"
+
+// Project include(s).
+#include "traccc/edm/measurement.hpp"
+#include "traccc/edm/track_parameters.hpp"
+#include "traccc/finding/candidate_link.hpp"
+#include "traccc/finding/device/build_unfitted_tracks.hpp"
+#include "traccc/finding/finding_config.hpp"
+
+namespace traccc::cuda::kernels {
+
+__global__ void build_unfitted_tracks(
+    device::build_unfitted_tracks_payload payload) {
+
+    device::build_unfitted_tracks(details::global_index1(), payload);
+}
+}  // namespace traccc::cuda::kernels
diff --git a/device/cuda/src/finding/kernels/build_unfitted_tracks.cuh b/device/cuda/src/finding/kernels/build_unfitted_tracks.cuh
new file mode 100644
index 0000000000..554ee276c0
--- /dev/null
+++ b/device/cuda/src/finding/kernels/build_unfitted_tracks.cuh
@@ -0,0 +1,18 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2023-2025 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+// Project include(s).
+#include "traccc/finding/device/build_unfitted_tracks.hpp"
+#include "traccc/finding/finding_config.hpp"
+
+namespace traccc::cuda::kernels {
+
+__global__ void build_unfitted_tracks(
+    device::build_unfitted_tracks_payload payload);
+}
diff --git a/device/cuda/src/fitting/kalman_fitting.cuh b/device/cuda/src/fitting/kalman_fitting.cuh
index fc558bf77d..56df38ae4d 100644
--- a/device/cuda/src/fitting/kalman_fitting.cuh
+++ b/device/cuda/src/fitting/kalman_fitting.cuh
@@ -12,6 +12,7 @@
 #include "../utils/global_index.hpp"
 #include "../utils/utils.hpp"
 #include "./kernels/fill_fitting_sort_keys.hpp"
+#include "./kernels/fill_fitting_state_sort_keys.hpp"
 #include "./kernels/fit_backward.hpp"
 #include "./kernels/fit_forward.hpp"
 #include "./kernels/fit_prelude.hpp"
@@ -34,7 +35,8 @@
 
 namespace traccc::cuda::details {
 
-/// Templated implementation of the CUDA track fitting algorithm.
+/// Templated implementation of the CUDA track fitting algorithm starting from
+/// track states.
 ///
 /// @tparam detector_t The (device) detector type to use
 /// @tparam bfield_t   The magnetic field type to use
@@ -50,32 +52,22 @@ namespace traccc::cuda::details {
 /// @return A container of the fitted track states
 ///
 template <typename detector_t, typename bfield_t>
-track_state_container_types::buffer kalman_fitting(
+track_state_container_types::buffer kalman_fitting_from_states(
     const typename detector_t::view_type& det_view, const bfield_t& field_view,
-    const typename edm::track_candidate_container<
-        typename detector_t::algebra_type>::const_view& track_candidates_view,
+    track_state_container_types::buffer&& track_states_buffer,
     const fitting_config& config, const memory_resource& mr, vecmem::copy& copy,
-    stream& str, unsigned int warp_size) {
+    stream& str, unsigned int warp_size, bool states_are_fit) {
 
     // Get a convenience variable for the stream that we'll be using.
     cudaStream_t stream = details::get_stream(str);
 
+    track_state_container_types::view track_states_view(track_states_buffer);
+
     // Get the number of tracks.
-    const edm::track_candidate_collection<
-        default_algebra>::const_device::size_type n_tracks =
-        copy.get_size(track_candidates_view.tracks);
+    const auto n_tracks = copy.get_size(track_states_view.headers);
 
     // Get the sizes of the track candidates in each track.
-    const std::vector<unsigned int> candidate_sizes =
-        copy.get_sizes(track_candidates_view.tracks);
-
-    // 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}};
-    copy.setup(track_states_buffer.headers)->ignore();
-    copy.setup(track_states_buffer.items)->ignore();
+    const auto candidate_sizes = copy.get_sizes(track_states_view.items);
 
     // Return early, if there are no tracks.
     if (n_tracks == 0) {
@@ -113,9 +105,8 @@ track_state_container_types::buffer kalman_fitting(
     const unsigned int nBlocks = (n_tracks + nThreads - 1) / nThreads;
 
     // Fill the keys and param_ids buffers.
-    fill_fitting_sort_keys(nBlocks, nThreads, stream,
-                           track_candidates_view.tracks, keys_buffer,
-                           param_ids_buffer);
+    fill_fitting_state_sort_keys(nBlocks, nThreads, stream, track_states_buffer,
+                                 keys_buffer, param_ids_buffer);
 
     // Sort the key to get the sorted parameter ids
     vecmem::device_vector<device::sort_key> keys_device(keys_buffer);
@@ -125,10 +116,6 @@ track_state_container_types::buffer kalman_fitting(
             .on(stream),
         keys_device.begin(), keys_device.end(), param_ids_device.begin());
 
-    // Run the fitting, using the sorted parameter IDs.
-    fit_prelude(nBlocks, nThreads, 0, stream, param_ids_buffer,
-                track_candidates_view, track_states_buffer,
-                param_liveness_buffer);
     TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
     str.synchronize();
 
@@ -144,9 +131,11 @@ track_state_container_types::buffer kalman_fitting(
 
     for (std::size_t i = 0; i < config.n_iterations; ++i) {
         // Run the track fitting
-        fit_forward<fitter_t>(nBlocks, nThreads, 0, stream, config,
-                              host_payload);
-        TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
+        if (i > 0 || !states_are_fit) {
+            fit_forward<fitter_t>(nBlocks, nThreads, 0, stream, config,
+                                  host_payload);
+            TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
+        }
         fit_backward<fitter_t>(nBlocks, nThreads, 0, stream, config,
                                host_payload);
         TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
@@ -156,4 +145,98 @@ track_state_container_types::buffer kalman_fitting(
     return track_states_buffer;
 }
 
+/// Templated implementation of the CUDA track fitting algorithm starting from
+/// track candidates.
+///
+/// @tparam detector_t The (device) detector type to use
+/// @tparam bfield_t   The magnetic field type to use
+///
+/// @param[in] det_view     A view of the detector geometry
+/// @param[in] field_view   A view of the magnetic field
+/// @param[in] track_candidates_view All track candidates to fit
+/// @param[in] config       The fitting configuration
+/// @param[in] mr           Memory resource(s) to use
+/// @param[in] copy         The copy object to use for memory transfers
+/// @param[in] queue        The Alpaka queue to use for execution
+///
+/// @return A container of the fitted track states
+///
+template <typename detector_t, typename bfield_t>
+track_state_container_types::buffer kalman_fitting_from_candidates(
+    const typename detector_t::view_type& det_view, const bfield_t& field_view,
+    const typename edm::track_candidate_container<
+        typename detector_t::algebra_type>::const_view& track_candidates_view,
+    const fitting_config& config, const memory_resource& mr, vecmem::copy& copy,
+    stream& str, unsigned int warp_size) {
+
+    // Get a convenience variable for the stream that we'll be using.
+    cudaStream_t stream = details::get_stream(str);
+
+    // Get the number of tracks.
+    const edm::track_candidate_collection<
+        default_algebra>::const_device::size_type n_tracks =
+        copy.get_size(track_candidates_view.tracks);
+
+    // Get the sizes of the track candidates in each track.
+    const std::vector<unsigned int> candidate_sizes =
+        copy.get_sizes(track_candidates_view.tracks);
+
+    // 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}};
+    copy.setup(track_states_buffer.headers)->ignore();
+    copy.setup(track_states_buffer.items)->ignore();
+
+    // Return early, if there are no tracks.
+    if (n_tracks == 0) {
+        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::data::vector_buffer<unsigned int> param_liveness_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);
+    vecmem::copy::event_type param_liveness_setup_event =
+        copy.setup(param_liveness_buffer);
+    keys_setup_event->ignore();
+    param_ids_setup_event->ignore();
+    param_liveness_setup_event->ignore();
+
+    // Launch parameters for all the kernels.
+    const unsigned int nThreads = warp_size * 4;
+    const unsigned int nBlocks = (n_tracks + nThreads - 1) / nThreads;
+
+    // Fill the keys and param_ids buffers.
+    fill_fitting_sort_keys(nBlocks, nThreads, stream,
+                           track_candidates_view.tracks, keys_buffer,
+                           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);
+    thrust::sort_by_key(
+        thrust::cuda::par_nosync(std::pmr::polymorphic_allocator(&mr.main))
+            .on(stream),
+        keys_device.begin(), keys_device.end(), param_ids_device.begin());
+
+    // Run the fitting, using the sorted parameter IDs.
+    fit_prelude(nBlocks, nThreads, 0, stream, param_ids_buffer,
+                track_candidates_view, track_states_buffer,
+                param_liveness_buffer);
+    TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
+    str.synchronize();
+
+    return kalman_fitting_from_states<detector_t>(
+        det_view, field_view, std::move(track_states_buffer), config, mr, copy,
+        str, warp_size, false);
+}
+
 }  // namespace traccc::cuda::details
diff --git a/device/cuda/src/fitting/kalman_fitting_algorithm_default_detector.cu b/device/cuda/src/fitting/kalman_fitting_algorithm_default_detector.cu
index 8cf34790a8..6ec3381a09 100644
--- a/device/cuda/src/fitting/kalman_fitting_algorithm_default_detector.cu
+++ b/device/cuda/src/fitting/kalman_fitting_algorithm_default_detector.cu
@@ -19,11 +19,13 @@ kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()(
 
     // Run the track fitting.
     if (field.is<const_bfield_backend_t<scalar>>()) {
-        return details::kalman_fitting<default_detector::device>(
+        return details::kalman_fitting_from_candidates<
+            default_detector::device>(
             det, field.as<const_bfield_backend_t<scalar>>(), track_candidates,
             m_config, m_mr, m_copy.get(), m_stream, m_warp_size);
     } else if (field.is<cuda::inhom_bfield_backend_t<scalar>>()) {
-        return details::kalman_fitting<default_detector::device>(
+        return details::kalman_fitting_from_candidates<
+            default_detector::device>(
             det, field.as<cuda::inhom_bfield_backend_t<scalar>>(),
             track_candidates, m_config, m_mr, m_copy.get(), m_stream,
             m_warp_size);
@@ -34,4 +36,26 @@ kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()(
     }
 }
 
+kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()(
+    const default_detector::view& det, const bfield& field,
+    track_state_container_types::buffer&& track_states) const {
+
+    // Run the track fitting.
+    if (field.is<const_bfield_backend_t<scalar>>()) {
+        return details::kalman_fitting_from_states<default_detector::device>(
+            det, field.as<const_bfield_backend_t<scalar>>(),
+            std::move(track_states), m_config, m_mr, m_copy.get(), m_stream,
+            m_warp_size, true);
+    } else if (field.is<cuda::inhom_bfield_backend_t<scalar>>()) {
+        return details::kalman_fitting_from_states<default_detector::device>(
+            det, field.as<cuda::inhom_bfield_backend_t<scalar>>(),
+            std::move(track_states), m_config, m_mr, m_copy.get(), m_stream,
+            m_warp_size, true);
+    } else {
+        throw std::invalid_argument(
+            "Unsupported b-field type received in "
+            "traccc::cuda::kalman_fitting_algorithm");
+    }
+}
+
 }  // namespace traccc::cuda
diff --git a/device/cuda/src/fitting/kalman_fitting_algorithm_telescope_detector.cu b/device/cuda/src/fitting/kalman_fitting_algorithm_telescope_detector.cu
index 9c124188db..d316496bb0 100644
--- a/device/cuda/src/fitting/kalman_fitting_algorithm_telescope_detector.cu
+++ b/device/cuda/src/fitting/kalman_fitting_algorithm_telescope_detector.cu
@@ -19,11 +19,13 @@ kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()(
 
     // Run the track fitting.
     if (field.is<const_bfield_backend_t<scalar>>()) {
-        return details::kalman_fitting<telescope_detector::device>(
+        return details::kalman_fitting_from_candidates<
+            telescope_detector::device>(
             det, field.as<const_bfield_backend_t<scalar>>(), track_candidates,
             m_config, m_mr, m_copy.get(), m_stream, m_warp_size);
     } else if (field.is<cuda::inhom_bfield_backend_t<scalar>>()) {
-        return details::kalman_fitting<telescope_detector::device>(
+        return details::kalman_fitting_from_candidates<
+            telescope_detector::device>(
             det, field.as<cuda::inhom_bfield_backend_t<scalar>>(),
             track_candidates, m_config, m_mr, m_copy.get(), m_stream,
             m_warp_size);
@@ -34,4 +36,26 @@ kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()(
     }
 }
 
+kalman_fitting_algorithm::output_type kalman_fitting_algorithm::operator()(
+    const telescope_detector::view& det, const bfield& field,
+    track_state_container_types::buffer&& track_states) const {
+
+    // Run the track fitting.
+    if (field.is<const_bfield_backend_t<scalar>>()) {
+        return details::kalman_fitting_from_states<telescope_detector::device>(
+            det, field.as<const_bfield_backend_t<scalar>>(),
+            std::move(track_states), m_config, m_mr, m_copy.get(), m_stream,
+            m_warp_size, true);
+    } else if (field.is<cuda::inhom_bfield_backend_t<scalar>>()) {
+        return details::kalman_fitting_from_states<telescope_detector::device>(
+            det, field.as<cuda::inhom_bfield_backend_t<scalar>>(),
+            std::move(track_states), m_config, m_mr, m_copy.get(), m_stream,
+            m_warp_size, true);
+    } else {
+        throw std::invalid_argument(
+            "Unsupported b-field type received in "
+            "traccc::cuda::kalman_fitting_algorithm");
+    }
+}
+
 }  // namespace traccc::cuda
diff --git a/device/cuda/src/fitting/kernels/fill_fitting_state_sort_keys.cu b/device/cuda/src/fitting/kernels/fill_fitting_state_sort_keys.cu
new file mode 100644
index 0000000000..8a264d79dd
--- /dev/null
+++ b/device/cuda/src/fitting/kernels/fill_fitting_state_sort_keys.cu
@@ -0,0 +1,41 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2022-2025 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+// Local include(s).
+#include "../../utils/cuda_error_handling.hpp"
+#include "../../utils/global_index.hpp"
+#include "fill_fitting_state_sort_keys.hpp"
+
+// Project include(s).
+#include "traccc/fitting/device/fill_fitting_state_sort_keys.hpp"
+
+namespace traccc::cuda {
+namespace kernels {
+
+__global__ void fill_fitting_state_sort_keys(
+    track_state_container_types::const_view track_states_view,
+    vecmem::data::vector_view<device::sort_key> keys_view,
+    vecmem::data::vector_view<unsigned int> ids_view) {
+
+    device::fill_fitting_state_sort_keys(
+        details::global_index1(), track_states_view, keys_view, ids_view);
+}
+
+}  // namespace kernels
+
+void fill_fitting_state_sort_keys(
+    const dim3& grid_size, const dim3& block_size, cudaStream_t stream,
+    track_state_container_types::const_view track_states_view,
+    vecmem::data::vector_view<device::sort_key> keys_view,
+    vecmem::data::vector_view<unsigned int> ids_view) {
+
+    kernels::fill_fitting_state_sort_keys<<<grid_size, block_size, 0, stream>>>(
+        track_states_view, keys_view, ids_view);
+    TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
+}
+
+}  // namespace traccc::cuda
diff --git a/device/cuda/src/fitting/kernels/fill_fitting_state_sort_keys.hpp b/device/cuda/src/fitting/kernels/fill_fitting_state_sort_keys.hpp
new file mode 100644
index 0000000000..db31654c01
--- /dev/null
+++ b/device/cuda/src/fitting/kernels/fill_fitting_state_sort_keys.hpp
@@ -0,0 +1,29 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2022-2025 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+// Project include(s).
+#include "traccc/edm/device/sort_key.hpp"
+#include "traccc/edm/track_state.hpp"
+
+// CUDA include(s).
+#include <cuda_runtime.h>
+
+// VecMem include(s).
+#include <vecmem/containers/data/vector_view.hpp>
+
+namespace traccc::cuda {
+
+/// Function calling a kernel for @c traccc::device::fill_fitting_sort_keys
+void fill_fitting_state_sort_keys(
+    const dim3& grid_size, const dim3& block_size, cudaStream_t stream,
+    track_state_container_types::const_view track_states_view,
+    vecmem::data::vector_view<device::sort_key> keys_view,
+    vecmem::data::vector_view<unsigned int> ids_view);
+
+}  // namespace traccc::cuda
diff --git a/examples/run/cuda/full_chain_algorithm.cpp b/examples/run/cuda/full_chain_algorithm.cpp
index 2689b01df3..f6af42688c 100644
--- a/examples/run/cuda/full_chain_algorithm.cpp
+++ b/examples/run/cuda/full_chain_algorithm.cpp
@@ -177,12 +177,13 @@ full_chain_algorithm::output_type full_chain_algorithm::operator()(
                                          m_seeding(spacepoints), m_field_vec);
 
         // Run the track finding (asynchronously).
-        const finding_algorithm::output_type track_candidates = m_finding(
-            m_device_detector_view, m_field, measurements, track_params);
+        finding_algorithm::fitted_output_type ckf_track_states =
+            m_finding(m_device_detector_view, m_field, measurements,
+                      track_params, traccc::device::finding_return_fitted{});
 
         // Run the track fitting (asynchronously).
-        const fitting_algorithm::output_type track_states = m_fitting(
-            m_device_detector_view, m_field, {track_candidates, measurements});
+        const auto track_states = m_fitting(m_device_detector_view, m_field,
+                                            std::move(ckf_track_states));
 
         // Copy a limited amount of result data back to the host.
         output_type result{&m_host_mr};
diff --git a/tests/cpu/test_ckf_combinatorics_telescope.cpp b/tests/cpu/test_ckf_combinatorics_telescope.cpp
index 2d907d3100..4faa485cb9 100644
--- a/tests/cpu/test_ckf_combinatorics_telescope.cpp
+++ b/tests/cpu/test_ckf_combinatorics_telescope.cpp
@@ -119,11 +119,13 @@ TEST_P(CpuCkfCombinatoricsTelescopeTests, Run) {
         std::numeric_limits<unsigned int>::max();
     cfg_no_limit.max_num_branches_per_surface = 10;
     cfg_no_limit.chi2_max = 30.f;
+    cfg_no_limit.duplicate_removal_minimum_length = 100u;
 
     traccc::finding_config cfg_limit;
     cfg_limit.max_num_branches_per_seed = 500;
     cfg_limit.max_num_branches_per_surface = 10;
     cfg_limit.chi2_max = 30.f;
+    cfg_limit.duplicate_removal_minimum_length = 100u;
 
     // Finding algorithm object
     traccc::host::combinatorial_kalman_filter_algorithm host_finding(