From 4504bb539d69f1385732ea0c7a745121b0dfed05 Mon Sep 17 00:00:00 2001 From: Swiftb0y <12380386+Swiftb0y@users.noreply.github.com> Date: Wed, 3 May 2023 17:14:19 +0200 Subject: [PATCH] feat(lib): add Kokkos::mdspan library --- CMakeLists.txt | 5 + lib/mdspan/LICENSE | 238 +++++++ lib/mdspan/README.md | 82 +++ .../__p0009_bits/compressed_pair.hpp | 195 ++++++ .../experimental/__p0009_bits/config.hpp | 274 ++++++++ .../__p0009_bits/default_accessor.hpp | 56 ++ .../__p0009_bits/dynamic_extent.hpp | 35 + .../experimental/__p0009_bits/extents.hpp | 597 +++++++++++++++++ .../__p0009_bits/full_extent_t.hpp | 26 + .../experimental/__p0009_bits/layout_left.hpp | 222 ++++++ .../__p0009_bits/layout_right.hpp | 223 +++++++ .../__p0009_bits/layout_stride.hpp | 495 ++++++++++++++ .../experimental/__p0009_bits/macros.hpp | 631 ++++++++++++++++++ .../experimental/__p0009_bits/mdspan.hpp | 421 ++++++++++++ .../__p0009_bits/no_unique_address.hpp | 97 +++ .../__p0009_bits/trait_backports.hpp | 132 ++++ .../experimental/__p0009_bits/type_list.hpp | 87 +++ .../experimental/__p1684_bits/mdarray.hpp | 490 ++++++++++++++ .../__p2630_bits/strided_slice.hpp | 49 ++ .../experimental/__p2630_bits/submdspan.hpp | 42 ++ .../__p2630_bits/submdspan_extents.hpp | 323 +++++++++ .../__p2630_bits/submdspan_mapping.hpp | 299 +++++++++ lib/mdspan/include/experimental/mdarray | 28 + lib/mdspan/include/experimental/mdspan | 39 ++ lib/mdspan/include/mdspan/mdarray.hpp | 31 + lib/mdspan/include/mdspan/mdspan.hpp | 41 ++ 26 files changed, 5158 insertions(+) create mode 100644 lib/mdspan/LICENSE create mode 100644 lib/mdspan/README.md create mode 100644 lib/mdspan/include/experimental/__p0009_bits/compressed_pair.hpp create mode 100644 lib/mdspan/include/experimental/__p0009_bits/config.hpp create mode 100644 lib/mdspan/include/experimental/__p0009_bits/default_accessor.hpp create mode 100644 lib/mdspan/include/experimental/__p0009_bits/dynamic_extent.hpp create mode 100644 lib/mdspan/include/experimental/__p0009_bits/extents.hpp create mode 100644 lib/mdspan/include/experimental/__p0009_bits/full_extent_t.hpp create mode 100644 lib/mdspan/include/experimental/__p0009_bits/layout_left.hpp create mode 100644 lib/mdspan/include/experimental/__p0009_bits/layout_right.hpp create mode 100644 lib/mdspan/include/experimental/__p0009_bits/layout_stride.hpp create mode 100644 lib/mdspan/include/experimental/__p0009_bits/macros.hpp create mode 100644 lib/mdspan/include/experimental/__p0009_bits/mdspan.hpp create mode 100644 lib/mdspan/include/experimental/__p0009_bits/no_unique_address.hpp create mode 100644 lib/mdspan/include/experimental/__p0009_bits/trait_backports.hpp create mode 100644 lib/mdspan/include/experimental/__p0009_bits/type_list.hpp create mode 100644 lib/mdspan/include/experimental/__p1684_bits/mdarray.hpp create mode 100644 lib/mdspan/include/experimental/__p2630_bits/strided_slice.hpp create mode 100644 lib/mdspan/include/experimental/__p2630_bits/submdspan.hpp create mode 100644 lib/mdspan/include/experimental/__p2630_bits/submdspan_extents.hpp create mode 100644 lib/mdspan/include/experimental/__p2630_bits/submdspan_mapping.hpp create mode 100644 lib/mdspan/include/experimental/mdarray create mode 100644 lib/mdspan/include/experimental/mdspan create mode 100644 lib/mdspan/include/mdspan/mdarray.hpp create mode 100644 lib/mdspan/include/mdspan/mdspan.hpp diff --git a/CMakeLists.txt b/CMakeLists.txt index b6b34db0abf..a5a00d6667b 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -2363,6 +2363,11 @@ target_link_libraries(mixxx-lib PUBLIC mixxx-proto) # https://github.com/rigtorp/SPSCQueue target_include_directories(mixxx-lib SYSTEM PUBLIC lib/rigtorp/SPSCQueue/include) +# Kokkos mdspan +# https://github.com/kokkos/mdspan +# TODO: obsolete after adopting C++23 +target_include_directories(mixxx-lib SYSTEM PUBLIC lib/mdspan/include) + # Qt if(QT6) find_package(QT 6.2 NAMES Qt6 COMPONENTS Core REQUIRED) diff --git a/lib/mdspan/LICENSE b/lib/mdspan/LICENSE new file mode 100644 index 00000000000..6572cc2db05 --- /dev/null +++ b/lib/mdspan/LICENSE @@ -0,0 +1,238 @@ + ************************************************************************ + + 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. + + + ============================================================================== + Kokkos is under the Apache License v2.0 with LLVM Exceptions: + ============================================================================== + + Apache License + Version 2.0, January 2004 + http://www.apache.org/licenses/ + + TERMS AND CONDITIONS FOR USE, REPRODUCTION, AND DISTRIBUTION + + 1. Definitions. + + "License" shall mean the terms and conditions for use, reproduction, + and distribution as defined by Sections 1 through 9 of this document. + + "Licensor" shall mean the copyright owner or entity authorized by + the copyright owner that is granting the License. + + "Legal Entity" shall mean the union of the acting entity and all + other entities that control, are controlled by, or are under common + control with that entity. For the purposes of this definition, + "control" means (i) the power, direct or indirect, to cause the + direction or management of such entity, whether by contract or + otherwise, or (ii) ownership of fifty percent (50%) or more of the + outstanding shares, or (iii) beneficial ownership of such entity. + + "You" (or "Your") shall mean an individual or Legal Entity + exercising permissions granted by this License. + + "Source" form shall mean the preferred form for making modifications, + including but not limited to software source code, documentation + source, and configuration files. + + "Object" form shall mean any form resulting from mechanical + transformation or translation of a Source form, including but + not limited to compiled object code, generated documentation, + and conversions to other media types. + + "Work" shall mean the work of authorship, whether in Source or + Object form, made available under the License, as indicated by a + copyright notice that is included in or attached to the work + (an example is provided in the Appendix below). + + "Derivative Works" shall mean any work, whether in Source or Object + form, that is based on (or derived from) the Work and for which the + editorial revisions, annotations, elaborations, or other modifications + represent, as a whole, an original work of authorship. For the purposes + of this License, Derivative Works shall not include works that remain + separable from, or merely link (or bind by name) to the interfaces of, + the Work and Derivative Works thereof. + + "Contribution" shall mean any work of authorship, including + the original version of the Work and any modifications or additions + to that Work or Derivative Works thereof, that is intentionally + submitted to Licensor for inclusion in the Work by the copyright owner + or by an individual or Legal Entity authorized to submit on behalf of + the copyright owner. For the purposes of this definition, "submitted" + means any form of electronic, verbal, or written communication sent + to the Licensor or its representatives, including but not limited to + communication on electronic mailing lists, source code control systems, + and issue tracking systems that are managed by, or on behalf of, the + Licensor for the purpose of discussing and improving the Work, but + excluding communication that is conspicuously marked or otherwise + designated in writing by the copyright owner as "Not a Contribution." + + "Contributor" shall mean Licensor and any individual or Legal Entity + on behalf of whom a Contribution has been received by Licensor and + subsequently incorporated within the Work. + + 2. Grant of Copyright License. Subject to the terms and conditions of + this License, each Contributor hereby grants to You a perpetual, + worldwide, non-exclusive, no-charge, royalty-free, irrevocable + copyright license to reproduce, prepare Derivative Works of, + publicly display, publicly perform, sublicense, and distribute the + Work and such Derivative Works in Source or Object form. + + 3. Grant of Patent License. Subject to the terms and conditions of + this License, each Contributor hereby grants to You a perpetual, + worldwide, non-exclusive, no-charge, royalty-free, irrevocable + (except as stated in this section) patent license to make, have made, + use, offer to sell, sell, import, and otherwise transfer the Work, + where such license applies only to those patent claims licensable + by such Contributor that are necessarily infringed by their + Contribution(s) alone or by combination of their Contribution(s) + with the Work to which such Contribution(s) was submitted. If You + institute patent litigation against any entity (including a + cross-claim or counterclaim in a lawsuit) alleging that the Work + or a Contribution incorporated within the Work constitutes direct + or contributory patent infringement, then any patent licenses + granted to You under this License for that Work shall terminate + as of the date such litigation is filed. + + 4. Redistribution. You may reproduce and distribute copies of the + Work or Derivative Works thereof in any medium, with or without + modifications, and in Source or Object form, provided that You + meet the following conditions: + + (a) You must give any other recipients of the Work or + Derivative Works a copy of this License; and + + (b) You must cause any modified files to carry prominent notices + stating that You changed the files; and + + (c) You must retain, in the Source form of any Derivative Works + that You distribute, all copyright, patent, trademark, and + attribution notices from the Source form of the Work, + excluding those notices that do not pertain to any part of + the Derivative Works; and + + (d) If the Work includes a "NOTICE" text file as part of its + distribution, then any Derivative Works that You distribute must + include a readable copy of the attribution notices contained + within such NOTICE file, excluding those notices that do not + pertain to any part of the Derivative Works, in at least one + of the following places: within a NOTICE text file distributed + as part of the Derivative Works; within the Source form or + documentation, if provided along with the Derivative Works; or, + within a display generated by the Derivative Works, if and + wherever such third-party notices normally appear. The contents + of the NOTICE file are for informational purposes only and + do not modify the License. You may add Your own attribution + notices within Derivative Works that You distribute, alongside + or as an addendum to the NOTICE text from the Work, provided + that such additional attribution notices cannot be construed + as modifying the License. + + You may add Your own copyright statement to Your modifications and + may provide additional or different license terms and conditions + for use, reproduction, or distribution of Your modifications, or + for any such Derivative Works as a whole, provided Your use, + reproduction, and distribution of the Work otherwise complies with + the conditions stated in this License. + + 5. Submission of Contributions. Unless You explicitly state otherwise, + any Contribution intentionally submitted for inclusion in the Work + by You to the Licensor shall be under the terms and conditions of + this License, without any additional terms or conditions. + Notwithstanding the above, nothing herein shall supersede or modify + the terms of any separate license agreement you may have executed + with Licensor regarding such Contributions. + + 6. Trademarks. This License does not grant permission to use the trade + names, trademarks, service marks, or product names of the Licensor, + except as required for reasonable and customary use in describing the + origin of the Work and reproducing the content of the NOTICE file. + + 7. Disclaimer of Warranty. Unless required by applicable law or + agreed to in writing, Licensor provides the Work (and each + Contributor provides its Contributions) on an "AS IS" BASIS, + WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or + implied, including, without limitation, any warranties or conditions + of TITLE, NON-INFRINGEMENT, MERCHANTABILITY, or FITNESS FOR A + PARTICULAR PURPOSE. You are solely responsible for determining the + appropriateness of using or redistributing the Work and assume any + risks associated with Your exercise of permissions under this License. + + 8. Limitation of Liability. In no event and under no legal theory, + whether in tort (including negligence), contract, or otherwise, + unless required by applicable law (such as deliberate and grossly + negligent acts) or agreed to in writing, shall any Contributor be + liable to You for damages, including any direct, indirect, special, + incidental, or consequential damages of any character arising as a + result of this License or out of the use or inability to use the + Work (including but not limited to damages for loss of goodwill, + work stoppage, computer failure or malfunction, or any and all + other commercial damages or losses), even if such Contributor + has been advised of the possibility of such damages. + + 9. Accepting Warranty or Additional Liability. While redistributing + the Work or Derivative Works thereof, You may choose to offer, + and charge a fee for, acceptance of support, warranty, indemnity, + or other liability obligations and/or rights consistent with this + License. However, in accepting such obligations, You may act only + on Your own behalf and on Your sole responsibility, not on behalf + of any other Contributor, and only if You agree to indemnify, + defend, and hold each Contributor harmless for any liability + incurred by, or claims asserted against, such Contributor by reason + of your accepting any such warranty or additional liability. + + END OF TERMS AND CONDITIONS Apache 2.0 + + ---- LLVM Exceptions to the Apache 2.0 License ---- + + As an exception, if, as a result of your compiling your source code, portions + of this Software are embedded into an Object form of such source code, you + may redistribute such embedded portions in such Object form without complying + with the conditions of Sections 4(a), 4(b) and 4(d) of the License. + + In addition, if you combine or link compiled forms of this Software with + software that is licensed under the GPLv2 ("Combined Software") and if a + court of competent jurisdiction determines that the patent provision (Section + 3), the indemnity provision (Section 9) or other Section of the License + conflicts with the conditions of the GPLv2, you may retroactively and + prospectively choose to deem waived or otherwise exclude such Section(s) of + the License, but only in their entirety and only with respect to the Combined + Software. + + ============================================================================== + Software from third parties included in Kokkos: + ============================================================================== + + Kokkos contains third party software which is under different license + terms. All such code will be identified clearly using at least one of two + mechanisms: + 1) It will be in a separate directory tree with its own `LICENSE.txt` or + `LICENSE` file at the top containing the specific license and restrictions + which apply to that software, or + 2) It will contain specific license and restriction terms at the top of every + file. + + + THIS SOFTWARE IS PROVIDED BY NTESS "AS IS" AND ANY + EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR + PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NTESS OR THE + CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, + EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, + PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR + PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF + LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING + NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS + SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + + Questions? Contact: + Christian R. Trott (crtrott@sandia.gov) and + Damien T. Lebrun-Grandie (lebrungrandt@ornl.gov) + + ************************************************************************ diff --git a/lib/mdspan/README.md b/lib/mdspan/README.md new file mode 100644 index 00000000000..7c4af5ed8a8 --- /dev/null +++ b/lib/mdspan/README.md @@ -0,0 +1,82 @@ +Reference `mdspan` implementation +========================================== + +The ISO-C++ proposal [P0009](https://wg21.link/p0009) will add support for non-owning multi-dimensional array references to the C++ standard library. This repository aims to provide a production-quality implementation of the proposal as written (with a few caveats, see below) in preparation for the addition of `mdspan` to the standard. Please feel free to use this, file bugs when it breaks, and let us know how it works for you :-) + +[Try it out on Godbolt](https://godbolt.org/z/Mxa7cej1a){: .btn } + +During the final leg of the ISO C++ committee review process a number of breaking changes were proposed and accepted (issue #136). These are now merged into the stable branch. + +Note: There is a tag mdspan-0.3.0 which reflects the status of P0009 before R17 - i.e. it does not have the integral type template parameter for `extents`. +Note: There is a tag mdspan-0.4.0 which reflects the status of P0009 before + +* renaming `pointer`, `data`, `is_contiguous` and `is_always_contiguous`; and before +* renaming `size_type` to `index_type` and introducing a new `size_type = make_unsigned_t` alias. + +Using `mdspan` +-------------- + +A [tutorial-style introduction](https://github.com/kokkos/mdspan/wiki/A-Gentle-Introduction-to-mdspan) to the basic usage of `mdspan` is provided on the project wiki. More advanced tutorials to come. + +Features in Addition To C++ Standard +------------------------------------ + +- C++17 backport (e.g., concepts not required) +- C++14 backport (e.g., fold expressions not required) + - Compile times of this backport will be substantially slower than the C++17 version +- Macros to enable, e.g., `__device__` marking of all functions for CUDA compatibility + +Building and Installation +------------------------- + +This implementation is header-only, with compiler features detected using feature test macros, so you can just use it directly with no building or installation. If you would like to run the included tests or benchmarks, you'll need CMake. + +### Running tests + +#### Configurations + +- clang-15 / cmake 3.23 + - Warning free with `-Wall -Wextra -pedantic` for C++23/20. In C++17 pedantic will give a few warnings, in C++14 Wextra will also give some. + - `cmake -DMDSPAN_ENABLE_TESTS=ON -DMDSPAN_ENABLE_BENCHMARKS=ON -DCMAKE_CXX_FLAGS="-Werror -Wall -Wextra -pedantic" -DCMAKE_CXX_STANDARD=23 -DMDSPAN_CXX_STANDARD=23 -DCMAKE_CXX_COMPILER=clang++` +- gcc-11 / cmake 3.23 + - Warning free with `-Wall -Wextra -pedantic` for C++23/20. In C++17 and C++14 pedantic will give a warning (note only with `CMAKE_CXX_EXTENSION=OFF`). + - `cmake -DMDSPAN_ENABLE_TESTS=ON -DMDSPAN_ENABLE_BENCHMARKS=ON -DCMAKE_CXX_FLAGS="-Werror -Wall -Wextra -pedantic" -DCMAKE_CXX_STANDARD=17 -DMDSPAN_CXX_STANDARD=17 -DCMAKE_CXX_COMPILER=g++ -DCMAKE_CXX_EXTENSIONS=OFF` +- CUDA 11.x / gcc 9.1 / cmake 3.23 + - has a few warnings in C++17 mostly in `mdarray` due to the use of `vector` as container for the tests. + - Note with CUDA 11.7 and GCC 11.1 as host compiler we observe some issues around CTAD, some of the layout tests won't compile + - CUDA 11.7 with GCC 9.1 works however + - `cmake -DMDSPAN_ENABLE_TESTS=ON -DMDSPAN_ENABLE_CUDA=ON -DMDSPAN_ENABLE_BENCHMARKS=ON -DCMAKE_CXX_STANDARD=17 -DCMAKE_CUDA_ARCHITECTURES=70 -DMDSPAN_CXX_STANDARD=17 -DCMAKE_CUDA_FLAGS="--expt-relaxed-constexpr --extended-lambda"` + + +### Running benchmarks + +TODO write this + +Caveats +------- + +This implementation is fully conforming with the version of `mdspan` voted into the C++23 draft standard in July 2022. +When not in C++23 mode the implementation deviates from the proposal as follows: + +### C++20 +- implements `operator()` not `operator[]` + - note you can control which operator is available with defining `MDSPAN_USE_BRACKET_OPERATOR=[0,1]` and `MDSPAN_USE_PAREN_OPERATOR=[0,1]` irrespective of whether multi dimensional subscript support is detected. + +### C++17 +- mdspan has a default constructor even in cases where it shouldn't (i.e. all static extents, and default constructible mapping/accessor) +- the conditional explicit markup is missing, making certain constructors implicit + - most notably you can implicitly convert from dynamic extent to static extent, which you can't in C++20 mode +- there is a constraint on `layout_left::mapping::stride()`, `layout_right::mapping::stride()` and `layout_stride::mapping::stride()` that `extents_type::rank() > 0` is `true`, which is not implemented in C++17 or C++14. + +### C++14 +- deduction guides don't exist +- submdspan (P2630) is not available - an earlier variant of submdspan is available up to release 0.5 in C++14 mode +- benchmarks are not available (they need submdspan) + + + +Acknowledgements +================ + +This work was undertaken as part of the [Kokkos project](https://github.com/kokkos/kokkos) at Sandia National Laboratories. Sandia National Laboratories is a multimission laboratory managed and operated by National Technology & Engineering Solutions of Sandia, LLC, a wholly owned subsidary of Honeywell International Inc., for the U. S. Department of Energy's National Nuclear Security Administration under contract DE-NA0003525. + diff --git a/lib/mdspan/include/experimental/__p0009_bits/compressed_pair.hpp b/lib/mdspan/include/experimental/__p0009_bits/compressed_pair.hpp new file mode 100644 index 00000000000..541e57876c4 --- /dev/null +++ b/lib/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 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() noexcept = default; + MDSPAN_INLINE_FUNCTION_DEFAULTED + constexpr __compressed_pair(__compressed_pair const &) noexcept = default; + MDSPAN_INLINE_FUNCTION_DEFAULTED + constexpr __compressed_pair(__compressed_pair &&) noexcept = default; + MDSPAN_INLINE_FUNCTION_DEFAULTED + _MDSPAN_CONSTEXPR_14_DEFAULTED __compressed_pair & + operator=(__compressed_pair const &) noexcept = default; + MDSPAN_INLINE_FUNCTION_DEFAULTED + _MDSPAN_CONSTEXPR_14_DEFAULTED __compressed_pair & + operator=(__compressed_pair &&) noexcept = default; + MDSPAN_INLINE_FUNCTION_DEFAULTED + ~__compressed_pair() noexcept = default; + template + 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 +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() noexcept = default; + MDSPAN_INLINE_FUNCTION_DEFAULTED + constexpr __compressed_pair(__compressed_pair const &) noexcept = default; + MDSPAN_INLINE_FUNCTION_DEFAULTED + constexpr __compressed_pair(__compressed_pair &&) noexcept = default; + MDSPAN_INLINE_FUNCTION_DEFAULTED + _MDSPAN_CONSTEXPR_14_DEFAULTED __compressed_pair & + operator=(__compressed_pair const &) noexcept = default; + MDSPAN_INLINE_FUNCTION_DEFAULTED + _MDSPAN_CONSTEXPR_14_DEFAULTED __compressed_pair & + operator=(__compressed_pair &&) noexcept = default; + MDSPAN_INLINE_FUNCTION_DEFAULTED + ~__compressed_pair() noexcept = default; + template + MDSPAN_INLINE_FUNCTION constexpr __compressed_pair(_T1Like &&__t1, _T2Like &&__t2) + : _T1((_T1Like &&) __t1), __t2_val((_T2Like &&) __t2) {} +}; + +// Second empty. +template +struct __compressed_pair< + _T1, _T2, + std::enable_if_t> + : 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() noexcept = default; + MDSPAN_INLINE_FUNCTION_DEFAULTED + constexpr __compressed_pair(__compressed_pair const &) noexcept = default; + MDSPAN_INLINE_FUNCTION_DEFAULTED + constexpr __compressed_pair(__compressed_pair &&) noexcept = default; + MDSPAN_INLINE_FUNCTION_DEFAULTED + _MDSPAN_CONSTEXPR_14_DEFAULTED __compressed_pair & + operator=(__compressed_pair const &) noexcept = default; + MDSPAN_INLINE_FUNCTION_DEFAULTED + _MDSPAN_CONSTEXPR_14_DEFAULTED __compressed_pair & + operator=(__compressed_pair &&) noexcept = default; + MDSPAN_INLINE_FUNCTION_DEFAULTED + ~__compressed_pair() noexcept = default; + + template + MDSPAN_INLINE_FUNCTION constexpr __compressed_pair(_T1Like &&__t1, _T2Like &&__t2) + : _T2((_T2Like &&) __t2), __t1_val((_T1Like &&) __t1) {} +}; + +// Both empty. +template +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() noexcept = default; + MDSPAN_INLINE_FUNCTION_DEFAULTED + constexpr __compressed_pair(__compressed_pair const &) noexcept = default; + MDSPAN_INLINE_FUNCTION_DEFAULTED + constexpr __compressed_pair(__compressed_pair &&) noexcept = default; + MDSPAN_INLINE_FUNCTION_DEFAULTED + _MDSPAN_CONSTEXPR_14_DEFAULTED __compressed_pair & + operator=(__compressed_pair const &) noexcept = default; + MDSPAN_INLINE_FUNCTION_DEFAULTED + _MDSPAN_CONSTEXPR_14_DEFAULTED __compressed_pair & + operator=(__compressed_pair &&) noexcept = default; + MDSPAN_INLINE_FUNCTION_DEFAULTED + ~__compressed_pair() noexcept = default; + template + 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/lib/mdspan/include/experimental/__p0009_bits/config.hpp b/lib/mdspan/include/experimental/__p0009_bits/config.hpp new file mode 100644 index 00000000000..d35e201cebd --- /dev/null +++ b/lib/mdspan/include/experimental/__p0009_bits/config.hpp @@ -0,0 +1,274 @@ +//@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() +# include +#else +# include +# include +#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 + +#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) + +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 +#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__ >= 11 && __CUDACC_VER_MINOR__ >= 7)) && \ + ((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 && !defined(_MDSPAN_COMPILER_MSVC) +# 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) +# define MDSPAN_USE_BRACKET_OPERATOR 1 +# 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/lib/mdspan/include/experimental/__p0009_bits/default_accessor.hpp b/lib/mdspan/include/experimental/__p0009_bits/default_accessor.hpp new file mode 100644 index 00000000000..ea0f537b2fe --- /dev/null +++ b/lib/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 // size_t + +namespace MDSPAN_IMPL_STANDARD_NAMESPACE { + +template +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) 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/lib/mdspan/include/experimental/__p0009_bits/dynamic_extent.hpp b/lib/mdspan/include/experimental/__p0009_bits/dynamic_extent.hpp new file mode 100644 index 00000000000..2e29da13d6a --- /dev/null +++ b/lib/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 +#endif + +#include // size_t +#include // 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::max(); +#endif +} // namespace MDSPAN_IMPL_STANDARD_NAMESPACE + +//============================================================================================================== diff --git a/lib/mdspan/include/experimental/__p0009_bits/extents.hpp b/lib/mdspan/include/experimental/__p0009_bits/extents.hpp new file mode 100644 index 00000000000..e3ef66bd98b --- /dev/null +++ b/lib/mdspan/include/experimental/__p0009_bits/extents.hpp @@ -0,0 +1,597 @@ +//@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" + +#ifdef __cpp_lib_span +#include +#endif +#include + +#include + +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 +static constexpr std::integral_constant __check_compatible_extents( + std::integral_constant, + std::integer_sequence, + std::integer_sequence) noexcept { + return {}; +} + +// This helper prevents ICE's on MSVC. +template +struct __compare_extent_compatible : std::integral_constant +{}; + +template +static constexpr std::integral_constant< + bool, _MDSPAN_FOLD_AND(__compare_extent_compatible::value)> +__check_compatible_extents( + std::integral_constant, + std::integer_sequence, + std::integer_sequence) noexcept { + return {}; +} + +// ------------------------------------------------------------------ +// ------------ 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 struct static_array_impl; + +template +struct static_array_impl { + MDSPAN_INLINE_FUNCTION + constexpr static T get(size_t r) { + if (r == R) + return FirstExt; + else + return static_array_impl::get(r); + } + template MDSPAN_INLINE_FUNCTION constexpr static T get() { +#if MDSPAN_HAS_CXX_17 + if constexpr (r == R) + return FirstExt; + else + return static_array_impl::template get(); +#else + get(r); +#endif + } +}; + +// End the recursion +template +struct static_array_impl { + MDSPAN_INLINE_FUNCTION + constexpr static T get(size_t) { return FirstExt; } + template MDSPAN_INLINE_FUNCTION constexpr static T get() { + return FirstExt; + } +}; + +// Don't start recursion if size 0 +template struct static_array_impl<0, T> { + MDSPAN_INLINE_FUNCTION + constexpr static T get(size_t) { return T(); } + template MDSPAN_INLINE_FUNCTION constexpr static T get() { + return T(); + } +}; + +// Static array, provides get(), get(r) and operator[r] +template 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() which return the sum of the first r-1 values. + +// Recursive implementation for get +template struct index_sequence_scan_impl; + +template +struct index_sequence_scan_impl { + MDSPAN_INLINE_FUNCTION + constexpr static size_t get(size_t r) { + if (r > R) + return FirstVal + index_sequence_scan_impl::get(r); + else + return 0; + } +}; + +template +struct index_sequence_scan_impl { +#if defined(__NVCC__) || defined(__NVCOMPILER) + // 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(R) > static_cast(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 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 struct possibly_empty_array { + 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 +struct maybe_static_array { + + static_assert(std::is_convertible::value, "maybe_static_array: TStatic must be convertible to TDynamic"); + static_assert(std::is_convertible::value, "maybe_static_array: TDynamic must be convertible to TStatic"); + +private: + // Static values member + using static_vals_t = static_array; + 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 + 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(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(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 &vals) { + for (size_t r = 0; r < N; r++) + m_dyn_vals[r] = static_cast(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 &) : m_dyn_vals{} {} + +#ifdef __cpp_lib_span + MDSPAN_TEMPLATE_REQUIRES(class T, size_t N, + /* requires */ (N == m_size_dynamic)) + MDSPAN_INLINE_FUNCTION + constexpr maybe_static_array(const std::span &vals) { + for (size_t r = 0; r < N; r++) + m_dyn_vals[r] = static_cast(vals[r]); + } +#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(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(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 &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(vals[r]); + } +// Precondition check +#ifdef _MDSPAN_DEBUG + else { + assert(static_cast(vals[r]) == + static_cast(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 &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(vals[r]); + } +#ifdef _MDSPAN_DEBUG + else { + assert(static_cast(vals[r]) == + static_cast(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(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 extents { +public: + // typedefs for integral types used + using index_type = IndexType; + using size_type = std::make_unsigned_t; + using rank_type = size_t; + + static_assert(std::is_integral::value && !std::is_same::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; + _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(dynvals)...) {} + + MDSPAN_TEMPLATE_REQUIRES( + class OtherIndexType, size_t N, + /* requires */ + ( + _MDSPAN_TRAIT(std::is_convertible, OtherIndexType, index_type) && + _MDSPAN_TRAIT(std::is_nothrow_constructible, index_type, + OtherIndexType) && + (N == m_rank || N == m_rank_dynamic))) + MDSPAN_INLINE_FUNCTION + MDSPAN_CONDITIONAL_EXPLICIT(N != m_rank_dynamic) + constexpr extents(const std::array &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, OtherIndexType, index_type) && + _MDSPAN_TRAIT(std::is_nothrow_constructible, index_type, OtherIndexType) && + (N == m_rank || N == m_rank_dynamic))) + MDSPAN_INLINE_FUNCTION + MDSPAN_CONDITIONAL_EXPLICIT(N != m_rank_dynamic) + constexpr extents(const std::span &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, + std::integral_constant, + const OtherExtents &exts, + DynamicValues... dynamic_values) noexcept { + return __construct_vals_from_extents( + std::integral_constant(), + std::integral_constant(), 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, + std::integral_constant, + const OtherExtents &exts, + DynamicValues... dynamic_values) noexcept { + return __construct_vals_from_extents( + std::integral_constant(), + std::integral_constant(), 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, + std::integral_constant, + const OtherExtents &, + DynamicValues... dynamic_values) noexcept { + return vals_t{static_cast(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( + std::integral_constant{}, + std::integer_sequence{}, + std::integer_sequence{}))::value)) + MDSPAN_INLINE_FUNCTION + MDSPAN_CONDITIONAL_EXPLICIT((((Extents != dynamic_extent) && + (OtherExtents == dynamic_extent)) || + ...) || + (std::numeric_limits::max() < + std::numeric_limits::max())) + constexpr extents(const extents &other) noexcept + : m_vals(__construct_vals_from_extents( + std::integral_constant(), + std::integral_constant(), other)) {} + + // Comparison operator + template + MDSPAN_INLINE_FUNCTION friend constexpr bool + operator==(const extents &lhs, + const extents &rhs) noexcept { + bool value = true; + for (size_type r = 0; r < m_rank; r++) + value &= rhs.extent(r) == lhs.extent(r); + return value; + } + +#if !(MDSPAN_HAS_CXX_20) + template + MDSPAN_INLINE_FUNCTION friend constexpr bool + operator!=(extents const &lhs, + extents const &rhs) noexcept { + return !(lhs == rhs); + } +#endif +}; + +// Recursive helper classes to implement dextents alias for extents +namespace detail { + +template > +struct __make_dextents; + +template +struct __make_dextents< + IndexType, Rank, ::MDSPAN_IMPL_STANDARD_NAMESPACE::extents> +{ + using type = typename __make_dextents< + IndexType, Rank - 1, + ::MDSPAN_IMPL_STANDARD_NAMESPACE::extents>::type; +}; + +template +struct __make_dextents< + IndexType, 0, ::MDSPAN_IMPL_STANDARD_NAMESPACE::extents> +{ + using type = ::MDSPAN_IMPL_STANDARD_NAMESPACE::extents; +}; + +} // end namespace detail + +// [mdspan.extents.dextents], alias template +template +using dextents = typename detail::__make_dextents::type; + +// Deduction guide for extents +#if defined(_MDSPAN_USE_CLASS_TEMPLATE_ARGUMENT_DEDUCTION) +template +extents(IndexTypes...) + -> extents; +#endif + +// Helper type traits for identifying a class as extents. +namespace detail { + +template struct __is_extents : ::std::false_type {}; + +template +struct __is_extents<::MDSPAN_IMPL_STANDARD_NAMESPACE::extents> + : ::std::true_type {}; + +template +#if MDSPAN_HAS_CXX_17 +inline +#else +static +#endif +constexpr bool __is_extents_v = __is_extents::value; + +} // namespace detail +} // namespace MDSPAN_IMPL_STANDARD_NAMESPACE diff --git a/lib/mdspan/include/experimental/__p0009_bits/full_extent_t.hpp b/lib/mdspan/include/experimental/__p0009_bits/full_extent_t.hpp new file mode 100644 index 00000000000..bd4b5c6a8ba --- /dev/null +++ b/lib/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/lib/mdspan/include/experimental/__p0009_bits/layout_left.hpp b/lib/mdspan/include/experimental/__p0009_bits/layout_left.hpp new file mode 100644 index 00000000000..af44494a98d --- /dev/null +++ b/lib/mdspan/include/experimental/__p0009_bits/layout_left.hpp @@ -0,0 +1,222 @@ +//@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" + +namespace MDSPAN_IMPL_STANDARD_NAMESPACE { + +//============================================================================== + +template +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, + MDSPAN_IMPL_STANDARD_NAMESPACE_STRING "::layout_left::mapping must be instantiated with a specialization of " MDSPAN_IMPL_STANDARD_NAMESPACE_STRING "::extents."); + + template + friend class mapping; + + // i0+(i1 + E(1)*(i2 + E(2)*i3)) + template + struct __rank_count {}; + + template + _MDSPAN_HOST_DEVICE + constexpr index_type __compute_offset( + __rank_count, const I& i, Indices... idx) const { + return __compute_offset(__rank_count(), idx...) * + __extents.extent(r) + i; + } + + template + _MDSPAN_HOST_DEVICE + constexpr index_type __compute_offset( + __rank_count, 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::value)) // needs two () due to comma + MDSPAN_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 + mapping(mapping 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::value)) // needs two () due to comma + MDSPAN_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 + mapping(layout_right::mapping 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) + ) + ) + MDSPAN_CONDITIONAL_EXPLICIT((extents_type::rank() > 0)) + MDSPAN_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 + mapping(layout_stride::mapping 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 !defined(_MDSPAN_HAS_CUDA) && !defined(_MDSPAN_HAS_HIP) && !defined(NDEBUG) + index_type stride = 1; + for(rank_type r=0; r<__extents.rank(); r++) { + if(stride != static_cast(other.stride(r))) { + // Note this throw will lead to a terminate if triggered since this function is marked noexcept + throw std::runtime_error("Assigning layout_stride to layout_left with invalid strides."); + } + stride *= __extents.extent(r); + } + #endif + } + + 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(), static_cast(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 constexpr bool is_unique() const noexcept { return true; } + MDSPAN_INLINE_FUNCTION constexpr bool is_exhaustive() const noexcept { return true; } + MDSPAN_INLINE_FUNCTION constexpr bool is_strided() const 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 + MDSPAN_INLINE_FUNCTION + friend constexpr bool operator==(mapping const& lhs, mapping const& rhs) noexcept { + return lhs.extents() == rhs.extents(); + } + + // In C++ 20 the not equal exists if equal is found +#if !(MDSPAN_HAS_CXX_20) + template + MDSPAN_INLINE_FUNCTION + friend constexpr bool operator!=(mapping const& lhs, mapping const& rhs) noexcept { + return lhs.extents() != rhs.extents(); + } +#endif + + // Not really public, but currently needed to implement fully constexpr useable submdspan: + template + constexpr index_type __get_stride(MDSPAN_IMPL_STANDARD_NAMESPACE::extents,std::integer_sequence) const { + return _MDSPAN_FOLD_TIMES_RIGHT((Idx():1),1); + } + template + constexpr index_type __stride() const noexcept { + return __get_stride(__extents, std::make_index_sequence()); + } + +private: + _MDSPAN_NO_UNIQUE_ADDRESS extents_type __extents{}; + +}; + + +} // end namespace MDSPAN_IMPL_STANDARD_NAMESPACE + diff --git a/lib/mdspan/include/experimental/__p0009_bits/layout_right.hpp b/lib/mdspan/include/experimental/__p0009_bits/layout_right.hpp new file mode 100644 index 00000000000..a0586484202 --- /dev/null +++ b/lib/mdspan/include/experimental/__p0009_bits/layout_right.hpp @@ -0,0 +1,223 @@ +//@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 +#include "layout_stride.hpp" + +namespace MDSPAN_IMPL_STANDARD_NAMESPACE { + +//============================================================================== +template +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, + MDSPAN_IMPL_STANDARD_NAMESPACE_STRING "::layout_right::mapping must be instantiated with a specialization of " MDSPAN_IMPL_STANDARD_NAMESPACE_STRING "::extents."); + + template + friend class mapping; + + // i0+(i1 + E(1)*(i2 + E(2)*i3)) + template + struct __rank_count {}; + + template + _MDSPAN_HOST_DEVICE + constexpr index_type __compute_offset( + index_type offset, __rank_count, const I& i, Indices... idx) const { + return __compute_offset(offset * __extents.extent(r) + i,__rank_count(), idx...); + } + + template + _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) const { + return static_cast(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::value)) // needs two () due to comma + MDSPAN_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 + mapping(mapping 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::value)) // needs two () due to comma + MDSPAN_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 + mapping(layout_left::mapping 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) + ) + ) + MDSPAN_CONDITIONAL_EXPLICIT((extents_type::rank() > 0)) + MDSPAN_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 + mapping(layout_stride::mapping 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 !defined(_MDSPAN_HAS_CUDA) && !defined(_MDSPAN_HAS_HIP) && !defined(NDEBUG) + index_type stride = 1; + for(rank_type r=__extents.rank(); r>0; r--) { + if(stride != static_cast(other.stride(r-1))) { + // Note this throw will lead to a terminate if triggered since this function is marked noexcept + throw std::runtime_error("Assigning layout_stride to layout_right with invalid strides."); + } + stride *= __extents.extent(r-1); + } + #endif + } + + 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()) && + _MDSPAN_FOLD_AND( + (_MDSPAN_TRAIT(std::is_convertible, Indices, index_type) && + _MDSPAN_TRAIT(std::is_nothrow_constructible, index_type, Indices)) + ) + ) + ) + _MDSPAN_HOST_DEVICE + constexpr index_type operator()(Indices... idxs) const noexcept { + return __compute_offset(__rank_count<0, extents_type::rank()>(), static_cast(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 constexpr bool is_unique() const noexcept { return true; } + MDSPAN_INLINE_FUNCTION constexpr bool is_exhaustive() const noexcept { return true; } + MDSPAN_INLINE_FUNCTION constexpr bool is_strided() const 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; + } + + template + MDSPAN_INLINE_FUNCTION + friend constexpr bool operator==(mapping const& lhs, mapping const& rhs) noexcept { + return lhs.extents() == rhs.extents(); + } + + // In C++ 20 the not equal exists if equal is found +#if !(MDSPAN_HAS_CXX_20) + template + MDSPAN_INLINE_FUNCTION + friend constexpr bool operator!=(mapping const& lhs, mapping const& rhs) noexcept { + return lhs.extents() != rhs.extents(); + } +#endif + + // Not really public, but currently needed to implement fully constexpr useable submdspan: + template + constexpr index_type __get_stride(MDSPAN_IMPL_STANDARD_NAMESPACE::extents,std::integer_sequence) const { + return _MDSPAN_FOLD_TIMES_RIGHT((Idx>N? __extents.template __extent():1),1); + } + template + constexpr index_type __stride() const noexcept { + return __get_stride(__extents, std::make_index_sequence()); + } + +private: + _MDSPAN_NO_UNIQUE_ADDRESS extents_type __extents{}; + +}; + +} // end namespace MDSPAN_IMPL_STANDARD_NAMESPACE + diff --git a/lib/mdspan/include/experimental/__p0009_bits/layout_stride.hpp b/lib/mdspan/include/experimental/__p0009_bits/layout_stride.hpp new file mode 100644 index 00000000000..030a494529b --- /dev/null +++ b/lib/mdspan/include/experimental/__p0009_bits/layout_stride.hpp @@ -0,0 +1,495 @@ +//@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" + +#if !defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS) +# include "no_unique_address.hpp" +#endif + +#include +#include +#include +#ifdef __cpp_lib_span +#include +#endif +#if defined(_MDSPAN_USE_CONCEPTS) && MDSPAN_HAS_CXX_20 && defined(__cpp_lib_concepts) +# include +#endif + +namespace MDSPAN_IMPL_STANDARD_NAMESPACE { + +struct layout_left { + template + class mapping; +}; +struct layout_right { + template + class mapping; +}; + +namespace detail { + template + constexpr bool __is_mapping_of = + std::is_same, Mapping>::value; + +#if defined(_MDSPAN_USE_CONCEPTS) && MDSPAN_HAS_CXX_20 +# if !defined(__cpp_lib_concepts) + namespace internal { + namespace detail { + template + concept __same_as = std::is_same_v<_Tp, _Up>; + } // namespace detail + template + concept __same_as = detail::__same_as && detail::__same_as; + } // namespace internal +# endif + + template + concept __layout_mapping_alike = requires { + requires __is_extents::value; +#if defined(__cpp_lib_concepts) + { M::is_always_strided() } -> std::same_as; + { M::is_always_exhaustive() } -> std::same_as; + { M::is_always_unique() } -> std::same_as; +#else + { M::is_always_strided() } -> internal::__same_as; + { M::is_always_exhaustive() } -> internal::__same_as; + { M::is_always_unique() } -> internal::__same_as; +#endif + std::bool_constant::value; + std::bool_constant::value; + std::bool_constant::value; + }; +#endif +} // namespace detail + +struct layout_stride { + template + class mapping +#if !defined(_MDSPAN_USE_ATTRIBUTE_NO_UNIQUE_ADDRESS) + : private detail::__no_unique_address_emulation< + detail::__compressed_pair< + Extents, + std::array + > + > +#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, + 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 = std::array; + using __member_pair_t = detail::__compressed_pair; + +#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 + _MDSPAN_HOST_DEVICE + constexpr index_type __get_size(::MDSPAN_IMPL_STANDARD_NAMESPACE::extents,std::integer_sequence) const { + return _MDSPAN_FOLD_TIMES_RIGHT( static_cast(extents().extent(Idx)), 1 ); + } + + //---------------------------------------------------------------------------- + + template + friend class mapping; + + //---------------------------------------------------------------------------- + + // Workaround for non-deducibility of the index sequence template parameter if it's given at the top level + template + struct __deduction_workaround; + + template + struct __deduction_workaround> + { + template + MDSPAN_INLINE_FUNCTION + static constexpr bool _eq_impl(mapping const& self, mapping const& other) noexcept { + return _MDSPAN_FOLD_AND((self.stride(Idxs) == other.stride(Idxs)) /* && ... */) + && _MDSPAN_FOLD_AND((self.extents().extent(Idxs) == other.extents().extent(Idxs)) /* || ... */); + } + template + MDSPAN_INLINE_FUNCTION + static constexpr bool _not_eq_impl(mapping const& self, mapping const& other) noexcept { + return _MDSPAN_FOLD_OR((self.stride(Idxs) != other.stride(Idxs)) /* || ... */) + || _MDSPAN_FOLD_OR((self.extents().extent(Idxs) != other.extents().extent(Idxs)) /* || ... */); + } + + template + 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() - 1)...) + 1; + } + + template + MDSPAN_INLINE_FUNCTION + static constexpr const __strides_storage_t fill_strides(const OtherMapping& map) { + return __strides_storage_t{static_cast(map.stride(Idxs))...}; + } + + MDSPAN_INLINE_FUNCTION + static constexpr const __strides_storage_t& fill_strides(const __strides_storage_t& s) { + return s; + } + + template + MDSPAN_INLINE_FUNCTION + static constexpr const __strides_storage_t fill_strides(const std::array& s) { + return __strides_storage_t{static_cast(s[Idxs])...}; + } + +#ifdef __cpp_lib_span + template + MDSPAN_INLINE_FUNCTION + static constexpr const __strides_storage_t fill_strides(const std::span& s) { + return __strides_storage_t{static_cast(s[Idxs])...}; + } +#endif + + template + MDSPAN_INLINE_FUNCTION + static constexpr size_t __return_zero() { return 0; } + + template + MDSPAN_INLINE_FUNCTION + static constexpr typename Mapping::index_type + __OFFSET(const Mapping& m) { return m(__return_zero()...); } + }; + + // Can't use defaulted parameter in the __deduction_workaround template because of a bug in MSVC warning C4348. + using __impl = __deduction_workaround>; + + + //---------------------------------------------------------------------------- + +#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_DEFAULTED constexpr mapping() noexcept = default; + 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&, typename Extents::index_type) && + _MDSPAN_TRAIT(std::is_nothrow_constructible, typename Extents::index_type, const std::remove_const_t&) + ) + ) + MDSPAN_INLINE_FUNCTION + constexpr + mapping( + extents_type const& e, + std::array 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. + */ + } + +#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&, typename Extents::index_type) && + _MDSPAN_TRAIT(std::is_nothrow_constructible, typename Extents::index_type, const std::remove_const_t&) + ) + ) + MDSPAN_INLINE_FUNCTION + constexpr + mapping( + extents_type const& e, + std::span 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 && + StridedLayoutMapping::is_always_unique() && + StridedLayoutMapping::is_always_strided() + ) + ) +#else + template + requires( + detail::__layout_mapping_alike && + _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::value) && + (detail::__is_mapping_of || + detail::__is_mapping_of || + detail::__is_mapping_of) + ) // 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 __strides_storage(); + } + + MDSPAN_INLINE_FUNCTION + constexpr index_type required_span_size() const noexcept { + index_type span_size = 1; + for(unsigned r = 0; r < extents_type::rank(); r++) { + // Return early if any of the extents are zero + if(extents().extent(r)==0) return 0; + span_size += ( static_cast(extents().extent(r) - 1 ) * __strides_storage()[r]); + } + return span_size; + } + + + MDSPAN_TEMPLATE_REQUIRES( + class... Indices, + /* requires */ ( + sizeof...(Indices) == Extents::rank() && + _MDSPAN_FOLD_AND(_MDSPAN_TRAIT(std::is_convertible, Indices, index_type) /*&& ...*/ ) && + _MDSPAN_FOLD_AND(_MDSPAN_TRAIT(std::is_nothrow_constructible, index_type, Indices) /*&& ...*/) + ) + ) + MDSPAN_FORCE_INLINE_FUNCTION + constexpr index_type operator()(Indices... idxs) const noexcept { + return static_cast(__impl::_call_op_impl(*this, static_cast(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; } + MDSPAN_INLINE_FUNCTION _MDSPAN_CONSTEXPR_14 bool is_exhaustive() const noexcept { + return required_span_size() == __get_size(extents(), std::make_index_sequence()); + } + MDSPAN_INLINE_FUNCTION static constexpr bool is_strided() noexcept { return true; } + + + MDSPAN_INLINE_FUNCTION + constexpr index_type stride(rank_type r) const noexcept +#if MDSPAN_HAS_CXX_20 + requires ( Extents::rank() > 0 ) +#endif + { + return __strides_storage()[r]; + } + +#if !(defined(_MDSPAN_USE_CONCEPTS) && MDSPAN_HAS_CXX_20) + MDSPAN_TEMPLATE_REQUIRES( + class StridedLayoutMapping, + /* requires */ ( + detail::__is_mapping_of && + (extents_type::rank() == StridedLayoutMapping::extents_type::rank()) && + StridedLayoutMapping::is_always_strided() + ) + ) +#else + template + requires( + detail::__layout_mapping_alike && + (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 { + bool strides_match = true; + for(rank_type r = 0; r < extents_type::rank(); r++) + strides_match = strides_match && (x.stride(r) == y.stride(r)); + return (x.extents() == y.extents()) && + (__impl::__OFFSET(y)== static_cast(0)) && + strides_match; + } + + // 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 const& rhs) noexcept { + return __impl::_eq_impl(lhs, rhs); + } + +#if !MDSPAN_HAS_CXX_20 + MDSPAN_TEMPLATE_REQUIRES( + class StridedLayoutMapping, + /* requires */ ( + detail::__is_mapping_of && + (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 not (x == y); + } + + MDSPAN_TEMPLATE_REQUIRES( + class OtherExtents, + /* requires */ ( + (extents_type::rank() == OtherExtents::rank()) + ) + ) + MDSPAN_INLINE_FUNCTION + friend constexpr bool operator!=(mapping const& lhs, mapping const& rhs) noexcept { + return __impl::_not_eq_impl(lhs, rhs); + } +#endif + + }; +}; + +} // end namespace MDSPAN_IMPL_STANDARD_NAMESPACE diff --git a/lib/mdspan/include/experimental/__p0009_bits/macros.hpp b/lib/mdspan/include/experimental/__p0009_bits/macros.hpp new file mode 100644 index 00000000000..3eeb39755c8 --- /dev/null +++ b/lib/mdspan/include/experimental/__p0009_bits/macros.hpp @@ -0,0 +1,631 @@ +//@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 // std::is_void + +#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 + +//============================================================================== +// {{{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) + +// end Preprocessor helpers }}}1 +//============================================================================== + +//============================================================================== +// {{{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 end Concept emulation }}}1 +//============================================================================== + +//============================================================================== +// {{{1 + +#ifdef _MDSPAN_USE_INLINE_VARIABLES +# define _MDSPAN_INLINE_VARIABLE inline +#else +# define _MDSPAN_INLINE_VARIABLE +#endif + +// end inline variables }}}1 +//============================================================================== + +//============================================================================== +// {{{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> \ + { 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 + +// end Return type deduction }}}1 +//============================================================================== + +//============================================================================== +// {{{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 +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 +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 +MDSPAN_FORCE_INLINE_FUNCTION +constexpr auto __fold_left_assign_impl(Arg1&& arg1) { + return (Arg1&&)arg1; +} + +template +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 +MDSPAN_FORCE_INLINE_FUNCTION +constexpr auto __fold_right_assign_impl(Arg1&& arg1) { + return (Arg1&&)arg1; +} + +template +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 +MDSPAN_FORCE_INLINE_FUNCTION +constexpr auto __fold_right_plus_impl(Arg1&& arg1) { + return (Arg1&&)arg1; +} + +template +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 +MDSPAN_FORCE_INLINE_FUNCTION +constexpr auto __fold_right_times_impl(Arg1&& arg1) { + return (Arg1&&)arg1; +} + +template +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 + +//------------------------------------------------------------------------------ +// {{{2 + +template +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 +struct __fold_right_and_impl_ { + using __next_t = __fold_right_and_impl_; + using __rv = decltype(std::declval() && std::declval()); + MDSPAN_FORCE_INLINE_FUNCTION + static constexpr __rv + __impl(Arg&& arg, Args&&... args) noexcept { + return ((Arg&&)arg) && __next_t::__impl((Args&&)args...); + } +}; + +template +MDSPAN_FORCE_INLINE_FUNCTION +constexpr typename __fold_right_and_impl_::__rv +__fold_right_and_impl(Args&&... args) { + return __fold_right_and_impl_::__impl((Args&&)args...); +} + +// end right and }}}2 +//------------------------------------------------------------------------------ + +//------------------------------------------------------------------------------ +// {{{2 + +template +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 +struct __fold_right_or_impl_ { + using __next_t = __fold_right_or_impl_; + using __rv = decltype(std::declval() || std::declval()); + MDSPAN_FORCE_INLINE_FUNCTION + static constexpr __rv + __impl(Arg&& arg, Args&&... args) noexcept { + return ((Arg&&)arg) || __next_t::__impl((Args&&)args...); + } +}; + +template +MDSPAN_FORCE_INLINE_FUNCTION +constexpr typename __fold_right_or_impl_::__rv +__fold_right_or_impl(Args&&... args) { + return __fold_right_or_impl_::__impl((Args&&)args...); +} + +// end right or }}}2 +//------------------------------------------------------------------------------ + +//------------------------------------------------------------------------------ +// {{{2 + +template +struct __fold_right_plus_impl_; +template +struct __fold_right_plus_impl_ { + using __rv = Arg&&; + MDSPAN_FORCE_INLINE_FUNCTION + static constexpr __rv + __impl(Arg&& arg) noexcept { + return (Arg&&)arg; + } +}; +template +struct __fold_right_plus_impl_ { + using __next_t = __fold_right_plus_impl_; + using __rv = decltype(std::declval() + std::declval()); + 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 +MDSPAN_FORCE_INLINE_FUNCTION +constexpr typename __fold_right_plus_impl_::__rv +__fold_right_plus_impl(Args&&... args) { + return __fold_right_plus_impl_::__impl((Args&&)args...); +} + +// end right plus }}}2 +//------------------------------------------------------------------------------ + +//------------------------------------------------------------------------------ +// {{{2 + +template +struct __fold_right_times_impl_; +template +struct __fold_right_times_impl_ { + using __rv = Arg&&; + MDSPAN_FORCE_INLINE_FUNCTION + static constexpr __rv + __impl(Arg&& arg) noexcept { + return (Arg&&)arg; + } +}; +template +struct __fold_right_times_impl_ { + using __next_t = __fold_right_times_impl_; + using __rv = decltype(std::declval() * std::declval()); + 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 +MDSPAN_FORCE_INLINE_FUNCTION +constexpr typename __fold_right_times_impl_::__rv +__fold_right_times_impl(Args&&... args) { + return __fold_right_times_impl_::__impl((Args&&)args...); +} + +// end right times }}}2 +//------------------------------------------------------------------------------ + +//------------------------------------------------------------------------------ +// {{{2 + +template +struct __fold_right_assign_impl_; +template +struct __fold_right_assign_impl_ { + using __rv = Arg&&; + MDSPAN_FORCE_INLINE_FUNCTION + static constexpr __rv + __impl(Arg&& arg) noexcept { + return (Arg&&)arg; + } +}; +template +struct __fold_right_assign_impl_ { + using __next_t = __fold_right_assign_impl_; + using __rv = decltype(std::declval() = std::declval()); + 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 +MDSPAN_FORCE_INLINE_FUNCTION +constexpr typename __fold_right_assign_impl_::__rv +__fold_right_assign_impl(Args&&... args) { + return __fold_right_assign_impl_::__impl((Args&&)args...); +} + +// end right assign }}}2 +//------------------------------------------------------------------------------ + +//------------------------------------------------------------------------------ +// {{{2 + +template +struct __fold_left_assign_impl_; +template +struct __fold_left_assign_impl_ { + using __rv = Arg&&; + MDSPAN_FORCE_INLINE_FUNCTION + static constexpr __rv + __impl(Arg&& arg) noexcept { + return (Arg&&)arg; + } +}; +template +struct __fold_left_assign_impl_ { + using __assign_result_t = decltype(std::declval() = std::declval()); + 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 +MDSPAN_FORCE_INLINE_FUNCTION +constexpr typename __fold_left_assign_impl_::__rv +__fold_left_assign_impl(Args&&... args) { + return __fold_left_assign_impl_::__impl((Args&&)args...); +} + +// end left assign }}}2 +//------------------------------------------------------------------------------ + +#endif + + +template +constexpr __mdspan_enable_fold_comma __fold_comma_impl(Args&&... args) noexcept { return { }; } + +template +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) + +#endif + +// end fold expressions }}}1 +//============================================================================== + +//============================================================================== +// {{{1 + +#if _MDSPAN_USE_VARIABLE_TEMPLATES +# define _MDSPAN_TRAIT(TRAIT, ...) TRAIT##_v<__VA_ARGS__> +#else +# define _MDSPAN_TRAIT(TRAIT, ...) TRAIT<__VA_ARGS__>::value +#endif + +// end Variable template compatibility }}}1 +//============================================================================== + +//============================================================================== +// {{{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 + +// end Pre-C++14 constexpr }}}1 +//============================================================================== diff --git a/lib/mdspan/include/experimental/__p0009_bits/mdspan.hpp b/lib/mdspan/include/experimental/__p0009_bits/mdspan.hpp new file mode 100644 index 00000000000..6febe300215 --- /dev/null +++ b/lib/mdspan/include/experimental/__p0009_bits/mdspan.hpp @@ -0,0 +1,421 @@ +//@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 +> +class mdspan +{ +private: + static_assert(detail::__is_extents_v, + MDSPAN_IMPL_STANDARD_NAMESPACE_STRING "::mdspan's Extents template parameter must be a specialization of " MDSPAN_IMPL_STANDARD_NAMESPACE_STRING "::extents."); + + // Workaround for non-deducibility of the index sequence template parameter if it's given at the top level + template + struct __deduction_workaround; + + template + struct __deduction_workaround> + { + 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 + MDSPAN_FORCE_INLINE_FUNCTION static constexpr + ReferenceType __callop(mdspan const& __self, const std::array& indices) noexcept { + return __self.__accessor_ref().access(__self.__ptr_ref(), __self.__mapping_ref()(indices[Idxs]...)); + } + }; + +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; + using element_type = ElementType; + using value_type = std::remove_cv_t; + 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>; + + using __map_acc_pair_t = detail::__compressed_pair; + +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 */ ( + _MDSPAN_FOLD_AND(_MDSPAN_TRAIT(std::is_convertible, SizeTypes, index_type) /* && ... */) && + _MDSPAN_FOLD_AND(_MDSPAN_TRAIT(std::is_nothrow_constructible, index_type, SizeTypes) /* && ... */) && + ((sizeof...(SizeTypes) == rank()) || (sizeof...(SizeTypes) == rank_dynamic())) && + _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(std::move(dynamic_extents))...)), accessor_type())) + { } + + MDSPAN_TEMPLATE_REQUIRES( + class SizeType, size_t N, + /* requires */ ( + _MDSPAN_TRAIT(std::is_convertible, SizeType, index_type) && + _MDSPAN_TRAIT(std::is_nothrow_constructible, index_type, 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& 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, SizeType, index_type) && + _MDSPAN_TRAIT(std::is_nothrow_constructible, index_type, 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 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, 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, typename OtherLayoutPolicy::template mapping) && + _MDSPAN_TRAIT(std::is_constructible, accessor_type, OtherAccessor) + ) + ) + MDSPAN_INLINE_FUNCTION + constexpr mdspan(const mdspan& 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(std::move(indices))...)); + } + #endif + + MDSPAN_TEMPLATE_REQUIRES( + class SizeType, + /* requires */ ( + _MDSPAN_TRAIT(std::is_convertible, SizeType, index_type) && + _MDSPAN_TRAIT(std::is_nothrow_constructible, index_type, SizeType) + ) + ) + MDSPAN_FORCE_INLINE_FUNCTION + constexpr reference operator[](const std::array< SizeType, rank()>& indices) const + { + return __impl::template __callop(*this, indices); + } + + #ifdef __cpp_lib_span + MDSPAN_TEMPLATE_REQUIRES( + class SizeType, + /* requires */ ( + _MDSPAN_TRAIT(std::is_convertible, SizeType, index_type) && + _MDSPAN_TRAIT(std::is_nothrow_constructible, index_type, SizeType) + ) + ) + MDSPAN_FORCE_INLINE_FUNCTION + constexpr reference operator[](std::span indices) const + { + return __impl::template __callop(*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(std::move(idx)))); + } + #endif + + #if MDSPAN_USE_PAREN_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) /* && ... */) && + extents_type::rank() == sizeof...(SizeTypes) + ) + ) + MDSPAN_FORCE_INLINE_FUNCTION + constexpr reference operator()(SizeTypes... indices) const + { + return __accessor_ref().access(__ptr_ref(), __mapping_ref()(static_cast(std::move(indices))...)); + } + + MDSPAN_TEMPLATE_REQUIRES( + class SizeType, + /* requires */ ( + _MDSPAN_TRAIT(std::is_convertible, SizeType, index_type) && + _MDSPAN_TRAIT(std::is_nothrow_constructible, index_type, SizeType) + ) + ) + MDSPAN_FORCE_INLINE_FUNCTION + constexpr reference operator()(const std::array& indices) const + { + return __impl::template __callop(*this, indices); + } + + #ifdef __cpp_lib_span + MDSPAN_TEMPLATE_REQUIRES( + class SizeType, + /* requires */ ( + _MDSPAN_TRAIT(std::is_convertible, SizeType, index_type) && + _MDSPAN_TRAIT(std::is_nothrow_constructible, index_type, SizeType) + ) + ) + MDSPAN_FORCE_INLINE_FUNCTION + constexpr reference operator()(std::span indices) const + { + return __impl::template __callop(*this, indices); + } + #endif // __cpp_lib_span + #endif // MDSPAN_USE_PAREN_OPERATOR + + MDSPAN_INLINE_FUNCTION constexpr size_t size() const noexcept { + return __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() 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 bool is_unique() const noexcept { return __mapping_ref().is_unique(); }; + MDSPAN_INLINE_FUNCTION constexpr bool is_exhaustive() const noexcept { return __mapping_ref().is_exhaustive(); }; + MDSPAN_INLINE_FUNCTION constexpr bool is_strided() const noexcept { 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 __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 + 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_integral, SizeTypes) /* && ... */) && + (sizeof...(SizeTypes) > 0) +) +MDSPAN_DEDUCTION_GUIDE explicit mdspan(ElementType*, SizeTypes...) + -> mdspan>; + +MDSPAN_TEMPLATE_REQUIRES( + class Pointer, + (_MDSPAN_TRAIT(std::is_pointer, std::remove_reference_t)) +) +MDSPAN_DEDUCTION_GUIDE mdspan(Pointer&&) -> mdspan>, extents>; + +MDSPAN_TEMPLATE_REQUIRES( + class CArray, + (_MDSPAN_TRAIT(std::is_array, CArray) && (std::rank_v == 1)) +) +MDSPAN_DEDUCTION_GUIDE mdspan(CArray&) -> mdspan, extents>>; + +template +MDSPAN_DEDUCTION_GUIDE mdspan(ElementType*, const ::std::array&) + -> mdspan>; + +#ifdef __cpp_lib_span +template +MDSPAN_DEDUCTION_GUIDE mdspan(ElementType*, ::std::span) + -> mdspan>; +#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 +MDSPAN_DEDUCTION_GUIDE mdspan(ElementType*, const extents&) + -> mdspan>; + +template +MDSPAN_DEDUCTION_GUIDE mdspan(ElementType*, const MappingType&) + -> mdspan; + +template +MDSPAN_DEDUCTION_GUIDE mdspan(const typename AccessorType::data_handle_type, const MappingType&, const AccessorType&) + -> mdspan; +#endif + +} // end namespace MDSPAN_IMPL_STANDARD_NAMESPACE diff --git a/lib/mdspan/include/experimental/__p0009_bits/no_unique_address.hpp b/lib/mdspan/include/experimental/__p0009_bits/no_unique_address.hpp new file mode 100644 index 00000000000..36e64ee24db --- /dev/null +++ b/lib/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 +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 +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/lib/mdspan/include/experimental/__p0009_bits/trait_backports.hpp b/lib/mdspan/include/experimental/__p0009_bits/trait_backports.hpp new file mode 100644 index 00000000000..4933dd9934e --- /dev/null +++ b/lib/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 +#include // integer_sequence + +//============================================================================== +// {{{1 + +#ifdef _MDSPAN_NEEDS_TRAIT_VARIABLE_TEMPLATE_BACKPORTS + +#if _MDSPAN_USE_VARIABLE_TEMPLATES +namespace MDSPAN_IMPL_STANDARD_NAMESPACE { + +#define _MDSPAN_BACKPORT_TRAIT(TRAIT) \ + template _MDSPAN_INLINE_VARIABLE constexpr auto TRAIT##_v = TRAIT::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 + +// end Variable template trait backports (e.g., is_void_v) }}}1 +//============================================================================== + +//============================================================================== +// {{{1 + +#if !defined(_MDSPAN_USE_INTEGER_SEQUENCE) || !_MDSPAN_USE_INTEGER_SEQUENCE + +namespace MDSPAN_IMPL_STANDARD_NAMESPACE { + +template +struct integer_sequence { + static constexpr size_t size() noexcept { return sizeof...(Vals); } + using value_type = T; +}; + +template +using index_sequence = std::integer_sequence; + +namespace __detail { + +template +struct __make_int_seq_impl; + +template +struct __make_int_seq_impl> +{ + using type = integer_sequence; +}; + +template +struct __make_int_seq_impl< + T, N, I, integer_sequence +> : __make_int_seq_impl> +{ }; + +} // end namespace __detail + +template +using make_integer_sequence = typename __detail::__make_int_seq_impl>::type; + +template +using make_index_sequence = typename __detail::__make_int_seq_impl>::type; + +template +using index_sequence_for = make_index_sequence; + +} // end namespace MDSPAN_IMPL_STANDARD_NAMESPACE + +#endif + +// end integer sequence (ugh...) }}}1 +//============================================================================== + +//============================================================================== +// {{{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 using TRAIT##_t = typename TRAIT::type; + +_MDSPAN_BACKPORT_TRAIT_ALIAS(remove_cv) +_MDSPAN_BACKPORT_TRAIT_ALIAS(remove_reference) + +template +using enable_if_t = typename enable_if<_B, _T>::type; + +#undef _MDSPAN_BACKPORT_TRAIT_ALIAS + +} // end namespace MDSPAN_IMPL_STANDARD_NAMESPACE + +#endif + +// end standard trait aliases }}}1 +//============================================================================== + +#endif //MDSPAN_INCLUDE_EXPERIMENTAL_BITS_TRAIT_BACKPORTS_HPP_ diff --git a/lib/mdspan/include/experimental/__p0009_bits/type_list.hpp b/lib/mdspan/include/experimental/__p0009_bits/type_list.hpp new file mode 100644 index 00000000000..deca7c15d09 --- /dev/null +++ b/lib/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 struct __type_list { static constexpr auto __size = sizeof...(_Ts); }; + +// Implementation of type_list at() that's heavily optimized for small typelists +template struct __type_at; +template > struct __type_at_large_impl; + +template +struct __type_at_entry { }; + +template +struct __type_at_assign_op_ignore_rest { + template + __type_at_assign_op_ignore_rest<_Result> operator=(_T&&); + using type = _Result; +}; + +struct __type_at_assign_op_impl { + template + __type_at_assign_op_impl operator=(__type_at_entry<_I, _Idx, _T>&&); + template + __type_at_assign_op_ignore_rest<_T> operator=(__type_at_entry<_I, _I, _T>&&); +}; + +template +struct __type_at_large_impl<_I, __type_list<_Ts...>, std::integer_sequence> + : decltype( + _MDSPAN_FOLD_ASSIGN_LEFT(__type_at_assign_op_impl{}, /* = ... = */ __type_at_entry<_I, _Idxs, _Ts>{}) + ) +{ }; + +template +struct __type_at<_I, __type_list<_Ts...>> + : __type_at_large_impl<_I, __type_list<_Ts...>> +{ }; + +template +struct __type_at<0, __type_list<_T0, _Ts...>> { + using type = _T0; +}; + +template +struct __type_at<1, __type_list<_T0, _T1, _Ts...>> { + using type = _T1; +}; + +template +struct __type_at<2, __type_list<_T0, _T1, _T2, _Ts...>> { + using type = _T2; +}; + +template +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/lib/mdspan/include/experimental/__p1684_bits/mdarray.hpp b/lib/mdspan/include/experimental/__p1684_bits/mdarray.hpp new file mode 100644 index 00000000000..3950273a83d --- /dev/null +++ b/lib/mdspan/include/experimental/__p1684_bits/mdarray.hpp @@ -0,0 +1,490 @@ +//@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 +#include + +namespace MDSPAN_IMPL_STANDARD_NAMESPACE { +namespace MDSPAN_IMPL_PROPOSED_NAMESPACE { + +namespace { + template + struct size_of_extents; + + template + struct size_of_extents> { + constexpr static size_t value() { + size_t size = 1; + for(size_t r=0; r::rank(); r++) + size *= extents::static_extent(r); + return size; + } + }; +} + +namespace { + template + struct container_is_array : std::false_type { + template + static constexpr C construct(const M& m) { return C(m.required_span_size()); } + }; + template + struct container_is_array> : std::true_type { + template + static constexpr std::array construct(const M&) { return std::array(); } + }; +} + +template < + class ElementType, + class Extents, + class LayoutPolicy = layout_right, + class Container = std::vector +> +class mdarray { +private: + static_assert(::MDSPAN_IMPL_STANDARD_NAMESPACE::detail::__is_extents_v, + 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; + using element_type = ElementType; + using mdspan_type = mdspan; + using const_mdspan_type = mdspan; + using value_type = std::remove_cv_t; + 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_FOLD_AND(_MDSPAN_TRAIT( std::is_convertible, SizeTypes, index_type) /* && ... */) && + _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::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::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::value) && + _MDSPAN_TRAIT( std::is_constructible, mapping_type, extents_type)) + ) : map_(exts), ctr_(container_is_array::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::value) + ) : map_(m), ctr_(container_is_array::construct(map_)) + { } + + // Constructors from container + MDSPAN_TEMPLATE_REQUIRES( + class... SizeTypes, + /* requires */ ( + _MDSPAN_FOLD_AND(_MDSPAN_TRAIT( std::is_convertible, SizeTypes, index_type) /* && ... */) && + _MDSPAN_TRAIT( std::is_constructible, extents_type, SizeTypes...) && + _MDSPAN_TRAIT( std::is_constructible, mapping_type, extents_type) + ) + ) + MDSPAN_INLINE_FUNCTION + explicit constexpr mdarray(const container_type& ctr, SizeTypes... dynamic_extents) + : map_(extents_type(dynamic_extents...)), ctr_(ctr) + { assert(ctr.size() >= static_cast(map_.required_span_size())); } + + + MDSPAN_FUNCTION_REQUIRES( + (MDSPAN_INLINE_FUNCTION constexpr), + mdarray, (const container_type& ctr, const extents_type& exts), , + /* requires */ (_MDSPAN_TRAIT( std::is_constructible, mapping_type, extents_type)) + ) : map_(exts), ctr_(ctr) + { assert(ctr.size() >= static_cast(map_.required_span_size())); } + + constexpr mdarray(const container_type& ctr, const mapping_type& m) + : map_(m), ctr_(ctr) + { assert(ctr.size() >= static_cast(map_.required_span_size())); } + + + // Constructors from container + MDSPAN_TEMPLATE_REQUIRES( + class... SizeTypes, + /* requires */ ( + _MDSPAN_FOLD_AND(_MDSPAN_TRAIT( std::is_convertible, SizeTypes, index_type) /* && ... */) && + _MDSPAN_TRAIT( std::is_constructible, extents_type, SizeTypes...) && + _MDSPAN_TRAIT( std::is_constructible, mapping_type, extents_type) + ) + ) + MDSPAN_INLINE_FUNCTION + explicit constexpr mdarray(container_type&& ctr, SizeTypes... dynamic_extents) + : map_(extents_type(dynamic_extents...)), ctr_(std::move(ctr)) + { assert(ctr_.size() >= static_cast(map_.required_span_size())); } + + + MDSPAN_FUNCTION_REQUIRES( + (MDSPAN_INLINE_FUNCTION constexpr), + mdarray, (container_type&& ctr, const extents_type& exts), , + /* requires */ (_MDSPAN_TRAIT( std::is_constructible, mapping_type, extents_type)) + ) : map_(exts), ctr_(std::move(ctr)) + { assert(ctr_.size() >= static_cast(map_.required_span_size())); } + + constexpr mdarray(container_type&& ctr, const mapping_type& m) + : map_(m), ctr_(std::move(ctr)) + { assert(ctr_.size() >= static_cast(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) && + _MDSPAN_TRAIT( std::is_constructible, container_type, OtherContainer) + ) + ) + MDSPAN_INLINE_FUNCTION + constexpr mdarray(const mdarray& other) + : map_(other.mapping()), ctr_(other.container()) + { + static_assert( std::is_constructible::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 container_type& ctr, const extents_type& exts, const Alloc& a) + : map_(exts), ctr_(ctr, a) + { assert(ctr_.size() >= static_cast(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 container_type& ctr, const mapping_type& map, const Alloc& a) + : map_(map), ctr_(ctr, a) + { assert(ctr_.size() >= static_cast(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(container_type&& ctr, const extents_type& exts, const Alloc& a) + : map_(exts), ctr_(std::move(ctr), a) + { assert(ctr_.size() >= static_cast(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(container_type&& ctr, const mapping_type& map, 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) && + _MDSPAN_TRAIT( std::is_constructible, container_type, OtherContainer, Alloc) + ) + ) + MDSPAN_INLINE_FUNCTION + constexpr mdarray(const mdarray& other, const Alloc& a) + : map_(other.mapping()), ctr_(other.container(), a) + { + static_assert( std::is_constructible::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(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(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& indices) const noexcept + { + return __impl::template __callop(*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& indices) noexcept + { + return __impl::template __callop(*this, indices); + } +#endif + + + #if MDSPAN_USE_PAREN_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(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(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& indices) const noexcept + { + return __impl::template __callop(*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& indices) noexcept + { + return __impl::template __callop(*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_type, + mdspan) + ) + ) + constexpr operator mdspan () { + return mdspan_type(data(), map_); + } + + MDSPAN_TEMPLATE_REQUIRES( + class OtherElementType, class OtherExtents, + class OtherLayoutType, class OtherAccessorType, + /* requires */ ( + _MDSPAN_TRAIT(std::is_assignable, const_mdspan_type, + mdspan) + ) + ) + constexpr operator mdspan () const { + return const_mdspan_type(data(), map_); + } + + MDSPAN_TEMPLATE_REQUIRES( + class OtherAccessorType = default_accessor, + /* requires */ ( + _MDSPAN_TRAIT(std::is_assignable, mdspan_type, + mdspan) + ) + ) + constexpr mdspan + to_mdspan(const OtherAccessorType& a = default_accessor()) { + return mdspan(data(), map_, a); + } + + MDSPAN_TEMPLATE_REQUIRES( + class OtherAccessorType = default_accessor, + /* requires */ ( + _MDSPAN_TRAIT(std::is_assignable, const_mdspan_type, + mdspan) + ) + ) + constexpr mdspan + to_mdspan(const OtherAccessorType& a = default_accessor()) const { + return mdspan(data(), map_, a); + } + +private: + mapping_type map_; + container_type ctr_; + + template + friend class mdarray; +}; + + +} // end namespace MDSPAN_IMPL_PROPOSED_NAMESPACE +} // end namespace MDSPAN_IMPL_STANDARD_NAMESPACE diff --git a/lib/mdspan/include/experimental/__p2630_bits/strided_slice.hpp b/lib/mdspan/include/experimental/__p2630_bits/strided_slice.hpp new file mode 100644 index 00000000000..58f38620ba1 --- /dev/null +++ b/lib/mdspan/include/experimental/__p2630_bits/strided_slice.hpp @@ -0,0 +1,49 @@ + +//@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 + +namespace MDSPAN_IMPL_STANDARD_NAMESPACE { +namespace MDSPAN_IMPL_PROPOSED_NAMESPACE { + +namespace { + template + struct __mdspan_is_integral_constant: std::false_type {}; + + template + struct __mdspan_is_integral_constant>: std::true_type {}; +} +// Slice Specifier allowing for strides and compile time extent +template +struct strided_slice { + using offset_type = OffsetType; + using extent_type = ExtentType; + using stride_type = StrideType; + + OffsetType offset; + ExtentType extent; + StrideType stride; + + static_assert(std::is_integral_v || __mdspan_is_integral_constant::value); + static_assert(std::is_integral_v || __mdspan_is_integral_constant::value); + static_assert(std::is_integral_v || __mdspan_is_integral_constant::value); +}; + +} // MDSPAN_IMPL_PROPOSED_NAMESPACE +} // MDSPAN_IMPL_STANDARD_NAMESPACE diff --git a/lib/mdspan/include/experimental/__p2630_bits/submdspan.hpp b/lib/mdspan/include/experimental/__p2630_bits/submdspan.hpp new file mode 100644 index 00000000000..b9672b7f9ac --- /dev/null +++ b/lib/mdspan/include/experimental/__p2630_bits/submdspan.hpp @@ -0,0 +1,42 @@ +//@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 { +namespace MDSPAN_IMPL_PROPOSED_NAMESPACE { +template +MDSPAN_INLINE_FUNCTION +constexpr auto +submdspan(const mdspan &src, + SliceSpecifiers... slices) { + const auto sub_mapping_offset = 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; + 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( + src.accessor().offset(src.data_handle(), sub_mapping_offset.offset), + sub_mapping_offset.mapping, + sub_accessor_t(src.accessor())); +} +} // namespace MDSPAN_IMPL_PROPOSED_NAMESPACE +} // namespace MDSPAN_IMPL_STANDARD_NAMESPACE diff --git a/lib/mdspan/include/experimental/__p2630_bits/submdspan_extents.hpp b/lib/mdspan/include/experimental/__p2630_bits/submdspan_extents.hpp new file mode 100644 index 00000000000..f56ce023f16 --- /dev/null +++ b/lib/mdspan/include/experimental/__p2630_bits/submdspan_extents.hpp @@ -0,0 +1,323 @@ +//@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 + +#include "strided_slice.hpp" +namespace MDSPAN_IMPL_STANDARD_NAMESPACE { +namespace MDSPAN_IMPL_PROPOSED_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 +MDSPAN_INLINE_FUNCTION +constexpr auto inv_map_rank(std::integral_constant, std::index_sequence) { + return std::index_sequence(); +} + +// specialization reducing rank by one (i.e., integral slice specifier) +template +MDSPAN_INLINE_FUNCTION +constexpr auto inv_map_rank(std::integral_constant, std::index_sequence, Slice, + SliceSpecifiers... slices) { + using next_idx_seq_t = std::conditional_t, + std::index_sequence, + std::index_sequence>; + + return inv_map_rank(std::integral_constant(), next_idx_seq_t(), + slices...); +} + +// Helper for identifying strided_slice +template struct is_strided_slice : std::false_type {}; + +template +struct is_strided_slice< + strided_slice> : std::true_type {}; + +// first_of(slice): getting begin of slice specifier range +MDSPAN_TEMPLATE_REQUIRES( + class Integral, + /* requires */(std::is_convertible_v) +) +MDSPAN_INLINE_FUNCTION +constexpr Integral first_of(const Integral &i) { + return i; +} + +MDSPAN_INLINE_FUNCTION +constexpr std::integral_constant +first_of(const ::MDSPAN_IMPL_STANDARD_NAMESPACE::full_extent_t &) { + return std::integral_constant(); +} + +MDSPAN_TEMPLATE_REQUIRES( + class Slice, + /* requires */(std::is_convertible_v>) +) +MDSPAN_INLINE_FUNCTION +constexpr auto first_of(const Slice &i) { + return std::get<0>(i); +} + +template +MDSPAN_INLINE_FUNCTION +constexpr OffsetType +first_of(const strided_slice &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) +) +MDSPAN_INLINE_FUNCTION +constexpr Integral + last_of(std::integral_constant, const Extents &, const Integral &i) { + return i; +} + +MDSPAN_TEMPLATE_REQUIRES( + size_t k, class Extents, class Slice, + /* requires */(std::is_convertible_v>) +) +MDSPAN_INLINE_FUNCTION +constexpr auto last_of(std::integral_constant, const Extents &, + const Slice &i) { + return std::get<1>(i); +} + +// 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 +MDSPAN_INLINE_FUNCTION +constexpr auto last_of(std::integral_constant, const Extents &ext, + ::MDSPAN_IMPL_STANDARD_NAMESPACE::full_extent_t) { + if constexpr (Extents::static_extent(k) == dynamic_extent) { + return ext.extent(k); + } else { + return std::integral_constant(); + } +#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 +MDSPAN_INLINE_FUNCTION +constexpr OffsetType +last_of(std::integral_constant, const Extents &, + const strided_slice &r) { + return r.extent; +} + +// get stride of slices +template +MDSPAN_INLINE_FUNCTION +constexpr auto stride_of(const T &) { + return std::integral_constant(); +} + +template +MDSPAN_INLINE_FUNCTION +constexpr auto +stride_of(const strided_slice &r) { + return r.stride; +} + +// divide which can deal with integral constant preservation +template +MDSPAN_INLINE_FUNCTION +constexpr auto divide(const T0 &v0, const T1 &v1) { + return IndexT(v0) / IndexT(v1); +} + +template +MDSPAN_INLINE_FUNCTION +constexpr auto divide(const std::integral_constant &, + const std::integral_constant &) { + // cutting short division by zero + // this is used for strided_slice with zero extent/stride + return std::integral_constant(); +} + +// multiply which can deal with integral constant preservation +template +MDSPAN_INLINE_FUNCTION +constexpr auto multiply(const T0 &v0, const T1 &v1) { + return IndexT(v0) * IndexT(v1); +} + +template +MDSPAN_INLINE_FUNCTION +constexpr auto multiply(const std::integral_constant &, + const std::integral_constant &) { + return std::integral_constant(); +} + +// compute new static extent from range, preserving static knowledge +template struct StaticExtentFromRange { + constexpr static size_t value = dynamic_extent; +}; + +template +struct StaticExtentFromRange, + std::integral_constant> { + constexpr static size_t value = val1 - val0; +}; + +// compute new static extent from strided_slice, preserving static +// knowledge +template struct StaticExtentFromStridedRange { + constexpr static size_t value = dynamic_extent; +}; + +template +struct StaticExtentFromStridedRange, + std::integral_constant> { + 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 +struct extents_constructor { + MDSPAN_TEMPLATE_REQUIRES( + class Slice, class... SlicesAndExtents, + /* requires */(!std::is_convertible_v && + !is_strided_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())), + decltype(last_of(std::integral_constant(), + std::declval(), + std::declval()))>::value; + + using next_t = + extents_constructor; + using index_t = typename Extents::index_type; + return next_t::next_extent( + ext, slices_and_extents..., + index_t(last_of(std::integral_constant(), ext, + sl)) - + index_t(first_of(sl))); + } + + MDSPAN_TEMPLATE_REQUIRES( + class Slice, class... SlicesAndExtents, + /* requires */ (std::is_convertible_v) + ) + MDSPAN_INLINE_FUNCTION + constexpr static auto next_extent(const Extents &ext, const Slice &, + SlicesAndExtents... slices_and_extents) { + using next_t = extents_constructor; + return next_t::next_extent(ext, slices_and_extents...); + } + + template + MDSPAN_INLINE_FUNCTION + constexpr static auto + next_extent(const Extents &ext, + const strided_slice &r, + SlicesAndExtents... slices_and_extents) { + using index_t = typename Extents::index_type; + using new_static_extent_t = + StaticExtentFromStridedRange; + if constexpr (new_static_extent_t::value == dynamic_extent) { + using next_t = + extents_constructor; + return next_t::next_extent( + ext, slices_and_extents..., + r.extent > 0 ? 1 + divide(r.extent - 1, r.stride) : 0); + } else { + constexpr size_t new_static_extent = new_static_extent_t::value; + using next_t = + extents_constructor; + return next_t::next_extent( + ext, slices_and_extents..., index_t(divide(ExtentType(), StrideType()))); + } + } +}; + +template +struct extents_constructor<0, Extents, NewStaticExtents...> { + + template + MDSPAN_INLINE_FUNCTION + constexpr static auto next_extent(const Extents &, NewExtents... new_exts) { + return extents( + new_exts...); + } +}; + +} // namespace detail + +// submdspan_extents creates new extents given src extents and submdspan slice +// specifiers +template +MDSPAN_INLINE_FUNCTION +constexpr auto submdspan_extents(const extents &src_exts, + SliceSpecifiers... slices) { + + using ext_t = extents; + return detail::extents_constructor::next_extent( + src_exts, slices...); +} +} // namespace MDSPAN_IMPL_PROPOSED_NAMESPACE +} // namespace MDSPAN_IMPL_STANDARD_NAMESPACE diff --git a/lib/mdspan/include/experimental/__p2630_bits/submdspan_mapping.hpp b/lib/mdspan/include/experimental/__p2630_bits/submdspan_mapping.hpp new file mode 100644 index 00000000000..48778d57e75 --- /dev/null +++ b/lib/mdspan/include/experimental/__p2630_bits/submdspan_mapping.hpp @@ -0,0 +1,299 @@ +//@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 +#include +#include +#include // index_sequence + +namespace MDSPAN_IMPL_STANDARD_NAMESPACE { +namespace MDSPAN_IMPL_PROPOSED_NAMESPACE { +//****************************************** +// Return type of submdspan_mapping overloads +//****************************************** +template struct mapping_offset { + Mapping mapping; + size_t offset; +}; +} // namespace MDSPAN_IMPL_PROPOSED_NAMESPACE + +namespace detail { +using MDSPAN_IMPL_PROPOSED_NAMESPACE::detail::first_of; +using MDSPAN_IMPL_PROPOSED_NAMESPACE::detail::stride_of; +using MDSPAN_IMPL_PROPOSED_NAMESPACE::detail::inv_map_rank; + +// constructs sub strides +template +MDSPAN_INLINE_FUNCTION +constexpr auto +construct_sub_strides(const SrcMapping &src_mapping, + std::index_sequence, + const std::tuple &slices_stride_factor) { + using index_type = typename SrcMapping::index_type; + return std::array{ + (static_cast(src_mapping.stride(InvMapIdxs)) * + static_cast(std::get(slices_stride_factor)))...}; +} +} // namespace detail + +//********************************** +// layout_left submdspan_mapping +//********************************* +namespace detail { + +// Figure out whether to preserve layout_left +template +struct preserve_layout_left_mapping; + +template +struct preserve_layout_left_mapping, SubRank, + SliceSpecifiers...> { + constexpr static bool value = + // Preserve layout for rank 0 + (SubRank == 0) || + ( + // Slice specifiers up to subrank need to be full_extent_t - except + // for the last one which could also be tuple but not a strided index + // range slice specifiers after subrank are integrals + ((Idx > SubRank - 1) || // these are only integral slice specifiers + (std::is_same_v) || + ((Idx == SubRank - 1) && + std::is_convertible_v>)) && + ...); +}; +} // namespace detail + +// 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 +// Actual submdspan mapping call +template +MDSPAN_INLINE_FUNCTION +constexpr auto +submdspan_mapping(const layout_left::mapping &src_mapping, + SliceSpecifiers... slices) { + using MDSPAN_IMPL_PROPOSED_NAMESPACE::submdspan_extents; + using MDSPAN_IMPL_PROPOSED_NAMESPACE::mapping_offset; + + // compute sub extents + using src_ext_t = Extents; + auto dst_ext = submdspan_extents(src_mapping.extents(), slices...); + using dst_ext_t = decltype(dst_ext); + + // figure out sub layout type + constexpr bool preserve_layout = detail::preserve_layout_left_mapping< + decltype(std::make_index_sequence()), dst_ext_t::rank(), + SliceSpecifiers...>::value; + using dst_layout_t = + std::conditional_t; + using dst_mapping_t = typename dst_layout_t::template mapping; + + if constexpr (std::is_same_v) { + // layout_left case + return mapping_offset{ + dst_mapping_t(dst_ext), + static_cast(src_mapping(detail::first_of(slices)...))}; + } else { + // layout_stride case + auto inv_map = detail::inv_map_rank( + std::integral_constant(), + std::index_sequence<>(), + slices...); + return mapping_offset{ + dst_mapping_t(dst_ext, detail::construct_sub_strides( + src_mapping, 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 + #if defined(_MDSPAN_HAS_HIP) || (defined(__NVCC__) && (__CUDACC_VER_MAJOR__ * 100 + __CUDACC_VER_MINOR__ * 10) < 1120) + std::tuple{detail::stride_of(slices)...})), + #else + std::tuple{detail::stride_of(slices)...})), + #endif + static_cast(src_mapping(detail::first_of(slices)...))}; + } +#if defined(__NVCC__) && !defined(__CUDA_ARCH__) && defined(__GNUC__) + __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 + +//********************************** +// layout_right submdspan_mapping +//********************************* +namespace detail { + +// Figure out whether to preserve layout_right +template +struct preserve_layout_right_mapping; + +template +struct preserve_layout_right_mapping, SubRank, + SliceSpecifiers...> { + constexpr static size_t SrcRank = sizeof...(SliceSpecifiers); + constexpr static bool value = + // Preserve layout for rank 0 + (SubRank == 0) || + ( + // The last subrank slice specifiers need to be full_extent_t - except + // for the srcrank-subrank one which could also be tuple but not a + // strided index range slice specifiers before srcrank-subrank are + // integrals + ((Idx < + SrcRank - SubRank) || // these are only integral slice specifiers + (std::is_same_v) || + ((Idx == SrcRank - SubRank) && + std::is_convertible_v>)) && + ...); +}; +} // namespace detail + +// 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 +MDSPAN_INLINE_FUNCTION +constexpr auto +submdspan_mapping(const layout_right::mapping &src_mapping, + SliceSpecifiers... slices) { + using MDSPAN_IMPL_PROPOSED_NAMESPACE::submdspan_extents; + using MDSPAN_IMPL_PROPOSED_NAMESPACE::mapping_offset; + + // get sub extents + using src_ext_t = Extents; + auto dst_ext = submdspan_extents(src_mapping.extents(), slices...); + using dst_ext_t = decltype(dst_ext); + + // determine new layout type + constexpr bool preserve_layout = detail::preserve_layout_right_mapping< + decltype(std::make_index_sequence()), dst_ext_t::rank(), + SliceSpecifiers...>::value; + using dst_layout_t = + std::conditional_t; + using dst_mapping_t = typename dst_layout_t::template mapping; + + if constexpr (std::is_same_v) { + // layout_right case + return mapping_offset{ + dst_mapping_t(dst_ext), + static_cast(src_mapping(detail::first_of(slices)...))}; + } else { + // layout_stride case + auto inv_map = detail::inv_map_rank( + std::integral_constant(), + std::index_sequence<>(), + slices...); + return mapping_offset{ + dst_mapping_t(dst_ext, detail::construct_sub_strides( + src_mapping, 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 + #if defined(_MDSPAN_HAS_HIP) || (defined(__NVCC__) && (__CUDACC_VER_MAJOR__ * 100 + __CUDACC_VER_MINOR__ * 10) < 1120) + std::tuple{detail::stride_of(slices)...})), + #else + std::tuple{detail::stride_of(slices)...})), + #endif + static_cast(src_mapping(detail::first_of(slices)...))}; + } +#if defined(__NVCC__) && !defined(__CUDA_ARCH__) && defined(__GNUC__) + __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 + +//********************************** +// layout_stride submdspan_mapping +//********************************* +template +MDSPAN_INLINE_FUNCTION +constexpr auto +submdspan_mapping(const layout_stride::mapping &src_mapping, + SliceSpecifiers... slices) { + using MDSPAN_IMPL_PROPOSED_NAMESPACE::submdspan_extents; + using MDSPAN_IMPL_PROPOSED_NAMESPACE::mapping_offset; + auto dst_ext = submdspan_extents(src_mapping.extents(), slices...); + using dst_ext_t = decltype(dst_ext); + auto inv_map = detail::inv_map_rank( + std::integral_constant(), + std::index_sequence<>(), + slices...); + using dst_mapping_t = typename layout_stride::template mapping; + return mapping_offset{ + dst_mapping_t(dst_ext, detail::construct_sub_strides( + src_mapping, 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 + #if defined(_MDSPAN_HAS_HIP) || (defined(__NVCC__) && (__CUDACC_VER_MAJOR__ * 100 + __CUDACC_VER_MINOR__ * 10) < 1120) + std::tuple(detail::stride_of(slices)...))), +#else + std::tuple(detail::stride_of(slices)...))), +#endif + static_cast(src_mapping(detail::first_of(slices)...))}; +} +} // namespace MDSPAN_IMPL_STANDARD_NAMESPACE diff --git a/lib/mdspan/include/experimental/mdarray b/lib/mdspan/include/experimental/mdarray new file mode 100644 index 00000000000..642d1f5ad9e --- /dev/null +++ b/lib/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/lib/mdspan/include/experimental/mdspan b/lib/mdspan/include/experimental/mdspan new file mode 100644 index 00000000000..e8ba715ec2f --- /dev/null +++ b/lib/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/lib/mdspan/include/mdspan/mdarray.hpp b/lib/mdspan/include/mdspan/mdarray.hpp new file mode 100644 index 00000000000..fd8f61c52f1 --- /dev/null +++ b/lib/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/lib/mdspan/include/mdspan/mdspan.hpp b/lib/mdspan/include/mdspan/mdspan.hpp new file mode 100644 index 00000000000..b440873526a --- /dev/null +++ b/lib/mdspan/include/mdspan/mdspan.hpp @@ -0,0 +1,41 @@ +//@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/__p2630_bits/submdspan.hpp" +#endif + +#endif // MDSPAN_HPP_