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

Author: Vakho Tsulaia

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

tests/unit_tests/device/cuda/detector_cuda.cpp

@@ -7,7 +7,11 @@
 
 // Detray test include(s)
 #include "detector_cuda_kernel.hpp"
+#include "detray/core/detail/alignment.hpp"
+#include "detray/definitions/detail/algebra.hpp"
+#include "detray/test/common/assert.hpp"
 #include "detray/test/utils/detectors/build_toy_detector.hpp"
+#include "detray/test/utils/types.hpp"
 
 // Vecmem include(s)
 #include <vecmem/memory/cuda/device_memory_resource.hpp>
@@ -96,3 +100,83 @@ TEST(detector_cuda, detector) {
         EXPECT_EQ(cylinders_host[i] == cylinders_device[i], true);
     }
 }
+
+TEST(detector_cuda, detector_alignment) {
+    // a few typedefs
+    using test_algebra = test::algebra;
+    using scalar = dscalar<test_algebra>;
+    using point3 = dpoint3D<test_algebra>;
+
+    // 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<test_algebra>(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 ---------
+
+    // 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<scalar>::mm, .2f * unit<scalar>::mm,
+                 .3f * unit<scalar>::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<detector_host_t>(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<transform_t> surfacexf_device_static(
+        det_host.surfaces().size(), &mng_mr);
+    vecmem::vector<transform_t> 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 translation_static = surfacexf_device_static[i].translation();
+        auto translation_aligned = surfacexf_device_aligned[i].translation();
+        auto translation_diff = translation_aligned - translation_static;
+        EXPECT_POINT3_NEAR(translation_diff, shift, 1e-4);
+    }
+}

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<transform_t> surfacexf_data_static,
+    vecmem::data::vector_view<transform_t> 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<transform_t> surfacexf_device_static(
+        surfacexf_data_static);
+    vecmem::device_vector<transform_t> 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<transform_t> surfacexf_data_static,
+    vecmem::data::vector_view<transform_t> surfacexf_data_aligned) {
+    constexpr int block_dim = 1u;
+    constexpr int thread_dim = 1u;
+
+    // run the test kernel
+    detector_alignment_test_kernel<<<block_dim, thread_dim>>>(
+        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

tests/unit_tests/device/cuda/detector_cuda_kernel.hpp

@@ -44,4 +44,11 @@ void detector_test(typename detector_host_t::view_type det_data,
                    vecmem::data::vector_view<disc_t> discs_data,
                    vecmem::data::vector_view<cylinder_t> 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<transform_t> surfacexf_data_static,
+    vecmem::data::vector_view<transform_t> surfacexf_data_aligned);
+
 }  // namespace detray

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)
@@ -24,6 +25,7 @@
 int main() {
 
     using algebra_t = detray::tutorial::algebra_t;
+    using scalar = detray::tutorial::scalar;
 
     // memory resource(s)
     vecmem::host_memory_resource host_mr;
@@ -81,10 +83,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);
@@ -99,11 +100,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<scalar>::mm,
+                                   .2f * detray::unit<scalar>::mm,
+                                   .3f * detray::unit<scalar>::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<host_detector_type>(
+            det_custom_buff, trf_buff_shifted);
+
+    std::cout << "\nCustom buffer setup (shifted):" << std::endl;
+    detray::tutorial::print(detector_view);
 }

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<mask_id::e_rectangle2>().size());
     printf("Number of trapezoids: %d\n",