ambiguity_resolution/kernels/count_shared_measurements.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 "count_shared_measurements.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
0020 namespace traccc::cuda::kernels {
0021
0022 __global__ void count_shared_measurements(
0023     device::count_shared_measurements_payload payload) {
0024   vecmem::device_vector<const unsigned int> accepted_ids(
0025       payload.accepted_ids_view);
0026
0027   auto globalIndex = threadIdx.x + blockIdx.x * blockDim.x;
0028
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::device_vector<const unsigned int> n_accepted_tracks_per_measurement(
0038       payload.n_accepted_tracks_per_measurement_view);
0039   vecmem::device_vector<unsigned int> n_shared(payload.n_shared_view);
0040
0041   const unsigned int id = accepted_ids.at(globalIndex);
0042
0043   for (const auto& meas_id : meas_ids[id]) {
0044     const auto unique_meas_idx = meas_id_to_unique_id.at(meas_id);
0045
0046     if (n_accepted_tracks_per_measurement.at(unique_meas_idx) > 1) {
0047       vecmem::device_atomic_ref<unsigned int>(n_shared.at(id)).fetch_add(1u);
0048     }
0049   }
0050 }
0051 }  // namespace traccc::cuda::kernels