Warning, /acts/Traccc/device/hip/src/seeding/seed_parameter_estimation_algorithm.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) 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/global_index.hpp"
0010 #include "../utils/hip_error_handling.hpp"
0011 #include "../utils/magnetic_field_types.hpp"
0012 #include "../utils/utils.hpp"
0013 #include "traccc/hip/seeding/seed_parameter_estimation_algorithm.hpp"
0014
0015 // Project include(s).
0016 #include "traccc/seeding/device/estimate_track_params.hpp"
0017
0018 namespace traccc::hip {
0019 namespace kernels {
0020
0021 /// HIP 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 hip::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 hipLaunchKernelGGL(kernels::estimate_track_params<bfield_view_t>,
0050 n_blocks, n_threads, 0,
0051 details::get_stream(stream()), payload.config,
0052 payload.measurements, payload.spacepoints,
0053 payload.seeds, bfield, payload.params);
0054 });
0055 TRACCC_HIP_ERROR_CHECK(hipGetLastError());
0056 }
0057
0058 } // namespace traccc::hip
