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,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/test/utils/types.hpp"
 #include "detray/test/utils/detectors/build_toy_detector.hpp"
 
 // Vecmem include(s)
@@ -96,3 +99,49 @@ 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::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 test_algebra = test::algebra;
+    using scalar = dscalar<test_algebra>;
+    using point3 = dpoint3D<test_algebra>;
+
+    typename decltype(det_host)::transform_container tf_store_aligned_host;
+    using tf_type = decltype(det_host)::transform_container::value_type;
+
+    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()[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<host_detector_type>(det_fixed_buff,
+                                                             tf_buff_shifted);
+}

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<detray::scalar>::mm,
+                                   .2f * detray::unit<detray::scalar>::mm,
+                                   .3f * detray::unit<detray::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",