ambiguity_resolution/kernels/fill_vectors.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_vectors.cuh"
0011
0012 // Project include(s).
0013 #include "traccc/ambiguity_resolution/ambiguity_resolution_config.hpp"
0014
0015 // VecMem include(s).
0016 #include <vecmem/containers/device_vector.hpp>
0017 #include <vecmem/containers/jagged_device_vector.hpp>
0018
0019 namespace traccc::cuda::kernels {
0020
0021 __global__ void fill_vectors(const ambiguity_resolution_config cfg,
0022                              device::fill_vectors_payload payload) {
0023   const edm::track_container<default_algebra>::const_device track_candidates(
0024       payload.tracks_view);
0025
0026   const auto globalIndex = details::global_index1();
0027   if (globalIndex >= track_candidates.tracks.size()) {
0028     return;
0029   }
0030
0031   const auto track = track_candidates.tracks.at(globalIndex);
0032
0033   vecmem::jagged_device_vector<measurement_id_type> meas_ids(
0034       payload.meas_ids_view);
0035   vecmem::device_vector<measurement_id_type> flat_meas_ids(
0036       payload.flat_meas_ids_view);
0037   vecmem::device_vector<traccc::scalar> pvals(payload.pvals_view);
0038   vecmem::device_vector<unsigned int> n_meas(payload.n_meas_view);
0039   vecmem::device_vector<int> status(payload.status_view);
0040
0041   pvals.at(globalIndex) = track.pval();
0042
0043   if (track.constituent_links().size() < cfg.min_meas_per_track) {
0044     status.at(globalIndex) = 0;
0045   } else {
0046     for (const auto& [type, meas_idx] :
0047          track_candidates.tracks.constituent_links().at(globalIndex)) {
0048       assert(type == edm::track_constituent_link::measurement);
0049       meas_ids.at(globalIndex)
0050           .push_back(track_candidates.measurements.at(meas_idx).identifier());
0051       flat_meas_ids.push_back(
0052           track_candidates.measurements.at(meas_idx).identifier());
0053     }
0054     n_meas.at(globalIndex) = track.constituent_links().size();
0055   }
0056 }
0057 }  // namespace traccc::cuda::kernels