device/cuda/propagator_cuda_kernel.cu

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/cuda_definitions.hpp"
0011
0012 // Detray test include(s)
0013 #include "propagator_cuda_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   int gid = threadIdx.x + blockIdx.x * blockDim.x;
0025   using detector_device_t =
0026       detector<typename detector_t::metadata, device_container_types>;
0027
0028   static_assert(std::is_same_v<typename detector_t::view_type,
0029                                typename detector_device_t::view_type>,
0030                 "Host and device detector views do not match");
0031
0032   detector_device_t det(det_data);
0033   vecmem::device_vector<test_track> tracks(tracks_data);
0034   vecmem::jagged_device_vector<detail::step_data<test_algebra>> steps(
0035       steps_data);
0036
0037   if (gid >= tracks.size()) {
0038     return;
0039   }
0040
0041   auto stepr = rk_stepper_t<covfie::field_view<bfield_bknd_t>>{};
0042   auto nav = navigator_t<detector_device_t>{};
0043
0044   // Create propagator
0045   using propagator_device_t =
0046       propagator<decltype(stepr), decltype(nav), actor_chain_device_t>;
0047
0048   propagator_device_t p{cfg};
0049
0050   // Create actor states
0051   step_tracer_device_t::state tracer_state(steps.at(gid));
0052   pathlimit_aborter_t::state aborter_state{cfg.stepping.path_limit};
0053   actor::parameter_updater_state<test_algebra> updater_state{cfg};
0054   actor::pointwise_material_interactor<test_algebra>::state interactor_state{};
0055
0056   // Create the actor states
0057   auto actor_states = ::detray::tie(tracer_state, aborter_state, updater_state,
0058                                     interactor_state);
0059   // Create the propagator state
0060   typename propagator_device_t::state state(tracks.at(gid), field_data, det);
0061
0062   auto& ptc = state.stepping().particle_hypothesis();
0063   state.set_particle(update_particle_hypothesis(ptc, tracks.at(gid)));
0064
0065   state.stepping().template set_constraint<step::constraint::e_accuracy>(
0066       cfg.stepping.step_constraint);
0067
0068   // Run propagation
0069   p.propagate(state, actor_states);
0070 }
0071
0072 /// Launch the device kernel
0073 template <typename bfield_bknd_t, typename detector_t>
0074 void propagator_test(
0075     typename detector_t::view_type det_view, const propagation::config& cfg,
0076     covfie::field_view<bfield_bknd_t> field_data,
0077     vecmem::data::vector_view<test_track>& tracks_data,
0078     vecmem::data::jagged_vector_view<detail::step_data<test_algebra>>&
0079         step_data) {
0080   constexpr int thread_dim = 2 * WARP_SIZE;
0081   int block_dim = tracks_data.size() / thread_dim + 1;
0082
0083   // run the test kernel
0084   propagator_test_kernel<bfield_bknd_t, detector_t><<<block_dim, thread_dim>>>(
0085       det_view, cfg, field_data, tracks_data, step_data);
0086
0087   // cuda error check
0088   DETRAY_CUDA_ERROR_CHECK(cudaGetLastError());
0089   DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0090 }
0091
0092 /// Explicit instantiation for a constant magnetic field
0093 template void
0094 propagator_test<bfield::const_bknd_t<dscalar<test_algebra>>,
0095                 detector<toy_metadata<test_algebra>, host_container_types>>(
0096     detector<toy_metadata<test_algebra>, host_container_types>::view_type,
0097     const propagation::config&,
0098     covfie::field_view<bfield::const_bknd_t<dscalar<test_algebra>>>,
0099     vecmem::data::vector_view<test_track>&,
0100     vecmem::data::jagged_vector_view<detail::step_data<test_algebra>>&);
0101
0102 /// Explicit instantiation for an inhomogeneous magnetic field
0103 template void
0104 propagator_test<bfield::cuda::inhom_bknd_t<dscalar<test_algebra>>,
0105                 detector<toy_metadata<test_algebra>, host_container_types>>(
0106     detector<toy_metadata<test_algebra>, host_container_types>::view_type,
0107     const propagation::config&,
0108     covfie::field_view<bfield::cuda::inhom_bknd_t<dscalar<test_algebra>>>,
0109     vecmem::data::vector_view<test_track>&,
0110     vecmem::data::jagged_vector_view<detail::step_data<test_algebra>>&);
0111
0112 }  // namespace detray