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 70f344c66d330..220ea73176d71 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_; }; @@ -584,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: @@ -592,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: @@ -606,25 +592,27 @@ 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*; - SOA_HOST_DEVICE SOA_INLINE T operator()(size_type index) const { return params_.addr_[index]; } + using ParamReturnType = const T&; + SOA_HOST_DEVICE SOA_INLINE T const& operator()(size_type index) const { return params_.addr_[index]; } private: SoAConstParametersImpl params_; }; // 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."); } @@ -634,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 operator()() const { return *params_.addr_; } - using NoParamReturnType = T; + 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."); } @@ -649,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 const typename T::Scalar* operator()() const { return params_.addr_; } - using NoParamReturnType = typename T::Scalar*; - //SOA_HOST_DEVICE SOA_INLINE T operator()(size_type index) const { return params_.addr_[index]; } + SOA_HOST_DEVICE SOA_INLINE typename T::Scalar const* operator()() const { return params_.addr_; } + using NoParamReturnType = typename T::Scalar const*; + using ParamReturnType = typename SoAValue::CMapType; + SOA_HOST_DEVICE SOA_INLINE ParamReturnType operator()(size_type index) const { + return SoAConstValue(index, params_)(); + } private: SoAConstParametersImpl params_; @@ -681,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/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 945626db03851..0099587a98180 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,30 @@ 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) { \ + 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) { \ + 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>(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 @@ -388,15 +410,31 @@ 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>::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 LOCAL_NAME(size_type index) const { \ + 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 { \ + 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>(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/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 c64e91cb9d0dc..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,6 +113,10 @@ class TestAlpakaAnalyzer : public edm::stream::EDAnalyzer<> { reinterpret_cast(view.metadata().addressOf_r()); } + 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.); for (int32_t i = 0; i < view.metadata().size(); ++i) {