ambiguity_resolution/kernels/fill_inverted_ids.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_inverted_ids.cuh"
0011
0012 // VecMem include(s).
0013 #include <vecmem/containers/device_vector.hpp>
0014
0015 namespace traccc::cuda::kernels {
0016
0017 __global__ void fill_inverted_ids(device::fill_inverted_ids_payload payload) {
0018   if (*(payload.terminate) == 1 || *(payload.n_updated_tracks) == 0) {
0019     return;
0020   }
0021
0022   vecmem::device_vector<const unsigned int> sorted_ids(payload.sorted_ids_view);
0023   vecmem::device_vector<unsigned int> inverted_ids(payload.inverted_ids_view);
0024
0025   auto globalIndex = threadIdx.x + blockIdx.x * blockDim.x;
0026   const unsigned int n_accepted = *(payload.n_accepted);
0027
0028   if (globalIndex >= n_accepted) {
0029     return;
0030   }
0031
0032   // Fill the inverted_ids vector which converts a track id to the index of
0033   // sorted ids
0034   inverted_ids[sorted_ids[globalIndex]] = globalIndex;
0035 }
0036
0037 }  // namespace traccc::cuda::kernels