Skip the holes for GPU CKF

core/CMakeLists.txt

@@ -58,7 +58,9 @@ traccc_add_library( traccc_core core TYPE SHARED
   "src/clusterization/clusterization_algorithm.cpp"
   # Finding algorithmic code
   "include/traccc/finding/candidate_link.hpp"
+  "include/traccc/finding/ckf_aborter.hpp"
   "include/traccc/finding/finding_algorithm.hpp"
+  "include/traccc/finding/finding_algorithm.ipp"
   "include/traccc/finding/finding_config.hpp"
   "include/traccc/finding/interaction_register.hpp"
   # Fitting algorithmic code

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;

core/include/traccc/finding/ckf_aborter.hpp

@@ -0,0 +1,54 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 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.5f;
+        /// Maximum step counts that track can make to reach the next surface
+        unsigned int max_count = 100;
+
+        bool success = false;
+        unsigned int count = 0;
+    };
+
+    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>;