Warning, /acts/Traccc/device/sycl/src/clusterization/clusterization_algorithm.sycl 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-2026 CERN for the benefit of the ACTS project
0004 *
0005 * Mozilla Public License Version 2.0
0006 */
0007
0008 // SYCL include(s).
0009 #include <sycl/sycl.hpp>
0010
0011 // Local include(s).
0012 #include "../sanity/contiguous_on.hpp"
0013 #include "../sanity/ordered_on.hpp"
0014 #include "../utils/barrier.hpp"
0015 #include "../utils/calculate1DimNdRange.hpp"
0016 #include "../utils/get_queue.hpp"
0017 #include "../utils/thread_id.hpp"
0018 #include "traccc/sycl/clusterization/clusterization_algorithm.hpp"
0019
0020 // Project include(s)
0021 #include "traccc/clusterization/device/ccl_kernel.hpp"
0022 #include "traccc/clusterization/device/reify_cluster_data.hpp"
0023 #include "traccc/utils/projections.hpp"
0024 #include "traccc/utils/relations.hpp"
0025
0026 // Vecmem include(s).
0027 #include <vecmem/utils/sycl/local_accessor.hpp>
0028
0029 // System include(s).
0030 #include <stdexcept>
0031
0032 namespace traccc::sycl {
0033 namespace kernels {
0034
0035 /// Class identifying the kernel running @c traccc::device::ccl_kernel
0036 class ccl_kernel;
0037 /// Class identifying the kernel running @c traccc::device::reify_cluster_data
0038 class reify_cluster_data;
0039
0040 } // namespace kernels
0041
0042 clusterization_algorithm::clusterization_algorithm(
0043 const traccc::memory_resource& mr, const vecmem::copy& copy,
0044 queue_wrapper& queue, const config_type& config,
0045 std::unique_ptr<const Logger> logger)
0046 : device::clusterization_algorithm(mr, copy, config, std::move(logger)),
0047 sycl::algorithm_base(queue) {}
0048
0049 bool clusterization_algorithm::input_is_contiguous(
0050 const edm::silicon_cell_collection::const_view& cells) const {
0051 return is_contiguous_on<edm::silicon_cell_collection::const_device>(
0052 cell_module_projection(), mr().main, copy(), details::get_queue(queue()),
0053 cells);
0054 }
0055
0056 bool clusterization_algorithm::input_is_sorted(
0057 const edm::silicon_cell_collection::const_view& cells) const {
0058 return is_ordered_on<edm::silicon_cell_collection::const_device>(
0059 channel0_major_cell_order_relation(), mr().main, copy(),
0060 details::get_queue(queue()), cells);
0061 }
0062
0063 void clusterization_algorithm::sort_cells_kernel(
0064 const unsigned int, const edm::silicon_cell_collection::const_view&,
0065 edm::silicon_cell_collection::view&,
0066 vecmem::data::vector_view<unsigned int>&, const bool) const {
0067 throw std::runtime_error(
0068 "Cell sorting is not yet implemented for the SYCL backend.");
0069 }
0070
0071 void clusterization_algorithm::ccl_kernel(
0072 const ccl_kernel_payload& payload) const {
0073 // Check if device is capable of allocating sufficient local memory
0074 assert(sizeof(device::details::index_t) * 2 *
0075 payload.config.max_partition_size() +
0076 3 * sizeof(std::size_t) <
0077 details::get_queue(queue())
0078 .get_device()
0079 .get_info<::sycl::info::device::local_mem_size>());
0080
0081 std::size_t num_blocks =
0082 (payload.n_cells + payload.config.target_partition_size() - 1) /
0083 payload.config.target_partition_size();
0084 const ::sycl::nd_range cclKernelRange(
0085 ::sycl::range<1>(num_blocks * payload.config.threads_per_partition),
0086 ::sycl::range<1>(payload.config.threads_per_partition));
0087 details::get_queue(queue())
0088 .submit([&](::sycl::handler& h) {
0089 // Allocate shared memory for the kernel.
0090 vecmem::sycl::local_accessor<std::size_t> shared_uint(3, h);
0091 vecmem::sycl::local_accessor<device::details::index_t> shared_idx(
0092 2 * payload.config.max_partition_size(), h);
0093
0094 // Launch the kernel.
0095 h.parallel_for<kernels::ccl_kernel>(
0096 cclKernelRange,
0097 [shared_uint, shared_idx, cells = payload.cells,
0098 det_descr = payload.det_descr, det_cond = payload.det_cond,
0099 measurements = payload.measurements, f_backup = payload.f_backup,
0100 gf_backup = payload.gf_backup, adjc_backup = payload.adjc_backup,
0101 adjv_backup = payload.adjv_backup,
0102 backup_mutex = payload.backup_mutex, config = payload.config,
0103 disjoint_set = payload.disjoint_set,
0104 cluster_sizes = payload.cluster_sizes](::sycl::nd_item<1> item) {
0105 // Construct more readable variable names.
0106 using vector_size_t = vecmem::data::vector_view<
0107 device::details::index_t>::size_type;
0108 vecmem::data::vector_view<device::details::index_t> f_view{
0109 static_cast<vector_size_t>(config.max_partition_size()),
0110 &shared_idx[0]};
0111 vecmem::data::vector_view<device::details::index_t> gf_view{
0112 static_cast<vector_size_t>(config.max_partition_size()),
0113 &shared_idx[config.max_partition_size()]};
0114 std::size_t& partition_start = shared_uint[0];
0115 std::size_t& partition_end = shared_uint[1];
0116 std::size_t& outi = shared_uint[2];
0117
0118 // Mutex for scratch space
0119 vecmem::device_atomic_ref<unsigned int> mutex(*backup_mutex);
0120
0121 // Barrier used in the algorithm.
0122 const details::barrier barrier{item};
0123 const details::thread_id thread_id{item};
0124
0125 // Run the algorithm for this thread.
0126 device::ccl_kernel(config, thread_id, cells, det_descr, det_cond,
0127 partition_start, partition_end, outi, f_view,
0128 gf_view, f_backup, gf_backup, adjc_backup,
0129 adjv_backup, mutex, disjoint_set,
0130 cluster_sizes, barrier, measurements);
0131 });
0132 })
0133 .wait_and_throw();
0134 }
0135
0136 void clusterization_algorithm::cluster_maker_kernel(
0137 unsigned int num_cells,
0138 const vecmem::data::vector_view<unsigned int>& disjoint_set,
0139 edm::silicon_cluster_collection::view& cluster_data,
0140 const vecmem::data::vector_view<const unsigned int>& permutation_map_view)
0141 const {
0142 const ::sycl::nd_range range =
0143 details::calculate1DimNdRange(num_cells, warp_size() * 8u);
0144 details::get_queue(queue())
0145 .submit([&](::sycl::handler& h) {
0146 h.parallel_for<kernels::reify_cluster_data>(
0147 range, [disjoint_set, permutation_map_view,
0148 cluster_data](::sycl::nd_item<1> item) {
0149 device::reify_cluster_data(
0150 static_cast<unsigned int>(item.get_global_linear_id()),
0151 disjoint_set, permutation_map_view, cluster_data);
0152 });
0153 })
0154 .wait_and_throw();
0155 }
0156
0157 } // namespace traccc::sycl
