device/cuda/propagation_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 #include "detray/definitions/detail/cuda_definitions.hpp"
0010 #include "propagation.hpp"
0011
0012 namespace detray::tutorial {
0013
0014 // Propagation configurations
0015 inline constexpr scalar path_limit{2.f * detray::unit<scalar>::m};
0016
0017 /// Kernel that runs the entire propagation loop
0018 __global__ void propagation_kernel(
0019     typename detray::tutorial::detector_host_t::view_type det_data,
0020     typename detray::tutorial::device_field_t::view_t field_data,
0021     const vecmem::data::vector_view<detray::tutorial::track_t> tracks_data) {
0022   int gid = threadIdx.x + blockIdx.x * blockDim.x;
0023
0024   // Setup device-side track collection
0025   vecmem::device_vector<detray::tutorial::track_t> tracks(tracks_data);
0026
0027   if (gid >= tracks.size()) {
0028     return;
0029   }
0030
0031   // Setup of the device-side detector
0032   detray::tutorial::detector_device_t det(det_data);
0033
0034   // Create propagator from a stepper and a navigator
0035   propagation::config cfg{};
0036   cfg.navigation.search_window = {3u, 3u};
0037   detray::tutorial::propagator_t p{cfg};
0038
0039   // Create actor states
0040   detray::actor::pathlimit_aborter<scalar>::state aborter_state{path_limit};
0041   detray::actor::parameter_updater_state<algebra_t> updater_state{cfg};
0042   detray::actor::pointwise_material_interactor<algebra_t>::state
0043       interactor_state{};
0044
0045   auto actor_states =
0046       detray::tie(aborter_state, updater_state, interactor_state);
0047
0048   // Create the propagator state for the track
0049   detray::tutorial::propagator_t::state state(tracks[gid], field_data, det);
0050
0051   // Run propagation
0052   p.propagate(state, actor_states);
0053 }
0054
0055 void propagation(typename detray::tutorial::detector_host_t::view_type det_data,
0056                  typename detray::tutorial::device_field_t::view_t field_data,
0057                  const vecmem::data::vector_view<track_t> tracks_data) {
0058   int thread_dim = 2 * WARP_SIZE;
0059   int block_dim = tracks_data.size() / thread_dim + 1;
0060
0061   // run the tutorial kernel
0062   propagation_kernel<<<block_dim, thread_dim>>>(det_data, field_data,
0063                                                 tracks_data);
0064
0065   // cuda error check
0066   DETRAY_CUDA_ERROR_CHECK(cudaGetLastError());
0067   DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0068 }
0069
0070 }  // namespace detray::tutorial