diff --git a/tests/benchmarks/include/detray/benchmarks/device/hip/propagation_benchmark.hip b/tests/benchmarks/include/detray/benchmarks/device/hip/propagation_benchmark.hip index e9ba88f01..e3b971048 100644 --- a/tests/benchmarks/include/detray/benchmarks/device/hip/propagation_benchmark.hip +++ b/tests/benchmarks/include/detray/benchmarks/device/hip/propagation_benchmark.hip @@ -28,14 +28,14 @@ __global__ void __launch_bounds__(256, 4) propagator_benchmark_kernel( using actor_chain_t = typename propagator_t::actor_chain_type; using propagator_device_t = propagator, actor_chain_t>; + navigator, actor_chain_t>; const detector_device_t det(det_view); const vecmem::device_vector> tracks( tracks_view); int gid = threadIdx.x + blockIdx.x * blockDim.x; - if (gid >= static_cast(tracks.size())) { + if (gid >= tracks.size()) { return; } @@ -61,7 +61,7 @@ __global__ void __launch_bounds__(256, 4) propagator_benchmark_kernel( if constexpr (kOPT == detray::benchmarks::propagation_opt::e_unsync) { p.propagate(p_state, actor_state_refs); } else if constexpr (kOPT == detray::benchmarks::propagation_opt::e_sync) { - /* Do nothing for now */ + p.propagate_sync(p_state, actor_state_refs); } } @@ -73,12 +73,13 @@ typename propagator_t::actor_chain_type::state_tuple *setup_actor_states( using actor_state_t = typename propagator_t::actor_chain_type::state_tuple; actor_state_t *device_actor_state_ptr{nullptr}; - [[maybe_unused]] hipError_t success = + hipError_t success = hipMalloc((void **)&device_actor_state_ptr, sizeof(actor_state_t)); assert(success == hipSuccess); + success = hipMemcpy(device_actor_state_ptr, input_actor_states, - sizeof(actor_state_t), hipMemcpyHostToDevice); + sizeof(actor_state_t), hipMemcpyHostToDevice); assert(success == hipSuccess); return device_actor_state_ptr; @@ -107,9 +108,9 @@ void run_propagation_kernel( int block_dim = (n_samples + thread_dim - 1) / thread_dim; // run the test kernel - hipLaunchKernelGGL((propagator_benchmark_kernel), - dim3(block_dim), dim3(thread_dim), 0, 0, cfg, det_view, - field_view, device_actor_state_ptr, tracks_view); + hipLaunchKernelGGL((propagator_benchmark_kernel), dim3(block_dim), dim3(thread_dim),0,0 , + cfg, det_view, field_view, device_actor_state_ptr, tracks_view); + // hip error check DETRAY_HIP_ERROR_CHECK(hipGetLastError()); @@ -117,28 +118,28 @@ void run_propagation_kernel( } /// Macro declaring the template instantiations for the different detector types -#define DECLARE_PROPAGATION_BENCHMARK(METADATA, CHAIN, FIELD, OPT) \ - \ - template void \ +#define DECLARE_PROPAGATION_BENCHMARK(METADATA, CHAIN, FIELD, OPT) \ + \ + template void \ run_propagation_kernel, OPT>( \ - const propagation::config &, detector::view_type, \ - covfie::field_view, \ + const propagation::config &, detector::view_type, \ + covfie::field_view, \ hip_propagator_type::actor_chain_type::state_tuple *, \ - vecmem::data::vector_view< \ - free_track_parameters::algebra_type>>, \ - const int); \ - \ + CHAIN>::actor_chain_type::state_tuple *, \ + vecmem::data::vector_view< \ + free_track_parameters::algebra_type>>, \ + const int); \ + \ template hip_propagator_type::actor_chain_type::state_tuple * \ + CHAIN>::actor_chain_type::state_tuple * \ setup_actor_states>( \ hip_propagator_type::actor_chain_type::state_tuple *); \ - \ - template void \ + CHAIN>::actor_chain_type::state_tuple *); \ + \ + template void \ release_actor_states>( \ hip_propagator_type::actor_chain_type::state_tuple *); + CHAIN>::actor_chain_type::state_tuple *); DECLARE_PROPAGATION_BENCHMARK(benchmarks::default_metadata, empty_chain, const_field_t, propagation_opt::e_unsync) diff --git a/tests/benchmarks/include/detray/benchmarks/device/hip/propagation_benchmark.hpp b/tests/benchmarks/include/detray/benchmarks/device/hip/propagation_benchmark.hpp index fb903a2ae..d99abfdac 100644 --- a/tests/benchmarks/include/detray/benchmarks/device/hip/propagation_benchmark.hpp +++ b/tests/benchmarks/include/detray/benchmarks/device/hip/propagation_benchmark.hpp @@ -9,12 +9,11 @@ // Project include(s) #include "detray/definitions/algebra.hpp" -#include "detray/navigation/caching_navigator.hpp" +#include "detray/navigation/navigator.hpp" #include "detray/propagator/actors.hpp" #include "detray/propagator/propagator.hpp" #include "detray/propagator/rk_stepper.hpp" #include "detray/tracks/tracks.hpp" -#include "detray/utils/logging.hpp" // Detray test include(s) #include "detray/test/common/bfield.hpp" @@ -59,7 +58,7 @@ template , typename detector::algebra_type>, - caching_navigator>, + navigator>, actor_chain_t::algebra_type>>; /// Launch the propagation kernelfor benchmarking @@ -162,9 +161,9 @@ struct hip_propagation_bm : public benchmark_base { m_cfg.propagation(), det_view, *bfield, device_actor_state_ptr, warmup_track_buffer, math::min(n_warmup, n_samples)); } else { - DETRAY_WARN_HOST( - "Running HIP benchmarks without warmup is " - "not recommended"); + std::cout << "WARNING: Running HIP benchmarks without warmup is " + "not recommended" + << std::endl; } // Calculate the propagation rate diff --git a/tests/include/detray/test/device/CMakeLists.txt b/tests/include/detray/test/device/CMakeLists.txt index 8f5daf47c..e258b03a0 100644 --- a/tests/include/detray/test/device/CMakeLists.txt +++ b/tests/include/detray/test/device/CMakeLists.txt @@ -1,6 +1,6 @@ # Detray library, part of the ACTS project (R&D line) # -# (c) 2024 CERN for the benefit of the ACTS project +# (c) 2025 CERN for the benefit of the ACTS project # # Mozilla Public License Version 2.0 @@ -22,3 +22,7 @@ target_link_libraries( if(DETRAY_BUILD_CUDA) add_subdirectory(cuda) endif() + +if(DETRAY_BUILD_HIP) + add_subdirectory(hip) +endif() diff --git a/tests/include/detray/test/device/hip/CMakeLists.txt b/tests/include/detray/test/device/hip/CMakeLists.txt new file mode 100644 index 000000000..94b2e9d28 --- /dev/null +++ b/tests/include/detray/test/device/hip/CMakeLists.txt @@ -0,0 +1,43 @@ +# Detray library, part of the ACTS project (R&D line) +# +# (c) 2025 CERN for the benefit of the ACTS project +# +# Mozilla Public License Version 2.0 + +# C++17 support for HIP requires CMake 3.21. +cmake_minimum_required(VERSION 3.21) # HIP langauge support requires minimum 3.21 + +find_package(HIPToolkit) + +# Enable HIP as a language. +enable_language(HIP) + +# Set the HIP build flags. +include(detray-compiler-options-hip) + +# Set up a test library, which the "new style" benchmarks and tests could use. +add_library( + detray_test_hip_${CMAKE_HIP_PLATFORM} + STATIC + "material_validation.hpp" + "material_validation.hip" + "navigation_validation.hpp" + "navigation_validation.hip" +) + +add_library( + detray::test_hip_${CMAKE_HIP_PLATFORM} + ALIAS detray_test_hip_${CMAKE_HIP_PLATFORM} +) + +target_link_libraries( + detray_test_hip_${CMAKE_HIP_PLATFORM} + PUBLIC + HIP::hiprt + vecmem::hip + #covfie::hip + detray::core_array + detray::test_device + detray::test_cpu + detray::validation_utils +) diff --git a/tests/include/detray/test/device/hip/material_validation.hip b/tests/include/detray/test/device/hip/material_validation.hip new file mode 100644 index 000000000..bd50cabe9 --- /dev/null +++ b/tests/include/detray/test/device/hip/material_validation.hip @@ -0,0 +1,125 @@ +/** Detray library, part of the ACTS project (R&D line) + * + * (c) 2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#include "detray/definitions/detail/hip_definitions.hpp" +#include "detray/propagator/actors.hpp" +#include "detray/propagator/line_stepper.hpp" +#include "material_validation.hpp" + +namespace detray::hip { + +template +__global__ void material_validation_kernel( + typename detector_t::view_type det_data, const propagation::config cfg, + vecmem::data::vector_view< + free_track_parameters> + tracks_view, + vecmem::data::vector_view< + material_validator::material_record> + mat_records_view, + vecmem::data::jagged_vector_view< + material_validator::material_params> + mat_steps_view) { + + using detector_device_t = + detector; + using algebra_t = typename detector_device_t::algebra_type; + using scalar_t = dscalar; + + using stepper_t = line_stepper; + using navigator_t = navigator; + // Propagator with full covariance transport, pathlimit aborter and + // material tracer + using material_tracer_t = + material_validator::material_tracer; + using pathlimit_aborter_t = pathlimit_aborter; + using actor_chain_t = + actor_chain, + parameter_resetter, + pointwise_material_interactor, + material_tracer_t>; + using propagator_t = propagator; + + detector_device_t det(det_data); + + vecmem::device_vector> tracks(tracks_view); + vecmem::device_vector + mat_records(mat_records_view); + vecmem::jagged_device_vector< + typename material_tracer_t::material_params_type> + mat_steps(mat_steps_view); + + int trk_id = threadIdx.x + blockIdx.x * blockDim.x; + if (trk_id >= tracks.size()) { + return; + } + + propagator_t p{cfg}; + + // Create the actor states + typename pathlimit_aborter_t::state aborter_state{cfg.stepping.path_limit}; + typename pointwise_material_interactor::state interactor_state{}; + typename material_tracer_t::state mat_tracer_state{mat_steps.at(trk_id)}; + + auto actor_states = + ::detray::tie(aborter_state, interactor_state, mat_tracer_state); + + // Run propagation + typename navigator_t::state::view_type nav_view{}; + typename propagator_t::state propagation(tracks[trk_id], det, nav_view); + + p.propagate(propagation, actor_states); + + // Record the accumulated material + assert(mat_records.size() == tracks.size()); + mat_records.at(trk_id) = mat_tracer_state.get_material_record(); +} + +/// Launch the device kernel +template +void material_validation_device( + typename detector_t::view_type det_view, const propagation::config &cfg, + vecmem::data::vector_view< + free_track_parameters> &tracks_view, + vecmem::data::vector_view< + material_validator::material_record> + &mat_records_view, + vecmem::data::jagged_vector_view< + material_validator::material_params> + &mat_steps_view) { + + constexpr int thread_dim = 2 * WARP_SIZE; + int block_dim = tracks_view.size() / thread_dim + 1; + + // run the test kernel + hipLaunchKernelGGL((material_validation_kernel) , dim3(block_dim) , dim3(thread_dim) , 0 , 0 , + det_view, cfg, tracks_view, mat_records_view, mat_steps_view); + + + // hip error check + DETRAY_HIP_ERROR_CHECK(hipGetLastError()); + DETRAY_HIP_ERROR_CHECK(hipDeviceSynchronize()); +} + +/// Macro declaring the template instantiations for the different detector types +#define DECLARE_MATERIAL_VALIDATION(METADATA) \ + \ + template void material_validation_device>( \ + typename detector::view_type, const propagation::config &, \ + vecmem::data::vector_view< \ + free_track_parameters::algebra_type>> \ + &, \ + vecmem::data::vector_view::scalar_type>> &, \ + vecmem::data::jagged_vector_view::scalar_type>> &); + +DECLARE_MATERIAL_VALIDATION(test::default_metadata) +DECLARE_MATERIAL_VALIDATION(test::toy_metadata) +DECLARE_MATERIAL_VALIDATION(test::default_telescope_metadata) + +} // namespace detray::hip diff --git a/tests/include/detray/test/device/hip/material_validation.hpp b/tests/include/detray/test/device/hip/material_validation.hpp new file mode 100644 index 000000000..5f4b9ea2c --- /dev/null +++ b/tests/include/detray/test/device/hip/material_validation.hpp @@ -0,0 +1,110 @@ +/** Detray library, part of the ACTS project (R&D line) + * + * (c) 2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s) +#include "detray/core/detector.hpp" +#include "detray/tracks/tracks.hpp" + +// Detray test include(s) +#include "detray/test/cpu/material_validation.hpp" +#include "detray/test/validation/material_validation_utils.hpp" + +// Vecmem include(s) +#include +#include +#include +#include + +// System include +#include + +namespace detray::hip { + +/// Launch the material validation kernel +/// +/// @param[in] det_view the detector vecmem view +/// @param[in] cfg the propagation configuration +/// @param[in] tracks_view the initial track parameter of every test track +/// @param[out] mat_records_view the accumulated material per track +template +void material_validation_device( + typename detector_t::view_type det_view, const propagation::config &cfg, + vecmem::data::vector_view< + free_track_parameters> &tracks_view, + vecmem::data::vector_view< + material_validator::material_record> + &mat_records_view, + vecmem::data::jagged_vector_view< + material_validator::material_params> + &mat_steps_view); + +/// Prepare data for device material trace run +struct run_material_validation { + + static constexpr std::string_view name = "hip"; + + template + auto operator()( + vecmem::memory_resource *host_mr, vecmem::memory_resource *dev_mr, + const detector_t &det, const propagation::config &cfg, + const vecmem::vector< + free_track_parameters> &tracks, + const std::vector &capacities) { + + using algebra_t = typename detector_t::algebra_type; + using scalar_t = dscalar; + using track_t = free_track_parameters; + using material_record_t = material_validator::material_record; + using material_params_t = material_validator::material_params; + + // Helper object for performing memory copies (to HIP devices) + vecmem::hip::copy hip_cpy; + + // Copy the detector to device and get its view + auto det_buffer = detray::get_buffer(det, *dev_mr, hip_cpy); + auto det_view = detray::get_data(det_buffer); + + // Move the track parameters to device + auto tracks_buffer = hip_cpy.to(vecmem::get_data(tracks), *dev_mr, + vecmem::copy::type::host_to_device); + vecmem::data::vector_view tracks_view = + vecmem::get_data(tracks_buffer); + + vecmem::data::vector_buffer mat_records_buffer( + static_cast(tracks.size()), *dev_mr, + vecmem::data::buffer_type::fixed_size); + hip_cpy.setup(mat_records_buffer)->wait(); + auto mat_records_view = vecmem::get_data(mat_records_buffer); + + // Buffer for the material parameters at every surface per track + vecmem::data::jagged_vector_buffer mat_steps_buffer( + capacities, *dev_mr, host_mr, vecmem::data::buffer_type::resizable); + hip_cpy.setup(mat_steps_buffer)->wait(); + auto mat_steps_view = vecmem::get_data(mat_steps_buffer); + + // Run the material tracing on device + material_validation_device( + det_view, cfg, tracks_view, mat_records_view, mat_steps_view); + + // Get the results back to the host and pass them on to be checked + vecmem::vector mat_records(host_mr); + hip_cpy(mat_records_buffer, mat_records)->wait(); + + vecmem::jagged_vector mat_steps(host_mr); + hip_cpy(mat_steps_buffer, mat_steps)->wait(); + + return std::make_tuple(mat_records, mat_steps); + } +}; + +template +using material_validation = detray::test::material_validation_impl< + detector_t, detray::hip::run_material_validation>; + +} // namespace detray::hip diff --git a/tests/include/detray/test/device/hip/navigation_validation.hip b/tests/include/detray/test/device/hip/navigation_validation.hip new file mode 100644 index 000000000..fcaf52849 --- /dev/null +++ b/tests/include/detray/test/device/hip/navigation_validation.hip @@ -0,0 +1,205 @@ +/** Detray library, part of the ACTS project (R&D line) + * + * (c) 2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#include "detray/definitions/detail/hip_definitions.hpp" +#include "navigation_validation.hpp" + +namespace detray::hip { + +template +__global__ void navigation_validation_kernel( + typename detector_t::view_type det_data, const propagation::config cfg, + pdg_particle ptc_hypo, + bfield_t field_data, + vecmem::data::jagged_vector_view + truth_intersection_traces_view, + vecmem::data::jagged_vector_view> + recorded_intersections_view, + vecmem::data::vector_view< + material_validator::material_record> + mat_records_view, + vecmem::data::jagged_vector_view< + material_validator::material_params> + mat_steps_view) { + + using detector_device_t = + detector; + using algebra_t = typename detector_device_t::algebra_type; + using scalar_t = dscalar; + + static_assert(std::is_same_v, + "Host and device detector view types do not match"); + + using hom_bfield_view_t = typename bfield::const_field_t::view_t; + using rk_stepper_t = rk_stepper; + using line_stepper_t = line_stepper; + // Use RK-stepper when a non-empty b-field was passed + static constexpr auto is_no_bfield{ + std::is_same_v}; + using stepper_t = + std::conditional_t; + + // Inspector that records all encountered surfaces + using intersection_t = typename intersection_record_t::intersection_type; + using object_tracer_t = + navigation::object_tracer; + // Navigation with inspection + using navigator_t = + navigator; + + // Propagator with pathlimit aborter + using material_tracer_t = + material_validator::material_tracer; + using pathlimit_aborter_t = pathlimit_aborter; + using actor_chain_t = actor_chain; + using propagator_t = propagator; + + detector_device_t det(det_data); + + vecmem::jagged_device_vector + truth_intersection_traces(truth_intersection_traces_view); + vecmem::jagged_device_vector< + navigation::detail::candidate_record> + recorded_intersections(recorded_intersections_view); + vecmem::device_vector + mat_records(mat_records_view); + vecmem::jagged_device_vector< + typename material_tracer_t::material_params_type> + mat_steps(mat_steps_view); + + // Check the memory setup + assert(truth_intersection_traces.size() == + recorded_intersections_view.size()); + + int trk_id = threadIdx.x + blockIdx.x * blockDim.x; + if (trk_id >= truth_intersection_traces.size()) { + return; + } + + propagator_t p{cfg}; + + // Create the actor states + typename pathlimit_aborter_t::state aborter_state{cfg.stepping.path_limit}; + typename material_tracer_t::state mat_tracer_state{mat_steps.at(trk_id)}; + auto actor_states = ::detray::tie(aborter_state, mat_tracer_state); + + // Get the initial track parameters + const auto &track = truth_intersection_traces[trk_id].front().track_param; + + // Save the initial intersection, since it is not recorded by the + // object tracer + assert(recorded_intersections.at(trk_id).empty()); + recorded_intersections.at(trk_id).push_back( + {track.pos(), track.dir(), + truth_intersection_traces[trk_id].front().intersection}); + // Did the insertion of an element work? + assert(recorded_intersections.at(trk_id).size() == 1); + + // Run propagation + if constexpr (is_no_bfield) { + typename propagator_t::state propagation( + track, det, + typename navigator_t::state::view_type{ + recorded_intersections_view.ptr()[trk_id]}); + propagation.set_particle(update_particle_hypothesis(ptc_hypo, track)); + + p.propagate(propagation, actor_states); + } else { + typename propagator_t::state propagation( + track, field_data, det, + typename navigator_t::state::view_type{ + recorded_intersections_view.ptr()[trk_id]}); + propagation.set_particle(update_particle_hypothesis(ptc_hypo, track)); + + p.propagate(propagation, actor_states); + } + + // Record the accumulated material + assert(truth_intersection_traces.size() == mat_records.size()); + mat_records.at(trk_id) = mat_tracer_state.get_material_record(); +} + +/// Launch the device kernel +template +void navigation_validation_device( + typename detector_t::view_type det_view, const propagation::config &cfg, + pdg_particle ptc_hypo, + bfield_t field_data, + vecmem::data::jagged_vector_view + &truth_intersection_traces_view, + vecmem::data::jagged_vector_view> + &recorded_intersections_view, + vecmem::data::vector_view< + material_validator::material_record> + &mat_records_view, + vecmem::data::jagged_vector_view< + material_validator::material_params> + &mat_steps_view) { + + constexpr int thread_dim = 2 * WARP_SIZE; + int block_dim = truth_intersection_traces_view.size() / thread_dim + 1; + + // run the test kernel + hipLaunchKernelGGL((navigation_validation_kernel) , + dim3(block_dim), dim3(thread_dim) , 0 , 0 , + det_view, cfg, ptc_hypo, field_data, truth_intersection_traces_view, + recorded_intersections_view, mat_records_view, mat_steps_view); + // hip error check + DETRAY_HIP_ERROR_CHECK(hipGetLastError()); + DETRAY_HIP_ERROR_CHECK(hipDeviceSynchronize()); +} + +/// Macro declaring the template instantiations for the different detector types +#define DECLARE_NAVIGATION_VALIDATION(METADATA) \ + \ + template void navigation_validation_device< \ + covfie::field_view< \ + bfield::const_bknd_t>>, \ + detector, detray::intersection_record>>( \ + typename detector::view_type, const propagation::config &, \ + pdg_particle::scalar_type>, \ + covfie::field_view< \ + bfield::const_bknd_t>>, \ + vecmem::data::jagged_vector_view< \ + const detray::intersection_record>> &, \ + vecmem::data::jagged_vector_view>::intersection_type>> &, \ + vecmem::data::vector_view::scalar_type>> &, \ + vecmem::data::jagged_vector_view::scalar_type>> &); \ + \ + template void navigation_validation_device< \ + detray::navigation_validator::empty_bfield, detector, \ + detray::intersection_record>>( \ + typename detector::view_type, const propagation::config &, \ + pdg_particle::scalar_type>, \ + detray::navigation_validator::empty_bfield, \ + vecmem::data::jagged_vector_view< \ + const detray::intersection_record>> &, \ + vecmem::data::jagged_vector_view>::intersection_type>> &, \ + vecmem::data::vector_view::scalar_type>> &, \ + vecmem::data::jagged_vector_view::scalar_type>> &); + +DECLARE_NAVIGATION_VALIDATION(test::default_metadata) +DECLARE_NAVIGATION_VALIDATION(test::toy_metadata) +DECLARE_NAVIGATION_VALIDATION(test::default_telescope_metadata) + +} // namespace detray::hip diff --git a/tests/include/detray/test/device/hip/navigation_validation.hpp b/tests/include/detray/test/device/hip/navigation_validation.hpp new file mode 100644 index 000000000..b1ce3002e --- /dev/null +++ b/tests/include/detray/test/device/hip/navigation_validation.hpp @@ -0,0 +1,504 @@ +/** Detray library, part of the ACTS project (R&D line) + * + * (c) 2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s) +#include "detray/core/detector.hpp" +#include "detray/definitions/pdg_particle.hpp" +#include "detray/propagator/line_stepper.hpp" +#include "detray/propagator/rk_stepper.hpp" +#include "detray/tracks/ray.hpp" +#include "detray/tracks/tracks.hpp" + +// Detray test include(s) +#include "detray/test/common/bfield.hpp" +#include "detray/test/framework/fixture_base.hpp" +#include "detray/test/framework/whiteboard.hpp" +#include "detray/test/utils/inspectors.hpp" +#include "detray/test/validation/detector_scan_utils.hpp" +#include "detray/test/validation/detector_scanner.hpp" +#include "detray/test/validation/material_validation_utils.hpp" +#include "detray/test/validation/navigation_validation_config.hpp" +#include "detray/test/validation/navigation_validation_utils.hpp" + +// Vecmem include(s) +#include +#include +#include +#include + +// System include(s) +#include +#include + +namespace detray::hip { + +/// Launch the navigation validation kernel +/// +/// @param[in] det_view the detector vecmem view +/// @param[in] cfg the propagation configuration +/// @param[in] field_data the magentic field view (maybe an empty field) +/// @param[in] truth_intersection_traces_view vecemem view of the truth data +/// @param[out] recorded_intersections_view vecemem view of the intersections +/// recorded by the navigator +template +void navigation_validation_device( + typename detector_t::view_type det_view, const propagation::config &cfg, + pdg_particle ptc_hypo, + bfield_t field_data, + vecmem::data::jagged_vector_view + &truth_intersection_traces_view, + vecmem::data::jagged_vector_view> + &recorded_intersections_view, + vecmem::data::vector_view< + material_validator::material_record> + &mat_records_view, + vecmem::data::jagged_vector_view< + material_validator::material_params> + &mat_steps_view); + +/// Prepare data for device navigation run +template +inline auto run_navigation_validation( + vecmem::memory_resource *host_mr, vecmem::memory_resource *dev_mr, + const detector_t &det, const propagation::config &cfg, + pdg_particle ptc_hypo, + bfield_t field_data, + const std::vector> + &truth_intersection_traces) { + + using scalar_t = dscalar; + using intersection_t = typename intersection_record_t::intersection_type; + using material_record_t = material_validator::material_record; + using material_params_t = material_validator::material_params; + + // Helper object for performing memory copies (to HIP devices) + vecmem::hip::copy hip_cpy; + + // Copy the detector to device and get its view + auto det_buffer = detray::get_buffer(det, *dev_mr, hip_cpy); + auto det_view = detray::get_data(det_buffer); + + // Move truth intersection traces data to device + auto truth_intersection_traces_data = + vecmem::get_data(truth_intersection_traces, host_mr); + auto truth_intersection_traces_buffer = + hip_cpy.to(truth_intersection_traces_data, *dev_mr, host_mr, + vecmem::copy::type::host_to_device); + vecmem::data::jagged_vector_view + truth_intersection_traces_view = + vecmem::get_data(truth_intersection_traces_buffer); + + // Buffer for the intersections recorded by the navigator + std::vector capacities; + for (const auto &trace : truth_intersection_traces) { + // Increase the capacity, in case the navigator finds more surfaces + // than the truth intersections (usually just one) + capacities.push_back(trace.size() + 10u); + } + + vecmem::data::jagged_vector_buffer< + navigation::detail::candidate_record> + recorded_intersections_buffer(capacities, *dev_mr, host_mr, + vecmem::data::buffer_type::resizable); + hip_cpy.setup(recorded_intersections_buffer)->wait(); + auto recorded_intersections_view = + vecmem::get_data(recorded_intersections_buffer); + + vecmem::data::vector_buffer mat_records_buffer( + static_cast(truth_intersection_traces_view.size()), + *dev_mr, vecmem::data::buffer_type::fixed_size); + hip_cpy.setup(mat_records_buffer)->wait(); + auto mat_records_view = vecmem::get_data(mat_records_buffer); + + // Buffer for the material parameters at every step per track + vecmem::data::jagged_vector_buffer mat_steps_buffer( + capacities, *dev_mr, host_mr, vecmem::data::buffer_type::resizable); + hip_cpy.setup(mat_steps_buffer)->wait(); + auto mat_steps_view = vecmem::get_data(mat_steps_buffer); + + // Run the navigation validation test on device + navigation_validation_device( + det_view, cfg, ptc_hypo, field_data, truth_intersection_traces_view, + recorded_intersections_view, mat_records_view, mat_steps_view); + + // Get the results back to the host and pass them on to the checking + vecmem::jagged_vector> + recorded_intersections(host_mr); + hip_cpy(recorded_intersections_buffer, recorded_intersections)->wait(); + + vecmem::vector mat_records(host_mr); + hip_cpy(mat_records_buffer, mat_records)->wait(); + + vecmem::jagged_vector mat_steps(host_mr); + hip_cpy(mat_steps_buffer, mat_steps)->wait(); + + return std::make_tuple(std::move(recorded_intersections), + std::move(mat_records), std::move(mat_steps)); +} + +/// @brief Test class that runs the navigation validation for a given detector +/// on device. +/// +/// @note The lifetime of the detector needs to be guaranteed outside this class +template class scan_type> +class navigation_validation : public test::fixture_base<> { + + using algebra_t = typename detector_t::algebra_type; + using scalar_t = dscalar; + using vector3_t = dvector3D; + using free_track_parameters_t = free_track_parameters; + using trajectory_type = typename scan_type::trajectory_type; + using intersection_trace_t = typename scan_type< + algebra_t>::template intersection_trace_type; + + /// Switch between rays and helices + static constexpr auto k_use_rays{ + std::is_same_v, trajectory_type>}; + + public: + using fixture_type = test::fixture_base<>; + using config = detray::test::navigation_validation_config; + + explicit navigation_validation( + const detector_t &det, const typename detector_t::name_map &names, + const config &cfg = {}, std::shared_ptr wb = nullptr, + const typename detector_t::geometry_context gctx = {}) + : m_cfg{cfg}, + m_gctx{gctx}, + m_det{det}, + m_names{names}, + m_whiteboard{wb} { + + if (!m_whiteboard) { + throw std::invalid_argument("No white board was passed to " + + m_cfg.name() + " test"); + } + + // Use ray or helix + const std::string det_name{m_det.name(m_names)}; + m_truth_data_name = k_use_rays ? det_name + "_ray_scan_for_hip" + : det_name + "_helix_scan_for_hip"; + + // Pin the data onto the whiteboard + if (!m_whiteboard->exists(m_truth_data_name) && + io::file_exists(m_cfg.intersection_file()) && + io::file_exists(m_cfg.track_param_file())) { + + // Name clash: Choose alternative name + if (m_whiteboard->exists(m_truth_data_name)) { + m_truth_data_name = io::alt_file_name(m_truth_data_name); + } + + std::vector intersection_traces; + + std::cout << "\nINFO: Reading data from file..." << std::endl; + + // Fill the intersection traces from file + detray::detector_scanner::read(m_cfg.intersection_file(), + m_cfg.track_param_file(), + intersection_traces); + + m_whiteboard->add(m_truth_data_name, + std::move(intersection_traces)); + } else if (m_whiteboard->exists(m_truth_data_name)) { + std::cout << "\nINFO: Fetching data from white board..." + << std::endl; + } else { + throw std::invalid_argument( + "Navigation validation: Could not find data files"); + } + + // Check that data is ready + if (!m_whiteboard->exists(m_truth_data_name)) { + throw std::invalid_argument( + "Data for navigation check is not on the whiteboard"); + } + } + + /// Run the check + void TestBody() override { + using namespace detray; + using namespace navigation; + + // Runge-Kutta stepper + using hom_bfield_t = bfield::const_field_t; + using bfield_view_t = + std::conditional_t; + using bfield_t = + std::conditional_t; + using intersection_t = + typename intersection_trace_t::value_type::intersection_type; + + bfield_t b_field{}; + if constexpr (!k_use_rays) { + b_field = create_const_field(m_cfg.B_vector()); + } + + // Fetch the truth data + auto &truth_intersection_traces = + m_whiteboard->template get>( + m_truth_data_name); + ASSERT_EQ(m_cfg.n_tracks(), truth_intersection_traces.size()); + + std::cout << "\nINFO: Running device navigation validation on: " + << m_det.name(m_names) << "...\n" + << std::endl; + + const std::string det_name{m_det.name(m_names)}; + const std::string prefix{k_use_rays ? det_name + "_ray_" + : det_name + "_helix_"}; + + std::ios_base::openmode io_mode = std::ios::trunc | std::ios::out; + const std::string debug_file_name{prefix + + "navigation_validation_hip.txt"}; + detray::io::file_handle debug_file{debug_file_name, io_mode}; + + // Run the propagation on device and record the navigation data + auto [recorded_intersections, mat_records, mat_steps] = + run_navigation_validation( + &m_host_mr, &m_dev_mr, m_det, m_cfg.propagation(), + m_cfg.ptc_hypothesis(), b_field, truth_intersection_traces); + + // Collect some statistics + std::size_t n_tracks{0u}; + std::size_t n_matching_error{0u}; + std::size_t n_fatal{0u}; + // Total number of encountered surfaces + navigation_validator::surface_stats n_surfaces{}; + // Missed by navigator + navigation_validator::surface_stats n_miss_nav{}; + // Missed by truth finder + navigation_validator::surface_stats n_miss_truth{}; + + std::vector>> + missed_intersections{}; + + EXPECT_EQ(recorded_intersections.size(), + truth_intersection_traces.size()); + + scalar_t min_pT{std::numeric_limits::max()}; + scalar_t max_pT{-std::numeric_limits::max()}; + for (std::size_t i = 0u; i < truth_intersection_traces.size(); ++i) { + auto &truth_trace = truth_intersection_traces[i]; + auto &recorded_trace = recorded_intersections[i]; + + if (n_tracks >= m_cfg.n_tracks()) { + break; + } + + // Get the original test trajectory (ray or helix) + const auto &start = truth_trace.front(); + const auto &trck_param = start.track_param; + trajectory_type test_traj = get_parametrized_trajectory(trck_param); + + const scalar q = start.charge; + const scalar pT{q == 0.f ? 1.f * unit::GeV + : trck_param.pT(q)}; + const scalar p{q == 0.f ? 1.f * unit::GeV + : trck_param.p(q)}; + + if (detray::detail::is_invalid_value(m_cfg.p_range()[0])) { + min_pT = std::min(min_pT, pT); + max_pT = std::max(max_pT, pT); + } else { + min_pT = m_cfg.p_range()[0]; + max_pT = m_cfg.p_range()[1]; + } + + // Recorded only the start position, which added by default + bool success{true}; + if (truth_trace.size() == 1) { + // Propagation did not succeed + success = false; + std::vector missed_inters{}; + missed_intersections.push_back( + std::make_pair(test_traj, missed_inters)); + + ++n_fatal; + } else { + // Adjust the track charge, which is unknown to the navigation + for (auto &record : recorded_trace) { + record.charge = q; + record.p_mag = p; + } + + // Compare truth and recorded data elementwise + auto [result, n_missed_nav, n_missed_truth, n_error, + missed_inters] = + navigation_validator::compare_traces( + m_cfg, truth_trace, recorded_trace, test_traj, n_tracks, + &(*debug_file)); + + missed_intersections.push_back( + std::make_pair(test_traj, std::move(missed_inters))); + + // Update statistics + success &= result; + n_miss_nav += n_missed_nav; + n_miss_truth += n_missed_truth; + n_matching_error += n_error; + } + + if (!success) { + detector_scanner::display_error( + m_gctx, m_det, m_names, m_cfg.name(), test_traj, + truth_trace, m_cfg.svg_style(), n_tracks, m_cfg.n_tracks(), + recorded_trace); + } + + EXPECT_TRUE(success) + << "\nINFO: Wrote navigation debugging data in: " + << debug_file_name; + + ++n_tracks; + + // After dummy records insertion, traces should have the same size + ASSERT_EQ(truth_trace.size(), recorded_trace.size()); + + // Count the number of different surface types on this trace + navigation_validator::surface_stats n_truth{}; + navigation_validator::surface_stats n_nav{}; + for (std::size_t j = 0; j < truth_trace.size(); ++j) { + const auto truth_desc = truth_trace[j].intersection.sf_desc; + const auto rec_desc = recorded_trace[j].intersection.sf_desc; + + // Exclude dummy records for missing surfaces + if (!truth_desc.barcode().is_invalid()) { + n_truth.count(truth_desc); + } + if (!rec_desc.barcode().is_invalid()) { + n_nav.count(rec_desc); + } + } + + // Take max count, since either trace might have skipped surfaces + const std::size_t n_portals{ + math::max(n_truth.n_portals, n_nav.n_portals)}; + const std::size_t n_sensitives{ + math::max(n_truth.n_sensitives, n_nav.n_sensitives)}; + const std::size_t n_passives{ + math::max(n_truth.n_passives, n_nav.n_passives)}; + const std::size_t n{n_portals + n_sensitives + n_passives}; + + // Cannot have less surfaces than truth intersections after matching + // (Don't count first entry, which records the initial track params) + ASSERT_TRUE(n >= (truth_trace.size() - 1u)); + + n_surfaces.n_portals += n_portals; + n_surfaces.n_sensitives += n_sensitives; + n_surfaces.n_passives += n_passives; + } + + // Calculate and display the result + navigation_validator::print_efficiency(n_tracks, n_surfaces, n_miss_nav, + n_miss_truth, n_fatal, + n_matching_error); + + // Print track positions for plotting + std::string momentum_str{""}; + if constexpr (!k_use_rays) { + momentum_str = + "_" + + std::to_string(std::floor(10. * static_cast(min_pT)) / + 10.) + + "_" + + std::to_string(std::ceil(10. * static_cast(max_pT)) / + 10.) + + "_GeV"; + } + + const auto data_path{ + std::filesystem::path{m_cfg.track_param_file()}.parent_path()}; + + // Create an output file path + auto make_path = [&data_path, &prefix, + &momentum_str](const std::string &name) { + return data_path / (prefix + name + momentum_str + ".csv"); + }; + + const auto truth_trk_path{make_path("truth_track_params_hip")}; + const auto trk_path{make_path("navigation_track_params_hip")}; + const auto truth_intr_path{make_path("truth_intersections_hip")}; + const auto intr_path{make_path("navigation_intersections_hip")}; + const auto mat_path{make_path("accumulated_material_hip")}; + const auto missed_path{make_path("missed_intersections_dists_hip")}; + + // Write the distance of the missed intersection local position + // to the surface boundaries to file for plotting + navigation_validator::write_dist_to_boundary( + m_det, m_names, missed_path.string(), missed_intersections); + detector_scanner::write_tracks(truth_trk_path.string(), + truth_intersection_traces); + navigation_validator::write_tracks(trk_path.string(), + recorded_intersections); + detector_scanner::write_intersections(truth_intr_path.string(), + truth_intersection_traces); + detector_scanner::write_intersections(intr_path.string(), + recorded_intersections); + material_validator::write_material(mat_path.string(), mat_records); + + std::cout + << "INFO: Wrote distance to boundary of missed intersections to: " + << missed_path << std::endl; + std::cout << "INFO: Wrote track states in: " << trk_path << std::endl; + std::cout << "INFO: Wrote truth intersections in: " << truth_intr_path + << std::endl; + std::cout << "INFO: Wrote track intersections in: " << intr_path + << std::endl; + std::cout << "INFO: Wrote accumulated material in: " << mat_path + << std::endl; + } + + private: + /// @returns either the helix or ray corresponding to the input track + /// parameters @param track + trajectory_type get_parametrized_trajectory( + const free_track_parameters_t &track) { + std::unique_ptr test_traj{nullptr}; + if constexpr (k_use_rays) { + test_traj = std::make_unique(track); + } else { + test_traj = + std::make_unique(track, m_cfg.B_vector()); + } + return *(test_traj.release()); + } + + /// Vecmem memory resource for the host allocations + vecmem::host_memory_resource m_host_mr{}; + /// Vecmem memory resource for the device allocations + vecmem::hip::device_memory_resource m_dev_mr{}; + /// The configuration of this test + config m_cfg; + /// Name of the truth data collection + std::string m_truth_data_name{""}; + /// The geometry context to check + typename detector_t::geometry_context m_gctx{}; + /// The detector to be checked + const detector_t &m_det; + /// Volume names + const typename detector_t::name_map &m_names; + /// Whiteboard to pin data + std::shared_ptr m_whiteboard{nullptr}; +}; + +template +using straight_line_navigation = + detray::hip::navigation_validation; + +template +using helix_navigation = + detray::hip::navigation_validation; + +} // namespace detray::hip diff --git a/tests/tools/src/hip/CMakeLists.txt b/tests/tools/src/hip/CMakeLists.txt index 7831d6b8f..eda2806ea 100644 --- a/tests/tools/src/hip/CMakeLists.txt +++ b/tests/tools/src/hip/CMakeLists.txt @@ -15,6 +15,40 @@ include(CMakeFindDependencyMacro) find_dependency(Boost COMPONENTS program_options REQUIRED) +if(DETRAY_BUILD_TESTING) + # Build the navigation validation executable. + detray_add_executable(navigation_validation_hip_${CMAKE_HIP_PLATFORM} + "navigation_validation_hip.cpp" + LINK_LIBRARIES HIP::hiprt GTest::gtest GTest::gtest_main + Boost::program_options detray::test_hip_${CMAKE_HIP_PLATFORM} detray::tools + ) + if( + ("${CMAKE_HIP_PLATFORM}" STREQUAL "hcc") + OR ("${CMAKE_HIP_PLATFORM}" STREQUAL "amd") + ) + set_source_files_properties( + navigation_validation_hip.cpp + PROPERTIES LANGUAGE HIP + ) + endif() + + # Build the material hip validation executable. + detray_add_executable(material_validation_hip_${CMAKE_HIP_PLATFORM} + "material_validation_hip.cpp" + LINK_LIBRARIES HIP::hiprt GTest::gtest GTest::gtest_main + Boost::program_options detray::test_hip_${CMAKE_HIP_PLATFORM} detray::tools + ) + if( + ("${CMAKE_HIP_PLATFORM}" STREQUAL "hcc") + OR ("${CMAKE_HIP_PLATFORM}" STREQUAL "amd") + ) + set_source_files_properties( + material_validation_hip.cpp + PROPERTIES LANGUAGE HIP + ) + endif() +endif() + if(DETRAY_BUILD_BENCHMARKS) # Build benchmarks for multiple algebra plugins # Currently vc and smatrix is not supported on device diff --git a/tests/tools/src/hip/material_validation_hip.cpp b/tests/tools/src/hip/material_validation_hip.cpp new file mode 100644 index 000000000..8fb733c38 --- /dev/null +++ b/tests/tools/src/hip/material_validation_hip.cpp @@ -0,0 +1,110 @@ +/** Detray library, part of the ACTS project (R&D line) + * + * (c) 2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Project include(s) +#include "detray/core/detector.hpp" +#include "detray/definitions/units.hpp" + +// Detray IO include(s) +#include "detray/io/frontend/detector_reader.hpp" + +// Detray test include(s) +#include "detray/options/detector_io_options.hpp" +#include "detray/options/parse_options.hpp" +#include "detray/options/propagation_options.hpp" +#include "detray/options/track_generator_options.hpp" +#include "detray/test/cpu/material_scan.hpp" +#include "detray/test/device/hip/material_validation.hpp" +#include "detray/test/framework/register_checks.hpp" + +// Vecmem include(s) +#include +#include + +// GTest include(s) +#include + +// Boost +#include "detray/options/boost_program_options.hpp" + +// System include(s) +#include +#include +#include + +namespace po = boost::program_options; + +using namespace detray; + +int main(int argc, char **argv) { + + // Use the most general type to be able to read in all detector files + using metadata_t = test::default_metadata; + using detector_t = detector; + + // Filter out the google test flags + ::testing::InitGoogleTest(&argc, argv); + + // Specific options for this test + po::options_description desc("\ndetray HIP material validation options"); + + desc.add_options()( + "data_dir", + po::value()->default_value("./validation_data/material"), + "Directory that contains the data files")( + "material_tol", po::value()->default_value(1.f), + "Tolerance for comparing the material traces [%]")( + "overlaps_tol", + po::value()->default_value(stepping::config{}.min_stepsize), + "Tolerance for considering surfaces to be overlapping [mm]"); + + // Configs to be filled + detray::io::detector_reader_config reader_cfg{}; + detray::test::material_validation::config mat_val_cfg{}; + test::material_scan::config mat_scan_cfg{}; + + po::variables_map vm = detray::options::parse_options( + desc, argc, argv, reader_cfg, mat_scan_cfg.track_generator(), + mat_val_cfg.propagation()); + + // General options + if (vm.count("material_tol")) { + mat_val_cfg.relative_error(vm["material_tol"].as() / 100.f); + } + if (vm.count("overlaps_tol")) { + mat_scan_cfg.overlaps_tol(vm["overlaps_tol"].as()); + } + const auto data_dir{vm["data_dir"].as()}; + + /// Vecmem memory resource for the device allocations + vecmem::hip::device_memory_resource dev_mr{}; + vecmem::host_memory_resource host_mr; + + const auto [det, names] = + detray::io::read_detector(host_mr, reader_cfg); + + auto white_board = std::make_shared(); + detector_t::geometry_context ctx{}; + + // Print the detector's material as recorded by a ray scan + mat_scan_cfg.name("material_scan_for_hip"); + mat_scan_cfg.track_generator().uniform_eta(true); + mat_scan_cfg.material_file(data_dir + "/material_scan_for_hip"); + detray::test::register_checks(det, names, mat_scan_cfg, + ctx, white_board); + + // Now trace the material during navigation and compare + mat_val_cfg.name("material_validation_hip"); + mat_val_cfg.material_file(data_dir + "/navigation_material_trace"); + mat_val_cfg.device_mr(&dev_mr); + + test::register_checks( + det, names, mat_val_cfg, ctx, white_board); + + // Run the checks + return RUN_ALL_TESTS(); +} diff --git a/tests/tools/src/hip/navigation_validation_hip.cpp b/tests/tools/src/hip/navigation_validation_hip.cpp new file mode 100644 index 000000000..9e9b9f162 --- /dev/null +++ b/tests/tools/src/hip/navigation_validation_hip.cpp @@ -0,0 +1,149 @@ +/** Detray library, part of the ACTS project (R&D line) + * + * (c) 2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Project include(s) +#include "detray/core/detector.hpp" +#include "detray/definitions/units.hpp" + +// Detray IO include(s) +#include "detray/io/frontend/detector_reader.hpp" + +// Detray test include(s) +#include "detray/options/detector_io_options.hpp" +#include "detray/options/parse_options.hpp" +#include "detray/options/propagation_options.hpp" +#include "detray/options/track_generator_options.hpp" +#include "detray/test/cpu/detector_scan.hpp" +#include "detray/test/device/hip/navigation_validation.hpp" +#include "detray/test/framework/register_checks.hpp" +#include "detray/test/framework/whiteboard.hpp" + +// Vecmem include(s) +#include + +// GTest include(s) +#include + +// Boost +#include "detray/options/boost_program_options.hpp" + +// System include(s) +#include +#include +#include + +namespace po = boost::program_options; + +using namespace detray; + +int main(int argc, char** argv) { + + // Use the most general type to be able to read in all detector files + using metadata_t = test::default_metadata; + using detector_t = detector; + using scalar = dscalar; + + // Filter out the google test flags + ::testing::InitGoogleTest(&argc, argv); + + // Specific options for this test + po::options_description desc("\ndetray detector HIP validation options"); + + desc.add_options()("write_scan_data", + "Write the ray/helix scan data to file")( + "data_dir", + boost::program_options::value()->default_value( + "./validation_data"), + "Directory that contains the data files")( + "overlaps_tol", + po::value()->default_value(stepping::config{}.min_stepsize), + "Tolerance for considering surfaces to be overlapping [mm]"); + + // Configs to be filled + detray::io::detector_reader_config reader_cfg{}; + reader_cfg.do_check(true); + detray::test::ray_scan::config ray_scan_cfg{}; + detray::test::helix_scan::config hel_scan_cfg{}; + detray::hip::straight_line_navigation::config str_nav_cfg{}; + detray::hip::helix_navigation::config hel_nav_cfg{}; + + po::variables_map vm = detray::options::parse_options( + desc, argc, argv, reader_cfg, hel_scan_cfg.track_generator(), + hel_nav_cfg.propagation()); + + const auto data_dir{vm["data_dir"].as()}; + + if (vm.count("overlaps_tol")) { + ray_scan_cfg.overlaps_tol(vm["overlaps_tol"].as()); + hel_scan_cfg.overlaps_tol(vm["overlaps_tol"].as()); + } + + // For now: Copy the options to the other tests + ray_scan_cfg.track_generator() = hel_scan_cfg.track_generator(); + str_nav_cfg.propagation() = hel_nav_cfg.propagation(); + str_nav_cfg.fail_on_diff(false); + + detector_t::geometry_context ctx{}; + vecmem::host_memory_resource host_mr; + + const auto [det, names] = + detray::io::read_detector(host_mr, reader_cfg); + const std::string& det_name = det.name(names); + + // Create the whiteboard for data transfer between the steps + auto white_board = std::make_shared(); + const std::string file_prefix{data_dir + "/" + det_name}; + ray_scan_cfg.name(det_name + "_ray_scan_for_hip"); + ray_scan_cfg.intersection_file(file_prefix + "_ray_scan_intersections"); + ray_scan_cfg.track_param_file(file_prefix + "_ray_scan_track_parameters"); + + hel_scan_cfg.name(det_name + "_helix_scan_for_hip"); + // Let the Newton algorithm dynamically choose tol. based on approx. error + hel_scan_cfg.mask_tolerance({detray::detail::invalid_value(), + detray::detail::invalid_value()}); + hel_scan_cfg.intersection_file(file_prefix + "_helix_scan_intersections"); + hel_scan_cfg.track_param_file(file_prefix + "_helix_scan_track_parameters"); + + // Navigation link consistency, discovered by ray intersection + detray::test::register_checks( + det, names, ray_scan_cfg, ctx, white_board); + + // Comparison of straight line navigation with ray scan + str_nav_cfg.name(det_name + "_straight_line_navigation_hip"); + // Number of tracks to check + str_nav_cfg.n_tracks(ray_scan_cfg.track_generator().n_tracks()); + // Ensure that the same mask tolerance is used + auto mask_tolerance = ray_scan_cfg.mask_tolerance(); + str_nav_cfg.propagation().navigation.min_mask_tolerance = + static_cast(mask_tolerance[0]); + str_nav_cfg.propagation().navigation.max_mask_tolerance = + static_cast(mask_tolerance[1]); + str_nav_cfg.intersection_file(ray_scan_cfg.intersection_file()); + str_nav_cfg.track_param_file(ray_scan_cfg.track_param_file()); + + detray::test::register_checks( + det, names, str_nav_cfg, ctx, white_board); + + // Navigation link consistency, discovered by helix intersection + detray::test::register_checks( + det, names, hel_scan_cfg, ctx, white_board); + + // Comparison of navigation in a constant B-field with helix + hel_nav_cfg.name(det_name + "_helix_navigation_hip"); + hel_nav_cfg.fail_on_diff(false); + // Number of tracks to check + hel_nav_cfg.n_tracks(hel_scan_cfg.track_generator().n_tracks()); + hel_nav_cfg.p_range(hel_scan_cfg.track_generator().mom_range()); + hel_nav_cfg.intersection_file(hel_scan_cfg.intersection_file()); + hel_nav_cfg.track_param_file(hel_scan_cfg.track_param_file()); + + detray::test::register_checks( + det, names, hel_nav_cfg, ctx, white_board); + + // Run the checks + return RUN_ALL_TESTS(); +} diff --git a/tests/tools/src/hip/propagation_benchmark_hip.cpp b/tests/tools/src/hip/propagation_benchmark_hip.cpp index 5903f9f35..d5e97362e 100644 --- a/tests/tools/src/hip/propagation_benchmark_hip.cpp +++ b/tests/tools/src/hip/propagation_benchmark_hip.cpp @@ -6,7 +6,7 @@ */ // Project include(s) -#include "detray/navigation/caching_navigator.hpp" +#include "detray/navigation/navigator.hpp" #include "detray/propagator/actors.hpp" #include "detray/propagator/rk_stepper.hpp" #include "detray/tracks/tracks.hpp" @@ -141,10 +141,8 @@ int main(int argc, char** argv) { dtuple<> empty_state{}; pointwise_material_interactor::state interactor_state{}; - parameter_resetter::state resetter_state{prop_cfg}; - auto actor_states = - detail::make_tuple(interactor_state, resetter_state); + auto actor_states = detail::make_tuple(interactor_state); // // Register benchmarks