ambiguity_resolution/kernels/scan_block_offsets.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 "scan_block_offsets.cuh"
0012
0013 // VecMem include(s).
0014 #include <vecmem/containers/device_vector.hpp>
0015
0016 namespace traccc::cuda::kernels {
0017
0018 __global__ void scan_block_offsets(device::scan_block_offsets_payload payload) {
0019   if (*(payload.terminate) == 1 || *(payload.n_updated_tracks) == 0) {
0020     return;
0021   }
0022
0023   extern __shared__ int shared_temp[];
0024
0025   vecmem::device_vector<const int> block_offsets(payload.block_offsets_view);
0026   vecmem::device_vector<int> scanned_block_offsets(
0027       payload.scanned_block_offsets_view);
0028
0029   // The number of blocks in the previous block_inclusive_scan = the nubmer of
0030   // threads of this kernel
0031   int n_blocks_prev = blockDim.x;
0032   auto threadIndex = threadIdx.x;
0033
0034   // 1. Load from global to shared
0035   shared_temp[threadIndex] = 0;
0036   if (threadIndex < n_blocks_prev) {
0037     shared_temp[threadIndex] = block_offsets[threadIndex];
0038   }
0039   __syncthreads();
0040
0041   // 2. Inclusive scan to caculated the scanned block offset which is the
0042   // prefix sum of block offsets
0043   for (int offset = 1; offset < n_blocks_prev; offset *= 2) {
0044     int temp = 0;
0045     if (threadIndex >= offset) {
0046       temp = shared_temp[threadIndex - offset];
0047     }
0048     __syncthreads();
0049     shared_temp[threadIndex] += temp;
0050     __syncthreads();
0051   }
0052
0053   // 3. Write back
0054   if (threadIndex < n_blocks_prev) {
0055     scanned_block_offsets[threadIndex] = shared_temp[threadIndex];
0056   }
0057 }
0058
0059 }  // namespace traccc::cuda::kernels