diff --git a/cmake/traccc-compiler-options-cuda.cmake b/cmake/traccc-compiler-options-cuda.cmake
index 42cf2ade48..eaf4f40bc6 100644
--- a/cmake/traccc-compiler-options-cuda.cmake
+++ b/cmake/traccc-compiler-options-cuda.cmake
@@ -29,7 +29,7 @@ traccc_add_flag( CMAKE_CUDA_FLAGS "--expt-relaxed-constexpr" )
 
 # Make CUDA generate debug symbols for the device code as well in a debug
 # build.
-traccc_add_flag( CMAKE_CUDA_FLAGS_DEBUG "-G --keep -src-in-ptx" )
+traccc_add_flag( CMAKE_CUDA_FLAGS_DEBUG "-G -src-in-ptx" )
 
 # Ensure that line information is embedded in debugging builds so that
 # profilers have access to line data.
diff --git a/core/include/traccc/finding/ckf_aborter.hpp b/core/include/traccc/finding/ckf_aborter.hpp
index 26dfedecb9..02eba824a0 100644
--- a/core/include/traccc/finding/ckf_aborter.hpp
+++ b/core/include/traccc/finding/ckf_aborter.hpp
@@ -11,6 +11,7 @@
 #include "detray/definitions/detail/qualifiers.hpp"
 #include "detray/propagator/base_actor.hpp"
 #include "detray/propagator/base_stepper.hpp"
+#include "traccc/definitions/primitives.hpp"
 
 // System include(s)
 #include <limits>
@@ -51,4 +52,4 @@ struct ckf_aborter : detray::actor {
     }
 };
 
-}  // namespace traccc
\ No newline at end of file
+}  // namespace traccc
diff --git a/device/common/include/traccc/finding/device/apply_interaction.hpp b/device/common/include/traccc/finding/device/apply_interaction.hpp
index dcd03f7df9..6212a28a3e 100644
--- a/device/common/include/traccc/finding/device/apply_interaction.hpp
+++ b/device/common/include/traccc/finding/device/apply_interaction.hpp
@@ -8,29 +8,31 @@
 #pragma once
 
 // Project include(s).
+#include "detray/navigation/navigator.hpp"
+#include "detray/propagator/actors/pointwise_material_interactor.hpp"
 #include "traccc/definitions/qualifiers.hpp"
 #include "traccc/finding/finding_config.hpp"
+#include "traccc/utils/particle.hpp"
 
 namespace traccc::device {
+template <typename detector_t>
+struct apply_interaction_payload {
+    typename detector_t::view_type det_data;
+    const int n_params;
+    bound_track_parameters_collection_types::view params_view;
+    vecmem::data::vector_view<const unsigned int> params_liveness_view;
+};
 
 /// Function applying the Pre material interaction to tracks spawned by bound
 /// track parameters
 ///
 /// @param[in] globalIndex     The index of the current thread
 /// @param[in] cfg             Track finding config object
-/// @param[in] det_data        Detector view object
-/// @param[in] n_params        The number of parameters (or tracks)
-/// @param[out] params_view    Collection of output bound track_parameters
-/// @param[in] params_liveness_view Vector of parameter liveness indicators
-///
+/// @param[inout] payload      The function call payload
 template <typename detector_t>
 TRACCC_DEVICE inline void apply_interaction(
     std::size_t globalIndex, const finding_config& cfg,
-    typename detector_t::view_type det_data, const int n_params,
-    bound_track_parameters_collection_types::view params_view,
-    vecmem::data::vector_view<const unsigned int> params_liveness_view);
-
+    const apply_interaction_payload<detector_t>& payload);
 }  // namespace traccc::device
 
-// Include the implementation.
-#include "traccc/finding/device/impl/apply_interaction.ipp"
+#include "./impl/apply_interaction.ipp"
diff --git a/device/common/include/traccc/finding/device/build_tracks.hpp b/device/common/include/traccc/finding/device/build_tracks.hpp
index 5719d39876..fb1b67ba8b 100644
--- a/device/common/include/traccc/finding/device/build_tracks.hpp
+++ b/device/common/include/traccc/finding/device/build_tracks.hpp
@@ -9,8 +9,22 @@
 
 // Project include(s).
 #include "traccc/definitions/qualifiers.hpp"
+#include "traccc/edm/measurement.hpp"
+#include "traccc/edm/track_candidate.hpp"
+#include "traccc/edm/track_parameters.hpp"
+#include "traccc/finding/candidate_link.hpp"
 
 namespace traccc::device {
+struct build_tracks_payload {
+    measurement_collection_types::const_view measurements_view;
+    bound_track_parameters_collection_types::const_view seeds_view;
+    vecmem::data::jagged_vector_view<const candidate_link> links_view;
+    vecmem::data::vector_view<const typename candidate_link::link_index_type>
+        tips_view;
+    track_candidate_container_types::view track_candidates_view;
+    vecmem::data::vector_view<unsigned int> valid_indices_view;
+    unsigned int& n_valid_tracks;
+};
 
 /// Function for building full tracks from the link container:
 /// The full tracks are built using the link container and tip link container.
@@ -19,28 +33,12 @@ namespace traccc::device {
 ///
 /// @param[in] globalIndex         The index of the current thread
 /// @param[in] cfg                    Track finding config object
-/// @param[in] measurements_view   Measurements container view
-/// @param[in] seeds_view          Seed container view
-/// @param[in] link_view           Link container view
-/// @param[in] param_to_link_view  Container for param index -> link index
-/// @param[in] tips_view           Tip link container view
-/// @param[out] track_candidates_view  Track candidate container view
-/// @param[out] valid_indices_view   Valid indices meeting criteria
-/// @param[out] n_valid_tracks       The number of valid tracks meeting criteria
-
+/// @param[inout] payload      The function call payload
 template <typename config_t>
-TRACCC_DEVICE inline void build_tracks(
-    std::size_t globalIndex, const config_t cfg,
-    measurement_collection_types::const_view measurements_view,
-    bound_track_parameters_collection_types::const_view seeds_view,
-    vecmem::data::jagged_vector_view<const candidate_link> links_view,
-    vecmem::data::vector_view<const typename candidate_link::link_index_type>
-        tips_view,
-    track_candidate_container_types::view track_candidates_view,
-    vecmem::data::vector_view<unsigned int> valid_indices_view,
-    unsigned int& n_valid_tracks);
+TRACCC_DEVICE inline void build_tracks(std::size_t globalIndex,
+                                       const config_t cfg,
+                                       const build_tracks_payload& payload);
 
 }  // namespace traccc::device
 
-// Include the implementation.
-#include "traccc/finding/device/impl/build_tracks.ipp"
+#include "./impl/build_tracks.ipp"
diff --git a/device/common/include/traccc/finding/device/fill_sort_keys.hpp b/device/common/include/traccc/finding/device/fill_sort_keys.hpp
index 53dcfaf7f5..deeb809919 100644
--- a/device/common/include/traccc/finding/device/fill_sort_keys.hpp
+++ b/device/common/include/traccc/finding/device/fill_sort_keys.hpp
@@ -12,21 +12,18 @@
 #include "traccc/edm/track_candidate.hpp"
 
 namespace traccc::device {
+struct fill_sort_keys_payload {
+    bound_track_parameters_collection_types::const_view params_view;
+    vecmem::data::vector_view<device::sort_key> keys_view;
+    vecmem::data::vector_view<unsigned int> ids_view;
+};
 
 /// Function used for fill key container
 ///
 /// @param[in] globalIndex   The index of the current thread
-/// @param[in] params_view   The input parameters
-/// @param[out] keys_view    The key values
-/// @param[out] ids_view     The param ids
-///
+/// @param[inout] payload      The function call payload
 TRACCC_HOST_DEVICE inline void fill_sort_keys(
-    std::size_t globalIndex,
-    bound_track_parameters_collection_types::const_view params_view,
-    vecmem::data::vector_view<device::sort_key> keys_view,
-    vecmem::data::vector_view<unsigned int> ids_view);
-
+    std::size_t globalIndex, const fill_sort_keys_payload& payload);
 }  // namespace traccc::device
 
-// Include the implementation.
-#include "traccc/finding/device/impl/fill_sort_keys.ipp"
+#include "./impl/fill_sort_keys.ipp"
diff --git a/device/common/include/traccc/finding/device/find_tracks.hpp b/device/common/include/traccc/finding/device/find_tracks.hpp
index 569fec29f3..e5943958b7 100644
--- a/device/common/include/traccc/finding/device/find_tracks.hpp
+++ b/device/common/include/traccc/finding/device/find_tracks.hpp
@@ -14,11 +14,38 @@
 #include "traccc/device/concepts/thread_id.hpp"
 #include "traccc/edm/measurement.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"
+#include "traccc/fitting/kalman_filter/gain_matrix_updater.hpp"
 
 // Thrust include(s)
 #include <thrust/binary_search.h>
 
 namespace traccc::device {
+template <typename detector_t>
+struct find_tracks_payload {
+    typename detector_t::view_type det_data;
+    measurement_collection_types::const_view measurements_view;
+    bound_track_parameters_collection_types::const_view in_params_view;
+    vecmem::data::vector_view<const unsigned int> in_params_liveness_view;
+    const unsigned int n_in_params;
+    vecmem::data::vector_view<const detray::geometry::barcode> barcodes_view;
+    vecmem::data::vector_view<const unsigned int> upper_bounds_view;
+    vecmem::data::vector_view<const candidate_link> prev_links_view;
+    const unsigned int step;
+    const unsigned int n_max_candidates;
+    bound_track_parameters_collection_types::view out_params_view;
+    vecmem::data::vector_view<unsigned int> out_params_liveness_view;
+    vecmem::data::vector_view<candidate_link> links_view;
+    unsigned int* n_total_candidates;
+};
+
+struct find_tracks_shared_payload {
+    unsigned int* shared_num_candidates;
+    std::pair<unsigned int, unsigned int>* shared_candidates;
+    unsigned int& shared_candidates_size;
+};
 
 /// Function for combinatorial finding.
 /// If the chi2 of the measurement < chi2_max, its measurement index and the
@@ -27,47 +54,14 @@ namespace traccc::device {
 /// @param[in] thread_id          A thread identifier object
 /// @param[in] barrier            A block-wide barrier
 /// @param[in] cfg                Track finding config object
-/// @param[in] det_data           Detector view object
-/// @param[in] measurements_view  Measurements container view
-/// @param[in] in_params_view     Input parameters
-/// @param[in] n_in_params        The number of input params
-/// @param[in] barcodes_view      View of a measurement -> barcode map
-/// @param[in] upper_bounds_view  Upper bounds of measurements unique w.r.t
-/// barcode
-/// @param[in] prev_links_view    link container from the previous step
-/// @param[in] prev_param_to_link_view  param_to_link container from the
-/// previous step
-/// @param[in] step               Step index
-/// @param[in] n_max_candidates   Number of maximum candidates
-/// @param[out] out_params_view   Output parameters
-/// @param[out] links_view        link container for the current step
-/// @param[out] n_total_candidates  The number of total candidates for the
-/// current step
-/// @param shared_num_candidates  Shared memory scratch space
-/// @param shared_candidates      Shared memory scratch space
-/// @param shared_candidates_size Shared memory scratch space
-///
+/// @param[inout] payload         The global memory payload
+/// @param[inout] shared_payload  The shared memory payload
 template <concepts::thread_id1 thread_id_t, concepts::barrier barrier_t,
           typename detector_t, typename config_t>
 TRACCC_DEVICE inline void find_tracks(
     thread_id_t& thread_id, barrier_t& barrier, const config_t cfg,
-    typename detector_t::view_type det_data,
-    measurement_collection_types::const_view measurements_view,
-    bound_track_parameters_collection_types::const_view in_params_view,
-    vecmem::data::vector_view<const unsigned int> in_params_liveness_view,
-    const unsigned int n_in_params,
-    vecmem::data::vector_view<const detray::geometry::barcode> barcodes_view,
-    vecmem::data::vector_view<const unsigned int> upper_bounds_view,
-    vecmem::data::vector_view<const candidate_link> prev_links_view,
-    const unsigned int step, const unsigned int& n_max_candidates,
-    bound_track_parameters_collection_types::view out_params_view,
-    vecmem::data::vector_view<unsigned int> out_params_liveness_view,
-    vecmem::data::vector_view<candidate_link> links_view,
-    unsigned int& n_total_candidates, unsigned int* shared_num_candidates,
-    std::pair<unsigned int, unsigned int>* shared_candidates,
-    unsigned int& shared_candidates_size);
-
+    const find_tracks_payload<detector_t>& payload,
+    const find_tracks_shared_payload& shared_payload);
 }  // namespace traccc::device
 
-// Include the implementation.
-#include "traccc/finding/device/impl/find_tracks.ipp"
+#include "./impl/find_tracks.ipp"
diff --git a/device/common/include/traccc/finding/device/impl/apply_interaction.ipp b/device/common/include/traccc/finding/device/impl/apply_interaction.ipp
index c74e12c1c0..bfbb587cc7 100644
--- a/device/common/include/traccc/finding/device/impl/apply_interaction.ipp
+++ b/device/common/include/traccc/finding/device/impl/apply_interaction.ipp
@@ -1,6 +1,6 @@
 /** TRACCC library, part of the ACTS project (R&D line)
  *
- * (c) 2023-2024 CERN for the benefit of the ACTS project
+ * (c) 2023 CERN for the benefit of the ACTS project
  *
  * Mozilla Public License Version 2.0
  */
@@ -8,35 +8,32 @@
 #pragma once
 
 // Project include(s).
-#include "traccc/definitions/math.hpp"
+#include "detray/navigation/navigator.hpp"
+#include "detray/propagator/actors/pointwise_material_interactor.hpp"
+#include "traccc/definitions/qualifiers.hpp"
+#include "traccc/finding/finding_config.hpp"
 #include "traccc/utils/particle.hpp"
 
-// Detray include(s).
-#include "detray/geometry/tracking_surface.hpp"
-#include "vecmem/containers/device_vector.hpp"
-
 namespace traccc::device {
 
 template <typename detector_t>
 TRACCC_DEVICE inline void apply_interaction(
     std::size_t globalIndex, const finding_config& cfg,
-    typename detector_t::view_type det_data, const int n_params,
-    bound_track_parameters_collection_types::view params_view,
-    vecmem::data::vector_view<const unsigned int> params_liveness_view) {
+    const apply_interaction_payload<detector_t>& payload) {
 
     // Type definitions
     using algebra_type = typename detector_t::algebra_type;
     using interactor_type = detray::pointwise_material_interactor<algebra_type>;
 
     // Detector
-    detector_t det(det_data);
+    detector_t det(payload.det_data);
 
     // in param
-    bound_track_parameters_collection_types::device params(params_view);
+    bound_track_parameters_collection_types::device params(payload.params_view);
     vecmem::device_vector<const unsigned int> params_liveness(
-        params_liveness_view);
+        payload.params_liveness_view);
 
-    if (globalIndex >= n_params) {
+    if (globalIndex >= payload.n_params) {
         return;
     }
 
@@ -57,5 +54,4 @@ TRACCC_DEVICE inline void apply_interaction(
             static_cast<int>(detray::navigation::direction::e_forward), sf);
     }
 }
-
 }  // 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_tracks.ipp
index ac8ec0b9f8..589cd7e8f0 100644
--- a/device/common/include/traccc/finding/device/impl/build_tracks.ipp
+++ b/device/common/include/traccc/finding/device/impl/build_tracks.ipp
@@ -7,33 +7,37 @@
 
 #pragma once
 
+// Project include(s).
+#include "traccc/definitions/qualifiers.hpp"
+#include "traccc/edm/measurement.hpp"
+#include "traccc/edm/track_candidate.hpp"
+#include "traccc/edm/track_parameters.hpp"
+#include "traccc/finding/candidate_link.hpp"
+
 namespace traccc::device {
 
 template <typename config_t>
-TRACCC_DEVICE inline void build_tracks(
-    std::size_t globalIndex, const config_t cfg,
-    measurement_collection_types::const_view measurements_view,
-    bound_track_parameters_collection_types::const_view seeds_view,
-    vecmem::data::jagged_vector_view<const candidate_link> links_view,
-    vecmem::data::vector_view<const typename candidate_link::link_index_type>
-        tips_view,
-    track_candidate_container_types::view track_candidates_view,
-    vecmem::data::vector_view<unsigned int> valid_indices_view,
-    unsigned int& n_valid_tracks) {
+TRACCC_DEVICE inline void build_tracks(std::size_t globalIndex,
+                                       const config_t cfg,
+                                       const build_tracks_payload& payload) {
 
-    measurement_collection_types::const_device measurements(measurements_view);
+    measurement_collection_types::const_device measurements(
+        payload.measurements_view);
 
-    bound_track_parameters_collection_types::const_device seeds(seeds_view);
+    bound_track_parameters_collection_types::const_device seeds(
+        payload.seeds_view);
 
-    vecmem::jagged_device_vector<const candidate_link> links(links_view);
+    vecmem::jagged_device_vector<const candidate_link> links(
+        payload.links_view);
 
     vecmem::device_vector<const typename candidate_link::link_index_type> tips(
-        tips_view);
+        payload.tips_view);
 
     track_candidate_container_types::device track_candidates(
-        track_candidates_view);
+        payload.track_candidates_view);
 
-    vecmem::device_vector<unsigned int> valid_indices(valid_indices_view);
+    vecmem::device_vector<unsigned int> valid_indices(
+        payload.valid_indices_view);
 
     if (globalIndex >= tips.size()) {
         return;
@@ -107,7 +111,7 @@ TRACCC_DEVICE inline void build_tracks(
         n_cands <= cfg.max_track_candidates_per_track) {
 
         vecmem::device_atomic_ref<unsigned int> num_valid_tracks(
-            n_valid_tracks);
+            payload.n_valid_tracks);
 
         const unsigned int pos = num_valid_tracks.fetch_add(1);
         valid_indices[pos] = globalIndex;
diff --git a/device/common/include/traccc/finding/device/impl/fill_sort_keys.ipp b/device/common/include/traccc/finding/device/impl/fill_sort_keys.ipp
index f95b42fbf7..e4d09fef1b 100644
--- a/device/common/include/traccc/finding/device/impl/fill_sort_keys.ipp
+++ b/device/common/include/traccc/finding/device/impl/fill_sort_keys.ipp
@@ -7,21 +7,23 @@
 
 #pragma once
 
+// Project include(s).
+#include "traccc/edm/device/sort_key.hpp"
+#include "traccc/edm/track_candidate.hpp"
+
 namespace traccc::device {
 
 TRACCC_HOST_DEVICE inline void fill_sort_keys(
-    std::size_t globalIndex,
-    bound_track_parameters_collection_types::const_view params_view,
-    vecmem::data::vector_view<device::sort_key> keys_view,
-    vecmem::data::vector_view<unsigned int> ids_view) {
+    std::size_t globalIndex, const fill_sort_keys_payload& payload) {
 
-    bound_track_parameters_collection_types::const_device params(params_view);
+    bound_track_parameters_collection_types::const_device params(
+        payload.params_view);
 
     // Keys
-    vecmem::device_vector<device::sort_key> keys_device(keys_view);
+    vecmem::device_vector<device::sort_key> keys_device(payload.keys_view);
 
     // Param id
-    vecmem::device_vector<unsigned int> ids_device(ids_view);
+    vecmem::device_vector<unsigned int> ids_device(payload.ids_view);
 
     if (globalIndex >= keys_device.size()) {
         return;
@@ -31,4 +33,4 @@ TRACCC_HOST_DEVICE inline void fill_sort_keys(
     ids_device.at(globalIndex) = globalIndex;
 }
 
-}  // namespace traccc::device
\ No newline at end of file
+}  // namespace traccc::device
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 921f17bf1a..5fb1cd87c0 100644
--- a/device/common/include/traccc/finding/device/impl/find_tracks.ipp
+++ b/device/common/include/traccc/finding/device/impl/find_tracks.ipp
@@ -8,14 +8,19 @@
 #pragma once
 
 // Project include(s).
+#include "traccc/definitions/primitives.hpp"
+#include "traccc/definitions/qualifiers.hpp"
 #include "traccc/device/concepts/barrier.hpp"
 #include "traccc/device/concepts/thread_id.hpp"
+#include "traccc/edm/measurement.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"
 #include "traccc/fitting/kalman_filter/gain_matrix_updater.hpp"
-#include "vecmem/containers/device_vector.hpp"
 
-// System include(s).
-#include <limits>
+// Thrust include(s)
+#include <thrust/binary_search.h>
 
 namespace traccc::device {
 
@@ -23,21 +28,8 @@ template <concepts::thread_id1 thread_id_t, concepts::barrier barrier_t,
           typename detector_t, typename config_t>
 TRACCC_DEVICE inline void find_tracks(
     thread_id_t& thread_id, barrier_t& barrier, const config_t cfg,
-    typename detector_t::view_type det_data,
-    measurement_collection_types::const_view measurements_view,
-    bound_track_parameters_collection_types::const_view in_params_view,
-    vecmem::data::vector_view<const unsigned int> in_params_liveness_view,
-    const unsigned int n_in_params,
-    vecmem::data::vector_view<const detray::geometry::barcode> barcodes_view,
-    vecmem::data::vector_view<const unsigned int> upper_bounds_view,
-    vecmem::data::vector_view<const candidate_link> prev_links_view,
-    const unsigned int step, const unsigned int& n_max_candidates,
-    bound_track_parameters_collection_types::view out_params_view,
-    vecmem::data::vector_view<unsigned int> out_params_liveness_view,
-    vecmem::data::vector_view<candidate_link> links_view,
-    unsigned int& n_total_candidates, unsigned int* shared_num_candidates,
-    std::pair<unsigned int, unsigned int>* shared_candidates,
-    unsigned int& shared_candidates_size) {
+    const find_tracks_payload<detector_t>& payload,
+    const find_tracks_shared_payload& shared_payload) {
 
     /*
      * Initialize the block-shared data; in particular, set the total size of
@@ -45,41 +37,46 @@ TRACCC_DEVICE inline void find_tracks(
      * each parameter to zero.
      */
     if (thread_id.getLocalThreadIdX() == 0) {
-        shared_candidates_size = 0;
+        shared_payload.shared_candidates_size = 0;
     }
 
-    shared_num_candidates[thread_id.getLocalThreadIdX()] = 0;
+    shared_payload.shared_num_candidates[thread_id.getLocalThreadIdX()] = 0;
 
     barrier.blockBarrier();
 
     /*
      * Initialize all of the device vectors from their vecmem views.
      */
-    detector_t det(det_data);
-    measurement_collection_types::const_device measurements(measurements_view);
+    detector_t det(payload.det_data);
+    measurement_collection_types::const_device measurements(
+        payload.measurements_view);
     bound_track_parameters_collection_types::const_device in_params(
-        in_params_view);
+        payload.in_params_view);
     vecmem::device_vector<const unsigned int> in_params_liveness(
-        in_params_liveness_view);
-    vecmem::device_vector<const candidate_link> prev_links(prev_links_view);
-    bound_track_parameters_collection_types::device out_params(out_params_view);
+        payload.in_params_liveness_view);
+    vecmem::device_vector<const candidate_link> prev_links(
+        payload.prev_links_view);
+    bound_track_parameters_collection_types::device out_params(
+        payload.out_params_view);
     vecmem::device_vector<unsigned int> out_params_liveness(
-        out_params_liveness_view);
-    vecmem::device_vector<candidate_link> links(links_view);
+        payload.out_params_liveness_view);
+    vecmem::device_vector<candidate_link> links(payload.links_view);
     vecmem::device_atomic_ref<unsigned int> num_total_candidates(
-        n_total_candidates);
+        *payload.n_total_candidates);
     vecmem::device_vector<const detray::geometry::barcode> barcodes(
-        barcodes_view);
-    vecmem::device_vector<const unsigned int> upper_bounds(upper_bounds_view);
+        payload.barcodes_view);
+    vecmem::device_vector<const unsigned int> upper_bounds(
+        payload.upper_bounds_view);
 
     /*
      * Compute the last step ID, using a sentinel value if the current step is
      * step 0.
      */
     const candidate_link::link_index_type::first_type previous_step =
-        (step == 0) ? std::numeric_limits<
-                          candidate_link::link_index_type::first_type>::max()
-                    : step - 1;
+        (payload.step == 0)
+            ? std::numeric_limits<
+                  candidate_link::link_index_type::first_type>::max()
+            : payload.step - 1;
 
     const unsigned int in_param_id = thread_id.getGlobalThreadIdX();
 
@@ -94,7 +91,8 @@ TRACCC_DEVICE inline void find_tracks(
     unsigned int init_meas = 0;
     unsigned int num_meas = 0;
 
-    if (in_param_id < n_in_params && in_params_liveness.at(in_param_id) > 0u) {
+    if (in_param_id < payload.n_in_params &&
+        in_params_liveness.at(in_param_id) > 0u) {
         /*
          * Get the barcode of this thread's parameters, then find the first
          * measurement that matches it.
@@ -142,8 +140,8 @@ TRACCC_DEVICE inline void find_tracks(
      * This loop keeps running until all threads have processed all of their
      * measurements.
      */
-    while (
-        barrier.blockOr(curr_meas < num_meas || shared_candidates_size > 0)) {
+    while (barrier.blockOr(curr_meas < num_meas ||
+                           shared_payload.shared_candidates_size > 0)) {
         /*
          * The outer loop consists of three general components. The first
          * components is that each thread starts to fill a shared buffer of
@@ -154,19 +152,19 @@ TRACCC_DEVICE inline void find_tracks(
          * either run out of measurements, or until the shared buffer is full.
          */
         for (; curr_meas < num_meas &&
-               shared_candidates_size < thread_id.getBlockDimX();
+               shared_payload.shared_candidates_size < thread_id.getBlockDimX();
              curr_meas++) {
-            unsigned int idx =
-                vecmem::device_atomic_ref<unsigned int>(shared_candidates_size)
-                    .fetch_add(1u);
+            unsigned int idx = vecmem::device_atomic_ref<unsigned int>(
+                                   shared_payload.shared_candidates_size)
+                                   .fetch_add(1u);
 
             /*
              * The buffer elemements are tuples of the measurement index and
              * the index of the thread that originally inserted that
              * measurement.
              */
-            shared_candidates[idx] = {init_meas + curr_meas,
-                                      thread_id.getLocalThreadIdX()};
+            shared_payload.shared_candidates[idx] = {
+                init_meas + curr_meas, thread_id.getLocalThreadIdX()};
         }
 
         barrier.blockBarrier();
@@ -175,9 +173,11 @@ TRACCC_DEVICE inline void find_tracks(
          * The shared buffer is now full; each thread picks out zero or one of
          * the measurements and processes it.
          */
-        if (thread_id.getLocalThreadIdX() < shared_candidates_size) {
+        if (thread_id.getLocalThreadIdX() <
+            shared_payload.shared_candidates_size) {
             const unsigned int owner_local_thread_id =
-                shared_candidates[thread_id.getLocalThreadIdX()].second;
+                shared_payload.shared_candidates[thread_id.getLocalThreadIdX()]
+                    .second;
             const unsigned int owner_global_thread_id =
                 owner_local_thread_id +
                 thread_id.getBlockDimX() * thread_id.getBlockIdX();
@@ -185,7 +185,8 @@ TRACCC_DEVICE inline void find_tracks(
             bound_track_parameters in_par =
                 in_params.at(owner_global_thread_id);
             const unsigned int meas_idx =
-                shared_candidates[thread_id.getLocalThreadIdX()].first;
+                shared_payload.shared_candidates[thread_id.getLocalThreadIdX()]
+                    .first;
 
             const auto& meas = measurements.at(meas_idx);
 
@@ -203,10 +204,10 @@ TRACCC_DEVICE inline void find_tracks(
                 // Add measurement candidates to link
                 const unsigned int l_pos = num_total_candidates.fetch_add(1);
 
-                if (l_pos >= n_max_candidates) {
-                    n_total_candidates = n_max_candidates;
+                if (l_pos >= payload.n_max_candidates) {
+                    *payload.n_total_candidates = payload.n_max_candidates;
                 } else {
-                    if (step == 0) {
+                    if (payload.step == 0) {
                         links.at(l_pos) = {
                             {previous_step, owner_global_thread_id},
                             meas_idx,
@@ -226,7 +227,8 @@ TRACCC_DEVICE inline void find_tracks(
                     // Increase the number of candidates (or branches) per input
                     // parameter
                     vecmem::device_atomic_ref<unsigned int>(
-                        shared_num_candidates[owner_local_thread_id])
+                        shared_payload
+                            .shared_num_candidates[owner_local_thread_id])
                         .fetch_add(1u);
 
                     out_params.at(l_pos) = trk_state.filtered();
@@ -242,15 +244,17 @@ TRACCC_DEVICE inline void find_tracks(
          * might end up having some spill-over; this spill-over should be moved
          * to the front of the buffer.
          */
-        shared_candidates[thread_id.getLocalThreadIdX()] =
-            shared_candidates[thread_id.getLocalThreadIdX() +
-                              thread_id.getBlockDimX()];
+        shared_payload.shared_candidates[thread_id.getLocalThreadIdX()] =
+            shared_payload.shared_candidates[thread_id.getLocalThreadIdX() +
+                                             thread_id.getBlockDimX()];
 
         if (thread_id.getLocalThreadIdX() == 0) {
-            if (shared_candidates_size >= thread_id.getBlockDimX()) {
-                shared_candidates_size -= thread_id.getBlockDimX();
+            if (shared_payload.shared_candidates_size >=
+                thread_id.getBlockDimX()) {
+                shared_payload.shared_candidates_size -=
+                    thread_id.getBlockDimX();
             } else {
-                shared_candidates_size = 0;
+                shared_payload.shared_candidates_size = 0;
             }
         }
     }
@@ -259,15 +263,17 @@ TRACCC_DEVICE inline void find_tracks(
      * Part three of the kernel inserts holes for parameters which did not
      * match any measurements.
      */
-    if (in_param_id < n_in_params && in_params_liveness.at(in_param_id) > 0u &&
-        shared_num_candidates[thread_id.getLocalThreadIdX()] == 0u) {
+    if (in_param_id < payload.n_in_params &&
+        in_params_liveness.at(in_param_id) > 0u &&
+        shared_payload.shared_num_candidates[thread_id.getLocalThreadIdX()] ==
+            0u) {
         // Add measurement candidates to link
         const unsigned int l_pos = num_total_candidates.fetch_add(1);
 
-        if (l_pos >= n_max_candidates) {
-            n_total_candidates = n_max_candidates;
+        if (l_pos >= payload.n_max_candidates) {
+            *payload.n_total_candidates = payload.n_max_candidates;
         } else {
-            if (step == 0) {
+            if (payload.step == 0) {
                 links.at(l_pos) = {{previous_step, in_param_id},
                                    std::numeric_limits<unsigned int>::max(),
                                    in_param_id,
diff --git a/device/common/include/traccc/finding/device/impl/make_barcode_sequence.ipp b/device/common/include/traccc/finding/device/impl/make_barcode_sequence.ipp
index 0fe6b5d27e..a70b6f6840 100644
--- a/device/common/include/traccc/finding/device/impl/make_barcode_sequence.ipp
+++ b/device/common/include/traccc/finding/device/impl/make_barcode_sequence.ipp
@@ -1,24 +1,24 @@
 /** TRACCC library, part of the ACTS project (R&D line)
  *
- * (c) 2023-2024 CERN for the benefit of the ACTS project
+ * (c) 2023 CERN for the benefit of the ACTS project
  *
  * Mozilla Public License Version 2.0
  */
 
 #pragma once
 
-// System include(s).
-#include <cassert>
+// Project include(s).
+#include "traccc/definitions/primitives.hpp"
+#include "traccc/definitions/qualifiers.hpp"
+#include "traccc/edm/measurement.hpp"
 
 namespace traccc::device {
 
 TRACCC_DEVICE inline void make_barcode_sequence(
-    std::size_t globalIndex,
-    measurement_collection_types::const_view uniques_view,
-    vecmem::data::vector_view<detray::geometry::barcode> barcodes_view) {
+    std::size_t globalIndex, const make_barcode_sequence_payload& payload) {
 
-    measurement_collection_types::const_device uniques(uniques_view);
-    vecmem::device_vector barcodes(barcodes_view);
+    measurement_collection_types::const_device uniques(payload.uniques_view);
+    vecmem::device_vector barcodes(payload.barcodes_view);
     assert(uniques.size() >= barcodes.size());
 
     if (globalIndex >= barcodes.size()) {
@@ -29,4 +29,4 @@ TRACCC_DEVICE inline void make_barcode_sequence(
     barcodes.at(globalIndex) = uniques.at(globalIndex).surface_link;
 }
 
-}  // namespace traccc::device
\ No newline at end of file
+}  // namespace traccc::device
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 2e89041be4..bb518215ce 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
@@ -1,44 +1,45 @@
 /** TRACCC library, part of the ACTS project (R&D line)
  *
- * (c) 2023-2024 CERN for the benefit of the ACTS project
+ * (c) 2023 CERN for the benefit of the ACTS project
  *
  * Mozilla Public License Version 2.0
  */
 
 #pragma once
 
-#include "vecmem/containers/device_vector.hpp"
+// Project include(s).
+#include "detray/core/detail/tuple_container.hpp"
+#include "detray/propagator/constrained_step.hpp"
+#include "detray/utils/tuple.hpp"
+#include "traccc/definitions/primitives.hpp"
+#include "traccc/definitions/qualifiers.hpp"
+#include "traccc/edm/measurement.hpp"
+#include "traccc/edm/track_parameters.hpp"
+#include "traccc/finding/candidate_link.hpp"
+#include "traccc/utils/particle.hpp"
+
 namespace traccc::device {
 
 template <typename propagator_t, typename bfield_t, typename config_t>
 TRACCC_DEVICE inline void propagate_to_next_surface(
     std::size_t globalIndex, const config_t cfg,
-    typename propagator_t::detector_type::view_type det_data,
-    bfield_t field_data,
-    bound_track_parameters_collection_types::view params_view,
-    vecmem::data::vector_view<unsigned int> params_liveness_view,
-    const vecmem::data::vector_view<const unsigned int>& param_ids_view,
-    vecmem::data::vector_view<const candidate_link> links_view,
-    const unsigned int step, const unsigned int n_in_params,
-    vecmem::data::vector_view<typename candidate_link::link_index_type>
-        tips_view,
-    vecmem::data::vector_view<unsigned int> n_tracks_per_seed_view) {
-
-    if (globalIndex >= n_in_params) {
+    const propagate_to_next_surface_payload<propagator_t, bfield_t>& payload) {
+
+    if (globalIndex >= payload.n_in_params) {
         return;
     }
 
     // Theta id
-    vecmem::device_vector<const unsigned int> param_ids(param_ids_view);
+    vecmem::device_vector<const unsigned int> param_ids(payload.param_ids_view);
 
     const unsigned int param_id = param_ids.at(globalIndex);
 
     // Number of tracks per seed
     vecmem::device_vector<unsigned int> n_tracks_per_seed(
-        n_tracks_per_seed_view);
+        payload.n_tracks_per_seed_view);
 
     // Links
-    vecmem::device_vector<const candidate_link> links(links_view);
+    vecmem::device_vector<const candidate_link> links(payload.links_view);
 
     // Seed id
     unsigned int orig_param_id = links.at(param_id).seed_idx;
@@ -48,7 +49,8 @@ TRACCC_DEVICE inline void propagate_to_next_surface(
         n_tracks_per_seed.at(orig_param_id));
 
     const unsigned int s_pos = num_tracks_per_seed.fetch_add(1);
-    vecmem::device_vector<unsigned int> params_liveness(params_liveness_view);
+    vecmem::device_vector<unsigned int> params_liveness(
+        payload.params_liveness_view);
 
     if (s_pos >= cfg.max_num_branches_per_seed) {
         params_liveness[param_id] = 0u;
@@ -57,19 +59,19 @@ TRACCC_DEVICE inline void propagate_to_next_surface(
 
     // tips
     vecmem::device_vector<typename candidate_link::link_index_type> tips(
-        tips_view);
+        payload.tips_view);
 
     if (links.at(param_id).n_skipped > cfg.max_num_skipping_per_cand) {
         params_liveness[param_id] = 0u;
-        tips.push_back({step, param_id});
+        tips.push_back({payload.step, param_id});
         return;
     }
 
     // Detector
-    typename propagator_t::detector_type det(det_data);
+    typename propagator_t::detector_type det(payload.det_data);
 
     // Parameters
-    bound_track_parameters_collection_types::device params(params_view);
+    bound_track_parameters_collection_types::device params(payload.params_view);
 
     if (params_liveness.at(param_id) == 0u) {
         return;
@@ -82,7 +84,7 @@ TRACCC_DEVICE inline void propagate_to_next_surface(
     propagator_t propagator(cfg.propagation);
 
     // Create propagator state
-    typename propagator_t::state propagation(in_par, field_data, det);
+    typename propagator_t::state propagation(in_par, payload.field_data, det);
     propagation.set_particle(
         detail::correct_particle_hypothesis(cfg.ptc_hypothesis, in_par));
     propagation._stepping
@@ -117,8 +119,8 @@ TRACCC_DEVICE inline void propagate_to_next_surface(
     if (s4.success) {
         params[param_id] = propagation._stepping._bound_params;
 
-        if (step == cfg.max_track_candidates_per_track - 1) {
-            tips.push_back({step, param_id});
+        if (payload.step == cfg.max_track_candidates_per_track - 1) {
+            tips.push_back({payload.step, param_id});
             params_liveness[param_id] = 0u;
         } else {
             params_liveness[param_id] = 1u;
@@ -126,8 +128,8 @@ TRACCC_DEVICE inline void propagate_to_next_surface(
     } else {
         params_liveness[param_id] = 0u;
 
-        if (step >= cfg.min_track_candidates_per_track - 1) {
-            tips.push_back({step, param_id});
+        if (payload.step >= cfg.min_track_candidates_per_track - 1) {
+            tips.push_back({payload.step, param_id});
         }
     }
 }
diff --git a/device/common/include/traccc/finding/device/impl/prune_tracks.ipp b/device/common/include/traccc/finding/device/impl/prune_tracks.ipp
index 6005568a62..d9979241f4 100644
--- a/device/common/include/traccc/finding/device/impl/prune_tracks.ipp
+++ b/device/common/include/traccc/finding/device/impl/prune_tracks.ipp
@@ -7,19 +7,22 @@
 
 #pragma once
 
+// Project include(s).
+#include "traccc/definitions/primitives.hpp"
+#include "traccc/definitions/qualifiers.hpp"
+#include "traccc/edm/track_candidate.hpp"
+
 namespace traccc::device {
 
-TRACCC_DEVICE inline void prune_tracks(
-    std::size_t globalIndex,
-    track_candidate_container_types::const_view track_candidates_view,
-    vecmem::data::vector_view<const unsigned int> valid_indices_view,
-    track_candidate_container_types::view prune_candidates_view) {
+TRACCC_DEVICE inline void prune_tracks(std::size_t globalIndex,
+                                       const prune_tracks_payload& payload) {
 
     track_candidate_container_types::const_device track_candidates(
-        track_candidates_view);
-    vecmem::device_vector<const unsigned int> valid_indices(valid_indices_view);
+        payload.track_candidates_view);
+    vecmem::device_vector<const unsigned int> valid_indices(
+        payload.valid_indices_view);
     track_candidate_container_types::device prune_candidates(
-        prune_candidates_view);
+        payload.prune_candidates_view);
 
     if (globalIndex >= prune_candidates.size()) {
         return;
@@ -40,4 +43,4 @@ TRACCC_DEVICE inline void prune_tracks(
     }
 }
 
-}  // namespace traccc::device
\ No newline at end of file
+}  // namespace traccc::device
diff --git a/device/common/include/traccc/finding/device/make_barcode_sequence.hpp b/device/common/include/traccc/finding/device/make_barcode_sequence.hpp
index 91f85fcdaf..aa848d4b87 100644
--- a/device/common/include/traccc/finding/device/make_barcode_sequence.hpp
+++ b/device/common/include/traccc/finding/device/make_barcode_sequence.hpp
@@ -13,19 +13,17 @@
 #include "traccc/edm/measurement.hpp"
 
 namespace traccc::device {
+struct make_barcode_sequence_payload {
+    measurement_collection_types::const_view uniques_view;
+    vecmem::data::vector_view<detray::geometry::barcode> barcodes_view;
+};
 
 /// Function filling the barcode sequence
 ///
 /// @param[in] globalIndex   The index of the current thread
-/// @param[in] uniques_view   Measurement container view object
-/// @param[out] barcodes_view   Unsorted module map of <module ID, header ID>
-///
+/// @param[inout] payload      The function call payload
 TRACCC_DEVICE inline void make_barcode_sequence(
-    std::size_t globalIndex,
-    measurement_collection_types::const_view uniques_view,
-    vecmem::data::vector_view<detray::geometry::barcode> barcodes_view);
-
+    std::size_t globalIndex, const make_barcode_sequence_payload& payload);
 }  // namespace traccc::device
 
-// Include the implementation.
-#include "traccc/finding/device/impl/make_barcode_sequence.ipp"
+#include "./impl/make_barcode_sequence.ipp"
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 fbb32645cf..299bf92a72 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
@@ -12,8 +12,24 @@
 #include "traccc/definitions/qualifiers.hpp"
 #include "traccc/edm/measurement.hpp"
 #include "traccc/edm/track_parameters.hpp"
+#include "traccc/finding/candidate_link.hpp"
+#include "traccc/utils/particle.hpp"
 
 namespace traccc::device {
+template <typename propagator_t, typename bfield_t>
+struct propagate_to_next_surface_payload {
+    typename propagator_t::detector_type::view_type det_data;
+    bfield_t field_data;
+    bound_track_parameters_collection_types::view params_view;
+    vecmem::data::vector_view<unsigned int> params_liveness_view;
+    const vecmem::data::vector_view<const unsigned int>& param_ids_view;
+    vecmem::data::vector_view<const candidate_link> links_view;
+    const unsigned int step;
+    const unsigned int n_in_params;
+    vecmem::data::vector_view<typename candidate_link::link_index_type>
+        tips_view;
+    vecmem::data::vector_view<unsigned int> n_tracks_per_seed_view;
+};
 
 /// Function for propagating the kalman-updated tracks to the next surface
 ///
@@ -24,33 +40,11 @@ namespace traccc::device {
 ///
 /// @param[in] globalIndex        The index of the current thread
 /// @param[in] cfg                Track finding config object
-/// @param[in] det_data           Detector view object
-/// @param[in] in_params_view     Input parameters
-/// @param[in] param_ids_view     Sorted param ids
-/// @param[in] links_view         Link container for the current step
-/// @param[in] step               Step index
-/// @param[in] n_in_params        The number of input parameters
-/// @param[out] out_params_view   Output parameters
-/// @param[out] param_to_link_view  Container for param index -> link index
-/// @param[out] tips_view         Tip link container for the current step
-/// @param[out] n_tracks_per_seed_view  Number of tracks per seed
-/// @param[out] n_out_params      The number of output parameters
-///
+/// @param[inout] payload      The function call payload
 template <typename propagator_t, typename bfield_t, typename config_t>
 TRACCC_DEVICE inline void propagate_to_next_surface(
     std::size_t globalIndex, const config_t cfg,
-    typename propagator_t::detector_type::view_type det_data,
-    bfield_t field_data,
-    bound_track_parameters_collection_types::view params_view,
-    vecmem::data::vector_view<unsigned int> params_liveness_view,
-    const vecmem::data::vector_view<const unsigned int>& param_ids_view,
-    vecmem::data::vector_view<const candidate_link> links_view,
-    const unsigned int step, const unsigned int n_in_params,
-    vecmem::data::vector_view<typename candidate_link::link_index_type>
-        tips_view,
-    vecmem::data::vector_view<unsigned int> n_tracks_per_seed_view);
-
+    const propagate_to_next_surface_payload<propagator_t, bfield_t>& payload);
 }  // namespace traccc::device
 
-// Include the implementation.
-#include "traccc/finding/device/impl/propagate_to_next_surface.ipp"
+#include "./impl/propagate_to_next_surface.ipp"
diff --git a/device/common/include/traccc/finding/device/prune_tracks.hpp b/device/common/include/traccc/finding/device/prune_tracks.hpp
index 5ff52e1ac1..85bd8b2f1d 100644
--- a/device/common/include/traccc/finding/device/prune_tracks.hpp
+++ b/device/common/include/traccc/finding/device/prune_tracks.hpp
@@ -14,19 +14,18 @@
 
 namespace traccc::device {
 
+struct prune_tracks_payload {
+    track_candidate_container_types::const_view track_candidates_view;
+    vecmem::data::vector_view<const unsigned int> valid_indices_view;
+    track_candidate_container_types::view prune_candidates_view;
+};
+
 /// Return a new track_candidates based on the criteria in configuration
 ///
 /// @param[in] globalIndex         The index of the current thread
-/// @param[in] track_candidates_view  Track candidate container view
-/// @param[in] valid_indices_view   Valid indices meeting criteria
-/// @param[out] prune_candidates_view  Track candidate container view
-TRACCC_DEVICE inline void prune_tracks(
-    std::size_t globalIndex,
-    track_candidate_container_types::const_view track_candidates_view,
-    vecmem::data::vector_view<const unsigned int> valid_indices_view,
-    track_candidate_container_types::view prune_candidates_view);
-
+/// @param[inout] payload      The function call payload
+TRACCC_DEVICE inline void prune_tracks(std::size_t globalIndex,
+                                       const prune_tracks_payload& payload);
 }  // namespace traccc::device
 
-// Include the implementation.
-#include "traccc/finding/device/impl/prune_tracks.ipp"
\ No newline at end of file
+#include "./impl/prune_tracks.ipp"
diff --git a/device/cuda/CMakeLists.txt b/device/cuda/CMakeLists.txt
index 30d01390e3..104e987a01 100644
--- a/device/cuda/CMakeLists.txt
+++ b/device/cuda/CMakeLists.txt
@@ -49,9 +49,25 @@ traccc_add_library( traccc_cuda cuda TYPE SHARED
   "src/clusterization/clusterization_algorithm.cu"
   "include/traccc/cuda/clusterization/measurement_sorting_algorithm.hpp"
   "src/clusterization/measurement_sorting_algorithm.cu"
+  "src/clusterization/kernels/ccl_kernel.cu"
+  "src/clusterization/kernels/ccl_kernel.cuh"
   # Finding
   "include/traccc/cuda/finding/finding_algorithm.hpp"
   "src/finding/finding_algorithm.cu"
+  "src/finding/kernels/make_barcode_sequence.cu"
+  "src/finding/kernels/make_barcode_sequence.cuh"
+  "src/finding/kernels/apply_interaction.cuh"
+  "src/finding/kernels/fill_sort_keys.cu"
+  "src/finding/kernels/fill_sort_keys.cuh"
+  "src/finding/kernels/prune_tracks.cu"
+  "src/finding/kernels/prune_tracks.cuh"
+  "src/finding/kernels/build_tracks.cu"
+  "src/finding/kernels/build_tracks.cuh"
+  "src/finding/kernels/find_tracks.cuh"
+  "src/finding/kernels/propagate_to_next_surface.cuh"
+  "src/finding/kernels/specializations/find_tracks_default_detector.cu"
+  "src/finding/kernels/specializations/apply_interaction_default_detector.cu"
+  "src/finding/kernels/specializations/propagate_to_next_surface_default_detector.cu"
   # Fitting
   "include/traccc/cuda/fitting/fitting_algorithm.hpp"
   "src/fitting/fitting_algorithm.cu")
diff --git a/device/cuda/src/clusterization/clusterization_algorithm.cu b/device/cuda/src/clusterization/clusterization_algorithm.cu
index 96983fb065..dd8df59d4e 100644
--- a/device/cuda/src/clusterization/clusterization_algorithm.cu
+++ b/device/cuda/src/clusterization/clusterization_algorithm.cu
@@ -11,6 +11,7 @@
 #include "../utils/barrier.hpp"
 #include "../utils/cuda_error_handling.hpp"
 #include "../utils/utils.hpp"
+#include "./kernels/ccl_kernel.cuh"
 #include "traccc/clusterization/clustering_config.hpp"
 #include "traccc/clusterization/device/ccl_kernel_definitions.hpp"
 #include "traccc/cuda/clusterization/clusterization_algorithm.hpp"
@@ -18,55 +19,12 @@
 #include "traccc/utils/projections.hpp"
 #include "traccc/utils/relations.hpp"
 
-// Project include(s)
-#include "traccc/clusterization/device/ccl_kernel.hpp"
-
 // Vecmem include(s).
 #include <cstring>
 #include <vecmem/utils/copy.hpp>
 
 namespace traccc::cuda {
 
-namespace kernels {
-
-/// CUDA kernel for running @c traccc::device::ccl_kernel
-__global__ void ccl_kernel(
-    const clustering_config cfg,
-    const edm::silicon_cell_collection::const_view cells_view,
-    const silicon_detector_description::const_view det_descr_view,
-    measurement_collection_types::view measurements_view,
-    vecmem::data::vector_view<unsigned int> cell_links,
-    vecmem::data::vector_view<device::details::index_t> f_backup_view,
-    vecmem::data::vector_view<device::details::index_t> gf_backup_view,
-    vecmem::data::vector_view<unsigned char> adjc_backup_view,
-    vecmem::data::vector_view<device::details::index_t> adjv_backup_view,
-    unsigned int* backup_mutex_ptr) {
-
-    __shared__ std::size_t partition_start, partition_end;
-    __shared__ std::size_t outi;
-    extern __shared__ device::details::index_t shared_v[];
-    vecmem::device_atomic_ref<unsigned int> backup_mutex(*backup_mutex_ptr);
-
-    using vector_size_t =
-        vecmem::data::vector_view<device::details::index_t>::size_type;
-
-    vecmem::data::vector_view<device::details::index_t> f_view{
-        static_cast<vector_size_t>(cfg.max_partition_size()), shared_v};
-    vecmem::data::vector_view<device::details::index_t> gf_view{
-        static_cast<vector_size_t>(cfg.max_partition_size()),
-        shared_v + cfg.max_partition_size()};
-    traccc::cuda::barrier barry_r;
-    const cuda::thread_id1 thread_id;
-
-    device::ccl_kernel(cfg, thread_id, cells_view, det_descr_view,
-                       partition_start, partition_end, outi, f_view, gf_view,
-                       f_backup_view, gf_backup_view, adjc_backup_view,
-                       adjv_backup_view, backup_mutex, barry_r,
-                       measurements_view, cell_links);
-}
-
-}  // namespace kernels
-
 clusterization_algorithm::clusterization_algorithm(
     const traccc::memory_resource& mr, vecmem::copy& copy, stream& str,
     const config_type& config)
diff --git a/device/cuda/src/clusterization/kernels/ccl_kernel.cu b/device/cuda/src/clusterization/kernels/ccl_kernel.cu
new file mode 100644
index 0000000000..ca10b1a534
--- /dev/null
+++ b/device/cuda/src/clusterization/kernels/ccl_kernel.cu
@@ -0,0 +1,65 @@
+/** 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
+ */
+
+// CUDA Library include(s).
+#include "../../sanity/contiguous_on.cuh"
+#include "../../sanity/ordered_on.cuh"
+#include "../../utils/barrier.hpp"
+#include "../../utils/cuda_error_handling.hpp"
+#include "../../utils/utils.hpp"
+#include "traccc/clusterization/clustering_config.hpp"
+#include "traccc/clusterization/device/ccl_kernel_definitions.hpp"
+#include "traccc/cuda/clusterization/clusterization_algorithm.hpp"
+#include "traccc/cuda/utils/thread_id.hpp"
+#include "traccc/utils/projections.hpp"
+#include "traccc/utils/relations.hpp"
+
+// Project include(s)
+#include "traccc/clusterization/device/ccl_kernel.hpp"
+
+// Vecmem include(s).
+#include <cstring>
+#include <vecmem/utils/copy.hpp>
+
+namespace traccc::cuda::kernels {
+
+/// CUDA kernel for running @c traccc::device::ccl_kernel
+__global__ void ccl_kernel(
+    const clustering_config cfg,
+    const edm::silicon_cell_collection::const_view cells_view,
+    const silicon_detector_description::const_view det_descr_view,
+    measurement_collection_types::view measurements_view,
+    vecmem::data::vector_view<unsigned int> cell_links,
+    vecmem::data::vector_view<device::details::index_t> f_backup_view,
+    vecmem::data::vector_view<device::details::index_t> gf_backup_view,
+    vecmem::data::vector_view<unsigned char> adjc_backup_view,
+    vecmem::data::vector_view<device::details::index_t> adjv_backup_view,
+    unsigned int* backup_mutex_ptr) {
+
+    __shared__ std::size_t partition_start, partition_end;
+    __shared__ std::size_t outi;
+    extern __shared__ device::details::index_t shared_v[];
+    vecmem::device_atomic_ref<unsigned int> backup_mutex(*backup_mutex_ptr);
+
+    using vector_size_t =
+        vecmem::data::vector_view<device::details::index_t>::size_type;
+
+    vecmem::data::vector_view<device::details::index_t> f_view{
+        static_cast<vector_size_t>(cfg.max_partition_size()), shared_v};
+    vecmem::data::vector_view<device::details::index_t> gf_view{
+        static_cast<vector_size_t>(cfg.max_partition_size()),
+        shared_v + cfg.max_partition_size()};
+    traccc::cuda::barrier barry_r;
+    const cuda::thread_id1 thread_id;
+
+    device::ccl_kernel(cfg, thread_id, cells_view, det_descr_view,
+                       partition_start, partition_end, outi, f_view, gf_view,
+                       f_backup_view, gf_backup_view, adjc_backup_view,
+                       adjv_backup_view, backup_mutex, barry_r,
+                       measurements_view, cell_links);
+}
+}  // namespace traccc::cuda::kernels
diff --git a/device/cuda/src/clusterization/kernels/ccl_kernel.cuh b/device/cuda/src/clusterization/kernels/ccl_kernel.cuh
new file mode 100644
index 0000000000..dc0cdcd341
--- /dev/null
+++ b/device/cuda/src/clusterization/kernels/ccl_kernel.cuh
@@ -0,0 +1,27 @@
+/** 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
+
+#include "traccc/clusterization/clustering_config.hpp"
+#include "traccc/clusterization/device/ccl_kernel_definitions.hpp"
+#include "traccc/cuda/clusterization/clusterization_algorithm.hpp"
+
+namespace traccc::cuda::kernels {
+
+__global__ void ccl_kernel(
+    const clustering_config cfg,
+    const edm::silicon_cell_collection::const_view cells_view,
+    const silicon_detector_description::const_view det_descr_view,
+    measurement_collection_types::view measurements_view,
+    vecmem::data::vector_view<unsigned int> cell_links,
+    vecmem::data::vector_view<device::details::index_t> f_backup_view,
+    vecmem::data::vector_view<device::details::index_t> gf_backup_view,
+    vecmem::data::vector_view<unsigned char> adjc_backup_view,
+    vecmem::data::vector_view<device::details::index_t> adjv_backup_view,
+    unsigned int* backup_mutex_ptr);
+}
diff --git a/device/cuda/src/finding/finding_algorithm.cu b/device/cuda/src/finding/finding_algorithm.cu
index 7c6fc54169..ee9de73c2c 100644
--- a/device/cuda/src/finding/finding_algorithm.cu
+++ b/device/cuda/src/finding/finding_algorithm.cu
@@ -10,19 +10,19 @@
 #include "../utils/barrier.hpp"
 #include "../utils/cuda_error_handling.hpp"
 #include "../utils/utils.hpp"
+#include "./kernels/apply_interaction.cuh"
+#include "./kernels/build_tracks.cuh"
+#include "./kernels/fill_sort_keys.cuh"
+#include "./kernels/find_tracks.cuh"
+#include "./kernels/make_barcode_sequence.cuh"
+#include "./kernels/propagate_to_next_surface.cuh"
+#include "./kernels/prune_tracks.cuh"
 #include "traccc/cuda/finding/finding_algorithm.hpp"
 #include "traccc/cuda/utils/thread_id.hpp"
 #include "traccc/definitions/primitives.hpp"
 #include "traccc/definitions/qualifiers.hpp"
 #include "traccc/edm/device/sort_key.hpp"
 #include "traccc/finding/candidate_link.hpp"
-#include "traccc/finding/device/apply_interaction.hpp"
-#include "traccc/finding/device/build_tracks.hpp"
-#include "traccc/finding/device/fill_sort_keys.hpp"
-#include "traccc/finding/device/find_tracks.hpp"
-#include "traccc/finding/device/make_barcode_sequence.hpp"
-#include "traccc/finding/device/propagate_to_next_surface.hpp"
-#include "traccc/finding/device/prune_tracks.hpp"
 #include "traccc/utils/projections.hpp"
 
 // detray include(s).
@@ -53,132 +53,6 @@
 #include <vector>
 
 namespace traccc::cuda {
-namespace kernels {
-
-/// CUDA kernel for running @c traccc::device::make_barcode_sequence
-__global__ void make_barcode_sequence(
-    measurement_collection_types::const_view measurements_view,
-    vecmem::data::vector_view<detray::geometry::barcode> barcodes_view) {
-
-    int gid = threadIdx.x + blockIdx.x * blockDim.x;
-
-    device::make_barcode_sequence(gid, measurements_view, barcodes_view);
-}
-
-/// CUDA kernel for running @c traccc::device::apply_interaction
-template <typename detector_t>
-__global__ void apply_interaction(
-    typename detector_t::view_type det_data, const finding_config cfg,
-    const int n_params,
-    bound_track_parameters_collection_types::view params_view,
-    vecmem::data::vector_view<const unsigned int> params_liveness_view) {
-
-    int gid = threadIdx.x + blockIdx.x * blockDim.x;
-
-    device::apply_interaction<detector_t>(gid, cfg, det_data, n_params,
-                                          params_view, params_liveness_view);
-}
-
-/// CUDA kernel for running @c traccc::device::find_tracks
-template <typename detector_t, typename config_t>
-__global__ void find_tracks(
-    const config_t cfg, typename detector_t::view_type det_data,
-    measurement_collection_types::const_view measurements_view,
-    bound_track_parameters_collection_types::const_view in_params_view,
-    vecmem::data::vector_view<const unsigned int> in_params_liveness_view,
-    const unsigned int n_in_params,
-    vecmem::data::vector_view<const detray::geometry::barcode> barcodes_view,
-    vecmem::data::vector_view<const unsigned int> upper_bounds_view,
-    vecmem::data::vector_view<const candidate_link> prev_links_view,
-    const unsigned int step, const unsigned int n_max_candidates,
-    bound_track_parameters_collection_types::view out_params_view,
-    vecmem::data::vector_view<unsigned int> out_params_liveness_view,
-    vecmem::data::vector_view<candidate_link> links_view,
-    unsigned int* n_candidates) {
-    __shared__ unsigned int shared_candidates_size;
-    extern __shared__ unsigned int s[];
-    unsigned int* shared_num_candidates = s;
-    std::pair<unsigned int, unsigned int>* shared_candidates =
-        reinterpret_cast<std::pair<unsigned int, unsigned int>*>(
-            &shared_num_candidates[blockDim.x]);
-
-    cuda::barrier barrier;
-    cuda::thread_id1 thread_id;
-
-    device::find_tracks<cuda::thread_id1, cuda::barrier, detector_t, config_t>(
-        thread_id, barrier, cfg, det_data, measurements_view, in_params_view,
-        in_params_liveness_view, n_in_params, barcodes_view, upper_bounds_view,
-        prev_links_view, step, n_max_candidates, out_params_view,
-        out_params_liveness_view, links_view, *n_candidates,
-        shared_num_candidates, shared_candidates, shared_candidates_size);
-}
-
-/// CUDA kernel for running @c traccc::device::fill_sort_keys
-__global__ void fill_sort_keys(
-    bound_track_parameters_collection_types::const_view params_view,
-    vecmem::data::vector_view<device::sort_key> keys_view,
-    vecmem::data::vector_view<unsigned int> ids_view) {
-
-    device::fill_sort_keys(threadIdx.x + blockIdx.x * blockDim.x, params_view,
-                           keys_view, ids_view);
-}
-
-/// CUDA kernel for running @c traccc::device::propagate_to_next_surface
-template <typename propagator_t, typename bfield_t, typename config_t>
-__global__ void propagate_to_next_surface(
-    const config_t cfg,
-    typename propagator_t::detector_type::view_type det_data,
-    bfield_t field_data,
-    bound_track_parameters_collection_types::view params_view,
-    vecmem::data::vector_view<unsigned int> params_liveness_view,
-    vecmem::data::vector_view<const unsigned int> param_ids_view,
-    vecmem::data::vector_view<const candidate_link> links_view,
-    const unsigned int step, const unsigned int n_candidates,
-    vecmem::data::vector_view<typename candidate_link::link_index_type>
-        tips_view,
-    vecmem::data::vector_view<unsigned int> n_tracks_per_seed_view) {
-
-    int gid = threadIdx.x + blockIdx.x * blockDim.x;
-
-    device::propagate_to_next_surface<propagator_t, bfield_t, config_t>(
-        gid, cfg, det_data, field_data, params_view, params_liveness_view,
-        param_ids_view, links_view, step, n_candidates, tips_view,
-        n_tracks_per_seed_view);
-}
-
-/// CUDA kernel for running @c traccc::device::build_tracks
-template <typename config_t>
-__global__ void build_tracks(
-    const config_t cfg,
-    measurement_collection_types::const_view measurements_view,
-    bound_track_parameters_collection_types::const_view seeds_view,
-    vecmem::data::jagged_vector_view<const candidate_link> links_view,
-    vecmem::data::vector_view<const typename candidate_link::link_index_type>
-        tips_view,
-    track_candidate_container_types::view track_candidates_view,
-    vecmem::data::vector_view<unsigned int> valid_indices_view,
-    unsigned int* n_valid_tracks) {
-
-    int gid = threadIdx.x + blockIdx.x * blockDim.x;
-
-    device::build_tracks(gid, cfg, measurements_view, seeds_view, links_view,
-                         tips_view, track_candidates_view, valid_indices_view,
-                         *n_valid_tracks);
-}
-
-/// CUDA kernel for running @c traccc::device::prune_tracks
-__global__ void prune_tracks(
-    track_candidate_container_types::const_view track_candidates_view,
-    vecmem::data::vector_view<const unsigned int> valid_indices_view,
-    track_candidate_container_types::view prune_candidates_view) {
-
-    int gid = threadIdx.x + blockIdx.x * blockDim.x;
-
-    device::prune_tracks(gid, track_candidates_view, valid_indices_view,
-                         prune_candidates_view);
-}
-
-}  // namespace kernels
 
 template <typename stepper_t, typename navigator_t>
 finding_algorithm<stepper_t, navigator_t>::finding_algorithm(
@@ -261,7 +135,7 @@ finding_algorithm<stepper_t, navigator_t>::operator()(
             (barcodes_buffer.size() + nThreads - 1) / nThreads;
 
         kernels::make_barcode_sequence<<<nBlocks, nThreads, 0, stream>>>(
-            uniques_buffer, barcodes_buffer);
+            {uniques_buffer, barcodes_buffer});
 
         TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
     }
@@ -312,10 +186,10 @@ finding_algorithm<stepper_t, navigator_t>::operator()(
             const unsigned int nBlocks =
                 (n_in_params + nThreads - 1) / nThreads;
 
-            kernels::apply_interaction<detector_type>
-                <<<nBlocks, nThreads, 0, stream>>>(det_view, m_cfg, n_in_params,
-                                                   in_params_buffer,
-                                                   param_liveness_buffer);
+            kernels::apply_interaction<std::decay_t<detector_type>>
+                <<<nBlocks, nThreads, 0, stream>>>(
+                    m_cfg, {det_view, static_cast<int>(n_in_params),
+                            in_params_buffer, param_liveness_buffer});
             TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
         }
 
@@ -358,17 +232,18 @@ finding_algorithm<stepper_t, navigator_t>::operator()(
             TRACCC_CUDA_ERROR_CHECK(cudaMemsetAsync(
                 n_candidates_device.get(), 0, sizeof(unsigned int), stream));
 
-            kernels::find_tracks<detector_type, config_type>
+            kernels::find_tracks<std::decay_t<detector_type>>
                 <<<nBlocks, nThreads,
                    nThreads * sizeof(unsigned int) +
                        2 * nThreads *
                            sizeof(std::pair<unsigned int, unsigned int>),
-                   stream>>>(m_cfg, det_view, measurements, in_params_buffer,
-                             param_liveness_buffer, n_in_params,
-                             barcodes_buffer, upper_bounds_buffer,
-                             link_map[prev_step], step, n_max_candidates,
-                             updated_params_buffer, updated_liveness_buffer,
-                             link_map[step], n_candidates_device.get());
+                   stream>>>(
+                    m_cfg, {det_view, measurements, in_params_buffer,
+                            param_liveness_buffer, n_in_params, barcodes_buffer,
+                            upper_bounds_buffer, link_map[prev_step], step,
+                            n_max_candidates, updated_params_buffer,
+                            updated_liveness_buffer, link_map[step],
+                            n_candidates_device.get()});
             TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
 
             std::swap(in_params_buffer, updated_params_buffer);
@@ -399,7 +274,7 @@ finding_algorithm<stepper_t, navigator_t>::operator()(
                 const unsigned int nBlocks =
                     (n_candidates + nThreads - 1) / nThreads;
                 kernels::fill_sort_keys<<<nBlocks, nThreads, 0, stream>>>(
-                    in_params_buffer, keys_buffer, param_ids_buffer);
+                    {in_params_buffer, keys_buffer, param_ids_buffer});
                 TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
 
                 // Sort the key and values
@@ -430,13 +305,13 @@ finding_algorithm<stepper_t, navigator_t>::operator()(
                 const unsigned int nThreads = m_warp_size * 2;
                 const unsigned int nBlocks =
                     (n_candidates + nThreads - 1) / nThreads;
-                kernels::propagate_to_next_surface<propagator_type, bfield_type,
-                                                   config_type>
+                kernels::propagate_to_next_surface<
+                    std::decay_t<propagator_type>, std::decay_t<bfield_type>>
                     <<<nBlocks, nThreads, 0, stream>>>(
-                        m_cfg, det_view, field_view, in_params_buffer,
-                        param_liveness_buffer, param_ids_buffer, link_map[step],
-                        step, n_candidates, tips_map[step],
-                        n_tracks_per_seed_buffer);
+                        m_cfg, {det_view, field_view, in_params_buffer,
+                                param_liveness_buffer, param_ids_buffer,
+                                link_map[step], step, n_candidates,
+                                tips_map[step], n_tracks_per_seed_buffer});
                 TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
 
                 m_stream.synchronize();
@@ -557,8 +432,8 @@ finding_algorithm<stepper_t, navigator_t>::operator()(
         const unsigned int nBlocks = (n_valid_tracks + nThreads - 1) / nThreads;
 
         kernels::prune_tracks<<<nBlocks, nThreads, 0, stream>>>(
-            track_candidates_buffer, valid_indices_buffer,
-            prune_candidates_buffer);
+            {track_candidates_buffer, valid_indices_buffer,
+             prune_candidates_buffer});
         TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
     }
 
diff --git a/device/cuda/src/finding/kernels/apply_interaction.cuh b/device/cuda/src/finding/kernels/apply_interaction.cuh
new file mode 100644
index 0000000000..cf2f1a5507
--- /dev/null
+++ b/device/cuda/src/finding/kernels/apply_interaction.cuh
@@ -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
+ */
+
+#pragma once
+
+#include "traccc/edm/track_parameters.hpp"
+#include "traccc/finding/device/apply_interaction.hpp"
+#include "traccc/geometry/detector.hpp"
+
+namespace traccc::cuda::kernels {
+
+template <typename detector_t>
+__global__ void apply_interaction(
+    const finding_config cfg,
+    device::apply_interaction_payload<detector_t> payload);
+
+extern template __global__ void
+apply_interaction<traccc::default_detector::device>(
+    const finding_config,
+    device::apply_interaction_payload<traccc::default_detector::device>);
+}  // namespace traccc::cuda::kernels
diff --git a/device/cuda/src/finding/kernels/build_tracks.cu b/device/cuda/src/finding/kernels/build_tracks.cu
new file mode 100644
index 0000000000..9280aba367
--- /dev/null
+++ b/device/cuda/src/finding/kernels/build_tracks.cu
@@ -0,0 +1,36 @@
+/** 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
+ */
+
+#include "build_tracks.cuh"
+#include "traccc/edm/measurement.hpp"
+#include "traccc/edm/track_candidate.hpp"
+#include "traccc/edm/track_parameters.hpp"
+#include "traccc/finding/candidate_link.hpp"
+#include "traccc/finding/device/build_tracks.hpp"
+#include "traccc/finding/finding_config.hpp"
+
+namespace traccc::cuda::kernels {
+
+__global__ void build_tracks(
+    const finding_config cfg,
+    measurement_collection_types::const_view measurements_view,
+    bound_track_parameters_collection_types::const_view seeds_view,
+    vecmem::data::jagged_vector_view<const candidate_link> links_view,
+    vecmem::data::vector_view<const typename candidate_link::link_index_type>
+        tips_view,
+    track_candidate_container_types::view track_candidates_view,
+    vecmem::data::vector_view<unsigned int> valid_indices_view,
+    unsigned int* n_valid_tracks) {
+
+    int gid = threadIdx.x + blockIdx.x * blockDim.x;
+
+    device::build_tracks(
+        gid, cfg,
+        {measurements_view, seeds_view, links_view, tips_view,
+         track_candidates_view, valid_indices_view, *n_valid_tracks});
+}
+}  // namespace traccc::cuda::kernels
diff --git a/device/cuda/src/finding/kernels/build_tracks.cuh b/device/cuda/src/finding/kernels/build_tracks.cuh
new file mode 100644
index 0000000000..0393692d51
--- /dev/null
+++ b/device/cuda/src/finding/kernels/build_tracks.cuh
@@ -0,0 +1,28 @@
+/** 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
+
+#include "traccc/edm/measurement.hpp"
+#include "traccc/edm/track_candidate.hpp"
+#include "traccc/edm/track_parameters.hpp"
+#include "traccc/finding/candidate_link.hpp"
+#include "traccc/finding/finding_config.hpp"
+
+namespace traccc::cuda::kernels {
+
+__global__ void build_tracks(
+    const finding_config cfg,
+    measurement_collection_types::const_view measurements_view,
+    bound_track_parameters_collection_types::const_view seeds_view,
+    vecmem::data::jagged_vector_view<const candidate_link> links_view,
+    vecmem::data::vector_view<const typename candidate_link::link_index_type>
+        tips_view,
+    track_candidate_container_types::view track_candidates_view,
+    vecmem::data::vector_view<unsigned int> valid_indices_view,
+    unsigned int* n_valid_tracks);
+}
diff --git a/device/cuda/src/finding/kernels/fill_sort_keys.cu b/device/cuda/src/finding/kernels/fill_sort_keys.cu
new file mode 100644
index 0000000000..4115841c0a
--- /dev/null
+++ b/device/cuda/src/finding/kernels/fill_sort_keys.cu
@@ -0,0 +1,18 @@
+/** 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
+ */
+
+#include "fill_sort_keys.cuh"
+#include "traccc/edm/track_parameters.hpp"
+#include "traccc/finding/device/fill_sort_keys.hpp"
+
+namespace traccc::cuda::kernels {
+
+__global__ void fill_sort_keys(device::fill_sort_keys_payload payload) {
+
+    device::fill_sort_keys(threadIdx.x + blockIdx.x * blockDim.x, payload);
+}
+}  // namespace traccc::cuda::kernels
diff --git a/device/cuda/src/finding/kernels/fill_sort_keys.cuh b/device/cuda/src/finding/kernels/fill_sort_keys.cuh
new file mode 100644
index 0000000000..5f9aedb22c
--- /dev/null
+++ b/device/cuda/src/finding/kernels/fill_sort_keys.cuh
@@ -0,0 +1,16 @@
+/** 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
+
+#include "traccc/edm/track_parameters.hpp"
+#include "traccc/finding/device/fill_sort_keys.hpp"
+
+namespace traccc::cuda::kernels {
+
+__global__ void fill_sort_keys(device::fill_sort_keys_payload payload);
+}
diff --git a/device/cuda/src/finding/kernels/find_tracks.cuh b/device/cuda/src/finding/kernels/find_tracks.cuh
new file mode 100644
index 0000000000..8ce32cf262
--- /dev/null
+++ b/device/cuda/src/finding/kernels/find_tracks.cuh
@@ -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
+ */
+
+#pragma once
+
+#include "../../utils/barrier.hpp"
+#include "traccc/cuda/utils/thread_id.hpp"
+#include "traccc/edm/track_parameters.hpp"
+#include "traccc/finding/device/find_tracks.hpp"
+#include "traccc/geometry/detector.hpp"
+
+namespace traccc::cuda::kernels {
+
+template <typename detector_t>
+__global__ void find_tracks(const finding_config cfg,
+                            device::find_tracks_payload<detector_t> payload);
+
+extern template __global__ void find_tracks<traccc::default_detector::device>(
+    const finding_config cfg,
+    device::find_tracks_payload<traccc::default_detector::device> payload);
+}  // namespace traccc::cuda::kernels
diff --git a/device/cuda/src/finding/kernels/make_barcode_sequence.cu b/device/cuda/src/finding/kernels/make_barcode_sequence.cu
new file mode 100644
index 0000000000..e6587b553d
--- /dev/null
+++ b/device/cuda/src/finding/kernels/make_barcode_sequence.cu
@@ -0,0 +1,21 @@
+/** 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
+ */
+
+#include "make_barcode_sequence.cuh"
+#include "traccc/edm/measurement.hpp"
+#include "traccc/finding/device/make_barcode_sequence.hpp"
+
+namespace traccc::cuda::kernels {
+
+__global__ void make_barcode_sequence(
+    device::make_barcode_sequence_payload payload) {
+
+    int gid = threadIdx.x + blockIdx.x * blockDim.x;
+
+    device::make_barcode_sequence(gid, payload);
+}
+}  // namespace traccc::cuda::kernels
diff --git a/device/cuda/src/finding/kernels/make_barcode_sequence.cuh b/device/cuda/src/finding/kernels/make_barcode_sequence.cuh
new file mode 100644
index 0000000000..13f147a047
--- /dev/null
+++ b/device/cuda/src/finding/kernels/make_barcode_sequence.cuh
@@ -0,0 +1,17 @@
+/** 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
+
+#include "traccc/edm/measurement.hpp"
+#include "traccc/finding/device/make_barcode_sequence.hpp"
+
+namespace traccc::cuda::kernels {
+
+__global__ void make_barcode_sequence(
+    device::make_barcode_sequence_payload payload);
+}
diff --git a/device/cuda/src/finding/kernels/propagate_to_next_surface.cuh b/device/cuda/src/finding/kernels/propagate_to_next_surface.cuh
new file mode 100644
index 0000000000..6854aebcd5
--- /dev/null
+++ b/device/cuda/src/finding/kernels/propagate_to_next_surface.cuh
@@ -0,0 +1,28 @@
+/** 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
+
+#include "./specializations/types.hpp"
+#include "traccc/finding/device/propagate_to_next_surface.hpp"
+#include "traccc/finding/finding_config.hpp"
+#include "traccc/geometry/detector.hpp"
+
+namespace traccc::cuda::kernels {
+
+template <typename propagator_t, typename bfield_t>
+__global__ void propagate_to_next_surface(
+    const finding_config cfg,
+    device::propagate_to_next_surface_payload<propagator_t, bfield_t> payload);
+
+extern template __global__ void
+propagate_to_next_surface<default_finding_algorithm::propagator_type,
+                          default_finding_algorithm::bfield_type>(
+    const finding_config, device::propagate_to_next_surface_payload<
+                              default_finding_algorithm::propagator_type,
+                              default_finding_algorithm::bfield_type>);
+}  // namespace traccc::cuda::kernels
diff --git a/device/cuda/src/finding/kernels/prune_tracks.cu b/device/cuda/src/finding/kernels/prune_tracks.cu
new file mode 100644
index 0000000000..f431676a4d
--- /dev/null
+++ b/device/cuda/src/finding/kernels/prune_tracks.cu
@@ -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
+ */
+
+#include "prune_tracks.cuh"
+#include "traccc/finding/device/prune_tracks.hpp"
+
+namespace traccc::cuda::kernels {
+
+__global__ void prune_tracks(device::prune_tracks_payload payload) {
+
+    int gid = threadIdx.x + blockIdx.x * blockDim.x;
+
+    device::prune_tracks(gid, payload);
+}
+}  // namespace traccc::cuda::kernels
diff --git a/device/cuda/src/finding/kernels/prune_tracks.cuh b/device/cuda/src/finding/kernels/prune_tracks.cuh
new file mode 100644
index 0000000000..2aaa23b6fa
--- /dev/null
+++ b/device/cuda/src/finding/kernels/prune_tracks.cuh
@@ -0,0 +1,16 @@
+/** 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
+
+#include "traccc/edm/track_parameters.hpp"
+#include "traccc/finding/device/prune_tracks.hpp"
+
+namespace traccc::cuda::kernels {
+
+__global__ void prune_tracks(device::prune_tracks_payload payload);
+}
diff --git a/device/cuda/src/finding/kernels/specializations/apply_interaction_default_detector.cu b/device/cuda/src/finding/kernels/specializations/apply_interaction_default_detector.cu
new file mode 100644
index 0000000000..5fa6f073b2
--- /dev/null
+++ b/device/cuda/src/finding/kernels/specializations/apply_interaction_default_detector.cu
@@ -0,0 +1,14 @@
+/** 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
+ */
+
+#include "apply_interaction_src.cuh"
+
+namespace traccc::cuda::kernels {
+template __global__ void apply_interaction<traccc::default_detector::device>(
+    const finding_config,
+    device::apply_interaction_payload<traccc::default_detector::device>);
+}
diff --git a/device/cuda/src/finding/kernels/specializations/apply_interaction_src.cuh b/device/cuda/src/finding/kernels/specializations/apply_interaction_src.cuh
new file mode 100644
index 0000000000..53b6b2e2a9
--- /dev/null
+++ b/device/cuda/src/finding/kernels/specializations/apply_interaction_src.cuh
@@ -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
+ */
+
+#pragma once
+
+#include "traccc/edm/track_parameters.hpp"
+#include "traccc/finding/device/apply_interaction.hpp"
+#include "traccc/geometry/detector.hpp"
+
+namespace traccc::cuda::kernels {
+
+template <typename detector_t>
+__global__ void apply_interaction(
+    const finding_config cfg,
+    device::apply_interaction_payload<detector_t> payload) {
+
+    int gid = threadIdx.x + blockIdx.x * blockDim.x;
+
+    device::apply_interaction<detector_t>(gid, cfg, payload);
+}
+}  // namespace traccc::cuda::kernels
diff --git a/device/cuda/src/finding/kernels/specializations/find_tracks_default_detector.cu b/device/cuda/src/finding/kernels/specializations/find_tracks_default_detector.cu
new file mode 100644
index 0000000000..c2bb3ba910
--- /dev/null
+++ b/device/cuda/src/finding/kernels/specializations/find_tracks_default_detector.cu
@@ -0,0 +1,14 @@
+/** 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
+ */
+
+#include "find_tracks_src.cuh"
+
+namespace traccc::cuda::kernels {
+template __global__ void find_tracks<traccc::default_detector::device>(
+    const finding_config cfg,
+    device::find_tracks_payload<traccc::default_detector::device> payload);
+}
diff --git a/device/cuda/src/finding/kernels/specializations/find_tracks_src.cuh b/device/cuda/src/finding/kernels/specializations/find_tracks_src.cuh
new file mode 100644
index 0000000000..11b2733a01
--- /dev/null
+++ b/device/cuda/src/finding/kernels/specializations/find_tracks_src.cuh
@@ -0,0 +1,36 @@
+/** 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
+
+#include "../../../utils/barrier.hpp"
+#include "traccc/cuda/utils/thread_id.hpp"
+#include "traccc/edm/track_parameters.hpp"
+#include "traccc/finding/device/find_tracks.hpp"
+#include "traccc/geometry/detector.hpp"
+
+namespace traccc::cuda::kernels {
+
+template <typename detector_t>
+__global__ void find_tracks(const finding_config cfg,
+                            device::find_tracks_payload<detector_t> payload) {
+    __shared__ unsigned int shared_candidates_size;
+    extern __shared__ unsigned int s[];
+    unsigned int* shared_num_candidates = s;
+    std::pair<unsigned int, unsigned int>* shared_candidates =
+        reinterpret_cast<std::pair<unsigned int, unsigned int>*>(
+            &shared_num_candidates[blockDim.x]);
+
+    cuda::barrier barrier;
+    cuda::thread_id1 thread_id;
+
+    device::find_tracks<cuda::thread_id1, cuda::barrier, detector_t,
+                        finding_config>(
+        thread_id, barrier, cfg, payload,
+        {shared_num_candidates, shared_candidates, shared_candidates_size});
+}
+}  // namespace traccc::cuda::kernels
diff --git a/device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_default_detector.cu b/device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_default_detector.cu
new file mode 100644
index 0000000000..c992a67e1d
--- /dev/null
+++ b/device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_default_detector.cu
@@ -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
+ */
+
+#include "./types.hpp"
+#include "propagate_to_next_surface_src.cuh"
+
+namespace traccc::cuda::kernels {
+
+template __global__ void
+propagate_to_next_surface<default_finding_algorithm::propagator_type,
+                          default_finding_algorithm::bfield_type>(
+    const finding_config, device::propagate_to_next_surface_payload<
+                              default_finding_algorithm::propagator_type,
+                              default_finding_algorithm::bfield_type>);
+}  // namespace traccc::cuda::kernels
diff --git a/device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_src.cuh b/device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_src.cuh
new file mode 100644
index 0000000000..d35724a757
--- /dev/null
+++ b/device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_src.cuh
@@ -0,0 +1,27 @@
+/** 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
+
+#include "traccc/finding/device/propagate_to_next_surface.hpp"
+#include "traccc/finding/finding_config.hpp"
+#include "traccc/geometry/detector.hpp"
+
+namespace traccc::cuda::kernels {
+
+template <typename propagator_t, typename bfield_t>
+__global__ void propagate_to_next_surface(
+    const finding_config cfg,
+    device::propagate_to_next_surface_payload<propagator_t, bfield_t> payload) {
+
+    int gid = threadIdx.x + blockIdx.x * blockDim.x;
+
+    device::propagate_to_next_surface<propagator_t, bfield_t, finding_config>(
+        gid, cfg, payload);
+}
+
+}  // namespace traccc::cuda::kernels
diff --git a/device/cuda/src/finding/kernels/specializations/types.hpp b/device/cuda/src/finding/kernels/specializations/types.hpp
new file mode 100644
index 0000000000..78aab7ac58
--- /dev/null
+++ b/device/cuda/src/finding/kernels/specializations/types.hpp
@@ -0,0 +1,35 @@
+/** 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
+
+#include "detray/detectors/bfield.hpp"
+#include "detray/propagator/actor_chain.hpp"
+#include "detray/propagator/actors/aborters.hpp"
+#include "detray/propagator/actors/parameter_resetter.hpp"
+#include "detray/propagator/actors/parameter_transporter.hpp"
+#include "detray/propagator/actors/pointwise_material_interactor.hpp"
+#include "detray/propagator/propagator.hpp"
+#include "detray/propagator/rk_stepper.hpp"
+#include "traccc/cuda/finding/finding_algorithm.hpp"
+#include "traccc/finding/ckf_aborter.hpp"
+#include "traccc/finding/interaction_register.hpp"
+#include "traccc/geometry/detector.hpp"
+
+namespace traccc::cuda::kernels {
+
+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,
+                       traccc::default_algebra, detray::constrained_step<>>;
+using default_navigator_type = detray::navigator<const default_detector_type>;
+
+using default_finding_algorithm =
+    finding_algorithm<default_stepper_type, default_navigator_type>;
+
+}  // namespace traccc::cuda::kernels