kernels/specializations/fit_backward_src.cuh

0001 /** TRACCC library, part of the ACTS project (R&D line)
0002  *
0003  * (c) 2025-2026 CERN for the benefit of the ACTS project
0004  *
0005  * Mozilla Public License Version 2.0
0006  */
0007
0008 #pragma once
0009
0010 // Local include(s).
0011 #include "../../../utils/global_index.hpp"
0012 #include "../fit_backward.hpp"
0013
0014 // Project include(s).
0015 #include "traccc/fitting/device/fit_backward.hpp"
0016
0017 namespace traccc::cuda {
0018 namespace kernels {
0019
0020 template <typename fitter_t>
0021 __global__ __launch_bounds__(128) void fit_backward(
0022     const fitting_config cfg, const device::fit_payload payload,
0023     const device::fit_tpayload<
0024         typename fitter_t::detector_type::const_view_type,
0025         typename fitter_t::bfield_type, typename fitter_t::surface_type>*
0026         tpayload) {
0027   device::fit_backward<fitter_t>(details::global_index1(), cfg, payload,
0028                                  *tpayload);
0029 }
0030
0031 }  // namespace kernels
0032
0033 template <typename fitter_t>
0034 void fit_backward(const dim3& grid_size, const dim3& block_size,
0035                   std::size_t shared_mem_size, const cudaStream_t& stream,
0036                   const fitting_config& cfg, const device::fit_payload& payload,
0037                   const device::fit_tpayload<
0038                       typename fitter_t::detector_type::const_view_type,
0039                       typename fitter_t::bfield_type,
0040                       typename fitter_t::surface_type>* tpayload) {
0041   kernels::fit_backward<fitter_t>
0042       <<<grid_size, block_size, shared_mem_size, stream>>>(cfg, payload,
0043                                                            tpayload);
0044 }
0045
0046 }  // namespace traccc::cuda