diff --git a/cmake/detray-compiler-options-hip.cmake b/cmake/detray-compiler-options-hip.cmake index b24877d564..e3ca1a2ddc 100644 --- a/cmake/detray-compiler-options-hip.cmake +++ b/cmake/detray-compiler-options-hip.cmake @@ -45,4 +45,4 @@ if(PROJECT_IS_TOP_LEVEL) detray_add_flag( CMAKE_HIP_FLAGS "-Werror all-warnings" ) endif() endif() -endif() +endif() \ No newline at end of file diff --git a/extern/covfie/CMakeLists.txt b/extern/covfie/CMakeLists.txt index d9e8f8c731..26fb435910 100644 --- a/extern/covfie/CMakeLists.txt +++ b/extern/covfie/CMakeLists.txt @@ -1,6 +1,6 @@ # Detray library, part of the ACTS project (R&D line) # -# (c) 2022-2023 CERN for the benefit of the ACTS project +# (c) 2022-2025 CERN for the benefit of the ACTS project # # Mozilla Public License Version 2.0 @@ -40,6 +40,12 @@ set(COVFIE_BUILD_TESTS OFF CACHE BOOL "Build covfie tests") set(COVFIE_BUILD_BENCHMARKS OFF CACHE BOOL "Build covfie benchmarks") set(COVFIE_PLATFORM_CPU ON CACHE BOOL "Enable covfie CPU platform") + +set(COVFIE_PLATFORM_HIP + ${DETRAY_BUILD_HIP} + CACHE BOOL + "Enable covfie HIP platform" +) set(COVFIE_PLATFORM_CUDA ${DETRAY_BUILD_CUDA} CACHE BOOL diff --git a/tests/include/detray/test/device/CMakeLists.txt b/tests/include/detray/test/device/CMakeLists.txt index b511b67d1d..9804415893 100644 --- a/tests/include/detray/test/device/CMakeLists.txt +++ b/tests/include/detray/test/device/CMakeLists.txt @@ -20,3 +20,5 @@ target_link_libraries( if(DETRAY_BUILD_CUDA) add_subdirectory(cuda) 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 0000000000..51abe4318c --- /dev/null +++ b/tests/include/detray/test/device/hip/CMakeLists.txt @@ -0,0 +1,41 @@ +# 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 + STATIC + "material_validation.hpp" + "material_validation.hip" + "navigation_validation.hpp" + "navigation_validation.hip" +) + +add_library(detray::test_hip ALIAS detray_test_hip) + +target_link_libraries( + detray_test_hip + PUBLIC + HIP::hiprt + vecmem::hip + covfie::hip + detray::core_array + detray::test_device + detray::test_cpu + detray::validation_utils + HIP::hiprt +) 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 0000000000..0ee2f33c7b --- /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) 2024 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 0000000000..5f4b9ea2c3 --- /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 0000000000..803afd5b57 --- /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) 2024 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 + hipLaunchKernel((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 0000000000..b1ce3002e3 --- /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/integration_tests/CMakeLists.txt b/tests/integration_tests/CMakeLists.txt index 5e4edd5cf4..58386a2c4b 100644 --- a/tests/integration_tests/CMakeLists.txt +++ b/tests/integration_tests/CMakeLists.txt @@ -1,6 +1,6 @@ # Detray library, part of the ACTS project (R&D line) # -# (c) 2021-2024 CERN for the benefit of the ACTS project +# (c) 2021-2025 CERN for the benefit of the ACTS project # # Mozilla Public License Version 2.0 @@ -13,6 +13,6 @@ if(DETRAY_BUILD_HOST) endif() # Set up all of the "device" tests. -if(DETRAY_BUILD_CUDA OR DETRAY_BUILD_SYCL) +if(DETRAY_BUILD_CUDA OR DETRAY_BUILD_SYCL OR DETRAY_BUILD_HIP) add_subdirectory(device) endif() diff --git a/tests/integration_tests/device/CMakeLists.txt b/tests/integration_tests/device/CMakeLists.txt index e5c5f99576..97cafe98a5 100644 --- a/tests/integration_tests/device/CMakeLists.txt +++ b/tests/integration_tests/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 @@ -11,3 +11,7 @@ endif() if(DETRAY_BUILD_SYCL) add_subdirectory(sycl) endif() + +if(DETRAY_BUILD_HIP) + add_subdirectory(hip) +endif() diff --git a/tests/integration_tests/device/hip/CMakeLists.txt b/tests/integration_tests/device/hip/CMakeLists.txt new file mode 100644 index 0000000000..81aaf7e37a --- /dev/null +++ b/tests/integration_tests/device/hip/CMakeLists.txt @@ -0,0 +1,63 @@ +# Detray library, part of the ACTS project (R&D line) +# +# (c) 2021-2025 CERN for the benefit of the ACTS project +# +# Mozilla Public License Version 2.0 + +message(STATUS "Building detray HIP integration tests") + +#HIP langauge support requires minimum 3.21 +cmake_minimum_required(VERSION 3.21) + +find_package(HIPToolkit) + +# Enable HIP as a language. +enable_language(HIP) + +# Set the HIP build flags. +include(detray-compiler-options-hip) + +# make unit tests for multiple algebras +# Currently vc and smatrix is not supported +set(algebras "array") + +if(DETRAY_EIGEN_PLUGIN) + list(APPEND algebras "eigen") +endif() + + + + +foreach(algebra ${algebras}) + # Unit tests for the selected algebra. + detray_add_integration_test(hip_${CMAKE_HIP_PLATFORM}_${algebra} + "propagator_hip_kernel.hpp" + "propagator_hip.cpp" + "propagator_hip_kernel.hip" + LINK_LIBRARIES GTest::gtest_main + HIP::hiprt vecmem::hip #covfie::hip + detray::core detray::algebra_${algebra} detray::test_device + detray::test_common + ) + # Make hipcc interpret .cpp files as .hip files to make linking work + # (only needed for the amd backend) + if( + ("${CMAKE_HIP_PLATFORM}" STREQUAL "hcc") + OR ("${CMAKE_HIP_PLATFORM}" STREQUAL "amd") + ) + set_source_files_properties(propagator_hip.cpp PROPERTIES LANGUAGE HIP) + endif() + + target_compile_definitions( + detray_integration_test_hip_${CMAKE_HIP_PLATFORM}_${algebra} + PRIVATE ${algebra}=${algebra} + ) + + set_tests_properties( + detray_integration_test_hip_${CMAKE_HIP_PLATFORM}_${algebra} + PROPERTIES + DEPENDS "detray_unit_test_hip_${CMAKE_HIP_PLATFORM}_${algebra}" + ) +endforeach() + + diff --git a/tests/integration_tests/device/hip/propagator_hip.cpp b/tests/integration_tests/device/hip/propagator_hip.cpp new file mode 100644 index 0000000000..230dda28bb --- /dev/null +++ b/tests/integration_tests/device/hip/propagator_hip.cpp @@ -0,0 +1,184 @@ +/** Detray library, part of the ACTS project (R&D line) + * + * (c) 2022-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Detray test include(s) +#include "detray/test/common/bfield.hpp" +#include "detray/test/common/build_toy_detector.hpp" +#include "propagator_hip_kernel.hpp" + +// Vecmem include(s) +#include +#include +#include +#include + +// GTest include +#include + +using namespace detray; + +class HipPropConstBFieldMng + : public ::testing::TestWithParam> {}; + +/// Propagation test using unified memory +TEST_P(HipPropConstBFieldMng, propagator) { + + // VecMem memory resource(s) + vecmem::hip::managed_memory_resource mng_mr; + + // Test configuration + propagator_test_config cfg{}; + cfg.track_generator.phi_steps(20).theta_steps(20); + cfg.track_generator.p_tot(10.f * unit::GeV); + cfg.track_generator.eta_range(-3.f, 3.f); + cfg.propagation.navigation.search_window = {3u, 3u}; + // Configuration for non-z-aligned B-fields + cfg.propagation.navigation.overstep_tolerance = std::get<0>(GetParam()); + cfg.propagation.stepping.step_constraint = std::get<1>(GetParam()); + + // Get the magnetic field + const vector3 B = std::get<2>(GetParam()); + auto field = create_const_field(B); + + // Create the toy geometry + auto [det, names] = build_toy_detector(mng_mr); + + run_propagation_test>( + &mng_mr, det, cfg, detray::get_data(det), std::move(field)); +} + +class HipPropConstBFieldCpy + : public ::testing::TestWithParam> {}; + +/// Propagation test using vecmem copy +TEST_P(HipPropConstBFieldCpy, propagator) { + + // VecMem memory resource(s) + vecmem::host_memory_resource host_mr; + vecmem::hip::managed_memory_resource mng_mr; + vecmem::hip::device_memory_resource dev_mr; + + vecmem::hip::copy hip_cpy; + + // Test configuration + propagator_test_config cfg{}; + cfg.track_generator.phi_steps(20u).theta_steps(20u); + cfg.track_generator.p_tot(10.f * unit::GeV); + cfg.track_generator.eta_range(-3.f, 3.f); + cfg.propagation.navigation.search_window = {3u, 3u}; + // Configuration for non-z-aligned B-fields + cfg.propagation.navigation.overstep_tolerance = std::get<0>(GetParam()); + cfg.propagation.stepping.step_constraint = std::get<1>(GetParam()); + + // Get the magnetic field + const vector3 B = std::get<2>(GetParam()); + auto field = create_const_field(B); + + // Create the toy geometry + auto [det, names] = build_toy_detector(host_mr); + + auto det_buff = detray::get_buffer(det, dev_mr, hip_cpy); + + run_propagation_test>( + &mng_mr, det, cfg, detray::get_data(det_buff), std::move(field)); +} + +INSTANTIATE_TEST_SUITE_P( + HipPropagatorValidation1, HipPropConstBFieldMng, + ::testing::Values(std::make_tuple(-100.f * unit::um, + std::numeric_limits::max(), + vector3{0.f * unit::T, + 0.f * unit::T, + 2.f * unit::T}))); + +INSTANTIATE_TEST_SUITE_P( + HipPropagatorValidation2, HipPropConstBFieldMng, + ::testing::Values(std::make_tuple(-400.f * unit::um, + std::numeric_limits::max(), + vector3{0.f * unit::T, + 1.f * unit::T, + 1.f * unit::T}))); + +INSTANTIATE_TEST_SUITE_P( + HipPropagatorValidation3, HipPropConstBFieldMng, + ::testing::Values(std::make_tuple(-400.f * unit::um, + std::numeric_limits::max(), + vector3{1.f * unit::T, + 0.f * unit::T, + 1.f * unit::T}))); + +INSTANTIATE_TEST_SUITE_P( + HIpPropagatorValidation4, HipPropConstBFieldMng, + ::testing::Values(std::make_tuple(-600.f * unit::um, + std::numeric_limits::max(), + vector3{1.f * unit::T, + 1.f * unit::T, + 1.f * unit::T}))); + +INSTANTIATE_TEST_SUITE_P( + HipPropagatorValidation5, HipPropConstBFieldCpy, + ::testing::Values(std::make_tuple(-100.f * unit::um, + std::numeric_limits::max(), + vector3{0.f * unit::T, + 0.f * unit::T, + 2.f * unit::T}))); + +INSTANTIATE_TEST_SUITE_P( + HIpPropagatorValidation6, HipPropConstBFieldCpy, + ::testing::Values(std::make_tuple(-400.f * unit::um, + std::numeric_limits::max(), + vector3{0.f * unit::T, + 1.f * unit::T, + 1.f * unit::T}))); + +INSTANTIATE_TEST_SUITE_P( + HipPropagatorValidation7, HipPropConstBFieldCpy, + ::testing::Values(std::make_tuple(-400.f * unit::um, + std::numeric_limits::max(), + vector3{1.f * unit::T, + 0.f * unit::T, + 1.f * unit::T}))); + +INSTANTIATE_TEST_SUITE_P( + HipPropagatorValidation8, HipPropConstBFieldCpy, + ::testing::Values(std::make_tuple(-600.f * unit::um, + std::numeric_limits::max(), + vector3{1.f * unit::T, + 1.f * unit::T, + 1.f * unit::T}))); + +/// This tests the device propagation in an inhomogenepus magnetic field +/* +TEST(HipPropagatorValidation9, inhomogeneous_bfield_cpy) { + + // VecMem memory resource(s) + vecmem::host_memory_resource host_mr; + vecmem::hip::managed_memory_resource mng_mr; + vecmem::hip::device_memory_resource dev_mr; + + vecmem::hip::copy hip_cpy; + + // Test configuration + propagator_test_config cfg{}; + cfg.track_generator.phi_steps(10u).theta_steps(10u); + cfg.track_generator.p_tot(10.f * unit::GeV); + cfg.track_generator.eta_range(-3.f, 3.f); + cfg.propagation.navigation.search_window = {3u, 3u}; + + // Get the magnetic field + auto field = create_inhom_field(); + + // Create the toy geometry with inhomogeneous bfield from file + auto [det, names] = build_toy_detector(host_mr); + + auto det_buff = detray::get_buffer(det, dev_mr, hip_cpy); + + //run_propagation_test( + // &mng_mr, det, cfg, detray::get_data(det_buff), std::move(field)); + +} +*/ \ No newline at end of file diff --git a/tests/integration_tests/device/hip/propagator_hip_kernel.hip b/tests/integration_tests/device/hip/propagator_hip_kernel.hip new file mode 100644 index 0000000000..12f7f1db3b --- /dev/null +++ b/tests/integration_tests/device/hip/propagator_hip_kernel.hip @@ -0,0 +1,111 @@ +/** Detray library, part of the ACTS project (R&D line) + * + * (c) 2022-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Project include(s) +#include "detray/definitions/detail/hip_definitions.hpp" + +// Detray test include(s) +#include "propagator_hip_kernel.hpp" + +namespace detray { + +template +__global__ void propagator_test_kernel( + typename detector_t::view_type det_data, const propagation::config cfg, + covfie::field_view field_data, + vecmem::data::vector_view tracks_data, + vecmem::data::jagged_vector_view> + steps_data) { + + int gid = threadIdx.x + blockIdx.x * blockDim.x; + using detector_device_t = + detector; + + static_assert(std::is_same_v, + "Host and device detector views do not match"); + + detector_device_t det(det_data); + vecmem::device_vector tracks(tracks_data); + vecmem::jagged_device_vector> steps( + steps_data); + + if (gid >= tracks.size()) { + return; + } + + auto stepr = rk_stepper_t>{}; + auto nav = navigator_t{}; + + // Create propagator + using propagator_device_t = + propagator; + + propagator_device_t p{cfg}; + + // Create actor states + step_tracer_device_t::state tracer_state(steps.at(gid)); + tracer_state.collect_only_on_surface(true); + pathlimit_aborter_t::state aborter_state{cfg.stepping.path_limit}; + pointwise_material_interactor::state interactor_state{}; + + // Create the actor states + auto actor_states = + ::detray::tie(tracer_state, aborter_state, interactor_state); + // Create the propagator state + typename propagator_device_t::state state(tracks[gid], field_data, det); + + state._stepping.template set_constraint( + cfg.stepping.step_constraint); + + // Run propagation + p.propagate(state, actor_states); +} + +/// Launch the device kernel +template +void propagator_test( + typename detector_t::view_type det_view, const propagation::config& cfg, + covfie::field_view field_data, + vecmem::data::vector_view& tracks_data, + vecmem::data::jagged_vector_view>& + step_data) { + + constexpr int thread_dim = 2 * WARP_SIZE; + int block_dim = tracks_data.size() / thread_dim + 1; + + // run the test kernel + hipLaunchKernelGGL((propagator_test_kernel), dim3(block_dim) , dim3(thread_dim), 0,0, + det_view, cfg, field_data, + tracks_data,step_data); + + // hip error check + DETRAY_HIP_ERROR_CHECK(hipGetLastError()); + DETRAY_HIP_ERROR_CHECK(hipDeviceSynchronize()); +} + +/// Explicit instantiation for a constant magnetic field +template void +propagator_test>, + detector, host_container_types>>( + detector, host_container_types>::view_type, + const propagation::config&, + covfie::field_view>>, + vecmem::data::vector_view&, + vecmem::data::jagged_vector_view>&); + +/// Explicit instantiation for an inhomogeneous magnetic field +/* +template void +propagator_test, host_container_types>>( + detector, host_container_types>::view_type, + const propagation::config&, covfie::field_view, + vecmem::data::vector_view&, + vecmem::data::jagged_vector_view>&); */ + +} // namespace detray diff --git a/tests/integration_tests/device/hip/propagator_hip_kernel.hpp b/tests/integration_tests/device/hip/propagator_hip_kernel.hpp new file mode 100644 index 0000000000..89906e3bb5 --- /dev/null +++ b/tests/integration_tests/device/hip/propagator_hip_kernel.hpp @@ -0,0 +1,123 @@ +/** Detray library, part of the ACTS project (R&D line) + * + * (c) 2023-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s) +#include "detray/detectors/toy_metadata.hpp" + +// Detray test include(s) +#include "detray/test/common/bfield.hpp" +#include "detray/test/device/propagator_test.hpp" + +// Vecmem include(s) +#include +#include + +// Covfie include(s) + +//we can include it after solving the linking problem - covfie::hip in the cake file +//#include // error + +namespace detray { + +using scalar = test::scalar; + + + +// this namespace - dependance with covfie -> cannot run it until solving the covfie::hip + +/*namespace bfield::hip { + +// Inhomogeneous field (hip) +using inhom_bknd_t = covfie::backend::affine, + covfie::backend::hip_device_array< + covfie::vector::vector_d>>>> + +} // namespace bfield::hip + +*/ + + + +/// Launch the propagation test kernel +template +void propagator_test( + typename detector_t::view_type, const propagation::config &, + covfie::field_view, vecmem::data::vector_view &, + vecmem::data::jagged_vector_view> &); + +/// Test function for propagator on the device +template +inline auto run_propagation_device( + vecmem::memory_resource *mr, const propagation::config &cfg, + typename detector_t::view_type det_view, + covfie::field_view field_data, dvector &tracks, + const vecmem::jagged_vector> &host_steps) + -> vecmem::jagged_vector> { + + // Helper object for performing memory copies. + vecmem::copy copy; + + // Get tracks data + auto tracks_data = vecmem::get_data(tracks); + + // Create vector buffer for track recording + std::vector sizes(tracks.size(), 0); + std::vector capacities; + for (auto &st : host_steps) { + // Add a few more elements for security (in case the device finds more + // surfaces) + capacities.push_back(st.size() + 10u); + } + + vecmem::data::jagged_vector_buffer> + steps_buffer(capacities, *mr, nullptr, + vecmem::data::buffer_type::resizable); + + copy.setup(steps_buffer)->wait(); + + // Run the propagator test for GPU device + propagator_test(det_view, cfg, field_data, + tracks_data, steps_buffer); + + vecmem::jagged_vector> steps(mr); + + copy(steps_buffer, steps)->wait(); + + return steps; +} + +/// Test chain for the propagator +template +inline auto run_propagation_test(vecmem::memory_resource *mr, detector_t &det, + const propagator_test_config &cfg, + typename detector_t::view_type det_view, + covfie::field &&field) { + + // Create the vector of initial track parameterizations + auto tracks_host = generate_tracks(mr, cfg.track_generator); + vecmem::vector tracks_device(tracks_host, mr); + + // Host propagation + auto host_steps = + run_propagation_host(mr, det, cfg.propagation, field, tracks_host); + + // Device propagation (device backend specific implementation) + covfie::field device_field(field); + auto device_steps = + run_propagation_device( + mr, cfg.propagation, det_view, device_field, tracks_device, + host_steps); + + // Check the results + compare_propagation_results(host_steps, device_steps); +} + +} // namespace detray diff --git a/tests/unit_tests/device/hip/CMakeLists.txt b/tests/unit_tests/device/hip/CMakeLists.txt index 22f974cfe4..c3a82bedca 100644 --- a/tests/unit_tests/device/hip/CMakeLists.txt +++ b/tests/unit_tests/device/hip/CMakeLists.txt @@ -5,6 +5,7 @@ # Mozilla Public License Version 2.0 message(STATUS "Building detray HIP unit tests") + cmake_minimum_required(VERSION 3.21) # HIP langauge support requires minimum 3.21 find_package(HIPToolkit) @@ -22,18 +23,28 @@ if(DETRAY_EIGEN_PLUGIN) list(APPEND algebras "eigen") endif() foreach(algebra ${algebras}) - detray_add_unit_test(hip_${algebra} + detray_add_unit_test(hip_${CMAKE_HIP_PLATFORM}_${algebra} "detector_hip_kernel.hpp" "detector_hip.cpp" "detector_hip_kernel.hip" - LINK_LIBRARIES GTest::gtest_main vecmem::hip detray::core - detray::algebra_${algebra} detray::test_common detray::test_utils HIP::hiprt + LINK_LIBRARIES GTest::gtest_main HIP::hiprt vecmem::hip detray::core + detray::algebra_${algebra} detray::test_common detray::test_utils + ) + + # Make hipcc interpret .cpp files as .hip files to make linking work + # (only needed for the amd backend) + if( + ("${CMAKE_HIP_PLATFORM}" STREQUAL "hcc") + OR ("${CMAKE_HIP_PLATFORM}" STREQUAL "amd") ) + set_source_files_properties(detector_hip.cpp PROPERTIES LANGUAGE HIP) + endif() + # Add definitions for HIP + algebra target_compile_definitions( - detray_unit_test_hip_${algebra} + detray_unit_test_hip_${CMAKE_HIP_PLATFORM}_${algebra} PRIVATE ${algebra}=${algebra} ) endforeach() diff --git a/tests/unit_tests/device/hip/detector_hip.cpp b/tests/unit_tests/device/hip/detector_hip.cpp index d0aabf2d58..375b2a56b5 100644 --- a/tests/unit_tests/device/hip/detector_hip.cpp +++ b/tests/unit_tests/device/hip/detector_hip.cpp @@ -68,7 +68,6 @@ TEST(detector_hip, detector) { // run the test code to copy the objects detector_test(toy_det_data, volumes_data, surfaces_data, transforms_data, rectangles_data, discs_data, cylinders_data); - // check if the same volume objects are copied for (unsigned int i = 0u; i < volumes_host.size(); i++) { EXPECT_EQ(volumes_host[i] == volumes_device[i], true);