From 4b8f023248f4e466dc77084e55012897f466025d Mon Sep 17 00:00:00 2001 From: Eric Cano Date: Fri, 9 Jun 2023 17:58:35 +0200 Subject: [PATCH 1/4] Adds intercompatibility between view with and without range checking. This change allows providing a range checked to the code, presumably to debug without changing the PortableCollections (which always provide non-range checking views). Range and non-range checking variants can be contructed from one another. --- DataFormats/SoATemplate/interface/SoACommon.h | 8 +++ DataFormats/SoATemplate/interface/SoAView.h | 49 +++++++++++++++++++ .../test/SoALayoutAndView_t.hip.cc | 8 +++ 3 files changed, 65 insertions(+) diff --git a/DataFormats/SoATemplate/interface/SoACommon.h b/DataFormats/SoATemplate/interface/SoACommon.h index 8ab80fc88c697..a53dcf65f04b6 100644 --- a/DataFormats/SoATemplate/interface/SoACommon.h +++ b/DataFormats/SoATemplate/interface/SoACommon.h @@ -132,6 +132,8 @@ namespace cms::soa { return reinterpret_cast(addr) % alignment; } + TupleOrPointerType tupleOrPointer() { return addr_; } + public: // scalar or column ValueType const* addr_ = nullptr; @@ -166,6 +168,8 @@ namespace cms::soa { return reinterpret_cast(addr) % alignment; } + TupleOrPointerType tupleOrPointer() { return {addr_, stride_}; } + public: // address and stride ScalarType const* addr_ = nullptr; @@ -201,6 +205,8 @@ namespace cms::soa { return reinterpret_cast(addr) % alignment; } + TupleOrPointerType tupleOrPointer() { return addr_; } + public: // scalar or column ValueType* addr_ = nullptr; @@ -234,6 +240,8 @@ namespace cms::soa { return reinterpret_cast(addr) % alignment; } + TupleOrPointerType tupleOrPointer() { return {addr_, stride_}; } + public: // address and stride ScalarType* addr_ = nullptr; diff --git a/DataFormats/SoATemplate/interface/SoAView.h b/DataFormats/SoATemplate/interface/SoAView.h index 8cf0307c71ad0..3333eae34428e 100644 --- a/DataFormats/SoATemplate/interface/SoAView.h +++ b/DataFormats/SoATemplate/interface/SoAView.h @@ -231,6 +231,15 @@ namespace cms::soa { #define _DECLARE_VIEW_MEMBER_LIST(R, DATA, LAYOUT_MEMBER_NAME) \ BOOST_PP_EXPAND(_DECLARE_VIEW_MEMBER_LIST_IMPL LAYOUT_MEMBER_NAME) +/** + * Generator of view member list. + */ +#define _DECLARE_VIEW_OTHER_MEMBER_LIST_IMPL(LAYOUT, MEMBER, NAME) \ + (const_cast_SoAParametersImpl(other.BOOST_PP_CAT(NAME, Parameters_)).tupleOrPointer()) + +#define _DECLARE_VIEW_OTHER_MEMBER_LIST(R, DATA, LAYOUT_MEMBER_NAME) \ + BOOST_PP_EXPAND(_DECLARE_VIEW_OTHER_MEMBER_LIST_IMPL LAYOUT_MEMBER_NAME) + /** * Generator of member initializer for copy constructor. */ @@ -535,6 +544,9 @@ namespace cms::soa { template \ using SoAConstValueWithConf = cms::soa::SoAConstValue; \ \ + template \ + friend struct VIEW; \ + \ /** \ * Helper/friend class allowing SoA introspection. \ */ \ @@ -582,6 +594,23 @@ namespace cms::soa { VIEW(VIEW const&) = default; \ VIEW& operator=(VIEW const&) = default; \ \ + /* Copy constructor for other parameters */ \ + template \ + VIEW(VIEW const& other): base_type{other.elements_, \ + _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_OTHER_MEMBER_LIST, BOOST_PP_EMPTY(), VALUE_LIST) \ + } {} \ + /* Copy operator for other parameters */ \ + template \ + VIEW& operator=(VIEW const& other) { static_cast(*this) = static_cast(other); } \ + \ /* Movable */ \ VIEW(VIEW &&) = default; \ VIEW& operator=(VIEW &&) = default; \ @@ -673,6 +702,9 @@ namespace cms::soa { template \ friend struct VIEW; \ \ + template \ + friend struct CONST_VIEW; \ + \ /* 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. \ @@ -739,6 +771,23 @@ namespace cms::soa { CONST_VIEW(CONST_VIEW const&) = default; \ CONST_VIEW& operator=(CONST_VIEW const&) = default; \ \ + /* Copy constructor for other parameters */ \ + template \ + CONST_VIEW(CONST_VIEW const& other): CONST_VIEW{other.elements_, \ + _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_OTHER_MEMBER_LIST, BOOST_PP_EMPTY(), VALUE_LIST) \ + } {} \ + /* Copy operator for other parameters */ \ + template \ + CONST_VIEW& operator=(CONST_VIEW const& other) { *this = other; } \ + \ /* Movable */ \ CONST_VIEW(CONST_VIEW &&) = default; \ CONST_VIEW& operator=(CONST_VIEW &&) = default; \ diff --git a/DataFormats/SoATemplate/test/SoALayoutAndView_t.hip.cc b/DataFormats/SoATemplate/test/SoALayoutAndView_t.hip.cc index 1965d0757d358..b29c43775ed01 100644 --- a/DataFormats/SoATemplate/test/SoALayoutAndView_t.hip.cc +++ b/DataFormats/SoATemplate/test/SoALayoutAndView_t.hip.cc @@ -37,6 +37,8 @@ GENERATE_SOA_LAYOUT(SoAHostDeviceLayoutTemplate, using SoAHostDeviceLayout = SoAHostDeviceLayoutTemplate<>; using SoAHostDeviceView = SoAHostDeviceLayout::View; +using SoAHostDeviceRangeCheckingView = + SoAHostDeviceLayout::ViewTemplate; using SoAHostDeviceConstView = SoAHostDeviceLayout::ConstView; GENERATE_SOA_LAYOUT(SoADeviceOnlyLayoutTemplate, @@ -128,6 +130,12 @@ int main(void) { hipCheck(hipHostMalloc((void**)&h_buf, hostDeviceSize)); SoAHostDeviceLayout h_soahdLayout(h_buf, numElements); SoAHostDeviceView h_soahd(h_soahdLayout); + + // Validation of range checking variants initialization + SoAHostDeviceRangeCheckingView h_soahdrc(h_soahdLayout); + [[maybe_unused]] SoAHostDeviceRangeCheckingView h_soahdrc2 = h_soahdLayout; + [[maybe_unused]] SoAHostDeviceRangeCheckingView h_soahdrc3{h_soahd}; + [[maybe_unused]] SoAHostDeviceRangeCheckingView h_soahdrc4 = h_soahd; SoAHostDeviceConstView h_soahd_c(h_soahdLayout); // Alocate buffer, stores and views on the device (single, shared buffer). From 9d89dbd27a5215803b9373ce408deeea26584677 Mon Sep 17 00:00:00 2001 From: Eric Cano Date: Fri, 9 Jun 2023 18:22:36 +0200 Subject: [PATCH 2/4] Fix SoA range checking tests Positivty tests were missing as the indices are of a signed type (for ROOT compatibility). --- DataFormats/SoATemplate/interface/SoAView.h | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/DataFormats/SoATemplate/interface/SoAView.h b/DataFormats/SoATemplate/interface/SoAView.h index 3333eae34428e..f219bd137a0cc 100644 --- a/DataFormats/SoATemplate/interface/SoAView.h +++ b/DataFormats/SoATemplate/interface/SoAView.h @@ -399,7 +399,7 @@ namespace cms::soa { template RestrictQualifier::ParamReturnType \ LOCAL_NAME(size_type _soa_impl_index) { \ if constexpr (rangeChecking == cms::soa::RangeChecking::enabled) { \ - if (_soa_impl_index >= base_type::elements_) \ + if (_soa_impl_index >= base_type::elements_ or _soa_impl_index < 0) \ SOA_THROW_OUT_OF_RANGE("Out of range index in mutable " #LOCAL_NAME "(size_type index)") \ } \ return typename cms::soa::SoAAccessors:: \ @@ -437,7 +437,7 @@ namespace cms::soa { template RestrictQualifier::ParamReturnType \ LOCAL_NAME(size_type _soa_impl_index) const { \ if constexpr (rangeChecking == cms::soa::RangeChecking::enabled) { \ - if (_soa_impl_index >= elements_) \ + if (_soa_impl_index >= elements_ or _soa_impl_index < 0) \ SOA_THROW_OUT_OF_RANGE("Out of range index in const " #LOCAL_NAME "(size_type index)") \ } \ return typename cms::soa::SoAAccessors:: \ @@ -649,7 +649,7 @@ namespace cms::soa { SOA_HOST_DEVICE SOA_INLINE \ element operator[](size_type _soa_impl_index) { \ if constexpr (rangeChecking == cms::soa::RangeChecking::enabled) { \ - if (_soa_impl_index >= base_type::elements_) \ + if (_soa_impl_index >= base_type::elements_ or _soa_impl_index < 0) \ SOA_THROW_OUT_OF_RANGE("Out of range index in " #VIEW "::operator[]") \ } \ return element{_soa_impl_index, _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_ELEMENT_CONSTR_CALL, ~, VALUE_LIST)}; \ @@ -810,7 +810,7 @@ namespace cms::soa { SOA_HOST_DEVICE SOA_INLINE \ const_element operator[](size_type _soa_impl_index) const { \ if constexpr (rangeChecking == cms::soa::RangeChecking::enabled) { \ - if (_soa_impl_index >= elements_) \ + if (_soa_impl_index >= elements_ or _soa_impl_index < 0) \ SOA_THROW_OUT_OF_RANGE("Out of range index in " #CONST_VIEW "::operator[]") \ } \ return const_element{ \ From 1f4e998e835b6c0003cb24681c6ab7c62d5eca34 Mon Sep 17 00:00:00 2001 From: Eric Cano Date: Tue, 13 Jun 2023 14:26:01 +0200 Subject: [PATCH 3/4] Adds complete tests for range checking view, both layout and view initialized. --- .../SoATemplate/test/SoALayoutAndView_t.cu | 71 +++++++++++++------ .../test/SoALayoutAndView_t.hip.cc | 63 ++++++++++------ 2 files changed, 92 insertions(+), 42 deletions(-) diff --git a/DataFormats/SoATemplate/test/SoALayoutAndView_t.cu b/DataFormats/SoATemplate/test/SoALayoutAndView_t.cu index d8a92a56bf37a..0e30f8301f099 100644 --- a/DataFormats/SoATemplate/test/SoALayoutAndView_t.cu +++ b/DataFormats/SoATemplate/test/SoALayoutAndView_t.cu @@ -35,6 +35,8 @@ GENERATE_SOA_LAYOUT(SoAHostDeviceLayoutTemplate, using SoAHostDeviceLayout = SoAHostDeviceLayoutTemplate<>; using SoAHostDeviceView = SoAHostDeviceLayout::View; +using SoAHostDeviceRangeCheckingView = + SoAHostDeviceLayout::ViewTemplate; using SoAHostDeviceConstView = SoAHostDeviceLayout::ConstView; GENERATE_SOA_LAYOUT(SoADeviceOnlyLayoutTemplate, @@ -126,6 +128,12 @@ int main(void) { cudaCheck(cudaMallocHost(&h_buf, hostDeviceSize)); SoAHostDeviceLayout h_soahdLayout(h_buf, numElements); SoAHostDeviceView h_soahd(h_soahdLayout); + + // Validation of range checking variants initialization + SoAHostDeviceRangeCheckingView h_soahdrc(h_soahdLayout); + [[maybe_unused]] SoAHostDeviceRangeCheckingView h_soahdrc2 = h_soahdLayout; + SoAHostDeviceRangeCheckingView h_soahdrc3{h_soahd}; + [[maybe_unused]] SoAHostDeviceRangeCheckingView h_soahdrc4 = h_soahd; SoAHostDeviceConstView h_soahd_c(h_soahdLayout); // Alocate buffer, stores and views on the device (single, shared buffer). @@ -248,29 +256,50 @@ int main(void) { } } - // Validation of range checking - 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 (operator[]) not caught on the host." << std::endl; - assert(false); - } catch (const std::out_of_range&) { - std::cout << "Pass: expected range-check exception (operator[]) successfully caught on the host." << std::endl; + { + // Get a view like the default, except for range checking (direct initialization from layout) + SoAHostDeviceRangeCheckingView soa1viewRangeChecking(h_soahdLayout); + try { + [[maybe_unused]] auto si = soa1viewRangeChecking[soa1viewRangeChecking.metadata().size()]; + std::cout << "Fail: expected range-check exception (view-level index access) not caught on the host (overflow)." + << std::endl; + assert(false); + } catch (const std::out_of_range&) { + } + try { + [[maybe_unused]] auto si = soa1viewRangeChecking[-1]; + std::cout << "Fail: expected range-check exception (view-level index access) not caught on the host (underflow)." + << std::endl; + assert(false); + } catch (const std::out_of_range&) { + } + [[maybe_unused]] auto si = soa1viewRangeChecking[soa1viewRangeChecking.metadata().size() - 1]; + [[maybe_unused]] auto si2 = soa1viewRangeChecking[0]; + std::cout << "Pass: expected range-check exceptions (view-level index access) successfully caught on the host " + "(layout initialization)." + << 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." + { + // Validation of view initialized range checking view initialization + try { + [[maybe_unused]] auto si = h_soahdrc3[h_soahdrc3.metadata().size()]; + std::cout << "Fail: expected range-check exception (view-level index access) not caught on the host (overflow)." + << std::endl; + assert(false); + } catch (const std::out_of_range&) { + } + try { + [[maybe_unused]] auto si = h_soahdrc3[-1]; + std::cout << "Fail: expected range-check exception (view-level index access) not caught on the host (underflow)." + << std::endl; + assert(false); + } catch (const std::out_of_range&) { + } + [[maybe_unused]] auto si = h_soahdrc3[h_soahdrc3.metadata().size() - 1]; + [[maybe_unused]] auto si2 = h_soahdrc3[0]; + std::cout << "Pass: expected range-check exceptions (view-level index access) successfully caught on the host " + "(view initialization)." << std::endl; } diff --git a/DataFormats/SoATemplate/test/SoALayoutAndView_t.hip.cc b/DataFormats/SoATemplate/test/SoALayoutAndView_t.hip.cc index b29c43775ed01..6a041248f3995 100644 --- a/DataFormats/SoATemplate/test/SoALayoutAndView_t.hip.cc +++ b/DataFormats/SoATemplate/test/SoALayoutAndView_t.hip.cc @@ -258,29 +258,50 @@ int main(void) { } } - // Validation of range checking - 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 (operator[]) not caught on the host." << std::endl; - assert(false); - } catch (const std::out_of_range&) { - std::cout << "Pass: expected range-check exception (operator[]) successfully caught on the host." << std::endl; + { + // Get a view like the default, except for range checking (direct initialization from layout) + SoAHostDeviceRangeCheckingView soa1viewRangeChecking(h_soahdLayout); + try { + [[maybe_unused]] auto si = soa1viewRangeChecking[soa1viewRangeChecking.metadata().size()]; + std::cout << "Fail: expected range-check exception (view-level index access) not caught on the host (overflow)." + << std::endl; + assert(false); + } catch (const std::out_of_range&) { + } + try { + [[maybe_unused]] auto si = soa1viewRangeChecking[-1]; + std::cout << "Fail: expected range-check exception (view-level index access) not caught on the host (underflow)." + << std::endl; + assert(false); + } catch (const std::out_of_range&) { + } + [[maybe_unused]] auto si = soa1viewRangeChecking[soa1viewRangeChecking.metadata().size() - 1]; + [[maybe_unused]] auto si2 = soa1viewRangeChecking[0]; + std::cout << "Pass: expected range-check exceptions (view-level index access) successfully caught on the host " + "(layout initialization)." + << 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." + { + // Validation of view initialized range checking view initialization + try { + [[maybe_unused]] auto si = h_soahdrc3[h_soahdrc3.metadata().size()]; + std::cout << "Fail: expected range-check exception (view-level index access) not caught on the host (overflow)." + << std::endl; + assert(false); + } catch (const std::out_of_range&) { + } + try { + [[maybe_unused]] auto si = h_soahdrc3[-1]; + std::cout << "Fail: expected range-check exception (view-level index access) not caught on the host (underflow)." + << std::endl; + assert(false); + } catch (const std::out_of_range&) { + } + [[maybe_unused]] auto si = h_soahdrc3[h_soahdrc3.metadata().size() - 1]; + [[maybe_unused]] auto si2 = h_soahdrc3[0]; + std::cout << "Pass: expected range-check exceptions (view-level index access) successfully caught on the host " + "(view initialization)." << std::endl; } From 2fe74705b4ab9590824299e11c39a08a1bb7a823 Mon Sep 17 00:00:00 2001 From: Eric Cano Date: Tue, 13 Jun 2023 15:01:02 +0200 Subject: [PATCH 4/4] Document the CMS_SOA_BYTE_SIZE_TYPE macro Explain that using a typedef as a template parameter causes the ROOT dicionary generation to fail. --- DataFormats/SoATemplate/interface/SoACommon.h | 2 ++ 1 file changed, 2 insertions(+) diff --git a/DataFormats/SoATemplate/interface/SoACommon.h b/DataFormats/SoATemplate/interface/SoACommon.h index a53dcf65f04b6..49f5049377b44 100644 --- a/DataFormats/SoATemplate/interface/SoACommon.h +++ b/DataFormats/SoATemplate/interface/SoACommon.h @@ -52,6 +52,8 @@ #define _VALUE_TYPE_EIGEN_COLUMN 2 /* The size type need to be "hardcoded" in the template parameters for classes serialized by ROOT */ +/* In practice, using a typedef as a template parameter to the Layout or its ViewTemplateFreeParams member + * declaration fails ROOT dictionary generation. */ #define CMS_SOA_BYTE_SIZE_TYPE std::size_t namespace cms::soa {