ambiguity_resolution/kernels/update_status.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 "update_status.cuh"
0011
0012 // VecMem include(s).
0013 #include <vecmem/containers/device_vector.hpp>
0014
0015 namespace traccc::cuda::kernels {
0016
0017 __global__ void update_status(device::update_status_payload payload) {
0018   if (*(payload.terminate) == 1) {
0019     return;
0020   }
0021
0022   vecmem::device_vector<const unsigned int> temp_sorted_ids(
0023       payload.temp_sorted_ids_view);
0024   vecmem::device_vector<unsigned int> sorted_ids(payload.sorted_ids_view);
0025   vecmem::device_vector<unsigned int> updated_tracks(
0026       payload.updated_tracks_view);
0027   vecmem::device_vector<int> is_updated(payload.is_updated_view);
0028   vecmem::device_vector<const unsigned int> n_shared(payload.n_shared_view);
0029
0030   auto globalIndex = threadIdx.x + blockIdx.x * blockDim.x;
0031   const unsigned int n_accepted = *(payload.n_accepted);
0032   const unsigned int n_updated = *(payload.n_updated_tracks);
0033
0034   /***********************
0035    * Update Max Shared
0036    ***********************/
0037
0038   unsigned int shared = 0;
0039
0040   if (globalIndex < n_accepted) {
0041     auto tid = sorted_ids[globalIndex];
0042     shared = n_shared[tid];
0043   }
0044
0045   for (int offset = 16; offset > 0; offset >>= 1) {
0046     unsigned int other_shared = __shfl_down_sync(0xffffffff, shared, offset);
0047
0048     if (other_shared > shared) {
0049       shared = other_shared;
0050     }
0051   }
0052
0053   if (threadIdx.x == 0) {
0054     atomicMax(payload.max_shared, shared);
0055   }
0056
0057   /***********************
0058    * Update Sorted Ids
0059    ***********************/
0060
0061   if (n_updated == 0) {
0062     return;
0063   }
0064
0065   // Reset is_updated vector
0066   if (globalIndex < n_updated) {
0067     is_updated[updated_tracks[globalIndex]] = 0;
0068   }
0069
0070   if (globalIndex < n_accepted) {
0071     auto tid = temp_sorted_ids[globalIndex];
0072     sorted_ids[globalIndex] = tid;
0073   }
0074 }
0075
0076 }  // namespace traccc::cuda::kernels