Cut by max_num_skipping_per_cand as soon as possible

fiedlerp · stephenswat · commit 3243e71c6614 · 2025-05-13T15:57:19.000+02:00
diff --git a/device/alpaka/src/finding/finding_algorithm.cpp b/device/alpaka/src/finding/finding_algorithm.cpp
@@ -276,6 +276,7 @@ finding_algorithm<stepper_t, navigator_t>::operator()(
                             .step = step,
                             .out_params_view = updated_params_buffer,
                             .out_params_liveness_view = updated_liveness_buffer,
+                            .tips_view = tips_buffer,
                             .n_tracks_per_seed_view = n_tracks_per_seed_buffer};
 
             auto bufAcc_payload =
@@ -355,7 +356,6 @@ finding_algorithm<stepper_t, navigator_t>::operator()(
                                 .params_view = in_params_buffer,
                                 .params_liveness_view = param_liveness_buffer,
                                 .param_ids_view = param_ids_buffer,
-                                .links_view = links_buffer,
                                 .prev_links_idx = step_to_link_idx_map[step],
                                 .step = step,
                                 .n_in_params = n_candidates,
diff --git a/device/common/include/traccc/finding/device/find_tracks.hpp b/device/common/include/traccc/finding/device/find_tracks.hpp
@@ -99,6 +99,11 @@ struct find_tracks_payload {
      */
     vecmem::data::vector_view<unsigned int> out_params_liveness_view;
 
+    /**
+     * @brief View object to the vector of tips
+     */
+    vecmem::data::vector_view<unsigned int> tips_view;
+
     /**
      * @brief View object to the vector of the number of tracks per initial
      * input seed
diff --git a/device/common/include/traccc/finding/device/impl/find_tracks.ipp b/device/common/include/traccc/finding/device/impl/find_tracks.ipp
@@ -71,6 +71,7 @@ TRACCC_HOST_DEVICE inline void find_tracks(
         payload.barcodes_view);
     vecmem::device_vector<const unsigned int> upper_bounds(
         payload.upper_bounds_view);
+    vecmem::device_vector<unsigned int> tips(payload.tips_view);
     vecmem::device_vector<unsigned int> n_tracks_per_seed(
         payload.n_tracks_per_seed_view);
 
@@ -308,25 +309,31 @@ TRACCC_HOST_DEVICE inline void find_tracks(
         const unsigned int s_pos = num_tracks_per_seed.fetch_add(1);
 
         if (s_pos < cfg.max_num_branches_per_seed) {
-            // Add measurement candidates to link
-            const unsigned int l_pos = links.bulk_append_implicit(1);
-
             const unsigned int n_skipped =
                 payload.step == 0 ? 0 : links.at(prev_link_idx).n_skipped;
 
-            links.at(l_pos) = {
-                .step = payload.step,
-                .previous_candidate_idx = prev_link_idx,
-                .meas_idx = std::numeric_limits<unsigned int>::max(),
-                .seed_idx = seed_idx,
-                .n_skipped = n_skipped + 1,
-                .chi2 = std::numeric_limits<traccc::scalar>::max()};
-
-            out_params.at(l_pos - payload.curr_links_idx) =
-                in_params.at(in_param_id);
-            out_params_liveness.at(l_pos - payload.curr_links_idx) = 1u;
+            if (n_skipped >= cfg.max_num_skipping_per_cand) {
+                // In case of max skipping being 0 and first step being skipped,
+                // the links are empty, and the tip has nowhere to point
+                assert(payload.step > 0);
+                tips.push_back(prev_link_idx);
+            } else {
+                // Add measurement candidates to link
+                const unsigned int l_pos = links.bulk_append_implicit(1);
+
+                links.at(l_pos) = {
+                    .step = payload.step,
+                    .previous_candidate_idx = prev_link_idx,
+                    .meas_idx = std::numeric_limits<unsigned int>::max(),
+                    .seed_idx = seed_idx,
+                    .n_skipped = n_skipped + 1,
+                    .chi2 = std::numeric_limits<traccc::scalar>::max()};
+
+                out_params.at(l_pos - payload.curr_links_idx) =
+                    in_params.at(in_param_id);
+                out_params_liveness.at(l_pos - payload.curr_links_idx) = 1u;
+            }
         }
     }
 }
-
 }  // 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
@@ -32,23 +32,13 @@ TRACCC_HOST_DEVICE inline void propagate_to_next_surface(
 
     const unsigned int param_id = param_ids.at(globalIndex);
 
-    // Links
-    vecmem::device_vector<const candidate_link> links(payload.links_view);
-
     // Parameter liveness
     vecmem::device_vector<unsigned int> params_liveness(
         payload.params_liveness_view);
 
     // tips
     vecmem::device_vector<unsigned int> tips(payload.tips_view);
 
-    if (links.at(payload.prev_links_idx + param_id).n_skipped >
-        cfg.max_num_skipping_per_cand) {
-        params_liveness[param_id] = 0u;
-        tips.push_back(payload.prev_links_idx + param_id);
-        return;
-    }
-
     // Detector
     typename propagator_t::detector_type det(payload.det_data);
 
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
@@ -50,11 +50,6 @@ struct propagate_to_next_surface_payload {
      */
     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 Index in the link vector at which the current step starts
      */
diff --git a/device/cuda/src/finding/finding_algorithm.cu b/device/cuda/src/finding/finding_algorithm.cu
@@ -282,6 +282,7 @@ finding_algorithm<stepper_t, navigator_t>::operator()(
                     .step = step,
                     .out_params_view = updated_params_buffer,
                     .out_params_liveness_view = updated_liveness_buffer,
+                    .tips_view = tips_buffer,
                     .n_tracks_per_seed_view = n_tracks_per_seed_buffer});
             TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
 
@@ -352,7 +353,6 @@ finding_algorithm<stepper_t, navigator_t>::operator()(
                         .params_view = in_params_buffer,
                         .params_liveness_view = param_liveness_buffer,
                         .param_ids_view = param_ids_buffer,
-                        .links_view = links_buffer,
                         .prev_links_idx = step_to_link_idx_map[step],
                         .step = step,
                         .n_in_params = n_candidates,
diff --git a/device/sycl/src/finding/find_tracks.hpp b/device/sycl/src/finding/find_tracks.hpp
@@ -290,6 +290,7 @@ track_candidate_container_types::buffer find_tracks(
                      updated_params = vecmem::get_data(updated_params_buffer),
                      updated_liveness =
                          vecmem::get_data(updated_liveness_buffer),
+                     tips = vecmem::get_data(tips_buffer),
                      n_tracks_per_seed =
                          vecmem::get_data(n_tracks_per_seed_buffer),
                      shared_candidates_size, shared_num_candidates,
@@ -305,7 +306,7 @@ track_candidate_container_types::buffer find_tracks(
                             {det, measurements, in_params, param_liveness,
                              n_in_params, barcodes, upper_bounds, links_view,
                              prev_links_idx, curr_links_idx, step,
-                             updated_params, updated_liveness,
+                             updated_params, updated_liveness, tips,
                              n_tracks_per_seed},
                             {&(shared_num_candidates[0]),
                              &(shared_candidates[0]),
@@ -392,7 +393,6 @@ track_candidate_container_types::buffer find_tracks(
                          param_liveness =
                              vecmem::get_data(param_liveness_buffer),
                          param_ids = vecmem::get_data(param_ids_buffer),
-                         links_view = vecmem::get_data(links_buffer),
                          prev_links_idx = step_to_link_idx_map[step], step,
                          n_candidates, tips = vecmem::get_data(tips_buffer)](
                             ::sycl::nd_item<1> item) {
@@ -401,8 +401,8 @@ track_candidate_container_types::buffer find_tracks(
                                 typename stepper_t::magnetic_field_type>(
                                 details::global_index(item), config,
                                 {det, field, in_params, param_liveness,
-                                 param_ids, links_view, prev_links_idx, step,
-                                 n_candidates, tips});
+                                 param_ids, prev_links_idx, step, n_candidates,
+                                 tips});
                         });
                 })
                 .wait_and_throw();
