ambiguity_resolution/kernels/fill_track_candidates.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_track_candidates.cuh"
0011
0012 // VecMem include(s).
0013 #include <vecmem/containers/device_vector.hpp>
0014
0015 namespace traccc::cuda::kernels {
0016
0017 __global__ void fill_track_candidates(
0018     device::fill_track_candidates_payload payload) {
0019   const auto globalIndex = details::global_index1();
0020   if (globalIndex >= payload.n_accepted) {
0021     return;
0022   }
0023
0024   // Set up the device objects.
0025   vecmem::device_vector<const unsigned int> sorted_ids(payload.sorted_ids_view);
0026   edm::track_collection<default_algebra>::const_device tracks(
0027       payload.tracks_view.tracks);
0028   edm::track_collection<default_algebra>::device res_tracks(
0029       payload.res_tracks_view.tracks);
0030
0031   // Copy the appropriate track candidate.
0032   res_tracks.at(globalIndex) = tracks.at(sorted_ids.at(globalIndex));
0033 }
0034
0035 }  // namespace traccc::cuda::kernels