diff --git a/libcudacxx/include/cuda/std/__mdspan/submdspan_helper.h b/libcudacxx/include/cuda/std/__mdspan/submdspan_helper.h
index c617480e79..2f90614f41 100644
--- a/libcudacxx/include/cuda/std/__mdspan/submdspan_helper.h
+++ b/libcudacxx/include/cuda/std/__mdspan/submdspan_helper.h
@@ -84,28 +84,17 @@ struct full_extent_t
 };
 _CCCL_GLOBAL_CONSTANT full_extent_t full_extent{};
 
-// [mdspan.submdspan.submdspan.mapping.result]
-template <class _LayoutMapping>
-struct submdspan_mapping_result
-{
-  static_assert(true, // __is_layout_mapping<_LayoutMapping>,
-                "[mdspan.submdspan.submdspan.mapping.result] shall meet the layout mapping requirements");
-
-  _CCCL_NO_UNIQUE_ADDRESS _LayoutMapping mapping{};
-  size_t offset{};
-};
-
 // [mdspan.submdspan.helpers]
 _CCCL_TEMPLATE(class _Tp)
 _CCCL_REQUIRES((!__integral_constant_like<_Tp>) )
-_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __de_ice(_Tp __val)
+_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __de_ice(_Tp __val) noexcept
 {
   return __val;
 }
 
 _CCCL_TEMPLATE(class _Tp)
 _CCCL_REQUIRES(__integral_constant_like<_Tp>)
-_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __de_ice(_Tp)
+_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr auto __de_ice(_Tp) noexcept
 {
   return _Tp::value;
 }
diff --git a/libcudacxx/include/cuda/std/__mdspan/submdspan_mapping.h b/libcudacxx/include/cuda/std/__mdspan/submdspan_mapping.h
new file mode 100644
index 0000000000..5f856f4196
--- /dev/null
+++ b/libcudacxx/include/cuda/std/__mdspan/submdspan_mapping.h
@@ -0,0 +1,332 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBCUDACXX___MDSPAN_SUBMDSPAN_MAPPING_H
+#define _LIBCUDACXX___MDSPAN_SUBMDSPAN_MAPPING_H
+
+#include <cuda/std/detail/__config>
+
+#include "cuda/std/__cccl/unreachable.h"
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/std/__fwd/mdspan.h>
+#include <cuda/std/__mdspan/concepts.h>
+#include <cuda/std/__mdspan/extents.h>
+#include <cuda/std/__mdspan/layout_left.h>
+#include <cuda/std/__mdspan/layout_right.h>
+#include <cuda/std/__mdspan/layout_stride.h>
+#include <cuda/std/__mdspan/mdspan.h>
+#include <cuda/std/__mdspan/submdspan_extents.h>
+#include <cuda/std/__mdspan/submdspan_helper.h>
+#include <cuda/std/__type_traits/make_unsigned.h>
+#include <cuda/std/__type_traits/remove_const.h>
+#include <cuda/std/__type_traits/type_list.h>
+#include <cuda/std/__utility/integer_sequence.h>
+#include <cuda/std/array>
+
+#if _CCCL_STD_VER >= 2014
+
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
+
+// [mdspan.sub.map]
+
+// [mdspan.submdspan.submdspan.mapping.result]
+template <class _LayoutMapping>
+struct submdspan_mapping_result
+{
+  static_assert(true, // __is_layout_mapping<_LayoutMapping>,
+                "[mdspan.submdspan.submdspan.mapping.result] shall meet the layout mapping requirements");
+
+  _CCCL_NO_UNIQUE_ADDRESS _LayoutMapping mapping{};
+  size_t offset{};
+};
+
+// [mdspan.sub.map.common]
+_CCCL_TEMPLATE(size_t _SliceIndex, class _LayoutMapping, class... _SliceSpecifiers)
+_CCCL_REQUIRES(__is_strided_slice<__get_slice_type<_SliceIndex, _SliceSpecifiers...>>)
+_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr auto
+__get_submdspan_strides(const _LayoutMapping& __mapping, _SliceSpecifiers... __slices) noexcept
+{
+  using _SliceType    = __get_slice_type<_SliceIndex, _SliceSpecifiers...>;
+  _SliceType& __slice = _CUDA_VSTD::__get_slice_at<_SliceIndex>(__slices...);
+
+  using __unsigned_stride = make_unsigned_t<typename _SliceType::stride_type>;
+  using __unsigned_extent = make_unsigned_t<typename _SliceType::extent_type>;
+  return __mapping.stride(_SliceIndex)
+       * (static_cast<__unsigned_stride>(__slice.stride) < static_cast<__unsigned_extent>(__slice.extent)
+            ? _CUDA_VSTD::__de_ice(__slice.stride)
+            : 1);
+}
+
+_CCCL_TEMPLATE(size_t _SliceIndex, class _LayoutMapping, class... _SliceSpecifiers)
+_CCCL_REQUIRES((!__is_strided_slice<__get_slice_type<_SliceIndex, _SliceSpecifiers...>>) )
+_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr auto
+__get_submdspan_strides(const _LayoutMapping& __mapping, _SliceSpecifiers...) noexcept
+{
+  return __mapping.stride(_SliceIndex);
+}
+
+template <class _LayoutMapping, class... _SliceSpecifiers, size_t... _SliceIndexes>
+_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr auto __submdspan_strides(
+  index_sequence<_SliceIndexes...>, const _LayoutMapping& __mapping, _SliceSpecifiers... __slices) noexcept
+{
+  using _Extents                                  = typename _LayoutMapping::extents_type;
+  using _IndexType                                = typename _Extents::index_type;
+  constexpr auto __map_rank_                      = _CUDA_VSTD::__map_rank<_IndexType, _SliceSpecifiers...>();
+  const array<_IndexType, _Extents::rank()> __arr = {
+    _CUDA_VSTD::__get_submdspan_strides<_SliceIndexes>(__mapping, __slices...)...};
+
+  using _SubExtent                            = __get_subextents_t<_Extents, _SliceSpecifiers...>;
+  array<_IndexType, _SubExtent::rank()> __res = {};
+  for (size_t __index = 0; __index != _SubExtent::rank(); ++__index)
+  {
+    if (__map_rank_[__index] != dynamic_extent)
+    {
+      __res[__map_rank_[__index]] = __arr[__index];
+    }
+  }
+  return __res;
+}
+
+template <class _LayoutMapping, class... _SliceSpecifiers>
+_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr auto
+__submdspan_strides(const _LayoutMapping& __mapping, _SliceSpecifiers... __slices)
+{
+  return _CUDA_VSTD::__submdspan_strides(_CUDA_VSTD::index_sequence_for<_SliceSpecifiers...>(), __mapping, __slices...);
+}
+
+// [mdspan.sub.map.common-8]
+template <class _LayoutMapping, class... _SliceSpecifiers, size_t... _SliceIndexes>
+_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr size_t
+__submdspan_offset(index_sequence<_SliceIndexes...>, const _LayoutMapping& __mapping, _SliceSpecifiers... __slices)
+{
+  using _Extents   = typename _LayoutMapping::extents_type;
+  using _IndexType = typename _Extents::index_type;
+  // If first_<index_type, k>(slices...)
+  const array<_IndexType, _Extents::rank()> __offsets = {
+    _CUDA_VSTD::__first_extent_from_slice<_IndexType, _SliceIndexes>(__slices...)...};
+
+  using _SubExtent = __get_subextents_t<_Extents, _SliceSpecifiers...>;
+  for (size_t __index = 0; __index != _SubExtent::rank(); ++__index)
+  {
+    // If first_<index_type, k>(slices...) equals extents().extent(k) for any rank index k of extents()
+    if (__offsets[__index] == __mapping.extents().extent(__index))
+    {
+      // then let offset be a value of type size_t equal to (*this).required_span_size()
+      return static_cast<size_t>(__mapping.required_span_size());
+    }
+  }
+  // Otherwise, let offset be a value of type size_t equal to (*this)(first_<index_type, P>(slices...)...).
+  return static_cast<size_t>(__mapping(__offsets[_SliceIndexes]...));
+}
+
+template <class _LayoutMapping, class... _SliceSpecifiers>
+_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr size_t
+__submdspan_offset(const _LayoutMapping& __mapping, _SliceSpecifiers... __slices)
+{
+  return _CUDA_VSTD::__submdspan_offset(_CUDA_VSTD::index_sequence_for<_SliceSpecifiers...>(), __mapping, __slices...);
+}
+
+// [mdspan.sub.map.common-9]
+// [mdspan.sub.map.common-9.1]
+template <class _SliceType>
+_CCCL_CONCEPT __is_strided_slice_stride_of_one = _CCCL_REQUIRES_EXPR((_SliceType))(
+  requires(__is_strided_slice<_SliceType>),
+  requires(__integral_constant_like<typename _SliceType::stride_type>),
+  requires(_SliceType::stride_type::value == 1));
+
+template <class _LayoutMapping, class _SliceType>
+_LIBCUDACXX_HIDE_FROM_ABI constexpr bool __is_unit_stride_slice()
+{
+  // [mdspan.sub.map.common-9.1]
+  if constexpr (__is_strided_slice_stride_of_one<_SliceType>)
+  {
+    return true;
+  }
+  // [mdspan.sub.map.common-9.2]
+  else if constexpr (__index_pair_like<_SliceType, typename _LayoutMapping::index_type>)
+  {
+    return true;
+  }
+  // [mdspan.sub.map.common-9.3]
+  else if constexpr (_CCCL_TRAIT(is_convertible, _SliceType, full_extent_t))
+  {
+    return true;
+  }
+  else
+  {
+    return false;
+  }
+  _CCCL_UNREACHABLE();
+}
+
+// [mdspan.sub.map.left]
+template <class _LayoutMapping, class _SubExtents, class... _SliceSpecifiers>
+_LIBCUDACXX_HIDE_FROM_ABI constexpr bool __can_layout_left()
+{
+  // [mdspan.sub.map.left-1.2]
+  if constexpr (_SubExtents::rank() == 0)
+  {
+    return true;
+  }
+  // [mdspan.sub.map.left-1.3.1]
+  // Note we can simplify metaprogramming here a bit because unit-stride slice is true if that condition holds
+  else if constexpr (_CCCL_FOLD_AND(_CCCL_TRAIT(is_convertible, _SliceSpecifiers, full_extent_t)))
+  {
+    return true;
+  }
+  else
+  {
+    // [mdspan.sub.map.left-1.3.2]
+    return _CUDA_VSTD::__is_unit_stride_slice<_LayoutMapping,
+                                              __type_index_c<_SubExtents::rank() - 1, _SliceSpecifiers...>>();
+  }
+  _CCCL_UNREACHABLE();
+}
+
+template <class _Extents, class... _SliceSpecifiers>
+_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr auto
+__submdspan_mapping_impl(const typename layout_left::mapping<_Extents>& __mapping, _SliceSpecifiers... __slices)
+{
+  // [mdspan.sub.map.left-1.1]
+  if constexpr (_Extents::rank() == 0)
+  {
+    return submdspan_mapping_result{__mapping, 0};
+  }
+
+  // [mdspan.sub.map.left-1.2]
+  // [mdspan.sub.map.left-1.3]
+  using _SubExtents    = __get_subextents_t<_Extents, _SliceSpecifiers...>;
+  const auto __sub_ext = _CUDA_VSTD::submdspan_extents(__mapping.extents(), __slices...);
+  const auto __offset  = _CUDA_VSTD::__submdspan_offset(__mapping, __slices...);
+  if constexpr (_CUDA_VSTD::
+                  __can_layout_left<typename layout_left::mapping<_Extents>, _SubExtents, _SliceSpecifiers...>())
+  {
+    return submdspan_mapping_result<layout_left::mapping<_SubExtents>>{layout_left::mapping{__sub_ext}, __offset};
+  }
+  // [mdspan.sub.map.left-1.4]
+  // TODO: Implement padded layouts
+  else
+  {
+    // [mdspan.sub.map.left-1.5]
+    const auto __sub_strides = _CUDA_VSTD::__submdspan_strides(__mapping, __slices...);
+    return submdspan_mapping_result<layout_stride::mapping<_SubExtents>>{
+      layout_stride::mapping{__sub_ext, __sub_strides}, __offset};
+  }
+  _CCCL_UNREACHABLE();
+}
+
+template <class _LayoutMapping, class _SubExtents, class... _SliceSpecifiers>
+_LIBCUDACXX_HIDE_FROM_ABI constexpr bool __can_layout_right()
+{
+  // [mdspan.sub.map.right-1.2]
+  if constexpr (_SubExtents::rank() == 0)
+  {
+    return true;
+  }
+  // [mdspan.sub.map.right-1.3.1]
+  // Note we can simplify metaprogramming here a bit because unit-stride slice is true if that condition holds
+  else if constexpr (_CCCL_FOLD_AND(_CCCL_TRAIT(is_convertible, _SliceSpecifiers, full_extent_t)))
+  {
+    return true;
+  }
+  else
+  {
+    // [mdspan.sub.map.right-1.3.2]
+    return _CUDA_VSTD::__is_unit_stride_slice<_LayoutMapping,
+                                              __type_index_c<_SubExtents::rank() - 1, _SliceSpecifiers...>>();
+  }
+  _CCCL_UNREACHABLE();
+}
+
+template <class _Extents, class... _SliceSpecifiers>
+_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr auto
+__submdspan_mapping_impl(const typename layout_right::mapping<_Extents>& __mapping, _SliceSpecifiers... __slices)
+{
+  // [mdspan.sub.map.right-1.1]
+  if constexpr (_Extents::rank() == 0)
+  {
+    return submdspan_mapping_result{__mapping, 0};
+  }
+  else
+  {
+    // [mdspan.sub.map.right-1.2]
+    // [mdspan.sub.map.right-1.3]
+    using _SubExtents    = __get_subextents_t<_Extents, _SliceSpecifiers...>;
+    const auto __sub_ext = _CUDA_VSTD::submdspan_extents(__mapping.extents(), __slices...);
+    const auto __offset  = _CUDA_VSTD::__submdspan_offset(__mapping, __slices...);
+    if constexpr (_CUDA_VSTD::
+                    __can_layout_right<typename layout_left::mapping<_Extents>, _SubExtents, _SliceSpecifiers...>())
+    {
+      return submdspan_mapping_result<layout_right::mapping<_SubExtents>>{layout_right::mapping{__sub_ext}, __offset};
+    }
+    // [mdspan.sub.map.right-1.4]
+    // TODO: Implement padded layouts
+    else
+    {
+      // [mdspan.sub.map.right-1.5]
+      const auto __sub_strides = _CUDA_VSTD::__submdspan_strides(__mapping, __slices...);
+      return submdspan_mapping_result<layout_stride::mapping<_SubExtents>>{
+        layout_stride::mapping{__sub_ext, __sub_strides}, __offset};
+    }
+  }
+  _CCCL_UNREACHABLE();
+}
+
+template <class _Extents, class... _SliceSpecifiers>
+_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr auto
+__submdspan_mapping_impl(const typename layout_stride::mapping<_Extents>& __mapping, _SliceSpecifiers... __slices)
+{
+  // [mdspan.sub.map.stride-1.1]
+  if constexpr (_Extents::rank() == 0)
+  {
+    return submdspan_mapping_result{__mapping, 0};
+  }
+  else
+  {
+    // [mdspan.sub.map.stride-1.2]
+    using _SubExtents        = __get_subextents_t<_Extents, _SliceSpecifiers...>;
+    const auto __sub_ext     = _CUDA_VSTD::submdspan_extents(__mapping.extents(), __slices...);
+    const auto __offset      = _CUDA_VSTD::__submdspan_offset(__mapping, __slices...);
+    const auto __sub_strides = _CUDA_VSTD::__submdspan_strides(__mapping, __slices...);
+    return submdspan_mapping_result{layout_stride::mapping{__sub_ext, __sub_strides}, __offset};
+  }
+}
+
+template <class _LayoutMapping, class... _SliceSpecifiers>
+_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr auto
+submdspan_mapping(const _LayoutMapping& __mapping, _SliceSpecifiers... __slices)
+{
+  return _CUDA_VSTD::__submdspan_mapping_impl(__mapping, __slices...);
+}
+
+_CCCL_TEMPLATE(class _Tp, class _Extents, class _Layout, class _Accessor, class... _SliceSpecifiers)
+_CCCL_REQUIRES(true)
+_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr auto
+submdspan(const mdspan<_Tp, _Extents, _Layout, _Accessor>& __src, _SliceSpecifiers... __slices)
+{
+  auto __sub_map_result = _CUDA_VSTD::submdspan_mapping(__src.mapping(), __slices...);
+  return mdspan(__src.accessor().offset(__src.data_handle(), __sub_map_result.offset),
+                __sub_map_result.mapping,
+                typename _Accessor::offset_policy(__src.accessor()));
+}
+
+_LIBCUDACXX_END_NAMESPACE_STD
+
+#endif // _CCCL_STD_VER >= 2014
+
+#endif // _LIBCUDACXX___MDSPAN_SUBMDSPAN_MAPPING_H
diff --git a/libcudacxx/include/cuda/std/mdspan b/libcudacxx/include/cuda/std/mdspan
index 75e4b8f370..ac75b2ac70 100644
--- a/libcudacxx/include/cuda/std/mdspan
+++ b/libcudacxx/include/cuda/std/mdspan
@@ -33,6 +33,7 @@ _CCCL_PUSH_MACROS
 #include <cuda/std/__mdspan/mdspan.h>
 #include <cuda/std/__mdspan/submdspan_extents.h>
 #include <cuda/std/__mdspan/submdspan_helper.h>
+#include <cuda/std/__mdspan/submdspan_mapping.h>
 #include <cuda/std/version>
 
 _CCCL_POP_MACROS
diff --git a/libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/helper.h b/libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/helper.h
new file mode 100644
index 0000000000..1a0c16db14
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/helper.h
@@ -0,0 +1,55 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef TEST_STD_CONTAINERS_VIEWS_MDSPAN_SUBMDSPAN_HELPER_H
+#define TEST_STD_CONTAINERS_VIEWS_MDSPAN_SUBMDSPAN_HELPER_H
+
+#include <cuda/std/mdspan>
+#include <cuda/std/type_traits>
+
+_CCCL_TEMPLATE(class MDSpan)
+_CCCL_REQUIRES((MDSpan::rank() == 0))
+__host__ __device__ constexpr bool equal_to(const MDSpan& mdspan, const char* expected)
+{
+  return mdspan.data_handle()[0] == expected[0];
+}
+
+_CCCL_TEMPLATE(class MDSpan)
+_CCCL_REQUIRES((MDSpan::rank() == 1))
+__host__ __device__ constexpr bool equal_to(const MDSpan& mdspan, const char* expected)
+{
+  for (size_t i = 0; i != mdspan.size(); ++i)
+  {
+    if (mdspan[i] != expected[i])
+    {
+      return false;
+    }
+  }
+  return true;
+}
+
+_CCCL_TEMPLATE(class MDSpan)
+_CCCL_REQUIRES((MDSpan::rank() == 2))
+__host__ __device__ constexpr bool equal_to(const MDSpan& mdspan, cuda::std::array<const char*, 2> expected)
+{
+  for (size_t i = 0; i != mdspan.extent(0); ++i)
+  {
+    for (size_t j = 0; j != mdspan.extent(1); ++j)
+    {
+      const cuda::std::array<size_t, 2> indices{i, j};
+      if (mdspan[indices] != expected[i][j])
+      {
+        return false;
+      }
+    }
+  }
+  return true;
+}
+
+#endif // TEST_STD_CONTAINERS_VIEWS_MDSPAN_SUBMDSPAN_HELPER_H
diff --git a/libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/layout_left.pass.cpp b/libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/layout_left.pass.cpp
new file mode 100644
index 0000000000..045d93ba2a
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/layout_left.pass.cpp
@@ -0,0 +1,284 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: c++11, c++14
+
+// <mdspan>
+
+// constexpr mdspan& operator=(const mdspan& rhs) = default;
+
+#include <cuda/std/cassert>
+#include <cuda/std/mdspan>
+#include <cuda/std/type_traits>
+
+#include "helper.h"
+#include "test_macros.h"
+
+__host__ __device__ constexpr bool test()
+{
+  constexpr char data[] = {'H', 'O', 'P', 'P', 'E', 'R'};
+
+  { // 1d mdspan
+    // ['H', 'O', 'P', 'P', 'E', 'R']
+    cuda::std::mdspan md{data, cuda::std::layout_left::mapping{cuda::std::dims<1>{6}}};
+    static_assert(md.rank() == 1);
+    static_assert(md.rank_dynamic() == 1);
+    assert(equal_to(md, "HOPPER"));
+
+    using mdspan_t = decltype(md);
+    static_assert(cuda::std::is_same_v<typename mdspan_t::layout_type, cuda::std::layout_left>);
+
+    { // full_extent
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [ x    x    x    x    x    x ]
+      cuda::std::mdspan sub = cuda::std::submdspan(md, cuda::std::full_extent);
+
+      static_assert(sub.rank() == 1);
+      static_assert(sub.rank_dynamic() == 1);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_left>);
+
+      assert(sub.stride(0) == 1);
+      assert(sub.extent(0) == 6);
+      assert(sub.size() == 6);
+      assert(equal_to(sub, "HOPPER"));
+    }
+
+    { // Slice of elements from start 0:4
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [ x    x    x    x           ]
+      cuda::std::mdspan sub = cuda::std::submdspan(md, cuda::std::pair{0, 4});
+
+      static_assert(sub.rank() == 1);
+      static_assert(sub.rank_dynamic() == 1);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_left>);
+
+      assert(sub.stride(0) == 1);
+      assert(sub.extent(0) == 4);
+      assert(sub.size() == 4);
+      assert(equal_to(sub, "HOPP"));
+    }
+
+    { // Slice of elements in the middle 2:5
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [           x    x    x      ]
+      cuda::std::mdspan sub = cuda::std::submdspan(md, cuda::std::pair{2, 5});
+
+      static_assert(sub.rank() == 1);
+      static_assert(sub.rank_dynamic() == 1);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_left>);
+
+      assert(sub.stride(0) == 1);
+      assert(sub.extent(0) == 3);
+      assert(sub.size() == 3);
+      assert(equal_to(sub, "PPE"));
+    }
+
+    { // Slice of elements in the end 3:6
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [                x    x    x ]
+      cuda::std::mdspan sub = cuda::std::submdspan(md, cuda::std::pair{3, 6});
+
+      static_assert(sub.rank() == 1);
+      static_assert(sub.rank_dynamic() == 1);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_left>);
+
+      assert(sub.stride(0) == 1);
+      assert(sub.extent(0) == 3);
+      assert(sub.size() == 3);
+      assert(equal_to(sub, "PER"));
+    }
+
+    { // Slice of elements with strided slice without offset, full size and stride 1
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [ x    x    x    x    x    x ] offset
+      // [ x    x    x    x    x    x ] size
+      // [ x    x    x    x    x    x ] stride
+      // [ x    x    x    x    x    x ]
+      cuda::std::strided_slice slice{0, md.extent(0), 1};
+      cuda::std::mdspan sub = cuda::std::submdspan(md, slice);
+
+      static_assert(sub.rank() == 1);
+      static_assert(sub.rank_dynamic() == 1);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_stride>);
+
+      assert(sub.stride(0) == 1);
+      assert(sub.extent(0) == 6);
+      assert(sub.size() == 6);
+      assert(equal_to(sub, "HOPPER"));
+    }
+
+    { // Slice of elements with strided slice with offset, full remaing size and stride 1
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [           x    x    x    x ] offset
+      // [           x    x    x    x ] size
+      // [           x    x    x    x ] stride
+      // [           x    x    x    x ]
+      cuda::std::strided_slice slice{2, md.extent(0) - 2, 1};
+      cuda::std::mdspan sub = cuda::std::submdspan(md, slice);
+
+      static_assert(sub.rank() == 1);
+      static_assert(sub.rank_dynamic() == 1);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_stride>);
+
+      assert(sub.stride(0) == 1);
+      assert(sub.extent(0) == 4);
+      assert(sub.size() == 4);
+      assert(equal_to(sub, "PPER"));
+    }
+
+    { // Slice of elements with strided slice with offset, smaller size and stride 1
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [           x    x    x    x ] offset
+      // [           x    x           ] size
+      // [           x    x           ] stride
+      // [           x    x           ]
+      cuda::std::strided_slice slice{2, md.extent(0) - 4, 1};
+      cuda::std::mdspan sub = cuda::std::submdspan(md, slice);
+
+      static_assert(sub.rank() == 1);
+      static_assert(sub.rank_dynamic() == 1);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_stride>);
+
+      assert(sub.stride(0) == 1);
+      assert(sub.extent(0) == 2);
+      assert(sub.size() == 2);
+      assert(equal_to(sub, "PP"));
+    }
+
+    { // Slice of elements with strided slice without offset, full size and stride 3
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [ x    x    x    x    x    x ] offset
+      // [ x    x    x    x    x    x ] size
+      // [ x              x           ] stride
+      // [ x              x           ]
+      cuda::std::strided_slice slice{0, md.extent(0), 3};
+      cuda::std::mdspan sub = cuda::std::submdspan(md, slice);
+
+      static_assert(sub.rank() == 1);
+      static_assert(sub.rank_dynamic() == 1);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_stride>);
+
+      assert(sub.stride(0) == 3);
+      assert(sub.extent(0) == 2);
+      assert(sub.size() == 2);
+      assert(equal_to(sub, "HP"));
+    }
+
+    { // Slice of elements with strided slice with offset, full size and stride 3
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [      x    x    x    x    x ] offset
+      // [      x    x    x    x    x ] size
+      // [      x              x      ] stride
+      // [      x              x      ]
+      cuda::std::strided_slice slice{1, md.extent(0) - 1, 3};
+      cuda::std::mdspan sub = cuda::std::submdspan(md, slice);
+
+      static_assert(sub.rank() == 1);
+      static_assert(sub.rank_dynamic() == 1);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_stride>);
+
+      assert(sub.stride(0) == 3);
+      assert(sub.extent(0) == 2);
+      assert(sub.size() == 2);
+      assert(equal_to(sub, "OE"));
+    }
+
+    { // Slice of elements with strided slice with offset, size less equal than stride
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [      x    x    x    x    x ] offset
+      // [      x    x    x           ] size
+      // [      x                     ] stride
+      // [      x                     ]
+      cuda::std::strided_slice slice{1, 3, 3};
+      cuda::std::mdspan sub = cuda::std::submdspan(md, slice);
+
+      static_assert(sub.rank() == 1);
+      static_assert(sub.rank_dynamic() == 1);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_stride>);
+
+      assert(sub.stride(0) == 1);
+      assert(sub.extent(0) == 1);
+      assert(sub.size() == 1);
+      assert(equal_to(sub, "O"));
+    }
+
+    { // Single element, with integral constant
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [                x           ]
+      cuda::std::mdspan sub = cuda::std::submdspan(md, cuda::std::integral_constant<size_t, 3>{});
+
+      static_assert(sub.rank() == 0);
+      static_assert(sub.rank_dynamic() == 0);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_left>);
+
+      assert(sub.size() == 1);
+      assert(equal_to(sub, "P"));
+    }
+  }
+
+  { // 2d mdspan
+    // ['H', 'P', 'E']
+    // ['O', 'P', 'R']
+    cuda::std::mdspan md{data, cuda::std::layout_left::mapping{cuda::std::dims<2>{2, 3}}};
+    static_assert(md.rank() == 2);
+    static_assert(md.rank_dynamic() == 2);
+    assert(equal_to(md, {"HPE", "OPR"}));
+
+    { // full_extent
+      // ['H', 'P', 'E']
+      // ['O', 'P', 'R']
+      // [ x    x    x ]
+      // [ x    x    x ]
+      cuda::std::mdspan sub = cuda::std::submdspan(md, cuda::std::full_extent, cuda::std::full_extent);
+
+      static_assert(sub.rank() == 2);
+      static_assert(sub.rank_dynamic() == 2);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_left>);
+
+      assert(sub.stride(0) == 1);
+      assert(sub.stride(1) == sub.extent(0));
+      assert(sub.extent(0) == md.extent(0));
+      assert(sub.extent(1) == md.extent(1));
+      assert(sub.size() == md.size());
+      assert(equal_to(sub, {"HPE", "OPR"}));
+    }
+  }
+  return true;
+}
+
+int main(int, char**)
+{
+  test();
+  // static_assert(test(), "");
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/layout_right.pass.cpp b/libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/layout_right.pass.cpp
new file mode 100644
index 0000000000..edfe478a96
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/layout_right.pass.cpp
@@ -0,0 +1,283 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: c++11, c++14
+
+// <mdspan>
+
+// constexpr mdspan& operator=(const mdspan& rhs) = default;
+
+#include <cuda/std/cassert>
+#include <cuda/std/mdspan>
+#include <cuda/std/type_traits>
+
+#include "helper.h"
+#include "test_macros.h"
+
+__host__ __device__ constexpr bool test()
+{
+  constexpr char data[] = {'H', 'O', 'P', 'P', 'E', 'R'};
+
+  { // 1d mdspan
+    // ['H', 'O', 'P', 'P', 'E', 'R']
+    cuda::std::mdspan md{data, cuda::std::layout_right::mapping{cuda::std::dims<1>{6}}};
+    static_assert(md.rank() == 1);
+    static_assert(md.rank_dynamic() == 1);
+    assert(equal_to(md, "HOPPER"));
+
+    using mdspan_t = decltype(md);
+    static_assert(cuda::std::is_same_v<typename mdspan_t::layout_type, cuda::std::layout_right>);
+
+    { // full_extent
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [ x    x    x    x    x    x ]
+      cuda::std::mdspan sub = cuda::std::submdspan(md, cuda::std::full_extent);
+
+      static_assert(sub.rank() == 1);
+      static_assert(sub.rank_dynamic() == 1);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_right>);
+
+      assert(sub.stride(0) == 1);
+      assert(sub.extent(0) == 6);
+      assert(sub.size() == 6);
+      assert(equal_to(sub, "HOPPER"));
+    }
+
+    { // Slice of elements from start 0:4
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [ x    x    x    x           ]
+      cuda::std::mdspan sub = cuda::std::submdspan(md, cuda::std::pair{0, 4});
+
+      static_assert(sub.rank() == 1);
+      static_assert(sub.rank_dynamic() == 1);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_right>);
+
+      assert(sub.stride(0) == 1);
+      assert(sub.extent(0) == 4);
+      assert(sub.size() == 4);
+      assert(equal_to(sub, "HOPP"));
+    }
+
+    { // Slice of elements in the middle 2:5
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [           x    x    x      ]
+      cuda::std::mdspan sub = cuda::std::submdspan(md, cuda::std::pair{2, 5});
+
+      static_assert(sub.rank() == 1);
+      static_assert(sub.rank_dynamic() == 1);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_right>);
+
+      assert(sub.stride(0) == 1);
+      assert(sub.extent(0) == 3);
+      assert(sub.size() == 3);
+      assert(equal_to(sub, "PPE"));
+    }
+
+    { // Slice of elements in the end 3:6
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [                x    x    x ]
+      cuda::std::mdspan sub = cuda::std::submdspan(md, cuda::std::pair{3, 6});
+
+      static_assert(sub.rank() == 1);
+      static_assert(sub.rank_dynamic() == 1);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_right>);
+
+      assert(sub.stride(0) == 1);
+      assert(sub.extent(0) == 3);
+      assert(sub.size() == 3);
+      assert(equal_to(sub, "PER"));
+    }
+
+    { // Slice of elements with strided slice without offset, full size and stride 1
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [ x    x    x    x    x    x ] offset
+      // [ x    x    x    x    x    x ] size
+      // [ x    x    x    x    x    x ] stride
+      // [ x    x    x    x    x    x ]
+      cuda::std::strided_slice slice{0, md.extent(0), 1};
+      cuda::std::mdspan sub = cuda::std::submdspan(md, slice);
+
+      static_assert(sub.rank() == 1);
+      static_assert(sub.rank_dynamic() == 1);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_stride>);
+
+      assert(sub.stride(0) == 1);
+      assert(sub.extent(0) == 6);
+      assert(sub.size() == 6);
+      assert(equal_to(sub, "HOPPER"));
+    }
+
+    { // Slice of elements with strided slice with offset, full remaing size and stride 1
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [           x    x    x    x ] offset
+      // [           x    x    x    x ] size
+      // [           x    x    x    x ] stride
+      // [           x    x    x    x ]
+      cuda::std::strided_slice slice{2, md.extent(0) - 2, 1};
+      cuda::std::mdspan sub = cuda::std::submdspan(md, slice);
+
+      static_assert(sub.rank() == 1);
+      static_assert(sub.rank_dynamic() == 1);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_stride>);
+
+      assert(sub.stride(0) == 1);
+      assert(sub.extent(0) == 4);
+      assert(sub.size() == 4);
+      assert(equal_to(sub, "PPER"));
+    }
+
+    { // Slice of elements with strided slice with offset, smaller size and stride 1
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [           x    x    x    x ] offset
+      // [           x    x           ] size
+      // [           x    x           ] stride
+      // [           x    x           ]
+      cuda::std::strided_slice slice{2, md.extent(0) - 4, 1};
+      cuda::std::mdspan sub = cuda::std::submdspan(md, slice);
+
+      static_assert(sub.rank() == 1);
+      static_assert(sub.rank_dynamic() == 1);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_stride>);
+
+      assert(sub.stride(0) == 1);
+      assert(sub.extent(0) == 2);
+      assert(sub.size() == 2);
+      assert(equal_to(sub, "PP"));
+    }
+
+    { // Slice of elements with strided slice without offset, full size and stride 3
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [ x    x    x    x    x    x ] offset
+      // [ x    x    x    x    x    x ] size
+      // [ x              x           ] stride
+      // [ x              x           ]
+      cuda::std::strided_slice slice{0, md.extent(0), 3};
+      cuda::std::mdspan sub = cuda::std::submdspan(md, slice);
+
+      static_assert(sub.rank() == 1);
+      static_assert(sub.rank_dynamic() == 1);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_stride>);
+
+      assert(sub.stride(0) == 3);
+      assert(sub.extent(0) == 2);
+      assert(sub.size() == 2);
+      assert(equal_to(sub, "HP"));
+    }
+
+    { // Slice of elements with strided slice with offset, full size and stride 3
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [      x    x    x    x    x ] offset
+      // [      x    x    x    x    x ] size
+      // [      x              x      ] stride
+      // [      x              x      ]
+      cuda::std::strided_slice slice{1, md.extent(0) - 1, 3};
+      cuda::std::mdspan sub = cuda::std::submdspan(md, slice);
+
+      static_assert(sub.rank() == 1);
+      static_assert(sub.rank_dynamic() == 1);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_stride>);
+
+      assert(sub.stride(0) == 3);
+      assert(sub.extent(0) == 2);
+      assert(sub.size() == 2);
+      assert(equal_to(sub, "OE"));
+    }
+
+    { // Slice of elements with strided slice with offset, size less equal than stride
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [      x    x    x    x    x ] offset
+      // [      x    x    x           ] size
+      // [      x                     ] stride
+      // [      x                     ]
+      cuda::std::strided_slice slice{1, 3, 3};
+      cuda::std::mdspan sub = cuda::std::submdspan(md, slice);
+
+      static_assert(sub.rank() == 1);
+      static_assert(sub.rank_dynamic() == 1);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_stride>);
+
+      assert(sub.stride(0) == 1);
+      assert(sub.extent(0) == 1);
+      assert(sub.size() == 1);
+      assert(equal_to(sub, "O"));
+    }
+
+    { // Single element, with integral constant
+      // ['H', 'O', 'P', 'P', 'E', 'R']
+      // [                x           ]
+      cuda::std::mdspan sub = cuda::std::submdspan(md, cuda::std::integral_constant<size_t, 3>{});
+
+      static_assert(sub.rank() == 0);
+      static_assert(sub.rank_dynamic() == 0);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_right>);
+
+      assert(sub.size() == 1);
+      assert(equal_to(sub, "P"));
+    }
+  }
+
+  { // 2d mdspan
+    // ['H', 'O', 'P']
+    // ['P', 'E', 'R']
+    cuda::std::mdspan md{data, cuda::std::layout_right::mapping{cuda::std::dims<2>{2, 3}}};
+    static_assert(md.rank() == 2);
+    static_assert(md.rank_dynamic() == 2);
+    assert(equal_to(md, {"HOP", "PER"}));
+
+    { // full_extent
+      // ['H', 'O', 'P'], ['P', 'E', 'R']
+      // [ x    x    x ], [ x    x    x ]
+      cuda::std::mdspan sub = cuda::std::submdspan(md, cuda::std::full_extent, cuda::std::full_extent);
+
+      static_assert(sub.rank() == 2);
+      static_assert(sub.rank_dynamic() == 2);
+
+      using submdspan_t = decltype(sub);
+      static_assert(cuda::std::is_same_v<typename submdspan_t::layout_type, cuda::std::layout_right>);
+
+      assert(sub.stride(0) == sub.extent(1));
+      assert(sub.stride(1) == 1);
+      assert(sub.extent(0) == md.extent(0));
+      assert(sub.extent(1) == md.extent(1));
+      assert(sub.size() == md.size());
+      assert(equal_to(sub, {"HOP", "PER"}));
+    }
+  }
+
+  return true;
+}
+
+int main(int, char**)
+{
+  test();
+  // static_assert(test(), "");
+  return 0;
+}