clusterization/kernels/ccl_kernel.cu

0001 /** TRACCC library, part of the ACTS project (R&D line)
0002  *
0003  * (c) 2022-2026 CERN for the benefit of the ACTS project
0004  *
0005  * Mozilla Public License Version 2.0
0006  */
0007
0008 // CUDA Library include(s).
0009 #include "../../utils/barrier.hpp"
0010 #include "../../utils/cuda_error_handling.hpp"
0011 #include "../../utils/thread_id.hpp"
0012 #include "traccc/clusterization/clustering_config.hpp"
0013 #include "traccc/clusterization/device/ccl_kernel_definitions.hpp"
0014
0015 // Project include(s)
0016 #include "traccc/clusterization/device/ccl_kernel.hpp"
0017
0018 namespace traccc::cuda::kernels {
0019
0020 /// CUDA kernel for running @c traccc::device::ccl_kernel
0021 __global__ void ccl_kernel(
0022     const clustering_config cfg,
0023     const edm::silicon_cell_collection::const_view cells_view,
0024     const detector_design_description::const_view det_desc_view,
0025     const detector_conditions_description::const_view det_cond_view,
0026     edm::measurement_collection::view measurements_view,
0027     vecmem::data::vector_view<device::details::fallback_index_t> f_backup_view,
0028     vecmem::data::vector_view<device::details::fallback_index_t> gf_backup_view,
0029     vecmem::data::vector_view<unsigned char> adjc_backup_view,
0030     vecmem::data::vector_view<device::details::fallback_index_t>
0031         adjv_backup_view,
0032     unsigned int* backup_mutex_ptr,
0033     vecmem::data::vector_view<unsigned int> disjoint_set_view,
0034     vecmem::data::vector_view<unsigned int> cluster_size_view) {
0035   __shared__ std::size_t partition_start, partition_end;
0036   __shared__ std::size_t outi;
0037   extern __shared__ device::details::index_t shared_v[];
0038   vecmem::device_atomic_ref<unsigned int> backup_mutex(*backup_mutex_ptr);
0039
0040   using vector_size_t =
0041       vecmem::data::vector_view<device::details::index_t>::size_type;
0042
0043   vecmem::data::vector_view<device::details::index_t> f_view{
0044       static_cast<vector_size_t>(cfg.max_partition_size()), shared_v};
0045   vecmem::data::vector_view<device::details::index_t> gf_view{
0046       static_cast<vector_size_t>(cfg.max_partition_size()),
0047       shared_v + cfg.max_partition_size()};
0048   traccc::cuda::barrier barry_r;
0049   const details::thread_id1 thread_id;
0050
0051   device::ccl_kernel(cfg, thread_id, cells_view, det_desc_view, det_cond_view,
0052                      partition_start, partition_end, outi, f_view, gf_view,
0053                      f_backup_view, gf_backup_view, adjc_backup_view,
0054                      adjv_backup_view, backup_mutex, disjoint_set_view,
0055                      cluster_size_view, barry_r, measurements_view);
0056 }
0057 }  // namespace traccc::cuda::kernels