From bf45c834b2aba6171d01e283deb23a4274525f0b Mon Sep 17 00:00:00 2001
From: Stephen Nicholas Swatman <stephen.nicholas.swatman@cern.ch>
Date: Wed, 16 Oct 2024 14:18:26 +0200
Subject: [PATCH] Split the CUDA CKF into different TUs

This commit splits the monstrously large CUDA track finding translation
unit up into smaller ones, one for each of the kernels. This should
speed up compilation times and decrease memory usage.

Also groups the payloads for each of the functions into convenient
structs, so we don't need to pass 20+ arguments for some of the kernel
calls.

Does not change the functionality of the code.
---
 .../traccc/finding/actors/ckf_aborter.hpp     |   3 +-
 .../finding/device/apply_interaction.hpp      |  40 +++-
 .../traccc/finding/device/build_tracks.hpp    |  70 +++++--
 .../traccc/finding/device/fill_sort_keys.hpp  |  31 ++-
 .../traccc/finding/device/find_tracks.hpp     | 146 ++++++++++----
 .../finding/device/impl/apply_interaction.ipp |  24 +--
 .../finding/device/impl/build_tracks.ipp      |  38 ++--
 .../finding/device/impl/fill_sort_keys.ipp    |  16 +-
 .../finding/device/impl/find_tracks.ipp       | 128 ++++++------
 .../device/impl/make_barcode_sequence.ipp     |  18 +-
 .../device/impl/propagate_to_next_surface.ipp |  56 +++---
 .../finding/device/impl/prune_tracks.ipp      |  21 +-
 .../finding/device/make_barcode_sequence.hpp  |  23 ++-
 .../device/propagate_to_next_surface.hpp      |  84 +++++---
 .../traccc/finding/device/prune_tracks.hpp    |  32 ++-
 device/cuda/CMakeLists.txt                    |  16 ++
 .../clusterization_algorithm.cu               |  44 +---
 .../src/clusterization/kernels/ccl_kernel.cu  |  65 ++++++
 .../src/clusterization/kernels/ccl_kernel.cuh |  27 +++
 device/cuda/src/finding/finding_algorithm.cu  | 189 +++---------------
 .../src/finding/kernels/apply_interaction.cuh |  20 ++
 .../cuda/src/finding/kernels/build_tracks.cu  |  25 +++
 .../cuda/src/finding/kernels/build_tracks.cuh |  21 ++
 .../src/finding/kernels/fill_sort_keys.cu     |  18 ++
 .../src/finding/kernels/fill_sort_keys.cuh    |  16 ++
 .../cuda/src/finding/kernels/find_tracks.cuh  |  21 ++
 .../finding/kernels/make_barcode_sequence.cu  |  21 ++
 .../finding/kernels/make_barcode_sequence.cuh |  17 ++
 .../kernels/propagate_to_next_surface.cuh     |  21 ++
 .../cuda/src/finding/kernels/prune_tracks.cu  |  19 ++
 .../cuda/src/finding/kernels/prune_tracks.cuh |  16 ++
 .../apply_interaction_default_detector.cu     |  14 ++
 .../specializations/apply_interaction_src.cuh |  25 +++
 .../find_tracks_default_detector.cu           |  14 ++
 .../specializations/find_tracks_src.cuh       |  37 ++++
 ...pagate_to_next_surface_default_detector.cu |  19 ++
 .../propagate_to_next_surface_src.cuh         |  27 +++
 .../finding/kernels/specializations/types.hpp |  35 ++++
 38 files changed, 986 insertions(+), 471 deletions(-)
 create mode 100644 device/cuda/src/clusterization/kernels/ccl_kernel.cu
 create mode 100644 device/cuda/src/clusterization/kernels/ccl_kernel.cuh
 create mode 100644 device/cuda/src/finding/kernels/apply_interaction.cuh
 create mode 100644 device/cuda/src/finding/kernels/build_tracks.cu
 create mode 100644 device/cuda/src/finding/kernels/build_tracks.cuh
 create mode 100644 device/cuda/src/finding/kernels/fill_sort_keys.cu
 create mode 100644 device/cuda/src/finding/kernels/fill_sort_keys.cuh
 create mode 100644 device/cuda/src/finding/kernels/find_tracks.cuh
 create mode 100644 device/cuda/src/finding/kernels/make_barcode_sequence.cu
 create mode 100644 device/cuda/src/finding/kernels/make_barcode_sequence.cuh
 create mode 100644 device/cuda/src/finding/kernels/propagate_to_next_surface.cuh
 create mode 100644 device/cuda/src/finding/kernels/prune_tracks.cu
 create mode 100644 device/cuda/src/finding/kernels/prune_tracks.cuh
 create mode 100644 device/cuda/src/finding/kernels/specializations/apply_interaction_default_detector.cu
 create mode 100644 device/cuda/src/finding/kernels/specializations/apply_interaction_src.cuh
 create mode 100644 device/cuda/src/finding/kernels/specializations/find_tracks_default_detector.cu
 create mode 100644 device/cuda/src/finding/kernels/specializations/find_tracks_src.cuh
 create mode 100644 device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_default_detector.cu
 create mode 100644 device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_src.cuh
 create mode 100644 device/cuda/src/finding/kernels/specializations/types.hpp

diff --git a/core/include/traccc/finding/actors/ckf_aborter.hpp b/core/include/traccc/finding/actors/ckf_aborter.hpp
index 26dfedecb9..02eba824a0 100644
--- a/core/include/traccc/finding/actors/ckf_aborter.hpp
+++ b/core/include/traccc/finding/actors/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..5650d129e2 100644
--- a/device/common/include/traccc/finding/device/apply_interaction.hpp
+++ b/device/common/include/traccc/finding/device/apply_interaction.hpp
@@ -8,29 +8,47 @@
 #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 {
+    /**
+     * @brief View object describing the tracking detector
+     */
+    typename detector_t::view_type det_data;
+
+    /**
+     * @brief Total number of input parameters (including non-live ones)
+     */
+    const int n_params;
+
+    /**
+     * @brief View object to the vector of bound track parameters
+     */
+    bound_track_parameters_collection_types::view params_view;
+
+    /**
+     * @brief View object to the vector of boolean-like integers describing
+     * whether each parameter is live. Has the same size as \ref 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..b364d3130f 100644
--- a/device/common/include/traccc/finding/device/build_tracks.hpp
+++ b/device/common/include/traccc/finding/device/build_tracks.hpp
@@ -9,8 +9,52 @@
 
 // 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 {
+    /**
+     * @brief View object to the vector of measurements
+     *
+     * @warning Measurements on the same surface must be adjacent
+     */
+    measurement_collection_types::const_view measurements_view;
+
+    /**
+     * @brief View object to the vector of measurements
+     */
+    bound_track_parameters_collection_types::const_view seeds_view;
+
+    /**
+     * @brief View object to the vector of candidate links
+     */
+    vecmem::data::jagged_vector_view<const candidate_link> links_view;
+
+    /**
+     * @brief View object to the vector of tips
+     */
+    vecmem::data::vector_view<const typename candidate_link::link_index_type>
+        tips_view;
+
+    /**
+     * @brief View object to the vector of track candidates
+     */
+    track_candidate_container_types::view track_candidates_view;
+
+    /**
+     * @brief View object to the vector of indices meeting the selection
+     * criteria
+     */
+    vecmem::data::vector_view<unsigned int> valid_indices_view;
+
+    /**
+     * @brief The number of valid tracks meeting criteria
+     */
+    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 +63,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..37d610ae9c 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,30 @@
 #include "traccc/edm/track_candidate.hpp"
 
 namespace traccc::device {
+struct fill_sort_keys_payload {
+    /**
+     * @brief View object to the vector of bound track parameters
+     */
+    bound_track_parameters_collection_types::const_view params_view;
+
+    /**
+     * @brief View object to the vector of sort keys
+     */
+    vecmem::data::vector_view<device::sort_key> keys_view;
+
+    /**
+     * @brief View object to the vector of parameter indices, which is the
+     * output to the algorithm
+     */
+    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..09dc3becbb 100644
--- a/device/common/include/traccc/finding/device/find_tracks.hpp
+++ b/device/common/include/traccc/finding/device/find_tracks.hpp
@@ -14,11 +14,114 @@
 #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 {
+    /**
+     * @brief View object to the tracking detector description
+     */
+    typename detector_t::view_type det_data;
+
+    /**
+     * @brief View object to the vector of bound track parameters
+     *
+     * @warning Measurements on the same surface must be adjacent
+     */
+    measurement_collection_types::const_view measurements_view;
+
+    /**
+     * @brief View object to the vector of track parameters
+     */
+    bound_track_parameters_collection_types::const_view in_params_view;
+
+    /**
+     * @brief View object to the vector of boolean-like integers describing the
+     * liveness of each parameter
+     */
+    vecmem::data::vector_view<const unsigned int> in_params_liveness_view;
+
+    /**
+     * @brief The total number of input parameters
+     */
+    const unsigned int n_in_params;
+
+    /**
+     * @brief View object to the vector of barcodes for each measurement
+     */
+    vecmem::data::vector_view<const detray::geometry::barcode> barcodes_view;
+
+    /**
+     * @brief View object to the vector of upper bounds of measurement indices
+     * per surface
+     */
+    vecmem::data::vector_view<const unsigned int> upper_bounds_view;
+
+    /**
+     * @brief View object to the link vector of the previous step
+     */
+    vecmem::data::vector_view<const candidate_link> prev_links_view;
+
+    /**
+     * @brief The current step identifier
+     */
+    const unsigned int step;
+
+    /**
+     * @brief The maximum number of new tracks to find
+     */
+    const unsigned int n_max_candidates;
+
+    /**
+     * @brief View object to the output track parameter vector
+     */
+    bound_track_parameters_collection_types::view out_params_view;
+
+    /**
+     * @brief View object to the output track parameter liveness vector
+     */
+    vecmem::data::vector_view<unsigned int> out_params_liveness_view;
+
+    /**
+     * @brief View object to the output candidate links
+     */
+    vecmem::data::vector_view<candidate_link> links_view;
+
+    /**
+     * @brief Pointer to the total of number of candidates; to be set to zero
+     * before launching the kernel
+     */
+    unsigned int* n_total_candidates;
+};
+
+struct find_tracks_shared_payload {
+    /**
+     * @brief Shared-memory vector with the number of measurements found per
+     * track
+     */
+    unsigned int* shared_num_candidates;
+
+    /**
+     * @brief Shared-memory vector of measurement candidats with ID and
+     * original track parameter identifier
+     *
+     * @note Length is always twice the block size
+     */
+    std::pair<unsigned int, unsigned int>* shared_candidates;
+
+    /**
+     * @brief Shared-memory atomic variable to track the size of
+     * \ref 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 +130,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 5db046851b..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 unsigned 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..149d9aa550 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 bd8a8c625e..23d207f32c 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;
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 789621f498..fe816bff8e 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);
 
@@ -202,10 +203,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,
@@ -225,7 +226,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();
@@ -241,15 +243,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;
             }
         }
     }
@@ -258,15 +262,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..f7d17c6ed2 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,24 @@
 #include "traccc/edm/measurement.hpp"
 
 namespace traccc::device {
+struct make_barcode_sequence_payload {
+    /**
+     * @brief View object to the vector of unique measurement indices
+     */
+    measurement_collection_types::const_view uniques_view;
+
+    /**
+     * @brief View object to the output vector of barcodes
+     */
+    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..e88639571d 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,64 @@
 #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 {
+    /**
+     * @brief View object to the tracking detector description
+     */
+    typename propagator_t::detector_type::view_type det_data;
+
+    /**
+     * @brief View object to the magnetic field
+     */
+    bfield_t field_data;
+
+    /**
+     * @brief View object to the vector of track parameters
+     */
+    bound_track_parameters_collection_types::view params_view;
+
+    /**
+     * @brief View object to the vector of track parameter liveness values
+     */
+    vecmem::data::vector_view<unsigned int> params_liveness_view;
+
+    /**
+     * @brief View object to the access order of parameters so they are sorted
+     */
+    const vecmem::data::vector_view<const unsigned int> param_ids_view;
+
+    /**
+     * @brief View object to the vector of candidate links
+     */
+    vecmem::data::vector_view<const candidate_link> links_view;
+
+    /**
+     * @brief Current CKF step number
+     */
+    const unsigned int step;
+
+    /**
+     * @brief Total number of input track parameters
+     */
+    const unsigned int n_in_params;
+
+    /**
+     * @brief View object to the vector of tips
+     */
+    vecmem::data::vector_view<typename candidate_link::link_index_type>
+        tips_view;
+
+    /**
+     * @brief View object to the vector of the number of tracks per initial
+     * input seed
+     */
+    vecmem::data::vector_view<unsigned int> n_tracks_per_seed_view;
+};
 
 /// Function for propagating the kalman-updated tracks to the next surface
 ///
@@ -24,33 +80,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..bc36925654 100644
--- a/device/common/include/traccc/finding/device/prune_tracks.hpp
+++ b/device/common/include/traccc/finding/device/prune_tracks.hpp
@@ -14,19 +14,29 @@
 
 namespace traccc::device {
 
+struct prune_tracks_payload {
+    /**
+     * @brief View object to the vector of track candidates
+     */
+    track_candidate_container_types::const_view track_candidates_view;
+
+    /**
+     * @brief View object to the vector containing the indices of valid tracks
+     */
+    vecmem::data::vector_view<const unsigned int> valid_indices_view;
+
+    /**
+     * @brief View object to the vector of pruned track candidates
+     */
+    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 feb85ed5d8..60a0848a86 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 503102c450..4fe8f760be 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 unsigned 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());
     }
@@ -313,10 +187,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());
         }
 
@@ -359,17 +233,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);
@@ -400,7 +275,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
@@ -426,13 +301,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_buffer,
-                        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_buffer,
+                                n_tracks_per_seed_buffer});
                 TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
 
                 m_stream.synchronize();
@@ -497,9 +372,9 @@ finding_algorithm<stepper_t, navigator_t>::operator()(
         const unsigned int nBlocks = (n_tips_total + nThreads - 1) / nThreads;
 
         kernels::build_tracks<<<nBlocks, nThreads, 0, stream>>>(
-            m_cfg, measurements, seeds_buffer, links_buffer, tips_buffer,
-            track_candidates_buffer, valid_indices_buffer,
-            n_valid_tracks_device.get());
+            m_cfg, {measurements, seeds_buffer, links_buffer, tips_buffer,
+                    track_candidates_buffer, valid_indices_buffer,
+                    n_valid_tracks_device.get()});
         TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
 
         // Global counter object: Device -> Host
@@ -525,8 +400,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..04ceefd305
--- /dev/null
+++ b/device/cuda/src/finding/kernels/apply_interaction.cuh
@@ -0,0 +1,20 @@
+/** 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);
+}  // 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..801bf118d6
--- /dev/null
+++ b/device/cuda/src/finding/kernels/build_tracks.cu
@@ -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
+ */
+
+#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,
+                             device::build_tracks_payload payload) {
+
+    int gid = threadIdx.x + blockIdx.x * blockDim.x;
+
+    device::build_tracks(gid, cfg, payload);
+}
+}  // 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..0cf59cd628
--- /dev/null
+++ b/device/cuda/src/finding/kernels/build_tracks.cuh
@@ -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
+ */
+
+#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/device/build_tracks.hpp"
+#include "traccc/finding/finding_config.hpp"
+
+namespace traccc::cuda::kernels {
+
+__global__ void build_tracks(const finding_config cfg,
+                             device::build_tracks_payload payload);
+}
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..9420e341f8
--- /dev/null
+++ b/device/cuda/src/finding/kernels/find_tracks.cuh
@@ -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
+ */
+
+#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);
+}  // 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..c5df625145
--- /dev/null
+++ b/device/cuda/src/finding/kernels/propagate_to_next_surface.cuh
@@ -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
+ */
+
+#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);
+}  // 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..1e9fbafc7a
--- /dev/null
+++ b/device/cuda/src/finding/kernels/specializations/find_tracks_src.cuh
@@ -0,0 +1,37 @@
+/** 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 "../propagate_to_next_surface.cuh"
+#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..2cb274b895
--- /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/actors/ckf_aborter.hpp"
+#include "traccc/finding/actors/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