Warning, /acts/Detray/tests/integration_tests/device/sycl/propagator_kernel.sycl is written in an unsupported language. File is not indexed.
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 // Detray test include(s)
0010 #include "detray/test/device/propagator_test.hpp"
0011 #include "propagator_sycl_kernel.hpp"
0012
0013 namespace detray {
0014
0015 /// Test function for propagator
0016 template <typename bfield_bknd_t, typename detector_t>
0017 void propagator_test(
0018 typename detector_t::view_type det_data, const propagation::config& cfg,
0019 covfie::field_view<bfield_bknd_t> field_data,
0020 vecmem::data::vector_view<test_track>& tracks_data,
0021 vecmem::data::jagged_vector_view<detail::step_data<test_algebra>>&
0022 steps_data,
0023 detray::sycl::queue_wrapper queue) {
0024
0025 unsigned int localSize = 64;
0026 unsigned int globalSize = tracks_data.size();
0027
0028 const auto num = (globalSize + localSize - 1) / localSize;
0029 const auto ndrange = ::sycl::nd_range<1>{::sycl::range<1>(num * localSize),
0030 ::sycl::range<1>(localSize)};
0031
0032 reinterpret_cast<::sycl::queue*>(queue.queue())
0033 ->submit([&](::sycl::handler& h) {
0034 h.parallel_for(ndrange, [det_data, cfg, field_data, tracks_data,
0035 steps_data](::sycl::nd_item<1> item) {
0036 using detector_device_t =
0037 detector<typename detector_t::metadata,
0038 device_container_types>;
0039
0040 static_assert(
0041 std::is_same_v<typename detector_t::view_type,
0042 typename detector_device_t::view_type>,
0043 "Host and device detector views do not match");
0044
0045 detector_device_t dev_det(det_data);
0046
0047 vecmem::device_vector<test_track> tracks(tracks_data);
0048 vecmem::jagged_device_vector<detail::step_data<test_algebra>>
0049 steps(steps_data);
0050
0051 unsigned int gid = item.get_global_linear_id();
0052
0053 if (gid >= tracks.size()) {
0054 return;
0055 }
0056
0057 auto stepr = rk_stepper_t<covfie::field_view<bfield_bknd_t>>{};
0058 auto nav = navigator_t<detector_device_t>{};
0059
0060 // Create propagator
0061 using propagator_device_t =
0062 propagator<decltype(stepr), decltype(nav),
0063 actor_chain_device_t>;
0064 propagator_device_t p{cfg};
0065
0066 // Create actor states
0067 step_tracer_device_t::state tracer_state(steps.at(gid));
0068 pathlimit_aborter_t::state aborter_state{
0069 cfg.stepping.path_limit};
0070 actor::parameter_updater_state<test_algebra> updater_state{cfg};
0071 actor::pointwise_material_interactor<test_algebra>::state
0072 interactor_state{};
0073
0074 // Create the actor states
0075 auto actor_states =
0076 ::detray::tie(tracer_state, aborter_state, updater_state,
0077 interactor_state);
0078 // Create the propagator state
0079 typename propagator_device_t::state state(tracks[gid],
0080 field_data, dev_det);
0081
0082 state.stepping()
0083 .template set_constraint<step::constraint::e_accuracy>(
0084 cfg.stepping.step_constraint);
0085
0086 p.propagate(state, actor_states);
0087 });
0088 })
0089 .wait_and_throw();
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 detray::sycl::queue_wrapper);
0102
0103 /// Explicit instantiation for an inhomogeneous magnetic field
0104 /*template void propagator_test<bfield::sycl::inhom_bknd_t,
0105 detector<toy_metadata<test_algebra>, host_container_types>>(
0106 detector<toy_metadata<test_algebra>, host_container_types>::view_type, const
0107 propagation::config&, covfie::field_view<bfield::sycl::inhom_bknd_t>,
0108 vecmem::data::vector_view<test_track>&,
0109 vecmem::data::jagged_vector_view<detail::step_data<test_algebra>>&,
0110 detray::sycl::queue_wrapper);*/
0111
0112 } // namespace detray
