From 9ee3668b30792ff34386288f0942851d0920732c Mon Sep 17 00:00:00 2001 From: Vakho Tsulaia Date: Sat, 30 Nov 2024 21:57:21 -0800 Subject: [PATCH] New function for creating misaligned detector views Introduced a new function `misaligned_detector_view()` that makes a misaligned detector view by combining "static" buffers of the detector (volumes, surfaces, etc.) with a buffer of potentially misaligned transforms. The new mechanism is tested by a new unit test `detector_align_cuda`. A demonstration of the usage of this mechanism added to the `cuda/detector_construction` tutorial code --- core/include/detray/core/detail/alignment.hpp | 34 ++++++++ .../unit_tests/device/cuda/detector_cuda.cpp | 81 +++++++++++++++++++ .../device/cuda/detector_cuda_kernel.cu | 54 +++++++++++++ .../device/cuda/detector_cuda_kernel.hpp | 7 ++ .../src/device/cuda/detector_construction.cpp | 37 ++++++++- .../src/device/cuda/detector_construction.cu | 4 + 6 files changed, 213 insertions(+), 4 deletions(-) create mode 100644 core/include/detray/core/detail/alignment.hpp 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..a573f6827 100644 --- a/tests/unit_tests/device/cuda/detector_cuda.cpp +++ b/tests/unit_tests/device/cuda/detector_cuda.cpp @@ -7,7 +7,10 @@ // Detray test include(s) #include "detector_cuda_kernel.hpp" +#include "detray/core/detail/alignment.hpp" +#include "detray/definitions/detail/algebra.hpp" #include "detray/test/utils/detectors/build_toy_detector.hpp" +#include "detray/test/utils/types.hpp" // Vecmem include(s) #include @@ -96,3 +99,81 @@ TEST(detector_cuda, detector) { EXPECT_EQ(cylinders_host[i] == cylinders_device[i], true); } } + +TEST(detector_cuda, detector_alignment) { + // memory resources + vecmem::host_memory_resource host_mr; + vecmem::cuda::device_memory_resource dev_mr; + vecmem::cuda::managed_memory_resource mng_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_buff_static = detray::get_buffer(det_host, dev_mr, cuda_cpy); + + // ---------- construct an "aligned" transform store --------- + // a few typedefs + using test_algebra = test::algebra; + using scalar = dscalar; + using point3 = dpoint3D; + + // build a vector of aligned transforms on the host + // for populating this vector take all transforms of the detector + // and shift them by the same translation + typename detector_host_t::transform_container tf_store_aligned_host; + + point3 shift{.1f * unit::mm, .2f * unit::mm, + .3f * unit::mm}; + + tf_store_aligned_host.reserve( + det_host.transform_store().size(), + typename decltype(det_host)::transform_container::context_type{}); + + for (const auto& tf : det_host.transform_store()) { + point3 shifted = tf.translation() + shift; + tf_store_aligned_host.push_back( + transform_t{shifted, tf.x(), tf.y(), tf.z()}); + } + + // copy the vector of aligned transforms to the device + // again, use synchronous copy and fixed size buffers + auto tf_buff_aligned = + get_buffer(tf_store_aligned_host, dev_mr, cuda_cpy, copy::sync, + vecmem::data::buffer_type::fixed_size); + + // Get the view of the aligned detector using the vector of aligned + // transforms and the static part of the detector copied to the device + // earlier + auto detector_view_aligned = + detail::misaligned_detector_view(det_buff_static, + tf_buff_aligned); + // Get the view of the static detector + auto detector_view_static = detray::get_data(det_buff_static); + + // make two vectors for surface transforms copied from device side + vecmem::vector surfacexf_device_static( + det_host.surfaces().size(), &mng_mr); + vecmem::vector surfacexf_device_aligned( + det_host.surfaces().size(), &mng_mr); + // views of the above vectors + auto surfacexf_data_static = vecmem::get_data(surfacexf_device_static); + auto surfacexf_data_aligned = vecmem::get_data(surfacexf_device_aligned); + + // run the test code to extract the surface transforms for the static + // and misaligned detector views and to store them into the vectors + detector_alignment_test(detector_view_static, detector_view_aligned, + surfacexf_data_static, surfacexf_data_aligned); + + // check that the relevant transforms have been properly shifted + for (unsigned int i = 0u; i < surfacexf_device_static.size(); i++) { + auto transdiff = surfacexf_device_aligned[i].translation() - + surfacexf_device_static[i].translation(); + EXPECT_EQ(transdiff == shift, true); + } +} diff --git a/tests/unit_tests/device/cuda/detector_cuda_kernel.cu b/tests/unit_tests/device/cuda/detector_cuda_kernel.cu index 28d0dea60..18b8122a8 100644 --- a/tests/unit_tests/device/cuda/detector_cuda_kernel.cu +++ b/tests/unit_tests/device/cuda/detector_cuda_kernel.cu @@ -6,6 +6,7 @@ */ #include "detray/definitions/detail/cuda_definitions.hpp" +#include "detray/geometry/tracking_surface.hpp" // Detray test include(s) #include "detector_cuda_kernel.hpp" @@ -107,4 +108,57 @@ void detector_test(typename detector_host_t::view_type det_data, DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize()); } +// cuda kernel to extract surface transforms from two detector views - static +// and misaligned - and to copy them into vectors +__global__ void detector_alignment_test_kernel( + typename detector_host_t::view_type det_data_static, + typename detector_host_t::view_type det_data_aligned, + vecmem::data::vector_view surfacexf_data_static, + vecmem::data::vector_view surfacexf_data_aligned) { + + auto ctx = typename detector_host_t::geometry_context{}; + + // two instances of device detectors + detector_device_t det_device_static(det_data_static); + detector_device_t det_device_aligned(det_data_aligned); + + // device vectors of surface transforms + vecmem::device_vector surfacexf_device_static( + surfacexf_data_static); + vecmem::device_vector surfacexf_device_aligned( + surfacexf_data_aligned); + + // copy surface transforms into relevant vectors + for (unsigned int i = 0u; i < det_device_static.surfaces().size(); i++) { + const auto sf = tracking_surface{det_device_static, + det_device_static.surfaces()[i]}; + surfacexf_device_static[i] = sf.transform(ctx); + } + + for (unsigned int i = 0u; i < det_device_aligned.surfaces().size(); i++) { + const auto sf = tracking_surface{det_device_aligned, + det_device_aligned.surfaces()[i]}; + surfacexf_device_aligned[i] = sf.transform(ctx); + } +} + +/// implementation of the alignment test function for detector +void detector_alignment_test( + typename detector_host_t::view_type det_data_static, + typename detector_host_t::view_type det_data_aligned, + vecmem::data::vector_view surfacexf_data_static, + vecmem::data::vector_view surfacexf_data_aligned) { + constexpr int block_dim = 1u; + constexpr int thread_dim = 1u; + + // run the test kernel + detector_alignment_test_kernel<<>>( + det_data_static, det_data_aligned, surfacexf_data_static, + surfacexf_data_aligned); + + // cuda error check + DETRAY_CUDA_ERROR_CHECK(cudaGetLastError()); + DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize()); +} + } // namespace detray diff --git a/tests/unit_tests/device/cuda/detector_cuda_kernel.hpp b/tests/unit_tests/device/cuda/detector_cuda_kernel.hpp index 2d1353cb9..ccb289a6a 100644 --- a/tests/unit_tests/device/cuda/detector_cuda_kernel.hpp +++ b/tests/unit_tests/device/cuda/detector_cuda_kernel.hpp @@ -41,4 +41,11 @@ void detector_test(typename detector_host_t::view_type det_data, vecmem::data::vector_view discs_data, vecmem::data::vector_view cylinders_data); +/// declaration of an alignment test function for detector +void detector_alignment_test( + typename detector_host_t::view_type det_data_static, + typename detector_host_t::view_type det_data_aligned, + vecmem::data::vector_view surfacexf_data_static, + vecmem::data::vector_view surfacexf_data_aligned); + } // namespace detray 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",