From dd6a41a8b76ac4dfa4fd222d538232550be45b86 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Fri, 9 Aug 2024 15:55:15 +0200 Subject: [PATCH] Implement `submdspan_mapping` --- .../cuda/std/__mdspan/submdspan_helper.h | 15 +- .../cuda/std/__mdspan/submdspan_mapping.h | 328 ++++++++++++++ libcudacxx/include/cuda/std/mdspan | 1 + .../views/mdspan/submdspan/helper.h | 55 +++ .../mdspan/submdspan/layout_left.pass.cpp | 413 +++++++++++++++++ .../mdspan/submdspan/layout_right.pass.cpp | 414 ++++++++++++++++++ 6 files changed, 1213 insertions(+), 13 deletions(-) create mode 100644 libcudacxx/include/cuda/std/__mdspan/submdspan_mapping.h create mode 100644 libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/helper.h create mode 100644 libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/layout_left.pass.cpp create mode 100644 libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/layout_right.pass.cpp diff --git a/libcudacxx/include/cuda/std/__mdspan/submdspan_helper.h b/libcudacxx/include/cuda/std/__mdspan/submdspan_helper.h index e317474a040..9ecbc679457 100644 --- a/libcudacxx/include/cuda/std/__mdspan/submdspan_helper.h +++ b/libcudacxx/include/cuda/std/__mdspan/submdspan_helper.h @@ -85,28 +85,17 @@ struct full_extent_t }; _CCCL_GLOBAL_CONSTANT full_extent_t full_extent{}; -// [mdspan.submdspan.submdspan.mapping.result] -template -struct submdspan_mapping_result -{ - static_assert(true, // __is_layout_mapping<_LayoutMapping>, - "[mdspan.submdspan.submdspan.mapping.result] shall meet the layout mapping requirements"); - - _CCCL_NO_UNIQUE_ADDRESS _LayoutMapping mapping{}; - size_t offset{}; -}; - // [mdspan.submdspan.helpers] _CCCL_TEMPLATE(class _Tp) _CCCL_REQUIRES((!__integral_constant_like<_Tp>) ) -_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __de_ice(_Tp __val) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __de_ice(_Tp __val) noexcept { return __val; } _CCCL_TEMPLATE(class _Tp) _CCCL_REQUIRES(__integral_constant_like<_Tp>) -_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __de_ice(_Tp) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr auto __de_ice(_Tp) noexcept { return _Tp::value; } diff --git a/libcudacxx/include/cuda/std/__mdspan/submdspan_mapping.h b/libcudacxx/include/cuda/std/__mdspan/submdspan_mapping.h new file mode 100644 index 00000000000..da40b8e433a --- /dev/null +++ b/libcudacxx/include/cuda/std/__mdspan/submdspan_mapping.h @@ -0,0 +1,328 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___MDSPAN_SUBMDSPAN_MAPPING_H +#define _LIBCUDACXX___MDSPAN_SUBMDSPAN_MAPPING_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#if _CCCL_STD_VER >= 2014 + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +// [mdspan.sub.map] + +// [mdspan.submdspan.submdspan.mapping.result] +template +struct submdspan_mapping_result +{ + static_assert(true, // __is_layout_mapping<_LayoutMapping>, + "[mdspan.submdspan.submdspan.mapping.result] shall meet the layout mapping requirements"); + + _CCCL_NO_UNIQUE_ADDRESS _LayoutMapping mapping{}; + size_t offset{}; +}; + +// [mdspan.sub.map.common] +template +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr auto +__get_submdspan_strides(const _LayoutMapping& __mapping, _Slices... __slices) noexcept +{ + using _SliceType = __get_slice_type<_SliceIndex, _Slices...>; + using _Extents = typename _LayoutMapping::extents_type; + using _IndexType = typename _Extents::index_type; + if constexpr (__is_strided_slice>) + { + _SliceType& __slice = _CUDA_VSTD::__get_slice_at<_SliceIndex>(__slices...); + using __unsigned_stride = make_unsigned_t; + using __unsigned_extent = make_unsigned_t; + return static_cast<_IndexType>( + __mapping.stride(_SliceIndex) + * (static_cast<__unsigned_stride>(__slice.stride) < static_cast<__unsigned_extent>(__slice.extent) + ? _CUDA_VSTD::__de_ice(__slice.stride) + : 1)); + } + else + { + return static_cast<_IndexType>(__mapping.stride(_SliceIndex)); + } +} + +template +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr auto +__submdspan_strides(index_sequence<_SliceIndices...>, const _LayoutMapping& __mapping, _Slices... __slices) noexcept +{ + using _Extents = typename _LayoutMapping::extents_type; + using _IndexType = typename _Extents::index_type; + using _SubExtents = __get_subextents_t<_Extents, _Slices...>; + return array<_IndexType, _SubExtents::rank()>{ + _CUDA_VSTD::__get_submdspan_strides<_SliceIndices>(__mapping, __slices...)...}; +} + +template +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr auto +__submdspan_strides(const _LayoutMapping& __mapping, _Slices... __slices) +{ + using _Extents = typename _LayoutMapping::extents_type; + using _IndexType = typename _Extents::index_type; + const auto __filtered_indices = __filter_slices_convertible_to_index<_IndexType, _Slices...>( + index_sequence<>{}, _CUDA_VSTD::index_sequence_for<_Slices...>()); + return _CUDA_VSTD::__submdspan_strides(__filtered_indices, __mapping, __slices...); +} + +// [mdspan.sub.map.common-8] +template +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr size_t +__submdspan_offset(index_sequence<_SliceIndices...>, const _LayoutMapping& __mapping, _Slices... __slices) +{ + using _Extents = typename _LayoutMapping::extents_type; + using _IndexType = typename _Extents::index_type; + // If first_(slices...) + const array<_IndexType, _Extents::rank()> __offsets = { + _CUDA_VSTD::__first_extent_from_slice<_IndexType, _SliceIndices>(__slices...)...}; + + using _SubExtents = __get_subextents_t<_Extents, _Slices...>; + for (size_t __index = 0; __index != _SubExtents::rank(); ++__index) + { + // If first_(slices...) equals extents().extent(k) for any rank index k of extents() + if (__offsets[__index] == __mapping.extents().extent(__index)) + { + // then let offset be a value of type size_t equal to (*this).required_span_size() + return static_cast(__mapping.required_span_size()); + } + } + // Otherwise, let offset be a value of type size_t equal to (*this)(first_(slices...)...). + return static_cast(__mapping(__offsets[_SliceIndices]...)); +} + +template +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr size_t +__submdspan_offset(const _LayoutMapping& __mapping, _Slices... __slices) +{ + return _CUDA_VSTD::__submdspan_offset(_CUDA_VSTD::index_sequence_for<_Slices...>(), __mapping, __slices...); +} + +// [mdspan.sub.map.common-9] +// [mdspan.sub.map.common-9.1] +template +_CCCL_CONCEPT __is_strided_slice_stride_of_one = _CCCL_REQUIRES_EXPR((_SliceType))( + requires(__is_strided_slice>), + requires(__integral_constant_like), + requires(_SliceType::stride_type::value == 1)); + +template +_LIBCUDACXX_HIDE_FROM_ABI constexpr bool __is_unit_stride_slice() +{ + // [mdspan.sub.map.common-9.1] + if constexpr (__is_strided_slice_stride_of_one<_SliceType>) + { + return true; + } + // [mdspan.sub.map.common-9.2] + else if constexpr (__index_pair_like<_SliceType, typename _LayoutMapping::index_type>) + { + return true; + } + // [mdspan.sub.map.common-9.3] + else if constexpr (_CCCL_TRAIT(is_convertible, _SliceType, full_extent_t)) + { + return true; + } + else + { + return false; + } + _CCCL_UNREACHABLE(); +} + +// [mdspan.sub.map.left] +template +_LIBCUDACXX_HIDE_FROM_ABI constexpr bool __can_layout_left() +{ + // [mdspan.sub.map.left-1.2] + if constexpr (_SubExtents::rank() == 0) + { + return true; + } + // [mdspan.sub.map.left-1.3.2] + else if constexpr (sizeof...(_OtherSlices) == 0) + { + return _CUDA_VSTD::__is_unit_stride_slice<_LayoutMapping, _Slice>(); + } + // [mdspan.sub.map.left-1.3.1] + else if constexpr (_CCCL_TRAIT(is_convertible, _Slice, full_extent_t)) + { + return _CUDA_VSTD::__can_layout_left<_LayoutMapping, _SubExtents, _OtherSlices...>(); + } + else + { + return false; + } + _CCCL_UNREACHABLE(); +} + +template +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr auto +__submdspan_mapping_impl(const typename layout_left::mapping<_Extents>& __mapping, _Slices... __slices) +{ + // [mdspan.sub.map.left-1.1] + if constexpr (_Extents::rank() == 0) + { + return submdspan_mapping_result{__mapping, 0}; + } + + // [mdspan.sub.map.left-1.2] + // [mdspan.sub.map.left-1.3] + using _SubExtents = __get_subextents_t<_Extents, _Slices...>; + const auto __sub_ext = _CUDA_VSTD::submdspan_extents(__mapping.extents(), __slices...); + const auto __offset = _CUDA_VSTD::__submdspan_offset(__mapping, __slices...); + if constexpr (_CUDA_VSTD::__can_layout_left, _SubExtents, _Slices...>()) + { + using __sub_mapping_t = layout_left::template mapping<_SubExtents>; + return submdspan_mapping_result<__sub_mapping_t>{__sub_mapping_t{__sub_ext}, __offset}; + } + // [mdspan.sub.map.left-1.4] + // TODO: Implement padded layouts + else + { + // [mdspan.sub.map.left-1.5] + using __sub_mapping_t = layout_stride::template mapping<_SubExtents>; + const auto __sub_strides = _CUDA_VSTD::__submdspan_strides(__mapping, __slices...); + return submdspan_mapping_result<__sub_mapping_t>{__sub_mapping_t{__sub_ext, __sub_strides}, __offset}; + } + _CCCL_UNREACHABLE(); +} + +template +_LIBCUDACXX_HIDE_FROM_ABI constexpr bool __can_layout_right() +{ + // [mdspan.sub.map.right-1.2] + if constexpr (_SubExtents::rank() == 0) + { + return true; + } + // [mdspan.sub.map.right-1.3.2] + else if constexpr (sizeof...(_OtherSlices) == 0) + { + return _CUDA_VSTD::__is_unit_stride_slice<_LayoutMapping, _Slice>(); + } + // [mdspan.sub.map.right-1.3.1] + else if constexpr (_CCCL_TRAIT(is_convertible, _Slice, full_extent_t)) + { + return _CUDA_VSTD::__can_layout_left<_LayoutMapping, _SubExtents, _OtherSlices...>(); + } + else + { + return false; + } + _CCCL_UNREACHABLE(); +} + +template +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr auto +__submdspan_mapping_impl(const typename layout_right::mapping<_Extents>& __mapping, _Slices... __slices) +{ + // [mdspan.sub.map.right-1.1] + if constexpr (_Extents::rank() == 0) + { + return submdspan_mapping_result{__mapping, 0}; + } + else + { + // [mdspan.sub.map.right-1.2] + // [mdspan.sub.map.right-1.3] + using _SubExtents = __get_subextents_t<_Extents, _Slices...>; + const auto __sub_ext = _CUDA_VSTD::submdspan_extents(__mapping.extents(), __slices...); + const auto __offset = _CUDA_VSTD::__submdspan_offset(__mapping, __slices...); + if constexpr (_CUDA_VSTD::__can_layout_right, _SubExtents, _Slices...>()) + { + using __sub_mapping_t = layout_right::template mapping<_SubExtents>; + return submdspan_mapping_result<__sub_mapping_t>{__sub_mapping_t{__sub_ext}, __offset}; + } + // [mdspan.sub.map.right-1.4] + // TODO: Implement padded layouts + else + { + // [mdspan.sub.map.right-1.5] + using __sub_mapping_t = layout_stride::template mapping<_SubExtents>; + const auto __sub_strides = _CUDA_VSTD::__submdspan_strides(__mapping, __slices...); + return submdspan_mapping_result<__sub_mapping_t>{__sub_mapping_t{__sub_ext, __sub_strides}, __offset}; + } + } + _CCCL_UNREACHABLE(); +} + +template +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr auto +__submdspan_mapping_impl(const typename layout_stride::mapping<_Extents>& __mapping, _Slices... __slices) +{ + // [mdspan.sub.map.stride-1.1] + if constexpr (_Extents::rank() == 0) + { + return submdspan_mapping_result{__mapping, 0}; + } + else + { + // [mdspan.sub.map.stride-1.2] + using _SubExtents = __get_subextents_t<_Extents, _Slices...>; + using __sub_mapping_t = layout_stride::template mapping<_SubExtents>; + const auto __sub_ext = _CUDA_VSTD::submdspan_extents(__mapping.extents(), __slices...); + const auto __offset = _CUDA_VSTD::__submdspan_offset(__mapping, __slices...); + const auto __sub_strides = _CUDA_VSTD::__submdspan_strides(__mapping, __slices...); + return submdspan_mapping_result<__sub_mapping_t>{__sub_mapping_t{__sub_ext, __sub_strides}, __offset}; + } +} + +template +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr auto +submdspan_mapping(const _LayoutMapping& __mapping, _Slices... __slices) +{ + return _CUDA_VSTD::__submdspan_mapping_impl(__mapping, __slices...); +} + +_CCCL_TEMPLATE(class _Tp, class _Extents, class _Layout, class _Accessor, class... _Slices) +_CCCL_REQUIRES(true) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr auto +submdspan(const mdspan<_Tp, _Extents, _Layout, _Accessor>& __src, _Slices... __slices) +{ + auto __sub_map_result = _CUDA_VSTD::submdspan_mapping(__src.mapping(), __slices...); + return mdspan(__src.accessor().offset(__src.data_handle(), __sub_map_result.offset), + __sub_map_result.mapping, + typename _Accessor::offset_policy(__src.accessor())); +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _CCCL_STD_VER >= 2014 + +#endif // _LIBCUDACXX___MDSPAN_SUBMDSPAN_MAPPING_H diff --git a/libcudacxx/include/cuda/std/mdspan b/libcudacxx/include/cuda/std/mdspan index 75e4b8f3705..ac75b2ac700 100644 --- a/libcudacxx/include/cuda/std/mdspan +++ b/libcudacxx/include/cuda/std/mdspan @@ -33,6 +33,7 @@ _CCCL_PUSH_MACROS #include #include #include +#include #include _CCCL_POP_MACROS diff --git a/libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/helper.h b/libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/helper.h new file mode 100644 index 00000000000..7d09d8eed44 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/helper.h @@ -0,0 +1,55 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef TEST_STD_CONTAINERS_VIEWS_MDSPAN_SUBMDSPAN_HELPER_H +#define TEST_STD_CONTAINERS_VIEWS_MDSPAN_SUBMDSPAN_HELPER_H + +#include +#include + +_CCCL_TEMPLATE(class MDSpan) +_CCCL_REQUIRES((MDSpan::rank() == 0)) +__host__ __device__ constexpr bool equal_to(const MDSpan& mdspan, const char* expected) +{ + return mdspan[cuda::std::array{}] == expected[0]; +} + +_CCCL_TEMPLATE(class MDSpan) +_CCCL_REQUIRES((MDSpan::rank() == 1)) +__host__ __device__ constexpr bool equal_to(const MDSpan& mdspan, const char* expected) +{ + for (size_t i = 0; i != mdspan.size(); ++i) + { + if (mdspan[i] != expected[i]) + { + return false; + } + } + return true; +} + +_CCCL_TEMPLATE(class MDSpan) +_CCCL_REQUIRES((MDSpan::rank() == 2)) +__host__ __device__ constexpr bool equal_to(const MDSpan& mdspan, cuda::std::array expected) +{ + for (size_t i = 0; i != mdspan.extent(0); ++i) + { + for (size_t j = 0; j != mdspan.extent(1); ++j) + { + const cuda::std::array indices{i, j}; + if (mdspan[indices] != expected[i][j]) + { + return false; + } + } + } + return true; +} + +#endif // TEST_STD_CONTAINERS_VIEWS_MDSPAN_SUBMDSPAN_HELPER_H diff --git a/libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/layout_left.pass.cpp b/libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/layout_left.pass.cpp new file mode 100644 index 00000000000..e2cecd5bcbd --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/layout_left.pass.cpp @@ -0,0 +1,413 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++11, c++14 + +// + +// constexpr mdspan& operator=(const mdspan& rhs) = default; + +#include +#include +#include + +#include "helper.h" +#include "test_macros.h" + +__host__ __device__ constexpr bool test() +{ + constexpr char data[] = {'H', 'O', 'P', 'P', 'E', 'R'}; + + { // 1d mdspan + // ['H', 'O', 'P', 'P', 'E', 'R'] + cuda::std::mdspan md{data, cuda::std::layout_left::mapping{cuda::std::dims<1>{6}}}; + static_assert(md.rank() == 1); + static_assert(md.rank_dynamic() == 1); + assert(equal_to(md, "HOPPER")); + + using mdspan_t = decltype(md); + static_assert(cuda::std::is_same_v); + + { // full_extent + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x x x x x x ] + cuda::std::mdspan sub = cuda::std::submdspan(md, cuda::std::full_extent); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 1); + assert(sub.extent(0) == 6); + assert(sub.size() == 6); + assert(equal_to(sub, "HOPPER")); + } + + { // Slice of elements from start 0:4 + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x x x x ] + const auto slice = cuda::std::pair{0, 4}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 1); + assert(sub.extent(0) == 4); + assert(sub.size() == 4); + assert(equal_to(sub, "HOPP")); + } + + { // Slice of elements in the middle 2:5 + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x x x ] + const auto slice = cuda::std::pair{2, 5}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 1); + assert(sub.extent(0) == 3); + assert(sub.size() == 3); + assert(equal_to(sub, "PPE")); + } + + { // Slice of elements in the end 3:6 + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x x x ] + const auto slice = cuda::std::pair{3, 6}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 1); + assert(sub.extent(0) == 3); + assert(sub.size() == 3); + assert(equal_to(sub, "PER")); + } + + { // Slice of elements with strided slice without offset, full size and stride 1 + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x x x x x x ] offset + // [ x x x x x x ] size + // [ x x x x x x ] stride + // [ x x x x x x ] + const cuda::std::strided_slice slice{0, md.extent(0), 1}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 1); + assert(sub.extent(0) == 6); + assert(sub.size() == 6); + assert(equal_to(sub, "HOPPER")); + } + + { // Slice of elements with strided slice with offset, full remaining size and stride 1 + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x x x x ] offset + // [ x x x x ] size + // [ x x x x ] stride + // [ x x x x ] + const cuda::std::strided_slice slice{2, md.extent(0) - 2, 1}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 1); + assert(sub.extent(0) == 4); + assert(sub.size() == 4); + assert(equal_to(sub, "PPER")); + } + + { // Slice of elements with strided slice with offset, smaller size and stride 1 + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x x x x ] offset + // [ x x ] size + // [ x x ] stride + // [ x x ] + const cuda::std::strided_slice slice{2, md.extent(0) - 4, 1}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 1); + assert(sub.extent(0) == 2); + assert(sub.size() == 2); + assert(equal_to(sub, "PP")); + } + + { // Slice of elements with strided slice without offset, full size and stride 3 + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x x x x x x ] offset + // [ x x x x x x ] size + // [ x x ] stride + // [ x x ] + const cuda::std::strided_slice slice{0, md.extent(0), 3}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 3); + assert(sub.extent(0) == 2); + assert(sub.size() == 2); + assert(equal_to(sub, "HP")); + } + + { // Slice of elements with strided slice with offset, full size and stride 3 + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x x x x x ] offset + // [ x x x x x ] size + // [ x x ] stride + // [ x x ] + const cuda::std::strided_slice slice{1, md.extent(0) - 1, 3}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 3); + assert(sub.extent(0) == 2); + assert(sub.size() == 2); + assert(equal_to(sub, "OE")); + } + + { // Slice of elements with strided slice with offset, size less equal than stride + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x x x x x ] offset + // [ x x x ] size + // [ x ] stride + // [ x ] + const cuda::std::strided_slice slice{1, 3, 3}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 1); + assert(sub.extent(0) == 1); + assert(sub.size() == 1); + assert(equal_to(sub, "O")); + } + + { // Single element, with integral constant + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x ] + const auto slice = cuda::std::integral_constant{}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice); + + static_assert(sub.rank() == 0); + static_assert(sub.rank_dynamic() == 0); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.size() == 1); + assert(equal_to(sub, "P")); + } + } + + { // 2d mdspan + // ['H', 'P', 'E'] + // ['O', 'P', 'R'] + cuda::std::mdspan md{data, cuda::std::layout_left::mapping{cuda::std::dims<2>{2, 3}}}; + static_assert(md.rank() == 2); + static_assert(md.rank_dynamic() == 2); + + assert(md.stride(0) == 1); + assert(md.stride(1) == md.extent(0)); + assert(md.extent(0) == 2); + assert(md.extent(1) == 3); + assert(md.size() == 6); + assert(equal_to(md, {"HPE", "OPR"})); + + { // full_extent + // ['H', 'P', 'E'] [ x ] [ x x x ] + // ['O', 'P', 'R'] [ x ] [ x x x ] + cuda::std::mdspan sub = cuda::std::submdspan(md, cuda::std::full_extent, cuda::std::full_extent); + + static_assert(sub.rank() == 2); + static_assert(sub.rank_dynamic() == 2); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 1); + assert(sub.stride(1) == md.extent(0)); + assert(sub.extent(0) == md.extent(0)); + assert(sub.extent(1) == md.extent(1)); + assert(sub.size() == md.size()); + assert(equal_to(sub, {"HPE", "OPR"})); + } + + { // full extent, then slice of elements from start 0:2 + // ['H', 'P', 'E'] [ x ] [ x x ] + // ['O', 'P', 'R'] [ x ] [ x x ] + const auto slice2 = cuda::std::pair{0, 2}; + cuda::std::mdspan sub = cuda::std::submdspan(md, cuda::std::full_extent, slice2); + + static_assert(sub.rank() == 2); + static_assert(sub.rank_dynamic() == 2); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == md.stride(0)); + assert(sub.stride(1) == md.stride(1)); + assert(sub.extent(0) == md.extent(0)); + assert(sub.extent(1) == 2); + assert(sub.size() == 4); + assert(equal_to(sub, {"HP", "OP"})); + } + + { // Slice of elements from start 0:1, then full extent + // ['H', 'P', 'E'] [ x ] [ x x x ] + // ['O', 'P', 'R'] [ ] [ ] + const auto slice1 = cuda::std::pair{0, 1}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice1, cuda::std::full_extent); + + static_assert(sub.rank() == 2); + static_assert(sub.rank_dynamic() == 2); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == md.stride(0)); + assert(sub.stride(1) == md.stride(1)); + assert(sub.extent(0) == 1); + assert(sub.extent(1) == md.extent(1)); + assert(sub.size() == 3); + assert(equal_to(sub, {"HPE", ""})); + } + + { // Slice of elements from middle 1:2, then strided_slice without offset, full size and stride 1 + // ['H', 'P', 'E'] [ ] [ ] + // ['O', 'P', 'R'] [ x ] [ x x x ] + const auto slice1 = cuda::std::pair{1, 2}; + const cuda::std::strided_slice slice2{0, md.extent(1), 1}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice1, slice2); + + static_assert(sub.rank() == 2); + static_assert(sub.rank_dynamic() == 2); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == md.stride(0)); + assert(sub.stride(1) == md.stride(1)); + assert(sub.extent(0) == 1); + assert(sub.extent(1) == md.extent(1)); + assert(sub.size() == 3); + assert(equal_to(sub, {"OPR", ""})); + } + + { // Slice of elements from middle 1:2, then strided_slice with offset, full size and stride 1 + // ['H', 'P', 'E'] [ ] [ ] + // ['O', 'P', 'R'] [ x ] [ x x ] + const auto slice1 = cuda::std::pair{1, 2}; + const cuda::std::strided_slice slice2{1, md.extent(1) - 1, 1}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice1, slice2); + + static_assert(sub.rank() == 2); + static_assert(sub.rank_dynamic() == 2); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == md.stride(0)); + assert(sub.stride(1) == md.stride(1)); + assert(sub.extent(0) == 1); + assert(sub.extent(1) == md.extent(1) - 1); + assert(sub.size() == 2); + assert(equal_to(sub, {"PR", ""})); + } + + { // Slice of elements from middle 1:2, then strided_slice without offset, full size and stride 2 + // ['H', 'P', 'E'] [ ] [ ] + // ['O', 'P', 'R'] [ x ] [ x x ] + const auto slice1 = cuda::std::pair{1, 2}; + const cuda::std::strided_slice slice2{0, md.extent(1), 2}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice1, slice2); + + static_assert(sub.rank() == 2); + static_assert(sub.rank_dynamic() == 2); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == md.stride(0)); + assert(sub.stride(1) == 2 * md.stride(1)); + assert(sub.extent(0) == 1); + assert(sub.extent(1) == md.extent(1) - 1); + assert(sub.size() == 2); + assert(equal_to(sub, {"OR", ""})); + } + + { // Slice of elements from middle 1:2, then index + // ['H', 'P', 'E'] [ ] [ ] + // ['O', 'P', 'R'] [ x ] [ x ] + const auto slice1 = cuda::std::pair{1, 2}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice1, 2); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == md.stride(0)); + assert(sub.extent(0) == 1); + assert(sub.size() == 1); + assert(equal_to(sub, "R")); + } + } + return true; +} + +int main(int, char**) +{ + test(); + static_assert(test(), ""); + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/layout_right.pass.cpp b/libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/layout_right.pass.cpp new file mode 100644 index 00000000000..e6aba8cad55 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/containers/views/mdspan/submdspan/layout_right.pass.cpp @@ -0,0 +1,414 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++11, c++14 + +// + +// constexpr mdspan& operator=(const mdspan& rhs) = default; + +#include +#include +#include + +#include "helper.h" +#include "test_macros.h" + +__host__ __device__ constexpr bool test() +{ + constexpr char data[] = {'H', 'O', 'P', 'P', 'E', 'R'}; + + { // 1d mdspan + // ['H', 'O', 'P', 'P', 'E', 'R'] + cuda::std::mdspan md{data, cuda::std::layout_right::mapping{cuda::std::dims<1>{6}}}; + static_assert(md.rank() == 1); + static_assert(md.rank_dynamic() == 1); + assert(equal_to(md, "HOPPER")); + + using mdspan_t = decltype(md); + static_assert(cuda::std::is_same_v); + + { // full_extent + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x x x x x x ] + cuda::std::mdspan sub = cuda::std::submdspan(md, cuda::std::full_extent); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 1); + assert(sub.extent(0) == 6); + assert(sub.size() == 6); + assert(equal_to(sub, "HOPPER")); + } + + { // Slice of elements from start 0:4 + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x x x x ] + const auto slice = cuda::std::pair{0, 4}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 1); + assert(sub.extent(0) == 4); + assert(sub.size() == 4); + assert(equal_to(sub, "HOPP")); + } + + { // Slice of elements in the middle 2:5 + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x x x ] + const auto slice = cuda::std::pair{2, 5}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 1); + assert(sub.extent(0) == 3); + assert(sub.size() == 3); + assert(equal_to(sub, "PPE")); + } + + { // Slice of elements in the end 3:6 + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x x x ] + const auto slice = cuda::std::pair{3, 6}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 1); + assert(sub.extent(0) == 3); + assert(sub.size() == 3); + assert(equal_to(sub, "PER")); + } + + { // Slice of elements with strided slice without offset, full size and stride 1 + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x x x x x x ] offset + // [ x x x x x x ] size + // [ x x x x x x ] stride + // [ x x x x x x ] + const cuda::std::strided_slice slice{0, md.extent(0), 1}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 1); + assert(sub.extent(0) == 6); + assert(sub.size() == 6); + assert(equal_to(sub, "HOPPER")); + } + + { // Slice of elements with strided slice with offset, full remaining size and stride 1 + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x x x x ] offset + // [ x x x x ] size + // [ x x x x ] stride + // [ x x x x ] + const cuda::std::strided_slice slice{2, md.extent(0) - 2, 1}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 1); + assert(sub.extent(0) == 4); + assert(sub.size() == 4); + assert(equal_to(sub, "PPER")); + } + + { // Slice of elements with strided slice with offset, smaller size and stride 1 + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x x x x ] offset + // [ x x ] size + // [ x x ] stride + // [ x x ] + const cuda::std::strided_slice slice{2, md.extent(0) - 4, 1}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 1); + assert(sub.extent(0) == 2); + assert(sub.size() == 2); + assert(equal_to(sub, "PP")); + } + + { // Slice of elements with strided slice without offset, full size and stride 3 + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x x x x x x ] offset + // [ x x x x x x ] size + // [ x x ] stride + // [ x x ] + const cuda::std::strided_slice slice{0, md.extent(0), 3}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 3); + assert(sub.extent(0) == 2); + assert(sub.size() == 2); + assert(equal_to(sub, "HP")); + } + + { // Slice of elements with strided slice with offset, full size and stride 3 + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x x x x x ] offset + // [ x x x x x ] size + // [ x x ] stride + // [ x x ] + const cuda::std::strided_slice slice{1, md.extent(0) - 1, 3}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 3); + assert(sub.extent(0) == 2); + assert(sub.size() == 2); + assert(equal_to(sub, "OE")); + } + + { // Slice of elements with strided slice with offset, size less equal than stride + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x x x x x ] offset + // [ x x x ] size + // [ x ] stride + // [ x ] + const cuda::std::strided_slice slice{1, 3, 3}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 1); + assert(sub.extent(0) == 1); + assert(sub.size() == 1); + assert(equal_to(sub, "O")); + } + + { // Single element, with integral constant + // ['H', 'O', 'P', 'P', 'E', 'R'] + // [ x ] + const auto slice = cuda::std::integral_constant{}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice); + + static_assert(sub.rank() == 0); + static_assert(sub.rank_dynamic() == 0); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.size() == 1); + assert(equal_to(sub, "P")); + } + } + + { // 2d mdspan + // ['H', 'O', 'P'] + // ['P', 'E', 'R'] + cuda::std::mdspan md{data, cuda::std::layout_right::mapping{cuda::std::dims<2>{2, 3}}}; + static_assert(md.rank() == 2); + static_assert(md.rank_dynamic() == 2); + + assert(md.stride(0) == md.extent(1)); + assert(md.stride(1) == 1); + assert(md.extent(0) == 2); + assert(md.extent(1) == 3); + assert(md.size() == 6); + assert(equal_to(md, {"HOP", "PER"})); + + { // full_extent + // ['H', 'O', 'P'] [ x ] [ x x x ] + // ['P', 'E', 'R'] [ x ] [ x x x ] + cuda::std::mdspan sub = cuda::std::submdspan(md, cuda::std::full_extent, cuda::std::full_extent); + + static_assert(sub.rank() == 2); + static_assert(sub.rank_dynamic() == 2); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == md.extent(1)); + assert(sub.stride(1) == 1); + assert(sub.extent(0) == md.extent(0)); + assert(sub.extent(1) == md.extent(1)); + assert(sub.size() == md.size()); + assert(equal_to(sub, {"HOP", "PER"})); + } + + { // full extent, then slice of elements from start 0:1 + // ['H', 'O', 'P'] [ x ] [ x ] + // ['P', 'E', 'R'] [ x ] [ x ] + const auto slice2 = cuda::std::pair{0, 1}; + cuda::std::mdspan sub = cuda::std::submdspan(md, cuda::std::full_extent, slice2); + + static_assert(sub.rank() == 2); + static_assert(sub.rank_dynamic() == 2); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == 1); + assert(sub.stride(1) == md.stride(1)); + assert(sub.extent(0) == md.extent(0)); + assert(sub.extent(1) == 1); + assert(sub.size() == 2); + assert(equal_to(sub, {"H", "O"})); + } + + { // Slice of elements from start 1:2, then full extent + // ['H', 'O', 'P'] [ ] [ ] + // ['P', 'E', 'R'] [ x ] [ x x x ] + const auto slice1 = cuda::std::pair{1, 2}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice1, cuda::std::full_extent); + + static_assert(sub.rank() == 2); + static_assert(sub.rank_dynamic() == 2); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == md.stride(0)); + assert(sub.stride(1) == md.stride(1)); + assert(sub.extent(0) == 1); + assert(sub.extent(1) == md.extent(1)); + assert(sub.size() == 3); + assert(equal_to(sub, {"PER", ""})); + } + + { // Slice of elements from middle 1:2, then strided_slice without offset, full size and stride 1 + // ['H', 'O', 'P'] [ ] [ ] + // ['P', 'E', 'R'] [ x ] [ x x x ] + const auto slice1 = cuda::std::pair{1, 2}; + const cuda::std::strided_slice slice2{0, md.extent(1), 1}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice1, slice2); + + static_assert(sub.rank() == 2); + static_assert(sub.rank_dynamic() == 2); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == md.stride(0)); + assert(sub.stride(1) == md.stride(1)); + assert(sub.extent(0) == 1); + assert(sub.extent(1) == md.extent(1)); + assert(sub.size() == 3); + assert(equal_to(sub, {"PER", ""})); + } + + { // Slice of elements from middle 1:2, then strided_slice with offset, full size and stride 1 + // ['H', 'O', 'P'] [ ] [ ] + // ['P', 'E', 'R'] [ x ] [ x x ] + const auto slice1 = cuda::std::pair{1, 2}; + const cuda::std::strided_slice slice2{1, md.extent(1) - 1, 1}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice1, slice2); + + static_assert(sub.rank() == 2); + static_assert(sub.rank_dynamic() == 2); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == md.stride(0)); + assert(sub.stride(1) == md.stride(1)); + assert(sub.extent(0) == 1); + assert(sub.extent(1) == md.extent(1) - 1); + assert(sub.size() == 2); + assert(equal_to(sub, {"ER", ""})); + } + + { // Slice of elements from middle 1:2, then strided_slice without offset, full size and stride 2 + // ['H', 'O', 'P'] [ ] [ ] + // ['P', 'E', 'R'] [ x ] [ x ] + const auto slice1 = cuda::std::pair{1, 2}; + const cuda::std::strided_slice slice2{0, 2, 2}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice1, slice2); + + static_assert(sub.rank() == 2); + static_assert(sub.rank_dynamic() == 2); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == md.stride(0)); + assert(sub.stride(1) == 1); + assert(sub.extent(0) == 1); + assert(sub.extent(1) == 1); + assert(sub.size() == 1); + assert(equal_to(sub, {"P", ""})); + } + + { // Slice of elements from middle 1:2, then index + // ['H', 'O', 'P'] [ ] [ ] + // ['P', 'E', 'R'] [ x ] [ x ] + const auto slice1 = cuda::std::pair{1, 2}; + cuda::std::mdspan sub = cuda::std::submdspan(md, slice1, 2); + + static_assert(sub.rank() == 1); + static_assert(sub.rank_dynamic() == 1); + + using submdspan_t = decltype(sub); + static_assert(cuda::std::is_same_v); + + assert(sub.stride(0) == md.stride(0)); + assert(sub.extent(0) == 1); + assert(sub.size() == 1); + assert(equal_to(sub, "R")); + } + } + + return true; +} + +int main(int, char**) +{ + test(); + // static_assert(test(), ""); + return 0; +}