device/sycl/propagator_sycl_kernel.hpp

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 #pragma once
0010
0011 // Project include(s)
0012 #include "detray/core/detector.hpp"
0013 #include "detray/detectors/toy_metadata.hpp"
0014
0015 // Detray test include(s)
0016 #include "detray/test/common/bfield.hpp"
0017 #include "detray/test/device/propagator_test.hpp"
0018 #include "queue_wrapper.hpp"
0019
0020 // Vecmem include(s)
0021 #include <vecmem/utils/sycl/copy.hpp>
0022
0023 namespace detray {
0024
0025 /// Launch the propagation test kernel
0026 template <typename bfield_bknd_t, typename detector_t>
0027 void propagator_test(
0028     typename detector_t::view_type, const propagation::config &,
0029     covfie::field_view<bfield_bknd_t>, vecmem::data::vector_view<test_track> &,
0030     vecmem::data::jagged_vector_view<detail::step_data<test_algebra>> &,
0031     sycl::queue_wrapper);
0032
0033 /// test function for propagator on the device
0034 template <typename bfield_bknd_t, typename detector_t>
0035 inline auto run_propagation_device(
0036     vecmem::memory_resource *mr, const propagation::config &cfg,
0037     typename detector_t::view_type det_view,
0038     covfie::field_view<bfield_bknd_t> field_data, sycl::queue_wrapper queue,
0039     dvector<test_track> &tracks,
0040     const vecmem::jagged_vector<detail::step_data<test_algebra>> &host_steps)
0041     -> vecmem::jagged_vector<detail::step_data<test_algebra>> {
0042   // Helper object for performing memory copies.
0043   vecmem::copy copy;
0044
0045   // Get tracks data
0046   auto tracks_data = vecmem::get_data(tracks);
0047
0048   // Create vector buffer for track recording
0049   std::vector<std::size_t> sizes(tracks.size(), 0);
0050   std::vector<std::size_t> capacities;
0051   for (auto &st : host_steps) {
0052     capacities.push_back(st.size());
0053   }
0054
0055   vecmem::data::jagged_vector_buffer<detail::step_data<test_algebra>>
0056       steps_buffer(capacities, *mr, nullptr,
0057                    vecmem::data::buffer_type::resizable);
0058
0059   copy.setup(steps_buffer)->wait();
0060
0061   // Run the propagator test for GPU device
0062   propagator_test<bfield_bknd_t, detector_t>(det_view, cfg, field_data,
0063                                              tracks_data, steps_buffer, queue);
0064
0065   vecmem::jagged_vector<detail::step_data<test_algebra>> steps(mr);
0066
0067   copy(steps_buffer, steps)->wait();
0068
0069   return steps;
0070 }
0071
0072 /// Test chain for the propagator
0073 template <typename device_bfield_bknd_t, typename host_bfield_bknd_t,
0074           typename detector_t>
0075 inline auto run_propagation_test(vecmem::memory_resource *mr, ::sycl::queue *q,
0076                                  detector_t &det,
0077                                  const propagator_test_config &cfg,
0078                                  typename detector_t::view_type det_view,
0079                                  covfie::field<host_bfield_bknd_t> &field) {
0080   // Create the vector of initial track parameterizations
0081   auto tracks_host = generate_tracks<generator_t>(mr, cfg.track_generator);
0082   vecmem::vector<test_track> tracks_device(tracks_host, mr);
0083
0084   // Host propagation
0085   auto host_steps =
0086       run_propagation_host(mr, det, cfg.propagation, field, tracks_host);
0087
0088   // Device propagation
0089   detray::sycl::queue_wrapper queue(q);
0090
0091   covfie::field<device_bfield_bknd_t> device_field(field);
0092   auto device_steps = run_propagation_device<device_bfield_bknd_t, detector_t>(
0093       mr, cfg.propagation, det_view, device_field, queue, tracks_device,
0094       host_steps);
0095
0096   // Check the results
0097   compare_propagation_results(host_steps, device_steps);
0098 }
0099
0100 }  // namespace detray