Skip holes and continue propagation for GPU CKF

core/include/traccc/finding/candidate_link.hpp

@@ -20,7 +20,7 @@ namespace traccc {
 struct candidate_link {
 
     // Type of index
-    using link_index_type = thrust::pair<unsigned int, unsigned int>;
+    using link_index_type = thrust::pair<int, unsigned int>;
 
     // Index of link from the previous step
     link_index_type previous;
@@ -30,9 +30,6 @@ struct candidate_link {
 
     // Index to the initial seed
     unsigned int seed_idx;
-
-    // How many times it skipped a surface
-    unsigned int n_skipped;
 };
 
 }  // namespace traccc

core/include/traccc/finding/ckf_aborter.hpp

@@ -0,0 +1,53 @@
+/** 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
+
+// Project include(s)
+#include "detray/definitions/detail/qualifiers.hpp"
+#include "detray/propagator/base_actor.hpp"
+#include "detray/propagator/base_stepper.hpp"
+
+// System include(s)
+#include <limits>
+
+namespace traccc {
+
+/// Aborter triggered when the next surface is reached
+struct ckf_aborter : detray::actor {
+    struct state {
+        // minimal step length to prevent from staying on the same surface
+        scalar min_step_length = 0.f;
+        bool success = false;
+
+        unsigned int count = 0;
+        unsigned int max_count = 100;
+    };
+
+    template <typename propagator_state_t>
+    DETRAY_HOST_DEVICE void operator()(state &abrt_state,
+                                       propagator_state_t &prop_state) const {
+
+        auto &navigation = prop_state._navigation;
+        auto &stepping = prop_state._stepping;
+
+        abrt_state.count++;
+
+        // Abort at the next sensitive surface
+        if (navigation.is_on_sensitive() &&
+            stepping._s > abrt_state.min_step_length) {
+            prop_state._heartbeat &= navigation.abort();
+            abrt_state.success = true;
+        }
+
+        if (abrt_state.count > abrt_state.max_count) {
+            prop_state._heartbeat &= navigation.abort();
+        }
+    }
+};
+
+}  // namespace traccc

core/include/traccc/finding/finding_algorithm.hpp

@@ -12,6 +12,7 @@
 #include "traccc/edm/measurement.hpp"
 #include "traccc/edm/track_candidate.hpp"
 #include "traccc/edm/track_state.hpp"
+#include "traccc/finding/ckf_aborter.hpp"
 #include "traccc/finding/finding_config.hpp"
 #include "traccc/finding/interaction_register.hpp"
 #include "traccc/fitting/kalman_filter/gain_matrix_updater.hpp"
@@ -62,7 +63,7 @@ class finding_algorithm
     using actor_type =
         detray::actor_chain<std::tuple, detray::pathlimit_aborter, transporter,
                             interaction_register<interactor>, interactor,
-                            detray::next_surface_aborter>;
+                            ckf_aborter>;
 
     using propagator_type =
         detray::propagator<stepper_t, navigator_t, actor_type>;