src/clusterization/clusterization_algorithm.hip

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 // Local include(s).
0009 #include "../utils/hip_error_handling.hpp"
0010 #include "../utils/utils.hpp"
0011 #include "./kernels/ccl_kernel.hpp"
0012 #include "./kernels/reify_cluster_data.hpp"
0013 #include "traccc/hip/clusterization/clusterization_algorithm.hpp"
0014
0015 // HIP include(s).
0016 #include <hip/hip_runtime.h>
0017
0018 // System include(s).
0019 #include <stdexcept>
0020
0021 namespace traccc::hip {
0022
0023 clusterization_algorithm::clusterization_algorithm(
0024     const traccc::memory_resource& mr, const vecmem::copy& copy,
0025     hip::stream& str, const config_type& config,
0026     std::unique_ptr<const Logger> logger)
0027     : device::clusterization_algorithm(mr, copy, config, std::move(logger)),
0028       hip::algorithm_base(str) {}
0029
0030 bool clusterization_algorithm::input_is_contiguous(
0031     const edm::silicon_cell_collection::const_view& /*cells*/) const {
0032   /// TODO: implement input checks
0033   return true;
0034 }
0035
0036 bool clusterization_algorithm::input_is_sorted(
0037     const edm::silicon_cell_collection::const_view& /*cells*/) const {
0038   /// TODO: implement input checks
0039   return true;
0040 }
0041
0042 void clusterization_algorithm::sort_cells_kernel(
0043     const unsigned int, const edm::silicon_cell_collection::const_view&,
0044     edm::silicon_cell_collection::view&,
0045     vecmem::data::vector_view<unsigned int>&, const bool) const {
0046   throw std::runtime_error(
0047       "Cell sorting is not yet implemented for the HIP backend.");
0048 }
0049
0050 void clusterization_algorithm::ccl_kernel(
0051     const ccl_kernel_payload& payload) const {
0052   const unsigned int num_blocks =
0053       (payload.n_cells + (payload.config.target_partition_size()) - 1) /
0054       payload.config.target_partition_size();
0055   hipLaunchKernelGGL(kernels::ccl_kernel, dim3{num_blocks},
0056                      dim3{payload.config.threads_per_partition},
0057                      2 * payload.config.max_partition_size() *
0058                          sizeof(device::details::index_t),
0059                      details::get_stream(stream()), payload.config,
0060                      payload.cells, payload.det_descr, payload.det_cond,
0061                      payload.measurements, payload.f_backup, payload.gf_backup,
0062                      payload.adjc_backup, payload.adjv_backup,
0063                      payload.backup_mutex, payload.disjoint_set,
0064                      payload.cluster_sizes);
0065   TRACCC_HIP_ERROR_CHECK(hipGetLastError());
0066 }
0067
0068 void clusterization_algorithm::cluster_maker_kernel(
0069     unsigned int num_cells,
0070     const vecmem::data::vector_view<unsigned int>& disjoint_set,
0071     edm::silicon_cluster_collection::view& cluster_data,
0072     const vecmem::data::vector_view<const unsigned int>& permutation_map_view)
0073     const {
0074   const unsigned int num_threads = warp_size() * 16u;
0075   const unsigned int num_blocks = (num_cells + num_threads - 1) / num_threads;
0076   hipLaunchKernelGGL(kernels::reify_cluster_data, num_blocks, num_threads, 0,
0077                      details::get_stream(stream()), disjoint_set,
0078                      permutation_map_view, cluster_data);
0079   TRACCC_HIP_ERROR_CHECK(hipGetLastError());
0080   // The base class destroys the input buffers right after this call, so
0081   // the kernel must finish before returning.
0082   stream().synchronize();
0083 }
0084
0085 }  // namespace traccc::hip