ambiguity_resolution/kernels/add_block_offset.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/barrier.hpp"
0010 #include "../../utils/global_index.hpp"
0011 #include "add_block_offset.cuh"
0012
0013 // VecMem include(s).
0014 #include <vecmem/containers/device_vector.hpp>
0015
0016 namespace traccc::cuda::kernels {
0017
0018 __global__ void add_block_offset(device::add_block_offset_payload payload) {
0019   if (*(payload.terminate) == 1 || *(payload.n_updated_tracks) == 0) {
0020     return;
0021   }
0022
0023   vecmem::device_vector<const int> block_offsets(payload.block_offsets_view);
0024   vecmem::device_vector<int> prefix_sums(payload.prefix_sums_view);
0025
0026   auto globalIndex = threadIdx.x + blockIdx.x * blockDim.x;
0027   const unsigned int n_accepted = *(payload.n_accepted);
0028
0029   if (globalIndex >= n_accepted || blockIdx.x == 0) {
0030     return;
0031   }
0032
0033   // Add the scanned block offsets to block-wise prefix sums of the number of
0034   // updated tracks.
0035   prefix_sums[globalIndex] += block_offsets[blockIdx.x - 1];
0036 }
0037
0038 }  // namespace traccc::cuda::kernels