New function: misaligned_detector_view (in core/detail/alignment.hpp)

As the name suggests, this function makes a misaligned detector view by combining "static"
buffers of the detector (volumes, surfaces, etc.) with a buffer of a potentially misaligned
transform store

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_trfstore_type>
+typename host_detector_type::view_type misaligned_detector_view(
+    typename host_detector_type::buffer_type& det_buffer,
+    typename host_trfstore_type::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

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);
@@ -102,4 +102,33 @@ int main() {
 
     std::cout << "\nCustom buffer setup:" << std::endl;
     detray::tutorial::print(detray::get_data(det_custom_buff));
+
+    // Construct an "aligned" transform store
+    typename decltype(det_host)::transform_container tfStore;
+    using tf_type = decltype(det_host)::transform_container::value_type;
+
+    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]};
+        tfStore.push_back(tf_type{shifted, tf.x(), tf.y(), tf.z()});
+    }
+
+    auto trf_buff_shifted =
+        detray::get_buffer(tfStore, dev_mr, cuda_cpy, detray::copy::sync,
+                           vecmem::data::buffer_type::fixed_size);
+
+    using host_detector_type = decltype(det_host);
+    using host_tf_container_type = decltype(det_host)::transform_container;
+    typename decltype(det_host)::view_type detector_view =
+        detray::detail::misaligned_detector_view<host_detector_type,
+                                                 host_tf_container_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",