diff --git a/src/alpaka/DataFormats/SoACommon.h b/src/alpaka/DataFormats/SoACommon.h new file mode 100644 index 000000000..e132a47f9 --- /dev/null +++ b/src/alpaka/DataFormats/SoACommon.h @@ -0,0 +1,578 @@ +/* + * Definitions of SoA common parameters for SoA class generators + */ + +#ifndef DataStructures_SoACommon_h +#define DataStructures_SoACommon_h + +#include "boost/preprocessor.hpp" +#include +#include +#include + +// CUDA attributes +#ifdef __CUDACC__ +#define SOA_HOST_ONLY ALPAKA_FN_HOST +#define SOA_DEVICE_ONLY ALPAKA_FN_ACC +#define SOA_HOST_DEVICE ALPAKA_FN_HOST_ACC +#define SOA_HOST_DEVICE_INLINE ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE +#else +#define SOA_HOST_ONLY +#define SOA_DEVICE_ONLY +#define SOA_HOST_DEVICE +#define SOA_HOST_DEVICE_INLINE inline +#endif + +// Exception throwing (or willful crash in kernels) +#if defined(__CUDACC__) && defined(__CUDA_ARCH__) +#define SOA_THROW_OUT_OF_RANGE(A) \ + { \ + printf(A "\n"); \ + *((char*)nullptr) = 0; \ + } +#else +#define SOA_THROW_OUT_OF_RANGE(A) \ + { throw std::out_of_range(A); } +#endif + +/* declare "scalars" (one value shared across the whole SoA) and "columns" (one value per element) */ +#define _VALUE_TYPE_SCALAR 0 +#define _VALUE_TYPE_COLUMN 1 +#define _VALUE_TYPE_EIGEN_COLUMN 2 + +namespace cms::soa { + + enum class SoAColumnType { + scalar = _VALUE_TYPE_SCALAR, + column = _VALUE_TYPE_COLUMN, + eigen = _VALUE_TYPE_EIGEN_COLUMN + }; + enum class RestrictQualify : bool { Enabled, Disabled, Default = Disabled }; + + enum class RangeChecking : bool { Enabled, Disabled, Default = Disabled }; + + template + struct add_restrict {}; + + template + struct add_restrict { + typedef T Value; + typedef T* __restrict__ Pointer; + typedef T& __restrict__ Reference; + typedef const T ConstValue; + typedef const T* __restrict__ PointerToConst; + typedef const T& __restrict__ ReferenceToConst; + }; + + template + struct add_restrict { + typedef T Value; + typedef T* Pointer; + typedef T& Reference; + typedef const T ConstValue; + typedef const T* PointerToConst; + typedef const T& ReferenceToConst; + }; + template + struct SoAParametersImpl; + + // Templated parameter sets for scalar columns and Eigen columns + template + struct SoAConstParametersImpl { + static const SoAColumnType columnType = COLUMN_TYPE; + typedef T ValueType; + typedef const ValueType* TupleOrPointerType; + const ValueType* addr_ = nullptr; + SOA_HOST_DEVICE_INLINE SoAConstParametersImpl(const ValueType* addr) : addr_(addr) {} + SOA_HOST_DEVICE_INLINE SoAConstParametersImpl(const SoAConstParametersImpl& o) { addr_ = o.addr_; } + SOA_HOST_DEVICE_INLINE SoAConstParametersImpl(const SoAParametersImpl& o) { + addr_ = o.addr_; + } + SOA_HOST_DEVICE_INLINE SoAConstParametersImpl() {} + static bool checkAlignement(ValueType* addr, size_t byteAlignment) { + return reinterpret_cast(addr) % byteAlignment; + } + }; + + template + struct SoAConstParametersImpl { + static const SoAColumnType columnType = SoAColumnType::eigen; + typedef T ValueType; + typedef typename T::Scalar ScalarType; + typedef std::tuple TupleOrPointerType; + const ScalarType* addr_ = nullptr; + size_t stride_ = 0; + SOA_HOST_DEVICE_INLINE SoAConstParametersImpl(const ScalarType* addr, size_t stride) + : addr_(addr), stride_(stride) {} + SOA_HOST_DEVICE_INLINE SoAConstParametersImpl(const TupleOrPointerType tuple) + : addr_(std::get<0>(tuple)), stride_(std::get<1>(tuple)) {} + SOA_HOST_DEVICE_INLINE SoAConstParametersImpl(const ScalarType* addr) : addr_(addr) {} + // Trick setter + return self-reference allowing commat-free 2-stage construction in macro contexts (in combination with the + // addr-only constructor. + SoAConstParametersImpl& setStride(size_t stride) { + stride_ = stride; + return *this; + } + SOA_HOST_DEVICE_INLINE SoAConstParametersImpl(const SoAConstParametersImpl& o) { + addr_ = o.addr_; + stride_ = o.stride_; + } + SOA_HOST_DEVICE_INLINE SoAConstParametersImpl(const SoAParametersImpl& o) { + addr_ = o.addr_; + stride_ = o.stride_; + } + SOA_HOST_DEVICE_INLINE SoAConstParametersImpl() {} + static bool checkAlignement(const TupleOrPointerType tuple, size_t byteAlignment) { + const auto& [addr, stride] = tuple; + return reinterpret_cast(addr) % byteAlignment; + } + }; + + // Matryoshka template to avoiding commas in macros + template + struct SoAConstParameters_ColumnType { + template + struct DataType : public SoAConstParametersImpl { + using SoAConstParametersImpl::SoAConstParametersImpl; + }; + }; + + // Templated parameter sets for scalar columns and Eigen columns + template + struct SoAParametersImpl { + static const SoAColumnType columnType = COLUMN_TYPE; + typedef T ValueType; + typedef const ValueType* TupleOrPointerType; + typedef SoAConstParametersImpl ConstType; + friend ConstType; + ValueType* addr_ = nullptr; + SOA_HOST_DEVICE_INLINE SoAParametersImpl(ValueType* addr) : addr_(addr) {} + SOA_HOST_DEVICE_INLINE SoAParametersImpl() {} + static bool checkAlignement(ValueType* addr, size_t byteAlignment) { + return reinterpret_cast(addr) % byteAlignment; + } + }; + + template + struct SoAParametersImpl { + static const SoAColumnType columnType = SoAColumnType::eigen; + typedef T ValueType; + typedef SoAConstParametersImpl ConstType; + friend ConstType; + typedef typename T::Scalar ScalarType; + typedef std::tuple TupleOrPointerType; + ScalarType* addr_ = nullptr; + size_t stride_ = 0; + SOA_HOST_DEVICE_INLINE SoAParametersImpl(ScalarType* addr, size_t stride) + : addr_(addr), stride_(stride) {} + SOA_HOST_DEVICE_INLINE SoAParametersImpl(const TupleOrPointerType tuple) + : addr_(std::get<0>(tuple)), stride_(std::get<1>(tuple)) {} + SOA_HOST_DEVICE_INLINE SoAParametersImpl() {} + SOA_HOST_DEVICE_INLINE SoAParametersImpl(ScalarType* addr) : addr_(addr) {} + // Trick setter + return self-reference allowing commat-free 2-stage construction in macro contexts (in combination with the + // addr-only constructor. + SoAParametersImpl& setStride(size_t stride) { + stride_ = stride; + return *this; + } + static bool checkAlignement(const TupleOrPointerType tuple, size_t byteAlignment) { + const auto& [addr, stride] = tuple; + return reinterpret_cast(addr) % byteAlignment; + } + }; + + // Matryoshka template to avoiding commas in macros + template + struct SoAParameters_ColumnType { + template + struct DataType : public SoAParametersImpl { + using SoAParametersImpl::SoAParametersImpl; + }; + }; + + // Helper template managing the value within it column + // The optional compile time alignment parameter enables informing the + // compiler of alignment (enforced by caller). + template + class SoAValue { + // Eigen is implemented in a specialization + static_assert(COLUMN_TYPE != SoAColumnType::eigen); + + public: + typedef add_restrict Restr; + typedef typename Restr::Value Val; + typedef typename Restr::Pointer Ptr; + typedef typename Restr::Reference Ref; + typedef typename Restr::PointerToConst PtrToConst; + typedef typename Restr::ReferenceToConst RefToConst; + SOA_HOST_DEVICE_INLINE SoAValue(size_t i, T* col) : idx_(i), col_(col) {} + SOA_HOST_DEVICE_INLINE SoAValue(size_t i, SoAParametersImpl params) + : idx_(i), col_(params.addr_) {} + /* SOA_HOST_DEVICE_INLINE operator T&() { return col_[idx_]; } */ + SOA_HOST_DEVICE_INLINE Ref operator()() { + // Ptr type will add the restrict qualifyer if needed + Ptr col = alignedCol(); + return col[idx_]; + } + SOA_HOST_DEVICE_INLINE RefToConst operator()() const { + // PtrToConst type will add the restrict qualifyer if needed + PtrToConst col = alignedCol(); + return col[idx_]; + } + SOA_HOST_DEVICE_INLINE Ptr operator&() { return &alignedCol()[idx_]; } + SOA_HOST_DEVICE_INLINE PtrToConst operator&() const { return &alignedCol()[idx_]; } + template + SOA_HOST_DEVICE_INLINE Ref operator=(const T2& v) { + return alignedCol()[idx_] = v; + } + typedef Val valueType; + static constexpr auto valueSize = sizeof(T); + + private: + SOA_HOST_DEVICE_INLINE Ptr alignedCol() const { + if constexpr (ALIGNMENT) { + return reinterpret_cast(__builtin_assume_aligned(col_, ALIGNMENT)); + } + return reinterpret_cast(col_); + } + size_t idx_; + T* col_; + }; + + // Helper template managing the value within it column + // TODO Create a const variant to avoid leaking mutable access. +#ifdef EIGEN_WORLD_VERSION + template + class SoAValue { + public: + typedef C Type; + typedef Eigen::Map> MapType; + typedef Eigen::Map> CMapType; + typedef add_restrict Restr; + typedef typename Restr::Value Val; + typedef typename Restr::Pointer Ptr; + typedef typename Restr::Reference Ref; + typedef typename Restr::PointerToConst PtrToConst; + typedef typename Restr::ReferenceToConst RefToConst; + SOA_HOST_DEVICE_INLINE SoAValue(size_t i, typename C::Scalar* col, size_t stride) + : val_(col + i, C::RowsAtCompileTime, C::ColsAtCompileTime, Eigen::InnerStride(stride)), + crCol_(col), + cVal_(crCol_ + i, C::RowsAtCompileTime, C::ColsAtCompileTime, Eigen::InnerStride(stride)), + stride_(stride) {} + SOA_HOST_DEVICE_INLINE SoAValue(size_t i, SoAParametersImpl params) + : val_(params.addr_ + i, + C::RowsAtCompileTime, + C::ColsAtCompileTime, + Eigen::InnerStride(params.stride_)), + crCol_(params.addr_), + cVal_(crCol_ + i, + C::RowsAtCompileTime, + C::ColsAtCompileTime, + Eigen::InnerStride(params.stride_)), + stride_(params.stride_) {} + SOA_HOST_DEVICE_INLINE MapType& operator()() { return val_; } + SOA_HOST_DEVICE_INLINE const CMapType& operator()() const { return cVal_; } + SOA_HOST_DEVICE_INLINE operator C() { return val_; } + SOA_HOST_DEVICE_INLINE operator const C() const { return cVal_; } + SOA_HOST_DEVICE_INLINE C* operator&() { return &val_; } + SOA_HOST_DEVICE_INLINE const C* operator&() const { return &cVal_; } + template + SOA_HOST_DEVICE_INLINE MapType& operator=(const C2& v) { + return val_ = v; + } + typedef typename C::Scalar ValueType; + static constexpr auto valueSize = sizeof(C::Scalar); + SOA_HOST_DEVICE_INLINE size_t stride() const { return stride_; } + + private: + MapType val_; + const Ptr crCol_; + CMapType cVal_; + size_t stride_; + }; +#else + template + class SoAValue { + // Eigen/Core should be pre-included before the SoA headers to enable support for Eigen columns. + static_assert(!sizeof(C), + "Eigen/Core should be pre-included before the SoA headers to enable support for Eigen columns."); + }; +#endif + // Helper template managing the value within it column + template + class SoAConstValue { + // Eigen is implemented in a specialization + static_assert(COLUMN_TYPE != SoAColumnType::eigen); + + public: + typedef add_restrict Restr; + typedef typename Restr::Value Val; + typedef typename Restr::Pointer Ptr; + typedef typename Restr::Reference Ref; + typedef typename Restr::PointerToConst PtrToConst; + typedef typename Restr::ReferenceToConst RefToConst; + typedef SoAParametersImpl Params; + typedef SoAConstParametersImpl ConstParams; + SOA_HOST_DEVICE_INLINE SoAConstValue(size_t i, const T* col) : idx_(i), col_(col) {} + SOA_HOST_DEVICE_INLINE SoAConstValue(size_t i, SoAParametersImpl params) + : idx_(i), col_(params.addr_) {} + SOA_HOST_DEVICE_INLINE SoAConstValue(size_t i, SoAConstParametersImpl params) + : idx_(i), col_(params.addr_) {} + /* SOA_HOST_DEVICE_INLINE operator T&() { return col_[idx_]; } */ + SOA_HOST_DEVICE_INLINE RefToConst operator()() const { + // Ptr type will add the restrict qualifyer if needed + PtrToConst col = alignedCol(); + return col[idx_]; + } + SOA_HOST_DEVICE_INLINE const T* operator&() const { return &alignedCol()[idx_]; } + typedef T valueType; + static constexpr auto valueSize = sizeof(T); + + private: + SOA_HOST_DEVICE_INLINE PtrToConst alignedCol() const { + if constexpr (ALIGNMENT) { + return reinterpret_cast(__builtin_assume_aligned(col_, ALIGNMENT)); + } + return reinterpret_cast(col_); + } + size_t idx_; + const T* col_; + }; + +#ifdef EIGEN_WORLD_VERSION + // Helper template managing the value within it column + // TODO Create a const variant to avoid leaking mutable access. + template + class SoAConstValue { + public: + typedef C Type; + typedef Eigen::Map> CMapType; + typedef CMapType& RefToConst; + typedef SoAConstParametersImpl ConstParams; + SOA_HOST_DEVICE_INLINE SoAConstValue(size_t i, typename C::Scalar* col, size_t stride) + : crCol_(col), + cVal_(crCol_ + i, C::RowsAtCompileTime, C::ColsAtCompileTime, Eigen::InnerStride(stride)), + stride_(stride) {} + SOA_HOST_DEVICE_INLINE SoAConstValue(size_t i, SoAConstParametersImpl params) + : crCol_(params.addr_), + cVal_(crCol_ + i, + C::RowsAtCompileTime, + C::ColsAtCompileTime, + Eigen::InnerStride(params.stride_)), + stride_(params.stride_) {} + SOA_HOST_DEVICE_INLINE const CMapType& operator()() const { return cVal_; } + SOA_HOST_DEVICE_INLINE operator const C() const { return cVal_; } + SOA_HOST_DEVICE_INLINE const C* operator&() const { return &cVal_; } + typedef typename C::Scalar ValueType; + static constexpr auto valueSize = sizeof(C::Scalar); + SOA_HOST_DEVICE_INLINE size_t stride() const { return stride_; } + + private: + const typename C::Scalar* __restrict__ crCol_; + CMapType cVal_; + size_t stride_; + }; +#else + template + class SoAConstValue { + // Eigen/Core should be pre-included before the SoA headers to enable support for Eigen columns. + static_assert(!sizeof(C), + "Eigen/Core should be pre-included before the SoA headers to enable support for Eigen columns."); + }; +#endif + + // Helper template to avoid commas in macro +#ifdef EIGEN_WORLD_VERSION + template + struct EigenConstMapMaker { + typedef Eigen::Map> Type; + class DataHolder { + public: + DataHolder(const typename C::Scalar* data) : data_(data) {} + EigenConstMapMaker::Type withStride(size_t stride) { + return EigenConstMapMaker::Type( + data_, C::RowsAtCompileTime, C::ColsAtCompileTime, Eigen::InnerStride(stride)); + } + + private: + const typename C::Scalar* const data_; + }; + static DataHolder withData(const typename C::Scalar* data) { return DataHolder(data); } + }; +#else + template + struct EigenConstMapMaker { + // Eigen/Core should be pre-included before the SoA headers to enable support for Eigen columns. + static_assert(!sizeof(C), + "Eigen/Core should be pre-included before the SoA headers to enable support for Eigen columns."); + }; +#endif + // Helper function to compute aligned size + inline size_t alignSize(size_t size, size_t alignment = 128) { + if (size) + return ((size - 1) / alignment + 1) * alignment; + else + return 0; + } + +} // namespace cms::soa + +#define SOA_SCALAR(TYPE, NAME) (_VALUE_TYPE_SCALAR, TYPE, NAME) +#define SOA_COLUMN(TYPE, NAME) (_VALUE_TYPE_COLUMN, TYPE, NAME) +#define SOA_EIGEN_COLUMN(TYPE, NAME) (_VALUE_TYPE_EIGEN_COLUMN, TYPE, NAME) + +/* Iterate on the macro MACRO and return the result as a comma separated list */ +#define _ITERATE_ON_ALL_COMMA(MACRO, DATA, ...) \ + BOOST_PP_TUPLE_ENUM(BOOST_PP_SEQ_TO_TUPLE(_ITERATE_ON_ALL(MACRO, DATA, __VA_ARGS__))) +/* Iterate MACRO on all elements */ +#define _ITERATE_ON_ALL(MACRO, DATA, ...) BOOST_PP_SEQ_FOR_EACH(MACRO, DATA, BOOST_PP_VARIADIC_TO_SEQ(__VA_ARGS__)) + +/* Switch on macros depending on scalar / column type */ +#define _SWITCH_ON_TYPE(VALUE_TYPE, IF_SCALAR, IF_COLUMN, IF_EIGEN_COLUMN) \ + BOOST_PP_IF( \ + BOOST_PP_EQUAL(VALUE_TYPE, _VALUE_TYPE_SCALAR), \ + IF_SCALAR, \ + BOOST_PP_IF( \ + BOOST_PP_EQUAL(VALUE_TYPE, _VALUE_TYPE_COLUMN), \ + IF_COLUMN, \ + BOOST_PP_IF(BOOST_PP_EQUAL(VALUE_TYPE, _VALUE_TYPE_EIGEN_COLUMN), IF_EIGEN_COLUMN, BOOST_PP_EMPTY()))) + +namespace cms::soa { + + /* Column accessors: templates implementing the global accesors (soa::x() and soa::x(index) */ + enum class SoAAccessType : bool { mutableAccess, constAccess }; + + template + struct SoAColumnAccessorsImpl {}; + + // Todo: add alignment support. + // Sfinae based const/non const variants. + // Column + template + struct SoAColumnAccessorsImpl { + //SOA_HOST_DEVICE_INLINE SoAColumnAccessorsImpl(T* baseAddress) : baseAddress_(baseAddress) {} + SOA_HOST_DEVICE_INLINE SoAColumnAccessorsImpl(const SoAParametersImpl& params) + : params_(params) {} + SOA_HOST_DEVICE_INLINE T* operator()() { return params_.addr_; } + typedef T* NoParamReturnType; + SOA_HOST_DEVICE_INLINE T& operator()(size_t index) { return params_.addr_[index]; } + + private: + SoAParametersImpl params_; + }; + + // Const column + template + struct SoAColumnAccessorsImpl { + SOA_HOST_DEVICE_INLINE + SoAColumnAccessorsImpl(const SoAConstParametersImpl& params) + : params_(params) {} + SOA_HOST_DEVICE_INLINE const T* operator()() const { return params_.addr_; } + typedef T* NoParamReturnType; + SOA_HOST_DEVICE_INLINE T operator()(size_t index) const { return params_.addr_[index]; } + + private: + SoAConstParametersImpl params_; + }; + + // Scalar + template + struct SoAColumnAccessorsImpl { + SOA_HOST_DEVICE_INLINE SoAColumnAccessorsImpl(const SoAParametersImpl& params) + : params_(params) {} + SOA_HOST_DEVICE_INLINE T& operator()() { return *params_.addr_; } + typedef T& NoParamReturnType; + SOA_HOST_DEVICE_INLINE void operator()(size_t index) const { + assert(false && "Indexed access impossible for SoA scalars."); + } + + private: + SoAParametersImpl params_; + }; + + // Const scalar + template + struct SoAColumnAccessorsImpl { + SOA_HOST_DEVICE_INLINE + SoAColumnAccessorsImpl(const SoAConstParametersImpl& params) + : params_(params) {} + SOA_HOST_DEVICE_INLINE T operator()() const { return *params_.addr_; } + typedef T NoParamReturnType; + SOA_HOST_DEVICE_INLINE void operator()(size_t index) const { + assert(false && "Indexed access impossible for SoA scalars."); + } + + private: + SoAConstParametersImpl params_; + }; + + template + struct SoAColumnAccessorsImpl { + //SOA_HOST_DEVICE_INLINE SoAColumnAccessorsImpl(T* baseAddress) : baseAddress_(baseAddress) {} + SOA_HOST_DEVICE_INLINE SoAColumnAccessorsImpl(const SoAParametersImpl& params) + : params_(params) {} + SOA_HOST_DEVICE_INLINE typename T::Scalar* operator()() { return params_.addr_; } + typedef typename T::Scalar* NoParamReturnType; + //SOA_HOST_DEVICE_INLINE T& operator()(size_t index) { return params_.addr_[index]; } + + private: + SoAParametersImpl params_; + }; + + // Const column + template + struct SoAColumnAccessorsImpl { + SOA_HOST_DEVICE_INLINE + SoAColumnAccessorsImpl(const SoAConstParametersImpl& params) + : params_(params) {} + SOA_HOST_DEVICE_INLINE const typename T::Scalar* operator()() const { return params_.addr_; } + typedef typename T::Scalar* NoParamReturnType; + //SOA_HOST_DEVICE_INLINE T operator()(size_t index) const { return params_.addr_[index]; } + + private: + SoAConstParametersImpl params_; + }; + + /* A helper template stager avoiding comma in macros */ + template + struct SoAAccessors { + template + struct ColumnType { + template + struct AccessType : public SoAColumnAccessorsImpl { + using SoAColumnAccessorsImpl::SoAColumnAccessorsImpl; + }; + }; + }; + + /* Enum parameters allowing templated control of layout/view behaviors */ + /* Alignement enforcement verifies every column is aligned, and + * hints the compiler that it can expect column pointers to be aligned */ + enum class AlignmentEnforcement : bool { Relaxed, Enforced }; + + struct CacheLineSize { + static constexpr size_t NvidiaGPU = 128; + static constexpr size_t IntelCPU = 64; + static constexpr size_t AMDCPU = 64; + static constexpr size_t ARMCPU = 64; + static constexpr size_t defaultSize = NvidiaGPU; + }; + + // An empty shell class to restrict the scope of tempalted operator<<(ostream, soa). + struct BaseLayout {}; +} // namespace cms::soa + +// Small wrapper for stream insertion of SoA printing +template ::value, SOA>::type> +SOA_HOST_ONLY std::ostream& operator<<(std::ostream& os, const SOA& soa) { + soa.toStream(os); + return os; +} +#endif // ndef DataStructures_SoACommon_h diff --git a/src/alpaka/DataFormats/SoALayout.h b/src/alpaka/DataFormats/SoALayout.h new file mode 100644 index 000000000..036f1cb4f --- /dev/null +++ b/src/alpaka/DataFormats/SoALayout.h @@ -0,0 +1,388 @@ +/* + * Structure-of-Arrays template with "columns" and "scalars", defined through preprocessor macros, + * with compile-time size and alignment, and accessors to the "rows" and "columns". + */ + +#ifndef DataStructures_SoALayout_h +#define DataStructures_SoALayout_h + +#include "SoACommon.h" + +#include +#include + +/* dump SoA fields information; these should expand to, for columns: + * Example: + * GENERATE_SOA_LAYOUT(SoA, + * // predefined static scalars + * // size_t size; + * // size_t alignment; + * + * // columns: one value per element + * SOA_COLUMN(double, x), + * SOA_COLUMN(double, y), + * SOA_COLUMN(double, z), + * SOA_EIGEN_COLUMN(Eigen::Vector3d, a), + * SOA_EIGEN_COLUMN(Eigen::Vector3d, b), + * SOA_EIGEN_COLUMN(Eigen::Vector3d, r), + * SOA_COLUMN(uint16_t, colour), + * SOA_COLUMN(int32_t, value), + * SOA_COLUMN(double *, py), + * SOA_COLUMN(uint32_t, count), + * SOA_COLUMN(uint32_t, anotherCount), + * + * // scalars: one value for the whole structure + * SOA_SCALAR(const char *, description), + * SOA_SCALAR(uint32_t, someNumber) + * ); + * + * dumps as: + * SoA(32, 64): + * sizeof(SoA): 152 + * Column x_ at offset 0 has size 256 and padding 0 + * Column y_ at offset 256 has size 256 and padding 0 + * Column z_ at offset 512 has size 256 and padding 0 + * Eigen value a_ at offset 768 has dimension (3 x 1) and per column size 256 and padding 0 + * Eigen value b_ at offset 1536 has dimension (3 x 1) and per column size 256 and padding 0 + * Eigen value r_ at offset 2304 has dimension (3 x 1) and per column size 256 and padding 0 + * Column colour_ at offset 3072 has size 64 and padding 0 + * Column value_ at offset 3136 has size 128 and padding 0 + * Column py_ at offset 3264 has size 256 and padding 0 + * Column count_ at offset 3520 has size 128 and padding 0 + * Column anotherCount_ at offset 3648 has size 128 and padding 0 + * Scalar description_ at offset 3776 has size 8 and padding 56 + * Scalar someNumber_ at offset 3840 has size 4 and padding 60 + * Final offset = 3904 computeDataSize(...): 3904 + * + */ + +// clang-format off +#define _DECLARE_SOA_STREAM_INFO_IMPL(VALUE_TYPE, CPP_TYPE, NAME) \ + _SWITCH_ON_TYPE( \ + VALUE_TYPE, \ + /* Dump scalar */ \ + os << " Scalar " BOOST_PP_STRINGIZE(NAME) " at offset " << offset << " has size " << sizeof(CPP_TYPE) \ + << " and padding " << ((sizeof(CPP_TYPE) - 1) / byteAlignment + 1) * byteAlignment - sizeof(CPP_TYPE) \ + << std::endl; \ + offset += ((sizeof(CPP_TYPE) - 1) / byteAlignment + 1) * byteAlignment; \ + , /* Dump column */ \ + os << " Column " BOOST_PP_STRINGIZE(NAME) " at offset " << offset << " has size " << sizeof(CPP_TYPE) * nElements_ \ + << " and padding " \ + << (((nElements_ * sizeof(CPP_TYPE) - 1) / byteAlignment) + 1) * byteAlignment - (sizeof(CPP_TYPE) * nElements_) \ + << std::endl; \ + offset += (((nElements_ * sizeof(CPP_TYPE) - 1) / byteAlignment) + 1) * byteAlignment; \ + , /* Dump Eigen column */ \ + os << " Eigen value " BOOST_PP_STRINGIZE(NAME) " at offset " << offset << " has dimension (" \ + << CPP_TYPE::RowsAtCompileTime << " x " << CPP_TYPE::ColsAtCompileTime \ + << ")" \ + << " and per column size " \ + << sizeof(CPP_TYPE::Scalar) * nElements_ \ + << " and padding " \ + << (((nElements_ * sizeof(CPP_TYPE::Scalar) - 1) / byteAlignment) + 1) * byteAlignment - \ + (sizeof(CPP_TYPE::Scalar) * nElements_) \ + << std::endl; \ + offset += (((nElements_ * sizeof(CPP_TYPE::Scalar) - 1) / byteAlignment) + 1) * byteAlignment * \ + CPP_TYPE::RowsAtCompileTime * CPP_TYPE::ColsAtCompileTime;) +// clang-format on + +#define _DECLARE_SOA_STREAM_INFO(R, DATA, TYPE_NAME) BOOST_PP_EXPAND(_DECLARE_SOA_STREAM_INFO_IMPL TYPE_NAME) + +/** + * SoAMetadata member computing column pitch + */ +// clang-format off +#define _DEFINE_METADATA_MEMBERS_IMPL(VALUE_TYPE, CPP_TYPE, NAME) \ + _SWITCH_ON_TYPE(VALUE_TYPE, \ + /* Scalar */ \ + size_t BOOST_PP_CAT(NAME, Pitch()) const { \ + return (((sizeof(CPP_TYPE) - 1) / ParentClass::byteAlignment) + 1) * ParentClass::byteAlignment; \ + } \ + typedef CPP_TYPE BOOST_PP_CAT(TypeOf_, NAME); \ + constexpr static cms::soa::SoAColumnType BOOST_PP_CAT(ColumnTypeOf_, NAME) = cms::soa::SoAColumnType::scalar; \ + SOA_HOST_DEVICE_INLINE \ + CPP_TYPE const* BOOST_PP_CAT(addressOf_, NAME)() const { \ + return parent_.soaMetadata().BOOST_PP_CAT(parametersOf_, NAME)().addr_; \ + } \ + typedef cms::soa::SoAParameters_ColumnType::DataType \ + BOOST_PP_CAT(ParametersTypeOf_, NAME); \ + SOA_HOST_DEVICE_INLINE \ + BOOST_PP_CAT(ParametersTypeOf_, NAME) BOOST_PP_CAT(parametersOf_, NAME)() const { \ + return BOOST_PP_CAT(ParametersTypeOf_, NAME) (parent_.BOOST_PP_CAT(NAME, _)); \ + } \ + SOA_HOST_DEVICE_INLINE \ + CPP_TYPE* BOOST_PP_CAT(addressOf_, NAME)() { \ + return parent_.soaMetadata().BOOST_PP_CAT(parametersOf_, NAME)().addr_; \ + }, \ + /* Column */ \ + typedef cms::soa::SoAParameters_ColumnType::DataType \ + BOOST_PP_CAT(ParametersTypeOf_, NAME); \ + SOA_HOST_DEVICE_INLINE \ + BOOST_PP_CAT(ParametersTypeOf_, NAME) BOOST_PP_CAT(parametersOf_, NAME)() const { \ + return BOOST_PP_CAT(ParametersTypeOf_, NAME) (parent_.BOOST_PP_CAT(NAME, _)); \ + } \ + SOA_HOST_DEVICE_INLINE \ + CPP_TYPE const* BOOST_PP_CAT(addressOf_, NAME)() const { \ + return parent_.soaMetadata().BOOST_PP_CAT(parametersOf_, NAME)().addr_; \ + } \ + SOA_HOST_DEVICE_INLINE \ + CPP_TYPE* BOOST_PP_CAT(addressOf_, NAME)() { \ + return parent_.soaMetadata().BOOST_PP_CAT(parametersOf_, NAME)().addr_; \ + } \ + SOA_HOST_DEVICE_INLINE \ + size_t BOOST_PP_CAT(NAME, Pitch()) const { \ + return (((parent_.nElements_ * sizeof(CPP_TYPE) - 1) / ParentClass::byteAlignment) + 1) * \ + ParentClass::byteAlignment; \ + } \ + typedef CPP_TYPE BOOST_PP_CAT(TypeOf_, NAME); \ + constexpr static cms::soa::SoAColumnType BOOST_PP_CAT(ColumnTypeOf_, NAME) = cms::soa::SoAColumnType::column;, \ + /* Eigen column */ \ + typedef cms::soa::SoAParameters_ColumnType::DataType \ + BOOST_PP_CAT(ParametersTypeOf_, NAME); \ + SOA_HOST_DEVICE_INLINE \ + BOOST_PP_CAT(ParametersTypeOf_, NAME) BOOST_PP_CAT(parametersOf_, NAME)() const { \ + return BOOST_PP_CAT(ParametersTypeOf_, NAME) ( \ + parent_.BOOST_PP_CAT(NAME, _), \ + parent_.BOOST_PP_CAT(NAME, Stride_)); \ + } \ + SOA_HOST_DEVICE_INLINE \ + size_t BOOST_PP_CAT(NAME, Pitch()) const { \ + return (((parent_.nElements_ * sizeof(CPP_TYPE::Scalar) - 1) / ParentClass::byteAlignment) + 1) * \ + ParentClass::byteAlignment * CPP_TYPE::RowsAtCompileTime * CPP_TYPE::ColsAtCompileTime; \ + } typedef CPP_TYPE BOOST_PP_CAT(TypeOf_, NAME); \ + constexpr static cms::soa::SoAColumnType BOOST_PP_CAT(ColumnTypeOf_, NAME) = cms::soa::SoAColumnType::eigen; \ + SOA_HOST_DEVICE_INLINE \ + CPP_TYPE::Scalar const* BOOST_PP_CAT(addressOf_, NAME)() const { \ + return parent_.soaMetadata().BOOST_PP_CAT(parametersOf_, NAME)().addr_; \ + } \ + SOA_HOST_DEVICE_INLINE \ + CPP_TYPE::Scalar* BOOST_PP_CAT(addressOf_, NAME)() { \ + return parent_.soaMetadata().BOOST_PP_CAT(parametersOf_, NAME)().addr_; \ + } \ +) +// clang-format on +#define _DEFINE_METADATA_MEMBERS(R, DATA, TYPE_NAME) _DEFINE_METADATA_MEMBERS_IMPL TYPE_NAME + +/** + * Member assignment for trivial constructor + */ +#define _DECLARE_MEMBER_TRIVIAL_CONSTRUCTION_IMPL(VALUE_TYPE, CPP_TYPE, NAME) \ + _SWITCH_ON_TYPE(VALUE_TYPE, /* Scalar */ \ + (BOOST_PP_CAT(NAME, _)(nullptr)), /* Column */ \ + (BOOST_PP_CAT(NAME, _)(nullptr)), /* Eigen column */ \ + (BOOST_PP_CAT(NAME, _)(nullptr))(BOOST_PP_CAT(NAME, Stride_)(0))) + +#define _DECLARE_MEMBER_TRIVIAL_CONSTRUCTION(R, DATA, TYPE_NAME) \ + BOOST_PP_EXPAND(_DECLARE_MEMBER_TRIVIAL_CONSTRUCTION_IMPL TYPE_NAME) +/** + * Computation of the column or scalar pointer location in the memory layout (at SoA construction time) + */ +#define _ASSIGN_SOA_COLUMN_OR_SCALAR_IMPL(VALUE_TYPE, CPP_TYPE, NAME) \ + _SWITCH_ON_TYPE(VALUE_TYPE, /* Scalar */ \ + BOOST_PP_CAT(NAME, _) = reinterpret_cast(curMem); \ + curMem += (((sizeof(CPP_TYPE) - 1) / byteAlignment) + 1) * byteAlignment; \ + , /* Column */ \ + BOOST_PP_CAT(NAME, _) = reinterpret_cast(curMem); \ + curMem += (((nElements_ * sizeof(CPP_TYPE) - 1) / byteAlignment) + 1) * byteAlignment; \ + , /* Eigen column */ \ + BOOST_PP_CAT(NAME, _) = reinterpret_cast(curMem); \ + curMem += (((nElements_ * sizeof(CPP_TYPE::Scalar) - 1) / byteAlignment) + 1) * byteAlignment * \ + CPP_TYPE::RowsAtCompileTime * CPP_TYPE::ColsAtCompileTime; \ + BOOST_PP_CAT(NAME, Stride_) = (((nElements_ * sizeof(CPP_TYPE::Scalar) - 1) / byteAlignment) + 1) * \ + byteAlignment / sizeof(CPP_TYPE::Scalar);) \ + if constexpr (alignmentEnforcement == AlignmentEnforcement::Enforced) \ + if (reinterpret_cast(BOOST_PP_CAT(NAME, _)) % byteAlignment) \ + throw std::out_of_range("In layout constructor: misaligned column: " #NAME); + +#define _ASSIGN_SOA_COLUMN_OR_SCALAR(R, DATA, TYPE_NAME) _ASSIGN_SOA_COLUMN_OR_SCALAR_IMPL TYPE_NAME + +/** + * Computation of the column or scalar size for SoA size computation + */ +#define _ACCUMULATE_SOA_ELEMENT_IMPL(VALUE_TYPE, CPP_TYPE, NAME) \ + _SWITCH_ON_TYPE(VALUE_TYPE, /* Scalar */ \ + ret += (((sizeof(CPP_TYPE) - 1) / byteAlignment) + 1) * byteAlignment; \ + , /* Column */ \ + ret += (((nElements * sizeof(CPP_TYPE) - 1) / byteAlignment) + 1) * byteAlignment; \ + , /* Eigen column */ \ + ret += (((nElements * sizeof(CPP_TYPE::Scalar) - 1) / byteAlignment) + 1) * byteAlignment * \ + CPP_TYPE::RowsAtCompileTime * CPP_TYPE::ColsAtCompileTime;) + +#define _ACCUMULATE_SOA_ELEMENT(R, DATA, TYPE_NAME) _ACCUMULATE_SOA_ELEMENT_IMPL TYPE_NAME + +/** + * Direct access to column pointer and indexed access + */ +#define _DECLARE_SOA_ACCESSOR_IMPL(VALUE_TYPE, CPP_TYPE, NAME) \ + _SWITCH_ON_TYPE( \ + VALUE_TYPE, /* Scalar */ \ + SOA_HOST_DEVICE_INLINE CPP_TYPE& NAME() { return *BOOST_PP_CAT(NAME, _); }, /* Column */ \ + SOA_HOST_DEVICE_INLINE CPP_TYPE* NAME() { \ + return BOOST_PP_CAT(NAME, _); \ + } SOA_HOST_DEVICE_INLINE CPP_TYPE& NAME(size_t index) { return BOOST_PP_CAT(NAME, _)[index]; }, \ + /* Eigen column */ /* Unsupported for the moment TODO */ \ + BOOST_PP_EMPTY()) + +#define _DECLARE_SOA_ACCESSOR(R, DATA, TYPE_NAME) BOOST_PP_EXPAND(_DECLARE_SOA_ACCESSOR_IMPL TYPE_NAME) + +/** + * Direct access to column pointer (const) and indexed access. + */ +#define _DECLARE_SOA_CONST_ACCESSOR_IMPL(VALUE_TYPE, CPP_TYPE, NAME) \ + _SWITCH_ON_TYPE( \ + VALUE_TYPE, /* Scalar */ \ + SOA_HOST_DEVICE_INLINE CPP_TYPE NAME() const { return *(BOOST_PP_CAT(NAME, _)); }, /* Column */ \ + SOA_HOST_DEVICE_INLINE CPP_TYPE const* NAME() \ + const { return BOOST_PP_CAT(NAME, _); } SOA_HOST_DEVICE_INLINE CPP_TYPE NAME(size_t index) \ + const { return *(BOOST_PP_CAT(NAME, _) + index); }, /* Eigen column */ \ + SOA_HOST_DEVICE_INLINE CPP_TYPE::Scalar const* NAME() \ + const { return BOOST_PP_CAT(NAME, _); } SOA_HOST_DEVICE_INLINE size_t BOOST_PP_CAT( \ + NAME, Stride)() { return BOOST_PP_CAT(NAME, Stride_); }) + +#define _DECLARE_SOA_CONST_ACCESSOR(R, DATA, TYPE_NAME) BOOST_PP_EXPAND(_DECLARE_SOA_CONST_ACCESSOR_IMPL TYPE_NAME) + +/** + * SoA class member declaration (column pointers). + */ +#define _DECLARE_SOA_DATA_MEMBER_IMPL(VALUE_TYPE, CPP_TYPE, NAME) \ + _SWITCH_ON_TYPE(VALUE_TYPE, /* Scalar */ \ + CPP_TYPE* BOOST_PP_CAT(NAME, _) = nullptr; \ + , /* Column */ \ + CPP_TYPE * BOOST_PP_CAT(NAME, _) = nullptr; \ + , /* Eigen column */ \ + CPP_TYPE::Scalar * BOOST_PP_CAT(NAME, _) = nullptr; \ + size_t BOOST_PP_CAT(NAME, Stride_) = 0;) + +#define _DECLARE_SOA_DATA_MEMBER(R, DATA, TYPE_NAME) BOOST_PP_EXPAND(_DECLARE_SOA_DATA_MEMBER_IMPL TYPE_NAME) + +#ifdef DEBUG +#define _DO_RANGECHECK true +#else +#define _DO_RANGECHECK false +#endif + +/* + * A macro defining a SoA layout (collection of scalars and columns of equal lengths) + */ +// clang-format off +#define GENERATE_SOA_LAYOUT(CLASS, ...) \ + template \ + struct CLASS: public cms::soa::BaseLayout { \ + /* these could be moved to an external type trait to free up the symbol names */ \ + using self_type = CLASS; \ + typedef cms::soa::AlignmentEnforcement AlignmentEnforcement; \ + \ + /* For CUDA applications, we align to the 128 bytes of the cache lines. \ + * See https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#global-memory-3-0 this is still valid \ + * up to compute capability 8.X. \ + */ \ + constexpr static size_t defaultAlignment = 128; \ + constexpr static size_t byteAlignment = ALIGNMENT; \ + constexpr static AlignmentEnforcement alignmentEnforcement = ALIGNMENT_ENFORCEMENT; \ + constexpr static size_t conditionalAlignment = \ + alignmentEnforcement == AlignmentEnforcement::Enforced ? byteAlignment : 0; \ + /* Those typedefs avoid having commas in macros (which is problematic) */ \ + template \ + using SoAValueWithConf = cms::soa::SoAValue; \ + \ + template \ + using SoAConstValueWithConf = cms::soa::SoAConstValue; \ + \ + /* dump the SoA internal structure */ \ + SOA_HOST_ONLY \ + void toStream(std::ostream & os) const { \ + os << #CLASS "(" << nElements_ << " elements, byte alignement= " << byteAlignment << ", @"<< mem_ <<"): " << std::endl; \ + os << " sizeof(" #CLASS "): " << sizeof(CLASS) << std::endl; \ + size_t offset = 0; \ + _ITERATE_ON_ALL(_DECLARE_SOA_STREAM_INFO, ~, __VA_ARGS__) \ + os << "Final offset = " << offset << " computeDataSize(...): " << computeDataSize(nElements_) \ + << std::endl; \ + os << std::endl; \ + } \ + \ + /* Helper function used by caller to externally allocate the storage */ \ + static size_t computeDataSize(size_t nElements) { \ + size_t ret = 0; \ + _ITERATE_ON_ALL(_ACCUMULATE_SOA_ELEMENT, ~, __VA_ARGS__) \ + return ret; \ + } \ + \ + /** \ + * Helper/friend class allowing SoA introspection. \ + */ \ + struct SoAMetadata { \ + friend CLASS; \ + SOA_HOST_DEVICE_INLINE size_t size() const { return parent_.nElements_; } \ + SOA_HOST_DEVICE_INLINE size_t byteSize() const { return parent_.byteSize_; } \ + SOA_HOST_DEVICE_INLINE size_t byteAlignment() const { return CLASS::byteAlignment; } \ + SOA_HOST_DEVICE_INLINE std::byte* data() { return parent_.mem_; } \ + SOA_HOST_DEVICE_INLINE const std::byte* data() const { return parent_.mem_; } \ + SOA_HOST_DEVICE_INLINE std::byte* nextByte() const { return parent_.mem_ + parent_.byteSize_; } \ + SOA_HOST_DEVICE_INLINE CLASS cloneToNewAddress(std::byte* addr) const { \ + return CLASS(addr, parent_.nElements_); \ + } \ + _ITERATE_ON_ALL(_DEFINE_METADATA_MEMBERS, ~, __VA_ARGS__) \ + \ + SoAMetadata& operator=(const SoAMetadata&) = delete; \ + SoAMetadata(const SoAMetadata&) = delete; \ + \ + private: \ + SOA_HOST_DEVICE_INLINE SoAMetadata(const CLASS& parent) : parent_(parent) {} \ + const CLASS& parent_; \ + typedef CLASS ParentClass; \ + }; \ + friend SoAMetadata; \ + SOA_HOST_DEVICE_INLINE const SoAMetadata soaMetadata() const { return SoAMetadata(*this); } \ + SOA_HOST_DEVICE_INLINE SoAMetadata soaMetadata() { return SoAMetadata(*this); } \ + \ + /* Trivial constuctor */ \ + CLASS() \ + : mem_(nullptr), \ + nElements_(0), \ + byteSize_(0), \ + _ITERATE_ON_ALL_COMMA(_DECLARE_MEMBER_TRIVIAL_CONSTRUCTION, ~, __VA_ARGS__) {} \ + \ + /* Constructor relying on user provided storage */ \ + SOA_HOST_ONLY CLASS(std::byte* mem, size_t nElements) : mem_(mem), nElements_(nElements), byteSize_(0) { \ + if constexpr (alignmentEnforcement == AlignmentEnforcement::Enforced) \ + if (reinterpret_cast(mem) % byteAlignment) \ + throw std::out_of_range("In " #CLASS "::" #CLASS ": misaligned buffer"); \ + auto curMem = mem_; \ + _ITERATE_ON_ALL(_ASSIGN_SOA_COLUMN_OR_SCALAR, ~, __VA_ARGS__) \ + /* Sanity check: we should have reached the computed size, only on host code */ \ + byteSize_ = computeDataSize(nElements_); \ + if (mem_ + byteSize_ != curMem) \ + throw std::out_of_range("In " #CLASS "::" #CLASS ": unexpected end pointer."); \ + } \ + \ + /* Constructor relying on user provided storage */ \ + SOA_DEVICE_ONLY CLASS(bool devConstructor, std::byte* mem, size_t nElements) : mem_(mem), nElements_(nElements) { \ + auto curMem = mem_; \ + _ITERATE_ON_ALL(_ASSIGN_SOA_COLUMN_OR_SCALAR, ~, __VA_ARGS__) \ + } \ + \ + /* dump the SoA internal structure */ \ + template \ + SOA_HOST_ONLY friend void dump(); \ + \ + private: \ + /* Range checker conditional to the macro _DO_RANGECHECK */ \ + SOA_HOST_DEVICE_INLINE \ + void rangeCheck(size_t index) const { \ + if constexpr (_DO_RANGECHECK) { \ + if (index >= nElements_) { \ + printf("In " #CLASS "::rangeCheck(): index out of range: %zu with nElements: %zu\n", index, nElements_); \ + assert(false); \ + } \ + } \ + } \ + \ + /* data members */ \ + std::byte* mem_; \ + size_t nElements_; \ + size_t byteSize_; \ + _ITERATE_ON_ALL(_DECLARE_SOA_DATA_MEMBER, ~, __VA_ARGS__) \ + }; +// clang-format on + +#endif // ndef DataStructures_SoALayout_h diff --git a/src/alpaka/DataFormats/SoAView.h b/src/alpaka/DataFormats/SoAView.h new file mode 100644 index 000000000..5b691c579 --- /dev/null +++ b/src/alpaka/DataFormats/SoAView.h @@ -0,0 +1,599 @@ +/* + * Structure-of-Arrays templates allowing access to a selection of scalars and columns from one + * or multiple SoA layouts or views. + * This template generator will allow handling subsets of columns from one or multiple SoA views or layouts. + */ + +#ifndef DataStructures_SoAView_h +#define DataStructures_SoAView_h + +#include "SoACommon.h" + +#define SOA_VIEW_LAYOUT(TYPE, NAME) (TYPE, NAME) + +#define SOA_VIEW_LAYOUT_LIST(...) __VA_ARGS__ + +#define SOA_VIEW_VALUE(LAYOUT_NAME, LAYOUT_MEMBER) (LAYOUT_NAME, LAYOUT_MEMBER, LAYOUT_MEMBER) + +#define SOA_VIEW_VALUE_RENAME(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) (LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) + +#define SOA_VIEW_VALUE_LIST(...) __VA_ARGS__ + +/* + * A macro defining a SoA view (collection of columns from multiple layouts or views.) + * + * Usage: + * GENERATE_SOA_VIEW(PixelXYView, + * SOA_VIEW_LAYOUT_LIST( + * SOA_VIEW_LAYOUT(PixelDigis, pixelDigis), + * SOA_VIEW_LAYOUT(PixelRecHitsLayout, pixelsRecHit) + * ), + * SOA_VIEW_VALUE_LIST( + * SOA_VIEW_VALUE_RENAME(pixelDigis, x, digisX), + * SOA_VIEW_VALUE_RENAME(pixelDigis, y, digisY), + * SOA_VIEW_VALUE_RENAME(pixelsRecHit, x, recHitsX), + * SOA_VIEW_VALUE_RENAME(pixelsRecHit, y, recHitsY) + * ) + * ); + * + */ + +namespace cms::soa { + + /* Traits for the different column type scenarios */ + /* Value traits passes the class as is in the case of column type and return + * an empty class with functions returning non-scalar as accessors. */ + template + struct ConstValueTraits : public C { + using C::C; + }; + + template + struct ConstValueTraits { + // Just take to SoAValue type to generate the right constructor. + SOA_HOST_DEVICE_INLINE ConstValueTraits(size_t, const typename C::valueType*) {} + SOA_HOST_DEVICE_INLINE ConstValueTraits(size_t, const typename C::Params&) {} + SOA_HOST_DEVICE_INLINE ConstValueTraits(size_t, const typename C::ConstParams&) {} + // Any attempt to do anything with the "scalar" value a const element will fail. + }; + +} // namespace cms::soa + +#include +/* + * Members definitions macros for viewa + */ + +/** + * Layout types aliasing for referencing by name + */ +#define _DECLARE_VIEW_LAYOUT_TYPE_ALIAS_IMPL(TYPE, NAME) typedef TYPE BOOST_PP_CAT(TypeOf_, NAME); + +#define _DECLARE_VIEW_LAYOUT_TYPE_ALIAS(R, DATA, TYPE_NAME) \ + BOOST_PP_EXPAND(_DECLARE_VIEW_LAYOUT_TYPE_ALIAS_IMPL TYPE_NAME) + +/** + * Member types aliasing for referencing by name + */ +#define _DECLARE_VIEW_MEMBER_TYPE_ALIAS_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME, DATA) \ + typedef typename BOOST_PP_CAT(TypeOf_, LAYOUT_NAME)::SoAMetadata::BOOST_PP_CAT(TypeOf_, LAYOUT_MEMBER) \ + BOOST_PP_CAT(TypeOf_, LOCAL_NAME); \ + typedef typename BOOST_PP_CAT(TypeOf_, LAYOUT_NAME)::SoAMetadata::BOOST_PP_CAT(ParametersTypeOf_, LAYOUT_MEMBER) \ + BOOST_PP_CAT(ParametersTypeOf_, LOCAL_NAME); \ + constexpr static cms::soa::SoAColumnType BOOST_PP_CAT(ColumnTypeOf_, LOCAL_NAME) = \ + BOOST_PP_CAT(TypeOf_, LAYOUT_NAME)::SoAMetadata::BOOST_PP_CAT(ColumnTypeOf_, LAYOUT_MEMBER); \ + SOA_HOST_DEVICE_INLINE DATA auto* BOOST_PP_CAT(addressOf_, LOCAL_NAME)() const { \ + return parent_.soaMetadata().BOOST_PP_CAT(parametersOf_, LOCAL_NAME)().addr_; \ + }; \ + SOA_HOST_DEVICE_INLINE \ + DATA BOOST_PP_CAT(ParametersTypeOf_, LOCAL_NAME) BOOST_PP_CAT(parametersOf_, LOCAL_NAME)() const { \ + return parent_.BOOST_PP_CAT(LOCAL_NAME, Parameters_); \ + }; + +#define _DECLARE_VIEW_MEMBER_TYPE_ALIAS(R, DATA, LAYOUT_MEMBER_NAME) \ + BOOST_PP_EXPAND(_DECLARE_VIEW_MEMBER_TYPE_ALIAS_IMPL BOOST_PP_TUPLE_PUSH_BACK(LAYOUT_MEMBER_NAME, DATA)) + +/** + * Generator of parameters (layouts/views) for constructor by layouts/views. + */ +#define _DECLARE_VIEW_CONSTRUCTION_PARAMETERS_IMPL(LAYOUT_TYPE, LAYOUT_NAME, DATA) (DATA LAYOUT_TYPE & LAYOUT_NAME) + +#define _DECLARE_VIEW_CONSTRUCTION_PARAMETERS(R, DATA, TYPE_NAME) \ + BOOST_PP_EXPAND(_DECLARE_VIEW_CONSTRUCTION_PARAMETERS_IMPL BOOST_PP_TUPLE_PUSH_BACK(TYPE_NAME, DATA)) + +/** + * Generator of parameters for constructor by column. + */ +#define _DECLARE_VIEW_CONSTRUCTION_BYCOLUMN_PARAMETERS_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME, DATA) \ + (DATA typename BOOST_PP_CAT(SoAMetadata::ParametersTypeOf_, LOCAL_NAME)::TupleOrPointerType LOCAL_NAME) + +#define _DECLARE_VIEW_CONSTRUCTION_BYCOLUMN_PARAMETERS(R, DATA, LAYOUT_MEMBER_NAME) \ + BOOST_PP_EXPAND( \ + _DECLARE_VIEW_CONSTRUCTION_BYCOLUMN_PARAMETERS_IMPL BOOST_PP_TUPLE_PUSH_BACK(LAYOUT_MEMBER_NAME, DATA)) + +/** + * Generator of member initialization from constructor. + * We use a lambda with auto return type to handle multiple possible return types. + */ +#define _DECLARE_VIEW_MEMBER_INITIALIZERS_IMPL(LAYOUT, MEMBER, NAME) \ + (BOOST_PP_CAT(NAME, Parameters_)([&]() -> auto { \ + auto params = LAYOUT.soaMetadata().BOOST_PP_CAT(parametersOf_, MEMBER)(); \ + if constexpr (alignmentEnforcement == AlignmentEnforcement::Enforced) \ + if (reinterpret_cast(params.addr_) % byteAlignment) \ + throw std::out_of_range("In constructor by layout: misaligned column: " #NAME); \ + return params; \ + }())) + +#define _DECLARE_VIEW_MEMBER_INITIALIZERS(R, DATA, LAYOUT_MEMBER_NAME) \ + BOOST_PP_EXPAND(_DECLARE_VIEW_MEMBER_INITIALIZERS_IMPL LAYOUT_MEMBER_NAME) + +/** + * Generator of size computation for constructor. + * This is the per-layout part of the lambda checking they all have the same size. + */ +#define _UPDATE_SIZE_OF_VIEW_IMPL(LAYOUT_TYPE, LAYOUT_NAME) \ + if (set) { \ + if (ret != LAYOUT_NAME.soaMetadata().size()) \ + throw std::out_of_range("In constructor by layout: different sizes from layouts."); \ + } else { \ + ret = LAYOUT_NAME.soaMetadata().size(); \ + set = true; \ + } + +#define _UPDATE_SIZE_OF_VIEW(R, DATA, TYPE_NAME) BOOST_PP_EXPAND(_UPDATE_SIZE_OF_VIEW_IMPL TYPE_NAME) + +/** + * Generator of member initialization from constructor. + * We use a lambda with auto return type to handle multiple possible return types. + */ +// clang-format off +#define _DECLARE_VIEW_MEMBER_INITIALIZERS_BYCOLUMN_IMPL(LAYOUT, MEMBER, NAME) \ + ( \ + BOOST_PP_CAT(NAME, Parameters_)([&]() -> auto { \ + if constexpr (alignmentEnforcement == AlignmentEnforcement::Enforced) \ + if (SoAMetadata:: BOOST_PP_CAT(ParametersTypeOf_, NAME)::checkAlignment(NAME, byteAlignment)) \ + throw std::out_of_range("In constructor by column: misaligned column: " #NAME); \ + return NAME; \ + }()) \ + ) +// clang-format on + +#define _DECLARE_VIEW_MEMBER_INITIALIZERS_BYCOLUMN(R, DATA, LAYOUT_MEMBER_NAME) \ + BOOST_PP_EXPAND(_DECLARE_VIEW_MEMBER_INITIALIZERS_BYCOLUMN_IMPL LAYOUT_MEMBER_NAME) + +/** + * Generator of element members initializer. + */ +#define _DECLARE_VIEW_ELEM_MEMBER_INIT_IMPL(LAYOUT, MEMBER, LOCAL_NAME, DATA) (LOCAL_NAME(DATA, LOCAL_NAME)) + +#define _DECLARE_VIEW_ELEM_MEMBER_INIT(R, DATA, LAYOUT_MEMBER_NAME) \ + BOOST_PP_EXPAND(_DECLARE_VIEW_ELEM_MEMBER_INIT_IMPL BOOST_PP_TUPLE_PUSH_BACK(LAYOUT_MEMBER_NAME, DATA)) + +/** + * Helper macro extracting the data type from metadata of a layout or view + */ +#define _COLUMN_TYPE(LAYOUT_NAME, LAYOUT_MEMBER) \ + typename std::remove_pointer::type + +/** + * Generator of parameters for (non-const) element subclass (expanded comma separated). + */ +#define _DECLARE_VIEW_ELEMENT_VALUE_ARG_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME, DATA) \ + (DATA typename BOOST_PP_CAT(SoAMetadata::ParametersTypeOf_, LOCAL_NAME) LOCAL_NAME) + +#define _DECLARE_VIEW_ELEMENT_VALUE_ARG(R, DATA, LAYOUT_MEMBER_NAME) \ + _DECLARE_VIEW_ELEMENT_VALUE_ARG_IMPL BOOST_PP_TUPLE_PUSH_BACK(LAYOUT_MEMBER_NAME, DATA) + +/** + * Generator of parameters for (const) element subclass (expanded comma separated). + */ +#define _DECLARE_CONST_VIEW_ELEMENT_VALUE_ARG_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME, DATA) \ + (DATA typename BOOST_PP_CAT(SoAMetadata::ParametersTypeOf_, LOCAL_NAME)::ConstType LOCAL_NAME) + +#define _DECLARE_CONST_VIEW_ELEMENT_VALUE_ARG(R, DATA, LAYOUT_MEMBER_NAME) \ + _DECLARE_CONST_VIEW_ELEMENT_VALUE_ARG_IMPL BOOST_PP_TUPLE_PUSH_BACK(LAYOUT_MEMBER_NAME, DATA) + +/** + * Generator of member initialization for constructor of element subclass + */ +#define _DECLARE_VIEW_CONST_ELEM_MEMBER_INIT_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME, DATA) \ + (BOOST_PP_CAT(LOCAL_NAME, _)(DATA, LOCAL_NAME)) + +/* declare AoS-like element value args for contructor; these should expand,for columns only */ +#define _DECLARE_VIEW_CONST_ELEM_MEMBER_INIT(R, DATA, LAYOUT_MEMBER_NAME) \ + BOOST_PP_EXPAND(_DECLARE_VIEW_CONST_ELEM_MEMBER_INIT_IMPL BOOST_PP_TUPLE_PUSH_BACK(LAYOUT_MEMBER_NAME, DATA)) + +/** + * Declaration of the members accessors of the const element subclass + */ +#define _DECLARE_VIEW_CONST_ELEMENT_ACCESSOR_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) \ + SOA_HOST_DEVICE_INLINE \ + typename SoAConstValueWithConf::RefToConst \ + LOCAL_NAME() const { \ + return BOOST_PP_CAT(LOCAL_NAME, _)(); \ + } + +#define _DECLARE_VIEW_CONST_ELEMENT_ACCESSOR(R, DATA, LAYOUT_MEMBER_NAME) \ + _DECLARE_VIEW_CONST_ELEMENT_ACCESSOR_IMPL LAYOUT_MEMBER_NAME + +/** + * Declaration of the private members of the const element subclass + */ +#define _DECLARE_VIEW_CONST_ELEMENT_VALUE_MEMBER_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) \ + const cms::soa::ConstValueTraits, \ + BOOST_PP_CAT(SoAMetadata::ColumnTypeOf_, LOCAL_NAME)> \ + BOOST_PP_CAT(LOCAL_NAME, _); + +#define _DECLARE_VIEW_CONST_ELEMENT_VALUE_MEMBER(R, DATA, LAYOUT_MEMBER_NAME) \ + _DECLARE_VIEW_CONST_ELEMENT_VALUE_MEMBER_IMPL LAYOUT_MEMBER_NAME + +/** + * Generator of the member-by-member copy operator of the element subclass. + */ +#define _DECLARE_VIEW_ELEMENT_VALUE_COPY_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) \ + if constexpr (SoAMetadata::BOOST_PP_CAT(ColumnTypeOf_, LOCAL_NAME) != cms::soa::SoAColumnType::scalar) \ + LOCAL_NAME() = other.LOCAL_NAME(); + +#define _DECLARE_VIEW_ELEMENT_VALUE_COPY(R, DATA, LAYOUT_MEMBER_NAME) \ + BOOST_PP_EXPAND(_DECLARE_VIEW_ELEMENT_VALUE_COPY_IMPL LAYOUT_MEMBER_NAME) + +/** + * Declaration of the private members of the const element subclass + */ +#define _DECLARE_VIEW_ELEMENT_VALUE_MEMBER_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) \ + SoAValueWithConf \ + LOCAL_NAME; + +#define _DECLARE_VIEW_ELEMENT_VALUE_MEMBER(R, DATA, LAYOUT_MEMBER_NAME) \ + _DECLARE_VIEW_ELEMENT_VALUE_MEMBER_IMPL LAYOUT_MEMBER_NAME + +/** + * Parameters passed to element subclass constructor in operator[] + */ +#define _DECLARE_VIEW_ELEMENT_CONSTR_CALL_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) \ + (BOOST_PP_CAT(LOCAL_NAME, Parameters_)) + +#define _DECLARE_VIEW_ELEMENT_CONSTR_CALL(R, DATA, LAYOUT_MEMBER_NAME) \ + BOOST_PP_EXPAND(_DECLARE_VIEW_ELEMENT_CONSTR_CALL_IMPL LAYOUT_MEMBER_NAME) + +/** + * Direct access to column pointer and indexed access + */ +#define _DECLARE_VIEW_SOA_ACCESSOR_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) \ + /* Column or scalar */ \ + SOA_HOST_DEVICE_INLINE \ + typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::mutableAccess>::NoParamReturnType \ + LOCAL_NAME() { \ + return typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::mutableAccess>(BOOST_PP_CAT(LOCAL_NAME, Parameters_))(); \ + } \ + SOA_HOST_DEVICE_INLINE auto& LOCAL_NAME(size_t index) { \ + return typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::mutableAccess>(BOOST_PP_CAT(LOCAL_NAME, Parameters_))(index); \ + } + +#define _DECLARE_VIEW_SOA_ACCESSOR(R, DATA, LAYOUT_MEMBER_NAME) \ + BOOST_PP_EXPAND(_DECLARE_VIEW_SOA_ACCESSOR_IMPL LAYOUT_MEMBER_NAME) + +/** + * Direct access to column pointer (const) and indexed access. + */ +#define _DECLARE_VIEW_SOA_CONST_ACCESSOR_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) \ + /* Column or scalar */ \ + SOA_HOST_DEVICE_INLINE auto LOCAL_NAME() const { \ + return typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::constAccess>(BOOST_PP_CAT(LOCAL_NAME, Parameters_))(); \ + } \ + SOA_HOST_DEVICE_INLINE auto LOCAL_NAME(size_t index) const { \ + return typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::constAccess>(BOOST_PP_CAT(LOCAL_NAME, Parameters_))(index); \ + } + +#define _DECLARE_VIEW_SOA_CONST_ACCESSOR(R, DATA, LAYOUT_MEMBER_NAME) \ + BOOST_PP_EXPAND(_DECLARE_VIEW_SOA_CONST_ACCESSOR_IMPL LAYOUT_MEMBER_NAME) + +/** + * SoA class member declaration (column pointers and parameters). + */ +#define _DECLARE_VIEW_SOA_MEMBER_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME, DATA) \ + typename BOOST_PP_CAT(SoAMetadata::ParametersTypeOf_, LOCAL_NAME) BOOST_PP_CAT(LOCAL_NAME, Parameters_); + +#define _DECLARE_VIEW_SOA_MEMBER(R, DATA, LAYOUT_MEMBER_NAME) \ + BOOST_PP_EXPAND(_DECLARE_VIEW_SOA_MEMBER_IMPL BOOST_PP_TUPLE_PUSH_BACK(LAYOUT_MEMBER_NAME, DATA)) + +/** + * Const SoA class member declaration (column pointers and parameters). + */ +#define _DECLARE_CONST_VIEW_SOA_MEMBER_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME, DATA) \ + typename BOOST_PP_CAT(SoAMetadata::ParametersTypeOf_, LOCAL_NAME)::ConstType BOOST_PP_CAT(LOCAL_NAME, Parameters_); + +#define _DECLARE_CONST_VIEW_SOA_MEMBER(R, DATA, LAYOUT_MEMBER_NAME) \ + BOOST_PP_EXPAND(_DECLARE_CONST_VIEW_SOA_MEMBER_IMPL BOOST_PP_TUPLE_PUSH_BACK(LAYOUT_MEMBER_NAME, DATA)) + +/* ---- MUTABLE VIEW -------------------------------------------------------------------------------------------------------------------- */ +// clang-format off +#define GENERATE_SOA_VIEW(CLASS, LAYOUTS_LIST, VALUE_LIST) \ + template \ + struct CLASS { \ + /* these could be moved to an external type trait to free up the symbol names */ \ + using self_type = CLASS; \ + typedef cms::soa::AlignmentEnforcement AlignmentEnforcement; \ + \ + /* For CUDA applications, we align to the 128 bytes of the cache lines. \ + * See https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#global-memory-3-0 this is still valid \ + * up to compute capability 8.X. \ + */ \ + constexpr static size_t defaultAlignment = cms::soa::CacheLineSize::defaultSize; \ + constexpr static size_t byteAlignment = ALIGNMENT; \ + constexpr static AlignmentEnforcement alignmentEnforcement = ALIGNMENT_ENFORCEMENT; \ + constexpr static size_t conditionalAlignment = \ + alignmentEnforcement == AlignmentEnforcement::Enforced ? byteAlignment : 0; \ + constexpr static cms::soa::RestrictQualify restrictQualify = RESTRICT_QUALIFY; \ + constexpr static cms::soa::RangeChecking rangeChecking = RANGE_CHECKING; \ + /* Those typedefs avoid having commas in macros (which is problematic) */ \ + template \ + using SoAValueWithConf = cms::soa::SoAValue; \ + \ + template \ + using SoAConstValueWithConf = cms::soa::SoAConstValue; \ + \ + /** \ + * Helper/friend class allowing SoA introspection. \ + */ \ + struct SoAMetadata { \ + friend CLASS; \ + SOA_HOST_DEVICE_INLINE size_t size() const { return parent_.nElements_; } \ + /* Alias layout or view types to name-derived identifyer to allow simpler definitions */ \ + _ITERATE_ON_ALL(_DECLARE_VIEW_LAYOUT_TYPE_ALIAS, ~, LAYOUTS_LIST) \ + \ + /* Alias member types to name-derived identifyer to allow simpler definitions */ \ + _ITERATE_ON_ALL(_DECLARE_VIEW_MEMBER_TYPE_ALIAS, BOOST_PP_EMPTY(), VALUE_LIST) \ + \ + /* Forbid copying to avoid const correctness evasion */ \ + SoAMetadata& operator=(const SoAMetadata&) = delete; \ + SoAMetadata(const SoAMetadata&) = delete; \ + \ + private: \ + SOA_HOST_DEVICE_INLINE SoAMetadata(const CLASS& parent) : parent_(parent) {} \ + const CLASS& parent_; \ + }; \ + friend SoAMetadata; \ + SOA_HOST_DEVICE_INLINE const SoAMetadata soaMetadata() const { return SoAMetadata(*this); } \ + SOA_HOST_DEVICE_INLINE SoAMetadata soaMetadata() { return SoAMetadata(*this); } \ + \ + /* Trivial constuctor */ \ + CLASS() {} \ + \ + /* Constructor relying on user provided layouts or views */ \ + SOA_HOST_ONLY CLASS(_ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_CONSTRUCTION_PARAMETERS, BOOST_PP_EMPTY(), LAYOUTS_LIST)) \ + : nElements_([&]() -> size_t { \ + bool set = false; \ + size_t ret = 0; \ + _ITERATE_ON_ALL(_UPDATE_SIZE_OF_VIEW, BOOST_PP_EMPTY(), LAYOUTS_LIST) \ + return ret; \ + }()), \ + _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_MEMBER_INITIALIZERS, ~, VALUE_LIST) {} \ + \ + /* Constructor relying on individually provided column addresses */ \ + SOA_HOST_ONLY CLASS(size_t nElements, \ + _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_CONSTRUCTION_BYCOLUMN_PARAMETERS, \ + BOOST_PP_EMPTY(), \ + VALUE_LIST)) \ + : nElements_(nElements), _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_MEMBER_INITIALIZERS_BYCOLUMN, ~, VALUE_LIST) {} \ + \ + struct const_element { \ + SOA_HOST_DEVICE_INLINE \ + const_element(size_t index, /* Declare parameters */ \ + _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_ELEMENT_VALUE_ARG, const, VALUE_LIST)) \ + : _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_CONST_ELEM_MEMBER_INIT, index, VALUE_LIST) {} \ + _ITERATE_ON_ALL(_DECLARE_VIEW_CONST_ELEMENT_ACCESSOR, ~, VALUE_LIST) \ + \ + private: \ + _ITERATE_ON_ALL(_DECLARE_VIEW_CONST_ELEMENT_VALUE_MEMBER, ~, VALUE_LIST) \ + }; \ + \ + struct element { \ + SOA_HOST_DEVICE_INLINE \ + element(size_t index, /* Declare parameters */ \ + _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_ELEMENT_VALUE_ARG, BOOST_PP_EMPTY(), VALUE_LIST)) \ + : _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_ELEM_MEMBER_INIT, index, VALUE_LIST) {} \ + SOA_HOST_DEVICE_INLINE \ + element& operator=(const element& other) { \ + _ITERATE_ON_ALL(_DECLARE_VIEW_ELEMENT_VALUE_COPY, ~, VALUE_LIST) \ + return *this; \ + } \ + _ITERATE_ON_ALL(_DECLARE_VIEW_ELEMENT_VALUE_MEMBER, ~, VALUE_LIST) \ + }; \ + \ + /* AoS-like accessor (non-const) */ \ + SOA_HOST_DEVICE_INLINE \ + element operator[](size_t index) { \ + if constexpr (rangeChecking == cms::soa::RangeChecking::Enabled) { \ + if (index >= nElements_) \ + SOA_THROW_OUT_OF_RANGE("Out of range index in " #CLASS "::operator[]") \ + } \ + return element(index, _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_ELEMENT_CONSTR_CALL, ~, VALUE_LIST)); \ + } \ + \ + /* AoS-like accessor (const) */ \ + SOA_HOST_DEVICE_INLINE \ + const_element operator[](size_t index) const { \ + if constexpr (rangeChecking == cms::soa::RangeChecking::Enabled) { \ + if (index >= nElements_) \ + SOA_THROW_OUT_OF_RANGE("Out of range index in " #CLASS "::operator[]") \ + } \ + return const_element(index, _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_ELEMENT_CONSTR_CALL, ~, VALUE_LIST)); \ + } \ + \ + /* accessors */ \ + _ITERATE_ON_ALL(_DECLARE_VIEW_SOA_ACCESSOR, ~, VALUE_LIST) \ + _ITERATE_ON_ALL(_DECLARE_VIEW_SOA_CONST_ACCESSOR, ~, VALUE_LIST) \ + \ + /* dump the SoA internal structure */ \ + template \ + SOA_HOST_ONLY friend void dump(); \ + \ + private: \ + size_t nElements_ = 0; \ + _ITERATE_ON_ALL(_DECLARE_VIEW_SOA_MEMBER, BOOST_PP_EMPTY(), VALUE_LIST) \ + }; +// clang-format on + +/* ---- CONST VIEW --------------------------------------------------------------------------------------------------------------------- */ +// clang-format off +#define GENERATE_SOA_CONST_VIEW(CLASS, LAYOUTS_LIST, VALUE_LIST) \ + template \ + struct CLASS { \ + /* these could be moved to an external type trait to free up the symbol names */ \ + using self_type = CLASS; \ + typedef cms::soa::AlignmentEnforcement AlignmentEnforcement; \ + \ + /* For CUDA applications, we align to the 128 bytes of the cache lines. \ + * See https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#global-memory-3-0 this is still valid \ + * up to compute capability 8.X. \ + */ \ + constexpr static size_t defaultAlignment = cms::soa::CacheLineSize::defaultSize; \ + constexpr static size_t byteAlignment = ALIGNMENT; \ + constexpr static AlignmentEnforcement alignmentEnforcement = ALIGNMENT_ENFORCEMENT; \ + constexpr static size_t conditionalAlignment = \ + alignmentEnforcement == AlignmentEnforcement::Enforced ? byteAlignment : 0; \ + constexpr static cms::soa::RestrictQualify restrictQualify = RESTRICT_QUALIFY; \ + constexpr static cms::soa::RangeChecking rangeChecking = RANGE_CHECKING; \ + /* Those typedefs avoid having commas in macros (which is problematic) */ \ + template \ + using SoAValueWithConf = cms::soa::SoAValue; \ + \ + template \ + using SoAConstValueWithConf = cms::soa::SoAConstValue; \ + /** \ + * Helper/friend class allowing SoA introspection. \ + */ \ + struct SoAMetadata { \ + friend CLASS; \ + SOA_HOST_DEVICE_INLINE size_t size() const { return parent_.nElements_; } \ + /* Alias layout/view types to name-derived identifyer to allow simpler definitions */ \ + _ITERATE_ON_ALL(_DECLARE_VIEW_LAYOUT_TYPE_ALIAS, ~, LAYOUTS_LIST) \ + \ + /* Alias member types to name-derived identifyer to allow simpler definitions */ \ + _ITERATE_ON_ALL(_DECLARE_VIEW_MEMBER_TYPE_ALIAS, const, VALUE_LIST) \ + \ + SoAMetadata& operator=(const SoAMetadata&) = delete; \ + SoAMetadata(const SoAMetadata&) = delete; \ + \ + private: \ + SOA_HOST_DEVICE_INLINE SoAMetadata(const CLASS& parent) : parent_(parent) {} \ + const CLASS& parent_; \ + }; \ + friend SoAMetadata; \ + SOA_HOST_DEVICE_INLINE const SoAMetadata soaMetadata() const { return SoAMetadata(*this); } \ + \ + /* Trivial constuctor */ \ + CLASS() {} \ + \ + /* Constructor relying on user provided layouts or views */ \ + SOA_HOST_ONLY CLASS(_ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_CONSTRUCTION_PARAMETERS, const, LAYOUTS_LIST)) \ + : nElements_([&]() -> size_t { \ + bool set = false; \ + size_t ret = 0; \ + _ITERATE_ON_ALL(_UPDATE_SIZE_OF_VIEW, BOOST_PP_EMPTY(), LAYOUTS_LIST) \ + return ret; \ + }()), \ + _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_MEMBER_INITIALIZERS, ~, VALUE_LIST) {} \ + \ + /* Constructor relying on individually provided column addresses */ \ + SOA_HOST_ONLY CLASS(size_t nElements, \ + _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_CONSTRUCTION_BYCOLUMN_PARAMETERS, const, VALUE_LIST)) \ + : nElements_(nElements), _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_MEMBER_INITIALIZERS_BYCOLUMN, ~, VALUE_LIST) {} \ + \ + struct const_element { \ + SOA_HOST_DEVICE_INLINE \ + const_element(size_t index, /* Declare parameters */ \ + _ITERATE_ON_ALL_COMMA(_DECLARE_CONST_VIEW_ELEMENT_VALUE_ARG, const, VALUE_LIST)) \ + : _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_CONST_ELEM_MEMBER_INIT, index, VALUE_LIST) {} \ + _ITERATE_ON_ALL(_DECLARE_VIEW_CONST_ELEMENT_ACCESSOR, ~, VALUE_LIST) \ + \ + private: \ + _ITERATE_ON_ALL(_DECLARE_VIEW_CONST_ELEMENT_VALUE_MEMBER, ~, VALUE_LIST) \ + }; \ + \ + /* AoS-like accessor (const) */ \ + SOA_HOST_DEVICE_INLINE \ + const_element operator[](size_t index) const { \ + if constexpr (rangeChecking == cms::soa::RangeChecking::Enabled) { \ + if (index >= nElements_) \ + SOA_THROW_OUT_OF_RANGE("Out of range index in " #CLASS "::operator[]") \ + } \ + return const_element(index, _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_ELEMENT_CONSTR_CALL, ~, VALUE_LIST)); \ + } \ + \ + /* accessors */ \ + _ITERATE_ON_ALL(_DECLARE_VIEW_SOA_CONST_ACCESSOR, ~, VALUE_LIST) \ + \ + /* dump the SoA internal structure */ \ + template \ + SOA_HOST_ONLY friend void dump(); \ + \ + private: \ + size_t nElements_ = 0; \ + _ITERATE_ON_ALL(_DECLARE_CONST_VIEW_SOA_MEMBER, const, VALUE_LIST) \ +}; +// clang-format on + +/** + * Helper macro turning layout field declaration into view field declaration. + */ +#define _VIEW_FIELD_FROM_LAYOUT_IMPL(VALUE_TYPE, CPP_TYPE, NAME, DATA) (DATA, NAME, NAME) + +#define _VIEW_FIELD_FROM_LAYOUT(R, DATA, VALUE_TYPE_NAME) \ + BOOST_PP_EXPAND((_VIEW_FIELD_FROM_LAYOUT_IMPL BOOST_PP_TUPLE_PUSH_BACK(VALUE_TYPE_NAME, DATA))) + +/** + * A macro defining both layout and view(s) in one go. + */ + +#define GENERATE_SOA_LAYOUT_VIEW_AND_CONST_VIEW(LAYOUT_NAME, VIEW_NAME, CONST_VIEW_NAME, ...) \ + GENERATE_SOA_LAYOUT(LAYOUT_NAME, __VA_ARGS__) \ + using BOOST_PP_CAT(LAYOUT_NAME, _default) = LAYOUT_NAME<>; \ + GENERATE_SOA_VIEW(VIEW_NAME, \ + SOA_VIEW_LAYOUT_LIST((BOOST_PP_CAT(LAYOUT_NAME, _default), BOOST_PP_CAT(instance_, LAYOUT_NAME))), \ + SOA_VIEW_VALUE_LIST(_ITERATE_ON_ALL_COMMA( \ + _VIEW_FIELD_FROM_LAYOUT, BOOST_PP_CAT(instance_, LAYOUT_NAME), __VA_ARGS__))) \ + GENERATE_SOA_CONST_VIEW( \ + CONST_VIEW_NAME, \ + SOA_VIEW_LAYOUT_LIST((BOOST_PP_CAT(LAYOUT_NAME, _default), BOOST_PP_CAT(instance_, LAYOUT_NAME))), \ + SOA_VIEW_VALUE_LIST( \ + _ITERATE_ON_ALL_COMMA(_VIEW_FIELD_FROM_LAYOUT, BOOST_PP_CAT(instance_, LAYOUT_NAME), __VA_ARGS__))) + +#define GENERATE_SOA_LAYOUT_AND_VIEW(LAYOUT_NAME, VIEW_NAME, ...) \ + GENERATE_SOA_LAYOUT(LAYOUT_NAME, __VA_ARGS__); \ + using BOOST_PP_CAT(LAYOUT_NAME, _default) = LAYOUT_NAME<>; \ + GENERATE_SOA_VIEW(VIEW_NAME, \ + SOA_VIEW_LAYOUT_LIST((BOOST_PP_CAT(LAYOUT_NAME, _default), BOOST_PP_CAT(instance_, LAYOUT_NAME))), \ + SOA_VIEW_VALUE_LIST(_ITERATE_ON_ALL_COMMA( \ + _VIEW_FIELD_FROM_LAYOUT, BOOST_PP_CAT(instance_, LAYOUT_NAME), __VA_ARGS__))) + +#define GENERATE_SOA_LAYOUT_AND_CONST_VIEW(LAYOUT_NAME, CONST_VIEW_NAME, ...) \ + GENERATE_SOA_LAYOUT(LAYOUT_NAME, __VA_ARGS__) \ + using BOOST_PP_CAT(LAYOUT_NAME, _default) = LAYOUT_NAME<>; \ + GENERATE_SOA_CONST_VIEW( \ + CONST_VIEW_NAME, \ + SOA_VIEW_LAYOUT_LIST((BOOST_PP_CAT(LAYOUT_NAME, _default), BOOST_PP_CAT(instance_, LAYOUT_NAME))), \ + SOA_VIEW_VALUE_LIST( \ + _ITERATE_ON_ALL_COMMA(_VIEW_FIELD_FROM_LAYOUT, BOOST_PP_CAT(instance_, LAYOUT_NAME), __VA_ARGS__))) + +#endif // ndef DataStructures_SoAView_h diff --git a/src/alpaka/test/alpaka/SoAStoreView_t.cc b/src/alpaka/test/alpaka/SoAStoreView_t.cc new file mode 100644 index 000000000..5114ab30a --- /dev/null +++ b/src/alpaka/test/alpaka/SoAStoreView_t.cc @@ -0,0 +1,300 @@ +#include + +#include "AlpakaCore/alpakaConfig.h" +#include "AlpakaCore/alpakaWorkDiv.h" +#include "Eigen/Geometry" +#include "DataFormats/SoALayout.h" +#include "DataFormats/SoAView.h" + +using namespace cms::alpakatools; +using namespace ALPAKA_ACCELERATOR_NAMESPACE; + +GENERATE_SOA_LAYOUT_AND_VIEW(SoAHostDeviceLayoutTemplate, + SoAHostDeviceViewTemplate, + // predefined static scalars + // size_t size; + // size_t alignment; + + // columns: one value per element + SOA_COLUMN(double, x), + SOA_COLUMN(double, y), + SOA_COLUMN(double, z), + SOA_EIGEN_COLUMN(Eigen::Vector3d, a), + SOA_EIGEN_COLUMN(Eigen::Vector3d, b), + SOA_EIGEN_COLUMN(Eigen::Vector3d, r), + // scalars: one value for the whole structure + SOA_SCALAR(const char*, description), + SOA_SCALAR(uint32_t, someNumber)) + +using SoAHostDeviceLayout = SoAHostDeviceLayoutTemplate<>; +using SoAHostDeviceView = + SoAHostDeviceViewTemplate; + +GENERATE_SOA_LAYOUT_AND_VIEW(SoADeviceOnlyLayoutTemplate, + SoADeviceOnlyViewTemplate, + SOA_COLUMN(uint16_t, color), + SOA_COLUMN(double, value), + SOA_COLUMN(double*, py), + SOA_COLUMN(uint32_t, count), + SOA_COLUMN(uint32_t, anotherCount)) + +using SoADeviceOnlyLayout = SoADeviceOnlyLayoutTemplate<>; +using SoADeviceOnlyView = + SoADeviceOnlyViewTemplate; + +// A 1 to 1 view of the store (except for unsupported types). +GENERATE_SOA_VIEW(SoAFullDeviceViewTemplate, + SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(SoAHostDeviceLayout, soaHD), + SOA_VIEW_LAYOUT(SoADeviceOnlyLayout, soaDO)), + SOA_VIEW_LAYOUT_LIST(SOA_VIEW_VALUE(soaHD, x), + SOA_VIEW_VALUE(soaHD, y), + SOA_VIEW_VALUE(soaHD, z), + SOA_VIEW_VALUE(soaDO, color), + SOA_VIEW_VALUE(soaDO, value), + SOA_VIEW_VALUE(soaDO, py), + SOA_VIEW_VALUE(soaDO, count), + SOA_VIEW_VALUE(soaDO, anotherCount), + SOA_VIEW_VALUE(soaHD, description), + SOA_VIEW_VALUE(soaHD, someNumber))) + +using SoAFullDeviceView = + SoAFullDeviceViewTemplate; + +// Eigen cross product kernel (on store) +struct crossProduct { + template + ALPAKA_FN_ACC void operator()(const T_Acc& acc, SoAHostDeviceView soa, const unsigned int numElements) const { + ::cms::alpakatools::for_each_element_in_grid(acc, numElements, [&](Idx i) { + auto si = soa[i]; + si.r() = si.a().cross(si.b()); + }); + } +}; + +// Device-only producer kernel +struct producerKernel { + template + ALPAKA_FN_ACC void operator()(const T_Acc& acc, SoAFullDeviceView soa, const unsigned int numElements) const { + ::cms::alpakatools::for_each_element_in_grid(acc, numElements, [&](Idx i) { + auto si = soa[i]; + si.color() &= 0x55 << i % (sizeof(si.color()) - sizeof(char)); + si.value() = sqrt(si.x() * si.x() + si.y() * si.y() + si.z() * si.z()); + }); + } +}; + +// Device-only consumer with result in host-device area +struct consumerKernel { + template + ALPAKA_FN_ACC void operator()(const T_Acc& acc, SoAFullDeviceView soa, const unsigned int numElements) const { + ::cms::alpakatools::for_each_element_in_grid(acc, numElements, [&](Idx i) { + auto si = soa[i]; + si.x() = si.color() * si.value(); + }); + } +}; + +// Get a view like the default, except for range checking +using RangeCheckingHostDeviceView = SoAHostDeviceViewTemplate; + +struct rangeCheckKernel { + template + ALPAKA_FN_ACC void operator()(const T_Acc& acc, RangeCheckingHostDeviceView soa) const { +#if defined(__CUDACC__) && defined(__CUDA_ARCH__) + printf("About to fail range check in CUDA thread: %d\n", threadIdx.x); +#endif + [[maybe_unused]] auto si = soa[soa.soaMetadata().size()]; + printf("We should not have reached here\n"); + } +}; + +template +Idx to_Idx(T v) { + return static_cast(v); +} + +int main(void) { + const DevHost host(alpaka::getDevByIdx(0u)); + const Device device(alpaka::getDevByIdx(0u)); + Queue queue(device); + + // Non-aligned number of elements to check alignment features. + constexpr unsigned int numElements = 65537; + + // Allocate buffer and store on host + Idx hostDeviceSize = SoAHostDeviceLayout::computeDataSize(numElements); + auto h_buf = alpaka::allocBuf(host, hostDeviceSize); + SoAHostDeviceLayout h_soahdLayout(alpaka::getPtrNative(h_buf), numElements); + SoAHostDeviceView h_soahd(h_soahdLayout); + + // Alocate buffer, stores and views on the device (single, shared buffer). + Idx deviceOnlySize = SoADeviceOnlyLayout::computeDataSize(numElements); + auto d_buf = alpaka::allocBuf(device, hostDeviceSize + deviceOnlySize); + SoAHostDeviceLayout d_soahdLayout(alpaka::getPtrNative(d_buf), numElements); + SoADeviceOnlyLayout d_soadoLayout(d_soahdLayout.soaMetadata().nextByte(), numElements); + SoAHostDeviceView d_soahdView(d_soahdLayout); + SoAFullDeviceView d_soaFullView(d_soahdLayout, d_soadoLayout); + + // Assert column alignments + assert(0 == reinterpret_cast(h_soahd.soaMetadata().addressOf_x()) % decltype(h_soahd)::byteAlignment); + assert(0 == reinterpret_cast(h_soahd.soaMetadata().addressOf_y()) % decltype(h_soahd)::byteAlignment); + assert(0 == reinterpret_cast(h_soahd.soaMetadata().addressOf_z()) % decltype(h_soahd)::byteAlignment); + assert(0 == reinterpret_cast(h_soahd.soaMetadata().addressOf_a()) % decltype(h_soahd)::byteAlignment); + assert(0 == reinterpret_cast(h_soahd.soaMetadata().addressOf_b()) % decltype(h_soahd)::byteAlignment); + assert(0 == reinterpret_cast(h_soahd.soaMetadata().addressOf_r()) % decltype(h_soahd)::byteAlignment); + assert(0 == + reinterpret_cast(h_soahd.soaMetadata().addressOf_description()) % decltype(h_soahd)::byteAlignment); + assert(0 == + reinterpret_cast(h_soahd.soaMetadata().addressOf_someNumber()) % decltype(h_soahd)::byteAlignment); + + assert(0 == reinterpret_cast(d_soahdLayout.soaMetadata().addressOf_x()) % + decltype(d_soahdLayout)::byteAlignment); + assert(0 == reinterpret_cast(d_soahdLayout.soaMetadata().addressOf_y()) % + decltype(d_soahdLayout)::byteAlignment); + assert(0 == reinterpret_cast(d_soahdLayout.soaMetadata().addressOf_z()) % + decltype(d_soahdLayout)::byteAlignment); + assert(0 == reinterpret_cast(d_soahdLayout.soaMetadata().addressOf_a()) % + decltype(d_soahdLayout)::byteAlignment); + assert(0 == reinterpret_cast(d_soahdLayout.soaMetadata().addressOf_b()) % + decltype(d_soahdLayout)::byteAlignment); + assert(0 == reinterpret_cast(d_soahdLayout.soaMetadata().addressOf_r()) % + decltype(d_soahdLayout)::byteAlignment); + assert(0 == reinterpret_cast(d_soahdLayout.soaMetadata().addressOf_description()) % + decltype(d_soahdLayout)::byteAlignment); + assert(0 == reinterpret_cast(d_soahdLayout.soaMetadata().addressOf_someNumber()) % + decltype(d_soahdLayout)::byteAlignment); + + assert(0 == reinterpret_cast(d_soadoLayout.soaMetadata().addressOf_color()) % + decltype(d_soadoLayout)::byteAlignment); + assert(0 == reinterpret_cast(d_soadoLayout.soaMetadata().addressOf_value()) % + decltype(d_soadoLayout)::byteAlignment); + assert(0 == reinterpret_cast(d_soadoLayout.soaMetadata().addressOf_py()) % + decltype(d_soadoLayout)::byteAlignment); + assert(0 == reinterpret_cast(d_soadoLayout.soaMetadata().addressOf_count()) % + decltype(d_soadoLayout)::byteAlignment); + assert(0 == reinterpret_cast(d_soadoLayout.soaMetadata().addressOf_anotherCount()) % + decltype(d_soadoLayout)::byteAlignment); + + // Views should get the same alignment as the stores they refer to + assert(0 == reinterpret_cast(d_soaFullView.soaMetadata().addressOf_x()) % + decltype(d_soaFullView)::byteAlignment); + assert(0 == reinterpret_cast(d_soaFullView.soaMetadata().addressOf_y()) % + decltype(d_soaFullView)::byteAlignment); + assert(0 == reinterpret_cast(d_soaFullView.soaMetadata().addressOf_z()) % + decltype(d_soaFullView)::byteAlignment); + // Limitation of views: we have to get scalar member addresses via metadata. + assert(0 == reinterpret_cast(d_soaFullView.soaMetadata().addressOf_description()) % + decltype(d_soaFullView)::byteAlignment); + assert(0 == reinterpret_cast(d_soaFullView.soaMetadata().addressOf_someNumber()) % + decltype(d_soaFullView)::byteAlignment); + assert(0 == reinterpret_cast(d_soaFullView.soaMetadata().addressOf_color()) % + decltype(d_soaFullView)::byteAlignment); + assert(0 == reinterpret_cast(d_soaFullView.soaMetadata().addressOf_value()) % + decltype(d_soaFullView)::byteAlignment); + assert(0 == reinterpret_cast(d_soaFullView.soaMetadata().addressOf_py()) % + decltype(d_soaFullView)::byteAlignment); + assert(0 == reinterpret_cast(d_soaFullView.soaMetadata().addressOf_count()) % + decltype(d_soaFullView)::byteAlignment); + assert(0 == reinterpret_cast(d_soaFullView.soaMetadata().addressOf_anotherCount()) % + decltype(d_soaFullView)::byteAlignment); + + // Initialize and fill the host buffer + std::memset(h_soahdLayout.soaMetadata().data(), 0, hostDeviceSize); + for (Idx i = 0; i < numElements; ++i) { + auto si = h_soahd[i]; + si.x() = si.a()(0) = si.b()(2) = 1.0 * i + 1.0; + si.y() = si.a()(1) = si.b()(1) = 2.0 * i; + si.z() = si.a()(2) = si.b()(0) = 3.0 * i - 1.0; + } + auto& sn = h_soahd.someNumber(); + sn = numElements + 2; + + // Push to device + alpaka::memcpy(queue, d_buf, h_buf, hostDeviceSize); + + // Process on device + const WorkDiv1D& workDivMaxNumModules = make_workdiv((numElements + 255) / 256, 256); + + alpaka::enqueue(queue, + alpaka::createTaskKernel<::ALPAKA_ACCELERATOR_NAMESPACE::Acc1D>( + workDivMaxNumModules, crossProduct(), d_soahdView, numElements)); + + // Paint the device only with 0xFF initially + alpaka::ViewSubView d_doSubBuf( + d_buf, + /* length */ Idx(d_soadoLayout.soaMetadata().byteSize()), + /* offset */ Idx(d_soahdLayout.soaMetadata().byteSize())); + alpaka::memset(queue, d_doSubBuf, 0xFF, Idx(d_soadoLayout.soaMetadata().byteSize())); + + // Produce to the device only area + alpaka::enqueue(queue, + alpaka::createTaskKernel<::ALPAKA_ACCELERATOR_NAMESPACE::Acc1D>( + workDivMaxNumModules, producerKernel(), d_soaFullView, numElements)); + + // Consume the device only area and generate a result on the host-device area + alpaka::enqueue(queue, + alpaka::createTaskKernel<::ALPAKA_ACCELERATOR_NAMESPACE::Acc1D>( + workDivMaxNumModules, consumerKernel(), d_soaFullView, numElements)); + + // Get result back + alpaka::memcpy(queue, h_buf, d_buf, hostDeviceSize); + + // Wait and validate. + alpaka::wait(queue); + for (Idx i = 0; i < numElements; ++i) { + auto si = h_soahd[i]; + assert(si.r() == si.a().cross(si.b())); + double initialX = 1.0 * i + 1.0; + double initialY = 2.0 * i; + double initialZ = 3.0 * i - 1.0; + uint16_t expectedColor = 0x55 << i % (sizeof(uint16_t) - sizeof(char)); + double expectedX = expectedColor * sqrt(initialX * initialX + initialY * initialY + initialZ * initialZ); + if (abs(si.x() - expectedX) / expectedX >= 2 * std::numeric_limits::epsilon()) { + std::cout << "X failed: for i=" << i << std::endl + << "initialX=" << initialX << " initialY=" << initialY << " initialZ=" << initialZ << std::endl + << "expectedX=" << expectedX << std::endl + << "resultX=" << si.x() << " resultY=" << si.y() << " resultZ=" << si.z() << std::endl + << "relativeDiff=" << abs(si.x() - expectedX) / expectedX + << " epsilon=" << std::numeric_limits::epsilon() << std::endl; + assert(false); + } + } + + // Validation of range checking + try { + // Get a view like the default, except for range checking + SoAHostDeviceViewTemplate + soa1viewRangeChecking(h_soahdLayout); + // This should throw an exception + [[maybe_unused]] auto si = soa1viewRangeChecking[soa1viewRangeChecking.soaMetadata().size()]; + assert(false); + } catch (const std::out_of_range&) { + } + + // Validation of range checking in a kernel + // Get a view like the default, except for range checking + RangeCheckingHostDeviceView soa1viewRangeChecking(d_soahdLayout); + // This should throw an exception in the kernel + try { + alpaka::enqueue(queue, + alpaka::createTaskKernel<::ALPAKA_ACCELERATOR_NAMESPACE::Acc1D>( + make_workdiv(1, 1), rangeCheckKernel(), soa1viewRangeChecking)); + } catch (const std::out_of_range&) { + std::cout << "Exception received in enqueue." << std::endl; + } + + // Wait and validate (that we failed). + try { + alpaka::wait(queue); + } catch (const std::runtime_error&) { + std::cout << "Exception received in wait." << std::endl; + } + + std::cout << "OK" << std::endl; +}