Introduced device::update_tip_length_buffer.

While modifying cuda::kernels::update_tip_length_buffer
to make use of that common function, and updating
device::build_tracks to be able to collaborate with the
slightly different data that update_tip_length_buffer is
now producing.

Author: krasznaa

device/alpaka/src/finding/combinatorial_kalman_filter.hpp

@@ -1,6 +1,6 @@
 /** TRACCC library, part of the ACTS project (R&D line)
  *
- * (c) 2024-2025 CERN for the benefit of the ACTS project
+ * (c) 2024-2026 CERN for the benefit of the ACTS project
  *
  * Mozilla Public License Version 2.0
  */
@@ -643,7 +643,7 @@ combinatorial_kalman_filter(
                 .links_view = links_buffer,
                 .tips_view = tips_buffer,
                 .tracks_view = track_candidates_buffer,
-                .tip_to_output_map = nullptr,
+                .tip_to_output_map = {},
                 .jacobian_ptr = jacobian_ptr.get(),
                 .link_predicted_parameter_view =
                     link_predicted_parameter_buffer,

device/common/CMakeLists.txt

@@ -77,13 +77,15 @@ traccc_add_library( traccc_device_common device_common
    "include/traccc/finding/device/gather_measurement_votes.hpp"
    "include/traccc/finding/device/fill_finding_propagation_sort_keys.hpp"
    "include/traccc/finding/device/propagate_to_next_surface.hpp"
+   "include/traccc/finding/device/update_tip_length_buffer.hpp"
    "include/traccc/finding/device/impl/apply_interaction.ipp"
    "include/traccc/finding/device/impl/build_tracks.ipp"
    "include/traccc/finding/device/impl/find_tracks.ipp"
    "include/traccc/finding/device/impl/gather_best_tips_per_measurement.ipp"
    "include/traccc/finding/device/impl/gather_measurement_votes.ipp"
    "include/traccc/finding/device/impl/fill_finding_propagation_sort_keys.ipp"
    "include/traccc/finding/device/impl/propagate_to_next_surface.ipp"
+   "include/traccc/finding/device/impl/update_tip_length_buffer.ipp"
    # Track fitting funtions(s).
    "include/traccc/fitting/device/fit.hpp"
    "include/traccc/fitting/device/fill_fitting_sort_keys.hpp"

device/common/include/traccc/finding/device/build_tracks.hpp

@@ -48,7 +48,7 @@ struct build_tracks_payload {
     /**
      * @brief Optional mapping from tip index to output index
      */
-    const unsigned int* tip_to_output_map = nullptr;
+    vecmem::data::vector_view<const unsigned int> tip_to_output_map;
     bound_matrix<default_algebra>* jacobian_ptr = nullptr;
     bound_track_parameters_collection_types::view link_predicted_parameter_view;
     bound_track_parameters_collection_types::view link_filtered_parameter_view;

device/common/include/traccc/finding/device/impl/build_tracks.ipp

@@ -45,8 +45,10 @@ TRACCC_HOST_DEVICE inline void build_tracks(
         return;
     }
 
-    const unsigned int output_idx = payload.tip_to_output_map != nullptr
-                                        ? payload.tip_to_output_map[globalIndex]
+    const vecmem::device_vector<const unsigned int> tip_to_output_map(
+        payload.tip_to_output_map);
+    const unsigned int output_idx = tip_to_output_map.size()
+                                        ? tip_to_output_map.at(globalIndex)
                                         : globalIndex;
 
     if (output_idx == std::numeric_limits<unsigned int>::max()) {

device/common/include/traccc/finding/device/impl/update_tip_length_buffer.ipp

@@ -0,0 +1,54 @@
+/** traccc library, part of the ACTS project (R&D line)
+ *
+ * (c) 2025-2026 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+// VecMem include(s).
+#include <vecmem/containers/device_vector.hpp>
+#include <vecmem/memory/device_atomic_ref.hpp>
+
+// System include(s).
+#include <cassert>
+#include <limits>
+
+namespace traccc::device {
+
+TRACCC_HOST_DEVICE inline void update_tip_length_buffer(
+    global_index_t thread_id, const update_tip_length_buffer_payload& payload) {
+
+    const vecmem::device_vector<const unsigned int> measurement_votes(
+        payload.measurement_votes);
+    if (thread_id >= measurement_votes.size()) {
+        return;
+    }
+
+    const vecmem::device_vector<const unsigned int> old_tip_length(
+        payload.old_tip_length);
+    vecmem::device_vector<unsigned int> new_tip_length(payload.new_tip_length);
+
+    const unsigned int total_measurements = old_tip_length.at(thread_id);
+    const unsigned int total_votes = measurement_votes.at(thread_id);
+
+    assert(total_votes <= total_measurements);
+
+    const float vote_fraction = static_cast<float>(total_votes) /
+                                static_cast<float>(total_measurements);
+
+    vecmem::device_vector<unsigned int> tip_to_output_map(
+        payload.tip_to_output_map);
+
+    if (vote_fraction < payload.min_measurement_voting_fraction) {
+        tip_to_output_map.at(thread_id) =
+            std::numeric_limits<unsigned int>::max();
+    } else {
+        const vecmem::device_vector<unsigned int>::size_type new_idx =
+            new_tip_length.push_back(total_measurements);
+        tip_to_output_map.at(thread_id) = new_idx;
+    }
+}
+
+}  // namespace traccc::device

device/common/include/traccc/finding/device/update_tip_length_buffer.hpp

@@ -0,0 +1,36 @@
+/** traccc library, part of the ACTS project (R&D line)
+ *
+ * (c) 2025-2026 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+// Local include(s).
+#include "traccc/device/global_index.hpp"
+
+// Project include(s).
+#include "traccc/definitions/qualifiers.hpp"
+
+// VecMem include(s).
+#include <vecmem/containers/data/vector_view.hpp>
+
+namespace traccc::device {
+
+/// Payload for the @c device::update_tip_length_buffer function
+struct update_tip_length_buffer_payload {
+    vecmem::data::vector_view<const unsigned int> old_tip_length;
+    vecmem::data::vector_view<unsigned int> new_tip_length;
+    vecmem::data::vector_view<const unsigned int> measurement_votes;
+    vecmem::data::vector_view<unsigned int> tip_to_output_map;
+    float min_measurement_voting_fraction;
+};
+
+TRACCC_HOST_DEVICE inline void update_tip_length_buffer(
+    global_index_t thread_id, const update_tip_length_buffer_payload& payload);
+
+}  // namespace traccc::device
+
+// Include the implementation.
+#include "traccc/finding/device/impl/update_tip_length_buffer.ipp"

device/cuda/src/finding/combinatorial_kalman_filter.cuh

@@ -569,7 +569,7 @@ combinatorial_kalman_filter(
     }
 
     vecmem::vector<unsigned int> tips_length_host(mr.host);
-    vecmem::unique_alloc_ptr<unsigned int[]> tip_to_output_map = nullptr;
+    vecmem::data::vector_buffer<unsigned int> tip_to_output_map;
 
     unsigned int n_tips_total_filtered = n_tips_total;
 
@@ -640,40 +640,36 @@ combinatorial_kalman_filter(
         }
 
         tip_to_output_map =
-            vecmem::make_unique_alloc<unsigned int[]>(mr.main, n_tips_total);
+            vecmem::data::vector_buffer<unsigned int>(n_tips_total, mr.main);
+        copy.setup(tip_to_output_map)->wait();
 
         {
             const unsigned int num_threads = 512;
             const unsigned int num_blocks =
                 (n_tips_total + num_threads - 1) / num_threads;
 
             vecmem::data::vector_buffer<unsigned int> new_tip_length_buffer{
-                n_tips_total, mr.main};
+                n_tips_total, mr.main, vecmem::data::buffer_type::resizable};
             copy.setup(new_tip_length_buffer)->wait();
 
-            auto tip_to_output_map_idx =
-                vecmem::make_unique_alloc<unsigned int>(mr.main);
-
-            TRACCC_CUDA_ERROR_CHECK(cudaMemsetAsync(
-                tip_to_output_map_idx.get(), 0, sizeof(unsigned int), stream));
-
             kernels::update_tip_length_buffer<<<num_blocks, num_threads, 0,
                                                 stream>>>(
-                tip_length_buffer, new_tip_length_buffer, votes_per_tip_buffer,
-                tip_to_output_map.get(), tip_to_output_map_idx.get(),
-                config.min_measurement_voting_fraction);
+                {tip_length_buffer, new_tip_length_buffer, votes_per_tip_buffer,
+                 tip_to_output_map, config.min_measurement_voting_fraction});
 
             TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
 
-            str.synchronize();
-
-            TRACCC_CUDA_ERROR_CHECK(cudaMemcpyAsync(
-                &n_tips_total_filtered, tip_to_output_map_idx.get(),
-                sizeof(unsigned int), cudaMemcpyDeviceToHost, stream));
+            if (mr.host) {
+                vecmem::async_size size =
+                    copy.get_size(tip_to_output_map, *(mr.host));
+                // Here we could give control back to the caller, once our code
+                // allows for it. (coroutines...)
+                n_tips_total_filtered = size.get();
+            } else {
+                n_tips_total_filtered = copy.get_size(tip_to_output_map);
+            }
 
             tip_length_buffer = std::move(new_tip_length_buffer);
-
-            str.synchronize();
         }
     }
 
@@ -709,7 +705,7 @@ combinatorial_kalman_filter(
             .links_view = links_buffer,
             .tips_view = tips_buffer,
             .tracks_view = {track_candidates_buffer},
-            .tip_to_output_map = tip_to_output_map.get(),
+            .tip_to_output_map = tip_to_output_map,
             .jacobian_ptr = jacobian_ptr.get(),
             .link_predicted_parameter_view = link_predicted_parameter_buffer,
             .link_filtered_parameter_view = link_filtered_parameter_buffer,

device/cuda/src/finding/kernels/update_tip_length_buffer.cu

@@ -1,52 +1,20 @@
 /** traccc library, part of the ACTS project (R&D line)
  *
- * (c) 2025 CERN for the benefit of the ACTS project
+ * (c) 2025-2026 CERN for the benefit of the ACTS project
  *
  * Mozilla Public License Version 2.0
  */
 
+// Local include(s).
+#include "../../utils/global_index.hpp"
 #include "update_tip_length_buffer.cuh"
 
 namespace traccc::cuda::kernels {
 
 __global__ void update_tip_length_buffer(
-    const vecmem::data::vector_view<const unsigned int> old_tip_length_view,
-    vecmem::data::vector_view<unsigned int> new_tip_length_view,
-    const vecmem::data::vector_view<const unsigned int> measurement_votes_view,
-    unsigned int* tip_to_output_map, unsigned int* tip_to_output_map_idx,
-    float min_measurement_voting_fraction) {
-    const unsigned int thread_idx = blockIdx.x * blockDim.x + threadIdx.x;
+    const device::update_tip_length_buffer_payload payload) {
 
-    assert(tip_to_output_map != nullptr);
-    assert(tip_to_output_map_idx != nullptr);
-
-    const vecmem::device_vector<const unsigned int> old_tip_length(
-        old_tip_length_view);
-    vecmem::device_vector<unsigned int> new_tip_length(new_tip_length_view);
-    const vecmem::device_vector<const unsigned int> measurement_votes(
-        measurement_votes_view);
-
-    if (thread_idx >= measurement_votes_view.size()) {
-        return;
-    }
-
-    const unsigned int total_measurements = old_tip_length.at(thread_idx);
-    const unsigned int total_votes = measurement_votes.at(thread_idx);
-
-    assert(total_votes <= total_measurements);
-
-    const scalar vote_fraction = static_cast<scalar>(total_votes) /
-                                 static_cast<scalar>(total_measurements);
-
-    if (vote_fraction < min_measurement_voting_fraction) {
-        tip_to_output_map[thread_idx] =
-            std::numeric_limits<unsigned int>::max();
-    } else {
-        const auto new_idx =
-            vecmem::device_atomic_ref(*tip_to_output_map_idx).fetch_add(1u);
-        new_tip_length.at(new_idx) = total_measurements;
-        tip_to_output_map[thread_idx] = new_idx;
-    }
+    device::update_tip_length_buffer(details::global_index1(), payload);
 }
 
 }  // namespace traccc::cuda::kernels

device/cuda/src/finding/kernels/update_tip_length_buffer.cuh

@@ -1,27 +1,18 @@
 /** traccc library, part of the ACTS project (R&D line)
  *
- * (c) 2025 CERN for the benefit of the ACTS project
+ * (c) 2025-2026 CERN for the benefit of the ACTS project
  *
  * Mozilla Public License Version 2.0
  */
 
-#include <vecmem/containers/device_vector.hpp>
-#include <vecmem/containers/vector.hpp>
+#pragma once
 
-#include "traccc/definitions/common.hpp"
-#include "traccc/device/array_insertion_mutex.hpp"
-#include "traccc/edm/measurement_collection.hpp"
-#include "traccc/finding/candidate_link.hpp"
-#include "traccc/finding/finding_config.hpp"
-#include "traccc/utils/prob.hpp"
+// Project include(s).
+#include "traccc/finding/device/update_tip_length_buffer.hpp"
 
 namespace traccc::cuda::kernels {
 
 __global__ void update_tip_length_buffer(
-    const vecmem::data::vector_view<const unsigned int> old_tip_length_view,
-    vecmem::data::vector_view<unsigned int> new_tip_length_view,
-    const vecmem::data::vector_view<const unsigned int> measurement_votes_view,
-    unsigned int* tip_to_output_map, unsigned int* tip_to_output_map_idx,
-    float min_measurement_voting_fraction);
+    const device::update_tip_length_buffer_payload payload);
 
 }  // namespace traccc::cuda::kernels