diff --git a/core/include/detray/core/detail/alignment.hpp b/core/include/detray/core/detail/alignment.hpp new file mode 100644 index 000000000..4f6632459 --- /dev/null +++ b/core/include/detray/core/detail/alignment.hpp @@ -0,0 +1,34 @@ +/** Detray library, part of the ACTS project (R&D line) + * + * (c) 2023-2024 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +namespace detray::detail { + +/// Creates detector view using "static" detector components and +/// a "misaligned" transform store +template +typename host_detector_type::view_type misaligned_detector_view( + typename host_detector_type::buffer_type& det_buffer, + typename host_detector_type::transform_container::buffer_type& trf_buffer) { + typename host_detector_type::view_type detview{ + detray::get_data( + detray::detail::get<0>(det_buffer.m_buffer)), // volumes + detray::get_data( + detray::detail::get<1>(det_buffer.m_buffer)), // surfaces + detray::get_data(trf_buffer), // transforms + detray::get_data(detray::detail::get<3>(det_buffer.m_buffer)), // masks + detray::get_data( + detray::detail::get<4>(det_buffer.m_buffer)), // materials + detray::get_data( + detray::detail::get<5>(det_buffer.m_buffer)), // accelerators + detray::get_data(detray::detail::get<6>( + det_buffer.m_buffer))}; // volume search grid + return detview; +} + +} // namespace detray::detail diff --git a/tests/unit_tests/device/cuda/detector_cuda.cpp b/tests/unit_tests/device/cuda/detector_cuda.cpp index a7d54f2b8..31bd51af8 100644 --- a/tests/unit_tests/device/cuda/detector_cuda.cpp +++ b/tests/unit_tests/device/cuda/detector_cuda.cpp @@ -7,6 +7,9 @@ // Detray test include(s) #include "detector_cuda_kernel.hpp" +#include "detray/core/detail/alignment.hpp" +#include "detray/definitions/detail/algebra.hpp" +#include "detray/plugins/algebra/array_definitions.hpp" #include "detray/test/utils/detectors/build_toy_detector.hpp" // Vecmem include(s) @@ -96,3 +99,46 @@ TEST(detector_cuda, detector) { EXPECT_EQ(cylinders_host[i] == cylinders_device[i], true); } } + +TEST(detector_align_cuda, detector) { + // Memory resources + vecmem::host_memory_resource host_mr; + vecmem::cuda::managed_memory_resource mng_mr; + vecmem::cuda::device_memory_resource dev_mr; + + // Helper object for performing memory copies to CUDA devices + vecmem::cuda::copy cuda_cpy; + + // Create toy geometry in host memory + auto [det_host, names_host] = build_toy_detector(host_mr); + + // Copy static detector data (including the initial set of transforms) to + // the device Use synchronous copy and fixed size buffers + auto det_fixed_buff = detray::get_buffer(det_host, dev_mr, cuda_cpy); + + // Construct an "aligned" transform store + using algebra_t = array; + using point3 = dpoint3D; + + typename decltype(det_host)::transform_container tf_store_aligned_host; + using tf_type = decltype(det_host)::transform_container::value_type; + + point3 shift{.1f * unit::mm, .2f * unit::mm, + .3f * unit::mm}; + for (const auto& tf : det_host.transform_store()) { + point3 shifted{tf.translation()[0] + shift[0], + tf.translation()[1] + shift[1], + tf.translation()[2] + shift[2]}; + tf_store_aligned_host.push_back( + tf_type{shifted, tf.x(), tf.y(), tf.z()}); + } + + auto tf_buff_shifted = + get_buffer(tf_store_aligned_host, dev_mr, cuda_cpy, copy::sync, + vecmem::data::buffer_type::fixed_size); + + using host_detector_type = decltype(det_host); + [[maybe_unused]] auto detector_view = + detail::misaligned_detector_view(det_fixed_buff, + tf_buff_shifted); +} diff --git a/tutorials/src/device/cuda/detector_construction.cpp b/tutorials/src/device/cuda/detector_construction.cpp index d503bd61e..8b6443abe 100644 --- a/tutorials/src/device/cuda/detector_construction.cpp +++ b/tutorials/src/device/cuda/detector_construction.cpp @@ -8,6 +8,7 @@ // Project include(s) #include "detector_construction.hpp" +#include "detray/core/detail/alignment.hpp" #include "detray/test/utils/detectors/build_toy_detector.hpp" // Vecmem include(s) @@ -77,10 +78,9 @@ int main() { auto sf_buff = detray::get_buffer(det_host.surfaces(), dev_mr, cuda_cpy, detray::copy::sync, vecmem::data::buffer_type::fixed_size); - // Use resizable buffer and asynchronous copy for alignment auto trf_buff = detray::get_buffer(det_host.transform_store(), dev_mr, - cuda_cpy, detray::copy::async, - vecmem::data::buffer_type::resizable); + cuda_cpy, detray::copy::sync, + vecmem::data::buffer_type::fixed_size); auto msk_buff = detray::get_buffer(det_host.mask_store(), dev_mr, cuda_cpy, detray::copy::sync, vecmem::data::buffer_type::fixed_size); @@ -95,11 +95,40 @@ int main() { vecmem::data::buffer_type::fixed_size); // Assemble the detector buffer - auto det_custom_buff = typename decltype(det_host)::buffer_type( + using host_detector_type = decltype(det_host); + auto det_custom_buff = typename host_detector_type::buffer_type( std::move(vol_buff), std::move(sf_buff), std::move(trf_buff), std::move(msk_buff), std::move(mat_buff), std::move(acc_buff), std::move(vgrid_buff)); std::cout << "\nCustom buffer setup:" << std::endl; detray::tutorial::print(detray::get_data(det_custom_buff)); + + // Construct an "aligned" transform store + using host_transform_type = + host_detector_type::transform_container::value_type; + + typename host_detector_type::transform_container host_aligned_transforms; + detray::tutorial::point3 shift{.1f * detray::unit::mm, + .2f * detray::unit::mm, + .3f * detray::unit::mm}; + + for (const auto& tf : det_host.transform_store()) { + detray::tutorial::point3 shifted{tf.translation()[0] + shift[0], + tf.translation()[1] + shift[1], + tf.translation()[2] + shift[2]}; + host_aligned_transforms.push_back( + host_transform_type{shifted, tf.x(), tf.y(), tf.z()}); + } + + auto trf_buff_shifted = detray::get_buffer( + host_aligned_transforms, dev_mr, cuda_cpy, detray::copy::sync, + vecmem::data::buffer_type::fixed_size); + + auto detector_view = + detray::detail::misaligned_detector_view( + det_custom_buff, trf_buff_shifted); + + std::cout << "\nCustom buffer setup (shifted):" << std::endl; + detray::tutorial::print(detector_view); } diff --git a/tutorials/src/device/cuda/detector_construction.cu b/tutorials/src/device/cuda/detector_construction.cu index 981d62f8e..1e41a7acd 100644 --- a/tutorials/src/device/cuda/detector_construction.cu +++ b/tutorials/src/device/cuda/detector_construction.cu @@ -25,6 +25,10 @@ __global__ void print_kernel( printf("Number of volumes: %d\n", det.volumes().size()); printf("Number of transforms: %d\n", det.transform_store().size()); + printf("First translation: {%f,%f,%f}\n", + det.transform_store().at(0).translation()[0], + det.transform_store().at(0).translation()[1], + det.transform_store().at(0).translation()[2]); printf("Number of rectangles: %d\n", det.mask_store().get().size()); printf("Number of trapezoids: %d\n",