src/seeding/seed_parameter_estimation_algorithm.cu

0001 /** TRACCC library, part of the ACTS project (R&D line)
0002  *
0003  * (c) 2021-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/cuda_error_handling.hpp"
0010 #include "../utils/global_index.hpp"
0011 #include "../utils/magnetic_field_types.hpp"
0012 #include "../utils/utils.hpp"
0013 #include "traccc/cuda/seeding/seed_parameter_estimation_algorithm.hpp"
0014
0015 // Project include(s).
0016 #include "traccc/seeding/device/estimate_track_params.hpp"
0017
0018 namespace traccc::cuda {
0019 namespace kernels {
0020
0021 /// CUDA kernel for running @c traccc::device::estimate_track_params
0022 template <typename bfield_t>
0023 __global__ void estimate_track_params(
0024     const track_params_estimation_config config,
0025     edm::measurement_collection::const_view measurements,
0026     edm::spacepoint_collection::const_view spacepoints,
0027     edm::seed_collection::const_view seeds, const bfield_t bfield,
0028     bound_track_parameters_collection_types::view params_view) {
0029   device::estimate_track_params(details::global_index1(), config, measurements,
0030                                 spacepoints, seeds, bfield, params_view);
0031 }
0032
0033 }  // namespace kernels
0034
0035 seed_parameter_estimation_algorithm::seed_parameter_estimation_algorithm(
0036     const track_params_estimation_config& config,
0037     const traccc::memory_resource& mr, const vecmem::copy& copy,
0038     const stream_wrapper& str, std::unique_ptr<const Logger> logger)
0039     : device::seed_parameter_estimation_algorithm(config, mr, copy,
0040                                                   std::move(logger)),
0041       cuda::algorithm_base(str) {}
0042
0043 void seed_parameter_estimation_algorithm::estimate_seed_params_kernel(
0044     const struct estimate_seed_params_kernel_payload& payload) const {
0045   const unsigned int n_threads = warp_size() * 4;
0046   const unsigned int n_blocks = (payload.n_seeds + n_threads - 1) / n_threads;
0047   magnetic_field_visitor<bfield_type_list<scalar>>(
0048       payload.bfield, [&]<typename bfield_view_t>(const bfield_view_t& bfield) {
0049         kernels::estimate_track_params<<<n_blocks, n_threads, 0,
0050                                          details::get_stream(stream())>>>(
0051             payload.config, payload.measurements, payload.spacepoints,
0052             payload.seeds, bfield, payload.params);
0053       });
0054   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
0055 }
0056
0057 }  // namespace traccc::cuda