src/clusterization/clusterization_algorithm.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 "../sanity/contiguous_on.cuh"
0010 #include "../sanity/ordered_on.cuh"
0011 #include "../utils/barrier.hpp"
0012 #include "../utils/cuda_error_handling.hpp"
0013 #include "../utils/thread_id.hpp"
0014 #include "./kernels/ccl_kernel.cuh"
0015 #include "./kernels/reify_cluster_data.cuh"
0016 #include "./kernels/sort_cells.cuh"
0017 #include "traccc/clusterization/device/ccl_kernel_definitions.hpp"
0018 #include "traccc/cuda/clusterization/clusterization_algorithm.hpp"
0019 #include "traccc/utils/projections.hpp"
0020 #include "traccc/utils/relations.hpp"
0021
0022 // Vecmem include(s).
0023 #include <cstring>
0024 #include <limits>
0025
0026 #include <vecmem/containers/device_vector.hpp>
0027 #include <vecmem/utils/copy.hpp>
0028
0029 namespace traccc::cuda {
0030
0031 clusterization_algorithm::clusterization_algorithm(
0032     const traccc::memory_resource& mr, const vecmem::copy& copy,
0033     const stream_wrapper& str, const config_type& config,
0034     std::unique_ptr<const Logger> logger)
0035     : device::clusterization_algorithm(mr, copy, config, std::move(logger)),
0036       cuda::algorithm_base(str) {}
0037
0038 bool clusterization_algorithm::input_is_contiguous(
0039     const edm::silicon_cell_collection::const_view& cells) const {
0040   return is_contiguous_on<edm::silicon_cell_collection::const_device>(
0041       cell_module_projection(), mr().main, copy(), stream(), cells);
0042 }
0043
0044 bool clusterization_algorithm::input_is_sorted(
0045     const edm::silicon_cell_collection::const_view& cells) const {
0046   return is_ordered_on<edm::silicon_cell_collection::const_device>(
0047       channel0_major_cell_order_relation(), mr().main, copy(), stream(), cells);
0048 }
0049
0050 void clusterization_algorithm::sort_cells_kernel(
0051     const unsigned int num_cells,
0052     const edm::silicon_cell_collection::const_view& cells,
0053     edm::silicon_cell_collection::view& new_cells,
0054     vecmem::data::vector_view<unsigned int>& permutation_map_view,
0055     const bool write_permutation) const {
0056   static constexpr std::size_t SORT_THREADS_PER_BLOCK = 512;
0057
0058   const unsigned blockSize = SORT_THREADS_PER_BLOCK;
0059   const unsigned int cellsPerThread = 4;
0060   const unsigned int numBlocks =
0061       (num_cells + (blockSize * cellsPerThread) - 1) /
0062       (blockSize * cellsPerThread);
0063
0064   kernels::sort_cells<SORT_THREADS_PER_BLOCK, 8>
0065       <<<numBlocks, blockSize, 0, details::get_stream(stream())>>>(
0066           cellsPerThread, num_cells, cells, new_cells, permutation_map_view,
0067           write_permutation);
0068   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
0069 }
0070
0071 void clusterization_algorithm::ccl_kernel(
0072     const ccl_kernel_payload& payload) const {
0073   const unsigned int num_blocks =
0074       (payload.n_cells + (payload.config.target_partition_size()) - 1) /
0075       payload.config.target_partition_size();
0076   kernels::ccl_kernel<<<num_blocks, payload.config.threads_per_partition,
0077                         2 * payload.config.max_partition_size() *
0078                             sizeof(device::details::index_t),
0079                         details::get_stream(stream())>>>(
0080       payload.config, payload.cells, payload.det_descr, payload.det_cond,
0081       payload.measurements, payload.f_backup, payload.gf_backup,
0082       payload.adjc_backup, payload.adjv_backup, payload.backup_mutex,
0083       payload.disjoint_set, payload.cluster_sizes);
0084   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
0085   // With sorting enabled, the kernel reads from sorted cell and permutation
0086   // map buffers that the base class destroys without any further
0087   // synchronization, so the kernel must finish before returning.
0088   if (payload.config.sort_cells) {
0089     stream().synchronize();
0090   }
0091 }
0092
0093 void clusterization_algorithm::cluster_maker_kernel(
0094     unsigned int num_cells,
0095     const vecmem::data::vector_view<unsigned int>& disjoint_set,
0096     edm::silicon_cluster_collection::view& cluster_data,
0097     const vecmem::data::vector_view<const unsigned int>& permutation_map_view)
0098     const {
0099   const unsigned int num_threads = warp_size() * 16u;
0100   const unsigned int num_blocks = (num_cells + num_threads - 1) / num_threads;
0101   kernels::reify_cluster_data<<<num_blocks, num_threads, 0,
0102                                 details::get_stream(stream())>>>(
0103       disjoint_set, permutation_map_view, cluster_data);
0104   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
0105   // The base class destroys the input buffers right after this call, so
0106   // the kernel must finish before returning.
0107   stream().synchronize();
0108 }
0109
0110 }  // namespace traccc::cuda