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,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 <vecmem/memory/cuda/device_memory_resource.hpp>
@@ -96,3 +99,80 @@ 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<test_algebra>;
+    using point3 = dpoint3D<test_algebra>;
+
+    // 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 transdiff = surfacexf_device_aligned[i].translation()
+		- surfacexf_device_static[i].translation();
+	EXPECT_EQ(transdiff == shift, true);
+    }
+}

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,53 @@ 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

@@ -41,4 +41,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)
@@ -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",