From 04882f8d6c5faefe04a1796e48c64f38f98f3e0d Mon Sep 17 00:00:00 2001 From: Christian Trott Date: Fri, 4 Oct 2024 10:50:51 -0600 Subject: [PATCH] Update mdspan corresponding to mdspan PR $360 --- .../experimental/__p0009_bits/layout_left.hpp | 3 + .../__p0009_bits/layout_right.hpp | 3 + .../__p0009_bits/layout_stride.hpp | 36 +++---- .../experimental/__p0009_bits/utility.hpp | 100 ++++++++++++++++++ .../__p2630_bits/submdspan_extents.hpp | 77 ++++++++++++-- .../__p2630_bits/submdspan_mapping.hpp | 55 +++++----- .../__p2642_bits/layout_padded.hpp | 9 +- .../__p2642_bits/layout_padded_fwd.hpp | 6 ++ 8 files changed, 234 insertions(+), 55 deletions(-) diff --git a/tpls/mdspan/include/experimental/__p0009_bits/layout_left.hpp b/tpls/mdspan/include/experimental/__p0009_bits/layout_left.hpp index 222fba7aa04..ed8aae020b6 100644 --- a/tpls/mdspan/include/experimental/__p0009_bits/layout_left.hpp +++ b/tpls/mdspan/include/experimental/__p0009_bits/layout_left.hpp @@ -237,10 +237,12 @@ class layout_left::mapping { // Not really public, but currently needed to implement fully constexpr useable submdspan: template + MDSPAN_INLINE_FUNCTION constexpr index_type __get_stride(MDSPAN_IMPL_STANDARD_NAMESPACE::extents,std::integer_sequence) const { return _MDSPAN_FOLD_TIMES_RIGHT((Idx():1),1); } template + MDSPAN_INLINE_FUNCTION constexpr index_type __stride() const noexcept { return __get_stride(__extents, std::make_index_sequence()); } @@ -255,6 +257,7 @@ class layout_left::mapping { SliceSpecifiers... slices) const; template + MDSPAN_INLINE_FUNCTION friend constexpr auto submdspan_mapping( const mapping& src, SliceSpecifiers... slices) { return src.submdspan_mapping_impl(slices...); diff --git a/tpls/mdspan/include/experimental/__p0009_bits/layout_right.hpp b/tpls/mdspan/include/experimental/__p0009_bits/layout_right.hpp index 284569f6533..26115e7a340 100644 --- a/tpls/mdspan/include/experimental/__p0009_bits/layout_right.hpp +++ b/tpls/mdspan/include/experimental/__p0009_bits/layout_right.hpp @@ -234,10 +234,12 @@ class layout_right::mapping { // Not really public, but currently needed to implement fully constexpr useable submdspan: template + MDSPAN_INLINE_FUNCTION constexpr index_type __get_stride(MDSPAN_IMPL_STANDARD_NAMESPACE::extents,std::integer_sequence) const { return _MDSPAN_FOLD_TIMES_RIGHT((Idx>N? __extents.template __extent():1),1); } template + MDSPAN_INLINE_FUNCTION constexpr index_type __stride() const noexcept { return __get_stride(__extents, std::make_index_sequence()); } @@ -252,6 +254,7 @@ class layout_right::mapping { SliceSpecifiers... slices) const; template + MDSPAN_INLINE_FUNCTION friend constexpr auto submdspan_mapping( const mapping& src, SliceSpecifiers... slices) { return src.submdspan_mapping_impl(slices...); diff --git a/tpls/mdspan/include/experimental/__p0009_bits/layout_stride.hpp b/tpls/mdspan/include/experimental/__p0009_bits/layout_stride.hpp index d6cdad2ab23..47ef2682d94 100644 --- a/tpls/mdspan/include/experimental/__p0009_bits/layout_stride.hpp +++ b/tpls/mdspan/include/experimental/__p0009_bits/layout_stride.hpp @@ -197,28 +197,22 @@ struct layout_stride { } template - MDSPAN_INLINE_FUNCTION static constexpr const __strides_storage_t fill_strides(const std::array& s) { return __strides_storage_t{static_cast(s[Idxs])...}; } MDSPAN_TEMPLATE_REQUIRES( class IntegralType, - // The is_convertible condition is added to make sfinae valid - // the extents_type::rank() > 0 is added to avoid use of non-standard zero length c-array - (std::is_convertible::value && (extents_type::rank() > 0)) + (std::is_convertible::value) ) MDSPAN_INLINE_FUNCTION - // despite the requirement some compilers still complain about zero length array during parsing - // making it length 1 now, but since the thing can't be instantiated due to requirement the actual - // instantiation of strides_storage will not fail despite mismatching length + // Need to avoid zero length c-array static constexpr const __strides_storage_t fill_strides(mdspan_non_standard_tag, const IntegralType (&s)[extents_type::rank()>0?extents_type::rank():1]) { return __strides_storage_t{static_cast(s[Idxs])...}; } #ifdef __cpp_lib_span template - MDSPAN_INLINE_FUNCTION static constexpr const __strides_storage_t fill_strides(const std::span& s) { return __strides_storage_t{static_cast(s[Idxs])...}; } @@ -242,10 +236,13 @@ struct layout_stride { // Can't use defaulted parameter in the __deduction_workaround template because of a bug in MSVC warning C4348. using __impl = __deduction_workaround>; + MDSPAN_FUNCTION static constexpr __strides_storage_t strides_storage(detail::with_rank<0>) { return {}; } + template + MDSPAN_FUNCTION static constexpr __strides_storage_t strides_storage(detail::with_rank) { __strides_storage_t s{}; @@ -273,7 +270,7 @@ struct layout_stride { //-------------------------------------------------------------------------------- - MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr mapping() noexcept + MDSPAN_INLINE_FUNCTION constexpr mapping() noexcept #if defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS) : __members{ #else @@ -299,7 +296,6 @@ struct layout_stride { _MDSPAN_TRAIT(std::is_nothrow_constructible, typename Extents::index_type, const std::remove_const_t&) ) ) - MDSPAN_INLINE_FUNCTION constexpr mapping( extents_type const& e, @@ -333,8 +329,7 @@ struct layout_stride { // MSVC 19.32 does not like using index_type here, requires the typename Extents::index_type // error C2641: cannot deduce template arguments for 'MDSPAN_IMPL_STANDARD_NAMESPACE::layout_stride::mapping' _MDSPAN_TRAIT(std::is_convertible, const std::remove_const_t&, typename Extents::index_type) && - _MDSPAN_TRAIT(std::is_nothrow_constructible, typename Extents::index_type, const std::remove_const_t&) && - (Extents::rank() > 0) + _MDSPAN_TRAIT(std::is_nothrow_constructible, typename Extents::index_type, const std::remove_const_t&) ) ) MDSPAN_INLINE_FUNCTION @@ -342,10 +337,8 @@ struct layout_stride { mapping( mdspan_non_standard_tag, extents_type const& e, - // despite the requirement some compilers still complain about zero length array during parsing - // making it length 1 now, but since the thing can't be instantiated due to requirement the actual - // instantiation of strides_storage will not fail despite mismatching length - IntegralTypes (&s)[extents_type::rank()>0?extents_type::rank():1] + // Need to avoid zero-length c-array + const IntegralTypes (&s)[extents_type::rank()>0?extents_type::rank():1] ) noexcept #if defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS) : __members{ @@ -379,7 +372,6 @@ struct layout_stride { _MDSPAN_TRAIT(std::is_nothrow_constructible, typename Extents::index_type, const std::remove_const_t&) ) ) - MDSPAN_INLINE_FUNCTION constexpr mapping( extents_type const& e, @@ -476,7 +468,8 @@ struct layout_stride { MDSPAN_INLINE_FUNCTION constexpr index_type required_span_size() const noexcept { index_type span_size = 1; - for(unsigned r = 0; r < extents_type::rank(); r++) { + // using int here to avoid warning about pointless comparison to 0 + for(int r = 0; r < static_cast(extents_type::rank()); r++) { // Return early if any of the extents are zero if(extents().extent(r)==0) return 0; span_size += ( static_cast(extents().extent(r) - 1 ) * __strides_storage()[r]); @@ -509,15 +502,18 @@ struct layout_stride { MDSPAN_INLINE_FUNCTION static constexpr bool is_unique() noexcept { return true; } private: + MDSPAN_INLINE_FUNCTION constexpr bool exhaustive_for_nonzero_span_size() const { return required_span_size() == __get_size(extents(), std::make_index_sequence()); } + MDSPAN_INLINE_FUNCTION constexpr bool is_exhaustive_impl(detail::with_rank<0>) const { return true; } + MDSPAN_INLINE_FUNCTION constexpr bool is_exhaustive_impl(detail::with_rank<1>) const { if (required_span_size() != static_cast(0)) { @@ -526,6 +522,7 @@ struct layout_stride { return stride(0) == 1; } template + MDSPAN_INLINE_FUNCTION constexpr bool is_exhaustive_impl(detail::with_rank) const { if (required_span_size() != static_cast(0)) { @@ -627,6 +624,7 @@ struct layout_stride { SliceSpecifiers... slices) const; template + MDSPAN_INLINE_FUNCTION friend constexpr auto submdspan_mapping( const mapping& src, SliceSpecifiers... slices) { return src.submdspan_mapping_impl(slices...); @@ -637,10 +635,12 @@ struct layout_stride { namespace detail { template +MDSPAN_INLINE_FUNCTION constexpr void validate_strides(with_rank<0>, Layout, const Extents&, const Mapping&) {} template +MDSPAN_INLINE_FUNCTION constexpr void validate_strides(with_rank, Layout, const Extents& ext, const Mapping& other) { static_assert(std::is_same::value && diff --git a/tpls/mdspan/include/experimental/__p0009_bits/utility.hpp b/tpls/mdspan/include/experimental/__p0009_bits/utility.hpp index e690cd6939b..f7f39d6024e 100644 --- a/tpls/mdspan/include/experimental/__p0009_bits/utility.hpp +++ b/tpls/mdspan/include/experimental/__p0009_bits/utility.hpp @@ -2,6 +2,8 @@ #include #include +#include +#include namespace MDSPAN_IMPL_STANDARD_NAMESPACE { namespace detail { @@ -64,6 +66,104 @@ constexpr struct } } stride; +// same as std::integral_constant but with __host__ __device__ annotations on +// the implicit conversion function and the call operator +template +struct integral_constant { + using value_type = T; + using type = integral_constant; + + static constexpr T value = v; + + MDSPAN_INLINE_FUNCTION_DEFAULTED + constexpr integral_constant() = default; + + // These interop functions work, because other than the value_type operator + // everything of std::integral_constant works on device (defaulted functions) + MDSPAN_FUNCTION + constexpr integral_constant(std::integral_constant) {}; + + MDSPAN_FUNCTION constexpr operator std::integral_constant() const noexcept { + return std::integral_constant{}; + } + + MDSPAN_FUNCTION constexpr operator value_type() const noexcept { + return value; + } + + MDSPAN_FUNCTION constexpr value_type operator()() const noexcept { + return value; + } +}; + +// The tuple implementation only comes in play when using capabilities +// such as submdspan which require C++17 anyway +#if MDSPAN_HAS_CXX_17 +template +struct tuple_member { + using type = T; + static constexpr size_t idx = Idx; + T val; + MDSPAN_FUNCTION constexpr T& get() { return val; } + MDSPAN_FUNCTION constexpr const T& get() const { return val; } +}; + +// A helper class which will be used via a fold expression to +// select the type with the correct Idx in a pack of tuple_member +template +struct tuple_idx_matcher { + using type = tuple_member; + template + MDSPAN_FUNCTION + constexpr auto operator | (Other v) const { + if constexpr (Idx == SearchIdx) { return *this; } + else { return v; } + } +}; + +template +struct tuple_impl; + +template +struct tuple_impl, Elements...>: public tuple_member ... { + + MDSPAN_FUNCTION + constexpr tuple_impl(Elements ... vals):tuple_member{vals}... {} + + template + MDSPAN_FUNCTION + constexpr auto& get() { + using base_t = decltype((tuple_idx_matcher() | ...) ); + return base_t::type::get(); + } + template + MDSPAN_FUNCTION + constexpr const auto& get() const { + using base_t = decltype((tuple_idx_matcher() | ...) ); + return base_t::type::get(); + } +}; + +// A simple tuple-like class for representing slices internally and is compatible with device code +// This doesn't support type access since we don't need it +// This is not meant as an external API +template +struct tuple: public tuple_impl()), Elements...> { + MDSPAN_FUNCTION + constexpr tuple(Elements ... vals):tuple_impl()), Elements ...>(vals ...) {} +}; + +template +MDSPAN_FUNCTION +constexpr auto& get(tuple& vals) { return vals.template get(); } + +template +MDSPAN_FUNCTION +constexpr const auto& get(const tuple& vals) { return vals.template get(); } + +template +tuple(Elements ...) -> tuple; +#endif } // namespace detail constexpr struct mdspan_non_standard_tag { diff --git a/tpls/mdspan/include/experimental/__p2630_bits/submdspan_extents.hpp b/tpls/mdspan/include/experimental/__p2630_bits/submdspan_extents.hpp index 779a9316951..4fe5dc6e29a 100644 --- a/tpls/mdspan/include/experimental/__p2630_bits/submdspan_extents.hpp +++ b/tpls/mdspan/include/experimental/__p2630_bits/submdspan_extents.hpp @@ -16,10 +16,11 @@ #pragma once -#include #include #include "strided_slice.hpp" +#include "../__p0009_bits/utility.hpp" + namespace MDSPAN_IMPL_STANDARD_NAMESPACE { namespace detail { @@ -68,6 +69,12 @@ struct index_pair_like, IndexType> { std::is_convertible_v; }; +template +struct index_pair_like, IndexType> { + static constexpr bool value = std::is_convertible_v && + std::is_convertible_v; +}; + template struct index_pair_like, IndexType> { static constexpr bool value = std::is_convertible_v; @@ -88,10 +95,16 @@ constexpr Integral first_of(const Integral &i) { return i; } +template +MDSPAN_INLINE_FUNCTION +constexpr Integral first_of(const std::integral_constant&) { + return integral_constant(); +} + MDSPAN_INLINE_FUNCTION -constexpr std::integral_constant +constexpr integral_constant first_of(const ::MDSPAN_IMPL_STANDARD_NAMESPACE::full_extent_t &) { - return std::integral_constant(); + return integral_constant(); } MDSPAN_TEMPLATE_REQUIRES( @@ -100,7 +113,24 @@ MDSPAN_TEMPLATE_REQUIRES( ) MDSPAN_INLINE_FUNCTION constexpr auto first_of(const Slice &i) { - return std::get<0>(i); + return get<0>(i); +} + +MDSPAN_TEMPLATE_REQUIRES( + class IdxT1, class IdxT2, + /* requires */ (index_pair_like, size_t>::value) + ) +constexpr auto first_of(const std::tuple& i) { + return get<0>(i); +} + +MDSPAN_TEMPLATE_REQUIRES( + class IdxT1, class IdxT2, + /* requires */ (index_pair_like, size_t>::value) + ) +MDSPAN_INLINE_FUNCTION +constexpr auto first_of(const std::pair& i) { + return i.first; } template @@ -137,7 +167,24 @@ MDSPAN_TEMPLATE_REQUIRES( MDSPAN_INLINE_FUNCTION constexpr auto last_of(std::integral_constant, const Extents &, const Slice &i) { - return std::get<1>(i); + return get<1>(i); +} + +MDSPAN_TEMPLATE_REQUIRES( + size_t k, class Extents, class IdxT1, class IdxT2, + /* requires */ (index_pair_like, size_t>::value) + ) +constexpr auto last_of(std::integral_constant, const Extents &, const std::tuple& i) { + return get<1>(i); +} + +MDSPAN_TEMPLATE_REQUIRES( + size_t k, class Extents, class IdxT1, class IdxT2, + /* requires */ (index_pair_like, size_t>::value) + ) +MDSPAN_INLINE_FUNCTION +constexpr auto last_of(std::integral_constant, const Extents &, const std::pair& i) { + return i.second; } template @@ -172,7 +219,7 @@ constexpr auto last_of(std::integral_constant, const Extents &ext, if constexpr (Extents::static_extent(k) == dynamic_extent) { return ext.extent(k); } else { - return std::integral_constant(); + return integral_constant(); } #if defined(__NVCC__) && !defined(__CUDA_ARCH__) && defined(__GNUC__) // Even with CUDA_ARCH protection this thing warns about calling host function @@ -204,7 +251,7 @@ last_of(std::integral_constant, const Extents &, template MDSPAN_INLINE_FUNCTION constexpr auto stride_of(const T &) { - return std::integral_constant(); + return integral_constant(); } template @@ -227,7 +274,7 @@ constexpr auto divide(const std::integral_constant &, const std::integral_constant &) { // cutting short division by zero // this is used for strided_slice with zero extent/stride - return std::integral_constant(); + return integral_constant(); } // multiply which can deal with integral constant preservation @@ -241,7 +288,7 @@ template MDSPAN_INLINE_FUNCTION constexpr auto multiply(const std::integral_constant &, const std::integral_constant &) { - return std::integral_constant(); + return integral_constant(); } // compute new static extent from range, preserving static knowledge @@ -255,6 +302,12 @@ struct StaticExtentFromRange, constexpr static size_t value = val1 - val0; }; +template +struct StaticExtentFromRange, + integral_constant> { + constexpr static size_t value = val1 - val0; +}; + // compute new static extent from strided_slice, preserving static // knowledge template struct StaticExtentFromStridedRange { @@ -267,6 +320,12 @@ struct StaticExtentFromStridedRange, constexpr static size_t value = val0 > 0 ? 1 + (val0 - 1) / val1 : 0; }; +template +struct StaticExtentFromStridedRange, + integral_constant> { + constexpr static size_t value = val0 > 0 ? 1 + (val0 - 1) / val1 : 0; +}; + // creates new extents through recursive calls to next_extent member function // next_extent has different overloads for different types of stride specifiers template diff --git a/tpls/mdspan/include/experimental/__p2630_bits/submdspan_mapping.hpp b/tpls/mdspan/include/experimental/__p2630_bits/submdspan_mapping.hpp index 69762e44e2c..2a2cdf76b92 100644 --- a/tpls/mdspan/include/experimental/__p2630_bits/submdspan_mapping.hpp +++ b/tpls/mdspan/include/experimental/__p2630_bits/submdspan_mapping.hpp @@ -17,9 +17,9 @@ #pragma once #include -#include #include #include // index_sequence +#include "../__p0009_bits/utility.hpp" // Suppress spurious warning with NVCC about no return statement. // This is a known issue in NVCC and NVC++ @@ -51,6 +51,7 @@ template struct submdspan_mapping_result { }; namespace detail { + // We use const Slice& and not Slice&& because the various // submdspan_mapping_impl overloads use their slices arguments // multiple times. This makes perfect forwarding not useful, but we @@ -84,14 +85,20 @@ any_slice_out_of_bounds(const extents &exts, } // constructs sub strides +template +struct sub_strides +{ + T values[N > 0 ? N : 1]; +}; + template MDSPAN_INLINE_FUNCTION constexpr auto construct_sub_strides( const SrcMapping &src_mapping, std::index_sequence, - const std::tuple &slices_stride_factor) { + const MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple &slices_stride_factor) { using index_type = typename SrcMapping::index_type; - return std::array{ + return sub_strides{{ (static_cast(src_mapping.stride(InvMapIdxs)) * - static_cast(std::get(slices_stride_factor)))...}; + static_cast(get(slices_stride_factor)))...}}; } template @@ -240,7 +247,7 @@ layout_left::mapping::submdspan_mapping_impl( auto inv_map = detail::inv_map_rank(std::integral_constant(), std::index_sequence<>(), slices...); return submdspan_mapping_result { - dst_mapping_t(dst_ext, + dst_mapping_t(mdspan_non_standard, dst_ext, detail::construct_sub_strides( *this, inv_map, // HIP needs deduction guides to have markups so we need to be explicit @@ -248,10 +255,10 @@ layout_left::mapping::submdspan_mapping_impl( // the issue But Clang-CUDA also doesn't accept the use of deduction guide so // disable it for CUDA altogether #if defined(_MDSPAN_HAS_HIP) || defined(_MDSPAN_HAS_CUDA) - std::tuple{ - detail::stride_of(slices)...})), + detail::tuple{ + detail::stride_of(slices)...}).values), #else - std::tuple{detail::stride_of(slices)...})), + detail::tuple{detail::stride_of(slices)...}).values), #endif offset }; @@ -318,7 +325,7 @@ MDSPAN_IMPL_PROPOSED_NAMESPACE::layout_left_padded::mapping(), slices...); using dst_mapping_t = typename layout_stride::template mapping; return submdspan_mapping_result { - dst_mapping_t(dst_ext, + dst_mapping_t(mdspan_non_standard, dst_ext, MDSPAN_IMPL_STANDARD_NAMESPACE::detail::construct_sub_strides( *this, inv_map, // HIP needs deduction guides to have markups so we need to be explicit @@ -326,10 +333,10 @@ MDSPAN_IMPL_PROPOSED_NAMESPACE::layout_left_padded::mapping{ - MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices)...})), + MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple{ + MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices)...}).values), #else - std::tuple{MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices)...})), + MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple{MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices)...}).values), #endif offset }; @@ -473,7 +480,7 @@ layout_right::mapping::submdspan_mapping_impl( auto inv_map = detail::inv_map_rank(std::integral_constant(), std::index_sequence<>(), slices...); return submdspan_mapping_result { - dst_mapping_t(dst_ext, + dst_mapping_t(mdspan_non_standard, dst_ext, detail::construct_sub_strides( *this, inv_map, // HIP needs deduction guides to have markups so we need to be explicit @@ -481,10 +488,10 @@ layout_right::mapping::submdspan_mapping_impl( // the issue But Clang-CUDA also doesn't accept the use of deduction guide so // disable it for CUDA altogether #if defined(_MDSPAN_HAS_HIP) || defined(_MDSPAN_HAS_CUDA) - std::tuple{ - detail::stride_of(slices)...})), + MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple{ + detail::stride_of(slices)...}).values), #else - std::tuple{detail::stride_of(slices)...})), + MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple{detail::stride_of(slices)...}).values), #endif offset }; @@ -543,7 +550,7 @@ MDSPAN_IMPL_PROPOSED_NAMESPACE::layout_right_padded::mapping(), slices...); using dst_mapping_t = typename layout_stride::template mapping; return submdspan_mapping_result { - dst_mapping_t(dst_ext, + dst_mapping_t(mdspan_non_standard, dst_ext, MDSPAN_IMPL_STANDARD_NAMESPACE::detail::construct_sub_strides( *this, inv_map, // HIP needs deduction guides to have markups so we need to be explicit @@ -551,10 +558,10 @@ MDSPAN_IMPL_PROPOSED_NAMESPACE::layout_right_padded::mapping{ - MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices)...})), + MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple{ + MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices)...}).values), #else - std::tuple{MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices)...})), + MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple{MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices)...}).values), #endif offset }; @@ -591,7 +598,7 @@ layout_stride::mapping::submdspan_mapping_impl( : this->operator()(detail::first_of(slices)...)); return submdspan_mapping_result { - dst_mapping_t(dst_ext, + dst_mapping_t(mdspan_non_standard, dst_ext, detail::construct_sub_strides( *this, inv_map, // HIP needs deduction guides to have markups so we need to be explicit @@ -600,10 +607,10 @@ layout_stride::mapping::submdspan_mapping_impl( #if defined(_MDSPAN_HAS_HIP) || \ (defined(__NVCC__) && \ (__CUDACC_VER_MAJOR__ * 100 + __CUDACC_VER_MINOR__ * 10) < 1120) - std::tuple( - detail::stride_of(slices)...))), + MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple( + detail::stride_of(slices)...).values)), #else - std::tuple(detail::stride_of(slices)...))), + MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple(detail::stride_of(slices)...)).values), #endif offset }; diff --git a/tpls/mdspan/include/experimental/__p2642_bits/layout_padded.hpp b/tpls/mdspan/include/experimental/__p2642_bits/layout_padded.hpp index 1502489adcd..a714090e448 100644 --- a/tpls/mdspan/include/experimental/__p2642_bits/layout_padded.hpp +++ b/tpls/mdspan/include/experimental/__p2642_bits/layout_padded.hpp @@ -23,6 +23,7 @@ #include "../__p0009_bits/layout_left.hpp" #include "../__p0009_bits/layout_right.hpp" #include "../__p0009_bits/layout_stride.hpp" +#include "../__p0009_bits/utility.hpp" namespace MDSPAN_IMPL_STANDARD_NAMESPACE { namespace MDSPAN_IMPL_PROPOSED_NAMESPACE { @@ -205,7 +206,7 @@ class layout_left_padded::mapping { public: #if !MDSPAN_HAS_CXX_20 || defined(__NVCC__) - MDSPAN_INLINE_FUNCTION_DEFAULTED + MDSPAN_INLINE_FUNCTION constexpr mapping() : mapping(extents_type{}) {} @@ -357,7 +358,7 @@ class layout_left_padded::mapping { return exts; } - MDSPAN_INLINE_FUNCTION constexpr std::array + constexpr std::array strides() const noexcept { if constexpr (extents_type::rank() == 0) { return {}; @@ -568,7 +569,7 @@ class layout_right_padded::mapping { public: #if !MDSPAN_HAS_CXX_20 || defined(__NVCC__) - MDSPAN_INLINE_FUNCTION_DEFAULTED + MDSPAN_INLINE_FUNCTION constexpr mapping() : mapping(extents_type{}) {} @@ -717,7 +718,7 @@ class layout_right_padded::mapping { return exts; } - MDSPAN_INLINE_FUNCTION constexpr std::array + constexpr std::array strides() const noexcept { if constexpr (extents_type::rank() == 0) { return {}; diff --git a/tpls/mdspan/include/experimental/__p2642_bits/layout_padded_fwd.hpp b/tpls/mdspan/include/experimental/__p2642_bits/layout_padded_fwd.hpp index 18daa28cc68..3f141ff08aa 100644 --- a/tpls/mdspan/include/experimental/__p2642_bits/layout_padded_fwd.hpp +++ b/tpls/mdspan/include/experimental/__p2642_bits/layout_padded_fwd.hpp @@ -85,12 +85,15 @@ struct is_layout_right_padded_mapping<_Mapping, template +MDSPAN_INLINE_FUNCTION constexpr void check_padded_layout_converting_constructor_mandates(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::with_rank<0>) {} template +MDSPAN_INLINE_FUNCTION constexpr void check_padded_layout_converting_constructor_mandates(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::with_rank<1>) {} template +MDSPAN_INLINE_FUNCTION constexpr void check_padded_layout_converting_constructor_mandates(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::with_rank) { using extents_type = typename _PaddedLayoutMappingType::extents_type; @@ -110,12 +113,15 @@ constexpr void check_padded_layout_converting_constructor_mandates(MDSPAN_IMPL_S } template +MDSPAN_INLINE_FUNCTION constexpr void check_padded_layout_converting_constructor_preconditions(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::with_rank<0>, const _OtherMapping&) {} template +MDSPAN_INLINE_FUNCTION constexpr void check_padded_layout_converting_constructor_preconditions(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::with_rank<1>, const _OtherMapping&) {} template +MDSPAN_INLINE_FUNCTION constexpr void check_padded_layout_converting_constructor_preconditions(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::with_rank, const _OtherMapping &other_mapping) { constexpr auto padded_stride_idx =