Hip integration test

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()

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

tests/include/detray/test/device/CMakeLists.txt

@@ -20,3 +20,5 @@ target_link_libraries(
 if(DETRAY_BUILD_CUDA)
     add_subdirectory(cuda)
 endif()
+
+

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
+)

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 <typename detector_t>
+__global__ void material_validation_kernel(
+    typename detector_t::view_type det_data, const propagation::config cfg,
+    vecmem::data::vector_view<
+        free_track_parameters<typename detector_t::algebra_type>>
+        tracks_view,
+    vecmem::data::vector_view<
+        material_validator::material_record<typename detector_t::scalar_type>>
+        mat_records_view,
+    vecmem::data::jagged_vector_view<
+        material_validator::material_params<typename detector_t::scalar_type>>
+        mat_steps_view) {
+
+    using detector_device_t =
+        detector<typename detector_t::metadata, device_container_types>;
+    using algebra_t = typename detector_device_t::algebra_type;
+    using scalar_t = dscalar<algebra_t>;
+
+    using stepper_t = line_stepper<algebra_t>;
+    using navigator_t = navigator<detector_device_t>;
+    // Propagator with full covariance transport, pathlimit aborter and
+    // material tracer
+    using material_tracer_t =
+        material_validator::material_tracer<scalar_t, vecmem::device_vector>;
+    using pathlimit_aborter_t = pathlimit_aborter<scalar_t>;
+    using actor_chain_t =
+        actor_chain<pathlimit_aborter_t, parameter_transporter<algebra_t>,
+                    parameter_resetter<algebra_t>,
+                    pointwise_material_interactor<algebra_t>,
+                    material_tracer_t>;
+    using propagator_t = propagator<stepper_t, navigator_t, actor_chain_t>;
+
+    detector_device_t det(det_data);
+
+    vecmem::device_vector<free_track_parameters<algebra_t>> tracks(tracks_view);
+    vecmem::device_vector<typename material_tracer_t::material_record_type>
+        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<algebra_t>::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 <typename detector_t>
+void material_validation_device(
+    typename detector_t::view_type det_view, const propagation::config &cfg,
+    vecmem::data::vector_view<
+        free_track_parameters<typename detector_t::algebra_type>> &tracks_view,
+    vecmem::data::vector_view<
+        material_validator::material_record<typename detector_t::scalar_type>>
+        &mat_records_view,
+    vecmem::data::jagged_vector_view<
+        material_validator::material_params<typename detector_t::scalar_type>>
+        &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<detector_t> , 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<detector<METADATA>>(             \
+        typename detector<METADATA>::view_type, const propagation::config &,  \
+        vecmem::data::vector_view<                                            \
+            free_track_parameters<typename detector<METADATA>::algebra_type>> \
+            &,                                                                \
+        vecmem::data::vector_view<material_validator::material_record<        \
+            typename detector<METADATA>::scalar_type>> &,                     \
+        vecmem::data::jagged_vector_view<material_validator::material_params< \
+            typename detector<METADATA>::scalar_type>> &);
+
+DECLARE_MATERIAL_VALIDATION(test::default_metadata)
+DECLARE_MATERIAL_VALIDATION(test::toy_metadata)
+DECLARE_MATERIAL_VALIDATION(test::default_telescope_metadata)
+
+}  // namespace detray::hip

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 <vecmem/memory/hip/device_memory_resource.hpp>
+#include <vecmem/memory/host_memory_resource.hpp>
+#include <vecmem/memory/memory_resource.hpp>
+#include <vecmem/utils/hip/copy.hpp>
+
+// System include
+#include <string_view>
+
+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 <typename detector_t>
+void material_validation_device(
+    typename detector_t::view_type det_view, const propagation::config &cfg,
+    vecmem::data::vector_view<
+        free_track_parameters<typename detector_t::algebra_type>> &tracks_view,
+    vecmem::data::vector_view<
+        material_validator::material_record<typename detector_t::scalar_type>>
+        &mat_records_view,
+    vecmem::data::jagged_vector_view<
+        material_validator::material_params<typename detector_t::scalar_type>>
+        &mat_steps_view);
+
+/// Prepare data for device material trace run
+struct run_material_validation {
+
+    static constexpr std::string_view name = "hip";
+
+    template <typename detector_t>
+    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<typename detector_t::algebra_type>> &tracks,
+        const std::vector<std::size_t> &capacities) {
+
+        using algebra_t = typename detector_t::algebra_type;
+        using scalar_t = dscalar<algebra_t>;
+        using track_t = free_track_parameters<algebra_t>;
+        using material_record_t = material_validator::material_record<scalar_t>;
+        using material_params_t = material_validator::material_params<scalar_t>;
+
+        // 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<track_t> tracks_view =
+            vecmem::get_data(tracks_buffer);
+
+        vecmem::data::vector_buffer<material_record_t> mat_records_buffer(
+            static_cast<unsigned int>(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<material_params_t> 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<detector_t>(
+            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<material_record_t> mat_records(host_mr);
+        hip_cpy(mat_records_buffer, mat_records)->wait();
+
+        vecmem::jagged_vector<material_params_t> mat_steps(host_mr);
+        hip_cpy(mat_steps_buffer, mat_steps)->wait();
+
+        return std::make_tuple(mat_records, mat_steps);
+    }
+};
+
+template <typename detector_t>
+using material_validation = detray::test::material_validation_impl<
+    detector_t, detray::hip::run_material_validation>;
+
+}  // namespace detray::hip