device/hip/propagator_hip_kernel.hip

0001 // This file is part of the ACTS project.
0002 //
0003 // Copyright (C) 2016 CERN for the benefit of the ACTS project
0004 //
0005 // This Source Code Form is subject to the terms of the Mozilla Public
0006 // License, v. 2.0. If a copy of the MPL was not distributed with this
0007 // file, You can obtain one at https://mozilla.org/MPL/2.0/.
0008
0009 // Project include(s)
0010 #include "detray/definitions/detail/hip_definitions.hpp"
0011
0012 // Detray test include(s)
0013 #include "propagator_hip_kernel.hpp"
0014
0015 namespace detray {
0016
0017 template <typename bfield_bknd_t, typename detector_t>
0018 __global__ void propagator_test_kernel(
0019     typename detector_t::view_type det_data, const propagation::config cfg,
0020     covfie::field_view<bfield_bknd_t> field_data,
0021     vecmem::data::vector_view<test_track> tracks_data,
0022     vecmem::data::jagged_vector_view<detail::step_data<test_algebra>>
0023         steps_data) {
0024
0025     int gid = threadIdx.x + blockIdx.x * blockDim.x;
0026     using detector_device_t =
0027         detector<typename detector_t::metadata, device_container_types>;
0028
0029     static_assert(std::is_same_v<typename detector_t::view_type,
0030                                  typename detector_device_t::view_type>,
0031                   "Host and device detector views do not match");
0032
0033     detector_device_t det(det_data);
0034     vecmem::device_vector<test_track> tracks(tracks_data);
0035     vecmem::jagged_device_vector<detail::step_data<test_algebra>> steps(
0036         steps_data);
0037
0038     if (gid >= static_cast<int>(tracks.size())) {
0039         return;
0040     }
0041
0042     auto stepr = rk_stepper_t<covfie::field_view<bfield_bknd_t>>{};
0043     auto nav = navigator_t<detector_device_t>{};
0044
0045     // Create propagator
0046     using propagator_device_t =
0047         propagator<decltype(stepr), decltype(nav), actor_chain_device_t>;
0048
0049     propagator_device_t p{cfg};
0050
0051     // Create actor states
0052     step_tracer_device_t::state tracer_state(steps.at(gid));
0053     pathlimit_aborter_t::state aborter_state{cfg.stepping.path_limit};
0054     actor::parameter_updater_state<test_algebra> updater_state{cfg};
0055     actor::pointwise_material_interactor<test_algebra>::state
0056         interactor_state{};
0057
0058     // Create the actor states
0059     auto actor_states = ::detray::tie(tracer_state, aborter_state,
0060                                       updater_state, interactor_state);
0061     // Create the propagator state
0062     typename propagator_device_t::state state(tracks.at(gid), field_data, det);
0063
0064     auto& ptc = state.stepping().particle_hypothesis();
0065     state.set_particle(update_particle_hypothesis(ptc, tracks.at(gid)));
0066
0067     state.stepping().template set_constraint<step::constraint::e_accuracy>(
0068         cfg.stepping.step_constraint);
0069
0070     // Run propagation
0071     p.propagate(state, actor_states);
0072 }
0073
0074 /// Launch the device kernel
0075 template <typename bfield_bknd_t, typename detector_t>
0076 void propagator_test(
0077     typename detector_t::view_type det_view, const propagation::config& cfg,
0078     covfie::field_view<bfield_bknd_t> field_data,
0079     vecmem::data::vector_view<test_track>& tracks_data,
0080     vecmem::data::jagged_vector_view<detail::step_data<test_algebra>>&
0081         step_data) {
0082
0083     constexpr int thread_dim = 2 * WARP_SIZE;
0084     int block_dim = tracks_data.size() / thread_dim + 1;
0085
0086     // run the test kernel
0087     hipLaunchKernelGGL((propagator_test_kernel<bfield_bknd_t, detector_t>),
0088                        dim3(block_dim), dim3(thread_dim), 0, 0, det_view, cfg,
0089                        field_data, tracks_data, step_data);
0090
0091     // hip error check
0092     DETRAY_HIP_ERROR_CHECK(hipGetLastError());
0093     DETRAY_HIP_ERROR_CHECK(hipDeviceSynchronize());
0094 }
0095
0096 /// Explicit instantiation for a constant magnetic field
0097 template void
0098 propagator_test<bfield::const_bknd_t<dscalar<test_algebra>>,
0099                 detector<toy_metadata<test_algebra>, host_container_types>>(
0100     detector<toy_metadata<test_algebra>, host_container_types>::view_type,
0101     const propagation::config&,
0102     covfie::field_view<bfield::const_bknd_t<dscalar<test_algebra>>>,
0103     vecmem::data::vector_view<test_track>&,
0104     vecmem::data::jagged_vector_view<detail::step_data<test_algebra>>&);
0105
0106 /// Explicit instantiation for an inhomogeneous magnetic field
0107 /*
0108 template void
0109 propagator_test<bfield::hip::inhom_bknd_t,
0110                 detector<toy_metadata<test_algebra>, host_container_types>>(
0111     detector<toy_metadata<test_algebra>, host_container_types>::view_type,
0112     const propagation::config&, covfie::field_view<bfield::hip::inhom_bknd_t>,
0113     vecmem::data::vector_view<test_track>&,
0114     vecmem::data::jagged_vector_view<detail::step_data<test_algebra>>&); */
0115
0116 }  // namespace detray