From 2004830b6e8428348eb52e9168b950f90044df75 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Fri, 28 Oct 2022 23:42:19 +0200 Subject: [PATCH 1/5] Fix the return type for const view of scalars Return const scalars by const reference instead of by value, which would break non-trivially copiable types. --- DataFormats/SoATemplate/interface/SoACommon.h | 10 +++++----- DataFormats/SoATemplate/interface/SoAView.h | 8 ++++++-- 2 files changed, 11 insertions(+), 7 deletions(-) diff --git a/DataFormats/SoATemplate/interface/SoACommon.h b/DataFormats/SoATemplate/interface/SoACommon.h index 70f344c66d330..01d8a14db8099 100644 --- a/DataFormats/SoATemplate/interface/SoACommon.h +++ b/DataFormats/SoATemplate/interface/SoACommon.h @@ -612,7 +612,7 @@ namespace cms::soa { : params_(params) {} SOA_HOST_DEVICE SOA_INLINE const T* operator()() const { return params_.addr_; } using NoParamReturnType = const T*; - SOA_HOST_DEVICE SOA_INLINE T operator()(size_type index) const { return params_.addr_[index]; } + SOA_HOST_DEVICE SOA_INLINE T const& operator()(size_type index) const { return params_.addr_[index]; } private: SoAConstParametersImpl params_; @@ -638,8 +638,8 @@ namespace cms::soa { struct SoAColumnAccessorsImpl { SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAConstParametersImpl& params) : params_(params) {} - SOA_HOST_DEVICE SOA_INLINE T operator()() const { return *params_.addr_; } - using NoParamReturnType = T; + SOA_HOST_DEVICE SOA_INLINE T const& operator()() const { return *params_.addr_; } + using NoParamReturnType = T const&; SOA_HOST_DEVICE SOA_INLINE void operator()(size_type index) const { assert(false && "Indexed access impossible for SoA scalars."); } @@ -667,8 +667,8 @@ namespace cms::soa { struct SoAColumnAccessorsImpl { SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAConstParametersImpl& params) : params_(params) {} - SOA_HOST_DEVICE SOA_INLINE const typename T::Scalar* operator()() const { return params_.addr_; } - using NoParamReturnType = typename T::Scalar*; + SOA_HOST_DEVICE SOA_INLINE typename T::Scalar const* operator()() const { return params_.addr_; } + using NoParamReturnType = typename T::Scalar const*; //SOA_HOST_DEVICE SOA_INLINE T operator()(size_type index) const { return params_.addr_[index]; } private: diff --git a/DataFormats/SoATemplate/interface/SoAView.h b/DataFormats/SoATemplate/interface/SoAView.h index 945626db03851..2cdb037d4d97d 100644 --- a/DataFormats/SoATemplate/interface/SoAView.h +++ b/DataFormats/SoATemplate/interface/SoAView.h @@ -388,12 +388,16 @@ namespace cms::soa { // clang-format off #define _DECLARE_VIEW_SOA_CONST_ACCESSOR_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) \ /* Column or scalar */ \ - SOA_HOST_DEVICE SOA_INLINE auto LOCAL_NAME() const { \ + SOA_HOST_DEVICE SOA_INLINE \ + typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::constAccess>::NoParamReturnType \ + 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 SOA_INLINE auto LOCAL_NAME(size_type index) const { \ + SOA_HOST_DEVICE SOA_INLINE auto const& LOCAL_NAME(size_type index) const { \ return typename cms::soa::SoAAccessors:: \ template ColumnType::template AccessType< \ cms::soa::SoAAccessType::constAccess>(BOOST_PP_CAT(LOCAL_NAME, Parameters_))(index); \ From 50cbaadbaecf3c51277c98ee3704a5afc3bfdf27 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Sat, 29 Oct 2022 00:09:54 +0200 Subject: [PATCH 2/5] Update the SoA tests to check the fields' address Update the SoA tests to check the various ways of taking the address of a column, scalar, or Eigen column fields. --- .../AlpakaTest/plugins/TestAlpakaAnalyzer.cc | 19 +++++++++++++++++++ 1 file changed, 19 insertions(+) diff --git a/HeterogeneousCore/AlpakaTest/plugins/TestAlpakaAnalyzer.cc b/HeterogeneousCore/AlpakaTest/plugins/TestAlpakaAnalyzer.cc index c64e91cb9d0dc..5d4a87492ef50 100644 --- a/HeterogeneousCore/AlpakaTest/plugins/TestAlpakaAnalyzer.cc +++ b/HeterogeneousCore/AlpakaTest/plugins/TestAlpakaAnalyzer.cc @@ -88,6 +88,25 @@ class TestAlpakaAnalyzer : public edm::stream::EDAnalyzer<> { reinterpret_cast(view.metadata().addressOf_r()); } + assert(view.metadata().addressOf_x() == view.x()); + assert(view.metadata().addressOf_x() == &view.x(0)); + assert(view.metadata().addressOf_x() == &view[0].x()); + assert(view.metadata().addressOf_y() == view.y()); + assert(view.metadata().addressOf_y() == &view.y(0)); + assert(view.metadata().addressOf_y() == &view[0].y()); + assert(view.metadata().addressOf_z() == view.z()); + assert(view.metadata().addressOf_z() == &view.z(0)); + assert(view.metadata().addressOf_z() == &view[0].z()); + assert(view.metadata().addressOf_id() == view.id()); + assert(view.metadata().addressOf_id() == &view.id(0)); + assert(view.metadata().addressOf_id() == &view[0].id()); + assert(view.metadata().addressOf_m() == view.m()); + //assert(view.metadata().addressOf_m() == &view.m(0).coeffRef(0,0)); // view.m(0) does not seem to be supported for Eigen columns ? + assert(view.metadata().addressOf_m() == &view[0].m().coeffRef(0, 0)); + assert(view.metadata().addressOf_r() == &view.r()); + //assert(view.metadata().addressOf_r() == &view.r(0)); // cannot access a scalar with an index + //assert(view.metadata().addressOf_r() == &view[0].r()); // cannot access a scalar via a SoA row-like accessor + const portabletest::Matrix matrix{{1, 2, 3, 4, 5, 6}, {2, 4, 6, 8, 10, 12}, {3, 6, 9, 12, 15, 18}}; assert(view.r() == 1.); for (int32_t i = 0; i < view.metadata().size(); ++i) { From 9a945419f75d8631ae3147f3cda07fec470c528d Mon Sep 17 00:00:00 2001 From: Eric Cano <37585813+ericcano@users.noreply.github.com> Date: Fri, 11 Nov 2022 12:48:51 +0100 Subject: [PATCH 3/5] Remove __builtin_assume_aligned compiler hints from SoA fields Remove the __builtin_assume_aligned compiler hints from the SoA fields, as it breaks the __restrict__ optimizations in nvcc. --- DataFormats/SoATemplate/README.md | 5 ++-- DataFormats/SoATemplate/interface/SoACommon.h | 28 +++++-------------- 2 files changed, 10 insertions(+), 23 deletions(-) diff --git a/DataFormats/SoATemplate/README.md b/DataFormats/SoATemplate/README.md index ba1b00e70233d..d8747a337c5b1 100644 --- a/DataFormats/SoATemplate/README.md +++ b/DataFormats/SoATemplate/README.md @@ -66,8 +66,9 @@ Serialization of Eigen data is not yet supported. The template shared by layouts and parameters are: - Byte aligment (defaulting to the nVidia GPU cache line size (128 bytes)) - Alignment enforcement (`relaxed` or `enforced`). When enforced, the alignment will be checked at construction - time, and the accesses are done with compiler hinting (using the widely supported `__builtin_assume_aligned` - intrinsic). + time.~~, and the accesses are done with compiler hinting (using the widely supported `__builtin_assume_aligned` + intrinsic).~~ It turned out that hinting `nvcc` for alignement removed the benefit of more important `__restrict__` + hinting. The `__builtin_assume_aligned` is hence currently not use. In addition, the views also provide access parameters: - Restrict qualify: add restrict hints to read accesses, so that the compiler knows it can relax accesses to the diff --git a/DataFormats/SoATemplate/interface/SoACommon.h b/DataFormats/SoATemplate/interface/SoACommon.h index 01d8a14db8099..d701f7935a81a 100644 --- a/DataFormats/SoATemplate/interface/SoACommon.h +++ b/DataFormats/SoATemplate/interface/SoACommon.h @@ -287,19 +287,19 @@ namespace cms::soa { SOA_HOST_DEVICE SOA_INLINE Ref operator()() { // Ptr type will add the restrict qualifyer if needed - Ptr col = alignedCol(); + Ptr col = col_; return col[idx_]; } SOA_HOST_DEVICE SOA_INLINE RefToConst operator()() const { // PtrToConst type will add the restrict qualifyer if needed - PtrToConst col = alignedCol(); + PtrToConst col = col_(); return col[idx_]; } - SOA_HOST_DEVICE SOA_INLINE Ptr operator&() { return &alignedCol()[idx_]; } + SOA_HOST_DEVICE SOA_INLINE Ptr operator&() { return &col_[idx_]; } - SOA_HOST_DEVICE SOA_INLINE PtrToConst operator&() const { return &alignedCol()[idx_]; } + SOA_HOST_DEVICE SOA_INLINE PtrToConst operator&() const { return &col_[idx_]; } /* This was an attempt to implement the syntax * @@ -318,7 +318,7 @@ namespace cms::soa { template SOA_HOST_DEVICE SOA_INLINE Ref operator=(const T2& v) { - return alignedCol()[idx_] = v; + return col_[idx_] = v; } */ @@ -327,13 +327,6 @@ namespace cms::soa { static constexpr auto valueSize = sizeof(T); private: - SOA_HOST_DEVICE SOA_INLINE Ptr alignedCol() const { - if constexpr (ALIGNMENT) { - return reinterpret_cast(__builtin_assume_aligned(col_, ALIGNMENT)); - } - return reinterpret_cast(col_); - } - size_type idx_; T* col_; }; @@ -437,11 +430,11 @@ namespace cms::soa { SOA_HOST_DEVICE SOA_INLINE RefToConst operator()() const { // Ptr type will add the restrict qualifyer if needed - PtrToConst col = alignedCol(); + PtrToConst col = col_; return col[idx_]; } - SOA_HOST_DEVICE SOA_INLINE const T* operator&() const { return &alignedCol()[idx_]; } + SOA_HOST_DEVICE SOA_INLINE const T* operator&() const { return &col_[idx_]; } /* This was an attempt to implement the syntax * @@ -461,13 +454,6 @@ namespace cms::soa { static constexpr auto valueSize = sizeof(T); private: - SOA_HOST_DEVICE SOA_INLINE PtrToConst alignedCol() const { - if constexpr (ALIGNMENT) { - return reinterpret_cast(__builtin_assume_aligned(col_, ALIGNMENT)); - } - return reinterpret_cast(col_); - } - size_type idx_; const T* col_; }; From 4f89eee8762aa44960de9c4211f4710decd628a2 Mon Sep 17 00:00:00 2001 From: Eric Cano <37585813+ericcano@users.noreply.github.com> Date: Tue, 15 Nov 2022 01:29:10 +0100 Subject: [PATCH 4/5] Add support for SoA view-level indexed access of Eigen columns Add support for SoA view-level indexed access of Eigen columns and enable the corresponding test. --- DataFormats/SoATemplate/interface/SoACommon.h | 53 ++++++++++++------- DataFormats/SoATemplate/interface/SoAView.h | 50 ++++++++++++----- .../AlpakaTest/plugins/TestAlpakaAnalyzer.cc | 2 +- 3 files changed, 73 insertions(+), 32 deletions(-) diff --git a/DataFormats/SoATemplate/interface/SoACommon.h b/DataFormats/SoATemplate/interface/SoACommon.h index d701f7935a81a..220ea73176d71 100644 --- a/DataFormats/SoATemplate/interface/SoACommon.h +++ b/DataFormats/SoATemplate/interface/SoACommon.h @@ -570,7 +570,7 @@ namespace cms::soa { /* Column accessors: templates implementing the global accesors (soa::x() and soa::x(index) */ enum class SoAAccessType : bool { mutableAccess, constAccess }; - template + template struct SoAColumnAccessorsImpl {}; // TODO from Eric Cano: @@ -578,13 +578,13 @@ namespace cms::soa { // - SFINAE-based const/non const variants // Column - template - struct SoAColumnAccessorsImpl { - //SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(T* baseAddress) : baseAddress_(baseAddress) {} + template + struct SoAColumnAccessorsImpl { SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAParametersImpl& params) : params_(params) {} SOA_HOST_DEVICE SOA_INLINE T* operator()() { return params_.addr_; } using NoParamReturnType = T*; + using ParamReturnType = T&; SOA_HOST_DEVICE SOA_INLINE T& operator()(size_type index) { return params_.addr_[index]; } private: @@ -592,12 +592,13 @@ namespace cms::soa { }; // Const column - template - struct SoAColumnAccessorsImpl { + template + struct SoAColumnAccessorsImpl { SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAConstParametersImpl& params) : params_(params) {} SOA_HOST_DEVICE SOA_INLINE const T* operator()() const { return params_.addr_; } using NoParamReturnType = const T*; + using ParamReturnType = const T&; SOA_HOST_DEVICE SOA_INLINE T const& operator()(size_type index) const { return params_.addr_[index]; } private: @@ -605,12 +606,13 @@ namespace cms::soa { }; // Scalar - template - struct SoAColumnAccessorsImpl { + template + struct SoAColumnAccessorsImpl { SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAParametersImpl& params) : params_(params) {} SOA_HOST_DEVICE SOA_INLINE T& operator()() { return *params_.addr_; } using NoParamReturnType = T&; + using ParamReturnType = void; SOA_HOST_DEVICE SOA_INLINE void operator()(size_type index) const { assert(false && "Indexed access impossible for SoA scalars."); } @@ -620,12 +622,13 @@ namespace cms::soa { }; // Const scalar - template - struct SoAColumnAccessorsImpl { + template + struct SoAColumnAccessorsImpl { SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAConstParametersImpl& params) : params_(params) {} SOA_HOST_DEVICE SOA_INLINE T const& operator()() const { return *params_.addr_; } using NoParamReturnType = T const&; + using ParamReturnType = void; SOA_HOST_DEVICE SOA_INLINE void operator()(size_type index) const { assert(false && "Indexed access impossible for SoA scalars."); } @@ -635,27 +638,32 @@ namespace cms::soa { }; // Eigen-type - template - struct SoAColumnAccessorsImpl { - //SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(T* baseAddress) : baseAddress_(baseAddress) {} + template + struct SoAColumnAccessorsImpl { SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAParametersImpl& params) : params_(params) {} SOA_HOST_DEVICE SOA_INLINE typename T::Scalar* operator()() { return params_.addr_; } using NoParamReturnType = typename T::Scalar*; - //SOA_HOST_DEVICE SOA_INLINE T& operator()(size_type index) { return params_.addr_[index]; } + using ParamReturnType = typename SoAValue::MapType; + SOA_HOST_DEVICE SOA_INLINE ParamReturnType operator()(size_type index) { + return SoAValue(index, params_)(); + } private: SoAParametersImpl params_; }; // Const Eigen-type - template - struct SoAColumnAccessorsImpl { + template + struct SoAColumnAccessorsImpl { SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAConstParametersImpl& params) : params_(params) {} SOA_HOST_DEVICE SOA_INLINE typename T::Scalar const* operator()() const { return params_.addr_; } using NoParamReturnType = typename T::Scalar const*; - //SOA_HOST_DEVICE SOA_INLINE T operator()(size_type index) const { return params_.addr_[index]; } + using ParamReturnType = typename SoAValue::CMapType; + SOA_HOST_DEVICE SOA_INLINE ParamReturnType operator()(size_type index) const { + return SoAConstValue(index, params_)(); + } private: SoAConstParametersImpl params_; @@ -667,8 +675,15 @@ namespace cms::soa { template struct ColumnType { template - struct AccessType : public SoAColumnAccessorsImpl { - using SoAColumnAccessorsImpl::SoAColumnAccessorsImpl; + struct AccessType { + template + struct Alignment { + template + struct RestrictQualifier + : public SoAColumnAccessorsImpl { + using SoAColumnAccessorsImpl::SoAColumnAccessorsImpl; + }; + }; }; }; }; diff --git a/DataFormats/SoATemplate/interface/SoAView.h b/DataFormats/SoATemplate/interface/SoAView.h index 2cdb037d4d97d..8ced878378f24 100644 --- a/DataFormats/SoATemplate/interface/SoAView.h +++ b/DataFormats/SoATemplate/interface/SoAView.h @@ -92,6 +92,16 @@ namespace cms::soa { typename BOOST_PP_CAT(TypeOf_, LAYOUT_NAME)::Metadata::BOOST_PP_CAT(ParametersTypeOf_, LAYOUT_MEMBER); \ constexpr static cms::soa::SoAColumnType BOOST_PP_CAT(ColumnTypeOf_, LOCAL_NAME) = \ BOOST_PP_CAT(TypeOf_, LAYOUT_NAME)::Metadata::BOOST_PP_CAT(ColumnTypeOf_, LAYOUT_MEMBER); \ + using BOOST_PP_CAT(ConstAccessorOf_, LOCAL_NAME) = \ + typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::constAccess>::template Alignment:: \ + template RestrictQualifier ; \ + using BOOST_PP_CAT(MutableAccessorOf_, LOCAL_NAME) = \ + typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::mutableAccess>::template Alignment:: \ + template RestrictQualifier ; \ SOA_HOST_DEVICE SOA_INLINE \ const auto BOOST_PP_CAT(parametersOf_, LOCAL_NAME)() const { \ return CAST(parent_.BOOST_PP_CAT(LOCAL_NAME, Parameters_)); \ @@ -364,18 +374,26 @@ namespace cms::soa { SOA_HOST_DEVICE SOA_INLINE \ typename cms::soa::SoAAccessors:: \ template ColumnType::template AccessType< \ - cms::soa::SoAAccessType::mutableAccess>::NoParamReturnType \ + cms::soa::SoAAccessType::mutableAccess>::template Alignment:: \ + template RestrictQualifier::NoParamReturnType \ LOCAL_NAME() { \ return typename cms::soa::SoAAccessors:: \ template ColumnType::template AccessType< \ - cms::soa::SoAAccessType::mutableAccess>(const_cast_SoAParametersImpl( \ - base_type:: BOOST_PP_CAT(LOCAL_NAME, Parameters_)))(); \ + cms::soa::SoAAccessType::mutableAccess>::template Alignment:: \ + template RestrictQualifier(const_cast_SoAParametersImpl( \ + base_type:: BOOST_PP_CAT(LOCAL_NAME, Parameters_)))(); \ } \ - SOA_HOST_DEVICE SOA_INLINE auto& LOCAL_NAME(size_type index) { \ - return typename cms::soa::SoAAccessors:: \ + SOA_HOST_DEVICE SOA_INLINE \ + typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::mutableAccess>::template Alignment:: \ + template RestrictQualifier::ParamReturnType \ + LOCAL_NAME(size_type index) { \ + return typename cms::soa::SoAAccessors:: \ template ColumnType::template AccessType< \ - cms::soa::SoAAccessType::mutableAccess>(const_cast_SoAParametersImpl( \ - base_type:: BOOST_PP_CAT(LOCAL_NAME, Parameters_)))(index); \ + cms::soa::SoAAccessType::mutableAccess>::template Alignment:: \ + template RestrictQualifier(const_cast_SoAParametersImpl( \ + base_type:: BOOST_PP_CAT(LOCAL_NAME, Parameters_)))(index); \ } // clang-format on @@ -391,16 +409,24 @@ namespace cms::soa { SOA_HOST_DEVICE SOA_INLINE \ typename cms::soa::SoAAccessors:: \ template ColumnType::template AccessType< \ - cms::soa::SoAAccessType::constAccess>::NoParamReturnType \ + cms::soa::SoAAccessType::constAccess>::template Alignment:: \ + template RestrictQualifier::NoParamReturnType \ LOCAL_NAME() const { \ return typename cms::soa::SoAAccessors:: \ template ColumnType::template AccessType< \ - cms::soa::SoAAccessType::constAccess>(BOOST_PP_CAT(LOCAL_NAME, Parameters_))(); \ + cms::soa::SoAAccessType::constAccess>::template Alignment:: \ + template RestrictQualifier(BOOST_PP_CAT(LOCAL_NAME, Parameters_))(); \ } \ - SOA_HOST_DEVICE SOA_INLINE auto const& LOCAL_NAME(size_type index) const { \ - return typename cms::soa::SoAAccessors:: \ + SOA_HOST_DEVICE SOA_INLINE \ + typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::constAccess>::template Alignment:: \ + template RestrictQualifier::ParamReturnType \ + LOCAL_NAME(size_type index) const { \ + return typename cms::soa::SoAAccessors:: \ template ColumnType::template AccessType< \ - cms::soa::SoAAccessType::constAccess>(BOOST_PP_CAT(LOCAL_NAME, Parameters_))(index); \ + cms::soa::SoAAccessType::constAccess>::template Alignment:: \ + template RestrictQualifier(BOOST_PP_CAT(LOCAL_NAME, Parameters_))(index); \ } // clang-format on diff --git a/HeterogeneousCore/AlpakaTest/plugins/TestAlpakaAnalyzer.cc b/HeterogeneousCore/AlpakaTest/plugins/TestAlpakaAnalyzer.cc index 5d4a87492ef50..1058af9c9e77e 100644 --- a/HeterogeneousCore/AlpakaTest/plugins/TestAlpakaAnalyzer.cc +++ b/HeterogeneousCore/AlpakaTest/plugins/TestAlpakaAnalyzer.cc @@ -101,7 +101,7 @@ class TestAlpakaAnalyzer : public edm::stream::EDAnalyzer<> { assert(view.metadata().addressOf_id() == &view.id(0)); assert(view.metadata().addressOf_id() == &view[0].id()); assert(view.metadata().addressOf_m() == view.m()); - //assert(view.metadata().addressOf_m() == &view.m(0).coeffRef(0,0)); // view.m(0) does not seem to be supported for Eigen columns ? + assert(view.metadata().addressOf_m() == &view.m(0).coeffRef(0, 0)); assert(view.metadata().addressOf_m() == &view[0].m().coeffRef(0, 0)); assert(view.metadata().addressOf_r() == &view.r()); //assert(view.metadata().addressOf_r() == &view.r(0)); // cannot access a scalar with an index From 067becf9bc2afa9257c83e4261d11714746ba786 Mon Sep 17 00:00:00 2001 From: Eric Cano <37585813+ericcano@users.noreply.github.com> Date: Wed, 16 Nov 2022 14:44:04 +0100 Subject: [PATCH 5/5] Add range checking to SoA view-level indexed accessors --- DataFormats/SoATemplate/interface/SoALayout.h | 11 ----- DataFormats/SoATemplate/interface/SoAView.h | 12 ++++- .../SoATemplate/test/SoALayoutAndView_t.cu | 19 ++++++-- .../AlpakaTest/plugins/TestAlpakaAnalyzer.cc | 46 +++++++++++-------- 4 files changed, 54 insertions(+), 34 deletions(-) diff --git a/DataFormats/SoATemplate/interface/SoALayout.h b/DataFormats/SoATemplate/interface/SoALayout.h index 9cee0272a3dff..ad262f48334a0 100644 --- a/DataFormats/SoATemplate/interface/SoALayout.h +++ b/DataFormats/SoATemplate/interface/SoALayout.h @@ -581,17 +581,6 @@ throw std::runtime_error("In " #CLASS "::" #CLASS ": unexpected end pointer."); \ } \ \ - /* Range checker conditional to the macro _DO_RANGECHECK */ \ - SOA_HOST_DEVICE SOA_INLINE \ - void rangeCheck(size_type index) const { \ - if constexpr (_DO_RANGECHECK) { \ - if (index >= elements_) { \ - printf("In " #CLASS "::rangeCheck(): index out of range: %zu with elements: %zu\n", index, elements_); \ - assert(false); \ - } \ - } \ - } \ - \ /* Data members */ \ std::byte* mem_; \ size_type elements_; \ diff --git a/DataFormats/SoATemplate/interface/SoAView.h b/DataFormats/SoATemplate/interface/SoAView.h index 8ced878378f24..0099587a98180 100644 --- a/DataFormats/SoATemplate/interface/SoAView.h +++ b/DataFormats/SoATemplate/interface/SoAView.h @@ -389,7 +389,11 @@ namespace cms::soa { cms::soa::SoAAccessType::mutableAccess>::template Alignment:: \ template RestrictQualifier::ParamReturnType \ LOCAL_NAME(size_type index) { \ - return typename cms::soa::SoAAccessors:: \ + if constexpr (rangeChecking == cms::soa::RangeChecking::enabled) { \ + if (index >= base_type::elements_) \ + SOA_THROW_OUT_OF_RANGE("Out of range index in mutable " #LOCAL_NAME "(size_type index)") \ + } \ + return typename cms::soa::SoAAccessors:: \ template ColumnType::template AccessType< \ cms::soa::SoAAccessType::mutableAccess>::template Alignment:: \ template RestrictQualifier(const_cast_SoAParametersImpl( \ @@ -423,7 +427,11 @@ namespace cms::soa { cms::soa::SoAAccessType::constAccess>::template Alignment:: \ template RestrictQualifier::ParamReturnType \ LOCAL_NAME(size_type index) const { \ - return typename cms::soa::SoAAccessors:: \ + if constexpr (rangeChecking == cms::soa::RangeChecking::enabled) { \ + if (index >= elements_) \ + SOA_THROW_OUT_OF_RANGE("Out of range index in const " #LOCAL_NAME "(size_type index)") \ + } \ + return typename cms::soa::SoAAccessors:: \ template ColumnType::template AccessType< \ cms::soa::SoAAccessType::constAccess>::template Alignment:: \ template RestrictQualifier(BOOST_PP_CAT(LOCAL_NAME, Parameters_))(index); \ diff --git a/DataFormats/SoATemplate/test/SoALayoutAndView_t.cu b/DataFormats/SoATemplate/test/SoALayoutAndView_t.cu index 00e3465a1304a..34dda3bd7a803 100644 --- a/DataFormats/SoATemplate/test/SoALayoutAndView_t.cu +++ b/DataFormats/SoATemplate/test/SoALayoutAndView_t.cu @@ -101,7 +101,7 @@ using RangeCheckingHostDeviceView = // We expect to just run one thread. __global__ void rangeCheckKernel(RangeCheckingHostDeviceView soa) { - printf("About to fail range-check in CUDA thread: %d\n", threadIdx.x); + printf("About to fail range-check (operator[]) in CUDA thread: %d\n", threadIdx.x); [[maybe_unused]] auto si = soa[soa.metadata().size()]; printf("Fail: range-check failure should have stopped the kernel.\n"); } @@ -250,10 +250,23 @@ int main(void) { soa1viewRangeChecking(h_soahdLayout); // This should throw an exception [[maybe_unused]] auto si = soa1viewRangeChecking[soa1viewRangeChecking.metadata().size()]; - std::cout << "Fail: expected range-check exception not caught on the host." << std::endl; + std::cout << "Fail: expected range-check exception (operator[]) not caught on the host." << std::endl; assert(false); } catch (const std::out_of_range&) { - std::cout << "Pass: expected range-check exception successfully caught on the host." << std::endl; + std::cout << "Pass: expected range-check exception (operator[]) successfully caught on the host." << std::endl; + } + + try { + // Get a view like the default, except for range checking + SoAHostDeviceLayout::ViewTemplate + soa1viewRangeChecking(h_soahdLayout); + // This should throw an exception + [[maybe_unused]] auto si = soa1viewRangeChecking[soa1viewRangeChecking.metadata().size()]; + std::cout << "Fail: expected range-check exception (view-level index access) not caught on the host." << std::endl; + assert(false); + } catch (const std::out_of_range&) { + std::cout << "Pass: expected range-check exception (view-level index access) successfully caught on the host." + << std::endl; } // Validation of range checking in a kernel diff --git a/HeterogeneousCore/AlpakaTest/plugins/TestAlpakaAnalyzer.cc b/HeterogeneousCore/AlpakaTest/plugins/TestAlpakaAnalyzer.cc index 1058af9c9e77e..2c6ed46c5d282 100644 --- a/HeterogeneousCore/AlpakaTest/plugins/TestAlpakaAnalyzer.cc +++ b/HeterogeneousCore/AlpakaTest/plugins/TestAlpakaAnalyzer.cc @@ -48,6 +48,29 @@ namespace { column.print(out); return out; } + + template + void checkViewAddresses(T const& view) { + assert(view.metadata().addressOf_x() == view.x()); + assert(view.metadata().addressOf_x() == &view.x(0)); + assert(view.metadata().addressOf_x() == &view[0].x()); + assert(view.metadata().addressOf_y() == view.y()); + assert(view.metadata().addressOf_y() == &view.y(0)); + assert(view.metadata().addressOf_y() == &view[0].y()); + assert(view.metadata().addressOf_z() == view.z()); + assert(view.metadata().addressOf_z() == &view.z(0)); + assert(view.metadata().addressOf_z() == &view[0].z()); + assert(view.metadata().addressOf_id() == view.id()); + assert(view.metadata().addressOf_id() == &view.id(0)); + assert(view.metadata().addressOf_id() == &view[0].id()); + assert(view.metadata().addressOf_m() == view.m()); + assert(view.metadata().addressOf_m() == &view.m(0).coeffRef(0, 0)); + assert(view.metadata().addressOf_m() == &view[0].m().coeffRef(0, 0)); + assert(view.metadata().addressOf_r() == &view.r()); + //assert(view.metadata().addressOf_r() == &view.r(0)); // cannot access a scalar with an index + //assert(view.metadata().addressOf_r() == &view[0].r()); // cannot access a scalar via a SoA row-like accessor + } + } // namespace class TestAlpakaAnalyzer : public edm::stream::EDAnalyzer<> { @@ -58,6 +81,8 @@ class TestAlpakaAnalyzer : public edm::stream::EDAnalyzer<> { void analyze(edm::Event const& event, edm::EventSetup const&) override { portabletest::TestHostCollection const& product = event.get(token_); auto const& view = product.const_view(); + auto& mview = product.view(); + auto const& cmview = product.view(); { edm::LogInfo msg("TestAlpakaAnalyzer"); @@ -88,24 +113,9 @@ class TestAlpakaAnalyzer : public edm::stream::EDAnalyzer<> { reinterpret_cast(view.metadata().addressOf_r()); } - assert(view.metadata().addressOf_x() == view.x()); - assert(view.metadata().addressOf_x() == &view.x(0)); - assert(view.metadata().addressOf_x() == &view[0].x()); - assert(view.metadata().addressOf_y() == view.y()); - assert(view.metadata().addressOf_y() == &view.y(0)); - assert(view.metadata().addressOf_y() == &view[0].y()); - assert(view.metadata().addressOf_z() == view.z()); - assert(view.metadata().addressOf_z() == &view.z(0)); - assert(view.metadata().addressOf_z() == &view[0].z()); - assert(view.metadata().addressOf_id() == view.id()); - assert(view.metadata().addressOf_id() == &view.id(0)); - assert(view.metadata().addressOf_id() == &view[0].id()); - assert(view.metadata().addressOf_m() == view.m()); - assert(view.metadata().addressOf_m() == &view.m(0).coeffRef(0, 0)); - assert(view.metadata().addressOf_m() == &view[0].m().coeffRef(0, 0)); - assert(view.metadata().addressOf_r() == &view.r()); - //assert(view.metadata().addressOf_r() == &view.r(0)); // cannot access a scalar with an index - //assert(view.metadata().addressOf_r() == &view[0].r()); // cannot access a scalar via a SoA row-like accessor + checkViewAddresses(view); + checkViewAddresses(mview); + checkViewAddresses(cmview); const portabletest::Matrix matrix{{1, 2, 3, 4, 5, 6}, {2, 4, 6, 8, 10, 12}, {3, 6, 9, 12, 15, 18}}; assert(view.r() == 1.);