Implement heterogeneous Association Map based on SoA blocks

DataFormats/SoATemplate/README.md

@@ -101,6 +101,11 @@ In addition to those fully parametrized templates, two further levels of paramet
 mirroring the structure of the underlying structs. The blocks are built via composition, 
 and access to individual layouts and views is provided by name.
 
+`SoABlocks` also have the possibility of generating methods for the `View` and `ConstView` classes using the macros `SOA_VIEW_METHODS`
+and `SOA_CONST_VIEW_METHODS`. Like the macros for the element methods, this can also be called only once, and if more methods
+have to be generated, they must be listed inside the same macro call. Since these methods can be called from the device,
+they must be prefixed with the `SOA_HOST_DEVICE` macro and, when possible, with the `constexpr` keyword.
+
 TODOs:
 - Add introspection utilities to print the structure and layout of a `SoABlocks` instance.
 
@@ -234,6 +239,43 @@ using SoAView = SoA::View;
 using SoAConstView = SoA::ConstView;
 ```
 
+as well as methods that operate on the View:
+
+```C++
+GENERATE_SOA_LAYOUT(PositionLayout,
+                    SOA_COLUMN(float, x),
+                    SOA_COLUMN(float, y),
+                    SOA_COLUMN(float, z))
+
+GENERATE_SOA_LAYOUT(VelocityLayout,
+                    SOA_COLUMN(float, vx),
+                    SOA_COLUMN(float, vy),
+                    SOA_COLUMN(float, vz))
+
+GENERATE_SOA_BLOCKS(PointsLayout,
+                    SOA_BLOCK(position, PositionLayout),
+                    SOA_BLOCK(velocity, VelocityLayout),
+                    SOA_VIEW_METHODS(
+                        SOA_HOST_DEVICE void update_position(uint32_t i, float time) {
+                            auto pos = this->position()[i];
+                            auto vel = this->velocity()[i];
+                            pos.x() += vel.vx() * time;
+                            pos.y() += vel.vy() * time;
+                            pos.z() += vel.vz() * time;
+                        }
+                    ),
+                    SOA_CONST_VIEW_METHODS(
+                        SOA_HOST_DEVICE auto distance2(uint32_t i, uint32_t j) const {
+                            auto pi = this->position()[i];
+                            auto pj = this->position()[j];
+                            return (pi.x() - pj.x()) * (pi.x() - pj.x()) + 
+                                   (pi.y() - pj.y()) * (pi.y() - pj.y()) + 
+                                   (pi.z() - pj.z()) * (pi.z() - pj.z());
+                        }
+                    )
+)
+```
+
 The buffer of the proper size is allocated, and the layout is populated with:
 
 ```C++

DataFormats/SoATemplate/interface/SoABlocks.h

@@ -466,6 +466,7 @@
                                                                                                                        \
       /* Accessors for the const views for each block */                                                               \
       _ITERATE_ON_ALL(_DECLARE_ACCESSORS_CONST_VIEW_BLOCKS, ~, __VA_ARGS__)                                            \
+      ENUM_IF_VALID(_ITERATE_ON_ALL(GENERATE_CONST_VIEW_METHODS, ~, __VA_ARGS__))                                      \
                                                                                                                        \
       private:                                                                                                         \
         _ITERATE_ON_ALL(_DECLARE_MEMBERS_CONST_VIEW_BLOCKS, ~, __VA_ARGS__)                                            \
@@ -548,6 +549,7 @@
                                                                                                                        \
       /* Accessors for the views for each block */                                                                     \
       _ITERATE_ON_ALL(_DECLARE_ACCESSORS_VIEW_BLOCKS, ~, __VA_ARGS__)                                                  \
+      ENUM_IF_VALID(_ITERATE_ON_ALL(GENERATE_VIEW_METHODS, ~, __VA_ARGS__))                                            \
                                                                                                                        \
        /* Data members inherited from the ConstView */                                                                 \
     };                                                                                                                 \

DataFormats/SoATemplate/test/SoACustomizedMethods_t.cc

@@ -30,7 +30,7 @@ TEST_CASE("SoACustomizedMethods") {
   }
   view.detectorType() = 42;
 
-  SECTION("ConstView methods") {
+  SECTION("ConstElement methods") {
     // arrays of norms
     std::array<float, elems> position_norms;
     std::array<double, elems> velocity_norms;
@@ -46,7 +46,7 @@ TEST_CASE("SoACustomizedMethods") {
     }
   }
 
-  SECTION("View methods") {
+  SECTION("Element methods") {
     // array of times
     std::array<double, elems> times;
 
@@ -71,4 +71,33 @@ TEST_CASE("SoACustomizedMethods") {
       REQUIRE_THAT(view[i].square_norm_velocity(), Catch::Matchers::WithinAbs(1., 1.e-9));
     }
   }
+
+  const auto points_sizes = std::array<cms::soa::size_type, 2>{{2, 2}};
+  const std::size_t blocksBufferSize = Points::computeDataSize(points_sizes);
+
+  std::unique_ptr<std::byte, decltype(std::free) *> points_buffer{
+      reinterpret_cast<std::byte *>(aligned_alloc(Points::alignment, blocksBufferSize)), std::free};
+
+  Points points(points_buffer.get(), points_sizes);
+  PointsView points_view{points};
+  PointsConstView points_const_view{points};
+  points_view.position()[0].x() = 2.f;
+  points_view.position()[0].y() = 4.f;
+  points_view.position()[0].z() = 3.f;
+  points_view.position()[1].x() = 1.f;
+  points_view.position()[1].y() = 1.f;
+  points_view.position()[1].z() = 1.f;
+
+  SECTION("View methods") { REQUIRE(points_const_view.distance2(0, 1) == 14.f); }
+
+  SECTION("ConstView methods") {
+    points_view.velocity()[0].vx() = 1.f;
+    points_view.velocity()[0].vy() = 3.f;
+    points_view.velocity()[0].vz() = 5.f;
+    const auto time = .5f;
+    points_view.update_position(0, time);
+    REQUIRE(points_view.position()[0].x() == 2.5f);
+    REQUIRE(points_view.position()[0].y() == 5.5f);
+    REQUIRE(points_view.position()[0].z() == 5.5f);
+  }
 }

DataFormats/SoATemplate/test/SoACustomizedMethods_t.cu

@@ -24,6 +24,15 @@ __global__ void checkNormalise(SoAView soaView, double* checkTimesFunction) {
   soaView[i].normalise();
 }
 
+__global__ void checkPointsDistance(PointsConstView view, bool* result) { *result &= (view.distance2(0, 1) == 14.f); }
+
+__global__ void checkPointsPositionUpdate(PointsView view, float time, bool* result) {
+  view.update_position(0, time);
+  *result &= (view.position()[0].x() == 2.5f);
+  *result &= (view.position()[0].y() == 5.5f);
+  *result &= (view.position()[0].z() == 5.5f);
+}
+
 TEST_CASE("SoACustomizedMethods CUDA", "[SoACustomizedMethods][cuda]") {
   // common number of elements for the SoAs
   const std::size_t elems = 10;
@@ -69,7 +78,7 @@ TEST_CASE("SoACustomizedMethods CUDA", "[SoACustomizedMethods][cuda]") {
   // Host → Device copy
   CUDA_CHECK(cudaMemcpy(d_buf, h_buf, bufferSize, cudaMemcpyHostToDevice));
 
-  SECTION("ConstView methods CUDA") {
+  SECTION("ConstElement methods CUDA") {
     calculateNorm<<<(elems + 255) / 256, 256>>>(d_Constview, d_position_norms, d_velocity_norms);
 
     CUDA_CHECK(cudaMemcpy(h_position_norms.data(), d_position_norms, elems * sizeof(float), cudaMemcpyDeviceToHost));
@@ -88,7 +97,7 @@ TEST_CASE("SoACustomizedMethods CUDA", "[SoACustomizedMethods][cuda]") {
     }
   }
 
-  SECTION("View methods CUDA") {
+  SECTION("Element methods CUDA") {
     std::array<double, elems> times;
 
     // Check for the correctness of the time() function
@@ -116,10 +125,55 @@ TEST_CASE("SoACustomizedMethods CUDA", "[SoACustomizedMethods][cuda]") {
     }
   }
 
+  const auto points_sizes = std::array<cms::soa::size_type, 2>{{2, 2}};
+  const std::size_t blocksBufferSize = Points::computeDataSize(points_sizes);
+
+  std::byte* points_buffer_host = nullptr;
+  CUDA_CHECK(cudaMallocHost(&points_buffer_host, blocksBufferSize));
+
+  Points h_points(points_buffer_host, points_sizes);
+  PointsView h_points_view{h_points};
+  h_points_view.position()[0].x() = 2.f;
+  h_points_view.position()[0].y() = 4.f;
+  h_points_view.position()[0].z() = 3.f;
+  h_points_view.position()[1].x() = 1.f;
+  h_points_view.position()[1].y() = 1.f;
+  h_points_view.position()[1].z() = 1.f;
+  h_points_view.velocity()[0].vx() = 1.f;
+  h_points_view.velocity()[0].vy() = 3.f;
+  h_points_view.velocity()[0].vz() = 5.f;
+
+  std::byte* points_buffer_device = nullptr;
+  CUDA_CHECK(cudaMalloc(&points_buffer_device, blocksBufferSize));
+
+  Points d_points(points_buffer_device, points_sizes);
+  PointsView d_points_view{d_points};
+  PointsConstView d_points_const_view{d_points};
+
+  CUDA_CHECK(cudaMemcpy(points_buffer_device, points_buffer_host, blocksBufferSize, cudaMemcpyHostToDevice));
+
+  bool* d_result = nullptr;
+  CUDA_CHECK(cudaMalloc(&d_result, sizeof(bool)));
+  CUDA_CHECK(cudaMemset(d_result, 1, sizeof(bool)));
+
+  SECTION("View methods") { checkPointsDistance<<<1, 1>>>(d_points_const_view, d_result); }
+
+  SECTION("ConstView methods") {
+    const auto time = .5f;
+    checkPointsPositionUpdate<<<1, 1>>>(d_points_view, time, d_result);
+  }
+
+  bool h_result;
+  CUDA_CHECK(cudaMemcpy(&h_result, d_result, sizeof(bool), cudaMemcpyDeviceToHost));
+  REQUIRE(h_result);
+
   // ===== cleanup =====
   CUDA_CHECK(cudaFree(d_position_norms));
   CUDA_CHECK(cudaFree(d_velocity_norms));
   CUDA_CHECK(cudaFree(d_times));
   CUDA_CHECK(cudaFree(d_buf));
+  CUDA_CHECK(cudaFree(points_buffer_device));
+  CUDA_CHECK(cudaFree(d_result));
   CUDA_CHECK(cudaFreeHost(h_buf));
+  CUDA_CHECK(cudaFreeHost(points_buffer_host));
 }

DataFormats/SoATemplate/test/SoACustomizedMethods_t.hip.cc

@@ -26,6 +26,15 @@ __global__ void checkNormalise(SoAView soaView, double* checkTimesFunction) {
   soaView[i].normalise();
 }
 
+__global__ void checkPointsDistance(PointsConstView view, bool* result) { *result &= (view.distance2(0, 1) == 14.f); }
+
+__global__ void checkPointsPositionUpdate(PointsView view, float time, bool* result) {
+  view.update_position(0, time);
+  *result &= (view.position()[0].x() == 2.5f);
+  *result &= (view.position()[0].y() == 5.5f);
+  *result &= (view.position()[0].z() == 5.5f);
+}
+
 TEST_CASE("SoACustomizedMethods hip", "[SoACustomizedMethods][hip]") {
   // common number of elements for the SoAs
   const std::size_t elems = 10;
@@ -71,7 +80,7 @@ TEST_CASE("SoACustomizedMethods hip", "[SoACustomizedMethods][hip]") {
   // Host → Device copy
   HIP_CHECK(hipMemcpy(d_buf, h_buf, bufferSize, hipMemcpyHostToDevice));
 
-  SECTION("ConstView methods HIP") {
+  SECTION("ConstElement methods HIP") {
     calculateNorm<<<(elems + 255) / 256, 256>>>(d_Constview, d_position_norms, d_velocity_norms);
 
     HIP_CHECK(hipMemcpy(h_position_norms.data(), d_position_norms, elems * sizeof(float), hipMemcpyDeviceToHost));
@@ -90,7 +99,7 @@ TEST_CASE("SoACustomizedMethods hip", "[SoACustomizedMethods][hip]") {
     }
   }
 
-  SECTION("View methods HIP") {
+  SECTION("Element methods HIP") {
     std::array<double, elems> times;
 
     // Check for the correctness of the time() function
@@ -118,10 +127,55 @@ TEST_CASE("SoACustomizedMethods hip", "[SoACustomizedMethods][hip]") {
     }
   }
 
+  const auto points_sizes = std::array<cms::soa::size_type, 2>{{2, 2}};
+  const std::size_t blocksBufferSize = Points::computeDataSize(points_sizes);
+
+  std::byte* points_buffer_host = nullptr;
+  HIP_CHECK(hipHostMalloc(&points_buffer_host, blocksBufferSize));
+
+  Points h_points(points_buffer_host, points_sizes);
+  PointsView h_points_view{h_points};
+  h_points_view.position()[0].x() = 2.f;
+  h_points_view.position()[0].y() = 4.f;
+  h_points_view.position()[0].z() = 3.f;
+  h_points_view.position()[1].x() = 1.f;
+  h_points_view.position()[1].y() = 1.f;
+  h_points_view.position()[1].z() = 1.f;
+  h_points_view.velocity()[0].vx() = 1.f;
+  h_points_view.velocity()[0].vy() = 3.f;
+  h_points_view.velocity()[0].vz() = 5.f;
+
+  std::byte* points_buffer_device = nullptr;
+  HIP_CHECK(hipMalloc(&points_buffer_device, blocksBufferSize));
+
+  Points d_points(points_buffer_device, points_sizes);
+  PointsView d_points_view{d_points};
+  PointsConstView d_points_const_view{d_points};
+
+  HIP_CHECK(hipMemcpy(points_buffer_device, points_buffer_host, blocksBufferSize, hipMemcpyHostToDevice));
+
+  bool* d_result = nullptr;
+  HIP_CHECK(hipMalloc(&d_result, sizeof(bool)));
+  HIP_CHECK(hipMemset(d_result, 1, sizeof(bool)));
+
+  SECTION("View methods") { checkPointsDistance<<<1, 1>>>(d_points_const_view, d_result); }
+
+  SECTION("ConstView methods") {
+    const auto time = .5f;
+    checkPointsPositionUpdate<<<1, 1>>>(d_points_view, time, d_result);
+  }
+
+  bool h_result;
+  HIP_CHECK(hipMemcpy(&h_result, d_result, sizeof(bool), hipMemcpyDeviceToHost));
+  REQUIRE(h_result);
+
   // ===== cleanup =====
   HIP_CHECK(hipFree(d_position_norms));
   HIP_CHECK(hipFree(d_velocity_norms));
   HIP_CHECK(hipFree(d_times));
   HIP_CHECK(hipFree(d_buf));
+  HIP_CHECK(hipFree(points_buffer_device));
+  HIP_CHECK(hipFree(d_result));
   HIP_CHECK(hipFreeHost(h_buf));
+  HIP_CHECK(hipFreeHost(points_buffer_host));
 }

DataFormats/SoATemplate/test/SoADefinition_CustomizedMethods.h

@@ -1,6 +1,7 @@
 #include <cmath>
 
 #include "DataFormats/SoATemplate/interface/SoALayout.h"
+#include "DataFormats/SoATemplate/interface/SoABlocks.h"
 
 GENERATE_SOA_LAYOUT(SoATemplate,
                     SOA_COLUMN(float, x),
@@ -18,21 +19,21 @@ GENERATE_SOA_LAYOUT(SoATemplate,
                             x() /= norm_position;
                             y() /= norm_position;
                             z() /= norm_position;
-                          };
+                          }
                           double norm_velocity = square_norm_velocity();
                           if (norm_velocity > 0.0f) {
                             v_x() /= norm_velocity;
                             v_y() /= norm_velocity;
                             v_z() /= norm_velocity;
-                          };
+                          }
                         }),
 
                     SOA_CONST_ELEMENT_METHODS(
                         SOA_HOST_DEVICE float square_norm_position()
-                            const { return sqrt(x() * x() + y() * y() + z() * z()); };
+                            const { return sqrt(x() * x() + y() * y() + z() * z()); }
 
                         SOA_HOST_DEVICE double square_norm_velocity()
-                            const { return sqrt(v_x() * v_x() + v_y() * v_y() + v_z() * v_z()); };
+                            const { return sqrt(v_x() * v_x() + v_y() * v_y() + v_z() * v_z()); }
 
                         template <typename T1, typename T2>
                         SOA_HOST_DEVICE static auto time(T1 pos, T2 vel) {
@@ -46,3 +47,42 @@ GENERATE_SOA_LAYOUT(SoATemplate,
 using SoA = SoATemplate<>;
 using SoAView = SoA::View;
 using SoAConstView = SoA::ConstView;
+
+// clang-format off
+GENERATE_SOA_LAYOUT(PositionLayout,
+                    SOA_COLUMN(float, x),
+                    SOA_COLUMN(float, y),
+                    SOA_COLUMN(float, z))
+
+GENERATE_SOA_LAYOUT(VelocityLayout,
+                    SOA_COLUMN(float, vx),
+                    SOA_COLUMN(float, vy),
+                    SOA_COLUMN(float, vz))
+
+GENERATE_SOA_BLOCKS(PointsLayout,
+                    SOA_BLOCK(position, PositionLayout),
+                    SOA_BLOCK(velocity, VelocityLayout),
+                    SOA_VIEW_METHODS(
+                        SOA_HOST_DEVICE void update_position(uint32_t i, float time) {
+                            auto pos = this->position()[i];
+                            auto vel = this->velocity()[i];
+                            pos.x() += vel.vx() * time;
+                            pos.y() += vel.vy() * time;
+                            pos.z() += vel.vz() * time;
+                        }
+                    ),
+                    SOA_CONST_VIEW_METHODS(
+                        SOA_HOST_DEVICE auto distance2(uint32_t i, uint32_t j) const {
+                            auto pi = this->position()[i];
+                            auto pj = this->position()[j];
+                            return (pi.x() - pj.x()) * (pi.x() - pj.x()) + 
+                                   (pi.y() - pj.y()) * (pi.y() - pj.y()) + 
+                                   (pi.z() - pj.z()) * (pi.z() - pj.z());
+                        }
+                    )
+)
+// clang-format on
+
+using Points = PointsLayout<>;
+using PointsView = Points::View;
+using PointsConstView = Points::ConstView;

DataFormats/TICL/BuildFile.xml

@@ -0,0 +1,3 @@
+<use name="alpaka"/>
+<use name="HeterogeneousCore/AlpakaInterface"/>
+<use name="DataFormats/SoATemplate"/>