summaryrefslogtreecommitdiff
path: root/ext/mdspan/include
diff options
context:
space:
mode:
authorTomi Valkeinen <tomi.valkeinen@ideasonboard.com>2025-02-05 19:09:43 +0200
committerTomi Valkeinen <tomi.valkeinen@ideasonboard.com>2025-03-26 15:44:00 +0200
commit6c49fe5b811464f59e3a31b869734071da0ec7c1 (patch)
tree3fb287472a670b1efe1866906db1ac67229d6d96 /ext/mdspan/include
parent9b2a7728b2b0b26065ba79cfbbd20f783f4a9988 (diff)
Add mdspan includes
From https://github.com/kokkos/mdspan Signed-off-by: Tomi Valkeinen <tomi.valkeinen@ideasonboard.com>
Diffstat (limited to 'ext/mdspan/include')
-rw-r--r--ext/mdspan/include/experimental/__p0009_bits/compressed_pair.hpp195
-rw-r--r--ext/mdspan/include/experimental/__p0009_bits/config.hpp290
-rw-r--r--ext/mdspan/include/experimental/__p0009_bits/default_accessor.hpp56
-rw-r--r--ext/mdspan/include/experimental/__p0009_bits/dynamic_extent.hpp35
-rw-r--r--ext/mdspan/include/experimental/__p0009_bits/extents.hpp693
-rw-r--r--ext/mdspan/include/experimental/__p0009_bits/full_extent_t.hpp26
-rw-r--r--ext/mdspan/include/experimental/__p0009_bits/layout_left.hpp269
-rw-r--r--ext/mdspan/include/experimental/__p0009_bits/layout_right.hpp265
-rw-r--r--ext/mdspan/include/experimental/__p0009_bits/layout_stride.hpp667
-rw-r--r--ext/mdspan/include/experimental/__p0009_bits/macros.hpp699
-rw-r--r--ext/mdspan/include/experimental/__p0009_bits/mdspan.hpp432
-rw-r--r--ext/mdspan/include/experimental/__p0009_bits/no_unique_address.hpp97
-rw-r--r--ext/mdspan/include/experimental/__p0009_bits/trait_backports.hpp132
-rw-r--r--ext/mdspan/include/experimental/__p0009_bits/type_list.hpp87
-rw-r--r--ext/mdspan/include/experimental/__p0009_bits/utility.hpp172
-rw-r--r--ext/mdspan/include/experimental/__p1684_bits/mdarray.hpp460
-rw-r--r--ext/mdspan/include/experimental/__p2389_bits/dims.hpp28
-rw-r--r--ext/mdspan/include/experimental/__p2630_bits/strided_slice.hpp48
-rw-r--r--ext/mdspan/include/experimental/__p2630_bits/submdspan.hpp40
-rw-r--r--ext/mdspan/include/experimental/__p2630_bits/submdspan_extents.hpp418
-rw-r--r--ext/mdspan/include/experimental/__p2630_bits/submdspan_mapping.hpp630
-rw-r--r--ext/mdspan/include/experimental/__p2642_bits/layout_padded.hpp869
-rw-r--r--ext/mdspan/include/experimental/__p2642_bits/layout_padded_fwd.hpp137
-rw-r--r--ext/mdspan/include/experimental/mdarray28
-rw-r--r--ext/mdspan/include/experimental/mdspan39
-rw-r--r--ext/mdspan/include/mdspan/mdarray.hpp31
-rw-r--r--ext/mdspan/include/mdspan/mdspan.hpp43
27 files changed, 6886 insertions, 0 deletions
diff --git a/ext/mdspan/include/experimental/__p0009_bits/compressed_pair.hpp b/ext/mdspan/include/experimental/__p0009_bits/compressed_pair.hpp
new file mode 100644
index 0000000..25389a2
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p0009_bits/compressed_pair.hpp
@@ -0,0 +1,195 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+#pragma once
+
+#include "macros.hpp"
+#include "trait_backports.hpp"
+
+#if !defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+# include "no_unique_address.hpp"
+#endif
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+namespace detail {
+
+// For no unique address emulation, this is the case taken when neither are empty.
+// For real `[[no_unique_address]]`, this case is always taken.
+template <class _T1, class _T2, class _Enable = void> struct __compressed_pair {
+ _MDSPAN_NO_UNIQUE_ADDRESS _T1 __t1_val{};
+ _MDSPAN_NO_UNIQUE_ADDRESS _T2 __t2_val{};
+ MDSPAN_FORCE_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 _T1 &__first() noexcept { return __t1_val; }
+ MDSPAN_FORCE_INLINE_FUNCTION constexpr _T1 const &__first() const noexcept {
+ return __t1_val;
+ }
+ MDSPAN_FORCE_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 _T2 &__second() noexcept { return __t2_val; }
+ MDSPAN_FORCE_INLINE_FUNCTION constexpr _T2 const &__second() const noexcept {
+ return __t2_val;
+ }
+
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ constexpr __compressed_pair() = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ constexpr __compressed_pair(__compressed_pair const &) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ constexpr __compressed_pair(__compressed_pair &&) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ _MDSPAN_CONSTEXPR_14_DEFAULTED __compressed_pair &
+ operator=(__compressed_pair const &) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ _MDSPAN_CONSTEXPR_14_DEFAULTED __compressed_pair &
+ operator=(__compressed_pair &&) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ ~__compressed_pair() = default;
+ template <class _T1Like, class _T2Like>
+ MDSPAN_INLINE_FUNCTION constexpr __compressed_pair(_T1Like &&__t1, _T2Like &&__t2)
+ : __t1_val((_T1Like &&) __t1), __t2_val((_T2Like &&) __t2) {}
+};
+
+#if !defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+
+// First empty.
+template <class _T1, class _T2>
+struct __compressed_pair<
+ _T1, _T2,
+ std::enable_if_t<_MDSPAN_TRAIT(std::is_empty, _T1) && !_MDSPAN_TRAIT(std::is_empty, _T2)>>
+ : private _T1 {
+ _T2 __t2_val{};
+ MDSPAN_FORCE_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 _T1 &__first() noexcept {
+ return *static_cast<_T1 *>(this);
+ }
+ MDSPAN_FORCE_INLINE_FUNCTION constexpr _T1 const &__first() const noexcept {
+ return *static_cast<_T1 const *>(this);
+ }
+ MDSPAN_FORCE_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 _T2 &__second() noexcept { return __t2_val; }
+ MDSPAN_FORCE_INLINE_FUNCTION constexpr _T2 const &__second() const noexcept {
+ return __t2_val;
+ }
+
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ constexpr __compressed_pair() = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ constexpr __compressed_pair(__compressed_pair const &) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ constexpr __compressed_pair(__compressed_pair &&) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ _MDSPAN_CONSTEXPR_14_DEFAULTED __compressed_pair &
+ operator=(__compressed_pair const &) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ _MDSPAN_CONSTEXPR_14_DEFAULTED __compressed_pair &
+ operator=(__compressed_pair &&) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ ~__compressed_pair() = default;
+ template <class _T1Like, class _T2Like>
+ MDSPAN_INLINE_FUNCTION constexpr __compressed_pair(_T1Like &&__t1, _T2Like &&__t2)
+ : _T1((_T1Like &&) __t1), __t2_val((_T2Like &&) __t2) {}
+};
+
+// Second empty.
+template <class _T1, class _T2>
+struct __compressed_pair<
+ _T1, _T2,
+ std::enable_if_t<!_MDSPAN_TRAIT(std::is_empty, _T1) && _MDSPAN_TRAIT(std::is_empty, _T2)>>
+ : private _T2 {
+ _T1 __t1_val{};
+ MDSPAN_FORCE_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 _T1 &__first() noexcept { return __t1_val; }
+ MDSPAN_FORCE_INLINE_FUNCTION constexpr _T1 const &__first() const noexcept {
+ return __t1_val;
+ }
+ MDSPAN_FORCE_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 _T2 &__second() noexcept {
+ return *static_cast<_T2 *>(this);
+ }
+ MDSPAN_FORCE_INLINE_FUNCTION constexpr _T2 const &__second() const noexcept {
+ return *static_cast<_T2 const *>(this);
+ }
+
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ constexpr __compressed_pair() = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ constexpr __compressed_pair(__compressed_pair const &) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ constexpr __compressed_pair(__compressed_pair &&) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ _MDSPAN_CONSTEXPR_14_DEFAULTED __compressed_pair &
+ operator=(__compressed_pair const &) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ _MDSPAN_CONSTEXPR_14_DEFAULTED __compressed_pair &
+ operator=(__compressed_pair &&) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ ~__compressed_pair() = default;
+
+ template <class _T1Like, class _T2Like>
+ MDSPAN_INLINE_FUNCTION constexpr __compressed_pair(_T1Like &&__t1, _T2Like &&__t2)
+ : _T2((_T2Like &&) __t2), __t1_val((_T1Like &&) __t1) {}
+};
+
+// Both empty.
+template <class _T1, class _T2>
+struct __compressed_pair<
+ _T1, _T2,
+ std::enable_if_t<_MDSPAN_TRAIT(std::is_empty, _T1) && _MDSPAN_TRAIT(std::is_empty, _T2)>>
+ // We need to use the __no_unique_address_emulation wrapper here to avoid
+ // base class ambiguities.
+#ifdef _MDSPAN_COMPILER_MSVC
+// MSVC doesn't allow you to access public static member functions of a type
+// when you *happen* to privately inherit from that type.
+ : protected __no_unique_address_emulation<_T1, 0>,
+ protected __no_unique_address_emulation<_T2, 1>
+#else
+ : private __no_unique_address_emulation<_T1, 0>,
+ private __no_unique_address_emulation<_T2, 1>
+#endif
+{
+ using __first_base_t = __no_unique_address_emulation<_T1, 0>;
+ using __second_base_t = __no_unique_address_emulation<_T2, 1>;
+
+ MDSPAN_FORCE_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 _T1 &__first() noexcept {
+ return this->__first_base_t::__ref();
+ }
+ MDSPAN_FORCE_INLINE_FUNCTION constexpr _T1 const &__first() const noexcept {
+ return this->__first_base_t::__ref();
+ }
+ MDSPAN_FORCE_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 _T2 &__second() noexcept {
+ return this->__second_base_t::__ref();
+ }
+ MDSPAN_FORCE_INLINE_FUNCTION constexpr _T2 const &__second() const noexcept {
+ return this->__second_base_t::__ref();
+ }
+
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ constexpr __compressed_pair() = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ constexpr __compressed_pair(__compressed_pair const &) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ constexpr __compressed_pair(__compressed_pair &&) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ _MDSPAN_CONSTEXPR_14_DEFAULTED __compressed_pair &
+ operator=(__compressed_pair const &) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ _MDSPAN_CONSTEXPR_14_DEFAULTED __compressed_pair &
+ operator=(__compressed_pair &&) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ ~__compressed_pair() = default;
+ template <class _T1Like, class _T2Like>
+ MDSPAN_INLINE_FUNCTION constexpr __compressed_pair(_T1Like &&__t1, _T2Like &&__t2) noexcept
+ : __first_base_t(_T1((_T1Like &&) __t1)),
+ __second_base_t(_T2((_T2Like &&) __t2))
+ { }
+};
+
+#endif // !defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+
+} // end namespace detail
+} // end namespace MDSPAN_IMPL_STANDARD_NAMESPACE
diff --git a/ext/mdspan/include/experimental/__p0009_bits/config.hpp b/ext/mdspan/include/experimental/__p0009_bits/config.hpp
new file mode 100644
index 0000000..a23d013
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p0009_bits/config.hpp
@@ -0,0 +1,290 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+#pragma once
+
+#ifndef __has_include
+# define __has_include(x) 0
+#endif
+
+#if __has_include(<version>)
+# include <version>
+#else
+# include <type_traits>
+# include <utility>
+#endif
+
+#ifdef _MSVC_LANG
+#define _MDSPAN_CPLUSPLUS _MSVC_LANG
+#else
+#define _MDSPAN_CPLUSPLUS __cplusplus
+#endif
+
+#define MDSPAN_CXX_STD_14 201402L
+#define MDSPAN_CXX_STD_17 201703L
+#define MDSPAN_CXX_STD_20 202002L
+// Note GCC has not updated this in version 13
+#ifdef __clang__
+#define MDSPAN_CXX_STD_23 202302L
+#else
+#define MDSPAN_CXX_STD_23 202100L
+#endif
+
+#define MDSPAN_HAS_CXX_14 (_MDSPAN_CPLUSPLUS >= MDSPAN_CXX_STD_14)
+#define MDSPAN_HAS_CXX_17 (_MDSPAN_CPLUSPLUS >= MDSPAN_CXX_STD_17)
+#define MDSPAN_HAS_CXX_20 (_MDSPAN_CPLUSPLUS >= MDSPAN_CXX_STD_20)
+#define MDSPAN_HAS_CXX_23 (_MDSPAN_CPLUSPLUS >= MDSPAN_CXX_STD_23)
+
+static_assert(_MDSPAN_CPLUSPLUS >= MDSPAN_CXX_STD_14, "mdspan requires C++14 or later.");
+
+#ifndef _MDSPAN_COMPILER_CLANG
+# if defined(__clang__)
+# define _MDSPAN_COMPILER_CLANG __clang__
+# endif
+#endif
+
+#if !defined(_MDSPAN_COMPILER_MSVC) && !defined(_MDSPAN_COMPILER_MSVC_CLANG)
+# if defined(_MSC_VER)
+# if !defined(_MDSPAN_COMPILER_CLANG)
+# define _MDSPAN_COMPILER_MSVC _MSC_VER
+# else
+# define _MDSPAN_COMPILER_MSVC_CLANG _MSC_VER
+# endif
+# endif
+#endif
+
+#ifndef _MDSPAN_COMPILER_INTEL
+# ifdef __INTEL_COMPILER
+# define _MDSPAN_COMPILER_INTEL __INTEL_COMPILER
+# endif
+#endif
+
+#ifndef _MDSPAN_COMPILER_APPLECLANG
+# ifdef __apple_build_version__
+# define _MDSPAN_COMPILER_APPLECLANG __apple_build_version__
+# endif
+#endif
+
+#ifndef _MDSPAN_HAS_CUDA
+# if defined(__CUDACC__)
+# define _MDSPAN_HAS_CUDA __CUDACC__
+# endif
+#endif
+
+#ifndef _MDSPAN_HAS_HIP
+# if defined(__HIPCC__)
+# define _MDSPAN_HAS_HIP __HIPCC__
+# endif
+#endif
+
+#ifndef _MDSPAN_HAS_SYCL
+# if defined(SYCL_LANGUAGE_VERSION)
+# define _MDSPAN_HAS_SYCL SYCL_LANGUAGE_VERSION
+# endif
+#endif
+
+#ifndef __has_cpp_attribute
+# define __has_cpp_attribute(x) 0
+#endif
+
+#ifndef _MDSPAN_PRESERVE_STANDARD_LAYOUT
+// Preserve standard layout by default, but we're not removing the old version
+// that turns this off until we're sure this doesn't have an unreasonable cost
+// to the compiler or optimizer.
+# define _MDSPAN_PRESERVE_STANDARD_LAYOUT 1
+#endif
+
+#if !defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+# if ((__has_cpp_attribute(no_unique_address) >= 201803L) && \
+ (!defined(__NVCC__) || MDSPAN_HAS_CXX_20) && \
+ (!defined(_MDSPAN_COMPILER_MSVC) || MDSPAN_HAS_CXX_20))
+# define _MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS 1
+# define _MDSPAN_NO_UNIQUE_ADDRESS [[no_unique_address]]
+# else
+# define _MDSPAN_NO_UNIQUE_ADDRESS
+# endif
+#endif
+
+// NVCC older than 11.6 chokes on the no-unique-address-emulation
+// so just pretend to use it (to avoid the full blown EBO workaround
+// which NVCC also doesn't like ...), and leave the macro empty
+#ifndef _MDSPAN_NO_UNIQUE_ADDRESS
+# if defined(__NVCC__)
+# define _MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS 1
+# define _MDSPAN_USE_FAKE_ATTRIBUTE_NO_UNIQUE_ADDRESS
+# endif
+# define _MDSPAN_NO_UNIQUE_ADDRESS
+#endif
+
+// AMDs HIP compiler seems to have issues with concepts
+// it pretends concepts exist, but doesn't ship <concept>
+#ifndef __HIPCC__
+#ifndef _MDSPAN_USE_CONCEPTS
+# if defined(__cpp_concepts) && __cpp_concepts >= 201507L
+# define _MDSPAN_USE_CONCEPTS 1
+# endif
+#endif
+#endif
+
+#ifndef _MDSPAN_USE_FOLD_EXPRESSIONS
+# if (defined(__cpp_fold_expressions) && __cpp_fold_expressions >= 201603L) \
+ || (!defined(__cpp_fold_expressions) && MDSPAN_HAS_CXX_17)
+# define _MDSPAN_USE_FOLD_EXPRESSIONS 1
+# endif
+#endif
+
+#ifndef _MDSPAN_USE_INLINE_VARIABLES
+# if defined(__cpp_inline_variables) && __cpp_inline_variables >= 201606L \
+ || (!defined(__cpp_inline_variables) && MDSPAN_HAS_CXX_17)
+# define _MDSPAN_USE_INLINE_VARIABLES 1
+# endif
+#endif
+
+#ifndef _MDSPAN_NEEDS_TRAIT_VARIABLE_TEMPLATE_BACKPORTS
+# if (!(defined(__cpp_lib_type_trait_variable_templates) && __cpp_lib_type_trait_variable_templates >= 201510L) \
+ || !MDSPAN_HAS_CXX_17)
+# if !(defined(_MDSPAN_COMPILER_APPLECLANG) && MDSPAN_HAS_CXX_17)
+# define _MDSPAN_NEEDS_TRAIT_VARIABLE_TEMPLATE_BACKPORTS 1
+# endif
+# endif
+#endif
+
+#ifndef _MDSPAN_USE_VARIABLE_TEMPLATES
+# if (defined(__cpp_variable_templates) && __cpp_variable_templates >= 201304 && MDSPAN_HAS_CXX_17) \
+ || (!defined(__cpp_variable_templates) && MDSPAN_HAS_CXX_17)
+# define _MDSPAN_USE_VARIABLE_TEMPLATES 1
+# endif
+#endif // _MDSPAN_USE_VARIABLE_TEMPLATES
+
+#ifndef _MDSPAN_USE_CONSTEXPR_14
+# if (defined(__cpp_constexpr) && __cpp_constexpr >= 201304) \
+ || (!defined(__cpp_constexpr) && MDSPAN_HAS_CXX_14) \
+ && (!(defined(__INTEL_COMPILER) && __INTEL_COMPILER <= 1700))
+# define _MDSPAN_USE_CONSTEXPR_14 1
+# endif
+#endif
+
+#ifndef _MDSPAN_USE_INTEGER_SEQUENCE
+# if defined(_MDSPAN_COMPILER_MSVC)
+# if (defined(__cpp_lib_integer_sequence) && __cpp_lib_integer_sequence >= 201304)
+# define _MDSPAN_USE_INTEGER_SEQUENCE 1
+# endif
+# endif
+#endif
+#ifndef _MDSPAN_USE_INTEGER_SEQUENCE
+# if (defined(__cpp_lib_integer_sequence) && __cpp_lib_integer_sequence >= 201304) \
+ || (!defined(__cpp_lib_integer_sequence) && MDSPAN_HAS_CXX_14) \
+ /* as far as I can tell, libc++ seems to think this is a C++11 feature... */ \
+ || (defined(__GLIBCXX__) && __GLIBCXX__ > 20150422 && __GNUC__ < 5 && !defined(__INTEL_CXX11_MODE__))
+ // several compilers lie about integer_sequence working properly unless the C++14 standard is used
+# define _MDSPAN_USE_INTEGER_SEQUENCE 1
+# elif defined(_MDSPAN_COMPILER_APPLECLANG) && MDSPAN_HAS_CXX_14
+ // appleclang seems to be missing the __cpp_lib_... macros, but doesn't seem to lie about C++14 making
+ // integer_sequence work
+# define _MDSPAN_USE_INTEGER_SEQUENCE 1
+# endif
+#endif
+
+#ifndef _MDSPAN_USE_RETURN_TYPE_DEDUCTION
+# if (defined(__cpp_return_type_deduction) && __cpp_return_type_deduction >= 201304) \
+ || (!defined(__cpp_return_type_deduction) && MDSPAN_HAS_CXX_14)
+# define _MDSPAN_USE_RETURN_TYPE_DEDUCTION 1
+# endif
+#endif
+
+#ifndef _MDSPAN_USE_CLASS_TEMPLATE_ARGUMENT_DEDUCTION
+# if (!defined(__NVCC__) || (__CUDACC_VER_MAJOR__ * 100 + __CUDACC_VER_MINOR__ * 10 >= 1170)) && \
+ ((defined(__cpp_deduction_guides) && __cpp_deduction_guides >= 201703) || \
+ (!defined(__cpp_deduction_guides) && MDSPAN_HAS_CXX_17))
+# define _MDSPAN_USE_CLASS_TEMPLATE_ARGUMENT_DEDUCTION 1
+# endif
+#endif
+
+#ifndef _MDSPAN_USE_STANDARD_TRAIT_ALIASES
+# if (defined(__cpp_lib_transformation_trait_aliases) && __cpp_lib_transformation_trait_aliases >= 201304) \
+ || (!defined(__cpp_lib_transformation_trait_aliases) && MDSPAN_HAS_CXX_14)
+# define _MDSPAN_USE_STANDARD_TRAIT_ALIASES 1
+# elif defined(_MDSPAN_COMPILER_APPLECLANG) && MDSPAN_HAS_CXX_14
+ // appleclang seems to be missing the __cpp_lib_... macros, but doesn't seem to lie about C++14
+# define _MDSPAN_USE_STANDARD_TRAIT_ALIASES 1
+# endif
+#endif
+
+#ifndef _MDSPAN_DEFAULTED_CONSTRUCTORS_INHERITANCE_WORKAROUND
+# ifdef __GNUC__
+# if __GNUC__ < 9
+# define _MDSPAN_DEFAULTED_CONSTRUCTORS_INHERITANCE_WORKAROUND 1
+# endif
+# endif
+#endif
+
+#ifndef MDSPAN_CONDITIONAL_EXPLICIT
+# if MDSPAN_HAS_CXX_20
+# define MDSPAN_CONDITIONAL_EXPLICIT(COND) explicit(COND)
+# else
+# define MDSPAN_CONDITIONAL_EXPLICIT(COND)
+# endif
+#endif
+
+#ifndef MDSPAN_USE_BRACKET_OPERATOR
+# if defined(__cpp_multidimensional_subscript)
+// The following if/else is necessary to workaround a clang issue
+// relative to using a parameter pack inside a bracket operator in C++2b/C++23 mode
+# if defined(_MDSPAN_COMPILER_CLANG) && \
+ ((__clang_major__ < 17) || \
+ (__clang_major__ == 17 && __clang_minor__ == 0 && \
+ __clang_patchlevel__ == 0))
+# define MDSPAN_USE_BRACKET_OPERATOR 0
+# else
+# define MDSPAN_USE_BRACKET_OPERATOR 1
+# endif
+# else
+# define MDSPAN_USE_BRACKET_OPERATOR 0
+# endif
+#endif
+
+#ifndef MDSPAN_USE_PAREN_OPERATOR
+# if !MDSPAN_USE_BRACKET_OPERATOR
+# define MDSPAN_USE_PAREN_OPERATOR 1
+# else
+# define MDSPAN_USE_PAREN_OPERATOR 0
+# endif
+#endif
+
+#if MDSPAN_USE_BRACKET_OPERATOR
+# define __MDSPAN_OP(mds,...) mds[__VA_ARGS__]
+// Corentins demo compiler for subscript chokes on empty [] call,
+// though I believe the proposal supports it?
+#ifdef MDSPAN_NO_EMPTY_BRACKET_OPERATOR
+# define __MDSPAN_OP0(mds) mds.accessor().access(mds.data_handle(),0)
+#else
+# define __MDSPAN_OP0(mds) mds[]
+#endif
+# define __MDSPAN_OP1(mds, a) mds[a]
+# define __MDSPAN_OP2(mds, a, b) mds[a,b]
+# define __MDSPAN_OP3(mds, a, b, c) mds[a,b,c]
+# define __MDSPAN_OP4(mds, a, b, c, d) mds[a,b,c,d]
+# define __MDSPAN_OP5(mds, a, b, c, d, e) mds[a,b,c,d,e]
+# define __MDSPAN_OP6(mds, a, b, c, d, e, f) mds[a,b,c,d,e,f]
+#else
+# define __MDSPAN_OP(mds,...) mds(__VA_ARGS__)
+# define __MDSPAN_OP0(mds) mds()
+# define __MDSPAN_OP1(mds, a) mds(a)
+# define __MDSPAN_OP2(mds, a, b) mds(a,b)
+# define __MDSPAN_OP3(mds, a, b, c) mds(a,b,c)
+# define __MDSPAN_OP4(mds, a, b, c, d) mds(a,b,c,d)
+# define __MDSPAN_OP5(mds, a, b, c, d, e) mds(a,b,c,d,e)
+# define __MDSPAN_OP6(mds, a, b, c, d, e, f) mds(a,b,c,d,e,f)
+#endif
diff --git a/ext/mdspan/include/experimental/__p0009_bits/default_accessor.hpp b/ext/mdspan/include/experimental/__p0009_bits/default_accessor.hpp
new file mode 100644
index 0000000..ea0f537
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p0009_bits/default_accessor.hpp
@@ -0,0 +1,56 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+#pragma once
+
+#include "macros.hpp"
+
+#include <cstddef> // size_t
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+
+template <class ElementType>
+struct default_accessor {
+
+ using offset_policy = default_accessor;
+ using element_type = ElementType;
+ using reference = ElementType&;
+ using data_handle_type = ElementType*;
+
+ MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr default_accessor() noexcept = default;
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherElementType,
+ /* requires */ (
+ _MDSPAN_TRAIT(std::is_convertible, OtherElementType(*)[], element_type(*)[])
+ )
+ )
+ MDSPAN_INLINE_FUNCTION
+ constexpr default_accessor(default_accessor<OtherElementType>) noexcept {}
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr data_handle_type
+ offset(data_handle_type p, size_t i) const noexcept {
+ return p + i;
+ }
+
+ MDSPAN_FORCE_INLINE_FUNCTION
+ constexpr reference access(data_handle_type p, size_t i) const noexcept {
+ return p[i];
+ }
+
+};
+
+} // end namespace MDSPAN_IMPL_STANDARD_NAMESPACE
diff --git a/ext/mdspan/include/experimental/__p0009_bits/dynamic_extent.hpp b/ext/mdspan/include/experimental/__p0009_bits/dynamic_extent.hpp
new file mode 100644
index 0000000..2e29da1
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p0009_bits/dynamic_extent.hpp
@@ -0,0 +1,35 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+#pragma once
+
+#include "macros.hpp"
+
+#if defined(__cpp_lib_span)
+#include <span>
+#endif
+
+#include <cstddef> // size_t
+#include <limits> // numeric_limits
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+#if defined(__cpp_lib_span)
+using std::dynamic_extent;
+#else
+_MDSPAN_INLINE_VARIABLE constexpr auto dynamic_extent = std::numeric_limits<size_t>::max();
+#endif
+} // namespace MDSPAN_IMPL_STANDARD_NAMESPACE
+
+//==============================================================================================================
diff --git a/ext/mdspan/include/experimental/__p0009_bits/extents.hpp b/ext/mdspan/include/experimental/__p0009_bits/extents.hpp
new file mode 100644
index 0000000..d58d377
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p0009_bits/extents.hpp
@@ -0,0 +1,693 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+
+#pragma once
+#include "dynamic_extent.hpp"
+#include "utility.hpp"
+
+#ifdef __cpp_lib_span
+#include <span>
+#endif
+#include <array>
+#include <type_traits>
+
+#include <cassert>
+#include <cinttypes>
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+namespace detail {
+
+// Function used to check compatibility of extents in converting constructor
+// can't be a private member function for some reason.
+template <size_t... Extents, size_t... OtherExtents>
+MDSPAN_INLINE_FUNCTION
+static constexpr std::integral_constant<bool, false> __check_compatible_extents(
+ std::integral_constant<bool, false>,
+ std::integer_sequence<size_t, Extents...>,
+ std::integer_sequence<size_t, OtherExtents...>) noexcept {
+ return {};
+}
+
+// This helper prevents ICE's on MSVC.
+template <size_t Lhs, size_t Rhs>
+struct __compare_extent_compatible : std::integral_constant<bool,
+ Lhs == dynamic_extent ||
+ Rhs == dynamic_extent ||
+ Lhs == Rhs>
+{};
+
+template <size_t... Extents, size_t... OtherExtents>
+MDSPAN_INLINE_FUNCTION
+static constexpr std::integral_constant<
+ bool, _MDSPAN_FOLD_AND(__compare_extent_compatible<Extents, OtherExtents>::value)>
+__check_compatible_extents(
+ std::integral_constant<bool, true>,
+ std::integer_sequence<size_t, Extents...>,
+ std::integer_sequence<size_t, OtherExtents...>) noexcept {
+ return {};
+}
+
+template<class IndexType, class ... Arguments>
+MDSPAN_INLINE_FUNCTION
+static constexpr bool are_valid_indices() {
+ return
+ _MDSPAN_FOLD_AND(std::is_convertible<Arguments, IndexType>::value) &&
+ _MDSPAN_FOLD_AND(std::is_nothrow_constructible<IndexType, Arguments>::value);
+}
+
+// ------------------------------------------------------------------
+// ------------ static_array ----------------------------------------
+// ------------------------------------------------------------------
+
+// array like class which provides an array of static values with get
+// function and operator [].
+
+// Implementation of Static Array with recursive implementation of get.
+template <size_t R, class T, T... Extents> struct static_array_impl;
+
+template <size_t R, class T, T FirstExt, T... Extents>
+struct static_array_impl<R, T, FirstExt, Extents...> {
+ MDSPAN_INLINE_FUNCTION
+ constexpr static T get(size_t r) {
+ if (r == R)
+ return FirstExt;
+ else
+ return static_array_impl<R + 1, T, Extents...>::get(r);
+ }
+ template <size_t r> MDSPAN_INLINE_FUNCTION constexpr static T get() {
+#if MDSPAN_HAS_CXX_17
+ if constexpr (r == R)
+ return FirstExt;
+ else
+ return static_array_impl<R + 1, T, Extents...>::template get<r>();
+#else
+ get(r);
+#endif
+ }
+};
+
+// End the recursion
+template <size_t R, class T, T FirstExt>
+struct static_array_impl<R, T, FirstExt> {
+ MDSPAN_INLINE_FUNCTION
+ constexpr static T get(size_t) { return FirstExt; }
+ template <size_t> MDSPAN_INLINE_FUNCTION constexpr static T get() {
+ return FirstExt;
+ }
+};
+
+// Don't start recursion if size 0
+template <class T> struct static_array_impl<0, T> {
+ MDSPAN_INLINE_FUNCTION
+ constexpr static T get(size_t) { return T(); }
+ template <size_t> MDSPAN_INLINE_FUNCTION constexpr static T get() {
+ return T();
+ }
+};
+
+// Static array, provides get<r>(), get(r) and operator[r]
+template <class T, T... Values> struct static_array:
+ public static_array_impl<0, T, Values...> {
+
+public:
+ using value_type = T;
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr static size_t size() { return sizeof...(Values); }
+};
+
+
+// ------------------------------------------------------------------
+// ------------ index_sequence_scan ---------------------------------
+// ------------------------------------------------------------------
+
+// index_sequence_scan takes compile time values and provides get(r)
+// and get<r>() which return the sum of the first r-1 values.
+
+// Recursive implementation for get
+template <size_t R, size_t... Values> struct index_sequence_scan_impl;
+
+template <size_t R, size_t FirstVal, size_t... Values>
+struct index_sequence_scan_impl<R, FirstVal, Values...> {
+ MDSPAN_INLINE_FUNCTION
+ constexpr static size_t get(size_t r) {
+ if (r > R)
+ return FirstVal + index_sequence_scan_impl<R + 1, Values...>::get(r);
+ else
+ return 0;
+ }
+};
+
+template <size_t R, size_t FirstVal>
+struct index_sequence_scan_impl<R, FirstVal> {
+#if defined(__NVCC__) || defined(__NVCOMPILER) || \
+ defined(_MDSPAN_COMPILER_INTEL)
+ // NVCC warns about pointless comparison with 0 for R==0 and r being const
+ // evaluatable and also 0.
+ MDSPAN_INLINE_FUNCTION
+ constexpr static size_t get(size_t r) {
+ return static_cast<int64_t>(R) > static_cast<int64_t>(r) ? FirstVal : 0;
+ }
+#else
+ MDSPAN_INLINE_FUNCTION
+ constexpr static size_t get(size_t r) { return R > r ? FirstVal : 0; }
+#endif
+};
+template <> struct index_sequence_scan_impl<0> {
+ MDSPAN_INLINE_FUNCTION
+ constexpr static size_t get(size_t) { return 0; }
+};
+
+// ------------------------------------------------------------------
+// ------------ possibly_empty_array -------------------------------
+// ------------------------------------------------------------------
+
+// array like class which provides get function and operator [], and
+// has a specialization for the size 0 case.
+// This is needed to make the maybe_static_array be truly empty, for
+// all static values.
+
+template <class T, size_t N> struct possibly_empty_array {
+ T vals[N]{};
+ MDSPAN_INLINE_FUNCTION
+ constexpr T &operator[](size_t r) { return vals[r]; }
+ MDSPAN_INLINE_FUNCTION
+ constexpr const T &operator[](size_t r) const { return vals[r]; }
+};
+
+template <class T> struct possibly_empty_array<T, 0> {
+ MDSPAN_INLINE_FUNCTION
+ constexpr T operator[](size_t) { return T(); }
+ MDSPAN_INLINE_FUNCTION
+ constexpr const T operator[](size_t) const { return T(); }
+};
+
+// ------------------------------------------------------------------
+// ------------ maybe_static_array ----------------------------------
+// ------------------------------------------------------------------
+
+// array like class which has a mix of static and runtime values but
+// only stores the runtime values.
+// The type of the static and the runtime values can be different.
+// The position of a dynamic value is indicated through a tag value.
+template <class TDynamic, class TStatic, TStatic dyn_tag, TStatic... Values>
+struct maybe_static_array {
+
+ static_assert(std::is_convertible<TStatic, TDynamic>::value, "maybe_static_array: TStatic must be convertible to TDynamic");
+ static_assert(std::is_convertible<TDynamic, TStatic>::value, "maybe_static_array: TDynamic must be convertible to TStatic");
+
+private:
+ // Static values member
+ using static_vals_t = static_array<TStatic, Values...>;
+ constexpr static size_t m_size = sizeof...(Values);
+ constexpr static size_t m_size_dynamic =
+ _MDSPAN_FOLD_PLUS_RIGHT((Values == dyn_tag), 0);
+
+ // Dynamic values member
+ _MDSPAN_NO_UNIQUE_ADDRESS possibly_empty_array<TDynamic, m_size_dynamic>
+ m_dyn_vals;
+
+ // static mapping of indices to the position in the dynamic values array
+ using dyn_map_t = index_sequence_scan_impl<0, static_cast<size_t>(Values == dyn_tag)...>;
+public:
+
+ // two types for static and dynamic values
+ using value_type = TDynamic;
+ using static_value_type = TStatic;
+ // tag value indicating dynamic value
+ constexpr static static_value_type tag_value = dyn_tag;
+
+ constexpr maybe_static_array() = default;
+
+ // constructor for all static values
+ // TODO: add precondition check?
+ MDSPAN_TEMPLATE_REQUIRES(class... Vals,
+ /* requires */ ((m_size_dynamic == 0) &&
+ (sizeof...(Vals) > 0)))
+ MDSPAN_INLINE_FUNCTION
+ constexpr maybe_static_array(Vals...) : m_dyn_vals{} {}
+
+ // constructors from dynamic values only
+ MDSPAN_TEMPLATE_REQUIRES(class... DynVals,
+ /* requires */ (sizeof...(DynVals) ==
+ m_size_dynamic &&
+ m_size_dynamic > 0))
+ MDSPAN_INLINE_FUNCTION
+ constexpr maybe_static_array(DynVals... vals)
+ : m_dyn_vals{static_cast<TDynamic>(vals)...} {}
+
+
+ MDSPAN_TEMPLATE_REQUIRES(class T, size_t N,
+ /* requires */ (N == m_size_dynamic && N > 0))
+ MDSPAN_INLINE_FUNCTION
+ constexpr maybe_static_array(const std::array<T, N> &vals) {
+ for (size_t r = 0; r < N; r++)
+ m_dyn_vals[r] = static_cast<TDynamic>(vals[r]);
+ }
+
+ MDSPAN_TEMPLATE_REQUIRES(class T, size_t N,
+ /* requires */ (N == m_size_dynamic && N == 0))
+ MDSPAN_INLINE_FUNCTION
+ constexpr maybe_static_array(const std::array<T, N> &) : m_dyn_vals{} {}
+
+#ifdef __cpp_lib_span
+ MDSPAN_TEMPLATE_REQUIRES(class T, size_t N,
+ /* requires */ (N == m_size_dynamic && N > 0))
+ MDSPAN_INLINE_FUNCTION
+ constexpr maybe_static_array(const std::span<T, N> &vals) {
+ for (size_t r = 0; r < N; r++)
+ m_dyn_vals[r] = static_cast<TDynamic>(vals[r]);
+ }
+
+ MDSPAN_TEMPLATE_REQUIRES(class T, size_t N,
+ /* requires */ (N == m_size_dynamic && N == 0))
+ MDSPAN_INLINE_FUNCTION
+ constexpr maybe_static_array(const std::span<T, N> &) : m_dyn_vals{} {}
+#endif
+
+ // constructors from all values
+ MDSPAN_TEMPLATE_REQUIRES(class... DynVals,
+ /* requires */ (sizeof...(DynVals) !=
+ m_size_dynamic &&
+ m_size_dynamic > 0))
+ MDSPAN_INLINE_FUNCTION
+ constexpr maybe_static_array(DynVals... vals)
+ : m_dyn_vals{} {
+ static_assert((sizeof...(DynVals) == m_size), "Invalid number of values.");
+ TDynamic values[m_size]{static_cast<TDynamic>(vals)...};
+ for (size_t r = 0; r < m_size; r++) {
+ TStatic static_val = static_vals_t::get(r);
+ if (static_val == dyn_tag) {
+ m_dyn_vals[dyn_map_t::get(r)] = values[r];
+ }
+// Precondition check
+#ifdef _MDSPAN_DEBUG
+ else {
+ assert(values[r] == static_cast<TDynamic>(static_val));
+ }
+#endif
+ }
+ }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class T, size_t N,
+ /* requires */ (N != m_size_dynamic && m_size_dynamic > 0))
+ MDSPAN_INLINE_FUNCTION
+ constexpr maybe_static_array(const std::array<T, N> &vals) {
+ static_assert((N == m_size), "Invalid number of values.");
+// Precondition check
+#ifdef _MDSPAN_DEBUG
+ assert(N == m_size);
+#endif
+ for (size_t r = 0; r < m_size; r++) {
+ TStatic static_val = static_vals_t::get(r);
+ if (static_val == dyn_tag) {
+ m_dyn_vals[dyn_map_t::get(r)] = static_cast<TDynamic>(vals[r]);
+ }
+// Precondition check
+#ifdef _MDSPAN_DEBUG
+ else {
+ assert(static_cast<TDynamic>(vals[r]) ==
+ static_cast<TDynamic>(static_val));
+ }
+#endif
+ }
+ }
+
+#ifdef __cpp_lib_span
+ MDSPAN_TEMPLATE_REQUIRES(
+ class T, size_t N,
+ /* requires */ (N != m_size_dynamic && m_size_dynamic > 0))
+ MDSPAN_INLINE_FUNCTION
+ constexpr maybe_static_array(const std::span<T, N> &vals) {
+ static_assert((N == m_size) || (m_size == dynamic_extent));
+#ifdef _MDSPAN_DEBUG
+ assert(N == m_size);
+#endif
+ for (size_t r = 0; r < m_size; r++) {
+ TStatic static_val = static_vals_t::get(r);
+ if (static_val == dyn_tag) {
+ m_dyn_vals[dyn_map_t::get(r)] = static_cast<TDynamic>(vals[r]);
+ }
+#ifdef _MDSPAN_DEBUG
+ else {
+ assert(static_cast<TDynamic>(vals[r]) ==
+ static_cast<TDynamic>(static_val));
+ }
+#endif
+ }
+ }
+#endif
+
+ // access functions
+ MDSPAN_INLINE_FUNCTION
+ constexpr static TStatic static_value(size_t r) { return static_vals_t::get(r); }
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr TDynamic value(size_t r) const {
+ TStatic static_val = static_vals_t::get(r);
+ return static_val == dyn_tag ? m_dyn_vals[dyn_map_t::get(r)]
+ : static_cast<TDynamic>(static_val);
+ }
+ MDSPAN_INLINE_FUNCTION
+ constexpr TDynamic operator[](size_t r) const { return value(r); }
+
+
+ // observers
+ MDSPAN_INLINE_FUNCTION
+ constexpr static size_t size() { return m_size; }
+ MDSPAN_INLINE_FUNCTION
+ constexpr static size_t size_dynamic() { return m_size_dynamic; }
+};
+
+} // namespace detail
+} // namespace MDSPAN_IMPL_STANDARD_NAMESPACE
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+
+// ------------------------------------------------------------------
+// ------------ extents ---------------------------------------------
+// ------------------------------------------------------------------
+
+// Class to describe the extents of a multi dimensional array.
+// Used by mdspan, mdarray and layout mappings.
+// See ISO C++ standard [mdspan.extents]
+
+template <class IndexType, size_t... Extents> class extents {
+public:
+ // typedefs for integral types used
+ using index_type = IndexType;
+ using size_type = std::make_unsigned_t<index_type>;
+ using rank_type = size_t;
+
+ static_assert(std::is_integral<index_type>::value && !std::is_same<index_type, bool>::value,
+ MDSPAN_IMPL_STANDARD_NAMESPACE_STRING "::extents::index_type must be a signed or unsigned integer type");
+private:
+ constexpr static rank_type m_rank = sizeof...(Extents);
+ constexpr static rank_type m_rank_dynamic =
+ _MDSPAN_FOLD_PLUS_RIGHT((Extents == dynamic_extent), /* + ... + */ 0);
+
+ // internal storage type using maybe_static_array
+ using vals_t =
+ detail::maybe_static_array<IndexType, size_t, dynamic_extent, Extents...>;
+ _MDSPAN_NO_UNIQUE_ADDRESS vals_t m_vals;
+
+public:
+ // [mdspan.extents.obs], observers of multidimensional index space
+ MDSPAN_INLINE_FUNCTION
+ constexpr static rank_type rank() noexcept { return m_rank; }
+ MDSPAN_INLINE_FUNCTION
+ constexpr static rank_type rank_dynamic() noexcept { return m_rank_dynamic; }
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr index_type extent(rank_type r) const noexcept { return m_vals.value(r); }
+ MDSPAN_INLINE_FUNCTION
+ constexpr static size_t static_extent(rank_type r) noexcept {
+ return vals_t::static_value(r);
+ }
+
+ // [mdspan.extents.cons], constructors
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ constexpr extents() noexcept = default;
+
+ // Construction from just dynamic or all values.
+ // Precondition check is deferred to maybe_static_array constructor
+ MDSPAN_TEMPLATE_REQUIRES(
+ class... OtherIndexTypes,
+ /* requires */ (
+ _MDSPAN_FOLD_AND(_MDSPAN_TRAIT(std::is_convertible, OtherIndexTypes,
+ index_type) /* && ... */) &&
+ _MDSPAN_FOLD_AND(_MDSPAN_TRAIT(std::is_nothrow_constructible, index_type,
+ OtherIndexTypes) /* && ... */) &&
+ (sizeof...(OtherIndexTypes) == m_rank ||
+ sizeof...(OtherIndexTypes) == m_rank_dynamic)))
+ MDSPAN_INLINE_FUNCTION
+ constexpr explicit extents(OtherIndexTypes... dynvals) noexcept
+ : m_vals(static_cast<index_type>(dynvals)...) {}
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherIndexType, size_t N,
+ /* requires */
+ (
+ _MDSPAN_TRAIT(std::is_convertible, const OtherIndexType&, index_type) &&
+ _MDSPAN_TRAIT(std::is_nothrow_constructible, index_type,
+ const OtherIndexType&) &&
+ (N == m_rank || N == m_rank_dynamic)))
+ MDSPAN_INLINE_FUNCTION
+ MDSPAN_CONDITIONAL_EXPLICIT(N != m_rank_dynamic)
+ constexpr extents(const std::array<OtherIndexType, N> &exts) noexcept
+ : m_vals(std::move(exts)) {}
+
+#ifdef __cpp_lib_span
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherIndexType, size_t N,
+ /* requires */
+ (_MDSPAN_TRAIT(std::is_convertible, const OtherIndexType&, index_type) &&
+ _MDSPAN_TRAIT(std::is_nothrow_constructible, index_type, const OtherIndexType&) &&
+ (N == m_rank || N == m_rank_dynamic)))
+ MDSPAN_INLINE_FUNCTION
+ MDSPAN_CONDITIONAL_EXPLICIT(N != m_rank_dynamic)
+ constexpr extents(const std::span<OtherIndexType, N> &exts) noexcept
+ : m_vals(std::move(exts)) {}
+#endif
+
+private:
+ // Function to construct extents storage from other extents.
+ // With C++ 17 the first two variants could be collapsed using if constexpr
+ // in which case you don't need all the requires clauses.
+ // in C++ 14 mode that doesn't work due to infinite recursion
+ MDSPAN_TEMPLATE_REQUIRES(
+ size_t DynCount, size_t R, class OtherExtents, class... DynamicValues,
+ /* requires */ ((R < m_rank) && (static_extent(R) == dynamic_extent)))
+ MDSPAN_INLINE_FUNCTION
+ constexpr
+ vals_t __construct_vals_from_extents(std::integral_constant<size_t, DynCount>,
+ std::integral_constant<size_t, R>,
+ const OtherExtents &exts,
+ DynamicValues... dynamic_values) noexcept {
+ return __construct_vals_from_extents(
+ std::integral_constant<size_t, DynCount + 1>(),
+ std::integral_constant<size_t, R + 1>(), exts, dynamic_values...,
+ exts.extent(R));
+ }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ size_t DynCount, size_t R, class OtherExtents, class... DynamicValues,
+ /* requires */ ((R < m_rank) && (static_extent(R) != dynamic_extent)))
+ MDSPAN_INLINE_FUNCTION
+ constexpr
+ vals_t __construct_vals_from_extents(std::integral_constant<size_t, DynCount>,
+ std::integral_constant<size_t, R>,
+ const OtherExtents &exts,
+ DynamicValues... dynamic_values) noexcept {
+ return __construct_vals_from_extents(
+ std::integral_constant<size_t, DynCount>(),
+ std::integral_constant<size_t, R + 1>(), exts, dynamic_values...);
+ }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ size_t DynCount, size_t R, class OtherExtents, class... DynamicValues,
+ /* requires */ ((R == m_rank) && (DynCount == m_rank_dynamic)))
+ MDSPAN_INLINE_FUNCTION
+ constexpr
+ vals_t __construct_vals_from_extents(std::integral_constant<size_t, DynCount>,
+ std::integral_constant<size_t, R>,
+ const OtherExtents &,
+ DynamicValues... dynamic_values) noexcept {
+ return vals_t{static_cast<index_type>(dynamic_values)...};
+ }
+
+public:
+
+ // Converting constructor from other extents specializations
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherIndexType, size_t... OtherExtents,
+ /* requires */
+ (
+ /* multi-stage check to protect from invalid pack expansion when sizes
+ don't match? */
+ decltype(detail::__check_compatible_extents(
+ // using: sizeof...(Extents) == sizeof...(OtherExtents) as the second argument fails with MSVC+NVCC with some obscure expansion error
+ // MSVC: 19.38.33133 NVCC: 12.0
+ std::integral_constant<bool, extents<int, Extents...>::rank() == extents<int, OtherExtents...>::rank()>{},
+ std::integer_sequence<size_t, Extents...>{},
+ std::integer_sequence<size_t, OtherExtents...>{}))::value
+ )
+ )
+ MDSPAN_INLINE_FUNCTION
+ MDSPAN_CONDITIONAL_EXPLICIT((((Extents != dynamic_extent) &&
+ (OtherExtents == dynamic_extent)) ||
+ ...) ||
+ (std::numeric_limits<index_type>::max() <
+ std::numeric_limits<OtherIndexType>::max()))
+ constexpr extents(const extents<OtherIndexType, OtherExtents...> &other) noexcept
+ : m_vals(__construct_vals_from_extents(
+ std::integral_constant<size_t, 0>(),
+ std::integral_constant<size_t, 0>(), other)) {}
+
+ // Comparison operator
+ template <class OtherIndexType, size_t... OtherExtents>
+ MDSPAN_INLINE_FUNCTION friend constexpr bool
+ operator==(const extents &lhs,
+ const extents<OtherIndexType, OtherExtents...> &rhs) noexcept {
+ return
+ rank() == extents<OtherIndexType, OtherExtents...>::rank() &&
+ detail::rankwise_equal(detail::with_rank<rank()>{}, rhs, lhs, detail::extent);
+ }
+
+#if !(MDSPAN_HAS_CXX_20)
+ template <class OtherIndexType, size_t... OtherExtents>
+ MDSPAN_INLINE_FUNCTION friend constexpr bool
+ operator!=(extents const &lhs,
+ extents<OtherIndexType, OtherExtents...> const &rhs) noexcept {
+ return !(lhs == rhs);
+ }
+#endif
+};
+
+// Recursive helper classes to implement dextents alias for extents
+namespace detail {
+
+template <class IndexType, size_t Rank,
+ class Extents = ::MDSPAN_IMPL_STANDARD_NAMESPACE::extents<IndexType>>
+struct __make_dextents;
+
+template <class IndexType, size_t Rank, size_t... ExtentsPack>
+struct __make_dextents<
+ IndexType, Rank, ::MDSPAN_IMPL_STANDARD_NAMESPACE::extents<IndexType, ExtentsPack...>>
+{
+ using type = typename __make_dextents<
+ IndexType, Rank - 1,
+ ::MDSPAN_IMPL_STANDARD_NAMESPACE::extents<IndexType,
+ ::MDSPAN_IMPL_STANDARD_NAMESPACE::dynamic_extent,
+ ExtentsPack...>>::type;
+};
+
+template <class IndexType, size_t... ExtentsPack>
+struct __make_dextents<
+ IndexType, 0, ::MDSPAN_IMPL_STANDARD_NAMESPACE::extents<IndexType, ExtentsPack...>>
+{
+ using type = ::MDSPAN_IMPL_STANDARD_NAMESPACE::extents<IndexType, ExtentsPack...>;
+};
+
+} // end namespace detail
+
+// [mdspan.extents.dextents], alias template
+template <class IndexType, size_t Rank>
+using dextents = typename detail::__make_dextents<IndexType, Rank>::type;
+
+// Deduction guide for extents
+#if defined(_MDSPAN_USE_CLASS_TEMPLATE_ARGUMENT_DEDUCTION)
+template <class... IndexTypes>
+extents(IndexTypes...)
+ -> extents<size_t,
+ ((void) sizeof(IndexTypes), ::MDSPAN_IMPL_STANDARD_NAMESPACE::dynamic_extent)...>;
+#endif
+
+// Helper type traits for identifying a class as extents.
+namespace detail {
+
+template <class T> struct __is_extents : ::std::false_type {};
+
+template <class IndexType, size_t... ExtentsPack>
+struct __is_extents<::MDSPAN_IMPL_STANDARD_NAMESPACE::extents<IndexType, ExtentsPack...>>
+ : ::std::true_type {};
+
+template <class T>
+#if MDSPAN_HAS_CXX_17
+inline
+#else
+static
+#endif
+constexpr bool __is_extents_v = __is_extents<T>::value;
+
+template<class InputIndexType, class ExtentsIndexType>
+MDSPAN_INLINE_FUNCTION
+constexpr void
+check_lower_bound(InputIndexType user_index,
+ ExtentsIndexType /* current_extent */,
+ std::true_type /* is_signed */)
+{
+ (void) user_index; // prevent unused variable warning
+#ifdef _MDSPAN_DEBUG
+ assert(static_cast<ExtentsIndexType>(user_index) >= 0);
+#endif
+}
+
+template<class InputIndexType, class ExtentsIndexType>
+MDSPAN_INLINE_FUNCTION
+constexpr void
+check_lower_bound(InputIndexType /* user_index */,
+ ExtentsIndexType /* current_extent */,
+ std::false_type /* is_signed */)
+{}
+
+template<class InputIndexType, class ExtentsIndexType>
+MDSPAN_INLINE_FUNCTION
+constexpr void
+check_upper_bound(InputIndexType user_index,
+ ExtentsIndexType current_extent)
+{
+ (void) user_index; // prevent unused variable warnings
+ (void) current_extent;
+#ifdef _MDSPAN_DEBUG
+ assert(static_cast<ExtentsIndexType>(user_index) < current_extent);
+#endif
+}
+
+// Returning true to use AND fold instead of comma
+// CPP14 mode doesn't like the use of void expressions
+// with the way the _MDSPAN_FOLD_AND is set up
+template<class InputIndex, class ExtentsIndexType>
+MDSPAN_INLINE_FUNCTION
+constexpr bool
+check_one_index(InputIndex user_index,
+ ExtentsIndexType current_extent)
+{
+ check_lower_bound(user_index, current_extent,
+ std::integral_constant<bool, std::is_signed<ExtentsIndexType>::value>{});
+ check_upper_bound(user_index, current_extent);
+ return true;
+}
+
+template<size_t ... RankIndices,
+ class ExtentsIndexType, size_t ... Exts,
+ class ... Indices>
+MDSPAN_INLINE_FUNCTION
+constexpr void
+check_all_indices_helper(std::index_sequence<RankIndices...>,
+ const extents<ExtentsIndexType, Exts...>& exts,
+ Indices... indices)
+{
+ // Suppress warning about statement has no effect
+ (void) _MDSPAN_FOLD_AND(
+ (check_one_index(indices, exts.extent(RankIndices)))
+ );
+}
+
+template<class ExtentsIndexType, size_t ... Exts,
+ class ... Indices>
+MDSPAN_INLINE_FUNCTION
+constexpr void
+check_all_indices(const extents<ExtentsIndexType, Exts...>& exts,
+ Indices... indices)
+{
+ check_all_indices_helper(std::make_index_sequence<sizeof...(Indices)>(),
+ exts, indices...);
+}
+
+} // namespace detail
+} // namespace MDSPAN_IMPL_STANDARD_NAMESPACE
diff --git a/ext/mdspan/include/experimental/__p0009_bits/full_extent_t.hpp b/ext/mdspan/include/experimental/__p0009_bits/full_extent_t.hpp
new file mode 100644
index 0000000..bd4b5c6
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p0009_bits/full_extent_t.hpp
@@ -0,0 +1,26 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+#pragma once
+
+#include "macros.hpp"
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+
+struct full_extent_t { explicit full_extent_t() = default; };
+
+_MDSPAN_INLINE_VARIABLE constexpr auto full_extent = full_extent_t{ };
+
+} // namespace MDSPAN_IMPL_STANDARD_NAMESPACE
diff --git a/ext/mdspan/include/experimental/__p0009_bits/layout_left.hpp b/ext/mdspan/include/experimental/__p0009_bits/layout_left.hpp
new file mode 100644
index 0000000..ed8aae0
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p0009_bits/layout_left.hpp
@@ -0,0 +1,269 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+#pragma once
+
+#include "macros.hpp"
+#include "trait_backports.hpp"
+#include "extents.hpp"
+#include "layout_stride.hpp"
+#include "utility.hpp"
+#if MDSPAN_HAS_CXX_17
+#include "../__p2642_bits/layout_padded_fwd.hpp"
+#endif
+#include <type_traits>
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+
+//==============================================================================
+
+template <class Extents>
+class layout_left::mapping {
+ public:
+ using extents_type = Extents;
+ using index_type = typename extents_type::index_type;
+ using size_type = typename extents_type::size_type;
+ using rank_type = typename extents_type::rank_type;
+ using layout_type = layout_left;
+ private:
+
+ static_assert(detail::__is_extents_v<extents_type>,
+ MDSPAN_IMPL_STANDARD_NAMESPACE_STRING "::layout_left::mapping must be instantiated with a specialization of " MDSPAN_IMPL_STANDARD_NAMESPACE_STRING "::extents.");
+
+ template <class>
+ friend class mapping;
+
+ // i0+(i1 + E(1)*(i2 + E(2)*i3))
+ template <size_t r, size_t Rank>
+ struct __rank_count {};
+
+ template <size_t r, size_t Rank, class I, class... Indices>
+ _MDSPAN_HOST_DEVICE
+ constexpr index_type __compute_offset(
+ __rank_count<r,Rank>, const I& i, Indices... idx) const {
+ return __compute_offset(__rank_count<r+1,Rank>(), idx...) *
+ __extents.extent(r) + i;
+ }
+
+ template<class I>
+ _MDSPAN_HOST_DEVICE
+ constexpr index_type __compute_offset(
+ __rank_count<extents_type::rank()-1,extents_type::rank()>, const I& i) const {
+ return i;
+ }
+
+ _MDSPAN_HOST_DEVICE
+ constexpr index_type __compute_offset(__rank_count<0,0>) const { return 0; }
+
+ public:
+
+ //--------------------------------------------------------------------------------
+
+ MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr mapping() noexcept = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr mapping(mapping const&) noexcept = default;
+
+ _MDSPAN_HOST_DEVICE
+ constexpr mapping(extents_type const& __exts) noexcept
+ :__extents(__exts)
+ { }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherExtents,
+ /* requires */ (
+ _MDSPAN_TRAIT(std::is_constructible, extents_type, OtherExtents)
+ )
+ )
+ MDSPAN_CONDITIONAL_EXPLICIT((!std::is_convertible<OtherExtents, extents_type>::value)) // needs two () due to comma
+ MDSPAN_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14
+ mapping(mapping<OtherExtents> const& other) noexcept // NOLINT(google-explicit-constructor)
+ :__extents(other.extents())
+ {
+ /*
+ * TODO: check precondition
+ * other.required_span_size() is a representable value of type index_type
+ */
+ }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherExtents,
+ /* requires */ (
+ _MDSPAN_TRAIT(std::is_constructible, extents_type, OtherExtents) &&
+ (extents_type::rank() <= 1)
+ )
+ )
+ MDSPAN_CONDITIONAL_EXPLICIT((!std::is_convertible<OtherExtents, extents_type>::value)) // needs two () due to comma
+ MDSPAN_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14
+ mapping(layout_right::mapping<OtherExtents> const& other) noexcept // NOLINT(google-explicit-constructor)
+ :__extents(other.extents())
+ {
+ /*
+ * TODO: check precondition
+ * other.required_span_size() is a representable value of type index_type
+ */
+ }
+
+#if MDSPAN_HAS_CXX_17
+ /**
+ * Converting constructor from `layout_left_padded::mapping`.
+ *
+ * This overload participates in overload resolution only if _Mapping is a layout_left_padded mapping and
+ * extents_type is constructible from _Mapping::extents_type.
+ *
+ * \note There is currently a difference from p2642r2, where this function is specified as taking
+ * `layout_left_padded< padding_value >::mapping< Extents>`. However, this makes `padding_value` non-deducible.
+ */
+ MDSPAN_TEMPLATE_REQUIRES(
+ class _Mapping,
+ /* requires */ (
+ MDSPAN_IMPL_PROPOSED_NAMESPACE::detail::is_layout_left_padded_mapping<_Mapping>::value
+ && std::is_constructible_v<extents_type, typename _Mapping::extents_type>
+ )
+ )
+ MDSPAN_CONDITIONAL_EXPLICIT((!std::is_convertible_v<typename _Mapping::extents_type, extents_type>))
+ mapping(const _Mapping& __other) noexcept
+ : __extents(__other.extents())
+ {
+ MDSPAN_IMPL_PROPOSED_NAMESPACE::detail::
+ check_padded_layout_converting_constructor_mandates<
+ extents_type, _Mapping>(detail::with_rank<extents_type::rank()>{});
+ MDSPAN_IMPL_PROPOSED_NAMESPACE::detail::
+ check_padded_layout_converting_constructor_preconditions<
+ extents_type>(detail::with_rank<extents_type::rank()>{}, __other);
+ }
+#endif
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherExtents,
+ /* requires */ (
+ _MDSPAN_TRAIT(std::is_constructible, extents_type, OtherExtents)
+ )
+ )
+ MDSPAN_CONDITIONAL_EXPLICIT((extents_type::rank() > 0))
+ MDSPAN_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14
+ mapping(layout_stride::mapping<OtherExtents> const& other) noexcept // NOLINT(google-explicit-constructor)
+ :__extents(other.extents())
+ {
+ /*
+ * TODO: check precondition
+ * other.required_span_size() is a representable value of type index_type
+ */
+ detail::validate_strides(detail::with_rank<extents_type::rank()>{}, layout_left{}, __extents, other);
+ }
+
+ MDSPAN_INLINE_FUNCTION_DEFAULTED _MDSPAN_CONSTEXPR_14_DEFAULTED mapping& operator=(mapping const&) noexcept = default;
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr const extents_type& extents() const noexcept {
+ return __extents;
+ }
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr index_type required_span_size() const noexcept {
+ index_type value = 1;
+ for(rank_type r=0; r<extents_type::rank(); r++) value*=__extents.extent(r);
+ return value;
+ }
+
+ //--------------------------------------------------------------------------------
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class... Indices,
+ /* requires */ (
+ (sizeof...(Indices) == extents_type::rank()) &&
+ (detail::are_valid_indices<index_type, Indices...>())
+ )
+ )
+ _MDSPAN_HOST_DEVICE
+ constexpr index_type operator()(Indices... idxs) const noexcept {
+#if ! defined(NDEBUG)
+ detail::check_all_indices(this->extents(), idxs...);
+#endif // ! NDEBUG
+ return __compute_offset(__rank_count<0, extents_type::rank()>(), static_cast<index_type>(idxs)...);
+ }
+
+
+
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_unique() noexcept { return true; }
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_exhaustive() noexcept { return true; }
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_strided() noexcept { return true; }
+
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_unique() noexcept { return true; }
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_exhaustive() noexcept { return true; }
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_strided() noexcept { return true; }
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr index_type stride(rank_type i) const noexcept
+#if MDSPAN_HAS_CXX_20
+ requires ( Extents::rank() > 0 )
+#endif
+ {
+ index_type value = 1;
+ for(rank_type r=0; r<i; r++) value*=__extents.extent(r);
+ return value;
+ }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherExtents,
+ /* requires */ ( Extents::rank() == OtherExtents::rank())
+ )
+ MDSPAN_INLINE_FUNCTION
+ friend constexpr bool operator==(mapping const& lhs, mapping<OtherExtents> const& rhs) noexcept {
+ return lhs.extents() == rhs.extents();
+ }
+
+ // In C++ 20 the not equal exists if equal is found
+#if !(MDSPAN_HAS_CXX_20)
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherExtents,
+ /* requires */ ( Extents::rank() == OtherExtents::rank())
+ )
+ MDSPAN_INLINE_FUNCTION
+ friend constexpr bool operator!=(mapping const& lhs, mapping<OtherExtents> const& rhs) noexcept {
+ return lhs.extents() != rhs.extents();
+ }
+#endif
+
+ // Not really public, but currently needed to implement fully constexpr useable submdspan:
+ template<size_t N, class SizeType, size_t ... E, size_t ... Idx>
+ MDSPAN_INLINE_FUNCTION
+ constexpr index_type __get_stride(MDSPAN_IMPL_STANDARD_NAMESPACE::extents<SizeType, E...>,std::integer_sequence<size_t, Idx...>) const {
+ return _MDSPAN_FOLD_TIMES_RIGHT((Idx<N? __extents.template __extent<Idx>():1),1);
+ }
+ template<size_t N>
+ MDSPAN_INLINE_FUNCTION
+ constexpr index_type __stride() const noexcept {
+ return __get_stride<N>(__extents, std::make_index_sequence<extents_type::rank()>());
+ }
+
+private:
+ _MDSPAN_NO_UNIQUE_ADDRESS extents_type __extents{};
+
+ // [mdspan.submdspan.mapping], submdspan mapping specialization
+ template<class... SliceSpecifiers>
+ MDSPAN_INLINE_FUNCTION
+ constexpr auto submdspan_mapping_impl(
+ SliceSpecifiers... slices) const;
+
+ template<class... SliceSpecifiers>
+ MDSPAN_INLINE_FUNCTION
+ friend constexpr auto submdspan_mapping(
+ const mapping& src, SliceSpecifiers... slices) {
+ return src.submdspan_mapping_impl(slices...);
+ }
+};
+
+
+} // end namespace MDSPAN_IMPL_STANDARD_NAMESPACE
+
diff --git a/ext/mdspan/include/experimental/__p0009_bits/layout_right.hpp b/ext/mdspan/include/experimental/__p0009_bits/layout_right.hpp
new file mode 100644
index 0000000..26115e7
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p0009_bits/layout_right.hpp
@@ -0,0 +1,265 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+#pragma once
+
+#include "macros.hpp"
+#include "trait_backports.hpp"
+#include "extents.hpp"
+#include "layout_stride.hpp"
+#include "utility.hpp"
+#if MDSPAN_HAS_CXX_17
+#include "../__p2642_bits/layout_padded_fwd.hpp"
+#endif
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+
+//==============================================================================
+template <class Extents>
+class layout_right::mapping {
+ public:
+ using extents_type = Extents;
+ using index_type = typename extents_type::index_type;
+ using size_type = typename extents_type::size_type;
+ using rank_type = typename extents_type::rank_type;
+ using layout_type = layout_right;
+ private:
+
+ static_assert(detail::__is_extents_v<extents_type>,
+ MDSPAN_IMPL_STANDARD_NAMESPACE_STRING "::layout_right::mapping must be instantiated with a specialization of " MDSPAN_IMPL_STANDARD_NAMESPACE_STRING "::extents.");
+
+ template <class>
+ friend class mapping;
+
+ // i0+(i1 + E(1)*(i2 + E(2)*i3))
+ template <size_t r, size_t Rank>
+ struct __rank_count {};
+
+ template <size_t r, size_t Rank, class I, class... Indices>
+ _MDSPAN_HOST_DEVICE
+ constexpr index_type __compute_offset(
+ index_type offset, __rank_count<r,Rank>, const I& i, Indices... idx) const {
+ return __compute_offset(offset * __extents.extent(r) + i,__rank_count<r+1,Rank>(), idx...);
+ }
+
+ template<class I, class ... Indices>
+ _MDSPAN_HOST_DEVICE
+ constexpr index_type __compute_offset(
+ __rank_count<0,extents_type::rank()>, const I& i, Indices... idx) const {
+ return __compute_offset(i,__rank_count<1,extents_type::rank()>(),idx...);
+ }
+
+ _MDSPAN_HOST_DEVICE
+ constexpr index_type __compute_offset(size_t offset, __rank_count<extents_type::rank(), extents_type::rank()>) const {
+ return static_cast<index_type>(offset);
+ }
+
+ _MDSPAN_HOST_DEVICE
+ constexpr index_type __compute_offset(__rank_count<0,0>) const { return 0; }
+
+ public:
+
+ //--------------------------------------------------------------------------------
+
+ MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr mapping() noexcept = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr mapping(mapping const&) noexcept = default;
+
+ _MDSPAN_HOST_DEVICE
+ constexpr mapping(extents_type const& __exts) noexcept
+ :__extents(__exts)
+ { }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherExtents,
+ /* requires */ (
+ _MDSPAN_TRAIT(std::is_constructible, extents_type, OtherExtents)
+ )
+ )
+ MDSPAN_CONDITIONAL_EXPLICIT((!std::is_convertible<OtherExtents, extents_type>::value)) // needs two () due to comma
+ MDSPAN_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14
+ mapping(mapping<OtherExtents> const& other) noexcept // NOLINT(google-explicit-constructor)
+ :__extents(other.extents())
+ {
+ /*
+ * TODO: check precondition
+ * other.required_span_size() is a representable value of type index_type
+ */
+ }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherExtents,
+ /* requires */ (
+ _MDSPAN_TRAIT(std::is_constructible, extents_type, OtherExtents) &&
+ (extents_type::rank() <= 1)
+ )
+ )
+ MDSPAN_CONDITIONAL_EXPLICIT((!std::is_convertible<OtherExtents, extents_type>::value)) // needs two () due to comma
+ MDSPAN_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14
+ mapping(layout_left::mapping<OtherExtents> const& other) noexcept // NOLINT(google-explicit-constructor)
+ :__extents(other.extents())
+ {
+ /*
+ * TODO: check precondition
+ * other.required_span_size() is a representable value of type index_type
+ */
+ }
+
+ /**
+ * Converting constructor from `layout_right_padded::mapping`.
+ *
+ * This overload participates in overload resolution only if _Mapping is a layout_right_padded mapping and
+ * extents_type is constructible from _Mapping::extents_type.
+ *
+ * \note There is currently a difference from p2642r2, where this function is specified as taking
+ * `layout_right_padded< padding_value >::mapping< Extents>`. However, this makes `padding_value` non-deducible.
+ */
+#if MDSPAN_HAS_CXX_17
+ MDSPAN_TEMPLATE_REQUIRES(
+ class _Mapping,
+ /* requires */ (
+ MDSPAN_IMPL_PROPOSED_NAMESPACE::detail::is_layout_right_padded_mapping<_Mapping>::value
+ && std::is_constructible_v<extents_type, typename _Mapping::extents_type>))
+ MDSPAN_CONDITIONAL_EXPLICIT((!std::is_convertible_v<typename _Mapping::extents_type, extents_type>))
+ mapping(const _Mapping &__other) noexcept
+ : __extents(__other.extents())
+ {
+ MDSPAN_IMPL_PROPOSED_NAMESPACE::detail::
+ check_padded_layout_converting_constructor_mandates<
+ extents_type, _Mapping>(detail::with_rank<extents_type::rank()>{});
+ MDSPAN_IMPL_PROPOSED_NAMESPACE::detail::
+ check_padded_layout_converting_constructor_preconditions<
+ extents_type>(detail::with_rank<extents_type::rank()>{}, __other);
+ }
+#endif
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherExtents,
+ /* requires */ (
+ _MDSPAN_TRAIT(std::is_constructible, extents_type, OtherExtents)
+ )
+ )
+ MDSPAN_CONDITIONAL_EXPLICIT((extents_type::rank() > 0))
+ MDSPAN_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14
+ mapping(layout_stride::mapping<OtherExtents> const& other) noexcept // NOLINT(google-explicit-constructor)
+ :__extents(other.extents())
+ {
+ /*
+ * TODO: check precondition
+ * other.required_span_size() is a representable value of type index_type
+ */
+ detail::validate_strides(detail::with_rank<extents_type::rank()>{}, layout_right{}, __extents, other);
+ }
+
+ MDSPAN_INLINE_FUNCTION_DEFAULTED _MDSPAN_CONSTEXPR_14_DEFAULTED mapping& operator=(mapping const&) noexcept = default;
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr const extents_type& extents() const noexcept {
+ return __extents;
+ }
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr index_type required_span_size() const noexcept {
+ index_type value = 1;
+ for(rank_type r=0; r != extents_type::rank(); ++r) value*=__extents.extent(r);
+ return value;
+ }
+
+ //--------------------------------------------------------------------------------
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class ... Indices,
+ /* requires */ (
+ (sizeof...(Indices) == extents_type::rank()) &&
+ (detail::are_valid_indices<index_type, Indices...>())
+ )
+ )
+ _MDSPAN_HOST_DEVICE
+ constexpr index_type operator()(Indices... idxs) const noexcept {
+#if ! defined(NDEBUG)
+ detail::check_all_indices(this->extents(), idxs...);
+#endif // ! NDEBUG
+ return __compute_offset(__rank_count<0, extents_type::rank()>(), static_cast<index_type>(idxs)...);
+ }
+
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_unique() noexcept { return true; }
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_exhaustive() noexcept { return true; }
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_strided() noexcept { return true; }
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_unique() noexcept { return true; }
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_exhaustive() noexcept { return true; }
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_strided() noexcept { return true; }
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr index_type stride(rank_type i) const noexcept
+#if MDSPAN_HAS_CXX_20
+ requires ( Extents::rank() > 0 )
+#endif
+ {
+ index_type value = 1;
+ for(rank_type r=extents_type::rank()-1; r>i; r--) value*=__extents.extent(r);
+ return value;
+ }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherExtents,
+ /* requires */ ( Extents::rank() == OtherExtents::rank())
+ )
+ MDSPAN_INLINE_FUNCTION
+ friend constexpr bool operator==(mapping const& lhs, mapping<OtherExtents> const& rhs) noexcept {
+ return lhs.extents() == rhs.extents();
+ }
+
+ // In C++ 20 the not equal exists if equal is found
+#if !(MDSPAN_HAS_CXX_20)
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherExtents,
+ /* requires */ (Extents::rank() == OtherExtents::rank())
+ )
+ MDSPAN_INLINE_FUNCTION
+ friend constexpr bool operator!=(mapping const& lhs, mapping<OtherExtents> const& rhs) noexcept {
+ return lhs.extents() != rhs.extents();
+ }
+#endif
+
+ // Not really public, but currently needed to implement fully constexpr useable submdspan:
+ template<size_t N, class SizeType, size_t ... E, size_t ... Idx>
+ MDSPAN_INLINE_FUNCTION
+ constexpr index_type __get_stride(MDSPAN_IMPL_STANDARD_NAMESPACE::extents<SizeType, E...>,std::integer_sequence<size_t, Idx...>) const {
+ return _MDSPAN_FOLD_TIMES_RIGHT((Idx>N? __extents.template __extent<Idx>():1),1);
+ }
+ template<size_t N>
+ MDSPAN_INLINE_FUNCTION
+ constexpr index_type __stride() const noexcept {
+ return __get_stride<N>(__extents, std::make_index_sequence<extents_type::rank()>());
+ }
+
+private:
+ _MDSPAN_NO_UNIQUE_ADDRESS extents_type __extents{};
+
+ // [mdspan.submdspan.mapping], submdspan mapping specialization
+ template<class... SliceSpecifiers>
+ MDSPAN_INLINE_FUNCTION
+ constexpr auto submdspan_mapping_impl(
+ SliceSpecifiers... slices) const;
+
+ template<class... SliceSpecifiers>
+ MDSPAN_INLINE_FUNCTION
+ friend constexpr auto submdspan_mapping(
+ const mapping& src, SliceSpecifiers... slices) {
+ return src.submdspan_mapping_impl(slices...);
+ }
+};
+
+} // end namespace MDSPAN_IMPL_STANDARD_NAMESPACE
+
diff --git a/ext/mdspan/include/experimental/__p0009_bits/layout_stride.hpp b/ext/mdspan/include/experimental/__p0009_bits/layout_stride.hpp
new file mode 100644
index 0000000..47ef268
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p0009_bits/layout_stride.hpp
@@ -0,0 +1,667 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+#pragma once
+
+#include "macros.hpp"
+#include "extents.hpp"
+#include "trait_backports.hpp"
+#include "compressed_pair.hpp"
+#include "utility.hpp"
+
+#if !defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+# include "no_unique_address.hpp"
+#endif
+
+#include <array>
+#include <type_traits>
+#include <utility>
+
+#ifdef __cpp_lib_span
+#include <span>
+#endif
+#if defined(_MDSPAN_USE_CONCEPTS) && MDSPAN_HAS_CXX_20 && defined(__cpp_lib_concepts)
+# include <concepts>
+#endif
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+
+struct layout_left {
+ template<class Extents>
+ class mapping;
+};
+struct layout_right {
+ template<class Extents>
+ class mapping;
+};
+
+namespace detail {
+ template<class Layout, class Mapping>
+ constexpr bool __is_mapping_of =
+ std::is_same<typename Layout::template mapping<typename Mapping::extents_type>, Mapping>::value;
+
+#if defined(_MDSPAN_USE_CONCEPTS) && MDSPAN_HAS_CXX_20
+# if !defined(__cpp_lib_concepts)
+ namespace internal {
+ namespace detail {
+ template <typename _Tp, typename _Up>
+ concept __same_as = std::is_same_v<_Tp, _Up>;
+ } // namespace detail
+ template <class T, class U>
+ concept __same_as = detail::__same_as<T, U> && detail::__same_as<U, T>;
+ } // namespace internal
+# endif
+
+ template<class M>
+ concept __layout_mapping_alike = requires {
+ requires __is_extents<typename M::extents_type>::value;
+#if defined(__cpp_lib_concepts)
+ { M::is_always_strided() } -> std::same_as<bool>;
+ { M::is_always_exhaustive() } -> std::same_as<bool>;
+ { M::is_always_unique() } -> std::same_as<bool>;
+#else
+ { M::is_always_strided() } -> internal::__same_as<bool>;
+ { M::is_always_exhaustive() } -> internal::__same_as<bool>;
+ { M::is_always_unique() } -> internal::__same_as<bool>;
+#endif
+ std::bool_constant<M::is_always_strided()>::value;
+ std::bool_constant<M::is_always_exhaustive()>::value;
+ std::bool_constant<M::is_always_unique()>::value;
+ };
+#endif
+
+} // namespace detail
+
+struct layout_stride {
+ template <class Extents>
+ class mapping
+#if !defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+ : private detail::__no_unique_address_emulation<
+ detail::__compressed_pair<
+ Extents,
+ detail::possibly_empty_array<typename Extents::index_type, Extents::rank()>
+ >
+ >
+#endif
+ {
+ public:
+ using extents_type = Extents;
+ using index_type = typename extents_type::index_type;
+ using size_type = typename extents_type::size_type;
+ using rank_type = typename extents_type::rank_type;
+ using layout_type = layout_stride;
+
+ // This could be a `requires`, but I think it's better and clearer as a `static_assert`.
+ static_assert(detail::__is_extents_v<Extents>,
+ MDSPAN_IMPL_STANDARD_NAMESPACE_STRING "::layout_stride::mapping must be instantiated with a specialization of " MDSPAN_IMPL_STANDARD_NAMESPACE_STRING "::extents.");
+
+
+ private:
+
+ //----------------------------------------------------------------------------
+
+ using __strides_storage_t = detail::possibly_empty_array<index_type, extents_type::rank()>;
+ using __member_pair_t = detail::__compressed_pair<extents_type, __strides_storage_t>;
+
+#if defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+ _MDSPAN_NO_UNIQUE_ADDRESS __member_pair_t __members;
+#else
+ using __base_t = detail::__no_unique_address_emulation<__member_pair_t>;
+#endif
+
+ MDSPAN_FORCE_INLINE_FUNCTION constexpr __strides_storage_t const&
+ __strides_storage() const noexcept {
+#if defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+ return __members.__second();
+#else
+ return this->__base_t::__ref().__second();
+#endif
+ }
+ MDSPAN_FORCE_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 __strides_storage_t&
+ __strides_storage() noexcept {
+#if defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+ return __members.__second();
+#else
+ return this->__base_t::__ref().__second();
+#endif
+ }
+
+ template<class SizeType, size_t ... Ep, size_t ... Idx>
+ _MDSPAN_HOST_DEVICE
+ constexpr index_type __get_size(::MDSPAN_IMPL_STANDARD_NAMESPACE::extents<SizeType, Ep...>,std::integer_sequence<size_t, Idx...>) const {
+ return _MDSPAN_FOLD_TIMES_RIGHT( static_cast<index_type>(extents().extent(Idx)), 1 );
+ }
+
+ //----------------------------------------------------------------------------
+
+ template <class>
+ friend class mapping;
+
+ //----------------------------------------------------------------------------
+
+ // Workaround for non-deducibility of the index sequence template parameter if it's given at the top level
+ template <class>
+ struct __deduction_workaround;
+
+ template <size_t... Idxs>
+ struct __deduction_workaround<std::index_sequence<Idxs...>>
+ {
+ template <class OtherExtents>
+ MDSPAN_INLINE_FUNCTION
+ static constexpr bool _eq_impl(mapping const& self, mapping<OtherExtents> const& other) noexcept {
+ using common_t = std::common_type_t<index_type, typename OtherExtents::index_type>;
+ return _MDSPAN_FOLD_AND((static_cast<common_t>(self.stride(Idxs)) == static_cast<common_t>(other.stride(Idxs))) /* && ... */)
+ && _MDSPAN_FOLD_AND((static_cast<common_t>(self.extents().extent(Idxs)) == static_cast<common_t>(other.extents().extent(Idxs))) /* || ... */);
+ }
+ template <class OtherExtents>
+ MDSPAN_INLINE_FUNCTION
+ static constexpr bool _not_eq_impl(mapping const& self, mapping<OtherExtents> const& other) noexcept {
+ using common_t = std::common_type_t<index_type, typename OtherExtents::index_type>;
+ return _MDSPAN_FOLD_OR((static_cast<common_t>(self.stride(Idxs)) != static_cast<common_t>(other.stride(Idxs))) /* || ... */)
+ || _MDSPAN_FOLD_OR((static_cast<common_t>(self.extents().extent(Idxs)) != static_cast<common_t>(other.extents().extent(Idxs))) /* || ... */);
+ }
+
+ template <class... Integral>
+ MDSPAN_FORCE_INLINE_FUNCTION
+ static constexpr size_t _call_op_impl(mapping const& self, Integral... idxs) noexcept {
+ return _MDSPAN_FOLD_PLUS_RIGHT((idxs * self.stride(Idxs)), /* + ... + */ 0);
+ }
+
+ MDSPAN_INLINE_FUNCTION
+ static constexpr size_t _req_span_size_impl(mapping const& self) noexcept {
+ // assumes no negative strides; not sure if I'm allowed to assume that or not
+ return __impl::_call_op_impl(self, (self.extents().template __extent<Idxs>() - 1)...) + 1;
+ }
+
+ template<class OtherMapping>
+ MDSPAN_INLINE_FUNCTION
+ static constexpr const __strides_storage_t fill_strides(const OtherMapping& map) {
+ return __strides_storage_t{static_cast<index_type>(map.stride(Idxs))...};
+ }
+
+ MDSPAN_INLINE_FUNCTION
+ static constexpr const __strides_storage_t& fill_strides(const __strides_storage_t& s) {
+ return s;
+ }
+
+ template<class IntegralType>
+ static constexpr const __strides_storage_t fill_strides(const std::array<IntegralType,extents_type::rank()>& s) {
+ return __strides_storage_t{static_cast<index_type>(s[Idxs])...};
+ }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class IntegralType,
+ (std::is_convertible<IntegralType, typename extents_type::index_type>::value)
+ )
+ MDSPAN_INLINE_FUNCTION
+ // Need to avoid zero length c-array
+ static constexpr const __strides_storage_t fill_strides(mdspan_non_standard_tag, const IntegralType (&s)[extents_type::rank()>0?extents_type::rank():1]) {
+ return __strides_storage_t{static_cast<index_type>(s[Idxs])...};
+ }
+
+#ifdef __cpp_lib_span
+ template<class IntegralType>
+ static constexpr const __strides_storage_t fill_strides(const std::span<IntegralType,extents_type::rank()>& s) {
+ return __strides_storage_t{static_cast<index_type>(s[Idxs])...};
+ }
+#endif
+
+ MDSPAN_INLINE_FUNCTION
+ static constexpr std::array<index_type, extents_type::rank()> return_strides(const __strides_storage_t& s) {
+ return std::array<index_type, extents_type::rank()>{s[Idxs]...};
+ }
+
+ template<size_t K>
+ MDSPAN_INLINE_FUNCTION
+ static constexpr size_t __return_zero() { return 0; }
+
+ template<class Mapping>
+ MDSPAN_INLINE_FUNCTION
+ static constexpr typename Mapping::index_type
+ __OFFSET(const Mapping& m) { return m(__return_zero<Idxs>()...); }
+ };
+
+ // Can't use defaulted parameter in the __deduction_workaround template because of a bug in MSVC warning C4348.
+ using __impl = __deduction_workaround<std::make_index_sequence<Extents::rank()>>;
+
+ MDSPAN_FUNCTION
+ static constexpr __strides_storage_t strides_storage(detail::with_rank<0>) {
+ return {};
+ }
+
+ template <std::size_t N>
+ MDSPAN_FUNCTION
+ static constexpr __strides_storage_t strides_storage(detail::with_rank<N>) {
+ __strides_storage_t s{};
+
+ extents_type e;
+ index_type stride = 1;
+ for(int r = static_cast<int>(extents_type::rank() - 1); r >= 0; r--) {
+ s[r] = stride;
+ stride *= e.extent(r);
+ }
+
+ return s;
+ }
+
+ //----------------------------------------------------------------------------
+
+#if defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+ MDSPAN_INLINE_FUNCTION constexpr explicit
+ mapping(__member_pair_t&& __m) : __members(::std::move(__m)) {}
+#else
+ MDSPAN_INLINE_FUNCTION constexpr explicit
+ mapping(__base_t&& __b) : __base_t(::std::move(__b)) {}
+#endif
+
+ public:
+
+ //--------------------------------------------------------------------------------
+
+ MDSPAN_INLINE_FUNCTION constexpr mapping() noexcept
+#if defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+ : __members{
+#else
+ : __base_t(__base_t{__member_pair_t(
+#endif
+ extents_type(),
+ __strides_storage_t(strides_storage(detail::with_rank<extents_type::rank()>{}))
+#if defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+ }
+#else
+ )})
+#endif
+ {}
+
+ MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr mapping(mapping const&) noexcept = default;
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class IntegralTypes,
+ /* requires */ (
+ // MSVC 19.32 does not like using index_type here, requires the typename Extents::index_type
+ // error C2641: cannot deduce template arguments for 'MDSPAN_IMPL_STANDARD_NAMESPACE::layout_stride::mapping'
+ _MDSPAN_TRAIT(std::is_convertible, const std::remove_const_t<IntegralTypes>&, typename Extents::index_type) &&
+ _MDSPAN_TRAIT(std::is_nothrow_constructible, typename Extents::index_type, const std::remove_const_t<IntegralTypes>&)
+ )
+ )
+ constexpr
+ mapping(
+ extents_type const& e,
+ std::array<IntegralTypes, extents_type::rank()> const& s
+ ) noexcept
+#if defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+ : __members{
+#else
+ : __base_t(__base_t{__member_pair_t(
+#endif
+ e, __strides_storage_t(__impl::fill_strides(s))
+#if defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+ }
+#else
+ )})
+#endif
+ {
+ /*
+ * TODO: check preconditions
+ * - s[i] > 0 is true for all i in the range [0, rank_ ).
+ * - REQUIRED-SPAN-SIZE(e, s) is a representable value of type index_type ([basic.fundamental]).
+ * - If rank_ is greater than 0, then there exists a permutation P of the integers in the
+ * range [0, rank_), such that s[ pi ] >= s[ pi − 1 ] * e.extent( pi − 1 ) is true for
+ * all i in the range [1, rank_ ), where pi is the ith element of P.
+ */
+ }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class IntegralTypes,
+ /* requires */ (
+ // MSVC 19.32 does not like using index_type here, requires the typename Extents::index_type
+ // error C2641: cannot deduce template arguments for 'MDSPAN_IMPL_STANDARD_NAMESPACE::layout_stride::mapping'
+ _MDSPAN_TRAIT(std::is_convertible, const std::remove_const_t<IntegralTypes>&, typename Extents::index_type) &&
+ _MDSPAN_TRAIT(std::is_nothrow_constructible, typename Extents::index_type, const std::remove_const_t<IntegralTypes>&)
+ )
+ )
+ MDSPAN_INLINE_FUNCTION
+ constexpr
+ mapping(
+ mdspan_non_standard_tag,
+ extents_type const& e,
+ // Need to avoid zero-length c-array
+ const IntegralTypes (&s)[extents_type::rank()>0?extents_type::rank():1]
+ ) noexcept
+#if defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+ : __members{
+#else
+ : __base_t(__base_t{__member_pair_t(
+#endif
+ e, __strides_storage_t(__impl::fill_strides(mdspan_non_standard, s))
+#if defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+ }
+#else
+ )})
+#endif
+ {
+ /*
+ * TODO: check preconditions
+ * - s[i] > 0 is true for all i in the range [0, rank_ ).
+ * - REQUIRED-SPAN-SIZE(e, s) is a representable value of type index_type ([basic.fundamental]).
+ * - If rank_ is greater than 0, then there exists a permutation P of the integers in the
+ * range [0, rank_), such that s[ pi ] >= s[ pi − 1 ] * e.extent( pi − 1 ) is true for
+ * all i in the range [1, rank_ ), where pi is the ith element of P.
+ */
+ }
+
+#ifdef __cpp_lib_span
+ MDSPAN_TEMPLATE_REQUIRES(
+ class IntegralTypes,
+ /* requires */ (
+ // MSVC 19.32 does not like using index_type here, requires the typename Extents::index_type
+ // error C2641: cannot deduce template arguments for 'MDSPAN_IMPL_STANDARD_NAMESPACE::layout_stride::mapping'
+ _MDSPAN_TRAIT(std::is_convertible, const std::remove_const_t<IntegralTypes>&, typename Extents::index_type) &&
+ _MDSPAN_TRAIT(std::is_nothrow_constructible, typename Extents::index_type, const std::remove_const_t<IntegralTypes>&)
+ )
+ )
+ constexpr
+ mapping(
+ extents_type const& e,
+ std::span<IntegralTypes, extents_type::rank()> const& s
+ ) noexcept
+#if defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+ : __members{
+#else
+ : __base_t(__base_t{__member_pair_t(
+#endif
+ e, __strides_storage_t(__impl::fill_strides(s))
+#if defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+ }
+#else
+ )})
+#endif
+ {
+ /*
+ * TODO: check preconditions
+ * - s[i] > 0 is true for all i in the range [0, rank_ ).
+ * - REQUIRED-SPAN-SIZE(e, s) is a representable value of type index_type ([basic.fundamental]).
+ * - If rank_ is greater than 0, then there exists a permutation P of the integers in the
+ * range [0, rank_), such that s[ pi ] >= s[ pi − 1 ] * e.extent( pi − 1 ) is true for
+ * all i in the range [1, rank_ ), where pi is the ith element of P.
+ */
+ }
+#endif // __cpp_lib_span
+
+#if !(defined(_MDSPAN_USE_CONCEPTS) && MDSPAN_HAS_CXX_20)
+ MDSPAN_TEMPLATE_REQUIRES(
+ class StridedLayoutMapping,
+ /* requires */ (
+ _MDSPAN_TRAIT(std::is_constructible, extents_type, typename StridedLayoutMapping::extents_type) &&
+ detail::__is_mapping_of<typename StridedLayoutMapping::layout_type, StridedLayoutMapping> &&
+ StridedLayoutMapping::is_always_unique() &&
+ StridedLayoutMapping::is_always_strided()
+ )
+ )
+#else
+ template<class StridedLayoutMapping>
+ requires(
+ detail::__layout_mapping_alike<StridedLayoutMapping> &&
+ _MDSPAN_TRAIT(std::is_constructible, extents_type, typename StridedLayoutMapping::extents_type) &&
+ StridedLayoutMapping::is_always_unique() &&
+ StridedLayoutMapping::is_always_strided()
+ )
+#endif
+ MDSPAN_CONDITIONAL_EXPLICIT(
+ !(std::is_convertible<typename StridedLayoutMapping::extents_type, extents_type>::value &&
+ (detail::__is_mapping_of<layout_left, StridedLayoutMapping> ||
+ detail::__is_mapping_of<layout_right, StridedLayoutMapping> ||
+ detail::__is_mapping_of<layout_stride, StridedLayoutMapping>))
+ ) // needs two () due to comma
+ MDSPAN_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14
+ mapping(StridedLayoutMapping const& other) noexcept // NOLINT(google-explicit-constructor)
+#if defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+ : __members{
+#else
+ : __base_t(__base_t{__member_pair_t(
+#endif
+ other.extents(), __strides_storage_t(__impl::fill_strides(other))
+#if defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+ }
+#else
+ )})
+#endif
+ {
+ /*
+ * TODO: check preconditions
+ * - other.stride(i) > 0 is true for all i in the range [0, rank_ ).
+ * - other.required_span_size() is a representable value of type index_type ([basic.fundamental]).
+ * - OFFSET(other) == 0
+ */
+ }
+
+ //--------------------------------------------------------------------------------
+
+ MDSPAN_INLINE_FUNCTION_DEFAULTED _MDSPAN_CONSTEXPR_14_DEFAULTED
+ mapping& operator=(mapping const&) noexcept = default;
+
+ MDSPAN_INLINE_FUNCTION constexpr const extents_type& extents() const noexcept {
+#if defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS)
+ return __members.__first();
+#else
+ return this->__base_t::__ref().__first();
+#endif
+ };
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr std::array< index_type, extents_type::rank() > strides() const noexcept {
+ return __impl::return_strides(__strides_storage());
+ }
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr index_type required_span_size() const noexcept {
+ index_type span_size = 1;
+ // using int here to avoid warning about pointless comparison to 0
+ for(int r = 0; r < static_cast<int>(extents_type::rank()); r++) {
+ // Return early if any of the extents are zero
+ if(extents().extent(r)==0) return 0;
+ span_size += ( static_cast<index_type>(extents().extent(r) - 1 ) * __strides_storage()[r]);
+ }
+ return span_size;
+ }
+
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class... Indices,
+ /* requires */ (
+ sizeof...(Indices) == Extents::rank() &&
+ (detail::are_valid_indices<index_type, Indices...>())
+ )
+ )
+ MDSPAN_FORCE_INLINE_FUNCTION
+ constexpr index_type operator()(Indices... idxs) const noexcept {
+#if ! defined(NDEBUG)
+ detail::check_all_indices(this->extents(), idxs...);
+#endif // ! NDEBUG
+ return static_cast<index_type>(__impl::_call_op_impl(*this, static_cast<index_type>(idxs)...));
+ }
+
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_unique() noexcept { return true; }
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_exhaustive() noexcept {
+ return false;
+ }
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_strided() noexcept { return true; }
+
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_unique() noexcept { return true; }
+
+ private:
+ MDSPAN_INLINE_FUNCTION
+ constexpr bool exhaustive_for_nonzero_span_size() const
+ {
+ return required_span_size() == __get_size(extents(), std::make_index_sequence<extents_type::rank()>());
+ }
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr bool is_exhaustive_impl(detail::with_rank<0>) const
+ {
+ return true;
+ }
+ MDSPAN_INLINE_FUNCTION
+ constexpr bool is_exhaustive_impl(detail::with_rank<1>) const
+ {
+ if (required_span_size() != static_cast<index_type>(0)) {
+ return exhaustive_for_nonzero_span_size();
+ }
+ return stride(0) == 1;
+ }
+ template <std::size_t N>
+ MDSPAN_INLINE_FUNCTION
+ constexpr bool is_exhaustive_impl(detail::with_rank<N>) const
+ {
+ if (required_span_size() != static_cast<index_type>(0)) {
+ return exhaustive_for_nonzero_span_size();
+ }
+
+ rank_type r_largest = 0;
+ for (rank_type r = 1; r < extents_type::rank(); r++) {
+ if (stride(r) > stride(r_largest)) {
+ r_largest = r;
+ }
+ }
+ for (rank_type r = 0; r < extents_type::rank(); r++) {
+ if (extents().extent(r) == 0 && r != r_largest) {
+ return false;
+ }
+ }
+ return true;
+ }
+
+ public:
+ MDSPAN_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 bool is_exhaustive() const noexcept {
+ return is_exhaustive_impl(detail::with_rank<extents_type::rank()>{});
+ }
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_strided() noexcept { return true; }
+
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr index_type stride(rank_type r) const noexcept {
+ return __strides_storage()[r];
+ }
+
+#if !(defined(_MDSPAN_USE_CONCEPTS) && MDSPAN_HAS_CXX_20)
+ MDSPAN_TEMPLATE_REQUIRES(
+ class StridedLayoutMapping,
+ /* requires */ (
+ detail::__is_mapping_of<typename StridedLayoutMapping::layout_type, StridedLayoutMapping> &&
+ (extents_type::rank() == StridedLayoutMapping::extents_type::rank()) &&
+ StridedLayoutMapping::is_always_strided()
+ )
+ )
+#else
+ template<class StridedLayoutMapping>
+ requires(
+ detail::__layout_mapping_alike<StridedLayoutMapping> &&
+ (extents_type::rank() == StridedLayoutMapping::extents_type::rank()) &&
+ StridedLayoutMapping::is_always_strided()
+ )
+#endif
+ MDSPAN_INLINE_FUNCTION
+ friend constexpr bool operator==(const mapping& x, const StridedLayoutMapping& y) noexcept {
+ return (x.extents() == y.extents()) &&
+ (__impl::__OFFSET(y) == static_cast<typename StridedLayoutMapping::index_type>(0)) &&
+ detail::rankwise_equal(detail::with_rank<extents_type::rank()>{}, x, y, detail::stride);
+ }
+
+ // This one is not technically part of the proposal. Just here to make implementation a bit more optimal hopefully
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherExtents,
+ /* requires */ (
+ (extents_type::rank() == OtherExtents::rank())
+ )
+ )
+ MDSPAN_INLINE_FUNCTION
+ friend constexpr bool operator==(mapping const& lhs, mapping<OtherExtents> const& rhs) noexcept {
+ return __impl::_eq_impl(lhs, rhs);
+ }
+
+#if !MDSPAN_HAS_CXX_20
+ MDSPAN_TEMPLATE_REQUIRES(
+ class StridedLayoutMapping,
+ /* requires */ (
+ detail::__is_mapping_of<typename StridedLayoutMapping::layout_type, StridedLayoutMapping> &&
+ (extents_type::rank() == StridedLayoutMapping::extents_type::rank()) &&
+ StridedLayoutMapping::is_always_strided()
+ )
+ )
+ MDSPAN_INLINE_FUNCTION
+ friend constexpr bool operator!=(const mapping& x, const StridedLayoutMapping& y) noexcept {
+ return !(x == y);
+ }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherExtents,
+ /* requires */ (
+ (extents_type::rank() == OtherExtents::rank())
+ )
+ )
+ MDSPAN_INLINE_FUNCTION
+ friend constexpr bool operator!=(mapping const& lhs, mapping<OtherExtents> const& rhs) noexcept {
+ return __impl::_not_eq_impl(lhs, rhs);
+ }
+#endif
+
+ // [mdspan.submdspan.mapping], submdspan mapping specialization
+ template<class... SliceSpecifiers>
+ MDSPAN_INLINE_FUNCTION
+ constexpr auto submdspan_mapping_impl(
+ SliceSpecifiers... slices) const;
+
+ template<class... SliceSpecifiers>
+ MDSPAN_INLINE_FUNCTION
+ friend constexpr auto submdspan_mapping(
+ const mapping& src, SliceSpecifiers... slices) {
+ return src.submdspan_mapping_impl(slices...);
+ }
+ };
+};
+
+namespace detail {
+
+template <class Layout, class Extents, class Mapping>
+MDSPAN_INLINE_FUNCTION
+constexpr void validate_strides(with_rank<0>, Layout, const Extents&, const Mapping&)
+{}
+
+template <std::size_t N, class Layout, class Extents, class Mapping>
+MDSPAN_INLINE_FUNCTION
+constexpr void validate_strides(with_rank<N>, Layout, const Extents& ext, const Mapping& other)
+{
+ static_assert(std::is_same<typename Mapping::layout_type, layout_stride>::value &&
+ (std::is_same<Layout, layout_left>::value ||
+ std::is_same<Layout, layout_right>::value)
+ , "This function is only intended to validate construction of "
+ "a layout_left or layout_right mapping from a layout_stride mapping.");
+
+ constexpr auto is_left = std::is_same<Layout, layout_left>::value;
+
+ typename Extents::index_type expected_stride = 1;
+
+ for (std::size_t r = 0; r < N; r++) {
+ const std::size_t s = is_left ? r : N - 1 - r;
+
+ MDSPAN_IMPL_PRECONDITION(common_integral_compare(expected_stride, other.stride(s))
+ && "invalid strides for layout_{left,right}");
+
+ expected_stride *= ext.extent(s);
+ }
+}
+
+} // namespace detail
+} // end namespace MDSPAN_IMPL_STANDARD_NAMESPACE
diff --git a/ext/mdspan/include/experimental/__p0009_bits/macros.hpp b/ext/mdspan/include/experimental/__p0009_bits/macros.hpp
new file mode 100644
index 0000000..b60c426
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p0009_bits/macros.hpp
@@ -0,0 +1,699 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+
+#pragma once
+
+#include "config.hpp"
+
+#include <cstdio>
+#include <cstdlib>
+#include <type_traits> // std::is_void
+#if defined(_MDSPAN_HAS_CUDA) || defined(_MDSPAN_HAS_HIP) || defined(_MDSPAN_HAS_SYCL)
+#include "assert.h"
+#endif
+
+#ifndef _MDSPAN_HOST_DEVICE
+# if defined(_MDSPAN_HAS_CUDA) || defined(_MDSPAN_HAS_HIP)
+# define _MDSPAN_HOST_DEVICE __host__ __device__
+# else
+# define _MDSPAN_HOST_DEVICE
+# endif
+#endif
+
+#ifndef MDSPAN_FORCE_INLINE_FUNCTION
+# ifdef _MDSPAN_COMPILER_MSVC // Microsoft compilers
+# define MDSPAN_FORCE_INLINE_FUNCTION __forceinline _MDSPAN_HOST_DEVICE
+# else
+# define MDSPAN_FORCE_INLINE_FUNCTION __attribute__((always_inline)) _MDSPAN_HOST_DEVICE
+# endif
+#endif
+
+#ifndef MDSPAN_INLINE_FUNCTION
+# define MDSPAN_INLINE_FUNCTION inline _MDSPAN_HOST_DEVICE
+#endif
+
+#ifndef MDSPAN_FUNCTION
+# define MDSPAN_FUNCTION _MDSPAN_HOST_DEVICE
+#endif
+
+#ifdef _MDSPAN_HAS_HIP
+# define MDSPAN_DEDUCTION_GUIDE _MDSPAN_HOST_DEVICE
+#else
+# define MDSPAN_DEDUCTION_GUIDE
+#endif
+
+// In CUDA defaulted functions do not need host device markup
+#ifndef MDSPAN_INLINE_FUNCTION_DEFAULTED
+# define MDSPAN_INLINE_FUNCTION_DEFAULTED
+#endif
+
+//==============================================================================
+// <editor-fold desc="Preprocessor helpers"> {{{1
+
+#define MDSPAN_PP_COUNT(...) \
+ _MDSPAN_PP_INTERNAL_EXPAND_ARGS_PRIVATE( \
+ _MDSPAN_PP_INTERNAL_ARGS_AUGMENTER(__VA_ARGS__) \
+ )
+
+#define _MDSPAN_PP_INTERNAL_ARGS_AUGMENTER(...) unused, __VA_ARGS__
+#define _MDSPAN_PP_INTERNAL_EXPAND(x) x
+#define _MDSPAN_PP_INTERNAL_EXPAND_ARGS_PRIVATE(...) \
+ _MDSPAN_PP_INTERNAL_EXPAND( \
+ _MDSPAN_PP_INTERNAL_COUNT_PRIVATE( \
+ __VA_ARGS__, 69, 68, 67, 66, 65, 64, 63, 62, 61, \
+ 60, 59, 58, 57, 56, 55, 54, 53, 52, 51, 50, 49, \
+ 48, 47, 46, 45, 44, 43, 42, 41, 40, 39, 38, 37, \
+ 36, 35, 34, 33, 32, 31, 30, 29, 28, 27, 26, 25, \
+ 24, 23, 22, 21, 20, 19, 18, 17, 16, 15, 14, 13, \
+ 12, 11, 10, 9, 8, 7, 6, 5, 4, 3, 2, 1, 0 \
+ ) \
+ )
+# define _MDSPAN_PP_INTERNAL_COUNT_PRIVATE( \
+ _1_, _2_, _3_, _4_, _5_, _6_, _7_, _8_, _9_, \
+ _10, _11, _12, _13, _14, _15, _16, _17, _18, _19, \
+ _20, _21, _22, _23, _24, _25, _26, _27, _28, _29, \
+ _30, _31, _32, _33, _34, _35, _36, _37, _38, _39, \
+ _40, _41, _42, _43, _44, _45, _46, _47, _48, _49, \
+ _50, _51, _52, _53, _54, _55, _56, _57, _58, _59, \
+ _60, _61, _62, _63, _64, _65, _66, _67, _68, _69, \
+ _70, count, ...) count \
+ /**/
+
+#define MDSPAN_PP_STRINGIFY_IMPL(x) #x
+#define MDSPAN_PP_STRINGIFY(x) MDSPAN_PP_STRINGIFY_IMPL(x)
+
+#define MDSPAN_PP_CAT_IMPL(x, y) x ## y
+#define MDSPAN_PP_CAT(x, y) MDSPAN_PP_CAT_IMPL(x, y)
+
+#define MDSPAN_PP_EVAL(X, ...) X(__VA_ARGS__)
+
+#define MDSPAN_PP_REMOVE_PARENS_IMPL(...) __VA_ARGS__
+#define MDSPAN_PP_REMOVE_PARENS(...) MDSPAN_PP_REMOVE_PARENS_IMPL __VA_ARGS__
+
+#define MDSPAN_IMPL_STANDARD_NAMESPACE_STRING MDSPAN_PP_STRINGIFY(MDSPAN_IMPL_STANDARD_NAMESPACE)
+#define MDSPAN_IMPL_PROPOSED_NAMESPACE_STRING MDSPAN_PP_STRINGIFY(MDSPAN_IMPL_STANDARD_NAMESPACE) "::" MDSPAN_PP_STRINGIFY(MDSPAN_IMPL_PROPOSED_NAMESPACE)
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+namespace detail {
+
+#if defined(_MDSPAN_HAS_CUDA) || defined(_MDSPAN_HAS_HIP)
+MDSPAN_FUNCTION inline void default_precondition_violation_handler(const char* cond, const char* file, unsigned line)
+{
+ printf("%s:%u: precondition failure: `%s`\n", file, line, cond);
+ assert(0);
+}
+#elif defined(_MDSPAN_HAS_SYCL)
+MDSPAN_FUNCTION inline void default_precondition_violation_handler(const char* cond, const char* file, unsigned line)
+{
+ sycl::ext::oneapi::experimental::printf("%s:%u: precondition failure: `%s`\n", file, line, cond);
+ assert(0);
+}
+#else
+MDSPAN_FUNCTION inline void default_precondition_violation_handler(const char* cond, const char* file, unsigned line)
+{
+ std::fprintf(stderr, "%s:%u: precondition failure: `%s`\n", file, line, cond);
+ std::abort();
+}
+#endif
+
+} // namespace detail
+} // namespace MDSPAN_IMPL_STANDARD_NAMESPACE
+
+#ifndef MDSPAN_IMPL_PRECONDITION_VIOLATION_HANDLER
+#define MDSPAN_IMPL_PRECONDITION_VIOLATION_HANDLER(cond, file, line) \
+ MDSPAN_IMPL_STANDARD_NAMESPACE::detail::default_precondition_violation_handler(cond, file, line)
+#endif
+
+#ifndef MDSPAN_IMPL_CHECK_PRECONDITION
+ #ifndef NDEBUG
+ #define MDSPAN_IMPL_CHECK_PRECONDITION 0
+ #else
+ #define MDSPAN_IMPL_CHECK_PRECONDITION 1
+ #endif
+#endif
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+namespace detail {
+
+template <bool check = MDSPAN_IMPL_CHECK_PRECONDITION>
+MDSPAN_FUNCTION constexpr void precondition(const char* cond, const char* file, unsigned line)
+{
+ if (!check) { return; }
+ // in case the macro doesn't use the arguments for custom macros
+ (void) cond;
+ (void) file;
+ (void) line;
+ MDSPAN_IMPL_PRECONDITION_VIOLATION_HANDLER(cond, file, line);
+}
+
+} // namespace detail
+} // namespace MDSPAN_IMPL_STANDARD_NAMESPACE
+
+#define MDSPAN_IMPL_PRECONDITION(...) \
+ do { \
+ if (!(__VA_ARGS__)) { \
+ MDSPAN_IMPL_STANDARD_NAMESPACE::detail::precondition(#__VA_ARGS__, __FILE__, __LINE__); \
+ } \
+ } while (0)
+
+// </editor-fold> end Preprocessor helpers }}}1
+//==============================================================================
+
+//==============================================================================
+// <editor-fold desc="Concept emulation"> {{{1
+
+// These compatibility macros don't help with partial ordering, but they should do the trick
+// for what we need to do with concepts in mdspan
+#ifdef _MDSPAN_USE_CONCEPTS
+# define MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) > requires REQ
+# define MDSPAN_FUNCTION_REQUIRES(PAREN_PREQUALS, FNAME, PAREN_PARAMS, QUALS, REQ) \
+ MDSPAN_PP_REMOVE_PARENS(PAREN_PREQUALS) FNAME PAREN_PARAMS QUALS requires REQ \
+ /**/
+#else
+# define MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) , typename ::std::enable_if<(REQ), int>::type = 0>
+# define MDSPAN_FUNCTION_REQUIRES(PAREN_PREQUALS, FNAME, PAREN_PARAMS, QUALS, REQ) \
+ MDSPAN_TEMPLATE_REQUIRES( \
+ class __function_requires_ignored=void, \
+ (std::is_void<__function_requires_ignored>::value && REQ) \
+ ) MDSPAN_PP_REMOVE_PARENS(PAREN_PREQUALS) FNAME PAREN_PARAMS QUALS \
+ /**/
+#endif
+
+#if defined(_MDSPAN_COMPILER_MSVC) && (!defined(_MSVC_TRADITIONAL) || _MSVC_TRADITIONAL)
+# define MDSPAN_TEMPLATE_REQUIRES(...) \
+ MDSPAN_PP_CAT( \
+ MDSPAN_PP_CAT(MDSPAN_TEMPLATE_REQUIRES_, MDSPAN_PP_COUNT(__VA_ARGS__))\
+ (__VA_ARGS__), \
+ ) \
+ /**/
+#else
+# define MDSPAN_TEMPLATE_REQUIRES(...) \
+ MDSPAN_PP_EVAL( \
+ MDSPAN_PP_CAT(MDSPAN_TEMPLATE_REQUIRES_, MDSPAN_PP_COUNT(__VA_ARGS__)), \
+ __VA_ARGS__ \
+ ) \
+ /**/
+#endif
+
+#define MDSPAN_TEMPLATE_REQUIRES_2(TP1, REQ) \
+ template<TP1 \
+ MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) \
+ /**/
+#define MDSPAN_TEMPLATE_REQUIRES_3(TP1, TP2, REQ) \
+ template<TP1, TP2 \
+ MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) \
+ /**/
+#define MDSPAN_TEMPLATE_REQUIRES_4(TP1, TP2, TP3, REQ) \
+ template<TP1, TP2, TP3 \
+ MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) \
+ /**/
+#define MDSPAN_TEMPLATE_REQUIRES_5(TP1, TP2, TP3, TP4, REQ) \
+ template<TP1, TP2, TP3, TP4 \
+ MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) \
+ /**/
+#define MDSPAN_TEMPLATE_REQUIRES_6(TP1, TP2, TP3, TP4, TP5, REQ) \
+ template<TP1, TP2, TP3, TP4, TP5 \
+ MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) \
+ /**/
+#define MDSPAN_TEMPLATE_REQUIRES_7(TP1, TP2, TP3, TP4, TP5, TP6, REQ) \
+ template<TP1, TP2, TP3, TP4, TP5, TP6 \
+ MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) \
+ /**/
+#define MDSPAN_TEMPLATE_REQUIRES_8(TP1, TP2, TP3, TP4, TP5, TP6, TP7, REQ) \
+ template<TP1, TP2, TP3, TP4, TP5, TP6, TP7 \
+ MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) \
+ /**/
+#define MDSPAN_TEMPLATE_REQUIRES_9(TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, REQ) \
+ template<TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8 \
+ MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) \
+ /**/
+#define MDSPAN_TEMPLATE_REQUIRES_10(TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, REQ) \
+ template<TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9 \
+ MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) \
+ /**/
+#define MDSPAN_TEMPLATE_REQUIRES_11(TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, TP10, REQ) \
+ template<TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, TP10 \
+ MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) \
+ /**/
+#define MDSPAN_TEMPLATE_REQUIRES_12(TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, TP10, TP11, REQ) \
+ template<TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, TP10, TP11 \
+ MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) \
+ /**/
+#define MDSPAN_TEMPLATE_REQUIRES_13(TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, TP10, TP11, TP12, REQ) \
+ template<TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, TP10, TP11, TP12 \
+ MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) \
+ /**/
+#define MDSPAN_TEMPLATE_REQUIRES_14(TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, TP10, TP11, TP12, TP13, REQ) \
+ template<TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, TP10, TP11, TP12, TP13 \
+ MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) \
+ /**/
+#define MDSPAN_TEMPLATE_REQUIRES_15(TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, TP10, TP11, TP12, TP13, TP14, REQ) \
+ template<TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, TP10, TP11, TP12, TP13, TP14 \
+ MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) \
+ /**/
+#define MDSPAN_TEMPLATE_REQUIRES_16(TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, TP10, TP11, TP12, TP13, TP14, TP15, REQ) \
+ template<TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, TP10, TP11, TP12, TP13, TP14, TP15 \
+ MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) \
+ /**/
+#define MDSPAN_TEMPLATE_REQUIRES_17(TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, TP10, TP11, TP12, TP13, TP14, TP15, TP16, REQ) \
+ template<TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, TP10, TP11, TP12, TP13, TP14, TP15, TP16 \
+ MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) \
+ /**/
+#define MDSPAN_TEMPLATE_REQUIRES_18(TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, TP10, TP11, TP12, TP13, TP14, TP15, TP16, TP17, REQ) \
+ template<TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, TP10, TP11, TP12, TP13, TP14, TP15, TP16, TP17 \
+ MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) \
+ /**/
+#define MDSPAN_TEMPLATE_REQUIRES_19(TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, TP10, TP11, TP12, TP13, TP14, TP15, TP16, TP17, TP18, REQ) \
+ template<TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, TP10, TP11, TP12, TP13, TP14, TP15, TP16, TP17, TP18 \
+ MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) \
+ /**/
+#define MDSPAN_TEMPLATE_REQUIRES_20(TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, TP10, TP11, TP12, TP13, TP14, TP15, TP16, TP17, TP18, TP19, REQ) \
+ template<TP1, TP2, TP3, TP4, TP5, TP6, TP7, TP8, TP9, TP10, TP11, TP12, TP13, TP14, TP15, TP16, TP17, TP18, TP19 \
+ MDSPAN_CLOSE_ANGLE_REQUIRES(REQ) \
+ /**/
+
+#define MDSPAN_INSTANTIATE_ONLY_IF_USED \
+ MDSPAN_TEMPLATE_REQUIRES( \
+ class __instantiate_only_if_used_tparam=void, \
+ ( _MDSPAN_TRAIT(std::is_void, __instantiate_only_if_used_tparam) ) \
+ ) \
+ /**/
+
+// </editor-fold> end Concept emulation }}}1
+//==============================================================================
+
+//==============================================================================
+// <editor-fold desc="inline variables"> {{{1
+
+#ifdef _MDSPAN_USE_INLINE_VARIABLES
+# define _MDSPAN_INLINE_VARIABLE inline
+#else
+# define _MDSPAN_INLINE_VARIABLE
+#endif
+
+// </editor-fold> end inline variables }}}1
+//==============================================================================
+
+//==============================================================================
+// <editor-fold desc="Return type deduction"> {{{1
+
+#if _MDSPAN_USE_RETURN_TYPE_DEDUCTION
+# define _MDSPAN_DEDUCE_RETURN_TYPE_SINGLE_LINE(SIGNATURE, BODY) \
+ auto MDSPAN_PP_REMOVE_PARENS(SIGNATURE) { return MDSPAN_PP_REMOVE_PARENS(BODY); }
+# define _MDSPAN_DEDUCE_DECLTYPE_AUTO_RETURN_TYPE_SINGLE_LINE(SIGNATURE, BODY) \
+ decltype(auto) MDSPAN_PP_REMOVE_PARENS(SIGNATURE) { return MDSPAN_PP_REMOVE_PARENS(BODY); }
+#else
+# define _MDSPAN_DEDUCE_RETURN_TYPE_SINGLE_LINE(SIGNATURE, BODY) \
+ auto MDSPAN_PP_REMOVE_PARENS(SIGNATURE) \
+ -> std::remove_cv_t<std::remove_reference_t<decltype(BODY)>> \
+ { return MDSPAN_PP_REMOVE_PARENS(BODY); }
+# define _MDSPAN_DEDUCE_DECLTYPE_AUTO_RETURN_TYPE_SINGLE_LINE(SIGNATURE, BODY) \
+ auto MDSPAN_PP_REMOVE_PARENS(SIGNATURE) \
+ -> decltype(BODY) \
+ { return MDSPAN_PP_REMOVE_PARENS(BODY); }
+
+#endif
+
+// </editor-fold> end Return type deduction }}}1
+//==============================================================================
+
+//==============================================================================
+// <editor-fold desc="fold expressions"> {{{1
+
+struct __mdspan_enable_fold_comma { };
+
+#ifdef _MDSPAN_USE_FOLD_EXPRESSIONS
+# define _MDSPAN_FOLD_AND(...) ((__VA_ARGS__) && ...)
+# define _MDSPAN_FOLD_AND_TEMPLATE(...) ((__VA_ARGS__) && ...)
+# define _MDSPAN_FOLD_OR(...) ((__VA_ARGS__) || ...)
+# define _MDSPAN_FOLD_ASSIGN_LEFT(INIT, ...) (INIT = ... = (__VA_ARGS__))
+# define _MDSPAN_FOLD_ASSIGN_RIGHT(PACK, ...) (PACK = ... = (__VA_ARGS__))
+# define _MDSPAN_FOLD_TIMES_RIGHT(PACK, ...) (PACK * ... * (__VA_ARGS__))
+# define _MDSPAN_FOLD_PLUS_RIGHT(PACK, ...) (PACK + ... + (__VA_ARGS__))
+# define _MDSPAN_FOLD_COMMA(...) ((__VA_ARGS__), ...)
+#else
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+
+namespace __fold_compatibility_impl {
+
+// We could probably be more clever here, but at the (small) risk of losing some compiler understanding. For the
+// few operations we need, it's not worth generalizing over the operation
+
+#if _MDSPAN_USE_RETURN_TYPE_DEDUCTION
+
+MDSPAN_FORCE_INLINE_FUNCTION
+constexpr decltype(auto) __fold_right_and_impl() {
+ return true;
+}
+
+template <class Arg, class... Args>
+MDSPAN_FORCE_INLINE_FUNCTION
+constexpr decltype(auto) __fold_right_and_impl(Arg&& arg, Args&&... args) {
+ return ((Arg&&)arg) && __fold_compatibility_impl::__fold_right_and_impl((Args&&)args...);
+}
+
+MDSPAN_FORCE_INLINE_FUNCTION
+constexpr decltype(auto) __fold_right_or_impl() {
+ return false;
+}
+
+template <class Arg, class... Args>
+MDSPAN_FORCE_INLINE_FUNCTION
+constexpr auto __fold_right_or_impl(Arg&& arg, Args&&... args) {
+ return ((Arg&&)arg) || __fold_compatibility_impl::__fold_right_or_impl((Args&&)args...);
+}
+
+template <class Arg1>
+MDSPAN_FORCE_INLINE_FUNCTION
+constexpr auto __fold_left_assign_impl(Arg1&& arg1) {
+ return (Arg1&&)arg1;
+}
+
+template <class Arg1, class Arg2, class... Args>
+MDSPAN_FORCE_INLINE_FUNCTION
+constexpr auto __fold_left_assign_impl(Arg1&& arg1, Arg2&& arg2, Args&&... args) {
+ return __fold_compatibility_impl::__fold_left_assign_impl((((Arg1&&)arg1) = ((Arg2&&)arg2)), (Args&&)args...);
+}
+
+template <class Arg1>
+MDSPAN_FORCE_INLINE_FUNCTION
+constexpr auto __fold_right_assign_impl(Arg1&& arg1) {
+ return (Arg1&&)arg1;
+}
+
+template <class Arg1, class Arg2, class... Args>
+MDSPAN_FORCE_INLINE_FUNCTION
+constexpr auto __fold_right_assign_impl(Arg1&& arg1, Arg2&& arg2, Args&&... args) {
+ return ((Arg1&&)arg1) = __fold_compatibility_impl::__fold_right_assign_impl((Arg2&&)arg2, (Args&&)args...);
+}
+
+template <class Arg1>
+MDSPAN_FORCE_INLINE_FUNCTION
+constexpr auto __fold_right_plus_impl(Arg1&& arg1) {
+ return (Arg1&&)arg1;
+}
+
+template <class Arg1, class Arg2, class... Args>
+MDSPAN_FORCE_INLINE_FUNCTION
+constexpr auto __fold_right_plus_impl(Arg1&& arg1, Arg2&& arg2, Args&&... args) {
+ return ((Arg1&&)arg1) + __fold_compatibility_impl::__fold_right_plus_impl((Arg2&&)arg2, (Args&&)args...);
+}
+
+template <class Arg1>
+MDSPAN_FORCE_INLINE_FUNCTION
+constexpr auto __fold_right_times_impl(Arg1&& arg1) {
+ return (Arg1&&)arg1;
+}
+
+template <class Arg1, class Arg2, class... Args>
+MDSPAN_FORCE_INLINE_FUNCTION
+constexpr auto __fold_right_times_impl(Arg1&& arg1, Arg2&& arg2, Args&&... args) {
+ return ((Arg1&&)arg1) * __fold_compatibility_impl::__fold_right_times_impl((Arg2&&)arg2, (Args&&)args...);
+}
+
+#else
+
+//------------------------------------------------------------------------------
+// <editor-fold desc="right and"> {{{2
+
+template <class... Args>
+struct __fold_right_and_impl_;
+template <>
+struct __fold_right_and_impl_<> {
+ using __rv = bool;
+ MDSPAN_FORCE_INLINE_FUNCTION
+ static constexpr __rv
+ __impl() noexcept {
+ return true;
+ }
+};
+template <class Arg, class... Args>
+struct __fold_right_and_impl_<Arg, Args...> {
+ using __next_t = __fold_right_and_impl_<Args...>;
+ using __rv = decltype(std::declval<Arg>() && std::declval<typename __next_t::__rv>());
+ MDSPAN_FORCE_INLINE_FUNCTION
+ static constexpr __rv
+ __impl(Arg&& arg, Args&&... args) noexcept {
+ return ((Arg&&)arg) && __next_t::__impl((Args&&)args...);
+ }
+};
+
+template <class... Args>
+MDSPAN_FORCE_INLINE_FUNCTION
+constexpr typename __fold_right_and_impl_<Args...>::__rv
+__fold_right_and_impl(Args&&... args) {
+ return __fold_right_and_impl_<Args...>::__impl((Args&&)args...);
+}
+
+// </editor-fold> end right and }}}2
+//------------------------------------------------------------------------------
+
+//------------------------------------------------------------------------------
+// <editor-fold desc="right or"> {{{2
+
+template <class... Args>
+struct __fold_right_or_impl_;
+template <>
+struct __fold_right_or_impl_<> {
+ using __rv = bool;
+ MDSPAN_FORCE_INLINE_FUNCTION
+ static constexpr __rv
+ __impl() noexcept {
+ return false;
+ }
+};
+template <class Arg, class... Args>
+struct __fold_right_or_impl_<Arg, Args...> {
+ using __next_t = __fold_right_or_impl_<Args...>;
+ using __rv = decltype(std::declval<Arg>() || std::declval<typename __next_t::__rv>());
+ MDSPAN_FORCE_INLINE_FUNCTION
+ static constexpr __rv
+ __impl(Arg&& arg, Args&&... args) noexcept {
+ return ((Arg&&)arg) || __next_t::__impl((Args&&)args...);
+ }
+};
+
+template <class... Args>
+MDSPAN_FORCE_INLINE_FUNCTION
+constexpr typename __fold_right_or_impl_<Args...>::__rv
+__fold_right_or_impl(Args&&... args) {
+ return __fold_right_or_impl_<Args...>::__impl((Args&&)args...);
+}
+
+// </editor-fold> end right or }}}2
+//------------------------------------------------------------------------------
+
+//------------------------------------------------------------------------------
+// <editor-fold desc="right plus"> {{{2
+
+template <class... Args>
+struct __fold_right_plus_impl_;
+template <class Arg>
+struct __fold_right_plus_impl_<Arg> {
+ using __rv = Arg&&;
+ MDSPAN_FORCE_INLINE_FUNCTION
+ static constexpr __rv
+ __impl(Arg&& arg) noexcept {
+ return (Arg&&)arg;
+ }
+};
+template <class Arg1, class Arg2, class... Args>
+struct __fold_right_plus_impl_<Arg1, Arg2, Args...> {
+ using __next_t = __fold_right_plus_impl_<Arg2, Args...>;
+ using __rv = decltype(std::declval<Arg1>() + std::declval<typename __next_t::__rv>());
+ MDSPAN_FORCE_INLINE_FUNCTION
+ static constexpr __rv
+ __impl(Arg1&& arg, Arg2&& arg2, Args&&... args) noexcept {
+ return ((Arg1&&)arg) + __next_t::__impl((Arg2&&)arg2, (Args&&)args...);
+ }
+};
+
+template <class... Args>
+MDSPAN_FORCE_INLINE_FUNCTION
+constexpr typename __fold_right_plus_impl_<Args...>::__rv
+__fold_right_plus_impl(Args&&... args) {
+ return __fold_right_plus_impl_<Args...>::__impl((Args&&)args...);
+}
+
+// </editor-fold> end right plus }}}2
+//------------------------------------------------------------------------------
+
+//------------------------------------------------------------------------------
+// <editor-fold desc="right times"> {{{2
+
+template <class... Args>
+struct __fold_right_times_impl_;
+template <class Arg>
+struct __fold_right_times_impl_<Arg> {
+ using __rv = Arg&&;
+ MDSPAN_FORCE_INLINE_FUNCTION
+ static constexpr __rv
+ __impl(Arg&& arg) noexcept {
+ return (Arg&&)arg;
+ }
+};
+template <class Arg1, class Arg2, class... Args>
+struct __fold_right_times_impl_<Arg1, Arg2, Args...> {
+ using __next_t = __fold_right_times_impl_<Arg2, Args...>;
+ using __rv = decltype(std::declval<Arg1>() * std::declval<typename __next_t::__rv>());
+ MDSPAN_FORCE_INLINE_FUNCTION
+ static constexpr __rv
+ __impl(Arg1&& arg, Arg2&& arg2, Args&&... args) noexcept {
+ return ((Arg1&&)arg) * __next_t::__impl((Arg2&&)arg2, (Args&&)args...);
+ }
+};
+
+template <class... Args>
+MDSPAN_FORCE_INLINE_FUNCTION
+constexpr typename __fold_right_times_impl_<Args...>::__rv
+__fold_right_times_impl(Args&&... args) {
+ return __fold_right_times_impl_<Args...>::__impl((Args&&)args...);
+}
+
+// </editor-fold> end right times }}}2
+//------------------------------------------------------------------------------
+
+//------------------------------------------------------------------------------
+// <editor-fold desc="right assign"> {{{2
+
+template <class... Args>
+struct __fold_right_assign_impl_;
+template <class Arg>
+struct __fold_right_assign_impl_<Arg> {
+ using __rv = Arg&&;
+ MDSPAN_FORCE_INLINE_FUNCTION
+ static constexpr __rv
+ __impl(Arg&& arg) noexcept {
+ return (Arg&&)arg;
+ }
+};
+template <class Arg1, class Arg2, class... Args>
+struct __fold_right_assign_impl_<Arg1, Arg2, Args...> {
+ using __next_t = __fold_right_assign_impl_<Arg2, Args...>;
+ using __rv = decltype(std::declval<Arg1>() = std::declval<typename __next_t::__rv>());
+ MDSPAN_FORCE_INLINE_FUNCTION
+ static constexpr __rv
+ __impl(Arg1&& arg, Arg2&& arg2, Args&&... args) noexcept {
+ return ((Arg1&&)arg) = __next_t::__impl((Arg2&&)arg2, (Args&&)args...);
+ }
+};
+
+template <class... Args>
+MDSPAN_FORCE_INLINE_FUNCTION
+constexpr typename __fold_right_assign_impl_<Args...>::__rv
+__fold_right_assign_impl(Args&&... args) {
+ return __fold_right_assign_impl_<Args...>::__impl((Args&&)args...);
+}
+
+// </editor-fold> end right assign }}}2
+//------------------------------------------------------------------------------
+
+//------------------------------------------------------------------------------
+// <editor-fold desc="left assign"> {{{2
+
+template <class... Args>
+struct __fold_left_assign_impl_;
+template <class Arg>
+struct __fold_left_assign_impl_<Arg> {
+ using __rv = Arg&&;
+ MDSPAN_FORCE_INLINE_FUNCTION
+ static constexpr __rv
+ __impl(Arg&& arg) noexcept {
+ return (Arg&&)arg;
+ }
+};
+template <class Arg1, class Arg2, class... Args>
+struct __fold_left_assign_impl_<Arg1, Arg2, Args...> {
+ using __assign_result_t = decltype(std::declval<Arg1>() = std::declval<Arg2>());
+ using __next_t = __fold_left_assign_impl_<__assign_result_t, Args...>;
+ using __rv = typename __next_t::__rv;
+ MDSPAN_FORCE_INLINE_FUNCTION
+ static constexpr __rv
+ __impl(Arg1&& arg, Arg2&& arg2, Args&&... args) noexcept {
+ return __next_t::__impl(((Arg1&&)arg) = (Arg2&&)arg2, (Args&&)args...);
+ }
+};
+
+template <class... Args>
+MDSPAN_FORCE_INLINE_FUNCTION
+constexpr typename __fold_left_assign_impl_<Args...>::__rv
+__fold_left_assign_impl(Args&&... args) {
+ return __fold_left_assign_impl_<Args...>::__impl((Args&&)args...);
+}
+
+// </editor-fold> end left assign }}}2
+//------------------------------------------------------------------------------
+
+#endif
+
+
+template <class... Args>
+constexpr __mdspan_enable_fold_comma __fold_comma_impl(Args&&...) noexcept { return { }; }
+
+template <bool... Bs>
+struct __bools;
+
+} // __fold_compatibility_impl
+
+} // end namespace MDSPAN_IMPL_STANDARD_NAMESPACE
+
+# define _MDSPAN_FOLD_AND(...) MDSPAN_IMPL_STANDARD_NAMESPACE::__fold_compatibility_impl::__fold_right_and_impl((__VA_ARGS__)...)
+# define _MDSPAN_FOLD_OR(...) MDSPAN_IMPL_STANDARD_NAMESPACE::__fold_compatibility_impl::__fold_right_or_impl((__VA_ARGS__)...)
+# define _MDSPAN_FOLD_ASSIGN_LEFT(INIT, ...) MDSPAN_IMPL_STANDARD_NAMESPACE::__fold_compatibility_impl::__fold_left_assign_impl(INIT, (__VA_ARGS__)...)
+# define _MDSPAN_FOLD_ASSIGN_RIGHT(PACK, ...) MDSPAN_IMPL_STANDARD_NAMESPACE::__fold_compatibility_impl::__fold_right_assign_impl((PACK)..., __VA_ARGS__)
+# define _MDSPAN_FOLD_TIMES_RIGHT(PACK, ...) MDSPAN_IMPL_STANDARD_NAMESPACE::__fold_compatibility_impl::__fold_right_times_impl((PACK)..., __VA_ARGS__)
+# define _MDSPAN_FOLD_PLUS_RIGHT(PACK, ...) MDSPAN_IMPL_STANDARD_NAMESPACE::__fold_compatibility_impl::__fold_right_plus_impl((PACK)..., __VA_ARGS__)
+# define _MDSPAN_FOLD_COMMA(...) MDSPAN_IMPL_STANDARD_NAMESPACE::__fold_compatibility_impl::__fold_comma_impl((__VA_ARGS__)...)
+
+# define _MDSPAN_FOLD_AND_TEMPLATE(...) \
+ _MDSPAN_TRAIT(std::is_same, __fold_compatibility_impl::__bools<(__VA_ARGS__)..., true>, __fold_compatibility_impl::__bools<true, (__VA_ARGS__)...>)
+
+#endif
+
+// </editor-fold> end fold expressions }}}1
+//==============================================================================
+
+//==============================================================================
+// <editor-fold desc="Variable template compatibility"> {{{1
+
+#if _MDSPAN_USE_VARIABLE_TEMPLATES
+# define _MDSPAN_TRAIT(TRAIT, ...) TRAIT##_v<__VA_ARGS__>
+#else
+# define _MDSPAN_TRAIT(TRAIT, ...) TRAIT<__VA_ARGS__>::value
+#endif
+
+// </editor-fold> end Variable template compatibility }}}1
+//==============================================================================
+
+//==============================================================================
+// <editor-fold desc="Pre-C++14 constexpr"> {{{1
+
+#if _MDSPAN_USE_CONSTEXPR_14
+# define _MDSPAN_CONSTEXPR_14 constexpr
+// Workaround for a bug (I think?) in EDG frontends
+# ifdef __EDG__
+# define _MDSPAN_CONSTEXPR_14_DEFAULTED
+# else
+# define _MDSPAN_CONSTEXPR_14_DEFAULTED constexpr
+# endif
+#else
+# define _MDSPAN_CONSTEXPR_14
+# define _MDSPAN_CONSTEXPR_14_DEFAULTED
+#endif
+
+// </editor-fold> end Pre-C++14 constexpr }}}1
+//==============================================================================
diff --git a/ext/mdspan/include/experimental/__p0009_bits/mdspan.hpp b/ext/mdspan/include/experimental/__p0009_bits/mdspan.hpp
new file mode 100644
index 0000000..23114aa
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p0009_bits/mdspan.hpp
@@ -0,0 +1,432 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+
+#pragma once
+
+#include "default_accessor.hpp"
+#include "layout_right.hpp"
+#include "extents.hpp"
+#include "trait_backports.hpp"
+#include "compressed_pair.hpp"
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+template <
+ class ElementType,
+ class Extents,
+ class LayoutPolicy = layout_right,
+ class AccessorPolicy = default_accessor<ElementType>
+>
+class mdspan
+{
+private:
+ static_assert(detail::__is_extents_v<Extents>,
+ MDSPAN_IMPL_STANDARD_NAMESPACE_STRING "::mdspan's Extents template parameter must be a specialization of " MDSPAN_IMPL_STANDARD_NAMESPACE_STRING "::extents.");
+ static_assert(std::is_same<ElementType, typename AccessorPolicy::element_type>::value,
+ MDSPAN_IMPL_STANDARD_NAMESPACE_STRING "::mdspan's ElementType template parameter must be the same as its AccessorPolicy::element_type.");
+
+ // Workaround for non-deducibility of the index sequence template parameter if it's given at the top level
+ template <class>
+ struct __deduction_workaround;
+
+ template <size_t... Idxs>
+ struct __deduction_workaround<std::index_sequence<Idxs...>>
+ {
+ MDSPAN_FORCE_INLINE_FUNCTION static constexpr
+ size_t __size(mdspan const& __self) noexcept {
+ return _MDSPAN_FOLD_TIMES_RIGHT((__self.__mapping_ref().extents().extent(Idxs)), /* * ... * */ size_t(1));
+ }
+ MDSPAN_FORCE_INLINE_FUNCTION static constexpr
+ bool __empty(mdspan const& __self) noexcept {
+ return (__self.rank()>0) && _MDSPAN_FOLD_OR((__self.__mapping_ref().extents().extent(Idxs)==index_type(0)));
+ }
+ template <class ReferenceType, class SizeType, size_t N>
+ MDSPAN_FORCE_INLINE_FUNCTION static constexpr
+ ReferenceType __callop(mdspan const& __self, const std::array<SizeType, N>& indices) noexcept {
+ return __self.__accessor_ref().access(__self.__ptr_ref(), __self.__mapping_ref()(indices[Idxs]...));
+ }
+#ifdef __cpp_lib_span
+ template <class ReferenceType, class SizeType, size_t N>
+ MDSPAN_FORCE_INLINE_FUNCTION static constexpr
+ ReferenceType __callop(mdspan const& __self, const std::span<SizeType, N>& indices) noexcept {
+ return __self.__accessor_ref().access(__self.__ptr_ref(), __self.__mapping_ref()(indices[Idxs]...));
+ }
+#endif
+ };
+
+public:
+
+ //--------------------------------------------------------------------------------
+ // Domain and codomain types
+
+ using extents_type = Extents;
+ using layout_type = LayoutPolicy;
+ using accessor_type = AccessorPolicy;
+ using mapping_type = typename layout_type::template mapping<extents_type>;
+ using element_type = ElementType;
+ using value_type = std::remove_cv_t<element_type>;
+ using index_type = typename extents_type::index_type;
+ using size_type = typename extents_type::size_type;
+ using rank_type = typename extents_type::rank_type;
+ using data_handle_type = typename accessor_type::data_handle_type;
+ using reference = typename accessor_type::reference;
+
+ MDSPAN_INLINE_FUNCTION static constexpr size_t rank() noexcept { return extents_type::rank(); }
+ MDSPAN_INLINE_FUNCTION static constexpr size_t rank_dynamic() noexcept { return extents_type::rank_dynamic(); }
+ MDSPAN_INLINE_FUNCTION static constexpr size_t static_extent(size_t r) noexcept { return extents_type::static_extent(r); }
+ MDSPAN_INLINE_FUNCTION constexpr index_type extent(size_t r) const noexcept { return __mapping_ref().extents().extent(r); };
+
+private:
+
+ // Can't use defaulted parameter in the __deduction_workaround template because of a bug in MSVC warning C4348.
+ using __impl = __deduction_workaround<std::make_index_sequence<extents_type::rank()>>;
+
+ using __map_acc_pair_t = detail::__compressed_pair<mapping_type, accessor_type>;
+
+public:
+
+ //--------------------------------------------------------------------------------
+ // [mdspan.basic.cons], mdspan constructors, assignment, and destructor
+
+#if !MDSPAN_HAS_CXX_20
+ MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr mdspan() = default;
+#else
+ MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr mdspan()
+ requires(
+ // nvhpc has a bug where using just rank_dynamic() here doesn't work ...
+ (extents_type::rank_dynamic() > 0) &&
+ _MDSPAN_TRAIT(std::is_default_constructible, data_handle_type) &&
+ _MDSPAN_TRAIT(std::is_default_constructible, mapping_type) &&
+ _MDSPAN_TRAIT(std::is_default_constructible, accessor_type)
+ ) = default;
+#endif
+ MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr mdspan(const mdspan&) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr mdspan(mdspan&&) = default;
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class... SizeTypes,
+ /* requires */ (
+ ((sizeof...(SizeTypes) == rank()) || (sizeof...(SizeTypes) == rank_dynamic())) &&
+ (detail::are_valid_indices<index_type, SizeTypes...>()) &&
+ _MDSPAN_TRAIT(std::is_constructible, mapping_type, extents_type) &&
+ _MDSPAN_TRAIT(std::is_default_constructible, accessor_type)
+ )
+ )
+ MDSPAN_INLINE_FUNCTION
+ explicit constexpr mdspan(data_handle_type p, SizeTypes... dynamic_extents)
+ // TODO @proposal-bug shouldn't I be allowed to do `move(p)` here?
+ : __members(std::move(p), __map_acc_pair_t(mapping_type(extents_type(static_cast<index_type>(std::move(dynamic_extents))...)), accessor_type()))
+ { }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class SizeType, size_t N,
+ /* requires */ (
+ _MDSPAN_TRAIT(std::is_convertible, const SizeType&, index_type) &&
+ _MDSPAN_TRAIT(std::is_nothrow_constructible, index_type, const SizeType&) &&
+ ((N == rank()) || (N == rank_dynamic())) &&
+ _MDSPAN_TRAIT(std::is_constructible, mapping_type, extents_type) &&
+ _MDSPAN_TRAIT(std::is_default_constructible, accessor_type)
+ )
+ )
+ MDSPAN_CONDITIONAL_EXPLICIT(N != rank_dynamic())
+ MDSPAN_INLINE_FUNCTION
+ constexpr mdspan(data_handle_type p, const std::array<SizeType, N>& dynamic_extents)
+ : __members(std::move(p), __map_acc_pair_t(mapping_type(extents_type(dynamic_extents)), accessor_type()))
+ { }
+
+#ifdef __cpp_lib_span
+ MDSPAN_TEMPLATE_REQUIRES(
+ class SizeType, size_t N,
+ /* requires */ (
+ _MDSPAN_TRAIT(std::is_convertible, const SizeType&, index_type) &&
+ _MDSPAN_TRAIT(std::is_nothrow_constructible, index_type, const SizeType&) &&
+ ((N == rank()) || (N == rank_dynamic())) &&
+ _MDSPAN_TRAIT(std::is_constructible, mapping_type, extents_type) &&
+ _MDSPAN_TRAIT(std::is_default_constructible, accessor_type)
+ )
+ )
+ MDSPAN_CONDITIONAL_EXPLICIT(N != rank_dynamic())
+ MDSPAN_INLINE_FUNCTION
+ constexpr mdspan(data_handle_type p, std::span<SizeType, N> dynamic_extents)
+ : __members(std::move(p), __map_acc_pair_t(mapping_type(extents_type(as_const(dynamic_extents))), accessor_type()))
+ { }
+#endif
+
+ MDSPAN_FUNCTION_REQUIRES(
+ (MDSPAN_INLINE_FUNCTION constexpr),
+ mdspan, (data_handle_type p, const extents_type& exts), ,
+ /* requires */ (_MDSPAN_TRAIT(std::is_default_constructible, accessor_type) &&
+ _MDSPAN_TRAIT(std::is_constructible, mapping_type, const extents_type&))
+ ) : __members(std::move(p), __map_acc_pair_t(mapping_type(exts), accessor_type()))
+ { }
+
+ MDSPAN_FUNCTION_REQUIRES(
+ (MDSPAN_INLINE_FUNCTION constexpr),
+ mdspan, (data_handle_type p, const mapping_type& m), ,
+ /* requires */ (_MDSPAN_TRAIT(std::is_default_constructible, accessor_type))
+ ) : __members(std::move(p), __map_acc_pair_t(m, accessor_type()))
+ { }
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr mdspan(data_handle_type p, const mapping_type& m, const accessor_type& a)
+ : __members(std::move(p), __map_acc_pair_t(m, a))
+ { }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherElementType, class OtherExtents, class OtherLayoutPolicy, class OtherAccessor,
+ /* requires */ (
+ _MDSPAN_TRAIT(std::is_constructible, mapping_type, const typename OtherLayoutPolicy::template mapping<OtherExtents>&) &&
+ _MDSPAN_TRAIT(std::is_constructible, accessor_type, const OtherAccessor&)
+ )
+ )
+ MDSPAN_CONDITIONAL_EXPLICIT(
+ !_MDSPAN_TRAIT(std::is_convertible, const typename OtherLayoutPolicy::template mapping<OtherExtents>&, mapping_type) ||
+ !_MDSPAN_TRAIT(std::is_convertible, const OtherAccessor&, accessor_type)
+ )
+ MDSPAN_INLINE_FUNCTION
+ constexpr mdspan(const mdspan<OtherElementType, OtherExtents, OtherLayoutPolicy, OtherAccessor>& other)
+ : __members(other.__ptr_ref(), __map_acc_pair_t(other.__mapping_ref(), other.__accessor_ref()))
+ {
+ static_assert(_MDSPAN_TRAIT(std::is_constructible, data_handle_type, typename OtherAccessor::data_handle_type),"Incompatible data_handle_type for mdspan construction");
+ static_assert(_MDSPAN_TRAIT(std::is_constructible, extents_type, OtherExtents),"Incompatible extents for mdspan construction");
+ /*
+ * TODO: Check precondition
+ * For each rank index r of extents_type, static_extent(r) == dynamic_extent || static_extent(r) == other.extent(r) is true.
+ */
+ }
+
+ /* Might need this on NVIDIA?
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ ~mdspan() = default;
+ */
+
+ MDSPAN_INLINE_FUNCTION_DEFAULTED _MDSPAN_CONSTEXPR_14_DEFAULTED mdspan& operator=(const mdspan&) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED _MDSPAN_CONSTEXPR_14_DEFAULTED mdspan& operator=(mdspan&&) = default;
+
+
+ //--------------------------------------------------------------------------------
+ // [mdspan.basic.mapping], mdspan mapping domain multidimensional index to access codomain element
+
+ #if MDSPAN_USE_BRACKET_OPERATOR
+ MDSPAN_TEMPLATE_REQUIRES(
+ class... SizeTypes,
+ /* requires */ (
+ _MDSPAN_FOLD_AND(_MDSPAN_TRAIT(std::is_convertible, SizeTypes, index_type) /* && ... */) &&
+ _MDSPAN_FOLD_AND(_MDSPAN_TRAIT(std::is_nothrow_constructible, index_type, SizeTypes) /* && ... */) &&
+ (rank() == sizeof...(SizeTypes))
+ )
+ )
+ MDSPAN_FORCE_INLINE_FUNCTION
+ constexpr reference operator[](SizeTypes... indices) const
+ {
+ return __accessor_ref().access(__ptr_ref(), __mapping_ref()(static_cast<index_type>(std::move(indices))...));
+ }
+ #endif
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class SizeType,
+ /* requires */ (
+ _MDSPAN_TRAIT(std::is_convertible, const SizeType&, index_type) &&
+ _MDSPAN_TRAIT(std::is_nothrow_constructible, index_type, const SizeType&)
+ )
+ )
+ MDSPAN_FORCE_INLINE_FUNCTION
+ constexpr reference operator[](const std::array< SizeType, rank()>& indices) const
+ {
+ return __impl::template __callop<reference>(*this, indices);
+ }
+
+ #ifdef __cpp_lib_span
+ MDSPAN_TEMPLATE_REQUIRES(
+ class SizeType,
+ /* requires */ (
+ _MDSPAN_TRAIT(std::is_convertible, const SizeType&, index_type) &&
+ _MDSPAN_TRAIT(std::is_nothrow_constructible, index_type, const SizeType&)
+ )
+ )
+ MDSPAN_FORCE_INLINE_FUNCTION
+ constexpr reference operator[](std::span<SizeType, rank()> indices) const
+ {
+ return __impl::template __callop<reference>(*this, indices);
+ }
+ #endif // __cpp_lib_span
+
+ #if !MDSPAN_USE_BRACKET_OPERATOR
+ MDSPAN_TEMPLATE_REQUIRES(
+ class Index,
+ /* requires */ (
+ _MDSPAN_TRAIT(std::is_convertible, Index, index_type) &&
+ _MDSPAN_TRAIT(std::is_nothrow_constructible, index_type, Index) &&
+ extents_type::rank() == 1
+ )
+ )
+ MDSPAN_FORCE_INLINE_FUNCTION
+ constexpr reference operator[](Index idx) const
+ {
+ return __accessor_ref().access(__ptr_ref(), __mapping_ref()(static_cast<index_type>(std::move(idx))));
+ }
+ #endif
+
+ #if MDSPAN_USE_PAREN_OPERATOR
+ MDSPAN_TEMPLATE_REQUIRES(
+ class... SizeTypes,
+ /* requires */ (
+ extents_type::rank() == sizeof...(SizeTypes) &&
+ (detail::are_valid_indices<index_type, SizeTypes...>())
+ )
+ )
+ MDSPAN_FORCE_INLINE_FUNCTION
+ constexpr reference operator()(SizeTypes... indices) const
+ {
+ return __accessor_ref().access(__ptr_ref(), __mapping_ref()(static_cast<index_type>(std::move(indices))...));
+ }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class SizeType,
+ /* requires */ (
+ _MDSPAN_TRAIT(std::is_convertible, const SizeType&, index_type) &&
+ _MDSPAN_TRAIT(std::is_nothrow_constructible, index_type, const SizeType&)
+ )
+ )
+ MDSPAN_FORCE_INLINE_FUNCTION
+ constexpr reference operator()(const std::array<SizeType, rank()>& indices) const
+ {
+ return __impl::template __callop<reference>(*this, indices);
+ }
+
+ #ifdef __cpp_lib_span
+ MDSPAN_TEMPLATE_REQUIRES(
+ class SizeType,
+ /* requires */ (
+ _MDSPAN_TRAIT(std::is_convertible, const SizeType&, index_type) &&
+ _MDSPAN_TRAIT(std::is_nothrow_constructible, index_type, const SizeType&)
+ )
+ )
+ MDSPAN_FORCE_INLINE_FUNCTION
+ constexpr reference operator()(std::span<SizeType, rank()> indices) const
+ {
+ return __impl::template __callop<reference>(*this, indices);
+ }
+ #endif // __cpp_lib_span
+ #endif // MDSPAN_USE_PAREN_OPERATOR
+
+ MDSPAN_INLINE_FUNCTION constexpr size_type size() const noexcept {
+ return static_cast<size_type>(__impl::__size(*this));
+ };
+
+ MDSPAN_INLINE_FUNCTION constexpr bool empty() const noexcept {
+ return __impl::__empty(*this);
+ };
+
+ MDSPAN_INLINE_FUNCTION
+ friend constexpr void swap(mdspan& x, mdspan& y) noexcept {
+ // can't call the std::swap inside on HIP
+ #if !defined(_MDSPAN_HAS_HIP) && !defined(_MDSPAN_HAS_CUDA)
+ using std::swap;
+ swap(x.__ptr_ref(), y.__ptr_ref());
+ swap(x.__mapping_ref(), y.__mapping_ref());
+ swap(x.__accessor_ref(), y.__accessor_ref());
+ #else
+ mdspan tmp = y;
+ y = x;
+ x = tmp;
+ #endif
+ }
+
+ //--------------------------------------------------------------------------------
+ // [mdspan.basic.domobs], mdspan observers of the domain multidimensional index space
+
+
+ MDSPAN_INLINE_FUNCTION constexpr const extents_type& extents() const noexcept { return __mapping_ref().extents(); };
+ MDSPAN_INLINE_FUNCTION constexpr const data_handle_type& data_handle() const noexcept { return __ptr_ref(); };
+ MDSPAN_INLINE_FUNCTION constexpr const mapping_type& mapping() const noexcept { return __mapping_ref(); };
+ MDSPAN_INLINE_FUNCTION constexpr const accessor_type& accessor() const noexcept { return __accessor_ref(); };
+
+ //--------------------------------------------------------------------------------
+ // [mdspan.basic.obs], mdspan observers of the mapping
+
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_unique() { return mapping_type::is_always_unique(); };
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_exhaustive() { return mapping_type::is_always_exhaustive(); };
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_strided() { return mapping_type::is_always_strided(); };
+
+ MDSPAN_INLINE_FUNCTION constexpr bool is_unique() const { return __mapping_ref().is_unique(); };
+ MDSPAN_INLINE_FUNCTION constexpr bool is_exhaustive() const { return __mapping_ref().is_exhaustive(); };
+ MDSPAN_INLINE_FUNCTION constexpr bool is_strided() const { return __mapping_ref().is_strided(); };
+ MDSPAN_INLINE_FUNCTION constexpr index_type stride(size_t r) const { return __mapping_ref().stride(r); };
+
+private:
+
+ detail::__compressed_pair<data_handle_type, __map_acc_pair_t> __members{};
+
+ MDSPAN_FORCE_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 data_handle_type& __ptr_ref() noexcept { return __members.__first(); }
+ MDSPAN_FORCE_INLINE_FUNCTION constexpr data_handle_type const& __ptr_ref() const noexcept { return __members.__first(); }
+ MDSPAN_FORCE_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 mapping_type& __mapping_ref() noexcept { return __members.__second().__first(); }
+ MDSPAN_FORCE_INLINE_FUNCTION constexpr mapping_type const& __mapping_ref() const noexcept { return __members.__second().__first(); }
+ MDSPAN_FORCE_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 accessor_type& __accessor_ref() noexcept { return __members.__second().__second(); }
+ MDSPAN_FORCE_INLINE_FUNCTION constexpr accessor_type const& __accessor_ref() const noexcept { return __members.__second().__second(); }
+
+ template <class, class, class, class>
+ friend class mdspan;
+
+};
+
+#if defined(_MDSPAN_USE_CLASS_TEMPLATE_ARGUMENT_DEDUCTION)
+MDSPAN_TEMPLATE_REQUIRES(
+ class ElementType, class... SizeTypes,
+ /* requires */ _MDSPAN_FOLD_AND(_MDSPAN_TRAIT(std::is_convertible, SizeTypes, size_t) /* && ... */) &&
+ (sizeof...(SizeTypes) > 0)
+)
+MDSPAN_DEDUCTION_GUIDE explicit mdspan(ElementType*, SizeTypes...)
+ -> mdspan<ElementType, ::MDSPAN_IMPL_STANDARD_NAMESPACE::dextents<size_t, sizeof...(SizeTypes)>>;
+
+MDSPAN_TEMPLATE_REQUIRES(
+ class Pointer,
+ (_MDSPAN_TRAIT(std::is_pointer, std::remove_reference_t<Pointer>))
+)
+MDSPAN_DEDUCTION_GUIDE mdspan(Pointer&&) -> mdspan<std::remove_pointer_t<std::remove_reference_t<Pointer>>, extents<size_t>>;
+
+MDSPAN_TEMPLATE_REQUIRES(
+ class CArray,
+ (_MDSPAN_TRAIT(std::is_array, CArray) && (std::rank_v<CArray> == 1))
+)
+MDSPAN_DEDUCTION_GUIDE mdspan(CArray&) -> mdspan<std::remove_all_extents_t<CArray>, extents<size_t, ::std::extent_v<CArray,0>>>;
+
+template <class ElementType, class SizeType, size_t N>
+MDSPAN_DEDUCTION_GUIDE mdspan(ElementType*, const ::std::array<SizeType, N>&)
+ -> mdspan<ElementType, ::MDSPAN_IMPL_STANDARD_NAMESPACE::dextents<size_t, N>>;
+
+#ifdef __cpp_lib_span
+template <class ElementType, class SizeType, size_t N>
+MDSPAN_DEDUCTION_GUIDE mdspan(ElementType*, ::std::span<SizeType, N>)
+ -> mdspan<ElementType, ::MDSPAN_IMPL_STANDARD_NAMESPACE::dextents<size_t, N>>;
+#endif
+
+// This one is necessary because all the constructors take `data_handle_type`s, not
+// `ElementType*`s, and `data_handle_type` is taken from `accessor_type::data_handle_type`, which
+// seems to throw off automatic deduction guides.
+template <class ElementType, class SizeType, size_t... ExtentsPack>
+MDSPAN_DEDUCTION_GUIDE mdspan(ElementType*, const extents<SizeType, ExtentsPack...>&)
+ -> mdspan<ElementType, ::MDSPAN_IMPL_STANDARD_NAMESPACE::extents<SizeType, ExtentsPack...>>;
+
+template <class ElementType, class MappingType>
+MDSPAN_DEDUCTION_GUIDE mdspan(ElementType*, const MappingType&)
+ -> mdspan<ElementType, typename MappingType::extents_type, typename MappingType::layout_type>;
+
+template <class MappingType, class AccessorType>
+MDSPAN_DEDUCTION_GUIDE mdspan(const typename AccessorType::data_handle_type, const MappingType&, const AccessorType&)
+ -> mdspan<typename AccessorType::element_type, typename MappingType::extents_type, typename MappingType::layout_type, AccessorType>;
+#endif
+
+} // end namespace MDSPAN_IMPL_STANDARD_NAMESPACE
diff --git a/ext/mdspan/include/experimental/__p0009_bits/no_unique_address.hpp b/ext/mdspan/include/experimental/__p0009_bits/no_unique_address.hpp
new file mode 100644
index 0000000..36e64ee
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p0009_bits/no_unique_address.hpp
@@ -0,0 +1,97 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+#pragma once
+
+#include "macros.hpp"
+#include "trait_backports.hpp"
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+namespace detail {
+
+//==============================================================================
+
+template <class _T, size_t _Disambiguator = 0, class _Enable = void>
+struct __no_unique_address_emulation {
+ using __stored_type = _T;
+ _T __v;
+ MDSPAN_FORCE_INLINE_FUNCTION constexpr _T const &__ref() const noexcept {
+ return __v;
+ }
+ MDSPAN_FORCE_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 _T &__ref() noexcept {
+ return __v;
+ }
+};
+
+// Empty case
+// This doesn't work if _T is final, of course, but we're not using anything
+// like that currently. That kind of thing could be added pretty easily though
+template <class _T, size_t _Disambiguator>
+struct __no_unique_address_emulation<
+ _T, _Disambiguator,
+ std::enable_if_t<_MDSPAN_TRAIT(std::is_empty, _T) &&
+ // If the type isn't trivially destructible, its destructor
+ // won't be called at the right time, so don't use this
+ // specialization
+ _MDSPAN_TRAIT(std::is_trivially_destructible, _T)>> :
+#ifdef _MDSPAN_COMPILER_MSVC
+ // MSVC doesn't allow you to access public static member functions of a type
+ // when you *happen* to privately inherit from that type.
+ protected
+#else
+ // But we still want this to be private if possible so that we don't accidentally
+ // access members of _T directly rather than calling __ref() first, which wouldn't
+ // work if _T happens to be stateful and thus we're using the unspecialized definition
+ // of __no_unique_address_emulation above.
+ private
+#endif
+ _T {
+ using __stored_type = _T;
+ MDSPAN_FORCE_INLINE_FUNCTION constexpr _T const &__ref() const noexcept {
+ return *static_cast<_T const *>(this);
+ }
+ MDSPAN_FORCE_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 _T &__ref() noexcept {
+ return *static_cast<_T *>(this);
+ }
+
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ constexpr __no_unique_address_emulation() noexcept = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ constexpr __no_unique_address_emulation(
+ __no_unique_address_emulation const &) noexcept = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ constexpr __no_unique_address_emulation(
+ __no_unique_address_emulation &&) noexcept = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ _MDSPAN_CONSTEXPR_14_DEFAULTED __no_unique_address_emulation &
+ operator=(__no_unique_address_emulation const &) noexcept = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ _MDSPAN_CONSTEXPR_14_DEFAULTED __no_unique_address_emulation &
+ operator=(__no_unique_address_emulation &&) noexcept = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ ~__no_unique_address_emulation() noexcept = default;
+
+ // Explicitly make this not a reference so that the copy or move
+ // constructor still gets called.
+ MDSPAN_INLINE_FUNCTION
+ explicit constexpr __no_unique_address_emulation(_T const& __v) noexcept : _T(__v) {}
+ MDSPAN_INLINE_FUNCTION
+ explicit constexpr __no_unique_address_emulation(_T&& __v) noexcept : _T(::std::move(__v)) {}
+};
+
+//==============================================================================
+
+} // end namespace detail
+} // end namespace MDSPAN_IMPL_STANDARD_NAMESPACE
diff --git a/ext/mdspan/include/experimental/__p0009_bits/trait_backports.hpp b/ext/mdspan/include/experimental/__p0009_bits/trait_backports.hpp
new file mode 100644
index 0000000..4933dd9
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p0009_bits/trait_backports.hpp
@@ -0,0 +1,132 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+#ifndef MDSPAN_INCLUDE_EXPERIMENTAL_BITS_TRAIT_BACKPORTS_HPP_
+#define MDSPAN_INCLUDE_EXPERIMENTAL_BITS_TRAIT_BACKPORTS_HPP_
+
+#include "macros.hpp"
+#include "config.hpp"
+
+#include <type_traits>
+#include <utility> // integer_sequence
+
+//==============================================================================
+// <editor-fold desc="Variable template trait backports (e.g., is_void_v)"> {{{1
+
+#ifdef _MDSPAN_NEEDS_TRAIT_VARIABLE_TEMPLATE_BACKPORTS
+
+#if _MDSPAN_USE_VARIABLE_TEMPLATES
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+
+#define _MDSPAN_BACKPORT_TRAIT(TRAIT) \
+ template <class... Args> _MDSPAN_INLINE_VARIABLE constexpr auto TRAIT##_v = TRAIT<Args...>::value;
+
+_MDSPAN_BACKPORT_TRAIT(is_assignable)
+_MDSPAN_BACKPORT_TRAIT(is_constructible)
+_MDSPAN_BACKPORT_TRAIT(is_convertible)
+_MDSPAN_BACKPORT_TRAIT(is_default_constructible)
+_MDSPAN_BACKPORT_TRAIT(is_trivially_destructible)
+_MDSPAN_BACKPORT_TRAIT(is_same)
+_MDSPAN_BACKPORT_TRAIT(is_empty)
+_MDSPAN_BACKPORT_TRAIT(is_void)
+
+#undef _MDSPAN_BACKPORT_TRAIT
+
+} // end namespace MDSPAN_IMPL_STANDARD_NAMESPACE
+
+#endif // _MDSPAN_USE_VARIABLE_TEMPLATES
+
+#endif // _MDSPAN_NEEDS_TRAIT_VARIABLE_TEMPLATE_BACKPORTS
+
+// </editor-fold> end Variable template trait backports (e.g., is_void_v) }}}1
+//==============================================================================
+
+//==============================================================================
+// <editor-fold desc="integer sequence (ugh...)"> {{{1
+
+#if !defined(_MDSPAN_USE_INTEGER_SEQUENCE) || !_MDSPAN_USE_INTEGER_SEQUENCE
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+
+template <class T, T... Vals>
+struct integer_sequence {
+ static constexpr size_t size() noexcept { return sizeof...(Vals); }
+ using value_type = T;
+};
+
+template <size_t... Vals>
+using index_sequence = std::integer_sequence<size_t, Vals...>;
+
+namespace __detail {
+
+template <class T, T N, T I, class Result>
+struct __make_int_seq_impl;
+
+template <class T, T N, T... Vals>
+struct __make_int_seq_impl<T, N, N, integer_sequence<T, Vals...>>
+{
+ using type = integer_sequence<T, Vals...>;
+};
+
+template <class T, T N, T I, T... Vals>
+struct __make_int_seq_impl<
+ T, N, I, integer_sequence<T, Vals...>
+> : __make_int_seq_impl<T, N, I+1, integer_sequence<T, Vals..., I>>
+{ };
+
+} // end namespace __detail
+
+template <class T, T N>
+using make_integer_sequence = typename __detail::__make_int_seq_impl<T, N, 0, integer_sequence<T>>::type;
+
+template <size_t N>
+using make_index_sequence = typename __detail::__make_int_seq_impl<size_t, N, 0, integer_sequence<size_t>>::type;
+
+template <class... T>
+using index_sequence_for = make_index_sequence<sizeof...(T)>;
+
+} // end namespace MDSPAN_IMPL_STANDARD_NAMESPACE
+
+#endif
+
+// </editor-fold> end integer sequence (ugh...) }}}1
+//==============================================================================
+
+//==============================================================================
+// <editor-fold desc="standard trait aliases"> {{{1
+
+#if !defined(_MDSPAN_USE_STANDARD_TRAIT_ALIASES) || !_MDSPAN_USE_STANDARD_TRAIT_ALIASES
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+
+#define _MDSPAN_BACKPORT_TRAIT_ALIAS(TRAIT) \
+ template <class... Args> using TRAIT##_t = typename TRAIT<Args...>::type;
+
+_MDSPAN_BACKPORT_TRAIT_ALIAS(remove_cv)
+_MDSPAN_BACKPORT_TRAIT_ALIAS(remove_reference)
+
+template <bool _B, class _T=void>
+using enable_if_t = typename enable_if<_B, _T>::type;
+
+#undef _MDSPAN_BACKPORT_TRAIT_ALIAS
+
+} // end namespace MDSPAN_IMPL_STANDARD_NAMESPACE
+
+#endif
+
+// </editor-fold> end standard trait aliases }}}1
+//==============================================================================
+
+#endif //MDSPAN_INCLUDE_EXPERIMENTAL_BITS_TRAIT_BACKPORTS_HPP_
diff --git a/ext/mdspan/include/experimental/__p0009_bits/type_list.hpp b/ext/mdspan/include/experimental/__p0009_bits/type_list.hpp
new file mode 100644
index 0000000..deca7c1
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p0009_bits/type_list.hpp
@@ -0,0 +1,87 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+#include "macros.hpp"
+
+#include "trait_backports.hpp" // make_index_sequence
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+
+//==============================================================================
+
+namespace detail {
+
+template <class... _Ts> struct __type_list { static constexpr auto __size = sizeof...(_Ts); };
+
+// Implementation of type_list at() that's heavily optimized for small typelists
+template <size_t, class> struct __type_at;
+template <size_t, class _Seq, class=std::make_index_sequence<_Seq::__size>> struct __type_at_large_impl;
+
+template <size_t _I, size_t _Idx, class _T>
+struct __type_at_entry { };
+
+template <class _Result>
+struct __type_at_assign_op_ignore_rest {
+ template <class _T>
+ __type_at_assign_op_ignore_rest<_Result> operator=(_T&&);
+ using type = _Result;
+};
+
+struct __type_at_assign_op_impl {
+ template <size_t _I, size_t _Idx, class _T>
+ __type_at_assign_op_impl operator=(__type_at_entry<_I, _Idx, _T>&&);
+ template <size_t _I, class _T>
+ __type_at_assign_op_ignore_rest<_T> operator=(__type_at_entry<_I, _I, _T>&&);
+};
+
+template <size_t _I, class... _Ts, size_t... _Idxs>
+struct __type_at_large_impl<_I, __type_list<_Ts...>, std::integer_sequence<size_t, _Idxs...>>
+ : decltype(
+ _MDSPAN_FOLD_ASSIGN_LEFT(__type_at_assign_op_impl{}, /* = ... = */ __type_at_entry<_I, _Idxs, _Ts>{})
+ )
+{ };
+
+template <size_t _I, class... _Ts>
+struct __type_at<_I, __type_list<_Ts...>>
+ : __type_at_large_impl<_I, __type_list<_Ts...>>
+{ };
+
+template <class _T0, class... _Ts>
+struct __type_at<0, __type_list<_T0, _Ts...>> {
+ using type = _T0;
+};
+
+template <class _T0, class _T1, class... _Ts>
+struct __type_at<1, __type_list<_T0, _T1, _Ts...>> {
+ using type = _T1;
+};
+
+template <class _T0, class _T1, class _T2, class... _Ts>
+struct __type_at<2, __type_list<_T0, _T1, _T2, _Ts...>> {
+ using type = _T2;
+};
+
+template <class _T0, class _T1, class _T2, class _T3, class... _Ts>
+struct __type_at<3, __type_list<_T0, _T1, _T2, _T3, _Ts...>> {
+ using type = _T3;
+};
+
+
+} // namespace detail
+
+//==============================================================================
+
+} // end namespace MDSPAN_IMPL_STANDARD_NAMESPACE
+
diff --git a/ext/mdspan/include/experimental/__p0009_bits/utility.hpp b/ext/mdspan/include/experimental/__p0009_bits/utility.hpp
new file mode 100644
index 0000000..a078eeb
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p0009_bits/utility.hpp
@@ -0,0 +1,172 @@
+#pragma once
+
+#include <cstddef>
+#include <type_traits>
+#include <array>
+#include <utility>
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+namespace detail {
+
+// type alias used for rank-based tag dispatch
+//
+// this is used to enable alternatives to constexpr if when building for C++14
+//
+template <std::size_t N>
+using with_rank = std::integral_constant<std::size_t, N>;
+
+template <class I1, class I2>
+MDSPAN_INLINE_FUNCTION
+constexpr bool common_integral_compare(I1 x, I2 y)
+{
+ static_assert(std::is_integral<I1>::value &&
+ std::is_integral<I2>::value, "");
+
+ using I = std::common_type_t<I1, I2>;
+ return static_cast<I>(x) == static_cast<I>(y);
+}
+
+template <class T1, class T2, class F>
+MDSPAN_INLINE_FUNCTION
+constexpr bool rankwise_equal(with_rank<0>, const T1&, const T2&, F)
+{
+ return true;
+}
+
+template <std::size_t N, class T1, class T2, class F>
+MDSPAN_INLINE_FUNCTION
+constexpr bool rankwise_equal(with_rank<N>, const T1& x, const T2& y, F func)
+{
+ bool match = true;
+
+ for (std::size_t r = 0; r < N; r++) {
+ match = match && common_integral_compare(func(x, r), func(y, r));
+ }
+
+ return match;
+}
+
+constexpr struct
+{
+ template <class T, class I>
+ MDSPAN_INLINE_FUNCTION
+ constexpr auto operator()(const T& x, I i) const
+ {
+ return x.extent(i);
+ }
+} extent;
+
+constexpr struct
+{
+ template <class T, class I>
+ MDSPAN_INLINE_FUNCTION
+ constexpr auto operator()(const T& x, I i) const
+ {
+ return x.stride(i);
+ }
+} stride;
+
+// same as std::integral_constant but with __host__ __device__ annotations on
+// the implicit conversion function and the call operator
+template <class T, T v>
+struct integral_constant {
+ using value_type = T;
+ using type = integral_constant<T, v>;
+
+ static constexpr T value = v;
+
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ constexpr integral_constant() = default;
+
+ // These interop functions work, because other than the value_type operator
+ // everything of std::integral_constant works on device (defaulted functions)
+ MDSPAN_FUNCTION
+ constexpr integral_constant(std::integral_constant<T,v>) {};
+
+ MDSPAN_FUNCTION constexpr operator std::integral_constant<T,v>() const noexcept {
+ return std::integral_constant<T,v>{};
+ }
+
+ MDSPAN_FUNCTION constexpr operator value_type() const noexcept {
+ return value;
+ }
+
+ MDSPAN_FUNCTION constexpr value_type operator()() const noexcept {
+ return value;
+ }
+};
+
+// The tuple implementation only comes in play when using capabilities
+// such as submdspan which require C++17 anyway
+#if MDSPAN_HAS_CXX_17
+template<class T, size_t Idx>
+struct tuple_member {
+ using type = T;
+ static constexpr size_t idx = Idx;
+ T val;
+ MDSPAN_FUNCTION constexpr T& get() { return val; }
+ MDSPAN_FUNCTION constexpr const T& get() const { return val; }
+};
+
+// A helper class which will be used via a fold expression to
+// select the type with the correct Idx in a pack of tuple_member
+template<size_t SearchIdx, size_t Idx, class T>
+struct tuple_idx_matcher {
+ using type = tuple_member<T, Idx>;
+ template<class Other>
+ MDSPAN_FUNCTION
+ constexpr auto operator | ([[maybe_unused]] Other v) const {
+ if constexpr (Idx == SearchIdx) { return *this; }
+ else { return v; }
+ }
+};
+
+template<class IdxSeq, class ... Elements>
+struct tuple_impl;
+
+template<size_t ... Idx, class ... Elements>
+struct tuple_impl<std::index_sequence<Idx...>, Elements...>: public tuple_member<Elements, Idx> ... {
+
+ MDSPAN_FUNCTION
+ constexpr tuple_impl(Elements ... vals):tuple_member<Elements, Idx>{vals}... {}
+
+ template<size_t N>
+ MDSPAN_FUNCTION
+ constexpr auto& get() {
+ using base_t = decltype((tuple_idx_matcher<N, Idx, Elements>() | ...) );
+ return base_t::type::get();
+ }
+ template<size_t N>
+ MDSPAN_FUNCTION
+ constexpr const auto& get() const {
+ using base_t = decltype((tuple_idx_matcher<N, Idx, Elements>() | ...) );
+ return base_t::type::get();
+ }
+};
+
+// A simple tuple-like class for representing slices internally and is compatible with device code
+// This doesn't support type access since we don't need it
+// This is not meant as an external API
+template<class ... Elements>
+struct tuple: public tuple_impl<decltype(std::make_index_sequence<sizeof...(Elements)>()), Elements...> {
+ MDSPAN_FUNCTION
+ constexpr tuple(Elements ... vals):tuple_impl<decltype(std::make_index_sequence<sizeof...(Elements)>()), Elements ...>(vals ...) {}
+};
+
+template<size_t Idx, class ... Args>
+MDSPAN_FUNCTION
+constexpr auto& get(tuple<Args...>& vals) { return vals.template get<Idx>(); }
+
+template<size_t Idx, class ... Args>
+MDSPAN_FUNCTION
+constexpr const auto& get(const tuple<Args...>& vals) { return vals.template get<Idx>(); }
+
+template<class ... Elements>
+tuple(Elements ...) -> tuple<Elements...>;
+#endif
+} // namespace detail
+
+constexpr struct mdspan_non_standard_tag {
+} mdspan_non_standard;
+
+} // namespace MDSPAN_IMPL_STANDARD_NAMESPACE
diff --git a/ext/mdspan/include/experimental/__p1684_bits/mdarray.hpp b/ext/mdspan/include/experimental/__p1684_bits/mdarray.hpp
new file mode 100644
index 0000000..bdc5925
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p1684_bits/mdarray.hpp
@@ -0,0 +1,460 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+
+#pragma once
+
+#include "../mdspan"
+#include <cassert>
+#include <vector>
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+namespace MDSPAN_IMPL_PROPOSED_NAMESPACE {
+
+namespace {
+ template<class Extents>
+ struct size_of_extents;
+
+ template<class IndexType, size_t ... Extents>
+ struct size_of_extents<extents<IndexType, Extents...>> {
+ constexpr static size_t value() {
+ size_t size = 1;
+ for(size_t r=0; r<extents<IndexType, Extents...>::rank(); r++)
+ size *= extents<IndexType, Extents...>::static_extent(r);
+ return size;
+ }
+ };
+}
+
+namespace {
+ template<class C>
+ struct container_is_array : std::false_type {
+ template<class M>
+ static constexpr C construct(const M& m) { return C(m.required_span_size()); }
+ };
+ template<class T, size_t N>
+ struct container_is_array<std::array<T,N>> : std::true_type {
+ template<class M>
+ static constexpr std::array<T,N> construct(const M&) { return std::array<T,N>(); }
+ };
+}
+
+template <
+ class ElementType,
+ class Extents,
+ class LayoutPolicy = layout_right,
+ class Container = std::vector<ElementType>
+>
+class mdarray {
+private:
+ static_assert(::MDSPAN_IMPL_STANDARD_NAMESPACE::detail::__is_extents_v<Extents>,
+ MDSPAN_IMPL_PROPOSED_NAMESPACE_STRING "::mdspan's Extents template parameter must be a specialization of " MDSPAN_IMPL_STANDARD_NAMESPACE_STRING "::extents.");
+
+public:
+
+ //--------------------------------------------------------------------------------
+ // Domain and codomain types
+
+ using extents_type = Extents;
+ using layout_type = LayoutPolicy;
+ using container_type = Container;
+ using mapping_type = typename layout_type::template mapping<extents_type>;
+ using element_type = ElementType;
+ using mdspan_type = mdspan<element_type, extents_type, layout_type>;
+ using const_mdspan_type = mdspan<const element_type, extents_type, layout_type>;
+ using value_type = std::remove_cv_t<element_type>;
+ using index_type = typename Extents::index_type;
+ using size_type = typename Extents::size_type;
+ using rank_type = typename Extents::rank_type;
+ using pointer = typename container_type::pointer;
+ using reference = typename container_type::reference;
+ using const_pointer = typename container_type::const_pointer;
+ using const_reference = typename container_type::const_reference;
+
+public:
+
+ //--------------------------------------------------------------------------------
+ // [mdspan.basic.cons], mdspan constructors, assignment, and destructor
+
+#if !(MDSPAN_HAS_CXX_20)
+ MDSPAN_FUNCTION_REQUIRES(
+ (MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr),
+ mdarray, (), ,
+ /* requires */ (extents_type::rank_dynamic()!=0)) {}
+#else
+ MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr mdarray() requires(extents_type::rank_dynamic()!=0) = default;
+#endif
+ MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr mdarray(const mdarray&) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr mdarray(mdarray&&) = default;
+
+ // Constructors for container types constructible from a size
+ MDSPAN_TEMPLATE_REQUIRES(
+ class... SizeTypes,
+ /* requires */ (
+ (::MDSPAN_IMPL_STANDARD_NAMESPACE::detail::are_valid_indices<index_type, SizeTypes...>()) &&
+ _MDSPAN_TRAIT( std::is_constructible, extents_type, SizeTypes...) &&
+ _MDSPAN_TRAIT( std::is_constructible, mapping_type, extents_type) &&
+ (_MDSPAN_TRAIT( std::is_constructible, container_type, size_t) ||
+ container_is_array<container_type>::value) &&
+ (extents_type::rank()>0 || extents_type::rank_dynamic()==0)
+ )
+ )
+ MDSPAN_INLINE_FUNCTION
+ explicit constexpr mdarray(SizeTypes... dynamic_extents)
+ : map_(extents_type(dynamic_extents...)), ctr_(container_is_array<container_type>::construct(map_))
+ { }
+
+ MDSPAN_FUNCTION_REQUIRES(
+ (MDSPAN_INLINE_FUNCTION constexpr),
+ mdarray, (const extents_type& exts), ,
+ /* requires */ ((_MDSPAN_TRAIT( std::is_constructible, container_type, size_t) ||
+ container_is_array<container_type>::value) &&
+ _MDSPAN_TRAIT( std::is_constructible, mapping_type, extents_type))
+ ) : map_(exts), ctr_(container_is_array<container_type>::construct(map_))
+ { }
+
+ MDSPAN_FUNCTION_REQUIRES(
+ (MDSPAN_INLINE_FUNCTION constexpr),
+ mdarray, (const mapping_type& m), ,
+ /* requires */ (_MDSPAN_TRAIT( std::is_constructible, container_type, size_t) ||
+ container_is_array<container_type>::value)
+ ) : map_(m), ctr_(container_is_array<container_type>::construct(map_))
+ { }
+
+ MDSPAN_FUNCTION_REQUIRES(
+ (MDSPAN_INLINE_FUNCTION constexpr),
+ mdarray, (const extents_type& exts, const container_type& ctr), ,
+ /* requires */ (_MDSPAN_TRAIT( std::is_constructible, mapping_type, extents_type))
+ ) : map_(exts), ctr_(ctr)
+ { assert(ctr.size() >= static_cast<size_t>(map_.required_span_size())); }
+
+ constexpr mdarray(const mapping_type& m, const container_type& ctr)
+ : map_(m), ctr_(ctr)
+ { assert(ctr.size() >= static_cast<size_t>(map_.required_span_size())); }
+
+ MDSPAN_FUNCTION_REQUIRES(
+ (MDSPAN_INLINE_FUNCTION constexpr),
+ mdarray, (const extents_type& exts, container_type&& ctr), ,
+ /* requires */ (_MDSPAN_TRAIT( std::is_constructible, mapping_type, extents_type))
+ ) : map_(exts), ctr_(std::move(ctr))
+ { assert(ctr_.size() >= static_cast<size_t>(map_.required_span_size())); }
+
+ constexpr mdarray(const mapping_type& m, container_type&& ctr)
+ : map_(m), ctr_(std::move(ctr))
+ { assert(ctr_.size() >= static_cast<size_t>(map_.required_span_size())); }
+
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherElementType, class OtherExtents, class OtherLayoutPolicy, class OtherContainer,
+ /* requires */ (
+ _MDSPAN_TRAIT( std::is_constructible, mapping_type, typename OtherLayoutPolicy::template mapping<OtherExtents>) &&
+ _MDSPAN_TRAIT( std::is_constructible, container_type, OtherContainer)
+ )
+ )
+ MDSPAN_INLINE_FUNCTION
+ constexpr mdarray(const mdarray<OtherElementType, OtherExtents, OtherLayoutPolicy, OtherContainer>& other)
+ : map_(other.mapping()), ctr_(other.container())
+ {
+ static_assert( std::is_constructible<extents_type, OtherExtents>::value, "");
+ }
+
+ // Constructors for container types constructible from a size and allocator
+ MDSPAN_TEMPLATE_REQUIRES(
+ class Alloc,
+ /* requires */ (_MDSPAN_TRAIT( std::is_constructible, container_type, size_t, Alloc) &&
+ _MDSPAN_TRAIT( std::is_constructible, mapping_type, extents_type))
+ )
+ MDSPAN_INLINE_FUNCTION
+ constexpr mdarray(const extents_type& exts, const Alloc& a)
+ : map_(exts), ctr_(map_.required_span_size(), a)
+ { }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class Alloc,
+ /* requires */ (_MDSPAN_TRAIT( std::is_constructible, container_type, size_t, Alloc))
+ )
+ MDSPAN_INLINE_FUNCTION
+ constexpr mdarray(const mapping_type& map, const Alloc& a)
+ : map_(map), ctr_(map_.required_span_size(), a)
+ { }
+
+ // Constructors for container types constructible from a container and allocator
+ MDSPAN_TEMPLATE_REQUIRES(
+ class Alloc,
+ /* requires */ (_MDSPAN_TRAIT( std::is_constructible, container_type, container_type, Alloc) &&
+ _MDSPAN_TRAIT( std::is_constructible, mapping_type, extents_type))
+ )
+ MDSPAN_INLINE_FUNCTION
+ constexpr mdarray(const extents_type& exts, const container_type& ctr, const Alloc& a)
+ : map_(exts), ctr_(ctr, a)
+ { assert(ctr_.size() >= static_cast<size_t>(map_.required_span_size())); }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class Alloc,
+ /* requires */ (_MDSPAN_TRAIT( std::is_constructible, container_type, size_t, Alloc))
+ )
+ MDSPAN_INLINE_FUNCTION
+ constexpr mdarray(const mapping_type& map, const container_type& ctr, const Alloc& a)
+ : map_(map), ctr_(ctr, a)
+ { assert(ctr_.size() >= static_cast<size_t>(map_.required_span_size())); }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class Alloc,
+ /* requires */ (_MDSPAN_TRAIT( std::is_constructible, container_type, container_type, Alloc) &&
+ _MDSPAN_TRAIT( std::is_constructible, mapping_type, extents_type))
+ )
+ MDSPAN_INLINE_FUNCTION
+ constexpr mdarray(const extents_type& exts, container_type&& ctr, const Alloc& a)
+ : map_(exts), ctr_(std::move(ctr), a)
+ { assert(ctr_.size() >= static_cast<size_t>(map_.required_span_size())); }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class Alloc,
+ /* requires */ (_MDSPAN_TRAIT( std::is_constructible, container_type, size_t, Alloc))
+ )
+ MDSPAN_INLINE_FUNCTION
+ constexpr mdarray(const mapping_type& map, container_type&& ctr, const Alloc& a)
+ : map_(map), ctr_(std::move(ctr), a)
+ { assert(ctr_.size() >= map_.required_span_size()); }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherElementType, class OtherExtents, class OtherLayoutPolicy, class OtherContainer, class Alloc,
+ /* requires */ (
+ _MDSPAN_TRAIT( std::is_constructible, mapping_type, typename OtherLayoutPolicy::template mapping<OtherExtents>) &&
+ _MDSPAN_TRAIT( std::is_constructible, container_type, OtherContainer, Alloc)
+ )
+ )
+ MDSPAN_INLINE_FUNCTION
+ constexpr mdarray(const mdarray<OtherElementType, OtherExtents, OtherLayoutPolicy, OtherContainer>& other, const Alloc& a)
+ : map_(other.mapping()), ctr_(other.container(), a)
+ {
+ static_assert( std::is_constructible<extents_type, OtherExtents>::value, "");
+ }
+
+ MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr mdarray& operator= (const mdarray&) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr mdarray& operator= (mdarray&&) = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ ~mdarray() = default;
+
+ //--------------------------------------------------------------------------------
+ // [mdspan.basic.mapping], mdspan mapping domain multidimensional index to access codomain element
+
+ #if MDSPAN_USE_BRACKET_OPERATOR
+ MDSPAN_TEMPLATE_REQUIRES(
+ class... SizeTypes,
+ /* requires */ (
+ _MDSPAN_FOLD_AND(_MDSPAN_TRAIT( std::is_convertible, SizeTypes, index_type) /* && ... */) &&
+ extents_type::rank() == sizeof...(SizeTypes)
+ )
+ )
+ MDSPAN_FORCE_INLINE_FUNCTION
+ constexpr const_reference operator[](SizeTypes... indices) const noexcept
+ {
+ return ctr_[map_(static_cast<index_type>(std::move(indices))...)];
+ }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class... SizeTypes,
+ /* requires */ (
+ _MDSPAN_FOLD_AND(_MDSPAN_TRAIT( std::is_convertible, SizeTypes, index_type) /* && ... */) &&
+ extents_type::rank() == sizeof...(SizeTypes)
+ )
+ )
+ MDSPAN_FORCE_INLINE_FUNCTION
+ constexpr reference operator[](SizeTypes... indices) noexcept
+ {
+ return ctr_[map_(static_cast<index_type>(std::move(indices))...)];
+ }
+ #endif
+
+#if 0
+ MDSPAN_TEMPLATE_REQUIRES(
+ class SizeType, size_t N,
+ /* requires */ (
+ _MDSPAN_TRAIT( std::is_convertible, SizeType, index_type) &&
+ N == extents_type::rank()
+ )
+ )
+ MDSPAN_FORCE_INLINE_FUNCTION
+ constexpr const_reference operator[](const std::array<SizeType, N>& indices) const noexcept
+ {
+ return __impl::template __callop<reference>(*this, indices);
+ }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class SizeType, size_t N,
+ /* requires */ (
+ _MDSPAN_TRAIT( std::is_convertible, SizeType, index_type) &&
+ N == extents_type::rank()
+ )
+ )
+ MDSPAN_FORCE_INLINE_FUNCTION
+ constexpr reference operator[](const std::array<SizeType, N>& indices) noexcept
+ {
+ return __impl::template __callop<reference>(*this, indices);
+ }
+#endif
+
+
+ #if MDSPAN_USE_PAREN_OPERATOR
+ MDSPAN_TEMPLATE_REQUIRES(
+ class... SizeTypes,
+ /* requires */ (
+ (::MDSPAN_IMPL_STANDARD_NAMESPACE::detail::are_valid_indices<index_type, SizeTypes...>()) &&
+ extents_type::rank() == sizeof...(SizeTypes)
+ )
+ )
+ MDSPAN_FORCE_INLINE_FUNCTION
+ constexpr const_reference operator()(SizeTypes... indices) const noexcept
+ {
+ return ctr_[map_(static_cast<index_type>(std::move(indices))...)];
+ }
+ MDSPAN_TEMPLATE_REQUIRES(
+ class... SizeTypes,
+ /* requires */ (
+ (::MDSPAN_IMPL_STANDARD_NAMESPACE::detail::are_valid_indices<index_type, SizeTypes...>()) &&
+ extents_type::rank() == sizeof...(SizeTypes)
+ )
+ )
+ MDSPAN_FORCE_INLINE_FUNCTION
+ constexpr reference operator()(SizeTypes... indices) noexcept
+ {
+ return ctr_[map_(static_cast<index_type>(std::move(indices))...)];
+ }
+
+#if 0
+ MDSPAN_TEMPLATE_REQUIRES(
+ class SizeType, size_t N,
+ /* requires */ (
+ _MDSPAN_TRAIT( std::is_convertible, SizeType, index_type) &&
+ N == extents_type::rank()
+ )
+ )
+ MDSPAN_FORCE_INLINE_FUNCTION
+ constexpr const_reference operator()(const std::array<SizeType, N>& indices) const noexcept
+ {
+ return __impl::template __callop<reference>(*this, indices);
+ }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class SizeType, size_t N,
+ /* requires */ (
+ _MDSPAN_TRAIT( std::is_convertible, SizeType, index_type) &&
+ N == extents_type::rank()
+ )
+ )
+ MDSPAN_FORCE_INLINE_FUNCTION
+ constexpr reference operator()(const std::array<SizeType, N>& indices) noexcept
+ {
+ return __impl::template __callop<reference>(*this, indices);
+ }
+#endif
+ #endif
+
+ MDSPAN_INLINE_FUNCTION constexpr pointer data() noexcept { return ctr_.data(); };
+ MDSPAN_INLINE_FUNCTION constexpr const_pointer data() const noexcept { return ctr_.data(); };
+ MDSPAN_INLINE_FUNCTION constexpr container_type& container() noexcept { return ctr_; };
+ MDSPAN_INLINE_FUNCTION constexpr const container_type& container() const noexcept { return ctr_; };
+
+ //--------------------------------------------------------------------------------
+ // [mdspan.basic.domobs], mdspan observers of the domain multidimensional index space
+
+ MDSPAN_INLINE_FUNCTION static constexpr rank_type rank() noexcept { return extents_type::rank(); }
+ MDSPAN_INLINE_FUNCTION static constexpr rank_type rank_dynamic() noexcept { return extents_type::rank_dynamic(); }
+ MDSPAN_INLINE_FUNCTION static constexpr size_t static_extent(size_t r) noexcept { return extents_type::static_extent(r); }
+
+ MDSPAN_INLINE_FUNCTION constexpr const extents_type& extents() const noexcept { return map_.extents(); };
+ MDSPAN_INLINE_FUNCTION constexpr index_type extent(size_t r) const noexcept { return map_.extents().extent(r); };
+ MDSPAN_INLINE_FUNCTION constexpr index_type size() const noexcept {
+// return __impl::__size(*this);
+ return ctr_.size();
+ };
+
+
+ //--------------------------------------------------------------------------------
+ // [mdspan.basic.obs], mdspan observers of the mapping
+
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_unique() noexcept { return mapping_type::is_always_unique(); };
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_exhaustive() noexcept { return mapping_type::is_always_exhaustive(); };
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_strided() noexcept { return mapping_type::is_always_strided(); };
+
+ MDSPAN_INLINE_FUNCTION constexpr const mapping_type& mapping() const noexcept { return map_; };
+ MDSPAN_INLINE_FUNCTION constexpr bool is_unique() const noexcept { return map_.is_unique(); };
+ MDSPAN_INLINE_FUNCTION constexpr bool is_exhaustive() const noexcept { return map_.is_exhaustive(); };
+ MDSPAN_INLINE_FUNCTION constexpr bool is_strided() const noexcept { return map_.is_strided(); };
+ MDSPAN_INLINE_FUNCTION constexpr index_type stride(size_t r) const { return map_.stride(r); };
+
+ // Converstion to mdspan
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherElementType, class OtherExtents,
+ class OtherLayoutType, class OtherAccessorType,
+ /* requires */ (
+ _MDSPAN_TRAIT(std::is_assignable,
+ mdspan<OtherElementType, OtherExtents, OtherLayoutType, OtherAccessorType>,
+ mdspan_type)
+ )
+ )
+ constexpr operator mdspan<OtherElementType, OtherExtents, OtherLayoutType, OtherAccessorType> () {
+ return mdspan_type(data(), map_);
+ }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherElementType, class OtherExtents,
+ class OtherLayoutType, class OtherAccessorType,
+ /* requires */ (
+ _MDSPAN_TRAIT(std::is_assignable,
+ mdspan<OtherElementType, OtherExtents, OtherLayoutType, OtherAccessorType>,
+ const_mdspan_type)
+ )
+ )
+ constexpr operator mdspan<OtherElementType, OtherExtents, OtherLayoutType, OtherAccessorType> () const {
+ return const_mdspan_type(data(), map_);
+ }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherAccessorType = default_accessor<element_type>,
+ /* requires */ (
+ _MDSPAN_TRAIT(std::is_assignable, mdspan_type,
+ mdspan<element_type, extents_type, layout_type, OtherAccessorType>)
+ )
+ )
+ constexpr mdspan<element_type, extents_type, layout_type, OtherAccessorType>
+ to_mdspan(const OtherAccessorType& a = default_accessor<element_type>()) {
+ return mdspan<element_type, extents_type, layout_type, OtherAccessorType>(data(), map_, a);
+ }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class OtherAccessorType = default_accessor<const element_type>,
+ /* requires */ (
+ _MDSPAN_TRAIT(std::is_assignable, const_mdspan_type,
+ mdspan<const element_type, extents_type, layout_type, OtherAccessorType>)
+ )
+ )
+ constexpr mdspan<const element_type, extents_type, layout_type, OtherAccessorType>
+ to_mdspan(const OtherAccessorType& a = default_accessor<const element_type>()) const {
+ return mdspan<const element_type, extents_type, layout_type, OtherAccessorType>(data(), map_, a);
+ }
+
+private:
+ mapping_type map_;
+ container_type ctr_;
+
+ template <class, class, class, class>
+ friend class mdarray;
+};
+
+
+} // end namespace MDSPAN_IMPL_PROPOSED_NAMESPACE
+} // end namespace MDSPAN_IMPL_STANDARD_NAMESPACE
diff --git a/ext/mdspan/include/experimental/__p2389_bits/dims.hpp b/ext/mdspan/include/experimental/__p2389_bits/dims.hpp
new file mode 100644
index 0000000..0004521
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p2389_bits/dims.hpp
@@ -0,0 +1,28 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+
+#pragma once
+
+// backward compatibility import into experimental
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+namespace MDSPAN_IMPL_PROPOSED_NAMESPACE {
+
+template< ::std::size_t Rank, class IndexType = std::size_t>
+using dims =
+ :: MDSPAN_IMPL_STANDARD_NAMESPACE :: dextents<IndexType, Rank>;
+
+} // namespace MDSPAN_IMPL_PROPOSED_NAMESPACE
+} // namespace MDSPAN_IMPL_STANDARD_NAMESPACE
diff --git a/ext/mdspan/include/experimental/__p2630_bits/strided_slice.hpp b/ext/mdspan/include/experimental/__p2630_bits/strided_slice.hpp
new file mode 100644
index 0000000..89ba820
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p2630_bits/strided_slice.hpp
@@ -0,0 +1,48 @@
+
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+
+#pragma once
+
+#include <type_traits>
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+
+namespace {
+ template<class T>
+ struct __mdspan_is_integral_constant: std::false_type {};
+
+ template<class T, T val>
+ struct __mdspan_is_integral_constant<std::integral_constant<T,val>>: std::true_type {};
+}
+
+// Slice Specifier allowing for strides and compile time extent
+template <class OffsetType, class ExtentType, class StrideType>
+struct strided_slice {
+ using offset_type = OffsetType;
+ using extent_type = ExtentType;
+ using stride_type = StrideType;
+
+ _MDSPAN_NO_UNIQUE_ADDRESS OffsetType offset{};
+ _MDSPAN_NO_UNIQUE_ADDRESS ExtentType extent{};
+ _MDSPAN_NO_UNIQUE_ADDRESS StrideType stride{};
+
+ static_assert(std::is_integral_v<OffsetType> || __mdspan_is_integral_constant<OffsetType>::value);
+ static_assert(std::is_integral_v<ExtentType> || __mdspan_is_integral_constant<ExtentType>::value);
+ static_assert(std::is_integral_v<StrideType> || __mdspan_is_integral_constant<StrideType>::value);
+};
+
+} // MDSPAN_IMPL_STANDARD_NAMESPACE
diff --git a/ext/mdspan/include/experimental/__p2630_bits/submdspan.hpp b/ext/mdspan/include/experimental/__p2630_bits/submdspan.hpp
new file mode 100644
index 0000000..abddd0b
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p2630_bits/submdspan.hpp
@@ -0,0 +1,40 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+
+#pragma once
+
+#include "submdspan_extents.hpp"
+#include "submdspan_mapping.hpp"
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+template <class ElementType, class Extents, class LayoutPolicy,
+ class AccessorPolicy, class... SliceSpecifiers>
+MDSPAN_INLINE_FUNCTION
+constexpr auto
+submdspan(const mdspan<ElementType, Extents, LayoutPolicy, AccessorPolicy> &src,
+ SliceSpecifiers... slices) {
+ const auto sub_submdspan_mapping_result = submdspan_mapping(src.mapping(), slices...);
+ // NVCC has a problem with the deduction so lets figure out the type
+ using sub_mapping_t = std::remove_cv_t<decltype(sub_submdspan_mapping_result.mapping)>;
+ using sub_extents_t = typename sub_mapping_t::extents_type;
+ using sub_layout_t = typename sub_mapping_t::layout_type;
+ using sub_accessor_t = typename AccessorPolicy::offset_policy;
+ return mdspan<ElementType, sub_extents_t, sub_layout_t, sub_accessor_t>(
+ src.accessor().offset(src.data_handle(), sub_submdspan_mapping_result.offset),
+ sub_submdspan_mapping_result.mapping,
+ sub_accessor_t(src.accessor()));
+}
+} // namespace MDSPAN_IMPL_STANDARD_NAMESPACE
diff --git a/ext/mdspan/include/experimental/__p2630_bits/submdspan_extents.hpp b/ext/mdspan/include/experimental/__p2630_bits/submdspan_extents.hpp
new file mode 100644
index 0000000..4fe5dc6
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p2630_bits/submdspan_extents.hpp
@@ -0,0 +1,418 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+
+#pragma once
+
+#include <complex>
+
+#include "strided_slice.hpp"
+#include "../__p0009_bits/utility.hpp"
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+namespace detail {
+
+// Mapping from submapping ranks to srcmapping ranks
+// InvMapRank is an index_sequence, which we build recursively
+// to contain the mapped indices.
+// end of recursion specialization containing the final index_sequence
+template <size_t Counter, size_t... MapIdxs>
+MDSPAN_INLINE_FUNCTION
+constexpr auto inv_map_rank(std::integral_constant<size_t, Counter>, std::index_sequence<MapIdxs...>) {
+ return std::index_sequence<MapIdxs...>();
+}
+
+// specialization reducing rank by one (i.e., integral slice specifier)
+template<size_t Counter, class Slice, class... SliceSpecifiers, size_t... MapIdxs>
+MDSPAN_INLINE_FUNCTION
+constexpr auto inv_map_rank(std::integral_constant<size_t, Counter>, std::index_sequence<MapIdxs...>, Slice,
+ SliceSpecifiers... slices) {
+ using next_idx_seq_t = std::conditional_t<std::is_convertible_v<Slice, size_t>,
+ std::index_sequence<MapIdxs...>,
+ std::index_sequence<MapIdxs..., Counter>>;
+
+ return inv_map_rank(std::integral_constant<size_t,Counter + 1>(), next_idx_seq_t(),
+ slices...);
+}
+
+// Helper for identifying strided_slice
+template <class T> struct is_strided_slice : std::false_type {};
+
+template <class OffsetType, class ExtentType, class StrideType>
+struct is_strided_slice<
+ strided_slice<OffsetType, ExtentType, StrideType>> : std::true_type {};
+
+// Helper for identifying valid pair like things
+template <class T, class IndexType> struct index_pair_like : std::false_type {};
+
+template <class IdxT1, class IdxT2, class IndexType>
+struct index_pair_like<std::pair<IdxT1, IdxT2>, IndexType> {
+ static constexpr bool value = std::is_convertible_v<IdxT1, IndexType> &&
+ std::is_convertible_v<IdxT2, IndexType>;
+};
+
+template <class IdxT1, class IdxT2, class IndexType>
+struct index_pair_like<std::tuple<IdxT1, IdxT2>, IndexType> {
+ static constexpr bool value = std::is_convertible_v<IdxT1, IndexType> &&
+ std::is_convertible_v<IdxT2, IndexType>;
+};
+
+template <class IdxT1, class IdxT2, class IndexType>
+struct index_pair_like<tuple<IdxT1, IdxT2>, IndexType> {
+ static constexpr bool value = std::is_convertible_v<IdxT1, IndexType> &&
+ std::is_convertible_v<IdxT2, IndexType>;
+};
+
+template <class IdxT, class IndexType>
+struct index_pair_like<std::complex<IdxT>, IndexType> {
+ static constexpr bool value = std::is_convertible_v<IdxT, IndexType>;
+};
+
+template <class IdxT, class IndexType>
+struct index_pair_like<std::array<IdxT, 2>, IndexType> {
+ static constexpr bool value = std::is_convertible_v<IdxT, IndexType>;
+};
+
+// first_of(slice): getting begin of slice specifier range
+MDSPAN_TEMPLATE_REQUIRES(
+ class Integral,
+ /* requires */(std::is_convertible_v<Integral, size_t>)
+)
+MDSPAN_INLINE_FUNCTION
+constexpr Integral first_of(const Integral &i) {
+ return i;
+}
+
+template<class Integral, Integral v>
+MDSPAN_INLINE_FUNCTION
+constexpr Integral first_of(const std::integral_constant<Integral, v>&) {
+ return integral_constant<Integral, v>();
+}
+
+MDSPAN_INLINE_FUNCTION
+constexpr integral_constant<size_t, 0>
+first_of(const ::MDSPAN_IMPL_STANDARD_NAMESPACE::full_extent_t &) {
+ return integral_constant<size_t, 0>();
+}
+
+MDSPAN_TEMPLATE_REQUIRES(
+ class Slice,
+ /* requires */(index_pair_like<Slice, size_t>::value)
+)
+MDSPAN_INLINE_FUNCTION
+constexpr auto first_of(const Slice &i) {
+ return get<0>(i);
+}
+
+MDSPAN_TEMPLATE_REQUIRES(
+ class IdxT1, class IdxT2,
+ /* requires */ (index_pair_like<std::tuple<IdxT1, IdxT2>, size_t>::value)
+ )
+constexpr auto first_of(const std::tuple<IdxT1, IdxT2>& i) {
+ return get<0>(i);
+}
+
+MDSPAN_TEMPLATE_REQUIRES(
+ class IdxT1, class IdxT2,
+ /* requires */ (index_pair_like<std::pair<IdxT1, IdxT2>, size_t>::value)
+ )
+MDSPAN_INLINE_FUNCTION
+constexpr auto first_of(const std::pair<IdxT1, IdxT2>& i) {
+ return i.first;
+}
+
+template<class T>
+MDSPAN_INLINE_FUNCTION
+constexpr auto first_of(const std::complex<T> &i) {
+ return i.real();
+}
+
+template <class OffsetType, class ExtentType, class StrideType>
+MDSPAN_INLINE_FUNCTION
+constexpr OffsetType
+first_of(const strided_slice<OffsetType, ExtentType, StrideType> &r) {
+ return r.offset;
+}
+
+// last_of(slice): getting end of slice specifier range
+// We need however not just the slice but also the extents
+// of the original view and which rank from the extents.
+// This is needed in the case of slice being full_extent_t.
+MDSPAN_TEMPLATE_REQUIRES(
+ size_t k, class Extents, class Integral,
+ /* requires */(std::is_convertible_v<Integral, size_t>)
+)
+MDSPAN_INLINE_FUNCTION
+constexpr Integral
+ last_of(std::integral_constant<size_t, k>, const Extents &, const Integral &i) {
+ return i;
+}
+
+MDSPAN_TEMPLATE_REQUIRES(
+ size_t k, class Extents, class Slice,
+ /* requires */(index_pair_like<Slice, size_t>::value)
+)
+MDSPAN_INLINE_FUNCTION
+constexpr auto last_of(std::integral_constant<size_t, k>, const Extents &,
+ const Slice &i) {
+ return get<1>(i);
+}
+
+MDSPAN_TEMPLATE_REQUIRES(
+ size_t k, class Extents, class IdxT1, class IdxT2,
+ /* requires */ (index_pair_like<std::tuple<IdxT1, IdxT2>, size_t>::value)
+ )
+constexpr auto last_of(std::integral_constant<size_t, k>, const Extents &, const std::tuple<IdxT1, IdxT2>& i) {
+ return get<1>(i);
+}
+
+MDSPAN_TEMPLATE_REQUIRES(
+ size_t k, class Extents, class IdxT1, class IdxT2,
+ /* requires */ (index_pair_like<std::pair<IdxT1, IdxT2>, size_t>::value)
+ )
+MDSPAN_INLINE_FUNCTION
+constexpr auto last_of(std::integral_constant<size_t, k>, const Extents &, const std::pair<IdxT1, IdxT2>& i) {
+ return i.second;
+}
+
+template<size_t k, class Extents, class T>
+MDSPAN_INLINE_FUNCTION
+constexpr auto last_of(std::integral_constant<size_t, k>, const Extents &, const std::complex<T> &i) {
+ return i.imag();
+}
+
+// Suppress spurious warning with NVCC about no return statement.
+// This is a known issue in NVCC and NVC++
+// Depending on the CUDA and GCC version we need both the builtin
+// and the diagnostic push. I tried really hard to find something shorter
+// but no luck ...
+#if defined __NVCC__
+ #ifdef __NVCC_DIAG_PRAGMA_SUPPORT__
+ #pragma nv_diagnostic push
+ #pragma nv_diag_suppress = implicit_return_from_non_void_function
+ #else
+ #ifdef __CUDA_ARCH__
+ #pragma diagnostic push
+ #pragma diag_suppress implicit_return_from_non_void_function
+ #endif
+ #endif
+#elif defined __NVCOMPILER
+ #pragma diagnostic push
+ #pragma diag_suppress = implicit_return_from_non_void_function
+#endif
+template <size_t k, class Extents>
+MDSPAN_INLINE_FUNCTION
+constexpr auto last_of(std::integral_constant<size_t, k>, const Extents &ext,
+ ::MDSPAN_IMPL_STANDARD_NAMESPACE::full_extent_t) {
+ if constexpr (Extents::static_extent(k) == dynamic_extent) {
+ return ext.extent(k);
+ } else {
+ return integral_constant<size_t, Extents::static_extent(k)>();
+ }
+#if defined(__NVCC__) && !defined(__CUDA_ARCH__) && defined(__GNUC__)
+ // Even with CUDA_ARCH protection this thing warns about calling host function
+ __builtin_unreachable();
+#endif
+}
+#if defined __NVCC__
+ #ifdef __NVCC_DIAG_PRAGMA_SUPPORT__
+ #pragma nv_diagnostic pop
+ #else
+ #ifdef __CUDA_ARCH__
+ #pragma diagnostic pop
+ #endif
+ #endif
+#elif defined __NVCOMPILER
+ #pragma diagnostic pop
+#endif
+
+template <size_t k, class Extents, class OffsetType, class ExtentType,
+ class StrideType>
+MDSPAN_INLINE_FUNCTION
+constexpr OffsetType
+last_of(std::integral_constant<size_t, k>, const Extents &,
+ const strided_slice<OffsetType, ExtentType, StrideType> &r) {
+ return r.extent;
+}
+
+// get stride of slices
+template <class T>
+MDSPAN_INLINE_FUNCTION
+constexpr auto stride_of(const T &) {
+ return integral_constant<size_t, 1>();
+}
+
+template <class OffsetType, class ExtentType, class StrideType>
+MDSPAN_INLINE_FUNCTION
+constexpr auto
+stride_of(const strided_slice<OffsetType, ExtentType, StrideType> &r) {
+ return r.stride;
+}
+
+// divide which can deal with integral constant preservation
+template <class IndexT, class T0, class T1>
+MDSPAN_INLINE_FUNCTION
+constexpr auto divide(const T0 &v0, const T1 &v1) {
+ return IndexT(v0) / IndexT(v1);
+}
+
+template <class IndexT, class T0, T0 v0, class T1, T1 v1>
+MDSPAN_INLINE_FUNCTION
+constexpr auto divide(const std::integral_constant<T0, v0> &,
+ const std::integral_constant<T1, v1> &) {
+ // cutting short division by zero
+ // this is used for strided_slice with zero extent/stride
+ return integral_constant<IndexT, v0 == 0 ? 0 : v0 / v1>();
+}
+
+// multiply which can deal with integral constant preservation
+template <class IndexT, class T0, class T1>
+MDSPAN_INLINE_FUNCTION
+constexpr auto multiply(const T0 &v0, const T1 &v1) {
+ return IndexT(v0) * IndexT(v1);
+}
+
+template <class IndexT, class T0, T0 v0, class T1, T1 v1>
+MDSPAN_INLINE_FUNCTION
+constexpr auto multiply(const std::integral_constant<T0, v0> &,
+ const std::integral_constant<T1, v1> &) {
+ return integral_constant<IndexT, v0 * v1>();
+}
+
+// compute new static extent from range, preserving static knowledge
+template <class Arg0, class Arg1> struct StaticExtentFromRange {
+ constexpr static size_t value = dynamic_extent;
+};
+
+template <class Integral0, Integral0 val0, class Integral1, Integral1 val1>
+struct StaticExtentFromRange<std::integral_constant<Integral0, val0>,
+ std::integral_constant<Integral1, val1>> {
+ constexpr static size_t value = val1 - val0;
+};
+
+template <class Integral0, Integral0 val0, class Integral1, Integral1 val1>
+struct StaticExtentFromRange<integral_constant<Integral0, val0>,
+ integral_constant<Integral1, val1>> {
+ constexpr static size_t value = val1 - val0;
+};
+
+// compute new static extent from strided_slice, preserving static
+// knowledge
+template <class Arg0, class Arg1> struct StaticExtentFromStridedRange {
+ constexpr static size_t value = dynamic_extent;
+};
+
+template <class Integral0, Integral0 val0, class Integral1, Integral1 val1>
+struct StaticExtentFromStridedRange<std::integral_constant<Integral0, val0>,
+ std::integral_constant<Integral1, val1>> {
+ constexpr static size_t value = val0 > 0 ? 1 + (val0 - 1) / val1 : 0;
+};
+
+template <class Integral0, Integral0 val0, class Integral1, Integral1 val1>
+struct StaticExtentFromStridedRange<integral_constant<Integral0, val0>,
+ integral_constant<Integral1, val1>> {
+ constexpr static size_t value = val0 > 0 ? 1 + (val0 - 1) / val1 : 0;
+};
+
+// creates new extents through recursive calls to next_extent member function
+// next_extent has different overloads for different types of stride specifiers
+template <size_t K, class Extents, size_t... NewExtents>
+struct extents_constructor {
+ MDSPAN_TEMPLATE_REQUIRES(
+ class Slice, class... SlicesAndExtents,
+ /* requires */(!std::is_convertible_v<Slice, size_t> &&
+ !is_strided_slice<Slice>::value)
+ )
+ MDSPAN_INLINE_FUNCTION
+ constexpr static auto next_extent(const Extents &ext, const Slice &sl,
+ SlicesAndExtents... slices_and_extents) {
+ constexpr size_t new_static_extent = StaticExtentFromRange<
+ decltype(first_of(std::declval<Slice>())),
+ decltype(last_of(std::integral_constant<size_t, Extents::rank() - K>(),
+ std::declval<Extents>(),
+ std::declval<Slice>()))>::value;
+
+ using next_t =
+ extents_constructor<K - 1, Extents, NewExtents..., new_static_extent>;
+ using index_t = typename Extents::index_type;
+ return next_t::next_extent(
+ ext, slices_and_extents...,
+ index_t(last_of(std::integral_constant<size_t, Extents::rank() - K>(), ext,
+ sl)) -
+ index_t(first_of(sl)));
+ }
+
+ MDSPAN_TEMPLATE_REQUIRES(
+ class Slice, class... SlicesAndExtents,
+ /* requires */ (std::is_convertible_v<Slice, size_t>)
+ )
+ MDSPAN_INLINE_FUNCTION
+ constexpr static auto next_extent(const Extents &ext, const Slice &,
+ SlicesAndExtents... slices_and_extents) {
+ using next_t = extents_constructor<K - 1, Extents, NewExtents...>;
+ return next_t::next_extent(ext, slices_and_extents...);
+ }
+
+ template <class OffsetType, class ExtentType, class StrideType,
+ class... SlicesAndExtents>
+ MDSPAN_INLINE_FUNCTION
+ constexpr static auto
+ next_extent(const Extents &ext,
+ const strided_slice<OffsetType, ExtentType, StrideType> &r,
+ SlicesAndExtents... slices_and_extents) {
+ using index_t = typename Extents::index_type;
+ using new_static_extent_t =
+ StaticExtentFromStridedRange<ExtentType, StrideType>;
+ if constexpr (new_static_extent_t::value == dynamic_extent) {
+ using next_t =
+ extents_constructor<K - 1, Extents, NewExtents..., dynamic_extent>;
+ return next_t::next_extent(
+ ext, slices_and_extents...,
+ r.extent > 0 ? 1 + divide<index_t>(r.extent - 1, r.stride) : 0);
+ } else {
+ constexpr size_t new_static_extent = new_static_extent_t::value;
+ using next_t =
+ extents_constructor<K - 1, Extents, NewExtents..., new_static_extent>;
+ return next_t::next_extent(
+ ext, slices_and_extents..., index_t(divide<index_t>(ExtentType(), StrideType())));
+ }
+ }
+};
+
+template <class Extents, size_t... NewStaticExtents>
+struct extents_constructor<0, Extents, NewStaticExtents...> {
+
+ template <class... NewExtents>
+ MDSPAN_INLINE_FUNCTION
+ constexpr static auto next_extent(const Extents &, NewExtents... new_exts) {
+ return extents<typename Extents::index_type, NewStaticExtents...>(
+ new_exts...);
+ }
+};
+
+} // namespace detail
+
+// submdspan_extents creates new extents given src extents and submdspan slice
+// specifiers
+template <class IndexType, size_t... Extents, class... SliceSpecifiers>
+MDSPAN_INLINE_FUNCTION
+constexpr auto submdspan_extents(const extents<IndexType, Extents...> &src_exts,
+ SliceSpecifiers... slices) {
+
+ using ext_t = extents<IndexType, Extents...>;
+ return detail::extents_constructor<ext_t::rank(), ext_t>::next_extent(
+ src_exts, slices...);
+}
+} // namespace MDSPAN_IMPL_STANDARD_NAMESPACE
diff --git a/ext/mdspan/include/experimental/__p2630_bits/submdspan_mapping.hpp b/ext/mdspan/include/experimental/__p2630_bits/submdspan_mapping.hpp
new file mode 100644
index 0000000..46ccbaa
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p2630_bits/submdspan_mapping.hpp
@@ -0,0 +1,630 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+
+#pragma once
+
+#include <array>
+#include <type_traits>
+#include <utility> // index_sequence
+#include "../__p0009_bits/utility.hpp"
+
+// Suppress spurious warning with NVCC about no return statement.
+// This is a known issue in NVCC and NVC++
+// Depending on the CUDA and GCC version we need both the builtin
+// and the diagnostic push. I tried really hard to find something shorter
+// but no luck ...
+#if defined __NVCC__
+#ifdef __NVCC_DIAG_PRAGMA_SUPPORT__
+#pragma nv_diagnostic push
+#pragma nv_diag_suppress = implicit_return_from_non_void_function
+#else
+#ifdef __CUDA_ARCH__
+#pragma diagnostic push
+#pragma diag_suppress implicit_return_from_non_void_function
+#endif
+#endif
+#elif defined __NVCOMPILER
+#pragma diagnostic push
+#pragma diag_suppress = implicit_return_from_non_void_function
+#endif
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+//******************************************
+// Return type of submdspan_mapping overloads
+//******************************************
+template <class LayoutMapping> struct submdspan_mapping_result {
+ _MDSPAN_NO_UNIQUE_ADDRESS LayoutMapping mapping{};
+ size_t offset;
+};
+
+namespace detail {
+
+// We use const Slice& and not Slice&& because the various
+// submdspan_mapping_impl overloads use their slices arguments
+// multiple times. This makes perfect forwarding not useful, but we
+// still don't want to pass those (possibly of size 64 x 3 bits)
+// objects by value.
+template <class IndexType, class Slice>
+MDSPAN_INLINE_FUNCTION constexpr bool
+one_slice_out_of_bounds(const IndexType &ext, const Slice &slice) {
+ using common_t =
+ std::common_type_t<decltype(detail::first_of(slice)), IndexType>;
+ return static_cast<common_t>(detail::first_of(slice)) ==
+ static_cast<common_t>(ext);
+}
+
+template <size_t... RankIndices, class IndexType, size_t... Exts,
+ class... Slices>
+MDSPAN_INLINE_FUNCTION constexpr bool
+any_slice_out_of_bounds_helper(std::index_sequence<RankIndices...>,
+ const extents<IndexType, Exts...> &exts,
+ const Slices &... slices) {
+ return _MDSPAN_FOLD_OR(
+ (one_slice_out_of_bounds(exts.extent(RankIndices), slices)));
+}
+
+template <class IndexType, size_t... Exts, class... Slices>
+MDSPAN_INLINE_FUNCTION constexpr bool
+any_slice_out_of_bounds(const extents<IndexType, Exts...> &exts,
+ const Slices &... slices) {
+ return any_slice_out_of_bounds_helper(
+ std::make_index_sequence<sizeof...(Slices)>(), exts, slices...);
+}
+
+// constructs sub strides
+template<class T, size_t N>
+struct sub_strides
+{
+ T values[N > 0 ? N : 1];
+};
+
+template <class SrcMapping, class... slice_strides, size_t... InvMapIdxs>
+MDSPAN_INLINE_FUNCTION constexpr auto construct_sub_strides(
+ const SrcMapping &src_mapping, std::index_sequence<InvMapIdxs...>,
+ const MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple<slice_strides...> &slices_stride_factor) {
+ using index_type = typename SrcMapping::index_type;
+ return sub_strides<typename SrcMapping::index_type, sizeof...(InvMapIdxs)>{{
+ (static_cast<index_type>(src_mapping.stride(InvMapIdxs)) *
+ static_cast<index_type>(get<InvMapIdxs>(slices_stride_factor)))...}};
+}
+
+template<class SliceSpecifier, class IndexType>
+struct is_range_slice {
+ constexpr static bool value =
+ std::is_same_v<SliceSpecifier, full_extent_t> ||
+ index_pair_like<SliceSpecifier, IndexType>::value;
+};
+
+template<class SliceSpecifier, class IndexType>
+constexpr bool is_range_slice_v = is_range_slice<SliceSpecifier, IndexType>::value;
+
+template<class SliceSpecifier, class IndexType>
+struct is_index_slice {
+ constexpr static bool value = std::is_convertible_v<SliceSpecifier, IndexType>;
+};
+
+template<class SliceSpecifier, class IndexType>
+constexpr bool is_index_slice_v = is_index_slice<SliceSpecifier, IndexType>::value;
+
+} // namespace detail
+
+//**********************************
+// layout_left submdspan_mapping
+//*********************************
+namespace detail {
+
+// Figure out whether to preserve layout_left
+template <class IndexType, size_t SubRank, class IndexSequence,
+ class... SliceSpecifiers>
+struct deduce_layout_left_submapping;
+
+template <class IndexType, size_t SubRank, size_t... Idx,
+ class... SliceSpecifiers>
+struct deduce_layout_left_submapping<
+ IndexType, SubRank, std::index_sequence<Idx...>, SliceSpecifiers...> {
+
+ using count_range = index_sequence_scan_impl<
+ 0u, (is_index_slice_v<SliceSpecifiers, IndexType> ? 0u : 1u)...>;
+
+ constexpr static int gap_len =
+ (((Idx > 0 && count_range::get(Idx) == 1 &&
+ is_index_slice_v<SliceSpecifiers, IndexType>)
+ ? 1
+ : 0) +
+ ... + 0);
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr static bool layout_left_value() {
+ // Use layout_left for rank 0
+ if constexpr (SubRank == 0) {
+ return true;
+ // Use layout_left for rank 1 result if leftmost slice specifier is range like
+ } else if constexpr (SubRank == 1) {
+ return ((Idx > 0 || is_range_slice_v<SliceSpecifiers, IndexType>)&&...);
+ } else {
+ // Preserve if leftmost SubRank-1 slices are full_extent_t and
+ // the slice at idx Subrank - 1 is a range and
+ // for idx > SubRank the slice is an index
+ return ((((Idx < SubRank - 1) && std::is_same_v<SliceSpecifiers, full_extent_t>) ||
+ ((Idx == SubRank - 1) && is_range_slice_v<SliceSpecifiers, IndexType>) ||
+ ((Idx > SubRank - 1) && is_index_slice_v<SliceSpecifiers, IndexType>)) && ...);
+ }
+#if defined(__NVCC__) && !defined(__CUDA_ARCH__) && defined(__GNUC__)
+ __builtin_unreachable();
+#endif
+ }
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr static bool layout_left_padded_value() {
+ // Technically could also keep layout_left_padded for SubRank==0
+ // and SubRank==1 with leftmost slice specifier being a contiguous range
+ // but we intercept these cases separately
+
+ // In all other cases:
+ // leftmost slice must be range
+ // then there can be a gap with index slices
+ // then SubRank - 2 full_extent slices
+ // then another range slice
+ // then more index slices
+ // e.g. R I I I F F F R I I for obtaining a rank-5 from a rank-10
+ return ((((Idx == 0) && is_range_slice_v<SliceSpecifiers, IndexType>) ||
+ ((Idx > 0 && Idx <= gap_len) && is_index_slice_v<SliceSpecifiers, IndexType>) ||
+ ((Idx > gap_len && Idx < gap_len + SubRank - 1) && std::is_same_v<SliceSpecifiers, full_extent_t>) ||
+ ((Idx == gap_len + SubRank - 1) && is_range_slice_v<SliceSpecifiers, IndexType>) ||
+ ((Idx > gap_len + SubRank - 1) && is_index_slice_v<SliceSpecifiers, IndexType>)) && ... );
+ }
+};
+
+// We are reusing the same thing for layout_left and layout_left_padded
+// For layout_left as source StaticStride is static_extent(0)
+template<class Extents, size_t NumGaps, size_t StaticStride>
+struct compute_s_static_layout_left {
+ // Neither StaticStride nor any of the provided extents can be zero.
+ // StaticStride can never be zero, the static_extents we are looking at are associated with
+ // integral slice specifiers - which wouldn't be valid for zero extent
+ template<size_t ... Idx>
+ MDSPAN_INLINE_FUNCTION
+ static constexpr size_t value(std::index_sequence<Idx...>) {
+ size_t val = ((Idx>0 && Idx<=NumGaps ? (Extents::static_extent(Idx) == dynamic_extent?0:Extents::static_extent(Idx)) : 1) * ... * (StaticStride == dynamic_extent?0:StaticStride));
+ return val == 0?dynamic_extent:val;
+ }
+};
+
+} // namespace detail
+
+// Actual submdspan mapping call
+template <class Extents>
+template <class... SliceSpecifiers>
+MDSPAN_INLINE_FUNCTION constexpr auto
+layout_left::mapping<Extents>::submdspan_mapping_impl(
+ SliceSpecifiers... slices) const {
+
+ // compute sub extents
+ using src_ext_t = Extents;
+ auto dst_ext = submdspan_extents(extents(), slices...);
+ using dst_ext_t = decltype(dst_ext);
+
+ // figure out sub layout type
+ using deduce_layout = detail::deduce_layout_left_submapping<
+ typename dst_ext_t::index_type, dst_ext_t::rank(),
+ std::make_index_sequence<src_ext_t::rank()>,
+ SliceSpecifiers...>;
+
+ // Figure out if any slice's lower bound equals the corresponding extent.
+ // If so, bypass evaluating the layout mapping. This fixes LWG Issue 4060.
+ const bool out_of_bounds =
+ detail::any_slice_out_of_bounds(this->extents(), slices...);
+ auto offset = static_cast<size_t>(
+ out_of_bounds ? this->required_span_size()
+ : this->operator()(detail::first_of(slices)...));
+
+ if constexpr (deduce_layout::layout_left_value()) {
+ // layout_left case
+ using dst_mapping_t = typename layout_left::template mapping<dst_ext_t>;
+ return submdspan_mapping_result<dst_mapping_t>{dst_mapping_t(dst_ext),
+ offset};
+ } else if constexpr (deduce_layout::layout_left_padded_value()) {
+ constexpr size_t S_static = MDSPAN_IMPL_STANDARD_NAMESPACE::detail::compute_s_static_layout_left<Extents, deduce_layout::gap_len, Extents::static_extent(0)>::value(std::make_index_sequence<Extents::rank()>());
+ using dst_mapping_t = typename MDSPAN_IMPL_PROPOSED_NAMESPACE::layout_left_padded<S_static>::template mapping<dst_ext_t>;
+ return submdspan_mapping_result<dst_mapping_t>{
+ dst_mapping_t(dst_ext, stride(1 + deduce_layout::gap_len)), offset};
+ } else {
+ // layout_stride case
+ using dst_mapping_t = typename layout_stride::mapping<dst_ext_t>;
+ auto inv_map = detail::inv_map_rank(std::integral_constant<size_t, 0>(),
+ std::index_sequence<>(), slices...);
+ return submdspan_mapping_result<dst_mapping_t> {
+ dst_mapping_t(mdspan_non_standard, dst_ext,
+ detail::construct_sub_strides(
+ *this, inv_map,
+// HIP needs deduction guides to have markups so we need to be explicit
+// NVCC 11.0 has a bug with deduction guide here, tested that 11.2 does not have
+// the issue but Clang-CUDA also doesn't accept the use of deduction guide so
+// disable it for CUDA altogether
+#if defined(_MDSPAN_HAS_HIP) || defined(_MDSPAN_HAS_CUDA)
+ detail::tuple<decltype(detail::stride_of(slices))...>{
+ detail::stride_of(slices)...}).values),
+#else
+ detail::tuple{detail::stride_of(slices)...}).values),
+#endif
+ offset
+ };
+ }
+#if defined(__NVCC__) && !defined(__CUDA_ARCH__) && defined(__GNUC__)
+ __builtin_unreachable();
+#endif
+}
+
+template <size_t PaddingValue>
+template <class Extents>
+template <class... SliceSpecifiers>
+MDSPAN_INLINE_FUNCTION constexpr auto
+MDSPAN_IMPL_PROPOSED_NAMESPACE::layout_left_padded<PaddingValue>::mapping<Extents>::submdspan_mapping_impl(
+ SliceSpecifiers... slices) const {
+
+ // compute sub extents
+ using src_ext_t = Extents;
+ auto dst_ext = submdspan_extents(extents(), slices...);
+ using dst_ext_t = decltype(dst_ext);
+
+ if constexpr (Extents::rank() == 0) { // rank-0 case
+ using dst_mapping_t = typename MDSPAN_IMPL_PROPOSED_NAMESPACE::layout_left_padded<PaddingValue>::template mapping<Extents>;
+ return submdspan_mapping_result<dst_mapping_t>{*this, 0};
+ } else {
+ const bool out_of_bounds =
+ MDSPAN_IMPL_STANDARD_NAMESPACE::detail::any_slice_out_of_bounds(this->extents(), slices...);
+ auto offset = static_cast<size_t>(
+ out_of_bounds ? this->required_span_size()
+ : this->operator()(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::first_of(slices)...));
+ if constexpr (dst_ext_t::rank() == 0) { // result rank-0
+ // The following for some reasons leads to compiler error later, while not using a typedef works:
+ // Compilers: CUDA 11.2 with GCC 9.1
+ //
+ // using dst_mapping_t = typename layout_left::template mapping<dst_ext_t>;
+ // return submdspan_mapping_result<dst_mapping_t>{dst_mapping_t{dst_ext}, offset};
+ //
+ // Error: submdspan_mapping.hpp:299:23: error: 'dst_mapping_t' does not name a type
+ // 299 | using dst_mapping_t = typename layout_left::template mapping<dst_ext_t>;
+ // The same error is given (about dst_mapping_t not naming type) when a different name is used in 299:
+ // using dst_mapping_t2 = typename layout_left::template mapping<dst_ext_t>;
+
+ return submdspan_mapping_result<typename layout_left::template mapping<dst_ext_t>>
+ {typename layout_left::template mapping<dst_ext_t>{dst_ext}, offset};
+ } else { // general case
+ // Figure out if any slice's lower bound equals the corresponding extent.
+ // If so, bypass evaluating the layout mapping. This fixes LWG Issue 4060.
+ // figure out sub layout type
+ using deduce_layout = MDSPAN_IMPL_STANDARD_NAMESPACE::detail::deduce_layout_left_submapping<
+ typename dst_ext_t::index_type, dst_ext_t::rank(),
+ decltype(std::make_index_sequence<src_ext_t::rank()>()),
+ SliceSpecifiers...>;
+
+ if constexpr (deduce_layout::layout_left_value() && dst_ext_t::rank() == 1) { // getting rank-1 from leftmost
+ using dst_mapping_t = typename layout_left::template mapping<dst_ext_t>;
+ return submdspan_mapping_result<dst_mapping_t>{dst_mapping_t{dst_ext}, offset};
+ } else if constexpr (deduce_layout::layout_left_padded_value()) { // can keep layout_left_padded
+ constexpr size_t S_static = MDSPAN_IMPL_STANDARD_NAMESPACE::detail::compute_s_static_layout_left<Extents, deduce_layout::gap_len, static_padding_stride>::value(std::make_index_sequence<Extents::rank()>());
+ using dst_mapping_t = typename MDSPAN_IMPL_PROPOSED_NAMESPACE::layout_left_padded<S_static>::template mapping<dst_ext_t>;
+ return submdspan_mapping_result<dst_mapping_t>{
+ dst_mapping_t(dst_ext, stride(1 + deduce_layout::gap_len)), offset};
+ } else { // layout_stride
+ auto inv_map = MDSPAN_IMPL_STANDARD_NAMESPACE::detail::inv_map_rank(std::integral_constant<size_t, 0>(),
+ std::index_sequence<>(), slices...);
+ using dst_mapping_t = typename layout_stride::template mapping<dst_ext_t>;
+ return submdspan_mapping_result<dst_mapping_t> {
+ dst_mapping_t(mdspan_non_standard, dst_ext,
+ MDSPAN_IMPL_STANDARD_NAMESPACE::detail::construct_sub_strides(
+ *this, inv_map,
+// HIP needs deduction guides to have markups so we need to be explicit
+// NVCC 11.0 has a bug with deduction guide here, tested that 11.2 does not have
+// the issue but Clang-CUDA also doesn't accept the use of deduction guide so
+// disable it for CUDA alltogether
+#if defined(_MDSPAN_HAS_HIP) || defined(_MDSPAN_HAS_CUDA)
+ MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple<decltype(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices))...>{
+ MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices)...}).values),
+#else
+ MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple{MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices)...}).values),
+#endif
+ offset
+ };
+ }
+ }
+ }
+
+
+#if defined(__NVCC__) && !defined(__CUDA_ARCH__) && defined(__GNUC__)
+ __builtin_unreachable();
+#endif
+}
+
+//**********************************
+// layout_right submdspan_mapping
+//*********************************
+namespace detail {
+
+// Figure out whether to preserve layout_right
+template <class IndexType, size_t SubRank, class IndexSequence,
+ class... SliceSpecifiers>
+struct deduce_layout_right_submapping;
+
+template <class IndexType, size_t SubRank, size_t... Idx,
+ class... SliceSpecifiers>
+struct deduce_layout_right_submapping<
+ IndexType, SubRank, std::index_sequence<Idx...>, SliceSpecifiers...> {
+
+ static constexpr size_t Rank = sizeof...(Idx);
+ using count_range = index_sequence_scan_impl<
+ 0u, (std::is_convertible_v<SliceSpecifiers, IndexType> ? 0u : 1u)...>;
+ //__static_partial_sums<!std::is_convertible_v<SliceSpecifiers,
+ // IndexType>...>;
+ constexpr static int gap_len =
+ (((Idx < Rank - 1 && count_range::get(Idx) == SubRank - 1 &&
+ std::is_convertible_v<SliceSpecifiers, IndexType>)
+ ? 1
+ : 0) +
+ ... + 0);
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr static bool layout_right_value() {
+ // Use layout_right for rank 0
+ if constexpr (SubRank == 0) {
+ return true;
+ // Use layout_right for rank 1 result if rightmost slice specifier is range like
+ } else if constexpr (SubRank == 1) {
+ return ((Idx < Rank - 1 || is_range_slice_v<SliceSpecifiers, IndexType>)&&...);
+ } else {
+ // Preserve if rightmost SubRank-1 slices are full_extent_t and
+ // the slice at idx Rank-Subrank is a range and
+ // for idx < Rank - SubRank the slice is an index
+ return ((((Idx >= Rank - SubRank) && std::is_same_v<SliceSpecifiers, full_extent_t>) ||
+ ((Idx == Rank - SubRank) && is_range_slice_v<SliceSpecifiers, IndexType>) ||
+ ((Idx < Rank - SubRank) && is_index_slice_v<SliceSpecifiers, IndexType>)) && ...);
+ }
+#if defined(__NVCC__) && !defined(__CUDA_ARCH__) && defined(__GNUC__)
+ __builtin_unreachable();
+#endif
+ }
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr static bool layout_right_padded_value() {
+ // Technically could also keep layout_right_padded for SubRank==0
+ // and SubRank==1 with rightmost slice specifier being a contiguous range
+ // but we intercept these cases separately
+
+ // In all other cases:
+ // rightmost slice must be range
+ // then there can be a gap with index slices
+ // then SubRank - 2 full_extent slices
+ // then another range slice
+ // then more index slices
+ // e.g. I I R F F F I I I R for obtaining a rank-5 from a rank-10
+ return ((((Idx == Rank - 1) && is_range_slice_v<SliceSpecifiers, IndexType>) ||
+ ((Idx >= Rank - gap_len - 1 && Idx < Rank - 1) && is_index_slice_v<SliceSpecifiers, IndexType>) ||
+ ((Idx > Rank - gap_len - SubRank && Idx < Rank - gap_len - 1) && std::is_same_v<SliceSpecifiers, full_extent_t>) ||
+ ((Idx == Rank - gap_len - SubRank) && is_range_slice_v<SliceSpecifiers, IndexType>) ||
+ ((Idx < Rank - gap_len - SubRank) && is_index_slice_v<SliceSpecifiers, IndexType>)) && ... );
+ }
+};
+
+// We are reusing the same thing for layout_right and layout_right_padded
+// For layout_right as source StaticStride is static_extent(Rank-1)
+template<class Extents, size_t NumGaps, size_t StaticStride>
+struct compute_s_static_layout_right {
+ // Neither StaticStride nor any of the provided extents can be zero.
+ // StaticStride can never be zero, the static_extents we are looking at are associated with
+ // integral slice specifiers - which wouldn't be valid for zero extent
+ template<size_t ... Idx>
+ MDSPAN_INLINE_FUNCTION
+ static constexpr size_t value(std::index_sequence<Idx...>) {
+ size_t val = ((Idx >= Extents::rank() - 1 - NumGaps && Idx < Extents::rank() - 1 ? (Extents::static_extent(Idx) == dynamic_extent?0:Extents::static_extent(Idx)) : 1) * ... * (StaticStride == dynamic_extent?0:StaticStride));
+ return val == 0?dynamic_extent:val;
+ }
+};
+
+} // namespace detail
+
+// Actual submdspan mapping call
+template <class Extents>
+template <class... SliceSpecifiers>
+MDSPAN_INLINE_FUNCTION constexpr auto
+layout_right::mapping<Extents>::submdspan_mapping_impl(
+ SliceSpecifiers... slices) const {
+
+ // compute sub extents
+ using src_ext_t = Extents;
+ auto dst_ext = submdspan_extents(extents(), slices...);
+ using dst_ext_t = decltype(dst_ext);
+
+ // figure out sub layout type
+ using deduce_layout = detail::deduce_layout_right_submapping<
+ typename dst_ext_t::index_type, dst_ext_t::rank(),
+ std::make_index_sequence<src_ext_t::rank()>,
+ SliceSpecifiers...>;
+
+ // Figure out if any slice's lower bound equals the corresponding extent.
+ // If so, bypass evaluating the layout mapping. This fixes LWG Issue 4060.
+ const bool out_of_bounds =
+ detail::any_slice_out_of_bounds(this->extents(), slices...);
+ auto offset = static_cast<size_t>(
+ out_of_bounds ? this->required_span_size()
+ : this->operator()(detail::first_of(slices)...));
+
+ if constexpr (deduce_layout::layout_right_value()) {
+ // layout_right case
+ using dst_mapping_t = typename layout_right::mapping<dst_ext_t>;
+ return submdspan_mapping_result<dst_mapping_t>{dst_mapping_t(dst_ext),
+ offset};
+ } else if constexpr (deduce_layout::layout_right_padded_value()) {
+ constexpr size_t S_static = MDSPAN_IMPL_STANDARD_NAMESPACE::detail::compute_s_static_layout_left<Extents, deduce_layout::gap_len, Extents::static_extent(Extents::rank() - 1)>::value(std::make_index_sequence<Extents::rank()>());
+ using dst_mapping_t = typename MDSPAN_IMPL_PROPOSED_NAMESPACE::layout_right_padded<S_static>::template mapping<dst_ext_t>;
+ return submdspan_mapping_result<dst_mapping_t>{
+ dst_mapping_t(dst_ext,
+ stride(src_ext_t::rank() - 2 - deduce_layout::gap_len)),
+ offset};
+ } else {
+ // layout_stride case
+ using dst_mapping_t = typename layout_stride::mapping<dst_ext_t>;
+ auto inv_map = detail::inv_map_rank(std::integral_constant<size_t, 0>(),
+ std::index_sequence<>(), slices...);
+ return submdspan_mapping_result<dst_mapping_t> {
+ dst_mapping_t(mdspan_non_standard, dst_ext,
+ detail::construct_sub_strides(
+ *this, inv_map,
+// HIP needs deduction guides to have markups so we need to be explicit
+// NVCC 11.0 has a bug with deduction guide here, tested that 11.2 does not have
+// the issue but Clang-CUDA also doesn't accept the use of deduction guide so
+// disable it for CUDA altogether
+#if defined(_MDSPAN_HAS_HIP) || defined(_MDSPAN_HAS_CUDA)
+ MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple<decltype(detail::stride_of(slices))...>{
+ detail::stride_of(slices)...}).values),
+#else
+ MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple{detail::stride_of(slices)...}).values),
+#endif
+ offset
+ };
+ }
+#if defined(__NVCC__) && !defined(__CUDA_ARCH__) && defined(__GNUC__)
+ __builtin_unreachable();
+#endif
+}
+
+template <size_t PaddingValue>
+template <class Extents>
+template <class... SliceSpecifiers>
+MDSPAN_INLINE_FUNCTION constexpr auto
+MDSPAN_IMPL_PROPOSED_NAMESPACE::layout_right_padded<PaddingValue>::mapping<Extents>::submdspan_mapping_impl(
+ SliceSpecifiers... slices) const {
+
+ // compute sub extents
+ using src_ext_t = Extents;
+ auto dst_ext = submdspan_extents(extents(), slices...);
+ using dst_ext_t = decltype(dst_ext);
+
+ if constexpr (Extents::rank() == 0) { // rank-0 case
+ using dst_mapping_t = typename MDSPAN_IMPL_PROPOSED_NAMESPACE::layout_right_padded<PaddingValue>::template mapping<Extents>;
+ return submdspan_mapping_result<dst_mapping_t>{*this, 0};
+ } else {
+ // Figure out if any slice's lower bound equals the corresponding extent.
+ // If so, bypass evaluating the layout mapping. This fixes LWG Issue 4060.
+ // figure out sub layout type
+ const bool out_of_bounds =
+ MDSPAN_IMPL_STANDARD_NAMESPACE::detail::any_slice_out_of_bounds(this->extents(), slices...);
+ auto offset = static_cast<size_t>(
+ out_of_bounds ? this->required_span_size()
+ : this->operator()(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::first_of(slices)...));
+ if constexpr (dst_ext_t::rank() == 0) { // result rank-0
+ // Same issue as in layout_left_padded: see comment there
+ // using dst_mapping_t = typename layout_right::template mapping<dst_ext_t>;
+ // return submdspan_mapping_result<dst_mapping_t>{dst_mapping_t{dst_ext}, offset};
+ return submdspan_mapping_result<typename layout_right::template mapping<dst_ext_t>>
+ {typename layout_right::template mapping<dst_ext_t>{dst_ext}, offset};
+ } else { // general case
+ using deduce_layout = MDSPAN_IMPL_STANDARD_NAMESPACE::detail::deduce_layout_right_submapping<
+ typename dst_ext_t::index_type, dst_ext_t::rank(),
+ decltype(std::make_index_sequence<src_ext_t::rank()>()),
+ SliceSpecifiers...>;
+
+ if constexpr (deduce_layout::layout_right_value() && dst_ext_t::rank() == 1) { // getting rank-1 from rightmost
+ using dst_mapping_t = typename layout_right::template mapping<dst_ext_t>;
+ return submdspan_mapping_result<dst_mapping_t>{dst_mapping_t{dst_ext}, offset};
+ } else if constexpr (deduce_layout::layout_right_padded_value()) { // can keep layout_right_padded
+ constexpr size_t S_static = MDSPAN_IMPL_STANDARD_NAMESPACE::detail::compute_s_static_layout_right<Extents, deduce_layout::gap_len, static_padding_stride>::value(std::make_index_sequence<Extents::rank()>());
+ using dst_mapping_t = typename MDSPAN_IMPL_PROPOSED_NAMESPACE::layout_right_padded<S_static>::template mapping<dst_ext_t>;
+ return submdspan_mapping_result<dst_mapping_t>{
+ dst_mapping_t(dst_ext, stride(Extents::rank() - 2 - deduce_layout::gap_len)), offset};
+ } else { // layout_stride
+ auto inv_map = MDSPAN_IMPL_STANDARD_NAMESPACE::detail::inv_map_rank(std::integral_constant<size_t, 0>(),
+ std::index_sequence<>(), slices...);
+ using dst_mapping_t = typename layout_stride::template mapping<dst_ext_t>;
+ return submdspan_mapping_result<dst_mapping_t> {
+ dst_mapping_t(mdspan_non_standard, dst_ext,
+ MDSPAN_IMPL_STANDARD_NAMESPACE::detail::construct_sub_strides(
+ *this, inv_map,
+// HIP needs deduction guides to have markups so we need to be explicit
+// NVCC 11.0 has a bug with deduction guide here, tested that 11.2 does not have
+// the issue but Clang-CUDA also doesn't accept the use of deduction guide so
+// disable it for CUDA alltogether
+#if defined(_MDSPAN_HAS_HIP) || defined(_MDSPAN_HAS_CUDA)
+ MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple<decltype(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices))...>{
+ MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices)...}).values),
+#else
+ MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple{MDSPAN_IMPL_STANDARD_NAMESPACE::detail::stride_of(slices)...}).values),
+#endif
+ offset
+ };
+ }
+ }
+ }
+
+
+#if defined(__NVCC__) && !defined(__CUDA_ARCH__) && defined(__GNUC__)
+ __builtin_unreachable();
+#endif
+}
+
+//**********************************
+// layout_stride submdspan_mapping
+//*********************************
+template <class Extents>
+template <class... SliceSpecifiers>
+MDSPAN_INLINE_FUNCTION constexpr auto
+layout_stride::mapping<Extents>::submdspan_mapping_impl(
+ SliceSpecifiers... slices) const {
+ auto dst_ext = submdspan_extents(extents(), slices...);
+ using dst_ext_t = decltype(dst_ext);
+ auto inv_map = detail::inv_map_rank(std::integral_constant<size_t, 0>(),
+ std::index_sequence<>(), slices...);
+ using dst_mapping_t = typename layout_stride::template mapping<dst_ext_t>;
+
+ // Figure out if any slice's lower bound equals the corresponding extent.
+ // If so, bypass evaluating the layout mapping. This fixes LWG Issue 4060.
+ const bool out_of_bounds =
+ detail::any_slice_out_of_bounds(this->extents(), slices...);
+ auto offset = static_cast<size_t>(
+ out_of_bounds ? this->required_span_size()
+ : this->operator()(detail::first_of(slices)...));
+
+ return submdspan_mapping_result<dst_mapping_t> {
+ dst_mapping_t(mdspan_non_standard, dst_ext,
+ detail::construct_sub_strides(
+ *this, inv_map,
+// HIP needs deduction guides to have markups so we need to be explicit
+// NVCC 11.0 has a bug with deduction guide here, tested that 11.2 does not have
+// the issue but Clang-CUDA also doesn't accept the use of deduction guide so
+// disable it for CUDA alltogether
+#if defined(_MDSPAN_HAS_HIP) || defined(_MDSPAN_HAS_CUDA)
+ MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple<decltype(detail::stride_of(slices))...>(
+ detail::stride_of(slices)...)).values),
+#else
+ MDSPAN_IMPL_STANDARD_NAMESPACE::detail::tuple(detail::stride_of(slices)...)).values),
+#endif
+ offset
+ };
+}
+
+} // namespace MDSPAN_IMPL_STANDARD_NAMESPACE
+
+#if defined __NVCC__
+#ifdef __NVCC_DIAG_PRAGMA_SUPPORT__
+#pragma nv_diagnostic pop
+#else
+#ifdef __CUDA_ARCH__
+#pragma diagnostic pop
+#endif
+#endif
+#elif defined __NVCOMPILER
+#pragma diagnostic pop
+#endif
diff --git a/ext/mdspan/include/experimental/__p2642_bits/layout_padded.hpp b/ext/mdspan/include/experimental/__p2642_bits/layout_padded.hpp
new file mode 100644
index 0000000..a714090
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p2642_bits/layout_padded.hpp
@@ -0,0 +1,869 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+//
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+#pragma once
+
+#include <cassert>
+#include "layout_padded_fwd.hpp"
+#include "../__p0009_bits/dynamic_extent.hpp"
+#include "../__p0009_bits/extents.hpp"
+#include "../__p0009_bits/mdspan.hpp"
+#include "../__p0009_bits/layout_left.hpp"
+#include "../__p0009_bits/layout_right.hpp"
+#include "../__p0009_bits/layout_stride.hpp"
+#include "../__p0009_bits/utility.hpp"
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+namespace MDSPAN_IMPL_PROPOSED_NAMESPACE {
+
+namespace detail {
+template<class _T>
+MDSPAN_INLINE_FUNCTION
+constexpr _T
+find_next_multiple(_T alignment, _T offset)
+{
+ if ( alignment == 0 ) {
+ return _T(0);
+ } else {
+ return ( ( offset + alignment - 1 ) / alignment) * alignment;
+ }
+}
+
+template <class _ExtentsType, size_t _PaddingValue, size_t _ExtentToPadIdx>
+MDSPAN_INLINE_FUNCTION constexpr size_t get_actual_static_padding_value() {
+ constexpr auto rank = _ExtentsType::rank();
+
+ if constexpr (rank <= typename _ExtentsType::rank_type(1)) {
+ return 0;
+ } else if constexpr (_PaddingValue != dynamic_extent &&
+ _ExtentsType::static_extent(_ExtentToPadIdx) !=
+ dynamic_extent) {
+ static_assert(
+ (_PaddingValue != 0) ||
+ (_ExtentsType::static_extent(_ExtentToPadIdx) == 0),
+ "padding stride can be 0 only if "
+ "extents_type::static_extent(extent-to-pad) is 0 or dynamic_extent");
+ return find_next_multiple(_PaddingValue,
+ _ExtentsType::static_extent(_ExtentToPadIdx));
+ } else {
+ return dynamic_extent;
+ }
+ // Missing return statement warning from NVCC and ICC
+#if (defined(__NVCC__) || defined(__INTEL_COMPILER)) && !defined(__NVCOMPILER)
+ return 0;
+#endif
+}
+
+template <size_t _PaddingValue, typename _Extents, size_t _ExtentToPadIdx, size_t _Rank, typename Enabled = void>
+struct static_array_type_for_padded_extent
+{
+ static constexpr size_t padding_value = _PaddingValue;
+ using index_type = typename _Extents::index_type;
+ using extents_type = _Extents;
+ using type = ::MDSPAN_IMPL_STANDARD_NAMESPACE::detail::maybe_static_array<
+ index_type, size_t, dynamic_extent,
+ ::MDSPAN_IMPL_STANDARD_NAMESPACE::MDSPAN_IMPL_PROPOSED_NAMESPACE::detail::get_actual_static_padding_value<extents_type, _PaddingValue,
+ _ExtentToPadIdx>()>;
+};
+
+template <size_t _PaddingValue, typename _Extents, size_t _ExtentToPadIdx, size_t Rank>
+struct static_array_type_for_padded_extent<_PaddingValue, _Extents,
+ _ExtentToPadIdx, Rank, std::enable_if_t<Rank <= 1>> {
+ using index_type = typename _Extents::index_type;
+ using extents_type = _Extents;
+ using type =
+ ::MDSPAN_IMPL_STANDARD_NAMESPACE::detail::maybe_static_array<
+ index_type, size_t, dynamic_extent, 0>;
+};
+
+template <size_t _PaddingValue, typename _Extents, size_t _ExtentToPadIdx>
+struct padded_extent {
+ static constexpr size_t padding_value = _PaddingValue;
+ using index_type = typename _Extents::index_type;
+ using extents_type = _Extents;
+ using static_array_type = typename static_array_type_for_padded_extent<
+ padding_value, _Extents, _ExtentToPadIdx, _Extents::rank()>::type;
+
+ MDSPAN_INLINE_FUNCTION
+ static constexpr auto static_value() { return static_array_type::static_value(0); }
+
+ MDSPAN_INLINE_FUNCTION
+ static constexpr static_array_type
+ init_padding(const _Extents &exts) {
+ if constexpr ((_Extents::rank() > 1) && (padding_value == dynamic_extent)) {
+ return {exts.extent(_ExtentToPadIdx)};
+ } else {
+ return init_padding(exts, padding_value);
+ }
+ // Missing return statement warning from NVCC and ICC
+#if (defined(__NVCC__) || defined(__INTEL_COMPILER)) && !defined(__NVCOMPILER)
+ return {};
+#endif
+ }
+
+ MDSPAN_INLINE_FUNCTION static constexpr static_array_type
+ init_padding([[maybe_unused]] const _Extents &exts,
+ [[maybe_unused]] index_type pv) {
+ if constexpr (_Extents::rank() > 1) {
+ return {find_next_multiple(pv,
+ exts.extent(_ExtentToPadIdx))};
+ } else {
+ return {};
+ }
+ // Missing return statement warning from NVCC and ICC
+#if (defined(__NVCC__) || defined(__INTEL_COMPILER)) && !defined(__NVCOMPILER)
+ return {};
+#endif
+ }
+
+ template <typename _Mapping, size_t _PaddingStrideIdx>
+ MDSPAN_INLINE_FUNCTION static constexpr static_array_type
+ init_padding([[maybe_unused]] const _Mapping &other_mapping,
+ std::integral_constant<size_t, _PaddingStrideIdx>) {
+ if constexpr (_Extents::rank() > 1) {
+ return {other_mapping.stride(_PaddingStrideIdx)};
+ } else {
+ return {};
+ }
+ // Missing return statement warning from NVCC and ICC
+#if (defined(__NVCC__) || defined(__INTEL_COMPILER)) && !defined(__NVCOMPILER)
+ return {};
+#endif
+ }
+};
+} // namespace detail
+
+template <size_t PaddingValue>
+template <class Extents>
+class layout_left_padded<PaddingValue>::mapping {
+public:
+ static constexpr size_t padding_value = PaddingValue;
+
+ using extents_type = Extents;
+ using index_type = typename extents_type::index_type;
+ using size_type = typename extents_type::size_type;
+ using rank_type = typename extents_type::rank_type;
+ using layout_type = layout_left_padded<padding_value>;
+
+#ifndef MDSPAN_INTERNAL_TEST
+private:
+#endif // MDSPAN_INTERNAL_TEST
+
+ static constexpr rank_type padded_stride_idx = detail::layout_padded_constants<layout_type, extents_type>::padded_stride_idx;
+ static constexpr rank_type extent_to_pad_idx = detail::layout_padded_constants<layout_type, extents_type>::extent_to_pad_idx;
+
+ static_assert((padding_value != 0)
+ || (extents_type::static_extent(extent_to_pad_idx) == 0)
+ || (extents_type::static_extent(extent_to_pad_idx) == dynamic_extent),
+ "out of bounds access for rank 0");
+
+ using padded_stride_type = detail::padded_extent< padding_value, extents_type, extent_to_pad_idx >;
+
+ static constexpr size_t static_padding_stride = padded_stride_type::static_value();
+
+ typename padded_stride_type::static_array_type padded_stride = {};
+ extents_type exts = {};
+
+ MDSPAN_INLINE_FUNCTION constexpr index_type
+ compute_offset(std::index_sequence<>) const {
+ return 0;
+ }
+
+ template <size_t Rank, class IndexOffset>
+ MDSPAN_INLINE_FUNCTION constexpr index_type
+ compute_offset(std::index_sequence<Rank>, IndexOffset index_offset) const {
+ return index_offset;
+ }
+
+ template <size_t... Ranks, class... IndexOffsets>
+ MDSPAN_INLINE_FUNCTION constexpr index_type
+ compute_offset(std::index_sequence<Ranks...>,
+ IndexOffsets... index_offsets) const {
+ index_type indices[] = {static_cast<index_type>(index_offsets)...};
+ // self-recursive fold trick from
+ // https://github.com/llvm/llvm-project/blob/96e1914aa2e6d8966acbfbe2f4d184201f1aa318/libcxx/include/mdspan/layout_left.h#L144
+ index_type res = 0;
+ ((res = indices[extents_type::rank() - 1 - Ranks] +
+ ((extents_type::rank() - 1 - Ranks) == extent_to_pad_idx
+ ? padded_stride.value(0)
+ : exts.extent(extents_type::rank() - 1 - Ranks)) *
+ res),
+ ...);
+ return res;
+ }
+
+public:
+#if !MDSPAN_HAS_CXX_20 || defined(__NVCC__)
+ MDSPAN_INLINE_FUNCTION
+ constexpr mapping()
+ : mapping(extents_type{})
+ {}
+#else
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ constexpr mapping()
+ requires(static_padding_stride != dynamic_extent) = default;
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr mapping()
+ requires(static_padding_stride == dynamic_extent)
+ : mapping(extents_type{})
+ {}
+#endif
+
+ MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr mapping(const mapping&) noexcept = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr mapping& operator=(const mapping&) noexcept = default;
+
+ /**
+ * Initializes the mapping with the given extents.
+ *
+ * \param ext the given extents
+ */
+ MDSPAN_INLINE_FUNCTION
+ constexpr mapping(const extents_type& ext)
+ : padded_stride(padded_stride_type::init_padding(ext)), exts(ext)
+ {}
+
+ /**
+ * Initializes the mapping with the given extents and the specified padding value.
+ *
+ * This overload participates in overload resolution only if `is_convertible_v<Size, index_type>`
+ * is `true` and `is_nothrow_constructible_v<index_type, Size>` is `true`
+ *
+ * \param ext the given extents
+ * \param padding_value the padding value
+ */
+ MDSPAN_TEMPLATE_REQUIRES(
+ class _Size,
+ /* requires */ (
+ std::is_convertible_v<_Size, index_type>
+ && std::is_nothrow_constructible_v<index_type, _Size>
+ )
+ )
+ MDSPAN_INLINE_FUNCTION
+ constexpr mapping(const extents_type &ext, _Size dynamic_padding_value)
+ : padded_stride(padded_stride_type::init_padding(ext, dynamic_padding_value)), exts(ext)
+ {
+ assert((padding_value == dynamic_extent) || (static_cast<index_type>(padding_value) == static_cast<index_type>(dynamic_padding_value)));
+ }
+
+ /**
+ * Converting constructor from `layout_left::mapping`.
+ *
+ * This overload participates in overload resolution only if
+ * `is_constructible_v<extents_type, OtherExtents>` is true. If
+ * `OtherExtents::rank() > 1` then one of `padding_value`, `static_extent(0)`,
+ * or `OtherExtents::static_extent(0)` must be `dynamic_extent`; otherwise,
+ * `OtherExtents::static_extent(0)` must be equal to the least multiple of
+ * `padding_value` greater than or equal to `extents_type::static_extent(0)`
+ */
+ MDSPAN_TEMPLATE_REQUIRES(
+ class _OtherExtents,
+ /* requires */ (std::is_constructible_v<extents_type, _OtherExtents>))
+ MDSPAN_CONDITIONAL_EXPLICIT(
+ (!std::is_convertible_v<_OtherExtents, extents_type>))
+ MDSPAN_INLINE_FUNCTION
+ constexpr mapping(const layout_left::mapping<_OtherExtents> &other_mapping)
+ : padded_stride(padded_stride_type::init_padding(
+ other_mapping,
+ std::integral_constant<size_t, padded_stride_idx>{})),
+ exts(other_mapping.extents()) {
+ static_assert(
+ (_OtherExtents::rank() > 1) ||
+ (static_padding_stride != dynamic_extent) ||
+ (_OtherExtents::static_extent(extent_to_pad_idx) != dynamic_extent) ||
+ (static_padding_stride ==
+ _OtherExtents::static_extent(extent_to_pad_idx)));
+ }
+
+ /**
+ * Converting constructor from `layout_stride::mapping`.
+ *
+ * This overload participates in overload resolution only if
+ * `is_constructible_v<extents_type, OtherExtents>` is true
+ */
+ MDSPAN_TEMPLATE_REQUIRES(
+ class _OtherExtents,
+ /* requires */ (std::is_constructible_v<extents_type, _OtherExtents>))
+ MDSPAN_CONDITIONAL_EXPLICIT((extents_type::rank() > 0))
+ MDSPAN_INLINE_FUNCTION
+ constexpr mapping(const layout_stride::mapping<_OtherExtents> &other_mapping)
+ : padded_stride(padded_stride_type::init_padding(
+ other_mapping,
+ std::integral_constant<size_t, padded_stride_idx>{})),
+ exts(other_mapping.extents()) {}
+
+ /**
+ * Converting constructor from `layout_left_padded::mapping`.
+ *
+ * This overload participates in overload resolution only if
+ * `is_constructible_v<extents_type, OtherExtents>` is true. Either
+ * `padding_value` or `OtherPaddingStride` must be `std::dynamic_extent`, or
+ * `padding_value == OtherPaddingStride`.
+ */
+ MDSPAN_TEMPLATE_REQUIRES(
+ class _Mapping,
+ /* requires */ (detail::is_layout_left_padded_mapping<_Mapping>::value
+ &&std::is_constructible_v<
+ extents_type, typename _Mapping::extents_type>))
+ MDSPAN_CONDITIONAL_EXPLICIT((extents_type::rank() > 1 &&
+ (padding_value == dynamic_extent ||
+ _Mapping::padding_value == dynamic_extent)))
+ MDSPAN_INLINE_FUNCTION
+ constexpr mapping(const _Mapping &other_mapping)
+ : padded_stride(padded_stride_type::init_padding(
+ other_mapping,
+ std::integral_constant<size_t, padded_stride_idx>{})),
+ exts(other_mapping.extents()) {
+ static_assert(padding_value == dynamic_extent ||
+ _Mapping::padding_value == dynamic_extent ||
+ padding_value == _Mapping::padding_value);
+ }
+
+ /**
+ * Converting constructor from `layout_right_padded::mapping`.
+ *
+ * This overload participates in overload resolution only if
+ * `extents_type::rank()` is 0 or 1 and `is_constructible_v<extents_type,
+ * OtherExtents>` is `true`.
+ */
+ MDSPAN_TEMPLATE_REQUIRES(
+ class _Mapping,
+ /* requires */ (detail::is_layout_right_padded_mapping<_Mapping>::value
+ &&extents_type::rank() <= 1 &&
+ std::is_constructible_v<extents_type,
+ typename _Mapping::extents_type>))
+ MDSPAN_CONDITIONAL_EXPLICIT(
+ (!std::is_convertible_v<typename _Mapping::extents_type, extents_type>))
+ MDSPAN_INLINE_FUNCTION
+ constexpr mapping(const _Mapping &other_mapping) noexcept
+ : padded_stride(padded_stride_type::init_padding(
+ static_cast<extents_type>(other_mapping.extents()),
+ other_mapping.extents().extent(extent_to_pad_idx))),
+ exts(other_mapping.extents()) {}
+
+ MDSPAN_INLINE_FUNCTION constexpr const extents_type &
+ extents() const noexcept {
+ return exts;
+ }
+
+ constexpr std::array<index_type, extents_type::rank()>
+ strides() const noexcept {
+ if constexpr (extents_type::rank() == 0) {
+ return {};
+ } else if constexpr (extents_type::rank() == 1) {
+ return {1};
+ } else {
+ index_type value = 1;
+ std::array<index_type, extents_type::rank()> s{};
+ s[extent_to_pad_idx] = value;
+ value *= padded_stride.value(0);
+ for (rank_type r = extent_to_pad_idx + 1; r < extents_type::rank() - 1;
+ ++r) {
+ s[r] = value;
+ value *= exts.extent(r);
+ }
+ s[extents_type::rank() - 1] = value;
+ return s;
+ }
+ }
+
+ MDSPAN_INLINE_FUNCTION constexpr index_type
+ required_span_size() const noexcept {
+ if constexpr (extents_type::rank() == 0) {
+ return 1;
+ } else if constexpr (extents_type::rank() == 1) {
+ return exts.extent(0);
+ } else {
+ index_type value = padded_stride.value(0);
+ for (rank_type r = 1; r < extents_type::rank(); ++r) {
+ value *= exts.extent(r);
+ }
+ return value;
+ }
+ }
+
+ /**
+ * Return the mapping given the provided indices per rank.
+ *
+ * This overload participates in overload resolution only if:
+ * - `sizeof...(Indices) == extents_type::rank()`,
+ * - `(is_convertible_v<Indices, index_type> && ...) is true`, and
+ * - (is_nothrow_constructible_v<index_type, Indices> && ...) is true.
+ */
+ MDSPAN_TEMPLATE_REQUIRES(
+ class... _Indices,
+ /* requires */ (sizeof...(_Indices) == extents_type::rank() &&
+ (::MDSPAN_IMPL_STANDARD_NAMESPACE::detail::
+ are_valid_indices<index_type, _Indices...>())))
+ MDSPAN_INLINE_FUNCTION constexpr size_t
+ operator()(_Indices... idxs) const noexcept {
+#if !defined(NDEBUG)
+ ::MDSPAN_IMPL_STANDARD_NAMESPACE::detail::check_all_indices(this->extents(),
+ idxs...);
+#endif // ! NDEBUG
+ return compute_offset(std::index_sequence_for<_Indices...>{}, idxs...);
+ }
+
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_unique() noexcept {
+ return true;
+ }
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_exhaustive() noexcept {
+ return (extents_type::rank() <= rank_type(1)) ||
+ (extents_type::static_extent(extent_to_pad_idx) != dynamic_extent &&
+ extents_type::static_extent(extent_to_pad_idx) ==
+ padded_stride_type::static_value());
+ }
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_strided() noexcept {
+ return true;
+ }
+
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_unique() noexcept {
+ return true;
+ }
+ MDSPAN_INLINE_FUNCTION constexpr bool is_exhaustive() const noexcept {
+ return (extents_type::rank() < 2) ||
+ (exts.extent(extent_to_pad_idx) == padded_stride.value(0));
+ }
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_strided() noexcept {
+ return true;
+ }
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr index_type stride(rank_type r) const noexcept {
+ assert(r < extents_type::rank());
+ if (r == 0)
+ return index_type(1);
+
+ index_type value = padded_stride.value(0);
+ for (rank_type k = 1; k < r; k++)
+ value *= exts.extent(k);
+
+ return value;
+ }
+
+ /**
+ * Equality operator between `layout_left_padded`s
+ *
+ * This overload only participates in overload resolution if
+ * `OtherExtents::rank() == extents_type::rank()`.
+ *
+ * \note There is currently a difference from p2642r2, where this function is
+ * specified as taking `layout_left_padded< padding_value >::mapping<
+ * Extents>`. However, this makes `padding_value` non-deducible.
+ */
+ MDSPAN_TEMPLATE_REQUIRES(
+ class _Mapping,
+ /* requires */ (detail::is_layout_left_padded_mapping<_Mapping>::value &&
+ (_Mapping::extents_type::rank() == extents_type::rank())))
+ MDSPAN_INLINE_FUNCTION friend constexpr bool
+ operator==(const mapping &left, const _Mapping &right) noexcept {
+ // Workaround for some compilers not short-circuiting properly with
+ // compile-time checks i.e. we can't access stride(_padding_stride_idx) of a
+ // rank 0 mapping
+ bool strides_equal = true;
+ if constexpr (extents_type::rank() > rank_type(1)) {
+ strides_equal =
+ left.stride(padded_stride_idx) == right.stride(padded_stride_idx);
+ }
+ return (left.extents() == right.extents()) && strides_equal;
+ }
+
+#if !MDSPAN_HAS_CXX_20
+ /**
+ * Inequality operator between `layout_left_padded`s
+ *
+ * This overload only participates in overload resolution if
+ * `OtherExtents::rank() == extents_type::rank()`.
+ */
+ MDSPAN_TEMPLATE_REQUIRES(
+ class _Mapping,
+ /* requires */ (detail::is_layout_left_padded_mapping<_Mapping>::value &&
+ (_Mapping::extents_type::rank() == extents_type::rank())))
+ MDSPAN_INLINE_FUNCTION friend constexpr bool
+ operator!=(const mapping &left, const _Mapping &right) noexcept {
+ return !(left == right);
+ }
+#endif
+
+ // [mdspan.submdspan.mapping], submdspan mapping specialization
+ template<class... SliceSpecifiers>
+ MDSPAN_INLINE_FUNCTION
+ constexpr auto submdspan_mapping_impl(
+ SliceSpecifiers... slices) const;
+
+ template<class... SliceSpecifiers>
+ MDSPAN_INLINE_FUNCTION
+ friend constexpr auto submdspan_mapping(
+ const mapping& src, SliceSpecifiers... slices) {
+ return src.submdspan_mapping_impl(slices...);
+ }
+};
+
+template <size_t PaddingValue>
+template <class Extents>
+class layout_right_padded<PaddingValue>::mapping {
+public:
+ static constexpr size_t padding_value = PaddingValue;
+
+ using extents_type = Extents;
+ using index_type = typename extents_type::index_type;
+ using size_type = typename extents_type::size_type;
+ using rank_type = typename extents_type::rank_type;
+ using layout_type = layout_right_padded<padding_value>;
+
+#ifndef MDSPAN_INTERNAL_TEST
+ private:
+#endif // MDSPAN_INTERNAL_TEST
+
+ static constexpr rank_type padded_stride_idx = detail::layout_padded_constants<layout_type, extents_type>::padded_stride_idx;
+ static constexpr rank_type extent_to_pad_idx = detail::layout_padded_constants<layout_type, extents_type>::extent_to_pad_idx;
+
+ static_assert((padding_value != 0)
+ || (extents_type::static_extent(extent_to_pad_idx) == 0)
+ || (extents_type::static_extent(extent_to_pad_idx) == dynamic_extent),
+ "if padding stride is 0, static_extent(extent-to-pad-rank) must also be 0 or dynamic_extent");
+
+ using padded_stride_type = detail::padded_extent< padding_value, extents_type, extent_to_pad_idx >;
+ static constexpr size_t static_padding_stride = padded_stride_type::static_value();
+
+ typename padded_stride_type::static_array_type padded_stride = {};
+ extents_type exts = {};
+
+ MDSPAN_INLINE_FUNCTION constexpr index_type
+ compute_offset(std::index_sequence<>) const {
+ return 0;
+ }
+
+ template <size_t Rank, class IndexOffset>
+ MDSPAN_INLINE_FUNCTION constexpr index_type
+ compute_offset(std::index_sequence<Rank>, IndexOffset index_offset) const {
+ return index_offset;
+ }
+
+ template <size_t... Ranks, class... IndexOffsets>
+ MDSPAN_INLINE_FUNCTION constexpr index_type
+ compute_offset(std::index_sequence<Ranks...>,
+ IndexOffsets... index_offsets) const {
+ // self-recursive fold trick from
+ // https://github.com/llvm/llvm-project/blob/4d9771741d40cc9cfcccb6b033f43689d36b705a/libcxx/include/mdspan/layout_right.h#L141
+ index_type res = 0;
+ ((res = static_cast<index_type>(index_offsets) +
+ (Ranks == extent_to_pad_idx ? padded_stride.value(0)
+ : exts.extent(Ranks)) *
+ res),
+ ...);
+ return res;
+ }
+
+public:
+#if !MDSPAN_HAS_CXX_20 || defined(__NVCC__)
+ MDSPAN_INLINE_FUNCTION
+ constexpr mapping()
+ : mapping(extents_type{})
+ {}
+#else
+ MDSPAN_INLINE_FUNCTION_DEFAULTED
+ constexpr mapping()
+ requires(static_padding_stride != dynamic_extent) = default;
+
+ MDSPAN_INLINE_FUNCTION
+ constexpr mapping()
+ requires(static_padding_stride == dynamic_extent)
+ : mapping(extents_type{})
+ {}
+#endif
+
+ MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr mapping(const mapping&) noexcept = default;
+ MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr mapping& operator=(const mapping&) noexcept = default;
+
+ /**
+ * Initializes the mapping with the given extents.
+ *
+ * \param ext the given extents
+ */
+ MDSPAN_INLINE_FUNCTION
+ constexpr mapping(const extents_type &ext)
+ : padded_stride(padded_stride_type::init_padding(ext)), exts(ext) {}
+
+ /**
+ * Initializes the mapping with the given extents and the specified padding value.
+ *
+ * This overload participates in overload resolution only if `is_convertible_v<Size, index_type>`
+ * is `true` and `is_nothrow_constructible_v<index_type, Size>` is `true`
+ *
+ * \param ext the given extents
+ * \param padding_value the padding value
+ */
+ MDSPAN_TEMPLATE_REQUIRES(
+ class _Size,
+ /* requires */ (
+ std::is_convertible_v<_Size, index_type>
+ && std::is_nothrow_constructible_v<index_type, _Size>
+ )
+ )
+ MDSPAN_INLINE_FUNCTION
+ constexpr mapping(const extents_type &ext, _Size dynamic_padding_value)
+ : padded_stride(padded_stride_type::init_padding(ext, static_cast<index_type>(dynamic_padding_value))),
+ exts(ext) {
+ assert((padding_value == dynamic_extent) ||
+ (static_cast<index_type>(padding_value) == static_cast<index_type>(dynamic_padding_value)));
+ }
+
+ /**
+ * Converting constructor from `layout_right::mapping`.
+ *
+ * This overload participates in overload resolution only if `is_constructible_v<extents_type, OtherExtents>` is true.
+ * If `OtherExtents::rank() > 1` then one of `padding_value`, `static_extent(0)`, or `OtherExtents::static_extent(0)` must be `dynamic_extent`;
+ * otherwise, `OtherExtents::static_extent(0)` must be equal to the least multiple of `padding_value` greater than or equal to `extents_type::static_extent(0)`
+ */
+ MDSPAN_TEMPLATE_REQUIRES(
+ class _OtherExtents,
+ /* requires */ (std::is_constructible_v<extents_type, _OtherExtents>))
+ MDSPAN_CONDITIONAL_EXPLICIT(
+ (!std::is_convertible_v<_OtherExtents, extents_type>))
+ MDSPAN_INLINE_FUNCTION
+ constexpr mapping(const layout_right::mapping<_OtherExtents> &other_mapping)
+ : padded_stride(padded_stride_type::init_padding(
+ other_mapping,
+ std::integral_constant<size_t, padded_stride_idx>{})),
+ exts(other_mapping.extents()) {
+ static_assert(
+ (_OtherExtents::rank() > 1) ||
+ (padded_stride_type::static_value() != dynamic_extent) ||
+ (_OtherExtents::static_extent(extent_to_pad_idx) != dynamic_extent) ||
+ (padded_stride_type::static_value() ==
+ _OtherExtents::static_extent(extent_to_pad_idx)));
+ }
+
+ /**
+ * Converting constructor from `layout_stride::mapping`.
+ *
+ * This overload participates in overload resolution only if
+ * `is_constructible_v<extents_type, OtherExtents>` is true
+ */
+ MDSPAN_TEMPLATE_REQUIRES(
+ class _OtherExtents,
+ /* requires */ (std::is_constructible_v<extents_type, _OtherExtents>))
+ MDSPAN_CONDITIONAL_EXPLICIT((extents_type::rank() > 0))
+ MDSPAN_INLINE_FUNCTION
+ constexpr mapping(const layout_stride::mapping<_OtherExtents> &other_mapping)
+ : padded_stride(padded_stride_type::init_padding(
+ other_mapping,
+ std::integral_constant<size_t, padded_stride_idx>{})),
+ exts(other_mapping.extents()) {}
+
+ /**
+ * Converting constructor from `layout_right_padded::mapping`.
+ *
+ * This overload participates in overload resolution only if
+ * `is_constructible_v<extents_type, OtherExtents>` is true. Either
+ * `padding_value` or `OtherPaddingStride` must be `std::dynamic_extent`, or
+ * `padding_value == OtherPaddingStride`.
+ */
+ MDSPAN_TEMPLATE_REQUIRES(
+ class _Mapping,
+ /* requires */ (detail::is_layout_right_padded_mapping<_Mapping>::value
+ &&std::is_constructible_v<
+ extents_type, typename _Mapping::extents_type>))
+ MDSPAN_CONDITIONAL_EXPLICIT((extents_type::rank() > 1 &&
+ (padding_value == dynamic_extent ||
+ _Mapping::padding_value == dynamic_extent)))
+ MDSPAN_INLINE_FUNCTION
+ constexpr mapping(const _Mapping &other_mapping)
+ : padded_stride(padded_stride_type::init_padding(
+ other_mapping,
+ std::integral_constant<size_t, padded_stride_idx>{})),
+ exts(other_mapping.extents()) {
+ static_assert(padding_value == dynamic_extent ||
+ _Mapping::padding_value == dynamic_extent ||
+ padding_value == _Mapping::padding_value);
+ }
+
+ /**
+ * Converting constructor from `layout_left_padded::mapping`.
+ *
+ * This overload participates in overload resolution only if
+ * `extents_type::rank()` is 0 or 1 and `is_constructible_v<extents_type,
+ * OtherExtents>` is `true`.
+ */
+ MDSPAN_TEMPLATE_REQUIRES(
+ class _Mapping,
+ /* requires */ (detail::is_layout_left_padded_mapping<_Mapping>::value
+ &&extents_type::rank() <= 1 &&
+ std::is_constructible_v<extents_type,
+ typename _Mapping::extents_type>))
+ MDSPAN_CONDITIONAL_EXPLICIT(
+ (!std::is_convertible_v<typename _Mapping::extents_type, extents_type>))
+ MDSPAN_INLINE_FUNCTION
+ constexpr mapping(const _Mapping &other_mapping) noexcept
+ : padded_stride(padded_stride_type::init_padding(
+ static_cast<extents_type>(other_mapping.extents()),
+ other_mapping.extents().extent(extent_to_pad_idx))),
+ exts(other_mapping.extents()) {}
+
+ MDSPAN_INLINE_FUNCTION constexpr const extents_type &
+ extents() const noexcept {
+ return exts;
+ }
+
+ constexpr std::array<index_type, extents_type::rank()>
+ strides() const noexcept {
+ if constexpr (extents_type::rank() == 0) {
+ return {};
+ } else if constexpr (extents_type::rank() == 1) {
+ return {1};
+ } else {
+ index_type value = 1;
+ std::array<index_type, extents_type::rank()> s{};
+ s[extent_to_pad_idx] = value;
+ value *= padded_stride.value(0);
+ for (rank_type r = extent_to_pad_idx - 1; r > 0; --r) {
+ s[r] = value;
+ value *= exts.extent(r);
+ }
+ s[0] = value;
+ return s;
+ }
+ }
+
+ MDSPAN_INLINE_FUNCTION constexpr index_type
+ required_span_size() const noexcept {
+ if constexpr (extents_type::rank() == 0) {
+ return 1;
+ } else if constexpr (extents_type::rank() == 1) {
+ return exts.extent(0);
+ } else {
+ index_type value = 1;
+ for (rank_type r = 0; r < extent_to_pad_idx; ++r) {
+ value *= exts.extent(r);
+ }
+ return value * padded_stride.value(0);
+ }
+ }
+
+ /**
+ * Return the mapping given the provided indices per rank.
+ *
+ * This overload participates in overload resolution only if:
+ * - `sizeof...(Indices) == extents_type::rank()`,
+ * - `(is_convertible_v<Indices, index_type> && ...) is true`, and
+ * - (is_nothrow_constructible_v<index_type, Indices> && ...) is true.
+ */
+ MDSPAN_TEMPLATE_REQUIRES(
+ class... _Indices,
+ /* requires */ (sizeof...(_Indices) == extents_type::rank() &&
+ (::MDSPAN_IMPL_STANDARD_NAMESPACE::detail::
+ are_valid_indices<index_type, _Indices...>())))
+ MDSPAN_INLINE_FUNCTION constexpr size_t
+ operator()(_Indices... idxs) const noexcept {
+ return compute_offset(std::index_sequence_for<_Indices...>{}, idxs...);
+ }
+
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_unique() noexcept {
+ return true;
+ }
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_exhaustive() noexcept {
+ return (extents_type::rank() <= rank_type(1)) ||
+ (extents_type::static_extent(extent_to_pad_idx) != dynamic_extent &&
+ extents_type::static_extent(extent_to_pad_idx) ==
+ padded_stride_type::static_value());
+ }
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_always_strided() noexcept {
+ return true;
+ }
+
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_unique() noexcept {
+ return true;
+ }
+ MDSPAN_INLINE_FUNCTION constexpr bool is_exhaustive() const noexcept {
+ return (extents_type::rank() < 2) ||
+ (exts.extent(extent_to_pad_idx) == padded_stride.value(0));
+ }
+ MDSPAN_INLINE_FUNCTION static constexpr bool is_strided() noexcept {
+ return true;
+ }
+
+ MDSPAN_INLINE_FUNCTION constexpr index_type
+ stride(rank_type r) const noexcept {
+ assert(r < extents_type::rank());
+ if (r == extents_type::rank() - 1)
+ return index_type(1);
+
+ index_type value = padded_stride.value(0);
+ for (rank_type k = extents_type::rank() - 2; k > r; k--)
+ value *= exts.extent(k);
+
+ return value;
+ }
+
+ /**
+ * Equality operator between `layout_right_padded`s
+ *
+ * This overload only participates in overload resolution if
+ * `OtherExtents::rank() == extents_type::rank()`.
+ *
+ * \note There is currently a difference from p2642r2, where this function is
+ * specified as taking `layout_right_padded< padding_value >::mapping<
+ * Extents>`. However, this makes `padding_value` non-deducible.
+ */
+ MDSPAN_TEMPLATE_REQUIRES(
+ class _Mapping,
+ /* requires */ (detail::is_layout_right_padded_mapping<_Mapping>::value &&
+ (_Mapping::extents_type::rank() == extents_type::rank())))
+ MDSPAN_INLINE_FUNCTION friend constexpr bool
+ operator==(const mapping &left, const _Mapping &right) noexcept {
+ // Workaround for some compilers not short-circuiting properly with
+ // compile-time checks i.e. we can't access stride(_padding_stride_idx) of a
+ // rank 0 mapping
+ bool strides_equal = true;
+ if constexpr (extents_type::rank() > rank_type(1)) {
+ strides_equal =
+ left.stride(padded_stride_idx) == right.stride(padded_stride_idx);
+ }
+ return (left.extents() == right.extents()) && strides_equal;
+ }
+
+#if !MDSPAN_HAS_CXX_20
+ /**
+ * Inequality operator between `layout_right_padded`s
+ *
+ * This overload only participates in overload resolution if
+ * `OtherExtents::rank() == extents_type::rank()`.
+ */
+ MDSPAN_TEMPLATE_REQUIRES(
+ class _Mapping,
+ /* requires */ (detail::is_layout_right_padded_mapping<_Mapping>::value &&
+ (_Mapping::extents_type::rank() == extents_type::rank())))
+ MDSPAN_INLINE_FUNCTION friend constexpr bool
+ operator!=(const mapping &left, const _Mapping &right) noexcept {
+ return !(left == right);
+ }
+#endif
+
+ // [mdspan.submdspan.mapping], submdspan mapping specialization
+ template<class... SliceSpecifiers>
+ MDSPAN_INLINE_FUNCTION
+ constexpr auto submdspan_mapping_impl(
+ SliceSpecifiers... slices) const;
+
+ template<class... SliceSpecifiers>
+ MDSPAN_INLINE_FUNCTION
+ friend constexpr auto submdspan_mapping(
+ const mapping& src, SliceSpecifiers... slices) {
+ return src.submdspan_mapping_impl(slices...);
+ }
+};
+}
+}
diff --git a/ext/mdspan/include/experimental/__p2642_bits/layout_padded_fwd.hpp b/ext/mdspan/include/experimental/__p2642_bits/layout_padded_fwd.hpp
new file mode 100644
index 0000000..3f141ff
--- /dev/null
+++ b/ext/mdspan/include/experimental/__p2642_bits/layout_padded_fwd.hpp
@@ -0,0 +1,137 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+//
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+#pragma once
+
+#include <cassert>
+#include "../__p0009_bits/dynamic_extent.hpp"
+#include "../__p0009_bits/utility.hpp"
+
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+namespace MDSPAN_IMPL_PROPOSED_NAMESPACE {
+
+template <size_t padding_value = dynamic_extent>
+struct layout_left_padded {
+ template <class _Extents>
+ class mapping;
+};
+
+template <size_t padding_value = dynamic_extent>
+struct layout_right_padded {
+ template <class _Extents>
+ class mapping;
+};
+
+namespace detail {
+// The layout_padded_constants structs are only useful if rank > 1, otherwise they may wrap
+template <class _Layout, class _ExtentsType>
+struct layout_padded_constants;
+
+template <class _ExtentsType, size_t _PaddingStride>
+struct layout_padded_constants<layout_left_padded<_PaddingStride>, _ExtentsType>
+{
+ using rank_type = typename _ExtentsType::rank_type;
+ static constexpr rank_type padded_stride_idx = 1;
+ static constexpr rank_type extent_to_pad_idx = 0;
+};
+
+template <class _ExtentsType, size_t _PaddingStride>
+struct layout_padded_constants<layout_right_padded<_PaddingStride>, _ExtentsType>
+{
+ using rank_type = typename _ExtentsType::rank_type;
+ static constexpr rank_type padded_stride_idx = _ExtentsType::rank() - 2;
+ static constexpr rank_type extent_to_pad_idx = _ExtentsType::rank() - 1;
+};
+
+template <class _Layout>
+struct is_layout_left_padded : std::false_type {};
+
+template <size_t _PaddingStride>
+struct is_layout_left_padded<layout_left_padded<_PaddingStride>> : std::true_type {};
+
+template <class _Mapping, class _Enabled = void>
+struct is_layout_left_padded_mapping : std::false_type {};
+
+template <class _Mapping>
+struct is_layout_left_padded_mapping<_Mapping,
+ std::enable_if_t<std::is_same<_Mapping, typename layout_left_padded<_Mapping::padding_value>::template mapping<typename _Mapping::extents_type>>::value>>
+ : std::true_type {};
+
+template <class _Layout>
+struct is_layout_right_padded : std::false_type {};
+
+template <size_t _PaddingStride>
+struct is_layout_right_padded<layout_right_padded<_PaddingStride>> : std::true_type {};
+
+template <class _Mapping, class _Enabled = void>
+struct is_layout_right_padded_mapping : std::false_type {};
+
+template <class _Mapping>
+struct is_layout_right_padded_mapping<_Mapping,
+ std::enable_if_t<std::is_same<_Mapping, typename layout_right_padded<_Mapping::padding_value>::template mapping<typename _Mapping::extents_type>>::value>>
+ : std::true_type {};
+
+
+template <class _LayoutExtentsType, class _PaddedLayoutMappingType>
+MDSPAN_INLINE_FUNCTION
+constexpr void check_padded_layout_converting_constructor_mandates(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::with_rank<0>) {}
+
+template <class _LayoutExtentsType, class _PaddedLayoutMappingType>
+MDSPAN_INLINE_FUNCTION
+constexpr void check_padded_layout_converting_constructor_mandates(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::with_rank<1>) {}
+
+template <class _LayoutExtentsType, class _PaddedLayoutMappingType, std::size_t N>
+MDSPAN_INLINE_FUNCTION
+constexpr void check_padded_layout_converting_constructor_mandates(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::with_rank<N>)
+{
+ using extents_type = typename _PaddedLayoutMappingType::extents_type;
+ constexpr auto padding_value = _PaddedLayoutMappingType::padding_value;
+ constexpr auto idx = layout_padded_constants<typename _PaddedLayoutMappingType::layout_type, _LayoutExtentsType >::extent_to_pad_idx;
+
+ constexpr auto statically_determinable =
+ (_LayoutExtentsType::static_extent(idx) != dynamic_extent) &&
+ (extents_type::static_extent(idx) != dynamic_extent) &&
+ (padding_value != dynamic_extent);
+
+ static_assert(!statically_determinable ||
+ (padding_value == 0
+ ? _LayoutExtentsType::static_extent(idx) == 0
+ : _LayoutExtentsType::static_extent(idx) % padding_value == 0),
+ "");
+}
+
+template <typename _ExtentsType, typename _OtherMapping>
+MDSPAN_INLINE_FUNCTION
+constexpr void check_padded_layout_converting_constructor_preconditions(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::with_rank<0>,
+ const _OtherMapping&) {}
+template <typename _ExtentsType, typename _OtherMapping>
+MDSPAN_INLINE_FUNCTION
+constexpr void check_padded_layout_converting_constructor_preconditions(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::with_rank<1>,
+ const _OtherMapping&) {}
+template <typename _ExtentsType, typename _OtherMapping, std::size_t N>
+MDSPAN_INLINE_FUNCTION
+constexpr void check_padded_layout_converting_constructor_preconditions(MDSPAN_IMPL_STANDARD_NAMESPACE::detail::with_rank<N>,
+ const _OtherMapping &other_mapping) {
+ constexpr auto padded_stride_idx =
+ layout_padded_constants<typename _OtherMapping::layout_type,
+ _ExtentsType>::padded_stride_idx;
+ constexpr auto extent_to_pad_idx = layout_padded_constants<typename _OtherMapping::layout_type, _ExtentsType>::extent_to_pad_idx;
+ MDSPAN_IMPL_PRECONDITION(other_mapping.stride(padded_stride_idx) == other_mapping.extents().extent(extent_to_pad_idx));
+}
+
+
+}
+}
+}
diff --git a/ext/mdspan/include/experimental/mdarray b/ext/mdspan/include/experimental/mdarray
new file mode 100644
index 0000000..642d1f5
--- /dev/null
+++ b/ext/mdspan/include/experimental/mdarray
@@ -0,0 +1,28 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+
+#pragma once
+
+#ifndef MDSPAN_IMPL_STANDARD_NAMESPACE
+ #define MDSPAN_IMPL_STANDARD_NAMESPACE std
+#endif
+
+#ifndef MDSPAN_IMPL_PROPOSED_NAMESPACE
+ #define MDSPAN_IMPL_PROPOSED_NAMESPACE experimental
+#endif
+
+#include "mdspan"
+#include "../mdspan/mdarray.hpp"
diff --git a/ext/mdspan/include/experimental/mdspan b/ext/mdspan/include/experimental/mdspan
new file mode 100644
index 0000000..e8ba715
--- /dev/null
+++ b/ext/mdspan/include/experimental/mdspan
@@ -0,0 +1,39 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+
+#pragma once
+
+#ifndef MDSPAN_IMPL_STANDARD_NAMESPACE
+ #define MDSPAN_IMPL_STANDARD_NAMESPACE std
+#endif
+
+#ifndef MDSPAN_IMPL_PROPOSED_NAMESPACE
+ #define MDSPAN_IMPL_PROPOSED_NAMESPACE experimental
+#endif
+
+#include "../mdspan/mdspan.hpp"
+
+// backward compatibility import into experimental
+namespace MDSPAN_IMPL_STANDARD_NAMESPACE {
+ namespace MDSPAN_IMPL_PROPOSED_NAMESPACE {
+ using ::MDSPAN_IMPL_STANDARD_NAMESPACE::mdspan;
+ using ::MDSPAN_IMPL_STANDARD_NAMESPACE::extents;
+ using ::MDSPAN_IMPL_STANDARD_NAMESPACE::layout_left;
+ using ::MDSPAN_IMPL_STANDARD_NAMESPACE::layout_right;
+ using ::MDSPAN_IMPL_STANDARD_NAMESPACE::layout_stride;
+ using ::MDSPAN_IMPL_STANDARD_NAMESPACE::default_accessor;
+ }
+}
diff --git a/ext/mdspan/include/mdspan/mdarray.hpp b/ext/mdspan/include/mdspan/mdarray.hpp
new file mode 100644
index 0000000..fd8f61c
--- /dev/null
+++ b/ext/mdspan/include/mdspan/mdarray.hpp
@@ -0,0 +1,31 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+
+#ifndef MDARRAY_HPP_
+#define MDARRAY_HPP_
+
+#ifndef MDSPAN_IMPL_STANDARD_NAMESPACE
+ #define MDSPAN_IMPL_STANDARD_NAMESPACE Kokkos
+#endif
+
+#ifndef MDSPAN_IMPL_PROPOSED_NAMESPACE
+ #define MDSPAN_IMPL_PROPOSED_NAMESPACE Experimental
+#endif
+
+#include "mdspan.hpp"
+#include "../experimental/__p1684_bits/mdarray.hpp"
+
+#endif // MDARRAY_HPP_
diff --git a/ext/mdspan/include/mdspan/mdspan.hpp b/ext/mdspan/include/mdspan/mdspan.hpp
new file mode 100644
index 0000000..4a0e354
--- /dev/null
+++ b/ext/mdspan/include/mdspan/mdspan.hpp
@@ -0,0 +1,43 @@
+//@HEADER
+// ************************************************************************
+//
+// Kokkos v. 4.0
+// Copyright (2022) National Technology & Engineering
+// Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+
+#ifndef MDSPAN_HPP_
+#define MDSPAN_HPP_
+
+#ifndef MDSPAN_IMPL_STANDARD_NAMESPACE
+ #define MDSPAN_IMPL_STANDARD_NAMESPACE Kokkos
+#endif
+
+#ifndef MDSPAN_IMPL_PROPOSED_NAMESPACE
+ #define MDSPAN_IMPL_PROPOSED_NAMESPACE Experimental
+#endif
+
+#include "../experimental/__p0009_bits/default_accessor.hpp"
+#include "../experimental/__p0009_bits/full_extent_t.hpp"
+#include "../experimental/__p0009_bits/mdspan.hpp"
+#include "../experimental/__p0009_bits/dynamic_extent.hpp"
+#include "../experimental/__p0009_bits/extents.hpp"
+#include "../experimental/__p0009_bits/layout_stride.hpp"
+#include "../experimental/__p0009_bits/layout_left.hpp"
+#include "../experimental/__p0009_bits/layout_right.hpp"
+#include "../experimental/__p0009_bits/macros.hpp"
+#if MDSPAN_HAS_CXX_17
+#include "../experimental/__p2642_bits/layout_padded.hpp"
+#include "../experimental/__p2630_bits/submdspan.hpp"
+#endif
+#include "../experimental/__p2389_bits/dims.hpp"
+
+#endif // MDSPAN_HPP_