diff --git a/include/experimental/__p0009_bits/layout_left.hpp b/include/experimental/__p0009_bits/layout_left.hpp
index 222fba7a..ed8aae02 100644
--- a/include/experimental/__p0009_bits/layout_left.hpp
+++ b/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<size_t N, class SizeType, size_t ... E, size_t ... Idx>
+    MDSPAN_INLINE_FUNCTION
     constexpr index_type __get_stride(MDSPAN_IMPL_STANDARD_NAMESPACE::extents<SizeType, E...>,std::integer_sequence<size_t, Idx...>) const {
       return _MDSPAN_FOLD_TIMES_RIGHT((Idx<N? __extents.template __extent<Idx>():1),1);
     }
     template<size_t N>
+    MDSPAN_INLINE_FUNCTION
     constexpr index_type __stride() const noexcept {
       return __get_stride<N>(__extents, std::make_index_sequence<extents_type::rank()>());
     }
@@ -255,6 +257,7 @@ class layout_left::mapping {
        SliceSpecifiers... slices) const;
 
    template<class... SliceSpecifiers>
+     MDSPAN_INLINE_FUNCTION
      friend constexpr auto submdspan_mapping(
        const mapping& src, SliceSpecifiers... slices) {
          return src.submdspan_mapping_impl(slices...);
diff --git a/include/experimental/__p0009_bits/layout_right.hpp b/include/experimental/__p0009_bits/layout_right.hpp
index 284569f6..26115e7a 100644
--- a/include/experimental/__p0009_bits/layout_right.hpp
+++ b/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<size_t N, class SizeType, size_t ... E, size_t ... Idx>
+    MDSPAN_INLINE_FUNCTION
     constexpr index_type __get_stride(MDSPAN_IMPL_STANDARD_NAMESPACE::extents<SizeType, E...>,std::integer_sequence<size_t, Idx...>) const {
       return _MDSPAN_FOLD_TIMES_RIGHT((Idx>N? __extents.template __extent<Idx>():1),1);
     }
     template<size_t N>
+    MDSPAN_INLINE_FUNCTION
     constexpr index_type __stride() const noexcept {
       return __get_stride<N>(__extents, std::make_index_sequence<extents_type::rank()>());
     }
@@ -252,6 +254,7 @@ class layout_right::mapping {
        SliceSpecifiers... slices) const;
 
    template<class... SliceSpecifiers>
+     MDSPAN_INLINE_FUNCTION
      friend constexpr auto submdspan_mapping(
        const mapping& src, SliceSpecifiers... slices) {
          return src.submdspan_mapping_impl(slices...);
diff --git a/include/experimental/__p0009_bits/layout_stride.hpp b/include/experimental/__p0009_bits/layout_stride.hpp
index d6cdad2a..47ef2682 100644
--- a/include/experimental/__p0009_bits/layout_stride.hpp
+++ b/include/experimental/__p0009_bits/layout_stride.hpp
@@ -197,28 +197,22 @@ struct layout_stride {
       }
 
       template<class IntegralType>
-      MDSPAN_INLINE_FUNCTION
       static constexpr const __strides_storage_t fill_strides(const std::array<IntegralType,extents_type::rank()>& s) {
         return __strides_storage_t{static_cast<index_type>(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<IntegralType, typename extents_type::index_type>::value && (extents_type::rank() > 0))
+        (std::is_convertible<IntegralType, typename extents_type::index_type>::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<index_type>(s[Idxs])...};
       }
 
 #ifdef __cpp_lib_span
       template<class IntegralType>
-      MDSPAN_INLINE_FUNCTION
       static constexpr const __strides_storage_t fill_strides(const std::span<IntegralType,extents_type::rank()>& s) {
         return __strides_storage_t{static_cast<index_type>(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<std::make_index_sequence<Extents::rank()>>;
 
+    MDSPAN_FUNCTION
     static constexpr __strides_storage_t strides_storage(detail::with_rank<0>) {
       return {};
     }
+
     template <std::size_t N>
+    MDSPAN_FUNCTION
     static constexpr __strides_storage_t strides_storage(detail::with_rank<N>) {
       __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<IntegralTypes>&)
       )
     )
-    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<IntegralTypes>&, typename Extents::index_type) &&
-        _MDSPAN_TRAIT(std::is_nothrow_constructible, typename Extents::index_type, const std::remove_const_t<IntegralTypes>&) &&
-        (Extents::rank() > 0)
+        _MDSPAN_TRAIT(std::is_nothrow_constructible, typename Extents::index_type, const std::remove_const_t<IntegralTypes>&)
       )
     )
     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<IntegralTypes>&)
       )
     )
-    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<int>(extents_type::rank()); r++) {
         // Return early if any of the extents are zero
         if(extents().extent(r)==0) return 0;
         span_size += ( static_cast<index_type>(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<extents_type::rank()>());
     }
 
+    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<index_type>(0)) {
@@ -526,6 +522,7 @@ struct layout_stride {
       return stride(0) == 1;
     }
     template <std::size_t N>
+    MDSPAN_INLINE_FUNCTION
     constexpr bool is_exhaustive_impl(detail::with_rank<N>) const
     {
       if (required_span_size() != static_cast<index_type>(0)) {
@@ -627,6 +624,7 @@ struct layout_stride {
        SliceSpecifiers... slices) const;
 
    template<class... SliceSpecifiers>
+    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 <class Layout, class Extents, class Mapping>
+MDSPAN_INLINE_FUNCTION
 constexpr void validate_strides(with_rank<0>, Layout, const Extents&, const Mapping&)
 {}
 
 template <std::size_t N, class Layout, class Extents, class Mapping>
+MDSPAN_INLINE_FUNCTION
 constexpr void validate_strides(with_rank<N>, Layout, const Extents& ext, const Mapping& other)
 {
   static_assert(std::is_same<typename Mapping::layout_type, layout_stride>::value &&
diff --git a/include/experimental/__p0009_bits/utility.hpp b/include/experimental/__p0009_bits/utility.hpp
index e690cd69..c5cdf9d8 100644
--- a/include/experimental/__p0009_bits/utility.hpp
+++ b/include/experimental/__p0009_bits/utility.hpp
@@ -2,6 +2,8 @@
 
 #include <cstddef>
 #include <type_traits>
+#include <array>
+#include <utility>
 
 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 <class T, T v>
+struct integral_constant {
+  using value_type         = T;
+  using type               = integral_constant<T, v>;
+
+  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<T,v>) {};
+  
+  MDSPAN_FUNCTION constexpr operator std::integral_constant<T,v>() const noexcept {
+    return std::integral_constant<T,v>{};
+  }
+
+  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<class T, size_t Idx>
+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<size_t SearchIdx, size_t Idx, class T>
+struct tuple_idx_matcher {
+  using type = tuple_member<T, Idx>;
+  template<class Other>
+  MDSPAN_FUNCTION
+  constexpr auto operator | (Other v) const {
+    if constexpr (Idx == SearchIdx) { return *this; }
+    else { return v; }
+  }
+};
+
+template<class IdxSeq, class ... Elements>
+struct tuple_impl;
+
+template<size_t ... Idx, class ... Elements>
+struct tuple_impl<std::index_sequence<Idx...>, Elements...>: public tuple_member<Elements, Idx> ... {
+
+  MDSPAN_FUNCTION
+  constexpr tuple_impl(Elements ... vals):tuple_member<Elements, Idx>{vals}... {}
+
+  template<size_t N>
+  MDSPAN_FUNCTION
+  constexpr auto& get() {
+    using base_t = decltype((tuple_idx_matcher<N, Idx, Elements>() | ...) );
+    return base_t::type::get();
+  }
+  template<size_t N>
+  MDSPAN_FUNCTION
+  constexpr const auto& get() const {
+    using base_t = decltype((tuple_idx_matcher<N, Idx, Elements>() | ...) );
+    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<class ... Elements>
+struct tuple: public tuple_impl<decltype(std::make_index_sequence<sizeof...(Elements)>()), Elements...> {
+  MDSPAN_FUNCTION
+  constexpr tuple(Elements ... vals):tuple_impl<decltype(std::make_index_sequence<sizeof...(Elements)>()), Elements ...>(vals ...) {}
+};
+
+template<size_t Idx, class ... Args>
+MDSPAN_FUNCTION
+constexpr auto& get(tuple<Args...>& vals) { return vals.template get<Idx>(); }
+
+template<size_t Idx, class ... Args>
+MDSPAN_FUNCTION
+constexpr const auto& get(const tuple<Args...>& vals) { return vals.template get<Idx>(); }
+
+template<class ... Elements>
+tuple(Elements ...) -> tuple<Elements...>;
+#endif
 } // namespace detail
 
 constexpr struct mdspan_non_standard_tag {
diff --git a/include/experimental/__p2630_bits/submdspan_extents.hpp b/include/experimental/__p2630_bits/submdspan_extents.hpp
index 779a9316..4fe5dc6e 100644
--- a/include/experimental/__p2630_bits/submdspan_extents.hpp
+++ b/include/experimental/__p2630_bits/submdspan_extents.hpp
@@ -16,10 +16,11 @@
 
 #pragma once
 
-#include <tuple>
 #include <complex>
 
 #include "strided_slice.hpp"
+#include "../__p0009_bits/utility.hpp"
+
 namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
 namespace detail {
 
@@ -68,6 +69,12 @@ struct index_pair_like<std::tuple<IdxT1, IdxT2>, IndexType> {
                                 std::is_convertible_v<IdxT2, IndexType>;
 };
 
+template <class IdxT1, class IdxT2, class IndexType>
+struct index_pair_like<tuple<IdxT1, IdxT2>, IndexType> {
+  static constexpr bool value = std::is_convertible_v<IdxT1, IndexType> &&
+                                std::is_convertible_v<IdxT2, IndexType>;
+};
+
 template <class IdxT, class IndexType>
 struct index_pair_like<std::complex<IdxT>, IndexType> {
   static constexpr bool value = std::is_convertible_v<IdxT, IndexType>;
@@ -88,10 +95,16 @@ constexpr Integral first_of(const Integral &i) {
   return i;
 }
 
+template<class Integral, Integral v>
+MDSPAN_INLINE_FUNCTION
+constexpr Integral first_of(const std::integral_constant<Integral, v>&) {
+  return integral_constant<Integral, v>();
+}
+
 MDSPAN_INLINE_FUNCTION
-constexpr std::integral_constant<size_t, 0>
+constexpr integral_constant<size_t, 0>
 first_of(const ::MDSPAN_IMPL_STANDARD_NAMESPACE::full_extent_t &) {
-  return std::integral_constant<size_t, 0>();
+  return integral_constant<size_t, 0>();
 }
 
 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<std::tuple<IdxT1, IdxT2>, size_t>::value)
+  )
+constexpr auto first_of(const std::tuple<IdxT1, IdxT2>& i) {
+  return get<0>(i);
+}
+
+MDSPAN_TEMPLATE_REQUIRES(
+  class IdxT1, class IdxT2,
+  /* requires */ (index_pair_like<std::pair<IdxT1, IdxT2>, size_t>::value)
+  )
+MDSPAN_INLINE_FUNCTION
+constexpr auto first_of(const std::pair<IdxT1, IdxT2>& i) {
+  return i.first;
 }
 
 template<class T>
@@ -137,7 +167,24 @@ MDSPAN_TEMPLATE_REQUIRES(
 MDSPAN_INLINE_FUNCTION
 constexpr auto last_of(std::integral_constant<size_t, k>, 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<std::tuple<IdxT1, IdxT2>, size_t>::value)
+  )
+constexpr auto last_of(std::integral_constant<size_t, k>, const Extents &, const std::tuple<IdxT1, IdxT2>& i) {
+  return get<1>(i);
+}
+
+MDSPAN_TEMPLATE_REQUIRES(
+  size_t k, class Extents, class IdxT1, class IdxT2,
+  /* requires */ (index_pair_like<std::pair<IdxT1, IdxT2>, size_t>::value)
+  )
+MDSPAN_INLINE_FUNCTION
+constexpr auto last_of(std::integral_constant<size_t, k>, const Extents &, const std::pair<IdxT1, IdxT2>& i) {
+  return i.second;
 }
 
 template<size_t k, class Extents, class T>
@@ -172,7 +219,7 @@ constexpr auto last_of(std::integral_constant<size_t, k>, const Extents &ext,
   if constexpr (Extents::static_extent(k) == dynamic_extent) {
     return ext.extent(k);
   } else {
-    return std::integral_constant<size_t, Extents::static_extent(k)>();
+    return integral_constant<size_t, Extents::static_extent(k)>();
   }
 #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<size_t, k>, const Extents &,
 template <class T>
 MDSPAN_INLINE_FUNCTION
 constexpr auto stride_of(const T &) {
-  return std::integral_constant<size_t, 1>();
+  return integral_constant<size_t, 1>();
 }
 
 template <class OffsetType, class ExtentType, class StrideType>
@@ -227,7 +274,7 @@ constexpr auto divide(const std::integral_constant<T0, v0> &,
                       const std::integral_constant<T1, v1> &) {
   // cutting short division by zero
   // this is used for strided_slice with zero extent/stride
-  return std::integral_constant<IndexT, v0 == 0 ? 0 : v0 / v1>();
+  return integral_constant<IndexT, v0 == 0 ? 0 : v0 / v1>();
 }
 
 // multiply which can deal with integral constant preservation
@@ -241,7 +288,7 @@ template <class IndexT, class T0, T0 v0, class T1, T1 v1>
 MDSPAN_INLINE_FUNCTION
 constexpr auto multiply(const std::integral_constant<T0, v0> &,
                         const std::integral_constant<T1, v1> &) {
-  return std::integral_constant<IndexT, v0 * v1>();
+  return integral_constant<IndexT, v0 * v1>();
 }
 
 // compute new static extent from range, preserving static knowledge
@@ -255,6 +302,12 @@ struct StaticExtentFromRange<std::integral_constant<Integral0, val0>,
   constexpr static size_t value = val1 - val0;
 };
 
+template <class Integral0, Integral0 val0, class Integral1, Integral1 val1>
+struct StaticExtentFromRange<integral_constant<Integral0, val0>,
+                             integral_constant<Integral1, val1>> {
+  constexpr static size_t value = val1 - val0;
+};
+
 // compute new static extent from strided_slice, preserving static
 // knowledge
 template <class Arg0, class Arg1> struct StaticExtentFromStridedRange {
@@ -267,6 +320,12 @@ struct StaticExtentFromStridedRange<std::integral_constant<Integral0, val0>,
   constexpr static size_t value = val0 > 0 ? 1 + (val0 - 1) / val1 : 0;
 };
 
+template <class Integral0, Integral0 val0, class Integral1, Integral1 val1>
+struct StaticExtentFromStridedRange<integral_constant<Integral0, val0>,
+                                    integral_constant<Integral1, val1>> {
+  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 <size_t K, class Extents, size_t... NewExtents>
diff --git a/include/experimental/__p2630_bits/submdspan_mapping.hpp b/include/experimental/__p2630_bits/submdspan_mapping.hpp
index 69762e44..2a2cdf76 100644
--- a/include/experimental/__p2630_bits/submdspan_mapping.hpp
+++ b/include/experimental/__p2630_bits/submdspan_mapping.hpp
@@ -17,9 +17,9 @@
 #pragma once
 
 #include <array>
-#include <tuple>
 #include <type_traits>
 #include <utility> // 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 <class LayoutMapping> 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<IndexType, Exts...> &exts,
 }
 
 // constructs sub strides
+template<class T, size_t N>
+struct sub_strides
+{
+  T values[N > 0 ? N : 1];
+};
+
 template <class SrcMapping, class... slice_strides, size_t... InvMapIdxs>
 MDSPAN_INLINE_FUNCTION constexpr auto construct_sub_strides(
     const SrcMapping &src_mapping, std::index_sequence<InvMapIdxs...>,
-    const std::tuple<slice_strides...> &slices_stride_factor) {
+    const MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple<slice_strides...> &slices_stride_factor) {
   using index_type = typename SrcMapping::index_type;
-  return std::array<typename SrcMapping::index_type, sizeof...(InvMapIdxs)>{
+  return sub_strides<typename SrcMapping::index_type, sizeof...(InvMapIdxs)>{{
       (static_cast<index_type>(src_mapping.stride(InvMapIdxs)) *
-       static_cast<index_type>(std::get<InvMapIdxs>(slices_stride_factor)))...};
+       static_cast<index_type>(get<InvMapIdxs>(slices_stride_factor)))...}};
 }
 
 template<class SliceSpecifier, class IndexType>
@@ -240,7 +247,7 @@ layout_left::mapping<Extents>::submdspan_mapping_impl(
     auto inv_map = detail::inv_map_rank(std::integral_constant<size_t, 0>(),
                                         std::index_sequence<>(), slices...);
     return submdspan_mapping_result<dst_mapping_t> {
-      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<Extents>::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<decltype(detail::stride_of(slices))...>{
-                            detail::stride_of(slices)...})),
+                        detail::tuple<decltype(detail::stride_of(slices))...>{
+                            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<PaddingValue>::mapping<Extent
                                         std::index_sequence<>(), slices...);
       using dst_mapping_t = typename layout_stride::template mapping<dst_ext_t>;
     return submdspan_mapping_result<dst_mapping_t> {
-      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<PaddingValue>::mapping<Extent
 // the issue But Clang-CUDA also doesn't accept the use of deduction guide so
 // disable it for CUDA alltogether
 #if defined(_MDSPAN_HAS_HIP) || defined(_MDSPAN_HAS_CUDA)
-                        std::tuple<decltype(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices))...>{
-                            MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices)...})),
+                        MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple<decltype(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices))...>{
+                            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<Extents>::submdspan_mapping_impl(
     auto inv_map = detail::inv_map_rank(std::integral_constant<size_t, 0>(),
                                         std::index_sequence<>(), slices...);
     return submdspan_mapping_result<dst_mapping_t> {
-      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<Extents>::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<decltype(detail::stride_of(slices))...>{
-                            detail::stride_of(slices)...})),
+                        MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple<decltype(detail::stride_of(slices))...>{
+                            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<PaddingValue>::mapping<Exten
                                         std::index_sequence<>(), slices...);
       using dst_mapping_t = typename layout_stride::template mapping<dst_ext_t>;
     return submdspan_mapping_result<dst_mapping_t> {
-      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<PaddingValue>::mapping<Exten
 // the issue But Clang-CUDA also doesn't accept the use of deduction guide so
 // disable it for CUDA alltogether
 #if defined(_MDSPAN_HAS_HIP) || defined(_MDSPAN_HAS_CUDA)
-                        std::tuple<decltype(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices))...>{
-                            MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices)...})),
+                        MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple<decltype(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices))...>{
+                            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<Extents>::submdspan_mapping_impl(
                     : this->operator()(detail::first_of(slices)...));
 
   return submdspan_mapping_result<dst_mapping_t> {
-    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<Extents>::submdspan_mapping_impl(
 #if defined(_MDSPAN_HAS_HIP) ||                                                \
     (defined(__NVCC__) &&                                                      \
      (__CUDACC_VER_MAJOR__ * 100 + __CUDACC_VER_MINOR__ * 10) < 1120)
-                      std::tuple<decltype(detail::stride_of(slices))...>(
-                          detail::stride_of(slices)...))),
+                      MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple<decltype(detail::stride_of(slices))...>(
+                          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/include/experimental/__p2642_bits/layout_padded.hpp b/include/experimental/__p2642_bits/layout_padded.hpp
index 1502489a..a714090e 100644
--- a/include/experimental/__p2642_bits/layout_padded.hpp
+++ b/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<PaddingValue>::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<PaddingValue>::mapping {
     return exts;
   }
 
-  MDSPAN_INLINE_FUNCTION constexpr std::array<index_type, extents_type::rank()>
+  constexpr std::array<index_type, extents_type::rank()>
   strides() const noexcept {
     if constexpr (extents_type::rank() == 0) {
       return {};
@@ -568,7 +569,7 @@ class layout_right_padded<PaddingValue>::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<PaddingValue>::mapping {
     return exts;
   }
 
-  MDSPAN_INLINE_FUNCTION constexpr std::array<index_type, extents_type::rank()>
+  constexpr std::array<index_type, extents_type::rank()>
   strides() const noexcept {
     if constexpr (extents_type::rank() == 0) {
       return {};
diff --git a/include/experimental/__p2642_bits/layout_padded_fwd.hpp b/include/experimental/__p2642_bits/layout_padded_fwd.hpp
index 18daa28c..3f141ff0 100644
--- a/include/experimental/__p2642_bits/layout_padded_fwd.hpp
+++ b/include/experimental/__p2642_bits/layout_padded_fwd.hpp
@@ -85,12 +85,15 @@ struct is_layout_right_padded_mapping<_Mapping,
 
 
 template <class _LayoutExtentsType, class _PaddedLayoutMappingType>
+MDSPAN_INLINE_FUNCTION
 constexpr void check_padded_layout_converting_constructor_mandates(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::with_rank<0>) {}
 
 template <class _LayoutExtentsType, class _PaddedLayoutMappingType>
+MDSPAN_INLINE_FUNCTION
 constexpr void check_padded_layout_converting_constructor_mandates(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::with_rank<1>) {}
 
 template <class _LayoutExtentsType, class _PaddedLayoutMappingType, std::size_t N>
+MDSPAN_INLINE_FUNCTION
 constexpr void check_padded_layout_converting_constructor_mandates(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::with_rank<N>)
 {
   using extents_type = typename _PaddedLayoutMappingType::extents_type;
@@ -110,12 +113,15 @@ constexpr void check_padded_layout_converting_constructor_mandates(MDSPAN_IMPL_S
 }
 
 template <typename _ExtentsType, typename _OtherMapping>
+MDSPAN_INLINE_FUNCTION
 constexpr void check_padded_layout_converting_constructor_preconditions(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::with_rank<0>,
                                                                         const _OtherMapping&) {}
 template <typename _ExtentsType, typename _OtherMapping>
+MDSPAN_INLINE_FUNCTION
 constexpr void check_padded_layout_converting_constructor_preconditions(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::with_rank<1>,
                                                                         const _OtherMapping&) {}
 template <typename _ExtentsType, typename _OtherMapping, std::size_t N>
+MDSPAN_INLINE_FUNCTION
 constexpr void check_padded_layout_converting_constructor_preconditions(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::with_rank<N>,
                                                                         const _OtherMapping &other_mapping) {
   constexpr auto padded_stride_idx =