From f98f43fada0a599f9a1d3e7219fde422286fe4e8 Mon Sep 17 00:00:00 2001
From: Michael Schellenberger Costa <miscco@nvidia.com>
Date: Mon, 14 Nov 2022 08:20:30 +0100
Subject: [PATCH] Port `std::span` and enable if for C++11 onwards to support
 mdspan

Notably this does not take the fairly outdated state of our libc++ fork, but upgrades
it to the actually approved state that became C++20.

The only missing pieces are ranges support, but we can happily do that once we support ranges.
---
 .../pairs/pairs.pair/default-sfinae.pass.cpp  |  16 +-
 .../views/views.span/span.cons/array.fail.cpp |  73 ++++
 .../views/views.span/span.cons/array.pass.cpp | 119 ++++++
 .../views.span/span.cons/assign.pass.cpp      | 263 ++++++++++++
 .../views/views.span/span.cons/copy.pass.cpp  |  67 +++
 .../views.span/span.cons/deduct.pass.cpp      |  99 +++++
 .../views.span/span.cons/default.fail.cpp     |  30 ++
 .../views.span/span.cons/default.pass.cpp     |  89 ++++
 .../span.cons/initializer_list.pass.cpp       |  49 +++
 .../span.cons/iterator_len.pass.cpp           |  76 ++++
 .../span.cons/iterator_len.verify.cpp         |  39 ++
 .../span.cons/iterator_sentinel.pass.cpp      |  65 +++
 .../span.cons/iterator_sentinel.verify.cpp    |  35 ++
 .../views/views.span/span.cons/range.pass.cpp | 144 +++++++
 .../views.span/span.cons/range.verify.cpp     | 110 +++++
 .../span.cons/span.dtor.compile.pass.cpp      |  35 ++
 .../views/views.span/span.cons/span.fail.cpp  | 115 +++++
 .../views/views.span/span.cons/span.pass.cpp  | 120 ++++++
 .../views.span/span.cons/stdarray.pass.cpp    | 124 ++++++
 .../views/views.span/span.elem/back.pass.cpp  |  83 ++++
 .../views/views.span/span.elem/data.pass.cpp  | 112 +++++
 .../views/views.span/span.elem/front.pass.cpp |  84 ++++
 .../views.span/span.elem/op_idx.pass.cpp      | 110 +++++
 .../views.span/span.iterators/begin.pass.cpp  |  98 +++++
 .../views.span/span.iterators/end.pass.cpp    | 102 +++++
 .../views.span/span.iterators/rbegin.pass.cpp | 103 +++++
 .../views.span/span.iterators/rend.pass.cpp   | 102 +++++
 .../span.objectrep/as_bytes.pass.cpp          |  73 ++++
 .../span.objectrep/as_writable_bytes.pass.cpp |  73 ++++
 .../as_writable_bytes.verify.cpp              |  47 +++
 .../span.obs/empty.nodiscard.verify.cpp       |  30 ++
 .../views/views.span/span.obs/empty.pass.cpp  |  65 +++
 .../views/views.span/span.obs/size.pass.cpp   |  84 ++++
 .../views.span/span.obs/size_bytes.pass.cpp   |  85 ++++
 .../views/views.span/span.sub/first.pass.cpp  | 110 +++++
 .../views.span/span.sub/first.verify.cpp      |  36 ++
 .../views/views.span/span.sub/last.pass.cpp   | 109 +++++
 .../views/views.span/span.sub/last.verify.cpp |  36 ++
 .../views.span/span.sub/subspan.pass.cpp      | 208 +++++++++
 .../views.span/span.sub/subspan.verify.cpp    |  42 ++
 .../trivially_copyable.compile.pass.cpp       |  24 ++
 .../test/std/views/views.span/types.pass.cpp  |  95 +++++
 .upstream-tests/test/support/test_macros.h    |  17 +
 .../cuda/std/detail/libcxx/include/iterator   |   4 +-
 include/cuda/std/detail/libcxx/include/span   | 394 ++++++++----------
 .../std/detail/libcxx/include/type_traits     |   1 +
 .../cuda/std/detail/libcxx/include/version    |   1 +
 include/cuda/std/span                         |  29 ++
 48 files changed, 3795 insertions(+), 230 deletions(-)
 create mode 100644 .upstream-tests/test/std/views/views.span/span.cons/array.fail.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.cons/array.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.cons/assign.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.cons/copy.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.cons/deduct.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.cons/default.fail.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.cons/default.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.cons/initializer_list.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.cons/iterator_len.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.cons/iterator_len.verify.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.cons/iterator_sentinel.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.cons/iterator_sentinel.verify.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.cons/range.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.cons/range.verify.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.cons/span.dtor.compile.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.cons/span.fail.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.cons/span.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.cons/stdarray.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.elem/back.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.elem/data.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.elem/front.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.elem/op_idx.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.iterators/begin.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.iterators/end.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.iterators/rbegin.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.iterators/rend.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.objectrep/as_bytes.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.objectrep/as_writable_bytes.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.objectrep/as_writable_bytes.verify.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.obs/empty.nodiscard.verify.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.obs/empty.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.obs/size.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.obs/size_bytes.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.sub/first.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.sub/first.verify.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.sub/last.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.sub/last.verify.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.sub/subspan.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/span.sub/subspan.verify.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/trivially_copyable.compile.pass.cpp
 create mode 100644 .upstream-tests/test/std/views/views.span/types.pass.cpp
 create mode 100644 include/cuda/std/span

diff --git a/.upstream-tests/test/std/utilities/utility/pairs/pairs.pair/default-sfinae.pass.cpp b/.upstream-tests/test/std/utilities/utility/pairs/pairs.pair/default-sfinae.pass.cpp
index a33c1b079a..13d948e919 100644
--- a/.upstream-tests/test/std/utilities/utility/pairs/pairs.pair/default-sfinae.pass.cpp
+++ b/.upstream-tests/test/std/utilities/utility/pairs/pairs.pair/default-sfinae.pass.cpp
@@ -22,14 +22,6 @@
 
 #include "test_macros.h"
 
-#if TEST_STD_VER > 11
-#define CONSTEXPR_CXX14 constexpr
-#define STATIC_ASSERT_CXX14(Pred) static_assert(Pred, "")
-#else
-#define CONSTEXPR_CXX14
-#define STATIC_ASSERT_CXX14(Pred) assert(Pred)
-#endif
-
 struct DeletedDefault {
     // A class with a deleted default constructor. Used to test the SFINAE
     // on cuda::std::pair's default constructor.
@@ -121,20 +113,20 @@ __host__ __device__ void test_illformed_default()
     {
     typedef cuda::std::pair<IllFormedDefault, int> P;
     static_assert((cuda::std::is_constructible<P, IllFormedDefault, int>::value), "");
-    CONSTEXPR_CXX14 P p(IllFormedDefault(42), -5);
+    TEST_CONSTEXPR_CXX14 P p(IllFormedDefault(42), -5);
     STATIC_ASSERT_CXX14(p.first.value == 42 && p.second == -5);
     }
     {
     typedef cuda::std::pair<int, IllFormedDefault> P;
     static_assert((cuda::std::is_constructible<P, int, IllFormedDefault>::value), "");
-    CONSTEXPR_CXX14 IllFormedDefault dd(-5);
-    CONSTEXPR_CXX14 P p(42, dd);
+    TEST_CONSTEXPR_CXX14 IllFormedDefault dd(-5);
+    TEST_CONSTEXPR_CXX14 P p(42, dd);
     STATIC_ASSERT_CXX14(p.first == 42 && p.second.value == -5);
     }
     {
     typedef cuda::std::pair<IllFormedDefault, IllFormedDefault> P;
     static_assert((cuda::std::is_constructible<P, IllFormedDefault, IllFormedDefault>::value), "");
-    CONSTEXPR_CXX14 P p(IllFormedDefault(42), IllFormedDefault(-5));
+    TEST_CONSTEXPR_CXX14 P p(IllFormedDefault(42), IllFormedDefault(-5));
     STATIC_ASSERT_CXX14(p.first.value == 42 && p.second.value == -5);
     }
 }
diff --git a/.upstream-tests/test/std/views/views.span/span.cons/array.fail.cpp b/.upstream-tests/test/std/views/views.span/span.cons/array.fail.cpp
new file mode 100644
index 0000000000..eab06786b7
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.cons/array.fail.cpp
@@ -0,0 +1,73 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+// UNSUPPORTED: nvrtc
+
+// <span>
+
+// template<size_t N>
+//     constexpr span(element_type (&arr)[N]) noexcept;
+// template<size_t N>
+//     constexpr span(array<value_type, N>& arr) noexcept;
+// template<size_t N>
+//     constexpr span(const array<value_type, N>& arr) noexcept;
+//
+// Remarks: These constructors shall not participate in overload resolution unless:
+//   — extent == dynamic_extent || N == extent is true, and
+//   — remove_pointer_t<decltype(data(arr))>(*)[] is convertible to ElementType(*)[].
+//
+
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+__device__                int   arr[] = {1,2,3};
+__device__ const          int  carr[] = {4,5,6};
+__device__       volatile int  varr[] = {7,8,9};
+__device__ const volatile int cvarr[] = {1,3,5};
+
+int main(int, char**)
+{
+//  Size wrong
+    {
+    cuda::std::span<int, 2>   s1(arr); // expected-error {{no matching constructor for initialization of 'cuda::std::span<int, 2>'}}
+    }
+
+//  Type wrong
+    {
+    cuda::std::span<float>    s1(arr);   // expected-error {{no matching constructor for initialization of 'cuda::std::span<float>'}}
+    cuda::std::span<float, 3> s2(arr);   // expected-error {{no matching constructor for initialization of 'cuda::std::span<float, 3>'}}
+    }
+
+//  CV wrong (dynamically sized)
+    {
+    cuda::std::span<               int> s1{ carr};    // expected-error {{no matching constructor for initialization of 'cuda::std::span<int>'}}
+    cuda::std::span<               int> s2{ varr};    // expected-error {{no matching constructor for initialization of 'cuda::std::span<int>'}}
+    cuda::std::span<               int> s3{cvarr};    // expected-error {{no matching constructor for initialization of 'cuda::std::span<int>'}}
+    cuda::std::span<const          int> s4{ varr};    // expected-error {{no matching constructor for initialization of 'cuda::std::span<const int>'}}
+    cuda::std::span<const          int> s5{cvarr};    // expected-error {{no matching constructor for initialization of 'cuda::std::span<const int>'}}
+    cuda::std::span<      volatile int> s6{ carr};    // expected-error {{no matching constructor for initialization of 'cuda::std::span<volatile int>'}}
+    cuda::std::span<      volatile int> s7{cvarr};    // expected-error {{no matching constructor for initialization of 'cuda::std::span<volatile int>'}}
+    }
+
+//  CV wrong (statically sized)
+    {
+    cuda::std::span<               int,3> s1{ carr};  // expected-error {{no matching constructor for initialization of 'cuda::std::span<int, 3>'}}
+    cuda::std::span<               int,3> s2{ varr};  // expected-error {{no matching constructor for initialization of 'cuda::std::span<int, 3>'}}
+    cuda::std::span<               int,3> s3{cvarr};  // expected-error {{no matching constructor for initialization of 'cuda::std::span<int, 3>'}}
+    cuda::std::span<const          int,3> s4{ varr};  // expected-error {{no matching constructor for initialization of 'cuda::std::span<const int, 3>'}}
+    cuda::std::span<const          int,3> s5{cvarr};  // expected-error {{no matching constructor for initialization of 'cuda::std::span<const int, 3>'}}
+    cuda::std::span<      volatile int,3> s6{ carr};  // expected-error {{no matching constructor for initialization of 'cuda::std::span<volatile int, 3>'}}
+    cuda::std::span<      volatile int,3> s7{cvarr};  // expected-error {{no matching constructor for initialization of 'cuda::std::span<volatile int, 3>'}}
+    }
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.cons/array.pass.cpp b/.upstream-tests/test/std/views/views.span/span.cons/array.pass.cpp
new file mode 100644
index 0000000000..dc9d9415ff
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.cons/array.pass.cpp
@@ -0,0 +1,119 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// template<size_t N>
+//     constexpr span(element_type (&arr)[N]) noexcept;
+//
+// Remarks: These constructors shall not participate in overload resolution unless:
+//   — extent == dynamic_extent || N == extent is true, and
+//   — remove_pointer_t<decltype(data(arr))>(*)[] is convertible to ElementType(*)[].
+//
+
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+__device__                int   arr[] = {1,2,3};
+__device__ const          int  carr[] = {4,5,6};
+__device__       volatile int  varr[] = {7,8,9};
+__device__ const volatile int cvarr[] = {1,3,5};
+
+__host__ __device__ void checkCV()
+{
+//  Types the same (dynamic sized)
+    {
+    cuda::std::span<               int> s1{  arr};    // a cuda::std::span<               int> pointing at int.
+    cuda::std::span<const          int> s2{ carr};    // a cuda::std::span<const          int> pointing at const int.
+    cuda::std::span<      volatile int> s3{ varr};    // a cuda::std::span<      volatile int> pointing at volatile int.
+    cuda::std::span<const volatile int> s4{cvarr};    // a cuda::std::span<const volatile int> pointing at const volatile int.
+    assert(s1.size() + s2.size() + s3.size() + s4.size() == 12);
+    }
+
+//  Types the same (static sized)
+    {
+    cuda::std::span<               int,3> s1{  arr};  // a cuda::std::span<               int> pointing at int.
+    cuda::std::span<const          int,3> s2{ carr};  // a cuda::std::span<const          int> pointing at const int.
+    cuda::std::span<      volatile int,3> s3{ varr};  // a cuda::std::span<      volatile int> pointing at volatile int.
+    cuda::std::span<const volatile int,3> s4{cvarr};  // a cuda::std::span<const volatile int> pointing at const volatile int.
+    assert(s1.size() + s2.size() + s3.size() + s4.size() == 12);
+    }
+
+
+//  types different (dynamic sized)
+    {
+    cuda::std::span<const          int> s1{ arr};     // a cuda::std::span<const          int> pointing at int.
+    cuda::std::span<      volatile int> s2{ arr};     // a cuda::std::span<      volatile int> pointing at int.
+    cuda::std::span<      volatile int> s3{ arr};     // a cuda::std::span<      volatile int> pointing at const int.
+    cuda::std::span<const volatile int> s4{ arr};     // a cuda::std::span<const volatile int> pointing at int.
+    cuda::std::span<const volatile int> s5{carr};     // a cuda::std::span<const volatile int> pointing at const int.
+    cuda::std::span<const volatile int> s6{varr};     // a cuda::std::span<const volatile int> pointing at volatile int.
+    assert(s1.size() + s2.size() + s3.size() + s4.size() + s5.size() + s6.size() == 18);
+    }
+
+//  types different (static sized)
+    {
+    cuda::std::span<const          int,3> s1{ arr};   // a cuda::std::span<const          int> pointing at int.
+    cuda::std::span<      volatile int,3> s2{ arr};   // a cuda::std::span<      volatile int> pointing at int.
+    cuda::std::span<      volatile int,3> s3{ arr};   // a cuda::std::span<      volatile int> pointing at const int.
+    cuda::std::span<const volatile int,3> s4{ arr};   // a cuda::std::span<const volatile int> pointing at int.
+    cuda::std::span<const volatile int,3> s5{carr};   // a cuda::std::span<const volatile int> pointing at const int.
+    cuda::std::span<const volatile int,3> s6{varr};   // a cuda::std::span<const volatile int> pointing at volatile int.
+    assert(s1.size() + s2.size() + s3.size() + s4.size() + s5.size() + s6.size() == 18);
+    }
+}
+
+template<class T>
+__host__ __device__ constexpr bool testSpan()
+{
+    T val[2] = {};
+
+    ASSERT_NOEXCEPT(cuda::std::span<T>{val});
+    ASSERT_NOEXCEPT(cuda::std::span<T, 2>{val});
+    ASSERT_NOEXCEPT(cuda::std::span<const T>{val});
+    ASSERT_NOEXCEPT(cuda::std::span<const T, 2>{val});
+
+    cuda::std::span<T> s1 = val;
+    cuda::std::span<T, 2> s2 = val;
+    cuda::std::span<const T> s3 = val;
+    cuda::std::span<const T, 2> s4 = val;
+    assert(s1.data() == val && s1.size() == 2);
+    assert(s2.data() == val && s2.size() == 2);
+    assert(s3.data() == val && s3.size() == 2);
+    assert(s4.data() == val && s4.size() == 2);
+
+    cuda::std::span<const int> s5 = {{1,2}};
+    cuda::std::span<const int, 2> s6 = {{1,2}};
+    assert(s5.size() == 2);  // and it dangles
+    assert(s6.size() == 2);  // and it dangles
+
+    return true;
+}
+
+
+struct A {};
+
+int main(int, char**)
+{
+    testSpan<int>();
+    testSpan<double>();
+    testSpan<A>();
+
+    static_assert(testSpan<int>(), "");
+    static_assert(testSpan<double>(), "");
+    static_assert(testSpan<A>(), "");
+
+    checkCV();
+
+    return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.cons/assign.pass.cpp b/.upstream-tests/test/std/views/views.span/span.cons/assign.pass.cpp
new file mode 100644
index 0000000000..2dfb25cd8b
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.cons/assign.pass.cpp
@@ -0,0 +1,263 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+//  constexpr span& operator=(const span& other) noexcept = default;
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+#include <cuda/std/iterator>
+#include <cuda/std/utility>
+
+#include "test_macros.h"
+
+using cuda::std::span;
+
+template <typename T>
+__host__ __device__
+constexpr bool doAssign(T lhs, T rhs)
+{
+    ASSERT_NOEXCEPT(cuda::std::declval<T&>() = rhs);
+    lhs = rhs;
+    return lhs.data() == rhs.data()
+     &&    lhs.size() == rhs.size();
+}
+
+struct A{};
+
+__device__ constexpr int carr1[] = {1,2,3,4};
+__device__ constexpr int carr2[] = {3,4,5};
+__device__ constexpr int carr3[] = {7,8};
+__device__           int   arr[] = {5,6,7,9};
+
+int main(int, char**)
+{
+
+//  constexpr dynamically sized assignment
+    {
+//  On systems where 'ptrdiff_t' is a synonym for 'int',
+//  the call span(ptr, 0) selects the (pointer, size_type) constructor.
+//  On systems where 'ptrdiff_t' is NOT a synonym for 'int',
+//  it is ambiguous, because of 0 also being convertible to a null pointer
+//  and so the compiler can't choose between:
+//      span(pointer, size_type)
+//  and span(pointer, pointer)
+//  We cast zero to std::ptrdiff_t to remove that ambiguity.
+//  Example:
+//      On darwin x86_64, ptrdiff_t is the same as long int.
+//      On darwin i386, ptrdiff_t is the same as int.
+        constexpr cuda::std::span<const int> spans[] = {
+            {},
+            {carr1, static_cast<cuda::std::size_t>(0)},
+            {carr1,     1U},
+            {carr1,     2U},
+            {carr1,     3U},
+            {carr1,     4U},
+            {carr2, static_cast<cuda::std::size_t>(0)},
+            {carr2,     1U},
+            {carr2,     2U},
+            {carr2,     3U},
+            {carr3, static_cast<cuda::std::size_t>(0)},
+            {carr3,     1U},
+            {carr3,     2U}
+            };
+        //static_assert(cuda::std::size(spans) == 13 );
+
+//  No for loops in constexpr land :-(
+        STATIC_ASSERT_CXX14(doAssign(spans[0], spans[0]));
+        STATIC_ASSERT_CXX14(doAssign(spans[0], spans[1]));
+        STATIC_ASSERT_CXX14(doAssign(spans[0], spans[2]));
+        STATIC_ASSERT_CXX14(doAssign(spans[0], spans[3]));
+        STATIC_ASSERT_CXX14(doAssign(spans[0], spans[4]));
+        STATIC_ASSERT_CXX14(doAssign(spans[0], spans[5]));
+        STATIC_ASSERT_CXX14(doAssign(spans[0], spans[6]));
+        STATIC_ASSERT_CXX14(doAssign(spans[0], spans[7]));
+        STATIC_ASSERT_CXX14(doAssign(spans[0], spans[8]));
+        STATIC_ASSERT_CXX14(doAssign(spans[0], spans[9]));
+        STATIC_ASSERT_CXX14(doAssign(spans[0], spans[10]));
+        STATIC_ASSERT_CXX14(doAssign(spans[0], spans[11]));
+        STATIC_ASSERT_CXX14(doAssign(spans[0], spans[12]));
+
+        STATIC_ASSERT_CXX14(doAssign(spans[1], spans[1]));
+        STATIC_ASSERT_CXX14(doAssign(spans[1], spans[2]));
+        STATIC_ASSERT_CXX14(doAssign(spans[1], spans[3]));
+        STATIC_ASSERT_CXX14(doAssign(spans[1], spans[4]));
+        STATIC_ASSERT_CXX14(doAssign(spans[1], spans[5]));
+        STATIC_ASSERT_CXX14(doAssign(spans[1], spans[6]));
+        STATIC_ASSERT_CXX14(doAssign(spans[1], spans[7]));
+        STATIC_ASSERT_CXX14(doAssign(spans[1], spans[8]));
+        STATIC_ASSERT_CXX14(doAssign(spans[1], spans[9]));
+        STATIC_ASSERT_CXX14(doAssign(spans[1], spans[10]));
+        STATIC_ASSERT_CXX14(doAssign(spans[1], spans[11]));
+        STATIC_ASSERT_CXX14(doAssign(spans[1], spans[12]));
+
+        STATIC_ASSERT_CXX14(doAssign(spans[2], spans[2]));
+        STATIC_ASSERT_CXX14(doAssign(spans[2], spans[3]));
+        STATIC_ASSERT_CXX14(doAssign(spans[2], spans[4]));
+        STATIC_ASSERT_CXX14(doAssign(spans[2], spans[5]));
+        STATIC_ASSERT_CXX14(doAssign(spans[2], spans[6]));
+        STATIC_ASSERT_CXX14(doAssign(spans[2], spans[7]));
+        STATIC_ASSERT_CXX14(doAssign(spans[2], spans[8]));
+        STATIC_ASSERT_CXX14(doAssign(spans[2], spans[9]));
+        STATIC_ASSERT_CXX14(doAssign(spans[2], spans[10]));
+        STATIC_ASSERT_CXX14(doAssign(spans[2], spans[11]));
+        STATIC_ASSERT_CXX14(doAssign(spans[2], spans[12]));
+
+        STATIC_ASSERT_CXX14(doAssign(spans[3], spans[3]));
+        STATIC_ASSERT_CXX14(doAssign(spans[3], spans[4]));
+        STATIC_ASSERT_CXX14(doAssign(spans[3], spans[4]));
+        STATIC_ASSERT_CXX14(doAssign(spans[3], spans[4]));
+        STATIC_ASSERT_CXX14(doAssign(spans[3], spans[4]));
+        STATIC_ASSERT_CXX14(doAssign(spans[3], spans[4]));
+        STATIC_ASSERT_CXX14(doAssign(spans[3], spans[4]));
+        STATIC_ASSERT_CXX14(doAssign(spans[3], spans[4]));
+        STATIC_ASSERT_CXX14(doAssign(spans[3], spans[4]));
+        STATIC_ASSERT_CXX14(doAssign(spans[3], spans[10]));
+        STATIC_ASSERT_CXX14(doAssign(spans[3], spans[11]));
+        STATIC_ASSERT_CXX14(doAssign(spans[3], spans[12]));
+
+        STATIC_ASSERT_CXX14(doAssign(spans[4], spans[4]));
+        STATIC_ASSERT_CXX14(doAssign(spans[4], spans[5]));
+        STATIC_ASSERT_CXX14(doAssign(spans[4], spans[6]));
+        STATIC_ASSERT_CXX14(doAssign(spans[4], spans[7]));
+        STATIC_ASSERT_CXX14(doAssign(spans[4], spans[8]));
+        STATIC_ASSERT_CXX14(doAssign(spans[4], spans[9]));
+        STATIC_ASSERT_CXX14(doAssign(spans[4], spans[10]));
+        STATIC_ASSERT_CXX14(doAssign(spans[4], spans[11]));
+        STATIC_ASSERT_CXX14(doAssign(spans[4], spans[12]));
+
+        STATIC_ASSERT_CXX14(doAssign(spans[5], spans[5]));
+        STATIC_ASSERT_CXX14(doAssign(spans[5], spans[6]));
+        STATIC_ASSERT_CXX14(doAssign(spans[5], spans[7]));
+        STATIC_ASSERT_CXX14(doAssign(spans[5], spans[8]));
+        STATIC_ASSERT_CXX14(doAssign(spans[5], spans[9]));
+        STATIC_ASSERT_CXX14(doAssign(spans[5], spans[10]));
+        STATIC_ASSERT_CXX14(doAssign(spans[5], spans[11]));
+        STATIC_ASSERT_CXX14(doAssign(spans[5], spans[12]));
+
+        STATIC_ASSERT_CXX14(doAssign(spans[6], spans[6]));
+        STATIC_ASSERT_CXX14(doAssign(spans[6], spans[7]));
+        STATIC_ASSERT_CXX14(doAssign(spans[6], spans[8]));
+        STATIC_ASSERT_CXX14(doAssign(spans[6], spans[9]));
+        STATIC_ASSERT_CXX14(doAssign(spans[6], spans[10]));
+        STATIC_ASSERT_CXX14(doAssign(spans[6], spans[11]));
+        STATIC_ASSERT_CXX14(doAssign(spans[6], spans[12]));
+
+        STATIC_ASSERT_CXX14(doAssign(spans[7], spans[7]));
+        STATIC_ASSERT_CXX14(doAssign(spans[7], spans[8]));
+        STATIC_ASSERT_CXX14(doAssign(spans[7], spans[9]));
+        STATIC_ASSERT_CXX14(doAssign(spans[7], spans[10]));
+        STATIC_ASSERT_CXX14(doAssign(spans[7], spans[11]));
+        STATIC_ASSERT_CXX14(doAssign(spans[7], spans[12]));
+
+        STATIC_ASSERT_CXX14(doAssign(spans[8], spans[8]));
+        STATIC_ASSERT_CXX14(doAssign(spans[8], spans[9]));
+        STATIC_ASSERT_CXX14(doAssign(spans[8], spans[10]));
+        STATIC_ASSERT_CXX14(doAssign(spans[8], spans[11]));
+        STATIC_ASSERT_CXX14(doAssign(spans[8], spans[12]));
+
+        STATIC_ASSERT_CXX14(doAssign(spans[9], spans[9]));
+        STATIC_ASSERT_CXX14(doAssign(spans[9], spans[10]));
+        STATIC_ASSERT_CXX14(doAssign(spans[9], spans[11]));
+        STATIC_ASSERT_CXX14(doAssign(spans[9], spans[12]));
+
+        STATIC_ASSERT_CXX14(doAssign(spans[10], spans[10]));
+        STATIC_ASSERT_CXX14(doAssign(spans[10], spans[11]));
+        STATIC_ASSERT_CXX14(doAssign(spans[10], spans[12]));
+
+        STATIC_ASSERT_CXX14(doAssign(spans[11], spans[11]));
+        STATIC_ASSERT_CXX14(doAssign(spans[11], spans[12]));
+
+        STATIC_ASSERT_CXX14(doAssign(spans[12], spans[12]));
+
+//      for (size_t i = 0; i < cuda::std::size(spans); ++i)
+//          for (size_t j = i; j < cuda::std::size(spans); ++j)
+//              static_assert(doAssign(spans[i], spans[j]));
+    }
+
+//  constexpr statically sized assignment
+    {
+        using spanType = cuda::std::span<const int,2>;
+        
+        constexpr spanType spans[] = {
+            spanType{carr1, 2},
+            spanType{carr1 + 1, 2},
+            spanType{carr1 + 2, 2},
+            spanType{carr2, 2},
+            spanType{carr2 + 1, 2},
+            spanType{carr3, 2}
+            };
+        STATIC_ASSERT_CXX14(cuda::std::size(spans) == 6 );
+
+        //  No for loops in constexpr land :-(
+        STATIC_ASSERT_CXX14(doAssign(spans[0], spans[0]));
+        STATIC_ASSERT_CXX14(doAssign(spans[0], spans[1]));
+        STATIC_ASSERT_CXX14(doAssign(spans[0], spans[2]));
+        STATIC_ASSERT_CXX14(doAssign(spans[0], spans[3]));
+        STATIC_ASSERT_CXX14(doAssign(spans[0], spans[4]));
+        STATIC_ASSERT_CXX14(doAssign(spans[0], spans[5]));
+
+        STATIC_ASSERT_CXX14(doAssign(spans[1], spans[1]));
+        STATIC_ASSERT_CXX14(doAssign(spans[1], spans[2]));
+        STATIC_ASSERT_CXX14(doAssign(spans[1], spans[3]));
+        STATIC_ASSERT_CXX14(doAssign(spans[1], spans[4]));
+        STATIC_ASSERT_CXX14(doAssign(spans[1], spans[5]));
+
+        STATIC_ASSERT_CXX14(doAssign(spans[2], spans[2]));
+        STATIC_ASSERT_CXX14(doAssign(spans[2], spans[3]));
+        STATIC_ASSERT_CXX14(doAssign(spans[2], spans[4]));
+        STATIC_ASSERT_CXX14(doAssign(spans[2], spans[5]));
+
+        STATIC_ASSERT_CXX14(doAssign(spans[3], spans[3]));
+        STATIC_ASSERT_CXX14(doAssign(spans[3], spans[4]));
+        STATIC_ASSERT_CXX14(doAssign(spans[3], spans[5]));
+
+        STATIC_ASSERT_CXX14(doAssign(spans[4], spans[4]));
+        STATIC_ASSERT_CXX14(doAssign(spans[4], spans[5]));
+
+        STATIC_ASSERT_CXX14(doAssign(spans[5], spans[5]));
+
+//      for (size_t i = 0; i < cuda::std::size(spans); ++i)
+//          for (size_t j = i; j < cuda::std::size(spans); ++j)
+//              static_assert(doAssign(spans[i], spans[j]));
+    }
+
+//  dynamically sized assignment
+    {
+        cuda::std::span<int> spans[] = {
+            {},
+            {arr,     arr + 1},
+            {arr,     arr + 2},
+            {arr,     arr + 3},
+            {arr + 1, arr + 3} // same size as s2
+            };
+
+        for (size_t i = 0; i < 5; ++i)
+            for (size_t j = i; j < 5; ++j)
+                assert((doAssign(spans[i], spans[j])));
+    }
+
+//  statically sized assignment
+    {
+        using spanType = cuda::std::span<int,2>;
+        spanType spans[] = {
+            spanType{arr,     arr + 2},
+            spanType{arr + 1, arr + 3},
+            spanType{arr + 2, arr + 4}
+            };
+
+        for (size_t i = 0; i < 3; ++i)
+            for (size_t j = i; j < 3; ++j)
+                assert((doAssign(spans[i], spans[j])));
+    }
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.cons/copy.pass.cpp b/.upstream-tests/test/std/views/views.span/span.cons/copy.pass.cpp
new file mode 100644
index 0000000000..748e0f3728
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.cons/copy.pass.cpp
@@ -0,0 +1,67 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+//  constexpr span(const span& other) noexcept = default;
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+template <typename T>
+__host__ __device__
+constexpr bool doCopy(const T &rhs)
+{
+    ASSERT_NOEXCEPT(T{rhs});
+    T lhs{rhs};
+    return lhs.data() == rhs.data()
+     &&    lhs.size() == rhs.size();
+}
+
+struct A{};
+
+template <typename T>
+__host__ __device__
+void testCV ()
+{
+    int  arr[] = {1,2,3};
+    assert((doCopy(cuda::std::span<T>  ()          )));
+    assert((doCopy(cuda::std::span<T,0>()          )));
+    assert((doCopy(cuda::std::span<T>  (&arr[0], 1))));
+    assert((doCopy(cuda::std::span<T,1>(&arr[0], 1))));
+    assert((doCopy(cuda::std::span<T>  (&arr[0], 2))));
+    assert((doCopy(cuda::std::span<T,2>(&arr[0], 2))));
+}
+
+__device__ constexpr int carr[] = {1,2,3};
+
+int main(int, char**)
+{
+
+    STATIC_ASSERT_CXX14(doCopy(cuda::std::span<      int>  ()));
+    STATIC_ASSERT_CXX14(doCopy(cuda::std::span<      int,0>()));
+    STATIC_ASSERT_CXX14(doCopy(cuda::std::span<const int>  (&carr[0], 1)));
+    STATIC_ASSERT_CXX14(doCopy(cuda::std::span<const int,1>(&carr[0], 1)));
+    STATIC_ASSERT_CXX14(doCopy(cuda::std::span<const int>  (&carr[0], 2)));
+    STATIC_ASSERT_CXX14(doCopy(cuda::std::span<const int,2>(&carr[0], 2)));
+
+    STATIC_ASSERT_CXX14(doCopy(cuda::std::span<long>()));
+    STATIC_ASSERT_CXX14(doCopy(cuda::std::span<double>()));
+    STATIC_ASSERT_CXX14(doCopy(cuda::std::span<A>()));
+
+    testCV<               int>();
+    testCV<const          int>();
+    testCV<      volatile int>();
+    testCV<const volatile int>();
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.cons/deduct.pass.cpp b/.upstream-tests/test/std/views/views.span/span.cons/deduct.pass.cpp
new file mode 100644
index 0000000000..c7ea5fe3e6
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.cons/deduct.pass.cpp
@@ -0,0 +1,99 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11, c++14
+
+// <span>
+
+//   template<class It, class EndOrSize>
+//     span(It, EndOrSize) -> span<remove_reference_t<iter_reference_t<_It>>>;
+//
+//   template<class T, size_t N>
+//     span(T (&)[N]) -> span<T, N>;
+//
+//   template<class T, size_t N>
+//     span(array<T, N>&) -> span<T, N>;
+//
+//   template<class T, size_t N>
+//     span(const array<T, N>&) -> span<const T, N>;
+//
+//   template<class R>
+//     span(R&&) -> span<remove_reference_t<ranges::range_reference_t<R>>>;
+
+
+#include <cuda/std/span>
+#include <cuda/std/array>
+#include <cuda/std/cassert>
+#include <cuda/std/iterator>
+
+#include "test_macros.h"
+
+
+__host__ __device__
+void test_iterator_sentinel() {
+  int arr[] = {1, 2, 3};
+  {
+  cuda::std::span s{cuda::std::begin(arr), cuda::std::end(arr)};
+  ASSERT_SAME_TYPE(decltype(s), cuda::std::span<int>);
+  assert(s.size() == cuda::std::size(arr));
+  assert(s.data() == cuda::std::data(arr));
+  }
+  {
+  cuda::std::span s{cuda::std::begin(arr), 3};
+  ASSERT_SAME_TYPE(decltype(s), cuda::std::span<int>);
+  assert(s.size() == cuda::std::size(arr));
+  assert(s.data() == cuda::std::data(arr));
+  }
+}
+
+__host__ __device__
+void test_c_array() {
+    {
+    int arr[] = {1, 2, 3};
+    cuda::std::span s{arr};
+    ASSERT_SAME_TYPE(decltype(s), cuda::std::span<int, 3>);
+    assert(s.size() == cuda::std::size(arr));
+    assert(s.data() == cuda::std::data(arr));
+    }
+
+    {
+    const int arr[] = {1,2,3};
+    cuda::std::span s{arr};
+    ASSERT_SAME_TYPE(decltype(s), cuda::std::span<const int, 3>);
+    assert(s.size() == cuda::std::size(arr));
+    assert(s.data() == cuda::std::data(arr));
+    }
+}
+
+__host__ __device__
+void test_std_array() {
+    {
+    cuda::std::array<double, 4> arr = {1.0, 2.0, 3.0, 4.0};
+    cuda::std::span s{arr};
+    ASSERT_SAME_TYPE(decltype(s), cuda::std::span<double, 4>);
+    assert(s.size() == arr.size());
+    assert(s.data() == arr.data());
+    }
+
+    {
+    const cuda::std::array<long, 5> arr = {4, 5, 6, 7, 8};
+    cuda::std::span s{arr};
+    ASSERT_SAME_TYPE(decltype(s), cuda::std::span<const long, 5>);
+    assert(s.size() == arr.size());
+    assert(s.data() == arr.data());
+    }
+}
+
+int main(int, char**)
+{
+  test_iterator_sentinel();
+  test_c_array();
+  test_std_array();
+ 
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.cons/default.fail.cpp b/.upstream-tests/test/std/views/views.span/span.cons/default.fail.cpp
new file mode 100644
index 0000000000..6532c7c8c5
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.cons/default.fail.cpp
@@ -0,0 +1,30 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+// UNSUPPORTED: nvrtc
+
+// <span>
+
+// constexpr span() noexcept;
+//
+//  Remarks: This constructor shall not participate in overload resolution
+//          unless Extent == 0 || Extent == dynamic_extent is true.
+
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+int main(int, char**)
+{
+  cuda::std::span<int, 2> s; // expected-error {{no matching constructor for initialization of 'std::span<int, 2>'}}
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.cons/default.pass.cpp b/.upstream-tests/test/std/views/views.span/span.cons/default.pass.cpp
new file mode 100644
index 0000000000..fcb4039431
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.cons/default.pass.cpp
@@ -0,0 +1,89 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// constexpr span() noexcept;
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+__host__ __device__
+void checkCV()
+{
+//  Types the same (dynamic sized)
+    {
+    cuda::std::span<               int> s1;
+    cuda::std::span<const          int> s2;
+    cuda::std::span<      volatile int> s3;
+    cuda::std::span<const volatile int> s4;
+    assert(s1.size() + s2.size() + s3.size() + s4.size() == 0);
+    }
+
+//  Types the same (static sized)
+    {
+    cuda::std::span<               int,0> s1;
+    cuda::std::span<const          int,0> s2;
+    cuda::std::span<      volatile int,0> s3;
+    cuda::std::span<const volatile int,0> s4;
+    assert(s1.size() + s2.size() + s3.size() + s4.size() == 0);
+    }
+}
+
+
+template <typename T>
+__host__ __device__
+constexpr bool testConstexprSpan()
+{
+    cuda::std::span<const T>    s1;
+    cuda::std::span<const T, 0> s2;
+    return
+        s1.data() == nullptr && s1.size() == 0
+    &&  s2.data() == nullptr && s2.size() == 0;
+}
+
+
+template <typename T>
+__host__ __device__
+void testRuntimeSpan()
+{
+    ASSERT_NOEXCEPT(T{});
+    cuda::std::span<const T>    s1;
+    cuda::std::span<const T, 0> s2;
+    assert(s1.data() == nullptr && s1.size() == 0);
+    assert(s2.data() == nullptr && s2.size() == 0);
+}
+
+
+struct A{};
+
+int main(int, char**)
+{
+    STATIC_ASSERT_CXX14(testConstexprSpan<int>());
+    STATIC_ASSERT_CXX14(testConstexprSpan<long>());
+    STATIC_ASSERT_CXX14(testConstexprSpan<double>());
+    STATIC_ASSERT_CXX14(testConstexprSpan<A>());
+
+    testRuntimeSpan<int>();
+    testRuntimeSpan<long>();
+    testRuntimeSpan<double>();
+    testRuntimeSpan<A>();
+
+    checkCV();
+
+    static_assert( cuda::std::is_default_constructible<cuda::std::span<int, cuda::std::dynamic_extent>>::value, "");
+    static_assert( cuda::std::is_default_constructible<cuda::std::span<int, 0>>::value, "");
+    static_assert(!cuda::std::is_default_constructible<cuda::std::span<int, 2>>::value, "");
+
+    return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.cons/initializer_list.pass.cpp b/.upstream-tests/test/std/views/views.span/span.cons/initializer_list.pass.cpp
new file mode 100644
index 0000000000..3d25243326
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.cons/initializer_list.pass.cpp
@@ -0,0 +1,49 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+#include <cuda/std/cstddef>
+
+#include "test_macros.h"
+struct Sink {
+    constexpr Sink() = default;
+    __host__ __device__
+    constexpr Sink(Sink*) {}
+};
+
+__host__ __device__
+constexpr cuda::std::size_t count(cuda::std::span<const Sink> sp) {
+    return sp.size();
+}
+
+template<int N>
+__host__ __device__
+constexpr cuda::std::size_t countn(cuda::std::span<const Sink, N> sp) {
+    return sp.size();
+}
+
+__host__ __device__
+constexpr bool test() {
+    Sink a[10] = {};
+    assert(count({a}) == 10);
+    assert(count({a, a+10}) == 10);
+    assert(countn<10>({a}) == 10);
+    return true;
+}
+
+int main(int, char**) {
+    test();
+    static_assert(test(), "");
+
+    return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.cons/iterator_len.pass.cpp b/.upstream-tests/test/std/views/views.span/span.cons/iterator_len.pass.cpp
new file mode 100644
index 0000000000..d975429dd5
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.cons/iterator_len.pass.cpp
@@ -0,0 +1,76 @@
+//===---------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===---------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// template <class It>
+// constexpr explicit(Extent != dynamic_extent) span(It first, size_type count);
+//  If Extent is not equal to dynamic_extent, then count shall be equal to Extent.
+//
+
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+#include <cuda/std/iterator>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <size_t Extent>
+__host__ __device__
+constexpr bool test_constructibility() {
+  struct Other {};
+  static_assert(cuda::std::is_constructible<cuda::std::span<int, Extent>, int*, size_t>::value, "");
+  static_assert(!cuda::std::is_constructible<cuda::std::span<int, Extent>, const int*, size_t>::value, "");
+  static_assert(cuda::std::is_constructible<cuda::std::span<const int, Extent>, int*, size_t>::value, "");
+  static_assert(cuda::std::is_constructible<cuda::std::span<const int, Extent>, const int*, size_t>::value, "");
+  static_assert(!cuda::std::is_constructible<cuda::std::span<int, Extent>, volatile int*, size_t>::value, "");
+  static_assert(!cuda::std::is_constructible<cuda::std::span<int, Extent>, const volatile int*, size_t>::value, "");
+  static_assert(!cuda::std::is_constructible<cuda::std::span<const int, Extent>, volatile int*, size_t>::value, "");
+  static_assert(!cuda::std::is_constructible<cuda::std::span<const int, Extent>, const volatile int*, size_t>::value, "");
+  static_assert(!cuda::std::is_constructible<cuda::std::span<volatile int, Extent>, const int*, size_t>::value, "");
+  static_assert(!cuda::std::is_constructible<cuda::std::span<volatile int, Extent>, const volatile int*, size_t>::value, "");
+  static_assert(
+      !cuda::std::is_constructible<cuda::std::span<int, Extent>, double*, size_t>::value, ""); // iterator type differs from span type
+  static_assert(!cuda::std::is_constructible<cuda::std::span<int, Extent>, size_t, size_t>::value, "");
+  static_assert(!cuda::std::is_constructible<cuda::std::span<int, Extent>, Other*, size_t>::value, ""); // unrelated iterator type
+
+  return true;
+}
+
+template <class T>
+__host__ __device__
+constexpr bool test_ctor() {
+  T val[2] = {};
+  auto s1 = cuda::std::span<T>(val, 2);
+  auto s2 = cuda::std::span<T, 2>(val, 2);
+  assert(s1.data() == val && s1.size() == 2);
+  assert(s2.data() == val && s2.size() == 2);
+  return true;
+}
+
+__host__ __device__
+constexpr bool test() {
+  test_constructibility<cuda::std::dynamic_extent>();
+  test_constructibility<3>();
+
+  struct A {};
+  test_ctor<int>();
+  test_ctor<A>();
+
+  return true;
+}
+
+int main(int, char**) {
+  test();
+  static_assert(test(), "");
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.cons/iterator_len.verify.cpp b/.upstream-tests/test/std/views/views.span/span.cons/iterator_len.verify.cpp
new file mode 100644
index 0000000000..3136891097
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.cons/iterator_len.verify.cpp
@@ -0,0 +1,39 @@
+//===---------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===---------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+// UNSUPPORTED: nvrtc
+
+// <span>
+
+// template <class It>
+// constexpr explicit(Extent != dynamic_extent) span(It first, size_type count);
+//  If Extent is not equal to dynamic_extent, then count shall be equal to Extent.
+//
+
+#include <cuda/std/span>
+#include <cuda/std/cstddef>
+
+template <class T, size_t extent>
+__host__ __device__
+cuda::std::span<T, extent> createImplicitSpan(T* ptr, size_t len) {
+  return {ptr, len}; // expected-error {{chosen constructor is explicit in copy-initialization}}
+}
+
+int main(int, char**) {
+  // explicit constructor necessary
+  int arr[] = {1, 2, 3};
+  createImplicitSpan<int, 1>(arr, 3);
+
+  cuda::std::span<int> sp = {0, 0}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<int>'}}
+  cuda::std::span<int, 2> sp2 = {0, 0}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<int, 2>'}}
+  cuda::std::span<const int> csp = {0, 0}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<const int>'}}
+  cuda::std::span<const int, 2> csp2 = {0, 0}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<const int, 2>'}}
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.cons/iterator_sentinel.pass.cpp b/.upstream-tests/test/std/views/views.span/span.cons/iterator_sentinel.pass.cpp
new file mode 100644
index 0000000000..fead7b1758
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.cons/iterator_sentinel.pass.cpp
@@ -0,0 +1,65 @@
+//===---------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===---------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <cuda/std/span>
+
+// template <class It, class End>
+// constexpr explicit(Extent != dynamic_extent) span(It first, End last);
+// Requires: [first, last) shall be a valid range.
+//   If Extent is not equal to dynamic_extent, then last - first shall be equal to Extent.
+//
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+#include "test_iterators.h"
+
+template <class T, class Sentinel>
+__host__ __device__ constexpr bool test_ctor() {
+  T val[2] = {};
+  auto s1 = cuda::std::span<T>(cuda::std::begin(val), Sentinel(cuda::std::end(val)));
+  auto s2 = cuda::std::span<T, 2>(cuda::std::begin(val), Sentinel(cuda::std::end(val)));
+  assert(s1.data() == cuda::std::data(val) && s1.size() == cuda::std::size(val));
+  assert(s2.data() == cuda::std::data(val) && s2.size() == cuda::std::size(val));
+  return true;
+}
+
+template <size_t Extent>
+__host__ __device__ constexpr void test_constructibility() {
+  static_assert(cuda::std::is_constructible_v<cuda::std::span<int, Extent>, int*, int*>, "");
+  static_assert(!cuda::std::is_constructible_v<cuda::std::span<int, Extent>, const int*, const int*>, "");
+  static_assert(!cuda::std::is_constructible_v<cuda::std::span<int, Extent>, volatile int*, volatile int*>, "");
+  static_assert(cuda::std::is_constructible_v<cuda::std::span<const int, Extent>, int*, int*>, "");
+  static_assert(cuda::std::is_constructible_v<cuda::std::span<const int, Extent>, const int*, const int*>, "");
+  static_assert(!cuda::std::is_constructible_v<cuda::std::span<const int, Extent>, volatile int*, volatile int*>, "");
+  static_assert(cuda::std::is_constructible_v<cuda::std::span<volatile int, Extent>, int*, int*>, "");
+  static_assert(!cuda::std::is_constructible_v<cuda::std::span<volatile int, Extent>, const int*, const int*>, "");
+  static_assert(cuda::std::is_constructible_v<cuda::std::span<volatile int, Extent>, volatile int*, volatile int*>, "");
+  static_assert(!cuda::std::is_constructible_v<cuda::std::span<int, Extent>, int*, float*>, ""); // types wrong
+}
+
+__host__ __device__ constexpr bool test() {
+  test_constructibility<cuda::std::dynamic_extent>();
+  test_constructibility<3>();
+  struct A {};
+  assert((test_ctor<int, int*>()));
+  //assert((test_ctor<int, sized_sentinel<int*>>()));
+  assert((test_ctor<A, A*>()));
+  //assert((test_ctor<A, sized_sentinel<A*>>()));
+  return true;
+}
+
+int main(int, char**) {
+  test();
+  static_assert(test(), "");
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.cons/iterator_sentinel.verify.cpp b/.upstream-tests/test/std/views/views.span/span.cons/iterator_sentinel.verify.cpp
new file mode 100644
index 0000000000..fd7e003cf2
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.cons/iterator_sentinel.verify.cpp
@@ -0,0 +1,35 @@
+//===---------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===---------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11, c++14, c++17
+
+// <span>
+
+// template <class It, class End>
+// constexpr explicit(Extent != dynamic_extent) span(It first, End last);
+// Requires: [first, last) shall be a valid range.
+//   If Extent is not equal to dynamic_extent, then last - first shall be equal to Extent.
+//
+
+#include <cuda/std/span>
+#include <cuda/std/iterator>
+
+#include "test_macros.h"
+
+template<class T, size_t Extent>
+cuda::std::span<T, Extent> createImplicitSpan(T* first, T* last) {
+    return {first, last}; // expected-error {{chosen constructor is explicit in copy-initialization}}
+}
+
+int main(int, char**) {
+  // explicit constructor necessary
+  int arr[] = {1, 2, 3};
+  createImplicitSpan<int, 1>(cuda::std::begin(arr), cuda::std::end(arr));
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.cons/range.pass.cpp b/.upstream-tests/test/std/views/views.span/span.cons/range.pass.cpp
new file mode 100644
index 0000000000..96c420053d
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.cons/range.pass.cpp
@@ -0,0 +1,144 @@
+//===---------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===---------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <cuda/std/span>
+
+//  template<class Container>
+//    constexpr explicit(Extent != dynamic_extent) span(Container&);
+//  template<class Container>
+//    constexpr explicit(Extent != dynamic_extent) span(Container const&);
+
+// This test checks for libc++'s non-conforming temporary extension to cuda::std::span
+// to support construction from containers that look like contiguous ranges.
+//
+// This extension is only supported when we don't ship <ranges>, and we can
+// remove it once we get rid of _LIBCUDACXX_HAS_NO_INCOMPLETE_RANGES.
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+//  Look ma - I'm a container!
+template <typename T>
+struct IsAContainer {
+    __host__ __device__ constexpr IsAContainer() : v_{} {}
+    __host__ __device__ constexpr size_t size() const {return 1;}
+    __host__ __device__ constexpr       T *data() {return &v_;}
+    __host__ __device__ constexpr const T *data() const {return &v_;}
+    __host__ __device__ constexpr       T *begin() {return &v_;}
+    __host__ __device__ constexpr const T *begin() const {return &v_;}
+    __host__ __device__ constexpr       T *end() {return &v_ + 1;}
+    __host__ __device__ constexpr const T *end() const {return &v_ + 1;}
+
+    __host__ __device__ constexpr T const *getV() const {return &v_;} // for checking
+    T v_;
+};
+
+template <typename T>
+__host__ __device__
+constexpr bool unused(T &&) {return true;}
+
+__host__ __device__ void checkCV()
+{
+    IsAContainer<int> v{};
+
+//  Types the same
+    {
+    cuda::std::span<               int> s1{v};    // a span<               int> pointing at int.
+    unused(s1);
+    }
+
+//  types different
+    {
+    cuda::std::span<const          int> s1{v};    // a span<const          int> pointing at int.
+    cuda::std::span<      volatile int> s2{v};    // a span<      volatile int> pointing at int.
+    cuda::std::span<      volatile int> s3{v};    // a span<      volatile int> pointing at const int.
+    cuda::std::span<const volatile int> s4{v};    // a span<const volatile int> pointing at int.
+    unused(s1);
+    unused(s2);
+    unused(s3);
+    unused(s4);
+    }
+
+//  Constructing a const view from a temporary
+    {
+    cuda::std::span<const int>    s1{IsAContainer<int>()};
+    unused(s1);
+    }
+}
+
+
+template <typename T>
+__host__ __device__ constexpr bool testConstexprSpan()
+{
+    constexpr IsAContainer<const T> val{};
+    cuda::std::span<const T> s1{val};
+    return s1.data() == val.getV() && s1.size() == 1;
+}
+
+template <typename T>
+__host__ __device__ constexpr bool testConstexprSpanStatic()
+{
+    constexpr IsAContainer<const T> val{};
+    cuda::std::span<const T, 1> s1{val};
+    return s1.data() == val.getV() && s1.size() == 1;
+}
+
+template <typename T>
+__host__ __device__ void testRuntimeSpan()
+{
+    IsAContainer<T> val{};
+    const IsAContainer<T> cVal;
+    cuda::std::span<T>       s1{val};
+    cuda::std::span<const T> s2{cVal};
+    assert(s1.data() == val.getV()  && s1.size() == 1);
+    assert(s2.data() == cVal.getV() && s2.size() == 1);
+}
+
+template <typename T>
+__host__ __device__ void testRuntimeSpanStatic()
+{
+    IsAContainer<T> val{};
+    const IsAContainer<T> cVal;
+    cuda::std::span<T, 1>       s1{val};
+    cuda::std::span<const T, 1> s2{cVal};
+    assert(s1.data() == val.getV()  && s1.size() == 1);
+    assert(s2.data() == cVal.getV() && s2.size() == 1);
+}
+
+struct A{};
+
+int main(int, char**)
+{
+    static_assert(testConstexprSpan<int>(),    "");
+    static_assert(testConstexprSpan<long>(),   "");
+    static_assert(testConstexprSpan<double>(), "");
+    static_assert(testConstexprSpan<A>(),      "");
+
+    static_assert(testConstexprSpanStatic<int>(),    "");
+    static_assert(testConstexprSpanStatic<long>(),   "");
+    static_assert(testConstexprSpanStatic<double>(), "");
+    static_assert(testConstexprSpanStatic<A>(),      "");
+
+    testRuntimeSpan<int>();
+    testRuntimeSpan<long>();
+    testRuntimeSpan<double>();
+    testRuntimeSpan<A>();
+
+    testRuntimeSpanStatic<int>();
+    testRuntimeSpanStatic<long>();
+    testRuntimeSpanStatic<double>();
+    testRuntimeSpanStatic<A>();
+
+    checkCV();
+
+    return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.cons/range.verify.cpp b/.upstream-tests/test/std/views/views.span/span.cons/range.verify.cpp
new file mode 100644
index 0000000000..3a93971393
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.cons/range.verify.cpp
@@ -0,0 +1,110 @@
+//===---------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===---------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+// UNSUPPORTED: nvrtc
+
+// <cuda/std/span>
+
+//  template<class Container>
+//    constexpr explicit(Extent != dynamic_extent) span(Container&);
+//  template<class Container>
+//    constexpr explicit(Extent != dynamic_extent) span(Container const&);
+
+// This test checks for libc++'s non-conforming temporary extension to cuda::std::span
+// to support construction from containers that look like contiguous ranges.
+//
+// This extension is only supported when we don't ship <ranges>, and we can
+// remove it once we get rid of _LIBCUDACXX_HAS_NO_INCOMPLETE_RANGES.
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+//  Look ma - I'm a container!
+template <typename T>
+struct IsAContainer {
+    __host__ __device__ constexpr IsAContainer() : v_{} {}
+    __host__ __device__ constexpr size_t size() const {return 1;}
+    __host__ __device__ constexpr       T *data() {return &v_;}
+    __host__ __device__ constexpr const T *data() const {return &v_;}
+
+    __host__ __device__ constexpr const T *getV() const {return &v_;} // for checking
+    T v_;
+};
+
+template <typename T>
+struct NotAContainerNoData {
+    __host__ __device__ size_t size() const {return 0;}
+};
+
+template <typename T>
+struct NotAContainerNoSize {
+    __host__ __device__ const T *data() const {return nullptr;}
+};
+
+template <typename T>
+struct NotAContainerPrivate {
+private:
+    __host__ __device__ size_t size() const {return 0;}
+    __host__ __device__ const T *data() const {return nullptr;}
+};
+
+template<class T, size_t extent, class container>
+__host__ __device__ cuda::std::span<T, extent> createImplicitSpan(container c) {
+    return {c}; // expected-error-re {{no matching constructor for initialization of 'cuda::std::span<{{.+}}>'}}
+}
+
+int main(int, char**)
+{
+
+//  Making non-const spans from const sources (a temporary binds to `const &`)
+    {
+    cuda::std::span<int>    s1{IsAContainer<int>()};          // expected-error {{no matching constructor for initialization of 'cuda::std::span<int>'}}
+    }
+
+//  Missing size and/or data
+    {
+    cuda::std::span<const int>    s1{NotAContainerNoData<int>()};   // expected-error {{no matching constructor for initialization of 'cuda::std::span<const int>'}}
+    cuda::std::span<const int>    s3{NotAContainerNoSize<int>()};   // expected-error {{no matching constructor for initialization of 'cuda::std::span<const int>'}}
+    cuda::std::span<const int>    s5{NotAContainerPrivate<int>()};  // expected-error {{no matching constructor for initialization of 'cuda::std::span<const int>'}}
+    }
+
+//  Not the same type
+    {
+    IsAContainer<int> c;
+    cuda::std::span<float>    s1{c};   // expected-error {{no matching constructor for initialization of 'cuda::std::span<float>'}}
+    }
+
+//  CV wrong
+    {
+    IsAContainer<const          int> c;
+    IsAContainer<const volatile int> cv;
+    IsAContainer<      volatile int> v;
+
+    cuda::std::span<               int> s1{c};    // expected-error {{no matching constructor for initialization of 'cuda::std::span<int>'}}
+    cuda::std::span<               int> s2{v};    // expected-error {{no matching constructor for initialization of 'cuda::std::span<int>'}}
+    cuda::std::span<               int> s3{cv};   // expected-error {{no matching constructor for initialization of 'cuda::std::span<int>'}}
+    cuda::std::span<const          int> s4{v};    // expected-error {{no matching constructor for initialization of 'cuda::std::span<const int>'}}
+    cuda::std::span<const          int> s5{cv};   // expected-error {{no matching constructor for initialization of 'cuda::std::span<const int>'}}
+    cuda::std::span<      volatile int> s6{c};    // expected-error {{no matching constructor for initialization of 'cuda::std::span<volatile int>'}}
+    cuda::std::span<      volatile int> s7{cv};   // expected-error {{no matching constructor for initialization of 'cuda::std::span<volatile int>'}}
+    }
+
+// explicit constructor necessary
+    {
+    IsAContainer<int> c;
+    const IsAContainer<int> cc;
+
+    createImplicitSpan<int, 1>(c);
+    createImplicitSpan<int, 1>(cc);
+    }
+
+    return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.cons/span.dtor.compile.pass.cpp b/.upstream-tests/test/std/views/views.span/span.cons/span.dtor.compile.pass.cpp
new file mode 100644
index 0000000000..524313b194
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.cons/span.dtor.compile.pass.cpp
@@ -0,0 +1,35 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// ~span() = default;
+
+#include <cuda/std/span>
+#include <cuda/std/type_traits>
+
+template <class T>
+__host__ __device__
+constexpr bool testDestructor() {
+  static_assert(cuda::std::is_nothrow_destructible<T>::value, "");
+  static_assert(cuda::std::is_trivially_destructible<T>::value, "");
+  return true;
+}
+
+__host__ __device__
+void test() {
+  testDestructor<cuda::std::span<int, 1>>();
+  testDestructor<cuda::std::span<int>>();
+}
+
+int main(int, char**)
+{
+    return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.cons/span.fail.cpp b/.upstream-tests/test/std/views/views.span/span.cons/span.fail.cpp
new file mode 100644
index 0000000000..4ab16a10b4
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.cons/span.fail.cpp
@@ -0,0 +1,115 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+// UNSUPPORTED: nvrtc
+
+// <span>
+
+// template<class OtherElementType, size_t OtherExtent>
+//    constexpr span(const span<OtherElementType, OtherExtent>& s) noexcept;
+//
+//  Remarks: This constructor shall not participate in overload resolution unless:
+//      Extent == dynamic_extent || Extent == OtherExtent is true, and
+//      OtherElementType(*)[] is convertible to ElementType(*)[].
+
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+template<class T, size_t extent, size_t otherExtent>
+cuda::std::span<T, extent> createImplicitSpan(cuda::std::span<T, otherExtent> s) {
+    return {s}; // expected-error {{chosen constructor is explicit in copy-initialization}}
+}
+
+void checkCV ()
+{
+//  cuda::std::span<               int>   sp;
+    cuda::std::span<const          int>  csp;
+    cuda::std::span<      volatile int>  vsp;
+    cuda::std::span<const volatile int> cvsp;
+
+//  cuda::std::span<               int, 0>   sp0;
+    cuda::std::span<const          int, 0>  csp0;
+    cuda::std::span<      volatile int, 0>  vsp0;
+    cuda::std::span<const volatile int, 0> cvsp0;
+
+//  Try to remove const and/or volatile (dynamic -> dynamic)
+    {
+    cuda::std::span<               int> s1{ csp}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<int>'}}
+    cuda::std::span<               int> s2{ vsp}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<int>'}}
+    cuda::std::span<               int> s3{cvsp}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<int>'}}
+
+    cuda::std::span<const          int> s4{ vsp}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<const int>'}}
+    cuda::std::span<const          int> s5{cvsp}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<const int>'}}
+
+    cuda::std::span<      volatile int> s6{ csp}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<volatile int>'}}
+    cuda::std::span<      volatile int> s7{cvsp}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<volatile int>'}}
+    }
+
+//  Try to remove const and/or volatile (static -> static)
+    {
+    cuda::std::span<               int, 0> s1{ csp0}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<int, 0>'}}
+    cuda::std::span<               int, 0> s2{ vsp0}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<int, 0>'}}
+    cuda::std::span<               int, 0> s3{cvsp0}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<int, 0>'}}
+
+    cuda::std::span<const          int, 0> s4{ vsp0}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<const int, 0>'}}
+    cuda::std::span<const          int, 0> s5{cvsp0}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<const int, 0>'}}
+
+    cuda::std::span<      volatile int, 0> s6{ csp0}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<volatile int, 0>'}}
+    cuda::std::span<      volatile int, 0> s7{cvsp0}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<volatile int, 0>'}}
+    }
+
+//  Try to remove const and/or volatile (static -> dynamic)
+    {
+    cuda::std::span<               int> s1{ csp0}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<int>'}}
+    cuda::std::span<               int> s2{ vsp0}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<int>'}}
+    cuda::std::span<               int> s3{cvsp0}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<int>'}}
+
+    cuda::std::span<const          int> s4{ vsp0}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<const int>'}}
+    cuda::std::span<const          int> s5{cvsp0}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<const int>'}}
+
+    cuda::std::span<      volatile int> s6{ csp0}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<volatile int>'}}
+    cuda::std::span<      volatile int> s7{cvsp0}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<volatile int>'}}
+    }
+
+//  Try to remove const and/or volatile (static -> static)
+    {
+    cuda::std::span<               int, 0> s1{ csp}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<int, 0>'}}
+    cuda::std::span<               int, 0> s2{ vsp}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<int, 0>'}}
+    cuda::std::span<               int, 0> s3{cvsp}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<int, 0>'}}
+
+    cuda::std::span<const          int, 0> s4{ vsp}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<const int, 0>'}}
+    cuda::std::span<const          int, 0> s5{cvsp}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<const int, 0>'}}
+
+    cuda::std::span<      volatile int, 0> s6{ csp}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<volatile int, 0>'}}
+    cuda::std::span<      volatile int, 0> s7{cvsp}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<volatile int, 0>'}}
+    }
+}
+
+int main(int, char**)
+{
+    cuda::std::span<int>      sp;
+    cuda::std::span<int, 0>   sp0;
+
+    cuda::std::span<float> s1{sp};    // expected-error {{no matching constructor for initialization of 'cuda::std::span<float>'}}
+    cuda::std::span<float> s2{sp0};   // expected-error {{no matching constructor for initialization of 'cuda::std::span<float>'}}
+    cuda::std::span<float, 0> s3{sp}; // expected-error {{no matching constructor for initialization of 'cuda::std::span<float, 0>'}}
+    cuda::std::span<float, 0> s4{sp0};    // expected-error {{no matching constructor for initialization of 'cuda::std::span<float, 0>'}}
+
+    checkCV();
+
+    // explicit constructor necessary
+    {
+    createImplicitSpan<int, 1>(sp);
+    }
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.cons/span.pass.cpp b/.upstream-tests/test/std/views/views.span/span.cons/span.pass.cpp
new file mode 100644
index 0000000000..4f8886bb9b
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.cons/span.pass.cpp
@@ -0,0 +1,120 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// template<class OtherElementType, size_t OtherExtent>
+//    constexpr span(const span<OtherElementType, OtherExtent>& s) noexcept;
+//
+//  Remarks: This constructor shall not participate in overload resolution unless:
+//      Extent == dynamic_extent || Extent == OtherExtent is true, and
+//      OtherElementType(*)[] is convertible to ElementType(*)[].
+
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+__host__ __device__
+void checkCV()
+{
+    cuda::std::span<               int>   sp;
+//  cuda::std::span<const          int>  csp;
+    cuda::std::span<      volatile int>  vsp;
+//  cuda::std::span<const volatile int> cvsp;
+
+    cuda::std::span<               int, 0>   sp0;
+//  cuda::std::span<const          int, 0>  csp0;
+    cuda::std::span<      volatile int, 0>  vsp0;
+//  cuda::std::span<const volatile int, 0> cvsp0;
+
+//  dynamic -> dynamic
+    {
+        cuda::std::span<const          int> s1{  sp}; // a cuda::std::span<const          int> pointing at int.
+        cuda::std::span<      volatile int> s2{  sp}; // a cuda::std::span<      volatile int> pointing at int.
+        cuda::std::span<const volatile int> s3{  sp}; // a cuda::std::span<const volatile int> pointing at int.
+        cuda::std::span<const volatile int> s4{ vsp}; // a cuda::std::span<const volatile int> pointing at volatile int.
+        assert(s1.size() + s2.size() + s3.size() + s4.size() == 0);
+    }
+
+//  static -> static
+    {
+        cuda::std::span<const          int, 0> s1{  sp0}; // a cuda::std::span<const          int> pointing at int.
+        cuda::std::span<      volatile int, 0> s2{  sp0}; // a cuda::std::span<      volatile int> pointing at int.
+        cuda::std::span<const volatile int, 0> s3{  sp0}; // a cuda::std::span<const volatile int> pointing at int.
+        cuda::std::span<const volatile int, 0> s4{ vsp0}; // a cuda::std::span<const volatile int> pointing at volatile int.
+        assert(s1.size() + s2.size() + s3.size() + s4.size() == 0);
+    }
+
+//  static -> dynamic
+    {
+        cuda::std::span<const          int> s1{  sp0};    // a cuda::std::span<const          int> pointing at int.
+        cuda::std::span<      volatile int> s2{  sp0};    // a cuda::std::span<      volatile int> pointing at int.
+        cuda::std::span<const volatile int> s3{  sp0};    // a cuda::std::span<const volatile int> pointing at int.
+        cuda::std::span<const volatile int> s4{ vsp0};    // a cuda::std::span<const volatile int> pointing at volatile int.
+        assert(s1.size() + s2.size() + s3.size() + s4.size() == 0);
+    }
+
+//  dynamic -> static (not allowed)
+}
+
+
+template <typename T>
+__host__ __device__
+constexpr bool testConstexprSpan()
+{
+    cuda::std::span<T>    s0{};
+    cuda::std::span<T, 0> s1{};
+    cuda::std::span<T>    s2(s1); // static -> dynamic
+    ASSERT_NOEXCEPT(cuda::std::span<T>   {s0});
+    ASSERT_NOEXCEPT(cuda::std::span<T, 0>{s1});
+    ASSERT_NOEXCEPT(cuda::std::span<T>   {s1});
+
+    return
+        s1.data() == nullptr && s1.size() == 0
+    &&  s2.data() == nullptr && s2.size() == 0;
+}
+
+
+template <typename T>
+__host__ __device__
+void testRuntimeSpan()
+{
+    cuda::std::span<T>    s0{};
+    cuda::std::span<T, 0> s1{};
+    cuda::std::span<T>    s2(s1); // static -> dynamic
+    ASSERT_NOEXCEPT(cuda::std::span<T>   {s0});
+    ASSERT_NOEXCEPT(cuda::std::span<T, 0>{s1});
+    ASSERT_NOEXCEPT(cuda::std::span<T>   {s1});
+
+    assert(s1.data() == nullptr && s1.size() == 0);
+    assert(s2.data() == nullptr && s2.size() == 0);
+}
+
+
+struct A{};
+
+int main(int, char**)
+{
+    STATIC_ASSERT_CXX14(testConstexprSpan<int>());
+    STATIC_ASSERT_CXX14(testConstexprSpan<long>());
+    STATIC_ASSERT_CXX14(testConstexprSpan<double>());
+    STATIC_ASSERT_CXX14(testConstexprSpan<A>());
+
+    testRuntimeSpan<int>();
+    testRuntimeSpan<long>();
+    testRuntimeSpan<double>();
+    testRuntimeSpan<A>();
+
+    checkCV();
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.cons/stdarray.pass.cpp b/.upstream-tests/test/std/views/views.span/span.cons/stdarray.pass.cpp
new file mode 100644
index 0000000000..18ad569007
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.cons/stdarray.pass.cpp
@@ -0,0 +1,124 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// template<size_t N>
+//     constexpr span(array<value_type, N>& arr) noexcept;
+// template<size_t N>
+//     constexpr span(const array<value_type, N>& arr) noexcept;
+//
+// Remarks: These constructors shall not participate in overload resolution unless:
+//   — extent == dynamic_extent || N == extent is true, and
+//   — remove_pointer_t<decltype(data(arr))>(*)[] is convertible to ElementType(*)[].
+//
+
+
+#include <cuda/std/span>
+#include <cuda/std/array>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+__host__ __device__
+void checkCV()
+{
+    cuda::std::array<int, 3> arr  = {1,2,3};
+//  STL says these are not cromulent
+//  std::array<const int,3> carr = {4,5,6};
+//  std::array<volatile int, 3> varr = {7,8,9};
+//  std::array<const volatile int, 3> cvarr = {1,3,5};
+
+//  Types the same (dynamic sized)
+    {
+    cuda::std::span<               int> s1{  arr};    // a cuda::std::span<               int> pointing at int.
+    }
+
+//  Types the same (static sized)
+    {
+    cuda::std::span<               int,3> s1{  arr};  // a cuda::std::span<               int> pointing at int.
+    }
+
+
+//  types different (dynamic sized)
+    {
+    cuda::std::span<const          int> s1{ arr};     // a cuda::std::span<const          int> pointing at int.
+    cuda::std::span<      volatile int> s2{ arr};     // a cuda::std::span<      volatile int> pointing at int.
+    cuda::std::span<      volatile int> s3{ arr};     // a cuda::std::span<      volatile int> pointing at const int.
+    cuda::std::span<const volatile int> s4{ arr};     // a cuda::std::span<const volatile int> pointing at int.
+    }
+
+//  types different (static sized)
+    {
+    cuda::std::span<const          int,3> s1{ arr};   // a cuda::std::span<const          int> pointing at int.
+    cuda::std::span<      volatile int,3> s2{ arr};   // a cuda::std::span<      volatile int> pointing at int.
+    cuda::std::span<      volatile int,3> s3{ arr};   // a cuda::std::span<      volatile int> pointing at const int.
+    cuda::std::span<const volatile int,3> s4{ arr};   // a cuda::std::span<const volatile int> pointing at int.
+    }
+}
+
+template <typename T, typename U>
+__host__ __device__
+TEST_CONSTEXPR_CXX17 bool testConstructorArray() {
+  cuda::std::array<U, 2> val = {U(), U()};
+  ASSERT_NOEXCEPT(cuda::std::span<T>{val});
+  ASSERT_NOEXCEPT(cuda::std::span<T, 2>{val});
+  cuda::std::span<T> s1{val};
+  cuda::std::span<T, 2> s2{val};
+  return s1.data() == &val[0] && s1.size() == 2 && s2.data() == &val[0] &&
+         s2.size() == 2;
+}
+
+template <typename T, typename U>
+__host__ __device__
+TEST_CONSTEXPR_CXX17 bool testConstructorConstArray() {
+  const cuda::std::array<U, 2> val = {U(), U()};
+  ASSERT_NOEXCEPT(cuda::std::span<const T>{val});
+  ASSERT_NOEXCEPT(cuda::std::span<const T, 2>{val});
+  cuda::std::span<const T> s1{val};
+  cuda::std::span<const T, 2> s2{val};
+  return s1.data() == &val[0] && s1.size() == 2 && s2.data() == &val[0] &&
+         s2.size() == 2;
+}
+
+template <typename T>
+__host__ __device__
+TEST_CONSTEXPR_CXX17 bool testConstructors() {
+  STATIC_ASSERT_CXX17((testConstructorArray<T, T>()));
+  STATIC_ASSERT_CXX17((testConstructorArray<const T, const T>()));
+  STATIC_ASSERT_CXX17((testConstructorArray<const T, T>()));
+  STATIC_ASSERT_CXX17((testConstructorConstArray<T, T>()));
+  STATIC_ASSERT_CXX17((testConstructorConstArray<const T, const T>()));
+  STATIC_ASSERT_CXX17((testConstructorConstArray<const T, T>()));
+
+  return testConstructorArray<T, T>() &&
+         testConstructorArray<const T, const T>() &&
+         testConstructorArray<const T, T>() &&
+         testConstructorConstArray<T, T>() &&
+         testConstructorConstArray<const T, const T>() &&
+         testConstructorConstArray<const T, T>();
+}
+
+struct A{};
+
+int main(int, char**)
+{
+    assert(testConstructors<int>());
+    assert(testConstructors<long>());
+    assert(testConstructors<double>());
+    assert(testConstructors<A>());
+
+    assert(testConstructors<int*>());
+    assert(testConstructors<const int*>());
+
+    checkCV();
+
+    return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.elem/back.pass.cpp b/.upstream-tests/test/std/views/views.span/span.elem/back.pass.cpp
new file mode 100644
index 0000000000..a59249df7e
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.elem/back.pass.cpp
@@ -0,0 +1,83 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// constexpr reference back() const noexcept;
+//   Expects: empty() is false.
+//   Effects: Equivalent to: return *(data() + (size() - 1));
+//
+
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+template <typename T>
+__host__ __device__
+constexpr bool unused(T &&) {return true;}
+
+template <typename Span>
+__host__ __device__
+constexpr bool testConstexprSpan(Span sp)
+{
+    ASSERT_NOEXCEPT(sp.back());
+    return &sp.back() == sp.data() + (sp.size() - 1);
+}
+
+template <typename Span>
+__host__ __device__
+void testRuntimeSpan(Span sp)
+{
+    ASSERT_NOEXCEPT(sp.back());
+    assert(&sp.back() == sp.data() + (sp.size() - 1));
+}
+
+template <typename Span>
+__host__ __device__
+void testEmptySpan(Span sp)
+{
+    if (!sp.empty())
+        unused(sp.back());
+}
+
+struct A{};
+__device__ constexpr int iArr1[] = { 0,  1,  2,  3,  4,  5,  6,  7,  8,  9};
+__device__           int iArr2[] = {10, 11, 12, 13, 14, 15, 16, 17, 18, 19};
+
+int main(int, char**)
+{
+    STATIC_ASSERT_CXX14(testConstexprSpan(cuda::std::span<const int>(iArr1, 1)));
+    STATIC_ASSERT_CXX14(testConstexprSpan(cuda::std::span<const int>(iArr1, 2)));
+    STATIC_ASSERT_CXX14(testConstexprSpan(cuda::std::span<const int>(iArr1, 3)));
+    STATIC_ASSERT_CXX14(testConstexprSpan(cuda::std::span<const int>(iArr1, 4)));
+
+    STATIC_ASSERT_CXX14(testConstexprSpan(cuda::std::span<const int, 1>(iArr1, 1)));
+    STATIC_ASSERT_CXX14(testConstexprSpan(cuda::std::span<const int, 2>(iArr1, 2)));
+    STATIC_ASSERT_CXX14(testConstexprSpan(cuda::std::span<const int, 3>(iArr1, 3)));
+    STATIC_ASSERT_CXX14(testConstexprSpan(cuda::std::span<const int, 4>(iArr1, 4)));
+
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 1));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 2));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 3));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 4));
+
+
+    testRuntimeSpan(cuda::std::span<int, 1>(iArr2, 1));
+    testRuntimeSpan(cuda::std::span<int, 2>(iArr2, 2));
+    testRuntimeSpan(cuda::std::span<int, 3>(iArr2, 3));
+    testRuntimeSpan(cuda::std::span<int, 4>(iArr2, 4));
+
+    cuda::std::span<int, 0> sp;
+    testEmptySpan(sp);
+
+    return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.elem/data.pass.cpp b/.upstream-tests/test/std/views/views.span/span.elem/data.pass.cpp
new file mode 100644
index 0000000000..c31e63519f
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.elem/data.pass.cpp
@@ -0,0 +1,112 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// constexpr pointer data() const noexcept;
+//
+
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+template <typename Span>
+__host__ __device__
+constexpr bool testConstexprSpan(Span sp, typename Span::pointer ptr)
+{
+    ASSERT_NOEXCEPT(sp.data());
+    return sp.data() == ptr;
+}
+
+
+template <typename Span>
+__host__ __device__
+void testRuntimeSpan(Span sp, typename Span::pointer ptr)
+{
+    ASSERT_NOEXCEPT(sp.data());
+    assert(sp.data() == ptr);
+}
+
+struct A{};
+__device__ constexpr int iArr1[] = { 0,  1,  2,  3,  4,  5,  6,  7,  8,  9};
+__device__           int iArr2[] = {10, 11, 12, 13, 14, 15, 16, 17, 18, 19};
+
+int main(int, char**)
+{
+
+//  dynamic size
+    static_assert(testConstexprSpan(cuda::std::span<int>(), nullptr),         "");
+    static_assert(testConstexprSpan(cuda::std::span<long>(), nullptr),        "");
+    static_assert(testConstexprSpan(cuda::std::span<double>(), nullptr),      "");
+    static_assert(testConstexprSpan(cuda::std::span<A>(), nullptr),           "");
+
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 1), iArr1), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 2), iArr1), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 3), iArr1), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 4), iArr1), "");
+
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1 + 1, 1), iArr1 + 1), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1 + 2, 2), iArr1 + 2), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1 + 3, 3), iArr1 + 3), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1 + 4, 4), iArr1 + 4), "");
+
+//  static size
+    static_assert(testConstexprSpan(cuda::std::span<int, 0>(), nullptr),         "");
+    static_assert(testConstexprSpan(cuda::std::span<long, 0>(), nullptr),        "");
+    static_assert(testConstexprSpan(cuda::std::span<double, 0>(), nullptr),      "");
+    static_assert(testConstexprSpan(cuda::std::span<A, 0>(), nullptr),           "");
+
+    static_assert(testConstexprSpan(cuda::std::span<const int, 1>(iArr1, 1), iArr1), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int, 2>(iArr1, 2), iArr1), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int, 3>(iArr1, 3), iArr1), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int, 4>(iArr1, 4), iArr1), "");
+
+    static_assert(testConstexprSpan(cuda::std::span<const int, 1>(iArr1 + 1, 1), iArr1 + 1), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int, 2>(iArr1 + 2, 2), iArr1 + 2), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int, 3>(iArr1 + 3, 3), iArr1 + 3), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int, 4>(iArr1 + 4, 4), iArr1 + 4), "");
+
+
+//  dynamic size
+    testRuntimeSpan(cuda::std::span<int>(), nullptr);
+    testRuntimeSpan(cuda::std::span<long>(), nullptr);
+    testRuntimeSpan(cuda::std::span<double>(), nullptr);
+    testRuntimeSpan(cuda::std::span<A>(), nullptr);
+
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 1), iArr2);
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 2), iArr2);
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 3), iArr2);
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 4), iArr2);
+
+    testRuntimeSpan(cuda::std::span<int>(iArr2 + 1, 1), iArr2 + 1);
+    testRuntimeSpan(cuda::std::span<int>(iArr2 + 2, 2), iArr2 + 2);
+    testRuntimeSpan(cuda::std::span<int>(iArr2 + 3, 3), iArr2 + 3);
+    testRuntimeSpan(cuda::std::span<int>(iArr2 + 4, 4), iArr2 + 4);
+
+//  static size
+    testRuntimeSpan(cuda::std::span<int, 0>(), nullptr);
+    testRuntimeSpan(cuda::std::span<long, 0>(), nullptr);
+    testRuntimeSpan(cuda::std::span<double, 0>(), nullptr);
+    testRuntimeSpan(cuda::std::span<A, 0>(), nullptr);
+
+    testRuntimeSpan(cuda::std::span<int, 1>(iArr2, 1), iArr2);
+    testRuntimeSpan(cuda::std::span<int, 2>(iArr2, 2), iArr2);
+    testRuntimeSpan(cuda::std::span<int, 3>(iArr2, 3), iArr2);
+    testRuntimeSpan(cuda::std::span<int, 4>(iArr2, 4), iArr2);
+
+    testRuntimeSpan(cuda::std::span<int, 1>(iArr2 + 1, 1), iArr2 + 1);
+    testRuntimeSpan(cuda::std::span<int, 2>(iArr2 + 2, 2), iArr2 + 2);
+    testRuntimeSpan(cuda::std::span<int, 3>(iArr2 + 3, 3), iArr2 + 3);
+    testRuntimeSpan(cuda::std::span<int, 4>(iArr2 + 4, 4), iArr2 + 4);
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.elem/front.pass.cpp b/.upstream-tests/test/std/views/views.span/span.elem/front.pass.cpp
new file mode 100644
index 0000000000..82db22d17e
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.elem/front.pass.cpp
@@ -0,0 +1,84 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// constexpr reference front() const noexcept;
+//   Expects: empty() is false.
+//   Effects: Equivalent to: return *data();
+//
+
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+template <typename T>
+__host__ __device__
+constexpr bool unused(T &&) {return true;}
+
+template <typename Span>
+__host__ __device__
+constexpr bool testConstexprSpan(Span sp)
+{
+    ASSERT_NOEXCEPT(sp.front());
+    return &sp.front() == sp.data();
+}
+
+
+template <typename Span>
+__host__ __device__
+void testRuntimeSpan(Span sp)
+{
+    ASSERT_NOEXCEPT(sp.front());
+    assert(&sp.front() == sp.data());
+}
+
+template <typename Span>
+__host__ __device__
+void testEmptySpan(Span sp)
+{
+    if (!sp.empty())
+        unused(sp.front());
+}
+
+struct A{};
+__device__ constexpr int iArr1[] = { 0,  1,  2,  3,  4,  5,  6,  7,  8,  9};
+__device__           int iArr2[] = {10, 11, 12, 13, 14, 15, 16, 17, 18, 19};
+
+int main(int, char**)
+{
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 1)), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 2)), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 3)), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 4)), "");
+
+    static_assert(testConstexprSpan(cuda::std::span<const int, 1>(iArr1, 1)), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int, 2>(iArr1, 2)), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int, 3>(iArr1, 3)), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int, 4>(iArr1, 4)), "");
+
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 1));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 2));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 3));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 4));
+
+
+    testRuntimeSpan(cuda::std::span<int, 1>(iArr2, 1));
+    testRuntimeSpan(cuda::std::span<int, 2>(iArr2, 2));
+    testRuntimeSpan(cuda::std::span<int, 3>(iArr2, 3));
+    testRuntimeSpan(cuda::std::span<int, 4>(iArr2, 4));
+
+    cuda::std::span<int, 0> sp;
+    testEmptySpan(sp);
+
+    return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.elem/op_idx.pass.cpp b/.upstream-tests/test/std/views/views.span/span.elem/op_idx.pass.cpp
new file mode 100644
index 0000000000..13aff0a8b9
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.elem/op_idx.pass.cpp
@@ -0,0 +1,110 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// constexpr reference operator[](size_type idx) const;
+//
+
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+template <typename Span>
+__host__ __device__
+constexpr bool testConstexprSpan(Span sp, size_t idx)
+{
+    ASSERT_NOEXCEPT(sp[idx]);
+
+    typename Span::reference r1 = sp[idx];
+    typename Span::reference r2 = *(sp.data() + idx);
+    return r1 == r2;
+}
+
+
+template <typename Span>
+__host__ __device__
+void testRuntimeSpan(Span sp, size_t idx)
+{
+    ASSERT_NOEXCEPT(sp[idx]);
+
+    typename Span::reference r1 = sp[idx];
+    typename Span::reference r2 = *(sp.data() + idx);
+    assert(r1 == r2);
+}
+
+struct A{};
+__device__ constexpr int iArr1[] = { 0,  1,  2,  3,  4,  5,  6,  7,  8,  9};
+__device__           int iArr2[] = {10, 11, 12, 13, 14, 15, 16, 17, 18, 19};
+
+int main(int, char**)
+{
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 1), 0), "");
+
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 2), 0), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 2), 1), "");
+
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 3), 0), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 3), 1), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 3), 2), "");
+
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 4), 0), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 4), 1), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 4), 2), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 4), 3), "");
+
+
+    static_assert(testConstexprSpan(cuda::std::span<const int, 1>(iArr1, 1), 0), "");
+
+    static_assert(testConstexprSpan(cuda::std::span<const int, 2>(iArr1, 2), 0), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int, 2>(iArr1, 2), 1), "");
+
+    static_assert(testConstexprSpan(cuda::std::span<const int, 3>(iArr1, 3), 0), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int, 3>(iArr1, 3), 1), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int, 3>(iArr1, 3), 2), "");
+
+    static_assert(testConstexprSpan(cuda::std::span<const int, 4>(iArr1, 4), 0), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int, 4>(iArr1, 4), 1), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int, 4>(iArr1, 4), 2), "");
+    static_assert(testConstexprSpan(cuda::std::span<const int, 4>(iArr1, 4), 3), "");
+
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 1), 0);
+
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 2), 0);
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 2), 1);
+
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 3), 0);
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 3), 1);
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 3), 2);
+
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 4), 0);
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 4), 1);
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 4), 2);
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 4), 3);
+
+
+    testRuntimeSpan(cuda::std::span<int, 1>(iArr2, 1), 0);
+
+    testRuntimeSpan(cuda::std::span<int, 2>(iArr2, 2), 0);
+    testRuntimeSpan(cuda::std::span<int, 2>(iArr2, 2), 1);
+
+    testRuntimeSpan(cuda::std::span<int, 3>(iArr2, 3), 0);
+    testRuntimeSpan(cuda::std::span<int, 3>(iArr2, 3), 1);
+    testRuntimeSpan(cuda::std::span<int, 3>(iArr2, 3), 2);
+
+    testRuntimeSpan(cuda::std::span<int, 4>(iArr2, 4), 0);
+    testRuntimeSpan(cuda::std::span<int, 4>(iArr2, 4), 1);
+    testRuntimeSpan(cuda::std::span<int, 4>(iArr2, 4), 2);
+    testRuntimeSpan(cuda::std::span<int, 4>(iArr2, 4), 3);
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.iterators/begin.pass.cpp b/.upstream-tests/test/std/views/views.span/span.iterators/begin.pass.cpp
new file mode 100644
index 0000000000..785616ccdf
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.iterators/begin.pass.cpp
@@ -0,0 +1,98 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// constexpr       iterator  begin() const noexcept;
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+template <typename Span>
+__host__ __device__
+constexpr bool testConstexprSpan(Span s)
+{
+    bool ret = true;
+    typename Span::iterator b = s.begin();
+
+    if (s.empty())
+    {
+        ret = ret && (b == s.end());
+    }
+    else
+    {
+        ret = ret && ( *b ==  s[0]);
+        ret = ret && (&*b == &s[0]);
+    }
+    return ret;
+}
+
+
+template <class Span>
+__host__ __device__
+void testRuntimeSpan(Span s)
+{
+    typename Span::iterator b = s.begin();
+
+    if (s.empty())
+    {
+        assert(b == s.end());
+    }
+    else
+    {
+        assert( *b ==  s[0]);
+        assert(&*b == &s[0]);
+    }
+}
+
+struct A{};
+__host__ __device__ bool operator==(A, A) {return true;}
+
+__device__ constexpr int iArr1[] = { 0,  1,  2,  3,  4,  5,  6,  7,  8,  9};
+__device__           int iArr2[] = {10, 11, 12, 13, 14, 15, 16, 17, 18, 19};
+
+int main(int, char**)
+{
+    static_assert(testConstexprSpan(cuda::std::span<int>()),            "");
+    static_assert(testConstexprSpan(cuda::std::span<long>()),           "");
+    static_assert(testConstexprSpan(cuda::std::span<double>()),         "");
+    static_assert(testConstexprSpan(cuda::std::span<A>()),              "");
+
+    static_assert(testConstexprSpan(cuda::std::span<int, 0>()),         "");
+    static_assert(testConstexprSpan(cuda::std::span<long, 0>()),        "");
+    static_assert(testConstexprSpan(cuda::std::span<double, 0>()),      "");
+    static_assert(testConstexprSpan(cuda::std::span<A, 0>()),           "");
+
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 1)),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 2)),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 3)),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 4)),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 5)),    "");
+
+    testRuntimeSpan(cuda::std::span<int>        ());
+    testRuntimeSpan(cuda::std::span<long>       ());
+    testRuntimeSpan(cuda::std::span<double>     ());
+    testRuntimeSpan(cuda::std::span<A>          ());
+
+    testRuntimeSpan(cuda::std::span<int, 0>     ());
+    testRuntimeSpan(cuda::std::span<long, 0>    ());
+    testRuntimeSpan(cuda::std::span<double, 0>  ());
+    testRuntimeSpan(cuda::std::span<A, 0>       ());
+
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 1));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 2));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 3));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 4));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 5));
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.iterators/end.pass.cpp b/.upstream-tests/test/std/views/views.span/span.iterators/end.pass.cpp
new file mode 100644
index 0000000000..b10c7e7e22
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.iterators/end.pass.cpp
@@ -0,0 +1,102 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// constexpr       iterator  end() const noexcept;
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+template <typename Span>
+__host__ __device__
+constexpr bool testConstexprSpan(Span s)
+{
+    bool ret = true;
+    typename Span::iterator e = s.end();
+    if (s.empty())
+    {
+        ret = ret &&  (e == s.begin());
+    }
+    else
+    {
+        typename Span::const_pointer last = &*(s.begin() + s.size() - 1);
+        ret = ret &&  (e != s.begin());
+        ret = ret &&  (&*( e-1) == last);
+    }
+
+    ret = ret &&  (static_cast<size_t>(e - s.begin()) == s.size());
+    return ret;
+}
+
+template <class Span>
+__host__ __device__
+void testRuntimeSpan(Span s)
+{
+    typename Span::iterator e = s.end();
+    if (s.empty())
+    {
+        assert(e == s.begin());
+    }
+    else
+    {
+        typename Span::const_pointer last = &*(s.begin() + s.size() - 1);
+        assert(e != s.begin());
+        assert(&*( e-1) == last);
+    }
+
+    assert(static_cast<size_t>(e - s.begin()) == s.size());
+}
+
+
+struct A{};
+__host__ __device__ bool operator==(A, A) {return true;}
+
+__device__ constexpr int iArr1[] = { 0,  1,  2,  3,  4,  5,  6,  7,  8,  9};
+__device__           int iArr2[] = {10, 11, 12, 13, 14, 15, 16, 17, 18, 19};
+
+int main(int, char**)
+{
+    static_assert(testConstexprSpan(cuda::std::span<int>()),            "");
+    static_assert(testConstexprSpan(cuda::std::span<long>()),           "");
+    static_assert(testConstexprSpan(cuda::std::span<double>()),         "");
+    static_assert(testConstexprSpan(cuda::std::span<A>()),              "");
+
+    static_assert(testConstexprSpan(cuda::std::span<int, 0>()),         "");
+    static_assert(testConstexprSpan(cuda::std::span<long, 0>()),        "");
+    static_assert(testConstexprSpan(cuda::std::span<double, 0>()),      "");
+    static_assert(testConstexprSpan(cuda::std::span<A, 0>()),           "");
+
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 1)),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 2)),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 3)),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 4)),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 5)),    "");
+
+    testRuntimeSpan(cuda::std::span<int>        ());
+    testRuntimeSpan(cuda::std::span<long>       ());
+    testRuntimeSpan(cuda::std::span<double>     ());
+    testRuntimeSpan(cuda::std::span<A>          ());
+
+    testRuntimeSpan(cuda::std::span<int, 0>        ());
+    testRuntimeSpan(cuda::std::span<long, 0>       ());
+    testRuntimeSpan(cuda::std::span<double, 0>     ());
+    testRuntimeSpan(cuda::std::span<A, 0>          ());
+
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 1));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 2));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 3));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 4));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 5));
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.iterators/rbegin.pass.cpp b/.upstream-tests/test/std/views/views.span/span.iterators/rbegin.pass.cpp
new file mode 100644
index 0000000000..b8a934dd2d
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.iterators/rbegin.pass.cpp
@@ -0,0 +1,103 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// constexpr       reverse_iterator  rbegin() const noexcept;
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+template <typename Span>
+__host__ __device__
+TEST_CONSTEXPR_CXX17 bool testConstexprSpan(Span s)
+{
+    bool ret = true;
+    typename Span::reverse_iterator b = s.rbegin();
+    if (s.empty())
+    {
+        ret = ret &&  ( b ==  s.rend());
+    }
+    else
+    {
+        const typename Span::size_type last = s.size() - 1;
+        ret = ret && ( *b ==  s[last]);
+        ret = ret && (&*b == &s[last]);
+    }
+    return ret;
+}
+
+
+template <class Span>
+__host__ __device__
+void testRuntimeSpan(Span s)
+{
+    typename Span::reverse_iterator b = s.rbegin();
+    if (s.empty())
+    {
+        assert(b == s.rend());
+    }
+    else
+    {
+        const typename Span::size_type last = s.size() - 1;
+        assert( *b ==  s[last]);
+        assert(&*b == &s[last]);
+    }
+}
+
+
+struct A{};
+__host__ __device__ bool operator==(A, A) {return true;}
+
+#if TEST_STD_VER > 14
+__device__ constexpr int iArr1[] = { 0,  1,  2,  3,  4,  5,  6,  7,  8,  9};
+#endif
+__device__           int iArr2[] = {10, 11, 12, 13, 14, 15, 16, 17, 18, 19};
+
+int main(int, char**)
+{
+#if TEST_STD_VER > 14
+    static_assert(testConstexprSpan(cuda::std::span<int>()),            "");
+    static_assert(testConstexprSpan(cuda::std::span<long>()),           "");
+    static_assert(testConstexprSpan(cuda::std::span<double>()),         "");
+    static_assert(testConstexprSpan(cuda::std::span<A>()),              "");
+
+    static_assert(testConstexprSpan(cuda::std::span<int, 0>()),         "");
+    static_assert(testConstexprSpan(cuda::std::span<long, 0>()),        "");
+    static_assert(testConstexprSpan(cuda::std::span<double, 0>()),      "");
+    static_assert(testConstexprSpan(cuda::std::span<A, 0>()),           "");
+
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 1)),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 2)),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 3)),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 4)),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 5)),    "");
+#endif // TEST_STD_VER > 14
+
+    testRuntimeSpan(cuda::std::span<int>        ());
+    testRuntimeSpan(cuda::std::span<long>       ());
+    testRuntimeSpan(cuda::std::span<double>     ());
+    testRuntimeSpan(cuda::std::span<A>          ());
+
+    testRuntimeSpan(cuda::std::span<int, 0>        ());
+    testRuntimeSpan(cuda::std::span<long, 0>       ());
+    testRuntimeSpan(cuda::std::span<double, 0>     ());
+    testRuntimeSpan(cuda::std::span<A, 0>          ());
+
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 1));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 2));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 3));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 4));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 5));
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.iterators/rend.pass.cpp b/.upstream-tests/test/std/views/views.span/span.iterators/rend.pass.cpp
new file mode 100644
index 0000000000..a749bc29f3
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.iterators/rend.pass.cpp
@@ -0,0 +1,102 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// constexpr       reverse_iterator  rend() const noexcept;
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+template <typename Span>
+__host__ __device__
+constexpr bool testConstexprSpan(Span s)
+{
+    bool ret = true;
+    typename Span::reverse_iterator e = s.rend();
+    if (s.empty())
+    {
+        ret = ret &&  (e == s.rbegin());
+    }
+    else
+    {
+        ret = ret &&  (e != s.rbegin());
+    }
+
+    ret = ret &&  (static_cast<size_t>(e - s.rbegin()) == s.size());
+    return ret;
+}
+
+template <class Span>
+__host__ __device__
+void testRuntimeSpan(Span s)
+{
+    typename Span::reverse_iterator e = s.rend();
+    if (s.empty())
+    {
+        assert(e == s.rbegin());
+    }
+    else
+    {
+        assert(e != s.rbegin());
+    }
+
+    assert(static_cast<size_t>(e - s.rbegin()) == s.size());
+}
+
+
+struct A{};
+__host__ __device__ bool operator==(A, A) {return true;}
+
+#if TEST_STD_VER > 14
+__device__ constexpr int iArr1[] = { 0,  1,  2,  3,  4,  5,  6,  7,  8,  9};
+#endif
+__device__           int iArr2[] = {10, 11, 12, 13, 14, 15, 16, 17, 18, 19};
+
+int main(int, char**)
+{
+#if TEST_STD_VER > 14
+    static_assert(testConstexprSpan(cuda::std::span<int>()),            "");
+    static_assert(testConstexprSpan(cuda::std::span<long>()),           "");
+    static_assert(testConstexprSpan(cuda::std::span<double>()),         "");
+    static_assert(testConstexprSpan(cuda::std::span<A>()),              "");
+
+    static_assert(testConstexprSpan(cuda::std::span<int, 0>()),         "");
+    static_assert(testConstexprSpan(cuda::std::span<long, 0>()),        "");
+    static_assert(testConstexprSpan(cuda::std::span<double, 0>()),      "");
+    static_assert(testConstexprSpan(cuda::std::span<A, 0>()),           "");
+
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 1)),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 2)),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 3)),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 4)),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 5)),    "");
+#endif // TEST_STD_VER > 14
+
+    testRuntimeSpan(cuda::std::span<int>        ());
+    testRuntimeSpan(cuda::std::span<long>       ());
+    testRuntimeSpan(cuda::std::span<double>     ());
+    testRuntimeSpan(cuda::std::span<A>          ());
+
+    testRuntimeSpan(cuda::std::span<int, 0>        ());
+    testRuntimeSpan(cuda::std::span<long, 0>       ());
+    testRuntimeSpan(cuda::std::span<double, 0>     ());
+    testRuntimeSpan(cuda::std::span<A, 0>          ());
+
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 1));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 2));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 3));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 4));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 5));
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.objectrep/as_bytes.pass.cpp b/.upstream-tests/test/std/views/views.span/span.objectrep/as_bytes.pass.cpp
new file mode 100644
index 0000000000..5b7d1a9ac1
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.objectrep/as_bytes.pass.cpp
@@ -0,0 +1,73 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// template <class ElementType, size_t Extent>
+//     span<const byte,
+//          Extent == dynamic_extent
+//              ? dynamic_extent
+//              : sizeof(ElementType) * Extent>
+//     as_bytes(span<ElementType, Extent> s) noexcept;
+
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+template <typename Span>
+__host__ __device__
+void testRuntimeSpan(Span sp)
+{
+    ASSERT_NOEXCEPT(cuda::std::as_bytes(sp));
+
+    auto spBytes = cuda::std::as_bytes(sp);
+    using SB = decltype(spBytes);
+    ASSERT_SAME_TYPE(const cuda::std::byte, typename SB::element_type);
+
+    if (sp.extent == cuda::std::dynamic_extent)
+        assert(spBytes.extent == cuda::std::dynamic_extent);
+    else
+        assert(spBytes.extent == sizeof(typename Span::element_type) * sp.extent);
+
+    assert((void *) spBytes.data() == (void *) sp.data());
+    assert(spBytes.size() == sp.size_bytes());
+}
+
+struct A{};
+__device__ int iArr2[] = { 0,  1,  2,  3,  4,  5,  6,  7,  8,  9};
+
+int main(int, char**)
+{
+    testRuntimeSpan(cuda::std::span<int>        ());
+    testRuntimeSpan(cuda::std::span<long>       ());
+    testRuntimeSpan(cuda::std::span<double>     ());
+    testRuntimeSpan(cuda::std::span<A>          ());
+
+    testRuntimeSpan(cuda::std::span<int, 0>        ());
+    testRuntimeSpan(cuda::std::span<long, 0>       ());
+    testRuntimeSpan(cuda::std::span<double, 0>     ());
+    testRuntimeSpan(cuda::std::span<A, 0>          ());
+
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 1));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 2));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 3));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 4));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 5));
+
+    testRuntimeSpan(cuda::std::span<int, 1>(iArr2 + 5, 1));
+    testRuntimeSpan(cuda::std::span<int, 2>(iArr2 + 4, 2));
+    testRuntimeSpan(cuda::std::span<int, 3>(iArr2 + 3, 3));
+    testRuntimeSpan(cuda::std::span<int, 4>(iArr2 + 2, 4));
+    testRuntimeSpan(cuda::std::span<int, 5>(iArr2 + 1, 5));
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.objectrep/as_writable_bytes.pass.cpp b/.upstream-tests/test/std/views/views.span/span.objectrep/as_writable_bytes.pass.cpp
new file mode 100644
index 0000000000..9e61ffe3f7
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.objectrep/as_writable_bytes.pass.cpp
@@ -0,0 +1,73 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// template <class ElementType, size_t Extent>
+//     span<byte,
+//          Extent == dynamic_extent
+//              ? dynamic_extent
+//              : sizeof(ElementType) * Extent>
+//     as_writable_bytes(span<ElementType, Extent> s) noexcept;
+
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+template<typename Span>
+__host__ __device__
+void testRuntimeSpan(Span sp)
+{
+    ASSERT_NOEXCEPT(cuda::std::as_writable_bytes(sp));
+
+    auto spBytes = cuda::std::as_writable_bytes(sp);
+    using SB = decltype(spBytes);
+    ASSERT_SAME_TYPE(cuda::std::byte, typename SB::element_type);
+
+    if (sp.extent == cuda::std::dynamic_extent)
+        assert(spBytes.extent == cuda::std::dynamic_extent);
+    else
+        assert(spBytes.extent == sizeof(typename Span::element_type) * sp.extent);
+
+    assert(static_cast<void*>(spBytes.data()) == static_cast<void*>(sp.data()));
+    assert(spBytes.size() == sp.size_bytes());
+}
+
+struct A{};
+__device__ int iArr2[] = { 0,  1,  2,  3,  4,  5,  6,  7,  8,  9};
+
+int main(int, char**)
+{
+    testRuntimeSpan(cuda::std::span<int>        ());
+    testRuntimeSpan(cuda::std::span<long>       ());
+    testRuntimeSpan(cuda::std::span<double>     ());
+    testRuntimeSpan(cuda::std::span<A>          ());
+
+    testRuntimeSpan(cuda::std::span<int, 0>        ());
+    testRuntimeSpan(cuda::std::span<long, 0>       ());
+    testRuntimeSpan(cuda::std::span<double, 0>     ());
+    testRuntimeSpan(cuda::std::span<A, 0>          ());
+
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 1));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 2));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 3));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 4));
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 5));
+
+    testRuntimeSpan(cuda::std::span<int, 1>(iArr2 + 5, 1));
+    testRuntimeSpan(cuda::std::span<int, 2>(iArr2 + 4, 2));
+    testRuntimeSpan(cuda::std::span<int, 3>(iArr2 + 3, 3));
+    testRuntimeSpan(cuda::std::span<int, 4>(iArr2 + 2, 4));
+    testRuntimeSpan(cuda::std::span<int, 5>(iArr2 + 1, 5));
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.objectrep/as_writable_bytes.verify.cpp b/.upstream-tests/test/std/views/views.span/span.objectrep/as_writable_bytes.verify.cpp
new file mode 100644
index 0000000000..bc205208b9
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.objectrep/as_writable_bytes.verify.cpp
@@ -0,0 +1,47 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+// UNSUPPORTED: nvrtc
+
+// <span>
+
+// template <class ElementType, size_t Extent>
+//     span<byte,
+//          Extent == dynamic_extent
+//              ? dynamic_extent
+//              : sizeof(ElementType) * Extent>
+//     as_writable_bytes(span<ElementType, Extent> s) noexcept;
+
+#include <cuda/std/span>
+
+#include "test_macros.h"
+
+__device__ constexpr int iArr2[] = { 0,  1,  2,  3,  4,  5,  6,  7,  8,  9};
+
+struct A {};
+
+__host__ __device__ void f() {
+  cuda::std::as_writable_bytes(cuda::std::span<const int>());            // expected-error {{no matching function for call to 'as_writable_bytes'}}
+  cuda::std::as_writable_bytes(cuda::std::span<const long>());           // expected-error {{no matching function for call to 'as_writable_bytes'}}
+  cuda::std::as_writable_bytes(cuda::std::span<const double>());         // expected-error {{no matching function for call to 'as_writable_bytes'}}
+  cuda::std::as_writable_bytes(cuda::std::span<const A>());              // expected-error {{no matching function for call to 'as_writable_bytes'}}
+
+  cuda::std::as_writable_bytes(cuda::std::span<const int, 0>());         // expected-error {{no matching function for call to 'as_writable_bytes'}}
+  cuda::std::as_writable_bytes(cuda::std::span<const long, 0>());        // expected-error {{no matching function for call to 'as_writable_bytes'}}
+  cuda::std::as_writable_bytes(cuda::std::span<const double, 0>());      // expected-error {{no matching function for call to 'as_writable_bytes'}}
+  cuda::std::as_writable_bytes(cuda::std::span<const A, 0>());           // expected-error {{no matching function for call to 'as_writable_bytes'}}
+
+  cuda::std::as_writable_bytes(cuda::std::span<const int>   (iArr2, 1));     // expected-error {{no matching function for call to 'as_writable_bytes'}}
+  cuda::std::as_writable_bytes(cuda::std::span<const int, 1>(iArr2 + 5, 1)); // expected-error {{no matching function for call to 'as_writable_bytes'}}
+}
+
+int main(int, char**)
+{
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.obs/empty.nodiscard.verify.cpp b/.upstream-tests/test/std/views/views.span/span.obs/empty.nodiscard.verify.cpp
new file mode 100644
index 0000000000..1f5b68c14e
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.obs/empty.nodiscard.verify.cpp
@@ -0,0 +1,30 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11, c++14, c++17
+// UNSUPPORTED: nvrtc
+
+// <span>
+
+// [[nodiscard]] constexpr bool empty() const noexcept;
+
+#include <cuda/std/span>
+
+void test() {
+  cuda::std::span<int> s1;
+  s1.empty(); // expected-warning {{ignoring return value of function declared with 'nodiscard' attribute}}
+
+  int arr[] = {1, 2, 3};
+  cuda::std::span<int, 3> s2{arr};
+  s2.empty(); // expected-warning {{ignoring return value of function declared with 'nodiscard' attribute}}
+}
+
+int main(int, char**)
+{
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.obs/empty.pass.cpp b/.upstream-tests/test/std/views/views.span/span.obs/empty.pass.cpp
new file mode 100644
index 0000000000..b1a6ee2764
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.obs/empty.pass.cpp
@@ -0,0 +1,65 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// [[nodiscard]] constexpr bool empty() const noexcept;
+//
+
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+struct A{};
+__device__ constexpr int iArr1[] = { 0,  1,  2,  3,  4,  5,  6,  7,  8,  9};
+__device__           int iArr2[] = {10, 11, 12, 13, 14, 15, 16, 17, 18, 19};
+
+int main(int, char**)
+{
+    static_assert( noexcept(cuda::std::span<int>   ().empty()), "");
+    static_assert( noexcept(cuda::std::span<int, 0>().empty()), "");
+
+
+    static_assert( cuda::std::span<int>().empty(),            "");
+    static_assert( cuda::std::span<long>().empty(),           "");
+    static_assert( cuda::std::span<double>().empty(),         "");
+    static_assert( cuda::std::span<A>().empty(),              "");
+
+    static_assert( cuda::std::span<int, 0>().empty(),         "");
+    static_assert( cuda::std::span<long, 0>().empty(),        "");
+    static_assert( cuda::std::span<double, 0>().empty(),      "");
+    static_assert( cuda::std::span<A, 0>().empty(),           "");
+
+    static_assert(!cuda::std::span<const int>(iArr1, 1).empty(), "");
+    static_assert(!cuda::std::span<const int>(iArr1, 2).empty(), "");
+    static_assert(!cuda::std::span<const int>(iArr1, 3).empty(), "");
+    static_assert(!cuda::std::span<const int>(iArr1, 4).empty(), "");
+    static_assert(!cuda::std::span<const int>(iArr1, 5).empty(), "");
+
+    assert( (cuda::std::span<int>().empty()           ));
+    assert( (cuda::std::span<long>().empty()          ));
+    assert( (cuda::std::span<double>().empty()        ));
+    assert( (cuda::std::span<A>().empty()             ));
+
+    assert( (cuda::std::span<int, 0>().empty()        ));
+    assert( (cuda::std::span<long, 0>().empty()       ));
+    assert( (cuda::std::span<double, 0>().empty()     ));
+    assert( (cuda::std::span<A, 0>().empty()          ));
+
+    assert(!(cuda::std::span<int, 1>(iArr2, 1).empty()));
+    assert(!(cuda::std::span<int, 2>(iArr2, 2).empty()));
+    assert(!(cuda::std::span<int, 3>(iArr2, 3).empty()));
+    assert(!(cuda::std::span<int, 4>(iArr2, 4).empty()));
+    assert(!(cuda::std::span<int, 5>(iArr2, 5).empty()));
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.obs/size.pass.cpp b/.upstream-tests/test/std/views/views.span/span.obs/size.pass.cpp
new file mode 100644
index 0000000000..1cb6e78db7
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.obs/size.pass.cpp
@@ -0,0 +1,84 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// constexpr size_type size() const noexcept;
+//
+
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+template <typename Span>
+__host__ __device__
+constexpr bool testConstexprSpan(Span sp, size_t sz)
+{
+    ASSERT_NOEXCEPT(sp.size());
+    return sp.size() == sz;
+}
+
+
+template <typename Span>
+__host__ __device__
+void testRuntimeSpan(Span sp, size_t sz)
+{
+    ASSERT_NOEXCEPT(sp.size());
+    assert(sp.size() == sz);
+}
+
+struct A{};
+__device__ constexpr int iArr1[] = { 0,  1,  2,  3,  4,  5,  6,  7,  8,  9};
+__device__           int iArr2[] = {10, 11, 12, 13, 14, 15, 16, 17, 18, 19};
+
+int main(int, char**)
+{
+    static_assert(testConstexprSpan(cuda::std::span<int>(), 0),            "");
+    static_assert(testConstexprSpan(cuda::std::span<long>(), 0),           "");
+    static_assert(testConstexprSpan(cuda::std::span<double>(), 0),         "");
+    static_assert(testConstexprSpan(cuda::std::span<A>(), 0),              "");
+
+    static_assert(testConstexprSpan(cuda::std::span<int, 0>(), 0),         "");
+    static_assert(testConstexprSpan(cuda::std::span<long, 0>(), 0),        "");
+    static_assert(testConstexprSpan(cuda::std::span<double, 0>(), 0),      "");
+    static_assert(testConstexprSpan(cuda::std::span<A, 0>(), 0),           "");
+
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 1), 1),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 2), 2),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 3), 3),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 4), 4),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 5), 5),    "");
+
+    testRuntimeSpan(cuda::std::span<int>        (), 0);
+    testRuntimeSpan(cuda::std::span<long>       (), 0);
+    testRuntimeSpan(cuda::std::span<double>     (), 0);
+    testRuntimeSpan(cuda::std::span<A>          (), 0);
+
+    testRuntimeSpan(cuda::std::span<int, 0>        (), 0);
+    testRuntimeSpan(cuda::std::span<long, 0>       (), 0);
+    testRuntimeSpan(cuda::std::span<double, 0>     (), 0);
+    testRuntimeSpan(cuda::std::span<A, 0>          (), 0);
+
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 1), 1);
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 2), 2);
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 3), 3);
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 4), 4);
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 5), 5);
+
+    testRuntimeSpan(cuda::std::span<int, 1>(iArr2 + 5, 1), 1);
+    testRuntimeSpan(cuda::std::span<int, 2>(iArr2 + 4, 2), 2);
+    testRuntimeSpan(cuda::std::span<int, 3>(iArr2 + 3, 3), 3);
+    testRuntimeSpan(cuda::std::span<int, 4>(iArr2 + 2, 4), 4);
+    testRuntimeSpan(cuda::std::span<int, 5>(iArr2 + 1, 5), 5);
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.obs/size_bytes.pass.cpp b/.upstream-tests/test/std/views/views.span/span.obs/size_bytes.pass.cpp
new file mode 100644
index 0000000000..d736a1ecc8
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.obs/size_bytes.pass.cpp
@@ -0,0 +1,85 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// constexpr size_type size_bytes() const noexcept;
+//
+//  Effects: Equivalent to: return size() * sizeof(element_type);
+
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+template <typename Span>
+__host__ __device__
+constexpr bool testConstexprSpan(Span sp, size_t sz)
+{
+    ASSERT_NOEXCEPT(sp.size_bytes());
+    return (size_t) sp.size_bytes() == sz * sizeof(typename Span::element_type);
+}
+
+
+template <typename Span>
+__host__ __device__
+void testRuntimeSpan(Span sp, size_t sz)
+{
+    ASSERT_NOEXCEPT(sp.size_bytes());
+    assert((size_t) sp.size_bytes() == sz * sizeof(typename Span::element_type));
+}
+
+struct A{};
+__device__ constexpr int iArr1[] = { 0,  1,  2,  3,  4,  5,  6,  7,  8,  9};
+__device__           int iArr2[] = {10, 11, 12, 13, 14, 15, 16, 17, 18, 19};
+
+int main(int, char**)
+{
+    static_assert(testConstexprSpan(cuda::std::span<int>(), 0),            "");
+    static_assert(testConstexprSpan(cuda::std::span<long>(), 0),           "");
+    static_assert(testConstexprSpan(cuda::std::span<double>(), 0),         "");
+    static_assert(testConstexprSpan(cuda::std::span<A>(), 0),              "");
+
+    static_assert(testConstexprSpan(cuda::std::span<int, 0>(), 0),         "");
+    static_assert(testConstexprSpan(cuda::std::span<long, 0>(), 0),        "");
+    static_assert(testConstexprSpan(cuda::std::span<double, 0>(), 0),      "");
+    static_assert(testConstexprSpan(cuda::std::span<A, 0>(), 0),           "");
+
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 1), 1),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 2), 2),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 3), 3),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 4), 4),    "");
+    static_assert(testConstexprSpan(cuda::std::span<const int>(iArr1, 5), 5),    "");
+
+    testRuntimeSpan(cuda::std::span<int>        (), 0);
+    testRuntimeSpan(cuda::std::span<long>       (), 0);
+    testRuntimeSpan(cuda::std::span<double>     (), 0);
+    testRuntimeSpan(cuda::std::span<A>          (), 0);
+
+    testRuntimeSpan(cuda::std::span<int, 0>        (), 0);
+    testRuntimeSpan(cuda::std::span<long, 0>       (), 0);
+    testRuntimeSpan(cuda::std::span<double, 0>     (), 0);
+    testRuntimeSpan(cuda::std::span<A, 0>          (), 0);
+
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 1), 1);
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 2), 2);
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 3), 3);
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 4), 4);
+    testRuntimeSpan(cuda::std::span<int>(iArr2, 5), 5);
+
+    testRuntimeSpan(cuda::std::span<int, 1>(iArr2 + 5, 1), 1);
+    testRuntimeSpan(cuda::std::span<int, 2>(iArr2 + 4, 2), 2);
+    testRuntimeSpan(cuda::std::span<int, 3>(iArr2 + 3, 3), 3);
+    testRuntimeSpan(cuda::std::span<int, 4>(iArr2 + 2, 4), 4);
+    testRuntimeSpan(cuda::std::span<int, 5>(iArr2 + 1, 5), 5);
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.sub/first.pass.cpp b/.upstream-tests/test/std/views/views.span/span.sub/first.pass.cpp
new file mode 100644
index 0000000000..fcd4019958
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.sub/first.pass.cpp
@@ -0,0 +1,110 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// template<size_t Count>
+//  constexpr span<element_type, Count> first() const;
+//
+// constexpr span<element_type, dynamic_extent> first(size_type count) const;
+//
+//  Requires: Count <= size().
+
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+template <typename Span, size_t Count>
+__host__ __device__
+constexpr bool testConstexprSpan(Span sp)
+{
+    ASSERT_NOEXCEPT(sp.template first<Count>());
+    ASSERT_NOEXCEPT(sp.first(Count));
+    auto s1 = sp.template first<Count>();
+    auto s2 = sp.first(Count);
+    using S1 = decltype(s1);
+    using S2 = decltype(s2);
+    ASSERT_SAME_TYPE(typename Span::value_type, typename S1::value_type);
+    ASSERT_SAME_TYPE(typename Span::value_type, typename S2::value_type);
+    static_assert(S1::extent == Count, "");
+    static_assert(S2::extent == cuda::std::dynamic_extent, "");
+    return
+        s1.data() == s2.data()
+     && s1.size() == s2.size();
+}
+
+
+template <typename Span, size_t Count>
+__host__ __device__
+void testRuntimeSpan(Span sp)
+{
+    ASSERT_NOEXCEPT(sp.template first<Count>());
+    ASSERT_NOEXCEPT(sp.first(Count));
+    auto s1 = sp.template first<Count>();
+    auto s2 = sp.first(Count);
+    using S1 = decltype(s1);
+    using S2 = decltype(s2);
+    ASSERT_SAME_TYPE(typename Span::value_type, typename S1::value_type);
+    ASSERT_SAME_TYPE(typename Span::value_type, typename S2::value_type);
+    static_assert(S1::extent == Count, "");
+    static_assert(S2::extent == cuda::std::dynamic_extent, "");
+    assert(s1.data() == s2.data());
+    assert(s1.size() == s2.size());
+}
+
+__device__ constexpr int carr1[] = {1,2,3,4};
+__device__           int   arr[] = {5,6,7};
+
+int main(int, char**)
+{
+    {
+    using Sp = cuda::std::span<const int>;
+    static_assert(testConstexprSpan<Sp, 0>(Sp{}), "");
+
+    static_assert(testConstexprSpan<Sp, 0>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 1>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 2>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 3>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 4>(Sp{carr1}), "");
+    }
+
+    {
+    using Sp = cuda::std::span<const int, 4>;
+
+    static_assert(testConstexprSpan<Sp, 0>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 1>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 2>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 3>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 4>(Sp{carr1}), "");
+    }
+    
+    {
+    using Sp = cuda::std::span<int>;
+    testRuntimeSpan<Sp, 0>(Sp{});
+
+    testRuntimeSpan<Sp, 0>(Sp{arr});
+    testRuntimeSpan<Sp, 1>(Sp{arr});
+    testRuntimeSpan<Sp, 2>(Sp{arr});
+    testRuntimeSpan<Sp, 3>(Sp{arr});
+    }
+
+    {
+    using Sp = cuda::std::span<int, 3>;
+
+    testRuntimeSpan<Sp, 0>(Sp{arr});
+    testRuntimeSpan<Sp, 1>(Sp{arr});
+    testRuntimeSpan<Sp, 2>(Sp{arr});
+    testRuntimeSpan<Sp, 3>(Sp{arr});
+    }
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.sub/first.verify.cpp b/.upstream-tests/test/std/views/views.span/span.sub/first.verify.cpp
new file mode 100644
index 0000000000..a500d74062
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.sub/first.verify.cpp
@@ -0,0 +1,36 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+// UNSUPPORTED: nvrtc
+
+// <span>
+
+// template<size_t Count>
+//  constexpr span<element_type, Count> first() const;
+//
+//  Requires: Count <= size().
+
+#include <cuda/std/span>
+#include <cuda/std/cstddef>
+
+void f() {
+  int array[] = {1, 2, 3, 4};
+  cuda::std::span<const int, 4> sp(array);
+
+  //  Count too large
+  [[maybe_unused]] auto s1 = sp.first<5>(); // expected-error@span:* {{span<T, N>::first<Count>(): Count out of range}}
+
+  //  Count numeric_limits
+  [[maybe_unused]] auto s2 = sp.first<cuda::std::size_t(-1)>(); // expected-error@span:* {{span<T, N>::first<Count>(): Count out of range}}
+}
+
+int main(int, char**)
+{
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.sub/last.pass.cpp b/.upstream-tests/test/std/views/views.span/span.sub/last.pass.cpp
new file mode 100644
index 0000000000..c10afc0b91
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.sub/last.pass.cpp
@@ -0,0 +1,109 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// template<size_t Count>
+//  constexpr span<element_type, Count> last() const;
+//
+// constexpr span<element_type, dynamic_extent> last(size_type count) const;
+//
+//  Requires: Count <= size().
+
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+template <typename Span, size_t Count>
+__host__ __device__
+constexpr bool testConstexprSpan(Span sp)
+{
+    assert((noexcept(sp.template last<Count>())));
+    assert((noexcept(sp.last(Count))));
+    auto s1 = sp.template last<Count>();
+    auto s2 = sp.last(Count);
+    using S1 = decltype(s1);
+    using S2 = decltype(s2);
+    ASSERT_SAME_TYPE(typename Span::value_type, typename S1::value_type);
+    ASSERT_SAME_TYPE(typename Span::value_type, typename S2::value_type);
+    static_assert(S1::extent == Count, "");
+    static_assert(S2::extent == cuda::std::dynamic_extent, "");
+    return
+        s1.data() == s2.data()
+     && s1.size() == s2.size();
+}
+
+template <typename Span, size_t Count>
+__host__ __device__
+void testRuntimeSpan(Span sp)
+{
+    assert((noexcept(sp.template last<Count>())));
+    assert((noexcept(sp.last(Count))));
+    auto s1 = sp.template last<Count>();
+    auto s2 = sp.last(Count);
+    using S1 = decltype(s1);
+    using S2 = decltype(s2);
+    ASSERT_SAME_TYPE(typename Span::value_type, typename S1::value_type);
+    ASSERT_SAME_TYPE(typename Span::value_type, typename S2::value_type);
+    static_assert(S1::extent == Count, "");
+    static_assert(S2::extent == cuda::std::dynamic_extent, "");
+    assert(s1.data() == s2.data());
+    assert(s1.size() == s2.size());
+}
+
+__device__ constexpr int carr1[] = {1,2,3,4};
+__device__           int   arr[] = {5,6,7};
+
+int main(int, char**)
+{
+    {
+    using Sp = cuda::std::span<const int>;
+    static_assert(testConstexprSpan<Sp, 0>(Sp{}), "");
+
+    static_assert(testConstexprSpan<Sp, 0>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 1>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 2>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 3>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 4>(Sp{carr1}), "");
+    }
+
+    {
+    using Sp = cuda::std::span<const int, 4>;
+
+    static_assert(testConstexprSpan<Sp, 0>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 1>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 2>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 3>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 4>(Sp{carr1}), "");
+    }
+    
+    {
+    using Sp = cuda::std::span<int>;
+    testRuntimeSpan<Sp, 0>(Sp{});
+
+    testRuntimeSpan<Sp, 0>(Sp{arr});
+    testRuntimeSpan<Sp, 1>(Sp{arr});
+    testRuntimeSpan<Sp, 2>(Sp{arr});
+    testRuntimeSpan<Sp, 3>(Sp{arr});
+    }
+
+    {
+    using Sp = cuda::std::span<int, 3>;
+
+    testRuntimeSpan<Sp, 0>(Sp{arr});
+    testRuntimeSpan<Sp, 1>(Sp{arr});
+    testRuntimeSpan<Sp, 2>(Sp{arr});
+    testRuntimeSpan<Sp, 3>(Sp{arr});
+    }
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.sub/last.verify.cpp b/.upstream-tests/test/std/views/views.span/span.sub/last.verify.cpp
new file mode 100644
index 0000000000..5deb1e7fc9
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.sub/last.verify.cpp
@@ -0,0 +1,36 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+// UNSUPPORTED: nvrtc
+
+// <span>
+
+// template<size_t Count>
+//  constexpr span<element_type, Count> last() const;
+//
+//  Requires: Count <= size().
+
+#include <cuda/std/span>
+#include <cuda/std/cstddef>
+
+void f() {
+  int array[] = {1, 2, 3, 4};
+  cuda::std::span<const int, 4> sp(array);
+
+  //  Count too large
+  [[maybe_unused]] auto s1 = sp.last<5>(); // expected-error@span:* {{span<T, N>::last<Count>(): Count out of range}}
+
+  //  Count numeric_limits
+  [[maybe_unused]] auto s2 = sp.last<cuda::std::size_t(-1)>(); // expected-error@span:* {{span<T, N>::last<Count>(): Count out of range}}
+}
+
+int main(int, char**)
+{
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.sub/subspan.pass.cpp b/.upstream-tests/test/std/views/views.span/span.sub/subspan.pass.cpp
new file mode 100644
index 0000000000..09247c83e0
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.sub/subspan.pass.cpp
@@ -0,0 +1,208 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// template<size_t Offset, size_t Count = dynamic_extent>
+//   constexpr span<element_type, see below> subspan() const;
+//
+// constexpr span<element_type, dynamic_extent> subspan(
+//   size_type offset, size_type count = dynamic_extent) const;
+//
+//  Requires: (0 <= Offset && Offset <= size())
+//      && (Count == dynamic_extent || Count >= 0 && Offset + Count <= size())
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+template <typename Span, size_t Offset, size_t Count>
+__host__ __device__
+constexpr bool testConstexprSpan(Span sp)
+{
+    ASSERT_NOEXCEPT(sp.template subspan<Offset, Count>());
+    ASSERT_NOEXCEPT(sp.subspan(Offset, Count));
+    auto s1 = sp.template subspan<Offset, Count>();
+    auto s2 = sp.subspan(Offset, Count);
+    using S1 = decltype(s1);
+    using S2 = decltype(s2);
+    ASSERT_SAME_TYPE(typename Span::value_type, typename S1::value_type);
+    ASSERT_SAME_TYPE(typename Span::value_type, typename S2::value_type);
+    static_assert(S1::extent == Count, "");
+    static_assert(S2::extent == cuda::std::dynamic_extent, "");
+    return
+        s1.data() == s2.data()
+     && s1.size() == s2.size();
+}
+
+template <typename Span, size_t Offset>
+__host__ __device__
+constexpr bool testConstexprSpan(Span sp)
+{
+    ASSERT_NOEXCEPT(sp.template subspan<Offset>());
+    ASSERT_NOEXCEPT(sp.subspan(Offset));
+    auto s1 = sp.template subspan<Offset>();
+    auto s2 = sp.subspan(Offset);
+    using S1 = decltype(s1);
+    using S2 = decltype(s2);
+    ASSERT_SAME_TYPE(typename Span::value_type, typename S1::value_type);
+    ASSERT_SAME_TYPE(typename Span::value_type, typename S2::value_type);
+    static_assert(S1::extent == (Span::extent == cuda::std::dynamic_extent ? cuda::std::dynamic_extent : Span::extent - Offset), "");
+    static_assert(S2::extent == cuda::std::dynamic_extent, "");
+    return
+        s1.data() == s2.data()
+     && s1.size() == s2.size();
+}
+
+
+template <typename Span, size_t Offset, size_t Count>
+__host__ __device__
+void testRuntimeSpan(Span sp)
+{
+    ASSERT_NOEXCEPT(sp.template subspan<Offset, Count>());
+    ASSERT_NOEXCEPT(sp.subspan(Offset, Count));
+    auto s1 = sp.template subspan<Offset, Count>();
+    auto s2 = sp.subspan(Offset, Count);
+    using S1 = decltype(s1);
+    using S2 = decltype(s2);
+    ASSERT_SAME_TYPE(typename Span::value_type, typename S1::value_type);
+    ASSERT_SAME_TYPE(typename Span::value_type, typename S2::value_type);
+    static_assert(S1::extent == Count, "");
+    static_assert(S2::extent == cuda::std::dynamic_extent, "");
+    assert(s1.data() == s2.data());
+    assert(s1.size() == s2.size());
+}
+
+
+template <typename Span, size_t Offset>
+__host__ __device__
+void testRuntimeSpan(Span sp)
+{
+    ASSERT_NOEXCEPT(sp.template subspan<Offset>());
+    ASSERT_NOEXCEPT(sp.subspan(Offset));
+    auto s1 = sp.template subspan<Offset>();
+    auto s2 = sp.subspan(Offset);
+    using S1 = decltype(s1);
+    using S2 = decltype(s2);
+    ASSERT_SAME_TYPE(typename Span::value_type, typename S1::value_type);
+    ASSERT_SAME_TYPE(typename Span::value_type, typename S2::value_type);
+    static_assert(S1::extent == (Span::extent == cuda::std::dynamic_extent ? cuda::std::dynamic_extent : Span::extent - Offset), "");
+    static_assert(S2::extent == cuda::std::dynamic_extent, "");
+    assert(s1.data() == s2.data());
+    assert(s1.size() == s2.size());
+}
+
+__device__ constexpr int carr1[] = {1,2,3,4};
+__device__           int  arr1[] = {5,6,7};
+
+int main(int, char**)
+{
+    {
+    using Sp = cuda::std::span<const int>;
+    static_assert(testConstexprSpan<Sp, 0>(Sp{}), "");
+
+    static_assert(testConstexprSpan<Sp, 0, 4>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 0, 3>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 0, 2>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 0, 1>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 0, 0>(Sp{carr1}), "");
+
+    static_assert(testConstexprSpan<Sp, 1, 3>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 2, 2>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 3, 1>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 4, 0>(Sp{carr1}), "");
+    }
+
+    {
+    using Sp = cuda::std::span<const int, 4>;
+
+    static_assert(testConstexprSpan<Sp, 0, 4>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 0, 3>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 0, 2>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 0, 1>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 0, 0>(Sp{carr1}), "");
+
+    static_assert(testConstexprSpan<Sp, 1, 3>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 2, 2>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 3, 1>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 4, 0>(Sp{carr1}), "");
+    }
+
+    {
+    using Sp = cuda::std::span<const int>;
+    static_assert(testConstexprSpan<Sp, 0>(Sp{}), "");
+
+    static_assert(testConstexprSpan<Sp, 0>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 1>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 2>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 3>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 4>(Sp{carr1}), "");
+    }
+
+    {
+    using Sp = cuda::std::span<const int, 4>;
+
+    static_assert(testConstexprSpan<Sp, 0>(Sp{carr1}), "");
+
+    static_assert(testConstexprSpan<Sp, 1>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 2>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 3>(Sp{carr1}), "");
+    static_assert(testConstexprSpan<Sp, 4>(Sp{carr1}), "");
+    }
+
+    {
+    using Sp = cuda::std::span<int>;
+    testRuntimeSpan<Sp, 0>(Sp{});
+
+    testRuntimeSpan<Sp, 0, 3>(Sp{arr1});
+    testRuntimeSpan<Sp, 0, 2>(Sp{arr1});
+    testRuntimeSpan<Sp, 0, 1>(Sp{arr1});
+    testRuntimeSpan<Sp, 0, 0>(Sp{arr1});
+
+    testRuntimeSpan<Sp, 1, 2>(Sp{arr1});
+    testRuntimeSpan<Sp, 2, 1>(Sp{arr1});
+    testRuntimeSpan<Sp, 3, 0>(Sp{arr1});
+    }
+
+    {
+    using Sp = cuda::std::span<int, 3>;
+
+    testRuntimeSpan<Sp, 0, 3>(Sp{arr1});
+    testRuntimeSpan<Sp, 0, 2>(Sp{arr1});
+    testRuntimeSpan<Sp, 0, 1>(Sp{arr1});
+    testRuntimeSpan<Sp, 0, 0>(Sp{arr1});
+
+    testRuntimeSpan<Sp, 1, 2>(Sp{arr1});
+    testRuntimeSpan<Sp, 2, 1>(Sp{arr1});
+    testRuntimeSpan<Sp, 3, 0>(Sp{arr1});
+    }
+
+    {
+    using Sp = cuda::std::span<int>;
+    testRuntimeSpan<Sp, 0>(Sp{});
+
+    testRuntimeSpan<Sp, 0>(Sp{arr1});
+    testRuntimeSpan<Sp, 1>(Sp{arr1});
+    testRuntimeSpan<Sp, 2>(Sp{arr1});
+    testRuntimeSpan<Sp, 3>(Sp{arr1});
+    }
+
+    {
+    using Sp = cuda::std::span<int, 3>;
+
+    testRuntimeSpan<Sp, 0>(Sp{arr1});
+    testRuntimeSpan<Sp, 1>(Sp{arr1});
+    testRuntimeSpan<Sp, 2>(Sp{arr1});
+    testRuntimeSpan<Sp, 3>(Sp{arr1});
+    }
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/views/views.span/span.sub/subspan.verify.cpp b/.upstream-tests/test/std/views/views.span/span.sub/subspan.verify.cpp
new file mode 100644
index 0000000000..9cf47fffe4
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/span.sub/subspan.verify.cpp
@@ -0,0 +1,42 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+// UNSUPPORTED: nvrtc
+
+// This test also generates spurious warnings when instantiating std::span
+// with a very large extent (like size_t(-2)) -- silence those.
+// ADDITIONAL_COMPILE_FLAGS: -Xclang -verify-ignore-unexpected=warning
+
+// <span>
+
+// template<size_t Offset, size_t Count = dynamic_extent>
+//   constexpr span<element_type, see below> subspan() const;
+//
+//  Requires: offset <= size() &&
+//            (count == dynamic_extent || count <= size() - offset)
+
+#include <cuda/std/span>
+#include <cuda/std/cstddef>
+
+void f() {
+  int array[] = {1, 2, 3, 4};
+  cuda::std::span<const int, 4> sp(array);
+
+  //  Offset too large templatized
+  [[maybe_unused]] auto s1 = sp.subspan<5>(); // expected-error@span:* {{span<T, N>::subspan<Offset, Count>(): Offset out of range}}
+
+  //  Count too large templatized
+  [[maybe_unused]] auto s2 = sp.subspan<0, 5>(); // expected-error@span:* {{span<T, N>::subspan<Offset, Count>(): Offset + Count out of range}}
+
+  //  Offset + Count too large templatized
+  [[maybe_unused]] auto s3 = sp.subspan<2, 3>(); // expected-error@span:* {{span<T, N>::subspan<Offset, Count>(): Offset + Count out of range}}
+
+  //  Offset + Count overflow templatized
+  [[maybe_unused]] auto s4 = sp.subspan<3, cuda::std::size_t(-2)>(); // expected-error@span:* {{span<T, N>::subspan<Offset, Count>(): Offset + Count out of range}}, expected-error-re@span:* {{array is too large{{(.* elements)}}}}
+}
diff --git a/.upstream-tests/test/std/views/views.span/trivially_copyable.compile.pass.cpp b/.upstream-tests/test/std/views/views.span/trivially_copyable.compile.pass.cpp
new file mode 100644
index 0000000000..aea36c7730
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/trivially_copyable.compile.pass.cpp
@@ -0,0 +1,24 @@
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: c++03, c++11
+
+// P2251 was voted into C++23, but is supported even in C++20 mode by all vendors.
+
+// <span>
+
+#include <cuda/std/span>
+#include <cuda/std/type_traits>
+
+static_assert(cuda::std::is_trivially_copyable<cuda::std::span<int>>::value, "");
+static_assert(cuda::std::is_trivially_copyable<cuda::std::span<int, 3>>::value, "");
+
+int main(int, char**)
+{
+    return 0;
+}
\ No newline at end of file
diff --git a/.upstream-tests/test/std/views/views.span/types.pass.cpp b/.upstream-tests/test/std/views/views.span/types.pass.cpp
new file mode 100644
index 0000000000..691de311dd
--- /dev/null
+++ b/.upstream-tests/test/std/views/views.span/types.pass.cpp
@@ -0,0 +1,95 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11
+
+// <span>
+
+// template<class ElementType, size_t Extent = dynamic_extent>
+// class span {
+// public:
+//  // constants and types
+//  using element_type           = ElementType;
+//  using value_type             = remove_cv_t<ElementType>;
+//  using size_type              = size_t;
+//  using difference_type        = ptrdiff_t;
+//  using pointer                = element_type *;
+//  using reference              = element_type &;
+//  using const_pointe           = const element_type *;
+//  using const_reference        = const element_type &;
+//  using iterator               = implementation-defined;
+//  using reverse_iterator       = std::reverse_iterator<iterator>;
+//
+//  static constexpr size_type extent = Extent;
+//
+
+#include <cuda/std/span>
+#include <cuda/std/cassert>
+#include <cuda/std/iterator>
+
+#include "test_macros.h"
+
+template <typename S, typename Iter>
+__host__ __device__
+void testIterator()
+{
+    typedef cuda::std::iterator_traits<Iter> ItT;
+
+    ASSERT_SAME_TYPE(typename ItT::iterator_category, cuda::std::random_access_iterator_tag);
+    ASSERT_SAME_TYPE(typename ItT::value_type,        typename S::value_type);
+    ASSERT_SAME_TYPE(typename ItT::reference,         typename S::reference);
+    ASSERT_SAME_TYPE(typename ItT::pointer,           typename S::pointer);
+    ASSERT_SAME_TYPE(typename ItT::difference_type,   typename S::difference_type);
+}
+
+template <typename S, typename ElementType, cuda::std::size_t Size>
+__host__ __device__
+void testSpan()
+{
+    ASSERT_SAME_TYPE(typename S::element_type,    ElementType);
+    ASSERT_SAME_TYPE(typename S::value_type,      typename cuda::std::remove_cv<ElementType>::type);
+    ASSERT_SAME_TYPE(typename S::size_type,       cuda::std::size_t);
+    ASSERT_SAME_TYPE(typename S::difference_type, cuda::std::ptrdiff_t);
+    ASSERT_SAME_TYPE(typename S::pointer,         ElementType *);
+    ASSERT_SAME_TYPE(typename S::const_pointer,   const ElementType *);
+    ASSERT_SAME_TYPE(typename S::reference,       ElementType &);
+    ASSERT_SAME_TYPE(typename S::const_reference, const ElementType &);
+
+    static_assert(S::extent == Size, ""); // check that it exists
+
+    testIterator<S, typename S::iterator>();
+    testIterator<S, typename S::reverse_iterator>();
+}
+
+
+template <typename T>
+__host__ __device__
+void test()
+{
+    testSpan<cuda::std::span<               T>,                T, cuda::std::dynamic_extent>();
+    testSpan<cuda::std::span<const          T>, const          T, cuda::std::dynamic_extent>();
+    testSpan<cuda::std::span<      volatile T>,       volatile T, cuda::std::dynamic_extent>();
+    testSpan<cuda::std::span<const volatile T>, const volatile T, cuda::std::dynamic_extent>();
+
+    testSpan<cuda::std::span<               T, 5>,                T, 5>();
+    testSpan<cuda::std::span<const          T, 5>, const          T, 5>();
+    testSpan<cuda::std::span<      volatile T, 5>,       volatile T, 5>();
+    testSpan<cuda::std::span<const volatile T, 5>, const volatile T, 5>();
+}
+
+struct A{};
+
+int main(int, char**)
+{
+    test<int>();
+    test<long>();
+    test<double>();
+    test<A>();
+
+    return 0;
+}
diff --git a/.upstream-tests/test/support/test_macros.h b/.upstream-tests/test/support/test_macros.h
index ed7306f2c5..74f3bc54d7 100644
--- a/.upstream-tests/test/support/test_macros.h
+++ b/.upstream-tests/test/support/test_macros.h
@@ -149,6 +149,11 @@
 # else
 #   define TEST_CONSTEXPR_CXX14
 # endif
+# if TEST_STD_VER >= 17
+#   define TEST_CONSTEXPR_CXX17 constexpr
+# else
+#   define TEST_CONSTEXPR_CXX17
+# endif
 # if TEST_STD_VER > 14
 #   define TEST_THROW_SPEC(...)
 # else
@@ -268,6 +273,18 @@
     static_assert(!noexcept(__VA_ARGS__), "Operation must NOT be noexcept")
 #endif
 
+#if TEST_STD_VER > 11
+#define STATIC_ASSERT_CXX14(Pred) static_assert(Pred, "")
+#else
+#define STATIC_ASSERT_CXX14(Pred) assert(Pred)
+#endif
+
+#if TEST_STD_VER > 14
+#define STATIC_ASSERT_CXX17(Pred) static_assert(Pred, "")
+#else
+#define STATIC_ASSERT_CXX17(Pred) assert(Pred)
+#endif
+
 /* Macros for testing libc++ specific behavior and extensions */
 #if defined(_LIBCUDACXX_VERSION)
 #define LIBCPP_ASSERT(...) assert(__VA_ARGS__)
diff --git a/include/cuda/std/detail/libcxx/include/iterator b/include/cuda/std/detail/libcxx/include/iterator
index 95f9fe4314..e91e76df1d 100644
--- a/include/cuda/std/detail/libcxx/include/iterator
+++ b/include/cuda/std/detail/libcxx/include/iterator
@@ -1856,7 +1856,7 @@ end(const _Cp& __c)
 
 #endif  // !defined(_LIBCUDACXX_CXX03_LANG)
 
-#if _LIBCUDACXX_STD_VER > 14
+#if _LIBCUDACXX_STD_VER > 11
 
 // #if _LIBCUDACXX_STD_VER > 11
 // template <>
@@ -1932,7 +1932,7 @@ constexpr _Tp* data(_Tp (&__array)[_Sz]) noexcept { return __array; }
 template <class _Ep>
 _LIBCUDACXX_INLINE_VISIBILITY
 constexpr const _Ep* data(::std::initializer_list<_Ep> __il) noexcept { return __il.begin(); }
-#endif
+#endif // _LIBCUDACXX_STD_VER > 11
 
 _LIBCUDACXX_END_NAMESPACE_STD
 
diff --git a/include/cuda/std/detail/libcxx/include/span b/include/cuda/std/detail/libcxx/include/span
index c1d2f9357f..595e835eaf 100644
--- a/include/cuda/std/detail/libcxx/include/span
+++ b/include/cuda/std/detail/libcxx/include/span
@@ -4,6 +4,7 @@
 // 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 //
 //===---------------------------------------------------------------------===//
 
@@ -22,6 +23,12 @@ inline constexpr size_t dynamic_extent = numeric_limits<size_t>::max();
 template <class ElementType, size_t Extent = dynamic_extent>
     class span;
 
+template<class ElementType, size_t Extent>
+  inline constexpr bool ranges::enable_view<span<ElementType, Extent>> = true;
+
+template<class ElementType, size_t Extent>
+    inline constexpr bool ranges::enable_borrowed_range<span<ElementType, Extent>> = true;
+
 // [span.objectrep], views of object representation
 template <class ElementType, size_t Extent>
     span<const byte, ((Extent == dynamic_extent) ? dynamic_extent :
@@ -32,42 +39,39 @@ template <class ElementType, size_t Extent>
         (sizeof(ElementType) * Extent))> as_writable_bytes(span<ElementType, Extent> s) noexcept;
 
 
-namespace std {
 template <class ElementType, size_t Extent = dynamic_extent>
 class span {
 public:
     // constants and types
     using element_type = ElementType;
     using value_type = remove_cv_t<ElementType>;
-    using index_type = size_t;
+    using size_type = size_t;
     using difference_type = ptrdiff_t;
     using pointer = element_type*;
     using const_pointer = const element_type*;
     using reference = element_type&;
     using const_reference = const element_type&;
     using iterator = implementation-defined;
-    using const_iterator = implementation-defined;
     using reverse_iterator = std::reverse_iterator<iterator>;
-    using const_reverse_iterator = std::reverse_iterator<const_iterator>;
-    static constexpr index_type extent = Extent;
+    static constexpr size_type extent = Extent;
 
     // [span.cons], span constructors, copy, assignment, and destructor
     constexpr span() noexcept;
-    constexpr span(pointer ptr, index_type count);
-    constexpr span(pointer firstElem, pointer lastElem);
+    template <class It>
+    constexpr explicit(Extent != dynamic_extent) span(It first, size_type count);
+    template <class It, class End>
+    constexpr explicit(Extent != dynamic_extent) span(It first, End last);
     template <size_t N>
-        constexpr span(element_type (&arr)[N]) noexcept;
+        constexpr span(type_identity_t<element_type> (&arr)[N]) noexcept;
     template <size_t N>
         constexpr span(array<value_type, N>& arr) noexcept;
     template <size_t N>
         constexpr span(const array<value_type, N>& arr) noexcept;
-    template <class Container>
-        constexpr span(Container& cont);
-    template <class Container>
-        constexpr span(const Container& cont);
+    template<class R>
+      constexpr explicit(Extent != dynamic_extent) span(R&& r);
     constexpr span(const span& other) noexcept = default;
     template <class OtherElementType, size_t OtherExtent>
-        constexpr span(const span<OtherElementType, OtherExtent>& s) noexcept;
+        constexpr explicit(Extent != dynamic_extent) span(const span<OtherElementType, OtherExtent>& s) noexcept;
     ~span() noexcept = default;
     constexpr span& operator=(const span& other) noexcept = default;
 
@@ -79,17 +83,17 @@ public:
     template <size_t Offset, size_t Count = dynamic_extent>
         constexpr span<element_type, see below> subspan() const;
 
-    constexpr span<element_type, dynamic_extent> first(index_type count) const;
-    constexpr span<element_type, dynamic_extent> last(index_type count) const;
-    constexpr span<element_type, dynamic_extent> subspan(index_type offset, index_type count = dynamic_extent) const;
+    constexpr span<element_type, dynamic_extent> first(size_type count) const;
+    constexpr span<element_type, dynamic_extent> last(size_type count) const;
+    constexpr span<element_type, dynamic_extent> subspan(size_type offset, size_type count = dynamic_extent) const;
 
     // [span.obs], span observers
-    constexpr index_type size() const noexcept;
-    constexpr index_type size_bytes() const noexcept;
-    constexpr bool empty() const noexcept;
+    constexpr size_type size() const noexcept;
+    constexpr size_type size_bytes() const noexcept;
+    [[nodiscard]] constexpr bool empty() const noexcept;
 
     // [span.elem], span element access
-    constexpr reference operator[](index_type idx) const;
+    constexpr reference operator[](size_type idx) const;
     constexpr reference front() const;
     constexpr reference back() const;
     constexpr pointer data() const noexcept;
@@ -97,18 +101,17 @@ public:
     // [span.iterators], span iterator support
     constexpr iterator begin() const noexcept;
     constexpr iterator end() const noexcept;
-    constexpr const_iterator cbegin() const noexcept;
-    constexpr const_iterator cend() const noexcept;
     constexpr reverse_iterator rbegin() const noexcept;
     constexpr reverse_iterator rend() const noexcept;
-    constexpr const_reverse_iterator crbegin() const noexcept;
-    constexpr const_reverse_iterator crend() const noexcept;
 
 private:
-    pointer data_;     // exposition only
-    index_type size_;  // exposition only
+    pointer data_;    // exposition only
+    size_type size_;  // exposition only
 };
 
+template<class It, class EndOrSize>
+    span(It, EndOrSize) -> span<remove_reference_t<iter_reference_t<_It>>>;
+
 template<class T, size_t N>
     span(T (&)[N]) -> span<T, N>;
 
@@ -118,22 +121,21 @@ template<class T, size_t N>
 template<class T, size_t N>
     span(const array<T, N>&) -> span<const T, N>;
 
-template<class Container>
-    span(Container&) -> span<typename Container::value_type>;
-
-template<class Container>
-    span(const Container&) -> span<const typename Container::value_type>;
+template<class R>
+    span(R&&) -> span<remove_reference_t<ranges::range_reference_t<R>>>;
 
 } // namespace std
 
 */
 
+#ifndef __cuda_std__
 #include <__config>
 #include <cstddef>      // for ptrdiff_t
 #include <iterator>     // for iterators
 #include <array>        // for array
 #include <type_traits>  // for remove_cv, etc
 #include <cstddef>      // for byte
+#endif // __cuda_std__
 
 #if defined(_LIBCUDACXX_USE_PRAGMA_GCC_SYSTEM_HEADER)
 #pragma GCC system_header
@@ -141,9 +143,9 @@ template<class Container>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
-#if _LIBCUDACXX_STD_VER > 17
+#if _LIBCUDACXX_STD_VER > 11
 
-inline constexpr size_t dynamic_extent = numeric_limits<size_t>::max();
+_LIBCUDACXX_INLINE_VAR constexpr size_t dynamic_extent = numeric_limits<size_t>::max();
 template <typename _Tp, size_t _Extent = dynamic_extent> class span;
 
 
@@ -172,74 +174,88 @@ template <class _Tp, class _ElementType>
 struct __is_span_compatible_container<_Tp, _ElementType,
         void_t<
         // is not a specialization of span
-            typename enable_if<!__is_span<_Tp>::value, nullptr_t>::type,
+            enable_if_t<!__is_span<_Tp>::value, nullptr_t>,
         // is not a specialization of array
-            typename enable_if<!__is_std_array<_Tp>::value, nullptr_t>::type,
+            enable_if_t<!__is_std_array<_Tp>::value, nullptr_t>,
         // is_array_v<Container> is false,
-            typename enable_if<!is_array_v<_Tp>, nullptr_t>::type,
+            enable_if_t<!is_array_v<_Tp>, nullptr_t>,
         // data(cont) and size(cont) are well formed
             decltype(data(declval<_Tp>())),
             decltype(size(declval<_Tp>())),
         // remove_pointer_t<decltype(data(cont))>(*)[] is convertible to ElementType(*)[]
-            typename enable_if<
-                is_convertible_v<remove_pointer_t<decltype(data(declval<_Tp &>()))>(*)[],
-                                 _ElementType(*)[]>,
-                nullptr_t>::type
+            enable_if_t<
+                is_convertible<remove_pointer_t<decltype(data(declval<_Tp &>()))>(*)[],
+                                 _ElementType(*)[]>::value,
+                nullptr_t>
         >>
     : public true_type {};
 
-
 template <typename _Tp, size_t _Extent>
 class _LIBCUDACXX_TEMPLATE_VIS span {
 public:
 //  constants and types
     using element_type           = _Tp;
     using value_type             = remove_cv_t<_Tp>;
-    using index_type             = size_t;
+    using size_type              = size_t;
     using difference_type        = ptrdiff_t;
     using pointer                = _Tp *;
     using const_pointer          = const _Tp *;
     using reference              = _Tp &;
     using const_reference        = const _Tp &;
     using iterator               =  __wrap_iter<pointer>;
-    using const_iterator         =  __wrap_iter<const_pointer>;
     using reverse_iterator       = _CUDA_VSTD::reverse_iterator<iterator>;
-    using const_reverse_iterator = _CUDA_VSTD::reverse_iterator<const_iterator>;
 
-    static constexpr index_type extent = _Extent;
+    static constexpr size_type extent = _Extent;
 
 // [span.cons], span constructors, copy, assignment, and destructor
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr span() noexcept : __data{nullptr}
-    { static_assert(_Extent == 0, "Can't default construct a statically sized span with size > 0"); }
+    template <size_t _Sz = _Extent, enable_if_t<_Sz == 0, nullptr_t> = nullptr>
+    _LIBCUDACXX_INLINE_VISIBILITY constexpr span() noexcept : __data{nullptr} {}
 
-    constexpr span           (const span&) noexcept = default;
-    constexpr span& operator=(const span&) noexcept = default;
+    span           (const span&) noexcept = default;
+    span& operator=(const span&) noexcept = default;
 
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr span(pointer __ptr, index_type __count) : __data{__ptr}
+    _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 span(pointer __ptr, size_type __count) : __data{__ptr}
         { (void)__count; _LIBCUDACXX_ASSERT(_Extent == __count, "size mismatch in span's constructor (ptr, len)"); }
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr span(pointer __f, pointer __l) : __data{__f}
+    _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 span(pointer __f, pointer __l) : __data{__f}
         { (void)__l;     _LIBCUDACXX_ASSERT(_Extent == distance(__f, __l), "size mismatch in span's constructor (ptr, ptr)"); }
-
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr span(element_type (&__arr)[_Extent])          noexcept : __data{__arr} {}
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr span(      array<value_type, _Extent>& __arr) noexcept : __data{__arr.data()} {}
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr span(const array<value_type, _Extent>& __arr) noexcept : __data{__arr.data()} {}
-
-    template <class _OtherElementType>
+#ifdef _LIBCUDACXX_COMPILER_NVRTC
+    template <size_t _Sz = _Extent, enable_if_t<_Sz != 0, nullptr_t> = nullptr>
+    _LIBCUDACXX_INLINE_VISIBILITY constexpr span(type_identity_t<element_type> (&__arr)[_Sz]) noexcept : __data{__arr} {}
+#else
+    _LIBCUDACXX_INLINE_VISIBILITY constexpr span(type_identity_t<element_type> (&__arr)[_Extent]) noexcept : __data{__arr} {}
+#endif
+    template <class _OtherElementType, enable_if_t<
+                is_convertible<_OtherElementType(*)[], element_type (*)[]>::value, nullptr_t> = nullptr>
+    _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX14 
+    span(      array<_OtherElementType, _Extent>& __arr) noexcept : __data{__arr.data()} {}
+    template <class _OtherElementType, enable_if_t<
+                is_convertible<_OtherElementType(*)[], element_type (*)[]>::value, nullptr_t> = nullptr>
+    _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX14 
+    span(const array<_OtherElementType, _Extent>& __arr) noexcept : __data{__arr.data()} {}
+
+    template <class _OtherElementType, enable_if_t<
+                is_convertible<_OtherElementType(*)[], element_type (*)[]>::value, nullptr_t> = nullptr>
     _LIBCUDACXX_INLINE_VISIBILITY
-        constexpr span(const span<_OtherElementType, _Extent>& __other,
-                       enable_if_t<
-                          is_convertible_v<_OtherElementType(*)[], element_type (*)[]>,
-                          nullptr_t> = nullptr)
+        constexpr span(const span<_OtherElementType, _Extent>& __other) noexcept
         : __data{__other.data()} {}
 
-    template <class _OtherElementType>
+    template <class _OtherElementType, enable_if_t<
+                is_convertible<_OtherElementType(*)[], element_type (*)[]>::value, nullptr_t> = nullptr>
     _LIBCUDACXX_INLINE_VISIBILITY
-        constexpr span(const span<_OtherElementType, dynamic_extent>& __other,
-                       enable_if_t<
-                          is_convertible_v<_OtherElementType(*)[], element_type (*)[]>,
-                          nullptr_t> = nullptr) noexcept
+        _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 span(const span<_OtherElementType, dynamic_extent>& __other) noexcept
         : __data{__other.data()} { _LIBCUDACXX_ASSERT(_Extent == __other.size(), "size mismatch in span's constructor (other span)"); }
 
+    template <class _Container, enable_if_t<
+                __is_span_compatible_container<_Container, _Tp>::value, nullptr_t> = nullptr>
+    _LIBCUDACXX_INLINE_VISIBILITY
+        constexpr span(      _Container& __c) noexcept(noexcept(_CUDA_VSTD::data(__c)))
+        : __data{_CUDA_VSTD::data(__c)} { _LIBCUDACXX_ASSERT(_Extent == _CUDA::VSTD::size(__c), "size mismatch in span's constructor (other span)"); }
+
+    template <class _Container, enable_if_t<
+                __is_span_compatible_container<_Container, _Tp>::value, nullptr_t> = nullptr>
+    _LIBCUDACXX_INLINE_VISIBILITY
+        constexpr span(const _Container& __c) noexcept(noexcept(_CUDA_VSTD::data(__c)))
+        : __data{_CUDA_VSTD::data(__c)} { _LIBCUDACXX_ASSERT(_Extent == _CUDA::VSTD::size(__c), "size mismatch in span's constructor (other span)"); }
 
 //  ~span() noexcept = default;
 
@@ -247,7 +263,7 @@ public:
     _LIBCUDACXX_INLINE_VISIBILITY
     constexpr span<element_type, _Count> first() const noexcept
     {
-        static_assert(_Count <= _Extent, "Count out of range in span::first()");
+        static_assert(_Count <= _Extent, "span<T, N>::first<Count>(): Count out of range");
         return {data(), _Count};
     }
 
@@ -255,21 +271,21 @@ public:
     _LIBCUDACXX_INLINE_VISIBILITY
     constexpr span<element_type, _Count> last() const noexcept
     {
-        static_assert(_Count <= _Extent, "Count out of range in span::last()");
+        static_assert(_Count <= _Extent, "span<T, N>::last<Count>(): Count out of range");
         return {data() + size() - _Count, _Count};
     }
 
     _LIBCUDACXX_INLINE_VISIBILITY
-    constexpr span<element_type, dynamic_extent> first(index_type __count) const noexcept
+    _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 span<element_type, dynamic_extent> first(size_type __count) const noexcept
     {
-        _LIBCUDACXX_ASSERT(__count <= size(), "Count out of range in span::first(count)");
+        _LIBCUDACXX_ASSERT(__count <= size(), "span<T, N>::first(count): count out of range");
         return {data(), __count};
     }
 
     _LIBCUDACXX_INLINE_VISIBILITY
-    constexpr span<element_type, dynamic_extent> last(index_type __count) const noexcept
+    _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 span<element_type, dynamic_extent> last(size_type __count) const noexcept
     {
-        _LIBCUDACXX_ASSERT(__count <= size(), "Count out of range in span::last(count)");
+        _LIBCUDACXX_ASSERT(__count <= size(), "span<T, N>::last(count): count out of range");
         return {data() + size() - __count, __count};
     }
 
@@ -278,238 +294,214 @@ public:
     constexpr auto subspan() const noexcept
         -> span<element_type, _Count != dynamic_extent ? _Count : _Extent - _Offset>
     {
-        static_assert(_Offset <= _Extent, "Offset out of range in span::subspan()");
-        return {data() + _Offset, _Count == dynamic_extent ? size() - _Offset : _Count};
+        static_assert(_Offset <= _Extent, "span<T, N>::subspan<Offset, Count>(): Offset out of range");
+        static_assert(_Count == dynamic_extent || _Count <= _Extent - _Offset, "span<T, N>::subspan<Offset, Count>(): Offset + Count out of range");
+
+        using _ReturnType = span<element_type, _Count != dynamic_extent ? _Count : _Extent - _Offset>;
+        return _ReturnType{data() + _Offset, _Count == dynamic_extent ? size() - _Offset : _Count};
     }
 
 
     _LIBCUDACXX_INLINE_VISIBILITY
-    constexpr span<element_type, dynamic_extent>
-       subspan(index_type __offset, index_type __count = dynamic_extent) const noexcept
+    _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 span<element_type, dynamic_extent>
+       subspan(size_type __offset, size_type __count = dynamic_extent) const noexcept
     {
-        _LIBCUDACXX_ASSERT(__offset <= size(), "Offset out of range in span::subspan(offset, count)");
-        _LIBCUDACXX_ASSERT(__count  <= size() || __count == dynamic_extent, "Count out of range in span::subspan(offset, count)");
+        _LIBCUDACXX_ASSERT(__offset <= size(), "span<T, N>::subspan(offset, count): offset out of range");
+        _LIBCUDACXX_ASSERT(__count  <= size() || __count == dynamic_extent, "span<T, N>::subspan(offset, count): count out of range");
         if (__count == dynamic_extent)
             return {data() + __offset, size() - __offset};
-        _LIBCUDACXX_ASSERT(__offset <= size() - __count, "count + offset out of range in span::subspan(offset, count)");
+        _LIBCUDACXX_ASSERT(__count <= size() - __offset, "span<T, N>::subspan(offset, count): offset + count out of range");
         return {data() + __offset, __count};
     }
 
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr index_type size()       const noexcept { return _Extent; }
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr index_type size_bytes() const noexcept { return _Extent * sizeof(element_type); }
+    _LIBCUDACXX_INLINE_VISIBILITY constexpr size_type size()       const noexcept { return _Extent; }
+    _LIBCUDACXX_INLINE_VISIBILITY constexpr size_type size_bytes() const noexcept { return _Extent * sizeof(element_type); }
     _LIBCUDACXX_INLINE_VISIBILITY constexpr bool empty()            const noexcept { return _Extent == 0; }
 
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr reference operator[](index_type __idx) const noexcept
+    _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 reference operator[](size_type __idx) const noexcept
     {
-        _LIBCUDACXX_ASSERT(__idx >= 0 && __idx < size(), "span<T,N>[] index out of bounds");
+        _LIBCUDACXX_ASSERT(__idx < size(), "span<T, N>::operator[](index): index out of range");
         return __data[__idx];
     }
 
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr reference front() const noexcept
+    _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 reference front() const noexcept
     {
-        static_assert(_Extent > 0, "span<T,N>[].front() on empty span");
+        _LIBCUDACXX_ASSERT(!empty(), "span<T, N>::front() on empty span");
         return __data[0];
     }
 
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr reference back() const noexcept
+    _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 reference back() const noexcept
     {
-        static_assert(_Extent > 0, "span<T,N>[].back() on empty span");
+        _LIBCUDACXX_ASSERT(!empty(), "span<T, N>::back() on empty span");
         return __data[size()-1];
     }
 
     _LIBCUDACXX_INLINE_VISIBILITY constexpr pointer data()                         const noexcept { return __data; }
 
 // [span.iter], span iterator support
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr iterator                 begin() const noexcept { return iterator(data()); }
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr iterator                   end() const noexcept { return iterator(data() + size()); }
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr const_iterator          cbegin() const noexcept { return const_iterator(data()); }
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr const_iterator            cend() const noexcept { return const_iterator(data() + size()); }
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr reverse_iterator        rbegin() const noexcept { return reverse_iterator(end()); }
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr reverse_iterator          rend() const noexcept { return reverse_iterator(begin()); }
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr const_reverse_iterator crbegin() const noexcept { return const_reverse_iterator(cend()); }
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr const_reverse_iterator   crend() const noexcept { return const_reverse_iterator(cbegin()); }
-
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr void swap(span &__other) noexcept
-    {
-        pointer __p = __data;
-        __data = __other.__data;
-        __other.__data = __p;
-    }
+    _LIBCUDACXX_INLINE_VISIBILITY constexpr iterator begin() const noexcept { return iterator(data()); }
+    _LIBCUDACXX_INLINE_VISIBILITY constexpr iterator end() const noexcept { return iterator(data() + size()); }
+    _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX14 reverse_iterator rbegin() const noexcept { return reverse_iterator(end()); }
+    _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX14 reverse_iterator rend() const noexcept { return reverse_iterator(begin()); }
 
     _LIBCUDACXX_INLINE_VISIBILITY span<const byte, _Extent * sizeof(element_type)> __as_bytes() const noexcept
-    { return {reinterpret_cast<const byte *>(data()), size_bytes()}; }
+    { return span<const byte, _Extent * sizeof(element_type)>{reinterpret_cast<const byte *>(data()), size_bytes()}; }
 
     _LIBCUDACXX_INLINE_VISIBILITY span<byte, _Extent * sizeof(element_type)> __as_writable_bytes() const noexcept
-    { return {reinterpret_cast<byte *>(data()), size_bytes()}; }
+    { return span<byte, _Extent * sizeof(element_type)>{reinterpret_cast<byte *>(data()), size_bytes()}; }
 
 private:
     pointer    __data;
-
 };
 
 
 template <typename _Tp>
 class _LIBCUDACXX_TEMPLATE_VIS span<_Tp, dynamic_extent> {
-private:
-
 public:
 //  constants and types
     using element_type           = _Tp;
     using value_type             = remove_cv_t<_Tp>;
-    using index_type             = size_t;
+    using size_type              = size_t;
     using difference_type        = ptrdiff_t;
     using pointer                = _Tp *;
     using const_pointer          = const _Tp *;
     using reference              = _Tp &;
     using const_reference        = const _Tp &;
     using iterator               =  __wrap_iter<pointer>;
-    using const_iterator         =  __wrap_iter<const_pointer>;
     using reverse_iterator       = _CUDA_VSTD::reverse_iterator<iterator>;
-    using const_reverse_iterator = _CUDA_VSTD::reverse_iterator<const_iterator>;
 
-    static constexpr index_type extent = dynamic_extent;
+    static constexpr size_type extent = dynamic_extent;
 
 // [span.cons], span constructors, copy, assignment, and destructor
     _LIBCUDACXX_INLINE_VISIBILITY constexpr span() noexcept : __data{nullptr}, __size{0} {}
 
-    constexpr span           (const span&) noexcept = default;
-    constexpr span& operator=(const span&) noexcept = default;
+    span           (const span&) noexcept = default;
+    span& operator=(const span&) noexcept = default;
 
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr span(pointer __ptr, index_type __count) : __data{__ptr}, __size{__count} {}
+    _LIBCUDACXX_INLINE_VISIBILITY constexpr span(pointer __ptr, size_type __count) : __data{__ptr}, __size{__count} {}
+#if _LIBCUDACXX_STD_VER > 14
     _LIBCUDACXX_INLINE_VISIBILITY constexpr span(pointer __f, pointer __l) : __data{__f}, __size{static_cast<size_t>(distance(__f, __l))} {}
+#else
+    _LIBCUDACXX_INLINE_VISIBILITY constexpr span(pointer __f, pointer __l) : __data{__f}, __size{static_cast<size_t>(__l - __f)} {}
+#endif
 
     template <size_t _Sz>
     _LIBCUDACXX_INLINE_VISIBILITY
-    constexpr span(element_type (&__arr)[_Sz])          noexcept : __data{__arr}, __size{_Sz} {}
+    constexpr span(type_identity_t<element_type> (&__arr)[_Sz]) noexcept : __data{__arr}, __size{_Sz} {}
 
-    template <size_t _Sz>
+    template <class _OtherElementType, size_t _Sz, enable_if_t<
+                is_convertible<_OtherElementType(*)[], element_type (*)[]>::value, nullptr_t> = nullptr>
     _LIBCUDACXX_INLINE_VISIBILITY
-    constexpr span(array<value_type, _Sz>& __arr)       noexcept : __data{__arr.data()}, __size{_Sz} {}
+    _LIBCUDACXX_CONSTEXPR_AFTER_CXX14 span(array<_OtherElementType, _Sz>& __arr)       noexcept : __data{__arr.data()}, __size{_Sz} {}
 
-    template <size_t _Sz>
+    template <class _OtherElementType, size_t _Sz, enable_if_t<
+                is_convertible<_OtherElementType(*)[], element_type (*)[]>::value, nullptr_t> = nullptr>
     _LIBCUDACXX_INLINE_VISIBILITY
-    constexpr span(const array<value_type, _Sz>& __arr) noexcept : __data{__arr.data()}, __size{_Sz} {}
+    _LIBCUDACXX_CONSTEXPR_AFTER_CXX14 span(const array<_OtherElementType, _Sz>& __arr) noexcept : __data{__arr.data()}, __size{_Sz} {}
 
-    template <class _Container>
+    template <class _Container, enable_if_t<
+                __is_span_compatible_container<_Container, _Tp>::value, nullptr_t> = nullptr>
     _LIBCUDACXX_INLINE_VISIBILITY
-        constexpr span(      _Container& __c,
-            enable_if_t<__is_span_compatible_container<_Container, _Tp>::value, nullptr_t> = nullptr)
-        : __data{_CUDA_VSTD::data(__c)}, __size{(index_type) _CUDA_VSTD::size(__c)} {}
+        constexpr span(      _Container& __c)
+        : __data{_CUDA_VSTD::data(__c)}, __size{(size_type) _CUDA_VSTD::size(__c)} {}
 
-    template <class _Container>
+    template <class _Container, enable_if_t<
+                __is_span_compatible_container<_Container, _Tp>::value, nullptr_t> = nullptr>
     _LIBCUDACXX_INLINE_VISIBILITY
-        constexpr span(const _Container& __c,
-            enable_if_t<__is_span_compatible_container<const _Container, _Tp>::value, nullptr_t> = nullptr)
-        : __data{_CUDA_VSTD::data(__c)}, __size{(index_type) _CUDA_VSTD::size(__c)} {}
+        constexpr span(const _Container& __c)
+        : __data{_CUDA_VSTD::data(__c)}, __size{(size_type) _CUDA_VSTD::size(__c)} {}
 
 
-    template <class _OtherElementType, size_t _OtherExtent>
+    template <class _OtherElementType, size_t _OtherExtent, enable_if_t<
+                is_convertible<_OtherElementType(*)[], element_type (*)[]>::value, nullptr_t> = nullptr>
     _LIBCUDACXX_INLINE_VISIBILITY
-        constexpr span(const span<_OtherElementType, _OtherExtent>& __other,
-                       enable_if_t<
-                          is_convertible_v<_OtherElementType(*)[], element_type (*)[]>,
-                          nullptr_t> = nullptr) noexcept
+        constexpr span(const span<_OtherElementType, _OtherExtent>& __other) noexcept
         : __data{__other.data()}, __size{__other.size()} {}
 
 //    ~span() noexcept = default;
 
     template <size_t _Count>
     _LIBCUDACXX_INLINE_VISIBILITY
-    constexpr span<element_type, _Count> first() const noexcept
+    _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 span<element_type, _Count> first() const noexcept
     {
-        _LIBCUDACXX_ASSERT(_Count <= size(), "Count out of range in span::first()");
-        return {data(), _Count};
+        _LIBCUDACXX_ASSERT(_Count <= size(), "span<T>::first<Count>(): Count out of range");
+        return span<element_type, _Count>{data(), _Count};
     }
 
     template <size_t _Count>
     _LIBCUDACXX_INLINE_VISIBILITY
-    constexpr span<element_type, _Count> last() const noexcept
+    _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 span<element_type, _Count> last() const noexcept
     {
-        _LIBCUDACXX_ASSERT(_Count <= size(), "Count out of range in span::last()");
-        return {data() + size() - _Count, _Count};
+        _LIBCUDACXX_ASSERT(_Count <= size(), "span<T>::last<Count>(): Count out of range");
+        return span<element_type, _Count>{data() + size() - _Count, _Count};
     }
 
     _LIBCUDACXX_INLINE_VISIBILITY
-    constexpr span<element_type, dynamic_extent> first(index_type __count) const noexcept
+    _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 span<element_type, dynamic_extent> first(size_type __count) const noexcept
     {
-        _LIBCUDACXX_ASSERT(__count <= size(), "Count out of range in span::first(count)");
+        _LIBCUDACXX_ASSERT(__count <= size(), "span<T>::first(count): count out of range");
         return {data(), __count};
     }
 
     _LIBCUDACXX_INLINE_VISIBILITY
-    constexpr span<element_type, dynamic_extent> last (index_type __count) const noexcept
+    _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 span<element_type, dynamic_extent> last (size_type __count) const noexcept
     {
-        _LIBCUDACXX_ASSERT(__count <= size(), "Count out of range in span::last(count)");
+        _LIBCUDACXX_ASSERT(__count <= size(), "span<T>::last(count): count out of range");
         return {data() + size() - __count, __count};
     }
 
     template <size_t _Offset, size_t _Count = dynamic_extent>
     _LIBCUDACXX_INLINE_VISIBILITY
-    constexpr span<_Tp, dynamic_extent> subspan() const noexcept
+    _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 span<element_type, _Count> subspan() const noexcept
     {
-        _LIBCUDACXX_ASSERT(_Offset <= size(), "Offset out of range in span::subspan()");
-        _LIBCUDACXX_ASSERT(_Count == dynamic_extent || _Offset + _Count <= size(), "Count out of range in span::subspan()");
-        return {data() + _Offset, _Count == dynamic_extent ? size() - _Offset : _Count};
+        _LIBCUDACXX_ASSERT(_Offset <= size(), "span<T>::subspan<Offset, Count>(): Offset out of range");
+        _LIBCUDACXX_ASSERT(_Count == dynamic_extent || _Count <= size() - _Offset, "span<T>::subspan<Offset, Count>(): Offset + Count out of range");
+        return span<element_type, _Count>{data() + _Offset, _Count == dynamic_extent ? size() - _Offset : _Count};
     }
 
-    constexpr span<element_type, dynamic_extent>
+    _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 span<element_type, dynamic_extent>
     _LIBCUDACXX_INLINE_VISIBILITY
-    subspan(index_type __offset, index_type __count = dynamic_extent) const noexcept
+    subspan(size_type __offset, size_type __count = dynamic_extent) const noexcept
     {
-        _LIBCUDACXX_ASSERT(__offset <= size(), "Offset out of range in span::subspan(offset, count)");
-        _LIBCUDACXX_ASSERT(__count  <= size() || __count == dynamic_extent, "count out of range in span::subspan(offset, count)");
+        _LIBCUDACXX_ASSERT(__offset <= size(), "span<T>::subspan(offset, count): offset out of range");
+        _LIBCUDACXX_ASSERT(__count <= size() || __count == dynamic_extent, "span<T>::subspan(offset, count): count out of range");
         if (__count == dynamic_extent)
             return {data() + __offset, size() - __offset};
-        _LIBCUDACXX_ASSERT(__offset <= size() - __count, "Offset + count out of range in span::subspan(offset, count)");
+        _LIBCUDACXX_ASSERT(__count <= size() - __offset, "span<T>::subspan(offset, count): offset + count out of range");
         return {data() + __offset, __count};
     }
 
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr index_type size()       const noexcept { return __size; }
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr index_type size_bytes() const noexcept { return __size * sizeof(element_type); }
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr bool empty()            const noexcept { return __size == 0; }
+    _LIBCUDACXX_INLINE_VISIBILITY constexpr size_type size()       const noexcept { return __size; }
+    _LIBCUDACXX_INLINE_VISIBILITY constexpr size_type size_bytes() const noexcept { return __size * sizeof(element_type); }
+    _LIBCUDACXX_INLINE_VISIBILITY constexpr bool empty()           const noexcept { return __size == 0; }
 
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr reference operator[](index_type __idx) const noexcept
+    _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 reference operator[](size_type __idx) const noexcept
     {
-        _LIBCUDACXX_ASSERT(__idx >= 0 && __idx < size(), "span<T>[] index out of bounds");
+        _LIBCUDACXX_ASSERT(__idx < size(), "span<T>::operator[](index): index out of range");
         return __data[__idx];
     }
 
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr reference front() const noexcept
+    _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 reference front() const noexcept
     {
-        _LIBCUDACXX_ASSERT(!empty(), "span<T>[].front() on empty span");
+        _LIBCUDACXX_ASSERT(!empty(), "span<T>::front() on empty span");
         return __data[0];
     }
 
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr reference back() const noexcept
+    _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 reference back() const noexcept
     {
-        _LIBCUDACXX_ASSERT(!empty(), "span<T>[].back() on empty span");
+        _LIBCUDACXX_ASSERT(!empty(), "span<T>::back() on empty span");
         return __data[size()-1];
     }
 
 
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr pointer data()                         const noexcept { return __data; }
+    _LIBCUDACXX_INLINE_VISIBILITY constexpr pointer data() const noexcept { return __data; }
 
 // [span.iter], span iterator support
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr iterator                 begin() const noexcept { return iterator(data()); }
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr iterator                   end() const noexcept { return iterator(data() + size()); }
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr const_iterator          cbegin() const noexcept { return const_iterator(data()); }
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr const_iterator            cend() const noexcept { return const_iterator(data() + size()); }
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr reverse_iterator        rbegin() const noexcept { return reverse_iterator(end()); }
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr reverse_iterator          rend() const noexcept { return reverse_iterator(begin()); }
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr const_reverse_iterator crbegin() const noexcept { return const_reverse_iterator(cend()); }
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr const_reverse_iterator   crend() const noexcept { return const_reverse_iterator(cbegin()); }
-
-    _LIBCUDACXX_INLINE_VISIBILITY constexpr void swap(span &__other) noexcept
-    {
-        pointer __p = __data;
-        __data = __other.__data;
-        __other.__data = __p;
-
-        index_type __sz = __size;
-        __size = __other.__size;
-        __other.__size = __sz;
-    }
-
+    _LIBCUDACXX_INLINE_VISIBILITY constexpr iterator begin() const noexcept { return iterator(data()); }
+    _LIBCUDACXX_INLINE_VISIBILITY constexpr iterator end() const noexcept { return iterator(data() + size()); }
+    _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX14 reverse_iterator rbegin() const noexcept { return reverse_iterator(end()); }
+    _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX14 reverse_iterator rend() const noexcept { return reverse_iterator(begin()); }
+    
     _LIBCUDACXX_INLINE_VISIBILITY span<const byte, dynamic_extent> __as_bytes() const noexcept
     { return {reinterpret_cast<const byte *>(data()), size_bytes()}; }
 
@@ -518,37 +510,9 @@ public:
 
 private:
     pointer    __data;
-    index_type __size;
+    size_type __size;
 };
 
-//  tuple interface
-template <class _Tp, size_t _Size>
-struct _LIBCUDACXX_TEMPLATE_VIS tuple_size<span<_Tp, _Size>>
-    : public integral_constant<size_t, _Size> {};
-
-template <class _Tp>
-struct _LIBCUDACXX_TEMPLATE_VIS tuple_size<span<_Tp, dynamic_extent>>; // declared but not defined
-
-
-template <size_t _Ip, class _Tp, size_t _Size>
-struct _LIBCUDACXX_TEMPLATE_VIS tuple_element<_Ip, span<_Tp, _Size>>
-{
-    static_assert( dynamic_extent != _Size, "std::tuple_element<> not supported for std::span<T, dynamic_extent>");
-    static_assert(_Ip < _Size, "Index out of bounds in std::tuple_element<> (std::span)");
-    typedef _Tp type;
-};
-
-template <size_t _Ip, class _Tp, size_t _Size>
-_LIBCUDACXX_INLINE_VISIBILITY constexpr
-_Tp&
-get(span<_Tp, _Size> __s) noexcept
-{
-    static_assert( dynamic_extent != _Size, "std::get<> not supported for std::span<T, dynamic_extent>");
-    static_assert(_Ip < _Size, "Index out of bounds in std::get<> (std::span)");
-    return __s[_Ip];
-}
-
-
 //  as_bytes & as_writable_bytes
 template <class _Tp, size_t _Extent>
 _LIBCUDACXX_INLINE_VISIBILITY
@@ -559,14 +523,12 @@ auto as_bytes(span<_Tp, _Extent> __s) noexcept
 template <class _Tp, size_t _Extent>
 _LIBCUDACXX_INLINE_VISIBILITY
 auto as_writable_bytes(span<_Tp, _Extent> __s) noexcept
--> enable_if_t<!is_const_v<_Tp>, decltype(__s.__as_writable_bytes())>
+-> enable_if_t<!is_const<_Tp>::value, decltype(__s.__as_writable_bytes())>
 { return __s.__as_writable_bytes(); }
 
-template <class _Tp, size_t _Extent>
-_LIBCUDACXX_INLINE_VISIBILITY
-constexpr void swap(span<_Tp, _Extent> &__lhs, span<_Tp, _Extent> &__rhs) noexcept
-{ __lhs.swap(__rhs); }
+#endif // _LIBCUDACXX_STD_VER > 11
 
+#if _LIBCUDACXX_STD_VER > 14
 
 //  Deduction guides
 template<class _Tp, size_t _Sz>
@@ -584,7 +546,7 @@ template<class _Container>
 template<class _Container>
     span(const _Container&) -> span<const typename _Container::value_type>;
 
-#endif // _LIBCUDACXX_STD_VER > 17
+#endif // _LIBCUDACXX_STD_VER > 14
 
 _LIBCUDACXX_END_NAMESPACE_STD
 
diff --git a/include/cuda/std/detail/libcxx/include/type_traits b/include/cuda/std/detail/libcxx/include/type_traits
index 44f7c654a7..52217b830b 100644
--- a/include/cuda/std/detail/libcxx/include/type_traits
+++ b/include/cuda/std/detail/libcxx/include/type_traits
@@ -4,6 +4,7 @@
 // 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 //
 //===----------------------------------------------------------------------===//
 
diff --git a/include/cuda/std/detail/libcxx/include/version b/include/cuda/std/detail/libcxx/include/version
index 736c4343f5..0871d45a9e 100644
--- a/include/cuda/std/detail/libcxx/include/version
+++ b/include/cuda/std/detail/libcxx/include/version
@@ -140,6 +140,7 @@ __cpp_lib_void_t                                        201411L <type_traits>
 # if !defined(_LIBCUDACXX_HAS_NO_THREADS)
 #   define __cpp_lib_shared_timed_mutex                 201402L
 # endif
+# define __cpp_lib_span                                 201902L
 # define __cpp_lib_string_udls                          201304L
 # define __cpp_lib_transformation_trait_aliases         201304L
 # define __cpp_lib_transparent_operators                201210L
diff --git a/include/cuda/std/span b/include/cuda/std/span
new file mode 100644
index 0000000000..12ec1e4721
--- /dev/null
+++ b/include/cuda/std/span
@@ -0,0 +1,29 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDA_SPAN
+#define _CUDA_SPAN
+
+#include "array"
+#include "cstddef"
+#include "type_traits"
+#include "version"
+#include "iterator"
+#include "utility"
+
+#include "detail/__config"
+
+#include "detail/__pragma_push"
+
+#include "detail/libcxx/include/span"
+
+#include "detail/__pragma_pop"
+
+#endif //_CUDA_SPAN