Skip to content

Hip validation #1046

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
Larb-YasmineBadrElhouda wants to merge 14 commits into
base: main
Choose a base branch
from
Open

Hip validation #1046

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
14 commits
Select commit Hold shift + click to select a range
e513d6f
hip benchmarks integration : cmake config + implementation
Larb-YasmineBadrElhouda Aug 8, 2025
75af892
hip benchmarks integration : cmake config + implementation
Larb-YasmineBadrElhouda Aug 13, 2025
fb8dcb7
hip benchmarks integration : cmake config + implementation
Larb-YasmineBadrElhouda Aug 13, 2025
19b75eb
hip benchmarks integration : cmake config + implementation
Larb-YasmineBadrElhouda Aug 13, 2025
077793f
hip benchmarks integration : cmake config + implementation
Larb-YasmineBadrElhouda Aug 13, 2025
0d1a9b1
hip benchmarks integration : cmake config + implementation
Larb-YasmineBadrElhouda Aug 14, 2025
7b3b342
hip benchmarks integration : cmake config + implementation
Larb-YasmineBadrElhouda Aug 14, 2025
7051196
Revert "hip benchmarks integration : cmake config + implementation"
Larb-YasmineBadrElhouda Aug 14, 2025
8273050
hip benchmarks integration : cmake config + implementation
Larb-YasmineBadrElhouda Aug 14, 2025
3782472
material & navigation & detector validation - HIP Integration
Larb-YasmineBadrElhouda Aug 14, 2025
8bb3c06
material & navigation & detector validation - HIP Integration
Larb-YasmineBadrElhouda Aug 14, 2025
4043114
material & navigation & detector validation - HIP Integration
Larb-YasmineBadrElhouda Aug 17, 2025
daeb578
pre-commit
Larb-YasmineBadrElhouda Aug 22, 2025
d5e06e6
Merge branch 'main' into hip-validation
Larb-YasmineBadrElhouda Jan 29, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
47 changes: 24 additions & 23 deletions tests/benchmarks/include/detray/benchmarks/device/hip/propagation_benchmark.hip
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<typename propagator_t::stepper_type,
caching_navigator<detector_device_t>, actor_chain_t>;
navigator<detector_device_t>, actor_chain_t>;

const detector_device_t det(det_view);
const vecmem::device_vector<free_track_parameters<algebra_t>> tracks(
tracks_view);

int gid = threadIdx.x + blockIdx.x * blockDim.x;
if (gid >= static_cast<int>(tracks.size())) {
if (gid >= tracks.size()) {
return;
}

Expand All @@ -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);
}
}

Expand All @@ -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;
Expand Down Expand Up @@ -107,38 +108,38 @@ void run_propagation_kernel(
int block_dim = (n_samples + thread_dim - 1) / thread_dim;

// run the test kernel
hipLaunchKernelGGL((propagator_benchmark_kernel<propagator_t, kOPT>),
dim3(block_dim), dim3(thread_dim), 0, 0, cfg, det_view,
field_view, device_actor_state_ptr, tracks_view);
hipLaunchKernelGGL((propagator_benchmark_kernel<propagator_t, kOPT>), 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());
DETRAY_HIP_ERROR_CHECK(hipDeviceSynchronize());
}

/// 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<hip_propagator_type<METADATA, FIELD, CHAIN>, OPT>( \
const propagation::config &, detector<METADATA>::view_type, \
covfie::field_view<FIELD>, \
const propagation::config &, detector<METADATA>::view_type, \
covfie::field_view<FIELD>, \
hip_propagator_type<METADATA, FIELD, \
CHAIN>::actor_chain_type::state_tuple *, \
vecmem::data::vector_view< \
free_track_parameters<detector<METADATA>::algebra_type>>, \
const int); \
\
CHAIN>::actor_chain_type::state_tuple *, \
vecmem::data::vector_view< \
free_track_parameters<detector<METADATA>::algebra_type>>, \
const int); \
\
template hip_propagator_type<METADATA, FIELD, \
CHAIN>::actor_chain_type::state_tuple * \
CHAIN>::actor_chain_type::state_tuple * \
setup_actor_states<hip_propagator_type<METADATA, FIELD, CHAIN>>( \
hip_propagator_type<METADATA, FIELD, \
CHAIN>::actor_chain_type::state_tuple *); \
\
template void \
CHAIN>::actor_chain_type::state_tuple *); \
\
template void \
release_actor_states<hip_propagator_type<METADATA, FIELD, CHAIN>>( \
hip_propagator_type<METADATA, FIELD, \
CHAIN>::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)
Expand Down
11 changes: 5 additions & 6 deletions tests/benchmarks/include/detray/benchmarks/device/hip/propagation_benchmark.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand Down Expand Up @@ -59,7 +58,7 @@ template <typename metadata_t, typename bfield_t,
using hip_propagator_type =
propagator<rk_stepper<covfie::field_view<bfield_t>,
typename detector<metadata_t>::algebra_type>,
caching_navigator<detector<metadata_t>>,
navigator<detector<metadata_t>>,
actor_chain_t<typename detector<metadata_t>::algebra_type>>;

/// Launch the propagation kernelfor benchmarking
Expand Down Expand Up @@ -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
Expand Down
6 changes: 5 additions & 1 deletion tests/include/detray/test/device/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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

Expand All @@ -22,3 +22,7 @@ target_link_libraries(
if(DETRAY_BUILD_CUDA)
add_subdirectory(cuda)
endif()

if(DETRAY_BUILD_HIP)
add_subdirectory(hip)
endif()
43 changes: 43 additions & 0 deletions tests/include/detray/test/device/hip/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
)
125 changes: 125 additions & 0 deletions tests/include/detray/test/device/hip/material_validation.hip
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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 <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
Loading
Loading