NextGenTrigger Task 1.7 Meeting

Thursday 13 Feb 2025, 09:30 → 10:30 Europe/Zurich

32/S-C22 (CERN)

nvcc bug

Overview

The goal is to develop a C++ library that allows to abstract the data layout of an array. Possible data layouts include aray of struct (AoS) and struct of array (SoA), see the following example.

constexpr std::size_t N = 42;

struct Point { int x, y, z; };
Point point_aos[N];  // data layout: AoS

template <std::size_t N>
struct PointSoA {
    int x[N];
    int y[N];
    int z[N];
};

Point<N> point_soa;  // data layout: SoA

We aim at writing a class that takes the struct Point, a data layout, and possibly more arguments. The class then allows for AoS access, but stores the data in a possibly different layout, thereby hiding the data layout.

New: Unit Tests using GTest

GTest was integrated into the CMake build and tests were added. We have 8 Tests in total. We test the following underlying containers in AoS and SoA layout:

• std::vector
• std::span
• pre-existing buffer of type char*
• std::pmr::vector

These are the first template parameter F of our wrapper class that is abstracting the data layout. Instances of this class allow an AoS-syntax access via the operator[].

template<
    template <class> class F,                  // container
    template <template <class> class> class S, // e.g. "Point"
    layout L                                   // data layout
>
struct wrapper;

New: CUDA Support

We can now use our wrapper on the host and in the device. See the following example.

#include "wrapper.h"

template <class T>
using pointer_type = T*;

template <
    template <class> class F,
    template <template <class> class> class S,
    wrapper::layout L
>
__global__ void add(int N, wrapper::wrapper<F, S, L> w) {
    for (int i = 0; i < N; ++i) w[i].y = w[i].x + w[i].y;
}

struct Point2D { double x, y; };

template <template <class> class F>
struct S {
    F<int> x;
    F<int> y;
    F<Point2D> point;
    F<double> identifier;
    __host__ __device__ void setX(int x_new) { x = x_new; }
};

void main() {
    int N = 8;
    wrapper::wrapper<pointer_type, S, wrapper::layout::aos> w; // change for soa
    cudaMallocManaged(&w.data, N * sizeof(S<wrapper::value>)); // change for soa

    for (int i = 0; i < N; ++i) {
        S<wrapper::reference> r = w[i];
        r.setX(1);
        r.y = 2;
        r.point = {0.5 * i, 0.5 * i};
        r.identifier = 0.1 * i;
    }

    add<<<1, 1>>>(N, w);
    cudaDeviceSynchronize();
    cudaFree(w.data);                                          // change for soa
}

Remarks

• Only three lines change if we use SoA instead of Aos, but they can be outsourced to factory function
  
• The kernel "add" can deduce the template parameters of wrapper
• In particular, the kernel does not depend on the container nor on the data layout
• Some CUDA C++20 features were not working properly, thus had to roll back to C++17.

New: Proposal for Summer Student Project

We have submitted a proposal for a summer student project. The project is about benchmarking different data layouts.

09:30 → 09:35 Project updates

Speakers: Dr Andrea Bocci (CERN), Attila Krasznahorkay (University of Massachusetts (US)), Axel Naumann (CERN)

09:40 → 10:20 Reports and updates

Speaker: All

10:25 → 10:30 AOB