device/cuda/navigator_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 #include "detray/definitions/detail/cuda_definitions.hpp"
0010
0011 // Detray test include(s)
0012 #include "navigator_cuda_kernel.hpp"
0013
0014 namespace detray {
0015
0016 __global__ void navigator_test_kernel(
0017     typename detector_host_t::view_type det_data,
0018     const propagation::config prop_cfg,
0019     vecmem::data::vector_view<free_track_parameters<test_algebra>> tracks_data,
0020     vecmem::data::jagged_vector_view<dindex> volume_records_data,
0021     vecmem::data::jagged_vector_view<point3> position_records_data) {
0022   int gid = threadIdx.x + blockIdx.x * blockDim.x;
0023
0024   detector_device_t det(det_data);
0025   vecmem::device_vector<free_track_parameters<test_algebra>> tracks(
0026       tracks_data);
0027   vecmem::jagged_device_vector<dindex> volume_records(volume_records_data);
0028   vecmem::jagged_device_vector<point3> position_records(position_records_data);
0029
0030   if (gid >= tracks.size()) {
0031     return;
0032   }
0033
0034   navigator_device_t nav;
0035
0036   auto& traj = tracks.at(gid);
0037   stepper_t stepper;
0038
0039   const navigation::config& nav_cfg = prop_cfg.navigation;
0040   const stepping::config& step_cfg = prop_cfg.stepping;
0041
0042   prop_state<navigator_device_t::state> propagation{
0043       stepper_t::state{traj}, navigator_device_t::state(det)};
0044
0045   navigator_device_t::state& navigation = propagation.navigation();
0046   stepper_t::state& stepping = propagation.stepping();
0047   const auto& ctx = propagation.context();
0048
0049   // Set initial volume
0050   navigation.set_volume(0u);
0051
0052   // Start propagation and record volume IDs
0053   nav.init(stepping(), navigation, nav_cfg, ctx);
0054   bool heartbeat = navigation.is_alive();
0055   bool do_reset{true};
0056
0057   while (heartbeat) {
0058     heartbeat =
0059         heartbeat && stepper.step(navigation(), stepping, step_cfg, do_reset);
0060
0061     navigation.set_high_trust();
0062
0063     do_reset = nav.update(stepping(), navigation, nav_cfg, prop_cfg.context);
0064     do_reset = do_reset || navigation.is_on_surface();
0065     heartbeat = heartbeat && navigation.is_alive();
0066
0067     // Record volume
0068     volume_records[gid].push_back(navigation.volume());
0069     position_records[gid].push_back(stepping().pos());
0070   }
0071 }
0072
0073 void navigator_test(
0074     typename detector_host_t::view_type det_data, propagation::config& prop_cfg,
0075     vecmem::data::vector_view<free_track_parameters<test_algebra>>& tracks_data,
0076     vecmem::data::jagged_vector_view<dindex>& volume_records_data,
0077     vecmem::data::jagged_vector_view<point3>& position_records_data) {
0078   constexpr int thread_dim = 2 * WARP_SIZE;
0079   constexpr int block_dim = theta_steps * phi_steps / thread_dim + 1;
0080
0081   // run the test kernel
0082   navigator_test_kernel<<<block_dim, thread_dim>>>(
0083       det_data, prop_cfg, tracks_data, volume_records_data,
0084       position_records_data);
0085
0086   // cuda error check
0087   DETRAY_CUDA_ERROR_CHECK(cudaGetLastError());
0088   DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0089 }
0090
0091 }  // namespace detray