diff --git a/tpls/mdspan/include/experimental/__p0009_bits/device_support.hpp b/tpls/mdspan/include/experimental/__p0009_bits/device_support.hpp
new file mode 100644
index 00000000000..ea6b4f5ac6b
--- /dev/null
+++ b/tpls/mdspan/include/experimental/__p0009_bits/device_support.hpp
@@ -0,0 +1,137 @@
+#pragma once
+#include <type_traits>
+#include "macros.hpp"
+#if !defined(_MDSPAN_HAS_CUDA) && !defined(_MDSPAN_HAS_HIP)
+#include <tuple>
+#endif
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+namespace detail {
+
+// 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>;
+
+  MDSPAN_INLINE_FUNCTION_DEFAULTED
+  constexpr integral_constant() = default;
+  
+  MDSPAN_INLINE_FUNCTION_DEFAULTED
+  constexpr integral_constant(std::integral_constant<T,v>) {};
+
+  static constexpr T value = v;
+  MDSPAN_INLINE_FUNCTION constexpr operator value_type() const noexcept {
+    return value;
+  }
+  MDSPAN_INLINE_FUNCTION constexpr value_type operator()() const noexcept {
+    return value;
+  }
+  MDSPAN_INLINE_FUNCTION constexpr operator std::integral_constant<T,v>() const noexcept {
+    return std::integral_constant<T,v>{};
+  }
+};
+
+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; }
+};
+
+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 SearchT, size_t Idx, class T>
+struct tuple_type_matcher {
+  using type = tuple_member<T, Idx>;
+  template<class Other>
+  MDSPAN_FUNCTION
+  constexpr auto operator + (Other v) const {
+    if constexpr (std::is_same_v<T, SearchT>) { 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<class T>
+  MDSPAN_FUNCTION
+  constexpr T& get() {
+    using base_t = decltype((tuple_type_matcher<T, Idx, Elements>() + ...) );
+    return base_t::type::get();
+  }
+  template<class T>
+  MDSPAN_FUNCTION
+  constexpr const T& get() const {
+    using base_t = decltype((tuple_type_matcher<T, Idx, Elements>() + ...) );
+    return base_t::type::get();
+  }
+
+  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();
+  }
+};
+
+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<class T, class ... Args>
+MDSPAN_FUNCTION
+constexpr auto& get(tuple<Args...>& vals) { return vals.template get<T>(); }
+
+template<class T, class ... Args>
+MDSPAN_FUNCTION
+constexpr const T& get(const tuple<Args...>& vals) { return vals.template get<T>(); }
+
+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...>;
+
+template<class T, size_t ... Idx>
+constexpr auto c_array_to_std(std::index_sequence<Idx...>, const T(&values)[sizeof...(Idx)]) {
+  return std::array{values[Idx]...};
+}
+template<class T, size_t N>
+constexpr auto c_array_to_std(const T(&values)[N]) {
+  return c_array_to_std(std::make_index_sequence<N>(), values);
+}
+}
+}
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<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/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<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/tpls/mdspan/include/experimental/__p0009_bits/layout_stride.hpp b/tpls/mdspan/include/experimental/__p0009_bits/layout_stride.hpp
index d6cdad2ab23..0e7ccc1523d 100644
--- a/tpls/mdspan/include/experimental/__p0009_bits/layout_stride.hpp
+++ b/tpls/mdspan/include/experimental/__p0009_bits/layout_stride.hpp
@@ -199,7 +199,15 @@ 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])...};
+        // avoid warning for use of host std::array operator[]
+  #if defined(_MDSPAN_HAS_CUDA) || defined(_MDSPAN_HAS_HIP)
+        const IntegralType* s_ptr = reinterpret_cast<const IntegralType*>(&s);
+  #else
+        const IntegralType *s_ptr = s.data();
+  #endif
+        // for rank == 0 the expansion is empty and s_ptr becomes unused
+        detail::maybe_unused_variable(s_ptr);
+        return __strides_storage_t{static_cast<index_type>(s_ptr[Idxs])...};
       }
 
       MDSPAN_TEMPLATE_REQUIRES(
@@ -218,7 +226,6 @@ struct layout_stride {
 
 #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 +249,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 +283,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 +309,7 @@ struct layout_stride {
         _MDSPAN_TRAIT(std::is_nothrow_constructible, typename Extents::index_type, const std::remove_const_t<IntegralTypes>&)
       )
     )
-    MDSPAN_INLINE_FUNCTION
+    MDSPAN_FUNCTION
     constexpr
     mapping(
       extents_type const& e,
@@ -379,7 +389,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 +485,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 +519,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 +539,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 +641,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 +652,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/tpls/mdspan/include/experimental/__p0009_bits/utility.hpp b/tpls/mdspan/include/experimental/__p0009_bits/utility.hpp
index e690cd6939b..96e2b5300f6 100644
--- a/tpls/mdspan/include/experimental/__p0009_bits/utility.hpp
+++ b/tpls/mdspan/include/experimental/__p0009_bits/utility.hpp
@@ -64,6 +64,10 @@ constexpr struct
   }
 } stride;
 
+template<class T>
+MDSPAN_INLINE_FUNCTION
+constexpr void maybe_unused_variable(const T&) {}
+
 } // 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..f43876e8f91 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 <tuple>
 #include <complex>
 
 #include "strided_slice.hpp"
+#include "../__p0009_bits/device_support.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/tpls/mdspan/include/experimental/__p2630_bits/submdspan_mapping.hpp b/tpls/mdspan/include/experimental/__p2630_bits/submdspan_mapping.hpp
index 69762e44e2c..97f3cd7466b 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 <array>
-#include <tuple>
 #include <type_traits>
 #include <utility> // index_sequence
+#include "../__p0009_bits/device_support.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
@@ -87,11 +88,11 @@ any_slice_out_of_bounds(const extents<IndexType, Exts...> &exts,
 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)>{
       (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>
@@ -248,10 +249,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::tuple<decltype(detail::stride_of(slices))...>{
                             detail::stride_of(slices)...})),
 #else
-                        std::tuple{detail::stride_of(slices)...})),
+                        detail::tuple{detail::stride_of(slices)...})),
 #endif
           offset
     };
@@ -326,10 +327,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::tuple<decltype(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices))...>{
                             MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices)...})),
 #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)...})),
 #endif
           offset
     };
@@ -481,10 +482,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))...>{
+                        MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple<decltype(detail::stride_of(slices))...>{
                             detail::stride_of(slices)...})),
 #else
-                        std::tuple{detail::stride_of(slices)...})),
+                        MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple{detail::stride_of(slices)...})),
 #endif
           offset
     };
@@ -551,10 +552,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::tuple<decltype(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices))...>{
                             MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices)...})),
 #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)...})),
 #endif
           offset
     };
@@ -600,10 +601,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))...>(
+                      MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple<decltype(detail::stride_of(slices))...>(
                           detail::stride_of(slices)...))),
 #else
-                      std::tuple(detail::stride_of(slices)...))),
+                      MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple(detail::stride_of(slices)...))),
 #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..8a5d9069d63 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/device_support.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{})
   {}
@@ -365,7 +366,7 @@ class layout_left_padded<PaddingValue>::mapping {
       return {1};
     } else {
       index_type value = 1;
-      std::array<index_type, extents_type::rank()> s{};
+      index_type s[extents_type::rank()];
       s[extent_to_pad_idx] = value;
       value *= padded_stride.value(0);
       for (rank_type r = extent_to_pad_idx + 1; r < extents_type::rank() - 1;
@@ -374,7 +375,7 @@ class layout_left_padded<PaddingValue>::mapping {
         value *= exts.extent(r);
       }
       s[extents_type::rank() - 1] = value;
-      return s;
+      return MDSPAN_IMPL_STANDARD_NAMESPACE::detail::c_array_to_std(s);
     }
   }
 
@@ -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{})
   {}
@@ -725,7 +726,7 @@ class layout_right_padded<PaddingValue>::mapping {
       return {1};
     } else {
       index_type value = 1;
-      std::array<index_type, extents_type::rank()> s{};
+      index_type s[extents_type::rank()];
       s[extent_to_pad_idx] = value;
       value *= padded_stride.value(0);
       for (rank_type r = extent_to_pad_idx - 1; r > 0; --r) {
@@ -733,7 +734,7 @@ class layout_right_padded<PaddingValue>::mapping {
         value *= exts.extent(r);
       }
       s[0] = value;
-      return s;
+      return MDSPAN_IMPL_STANDARD_NAMESPACE::detail::c_array_to_std(s);
     }
   }
 
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 <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 =