ambiguity_resolution/kernels/fill_tracks_per_measurement.cu

0001 /** TRACCC library, part of the ACTS project (R&D line)
0002  *
0003  * (c) 2025 CERN for the benefit of the ACTS project
0004  *
0005  * Mozilla Public License Version 2.0
0006  */
0007
0008 // Local include(s).
0009 #include "../../utils/global_index.hpp"
0010 #include "fill_tracks_per_measurement.cuh"
0011
0012 // VecMem include(s).
0013 #include <vecmem/containers/device_vector.hpp>
0014 #include <vecmem/containers/jagged_device_vector.hpp>
0015
0016 // Thrust include(s).
0017 #include <thrust/binary_search.h>
0018 #include <thrust/execution_policy.h>
0019 #include <thrust/find.h>
0020
0021 namespace traccc::cuda::kernels {
0022
0023 __global__ void fill_tracks_per_measurement(
0024     device::fill_tracks_per_measurement_payload payload) {
0025   vecmem::device_vector<const unsigned int> accepted_ids(
0026       payload.accepted_ids_view);
0027
0028   const auto globalIndex = details::global_index1();
0029   if (globalIndex >= accepted_ids.size()) {
0030     return;
0031   }
0032
0033   vecmem::jagged_device_vector<const measurement_id_type> meas_ids(
0034       payload.meas_ids_view);
0035   vecmem::device_vector<const unsigned int> meas_id_to_unique_id(
0036       payload.meas_id_to_unique_id_view);
0037   vecmem::jagged_device_vector<unsigned int> tracks_per_measurement(
0038       payload.tracks_per_measurement_view);
0039   vecmem::jagged_device_vector<int> track_status_per_measurement(
0040       payload.track_status_per_measurement_view);
0041   vecmem::device_vector<unsigned int> n_accepted_tracks_per_measurement(
0042       payload.n_accepted_tracks_per_measurement_view);
0043
0044   const unsigned int id = accepted_ids.at(globalIndex);
0045
0046   for (unsigned int i = 0; i < meas_ids[id].size(); i++) {
0047     auto meas_id = meas_ids[id][i];
0048
0049     if (thrust::find(thrust::seq, meas_ids[id].begin(),
0050                      meas_ids[id].begin() + i,
0051                      meas_id) != (meas_ids[id].begin() + i)) {
0052       continue;
0053     }
0054
0055     const auto unique_meas_idx = meas_id_to_unique_id.at(meas_id);
0056
0057     auto tracks = tracks_per_measurement.at(unique_meas_idx);
0058
0059     tracks_per_measurement.at(unique_meas_idx).push_back(id);
0060     track_status_per_measurement.at(unique_meas_idx).push_back(1);
0061
0062     vecmem::device_atomic_ref<unsigned int> n_accepted(
0063         n_accepted_tracks_per_measurement.at(
0064             static_cast<unsigned int>(unique_meas_idx)));
0065     n_accepted.fetch_add(1u);
0066   }
0067 }
0068 }  // namespace traccc::cuda::kernels