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

.github/copilot-instructions.md

@@ -0,0 +1,80 @@
+This project is a heterogeneous demonstrator for tracking algorithms (CPU, CUDA,
+SYCL, Alpaka, Kokkos, Futhark).
+
+Quick context for an AI coding agent
+- Big picture: `core/` contains the event-data-model, code shared by all
+  architectures, and the host algorithms. Device-specific implementations
+  live under `device/` in subfolders `cuda/`, `sycl/`, `alpaka/`, etc.
+  Examples and executables are built under `examples/` and top-level CMake
+  presets produce binaries in the build's `bin/` directory.
+- Primary build system: CMake (C++20). See `CMakeLists.txt`, `CMakePresets.json`
+  and `cmake/` for presets and helper macros. Important per-module
+  CMake: `core/CMakeLists.txt`, `device/*/CMakeLists.txt`, `extern/` for
+  external dependencies.
+
+What to change and where (concrete patterns)
+- New device specific algorithms are to be put into `device/<backend>/src/` with
+  headers in `device/<backend>/include/traccc/<backend>/`. Follow existing
+  patterns for kernel launches and host-device separation. See
+  `device/cuda/src/clusterization/` for examples.
+
+Build & test workflows (exact commands)
+- Common build (from repo root):
+  cmake --preset <preset> -S traccc/ -B build
+  cmake --build build/
+  Common presets: `host-fp32`, `cuda-fp32`, `sycl-fp32`, `alpaka-fp32` (see
+  `README.md`, build presets and `CMakePresets.json`).
+- Run unit tests (if built): use CTest from the build directory:
+  `ctest --output-on-failure -C <config>`. The project uses GoogleTest;
+  tests are controlled by `TRACCC_BUILD_TESTING`.
+
+Project-specific conventions and gotchas
+- C++ standard: project requires C++20 for host and device code. Many CMake
+  checks will fail early if not satisfied. See `CMakeLists.txt` for explicit
+  checks.
+- Templated inline code: `.ipp` files contain template implementations. Keep
+  header/impl separation consistent to avoid ODR/link errors.
+- VecMem/Eigen/Thrust/Acts and other dependencies are vendored under `extern/`
+  but can be toggled with `TRACCC_USE_SYSTEM_*` flags. Prefer following the repo
+  convention: use vendored versions unless the developer wants to link
+  system/Spack libs.
+
+Where to look for architecture decisions and data flow
+- High-level diagram and narratives: `README.md` (top-level) includes a mermaid
+  dataflow graph linking clusterization -> measurements -> spacepoints -> seeding -> finding -> fitting.
+- Core algorithm grouping and examples: `core/CMakeLists.txt` lists the major
+  components and files that implement clusterization, seeding, finding, and
+  fitting — good entry points when adding or modifying an algorithm.
+- Device backends: `device/` mirrors core algorithms for each backend. Use these
+  folders to understand how host/device separation is organized.
+
+Testing, benchmarking and performance
+- Examples for running simulated data are in the README
+  (e.g. `traccc_simulate_telescope` and `traccc_throughput_mt`).
+- For profiling builds, the repo includes support for collecting build profiling
+  logs when CTest launchers are enabled (see `CMakeLists.txt` section that
+  configures `traccc-ctest.sh`).
+
+Editing guidance for AI agents (practical rules)
+- Preserve header/implementation split; add `.ipp` files for template
+  implementations. Match namespaces: `traccc::...` and directory
+  structure -> namespace mapping.
+- Add compile-time feature guards using the module's CMake options when touching
+  device/backends (e.g. `TRACCC_BUILD_CUDA`). Update `CMakeLists.txt` entries if
+  new files are added to a target.
+- When adding new public APIs, update `core/CMakeLists.txt` or relevant module
+  CMake to include headers and sources in the `traccc_add_library`/target
+  definitions.
+- Keep changes minimal and provide a small test or example in `examples/` or add
+  a unit test under `tests/` to exercise the new code.
+- Add copyright notices to new files following existing patterns.
+
+Files to reference during development
+- Top-level: `README.md`, `CMakeLists.txt`, `CMakePresets.json`
+- Core: `core/CMakeLists.txt`, `core/include/traccc/edm/`, `core/include/traccc/seeding/`, `core/include/traccc/finding/`, `core/include/traccc/fitting/`
+- Device: `device/cuda/`, `device/sycl/`, `device/alpaka/` (look at `include/traccc/<backend>` and `src/` for wrappers)
+- Vendored deps and glue: `extern/`, `cmake/` helpers (e.g. `traccc-functions.cmake`)
+
+If anything is unclear or you'd like more detail (examples of header/ipp
+patterns, a unit test scaffold, or preferred presets for CI), ask and I'll
+extend this file with short examples.

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

@@ -75,13 +75,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