Warning, /acts/Traccc/device/hip/src/clusterization/kernels/ccl_kernel.hip is written in an unsupported language. File is not indexed.
0001 /** TRACCC library, part of the ACTS project (R&D line)
0002 *
0003 * (c) 2022-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/thread_id.hpp"
0011
0012 // Project include(s).
0013 #include "traccc/clusterization/device/ccl_kernel.hpp"
0014
0015 namespace traccc::hip::kernels {
0016
0017 /// HIP kernel for running @c traccc::device::ccl_kernel
0018 __global__ void ccl_kernel(
0019 const clustering_config cfg,
0020 const edm::silicon_cell_collection::const_view cells_view,
0021 const detector_design_description::const_view det_desc_view,
0022 const detector_conditions_description::const_view det_cond_view,
0023 edm::measurement_collection::view measurements_view,
0024 vecmem::data::vector_view<device::details::fallback_index_t> f_backup_view,
0025 vecmem::data::vector_view<device::details::fallback_index_t> gf_backup_view,
0026 vecmem::data::vector_view<unsigned char> adjc_backup_view,
0027 vecmem::data::vector_view<device::details::fallback_index_t>
0028 adjv_backup_view,
0029 unsigned int* backup_mutex_ptr,
0030 vecmem::data::vector_view<unsigned int> disjoint_set_view,
0031 vecmem::data::vector_view<unsigned int> cluster_size_view) {
0032 __shared__ std::size_t partition_start, partition_end;
0033 __shared__ std::size_t outi;
0034 extern __shared__ device::details::index_t shared_v[];
0035 vecmem::device_atomic_ref<unsigned int> backup_mutex(*backup_mutex_ptr);
0036
0037 using vector_size_t =
0038 vecmem::data::vector_view<device::details::index_t>::size_type;
0039
0040 vecmem::data::vector_view<device::details::index_t> f_view{
0041 static_cast<vector_size_t>(cfg.max_partition_size()), shared_v};
0042 vecmem::data::vector_view<device::details::index_t> gf_view{
0043 static_cast<vector_size_t>(cfg.max_partition_size()),
0044 shared_v + cfg.max_partition_size()};
0045 traccc::hip::barrier barry_r;
0046 const details::thread_id1 thread_id;
0047
0048 device::ccl_kernel(cfg, thread_id, cells_view, det_desc_view, det_cond_view,
0049 partition_start, partition_end, outi, f_view, gf_view,
0050 f_backup_view, gf_backup_view, adjc_backup_view,
0051 adjv_backup_view, backup_mutex, disjoint_set_view,
0052 cluster_size_view, barry_r, measurements_view);
0053 }
0054
0055 } // namespace traccc::hip::kernels
