diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/CMakeLists.txt b/libcudacxx/include/cuda/std/detail/libcxx/include/CMakeLists.txt index 067a9e03920..5420d5d8689 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/CMakeLists.txt +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/CMakeLists.txt @@ -5,6 +5,10 @@ set(files __algorithm/any_of.h __algorithm/comp_ref_type.h __algorithm/comp.h + __algorithm/copy_backward.h + __algorithm/copy_if.h + __algorithm/copy_n.h + __algorithm/copy.h __algorithm/count.h __algorithm/count_if.h __algorithm/equal.h @@ -17,14 +21,27 @@ set(files __algorithm/find.h __algorithm/for_each_n.h __algorithm/for_each.h + __algorithm/generate_n.h + __algorithm/generate.h __algorithm/half_positive.h __algorithm/is_permutation.h __algorithm/lexicographical_compare.h __algorithm/mismatch.h + __algorithm/move_backward.h + __algorithm/move.h __algorithm/none_of.h + __algorithm/remove_copy_if.h + __algorithm/remove_copy.h + __algorithm/remove_if.h + __algorithm/remove.h + __algorithm/replace_copy_if.h + __algorithm/replace_copy.h + __algorithm/replace_if.h + __algorithm/replace.h __algorithm/search_n.h __algorithm/search.h __algorithm/swap_ranges.h + __algorithm/transform.h __availability __bit_reference __bsd_locale_defaults.h diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/copy.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/copy.h new file mode 100644 index 00000000000..dd600cfc9ee --- /dev/null +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/copy.h @@ -0,0 +1,102 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___ALGORITHM_COPY_H +#define _LIBCUDACXX___ALGORITHM_COPY_H + +#ifndef __cuda_std__ +# include <__config> +#endif // __cuda_std__ + +#include "../__algorithm/unwrap_iter.h" +#include "../__type_traits/enable_if.h" +#include "../__type_traits/is_constant_evaluated.h" +#include "../__type_traits/is_same.h" +#include "../__type_traits/is_trivially_copy_assignable.h" +#include "../__type_traits/remove_const.h" +#include "../cstdlib" +#include "../cstring" + +#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 + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _OutputIterator +__copy(_InputIterator __first, _InputIterator __last, _OutputIterator __result) +{ + for (; __first != __last; ++__first, (void) ++__result) + { + *__result = *__first; + } + return __result; +} + +template +inline _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 bool +__dispatch_memmove(_Up* __result, _Tp* __first, const size_t __n) +{ +#if _LIBCUDACXX_STD_VER >= 20 + if (_CUDA_VSTD::is_constant_evaluated()) + { + return false; + } + else +#endif // _LIBCUDACXX_STD_VER >= 20 + { + // For now, we only ever use memmove on host + // clang-format off + NV_IF_ELSE_TARGET(NV_IS_HOST, ( + _CUDA_VSTD::memmove(__result, __first, __n * sizeof(_Up)); + return true; + ),( + return false; + )) + // clang-format on + } +} + +template , _Up), int> = 0, + __enable_if_t<_LIBCUDACXX_TRAIT(is_trivially_copy_assignable, _Up), int> = 0> +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 _Up* +__copy(_Tp* __first, _Tp* __last, _Up* __result) +{ + const ptrdiff_t __n = __last - __first; + if (__n > 0) + { + if (__dispatch_memmove(__result, __first, __n)) + { + return __result + __n; + } + for (ptrdiff_t __i = 0; __i < __n; ++__i) + { + *(__result + __i) = *(__first + __i); + } + } + return __result + __n; +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 _OutputIterator +copy(_InputIterator __first, _InputIterator __last, _OutputIterator __result) +{ + return _CUDA_VSTD::__copy(__unwrap_iter(__first), __unwrap_iter(__last), __unwrap_iter(__result)); +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___ALGORITHM_COPY_H diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/copy_backward.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/copy_backward.h new file mode 100644 index 00000000000..249b085cf0b --- /dev/null +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/copy_backward.h @@ -0,0 +1,76 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___ALGORITHM_COPY_BACKWARD_H +#define _LIBCUDACXX___ALGORITHM_COPY_BACKWARD_H + +#ifndef __cuda_std__ +# include <__config> +#endif // __cuda_std__ + +#include "../__algorithm/copy.h" +#include "../__algorithm/unwrap_iter.h" +#include "../__type_traits/enable_if.h" +#include "../__type_traits/is_same.h" +#include "../__type_traits/is_trivially_copy_assignable.h" +#include "../__type_traits/remove_const.h" + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _OutputIterator +__copy_backward(_BidirectionalIterator __first, _BidirectionalIterator __last, _OutputIterator __result) +{ + while (__first != __last) + { + *--__result = *--__last; + } + return __result; +} + +template , _Up), int> = 0, + __enable_if_t<_LIBCUDACXX_TRAIT(is_trivially_copy_assignable, _Up), int> = 0> +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 _Up* +__copy_backward(_Tp* __first, _Tp* __last, _Up* __result) +{ + const ptrdiff_t __n = __last - __first; + if (__n > 0) + { + if (__dispatch_memmove(__result - __n, __first, __n)) + { + return __result - __n; + } + for (ptrdiff_t __i = 1; __i <= __n; ++__i) + { + *(__result - __i) = *(__last - __i); + } + } + return __result - __n; +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 _BidirectionalIterator2 +copy_backward(_BidirectionalIterator1 __first, _BidirectionalIterator1 __last, _BidirectionalIterator2 __result) +{ + return _CUDA_VSTD::__copy_backward(__unwrap_iter(__first), __unwrap_iter(__last), __unwrap_iter(__result)); +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___ALGORITHM_COPY_BACKWARD_H diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/copy_if.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/copy_if.h new file mode 100644 index 00000000000..47be01ba0f0 --- /dev/null +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/copy_if.h @@ -0,0 +1,45 @@ +//===----------------------------------------------------------------------===// +// +// 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) 2023 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___ALGORITHM_COPY_IF_H +#define _LIBCUDACXX___ALGORITHM_COPY_IF_H + +#ifndef __cuda_std__ +# include <__config> +#endif // __cuda_std__ + +#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 + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _OutputIterator +copy_if(_InputIterator __first, _InputIterator __last, + _OutputIterator __result, _Predicate __pred) +{ + for (; __first != __last; ++__first) + { + if (__pred(*__first)) + { + *__result = *__first; + ++__result; + } + } + return __result; +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___ALGORITHM_COPY_IF_H diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/copy_n.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/copy_n.h new file mode 100644 index 00000000000..40f56530bb1 --- /dev/null +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/copy_n.h @@ -0,0 +1,70 @@ +//===----------------------------------------------------------------------===// +// +// 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) 2023 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___ALGORITHM_COPY_N_H +#define _LIBCUDACXX___ALGORITHM_COPY_N_H + +#ifndef __cuda_std__ +# include <__config> +#endif // __cuda_std__ + +#include "../__algorithm/copy.h" +#include "../__iterator/iterator_traits.h" +#include "../__type_traits/enable_if.h" +#include "../__utility/convert_to_integral.h" + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +template ::value, int> = 0, + __enable_if_t::value, int> = 0> +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 _OutputIterator +copy_n(_InputIterator __first, _Size __orig_n, _OutputIterator __result) +{ + using _IntegralSize = decltype(__convert_to_integral(__orig_n)); + _IntegralSize __n = static_cast<_IntegralSize>(__orig_n); + if (__n > 0) + { + *__result = *__first; + ++__result; + for (--__n; __n > 0; --__n) + { + ++__first; + *__result = *__first; + ++__result; + } + } + return __result; +} + +template ::value, int> = 0> +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 _OutputIterator +copy_n(_InputIterator __first, _Size __orig_n, _OutputIterator __result) +{ + using _IntegralSize = decltype(__convert_to_integral(__orig_n)); + _IntegralSize __n = static_cast<_IntegralSize>(__orig_n); + return _CUDA_VSTD::copy(__first, __first + __n, __result); +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___ALGORITHM_COPY_N_H diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/generate.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/generate.h new file mode 100644 index 00000000000..5d60aa017af --- /dev/null +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/generate.h @@ -0,0 +1,39 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___ALGORITHM_GENERATE_H +#define _LIBCUDACXX___ALGORITHM_GENERATE_H + +#ifndef __cuda_std__ +# include <__config> +#endif // __cuda_std__ + +#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 + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 void +generate(_ForwardIterator __first, _ForwardIterator __last, _Generator __gen) +{ + for (; __first != __last; ++__first) + { + *__first = __gen(); + } +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___ALGORITHM_GENERATE_H diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/generate_n.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/generate_n.h new file mode 100644 index 00000000000..ef9e797cbca --- /dev/null +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/generate_n.h @@ -0,0 +1,44 @@ +//===----------------------------------------------------------------------===// +// +// 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) 2023 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___ALGORITHM_GENERATE_N_H +#define _LIBCUDACXX___ALGORITHM_GENERATE_N_H + +#ifndef __cuda_std__ +# include <__config> +#endif // __cuda_std__ + +#include "../__utility/convert_to_integral.h" + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _OutputIterator +generate_n(_OutputIterator __first, _Size __orig_n, _Generator __gen) +{ + using _IntegralSize = decltype(__convert_to_integral(__orig_n)); + _IntegralSize __n = static_cast<_IntegralSize>(__orig_n); + for (; __n > 0; ++__first, (void) --__n) + { + *__first = __gen(); + } + return __first; +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___ALGORITHM_GENERATE_N_H diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/move.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/move.h new file mode 100644 index 00000000000..31a1cae0de3 --- /dev/null +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/move.h @@ -0,0 +1,73 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___ALGORITHM_MOVE_H +#define _LIBCUDACXX___ALGORITHM_MOVE_H + +#ifndef __cuda_std__ +# include <__config> +#endif // __cuda_std__ + +#include "../__algorithm/copy.h" +#include "../__algorithm/unwrap_iter.h" +#include "../__type_traits/enable_if.h" +#include "../__type_traits/is_same.h" +#include "../__type_traits/is_trivially_move_assignable.h" +#include "../__type_traits/remove_const.h" + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 _OutputIterator +__move(_InputIterator __first, _InputIterator __last, _OutputIterator __result) +{ + for (; __first != __last; ++__first, (void) ++__result) + *__result = _CUDA_VSTD::move(*__first); + return __result; +} + +template , _Up), int> = 0, + __enable_if_t<_LIBCUDACXX_TRAIT(is_trivially_move_assignable, _Up), int> = 0> +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 _Up* +__move(_Tp* __first, _Tp* __last, _Up* __result) +{ + const ptrdiff_t __n = __last - __first; + if (__n > 0) { + if (__dispatch_memmove(__result, __first, __n)) + { + return __result + __n; + } + for (ptrdiff_t __i = 0; __i < __n; ++__i) + { + *(__result + __i) = _CUDA_VSTD::move(*(__first + __i)); + } + } + return __result + __n; +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 _OutputIterator +move(_InputIterator __first, _InputIterator __last, _OutputIterator __result) +{ + return _CUDA_VSTD::__move(__unwrap_iter(__first), __unwrap_iter(__last), __unwrap_iter(__result)); +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___ALGORITHM_MOVE_H diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/move_backward.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/move_backward.h new file mode 100644 index 00000000000..3a03e5e0541 --- /dev/null +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/move_backward.h @@ -0,0 +1,76 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___ALGORITHM_MOVE_BACKWARD_H +#define _LIBCUDACXX___ALGORITHM_MOVE_BACKWARD_H + +#ifndef __cuda_std__ +# include <__config> +#endif // __cuda_std__ + +#include "../__algorithm/copy.h" +#include "../__algorithm/unwrap_iter.h" +#include "../__type_traits/enable_if.h" +#include "../__type_traits/is_same.h" +#include "../__type_traits/is_trivially_move_assignable.h" +#include "../__type_traits/remove_const.h" + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _OutputIterator +__move_backward(_BidirectionalIterator __first, _BidirectionalIterator __last, _OutputIterator __result) +{ + while (__first != __last) + { + *--__result = _CUDA_VSTD::move(*--__last); + } + return __result; +} + +template , _Up), int> = 0, + __enable_if_t<_LIBCUDACXX_TRAIT(is_trivially_move_assignable, _Up), int> = 0> +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 _Up* +__move_backward(_Tp* __first, _Tp* __last, _Up* __result) +{ + const ptrdiff_t __n = __last - __first; + if (__n > 0) + { + if (__dispatch_memmove(__result - __n, __first, __n)) + { + return __result - __n; + } + for (ptrdiff_t __i = 1; __i <= __n; ++__i) + { + *(__result - __i) = _CUDA_VSTD::move(*(__last - __i)); + } + } + return __result - __n; +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 _BidirectionalIterator2 +move_backward(_BidirectionalIterator1 __first, _BidirectionalIterator1 __last, _BidirectionalIterator2 __result) +{ + return _CUDA_VSTD::__move_backward(__unwrap_iter(__first), __unwrap_iter(__last), __unwrap_iter(__result)); +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___ALGORITHM_MOVE_BACKWARD_H diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/remove.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/remove.h new file mode 100644 index 00000000000..dd05d640ccf --- /dev/null +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/remove.h @@ -0,0 +1,53 @@ +//===----------------------------------------------------------------------===// +// +// 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) 2023 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___ALGORITHM_REMOVE_H +#define _LIBCUDACXX___ALGORITHM_REMOVE_H + +#ifndef __cuda_std__ +# include <__config> +#endif // __cuda_std__ + +#include "../__algorithm/find.h" +#include "../__utility/move.h" + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +template +_LIBCUDACXX_NODISCARD_EXT inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 + _ForwardIterator + remove(_ForwardIterator __first, _ForwardIterator __last, const _Tp& __value_) +{ + __first = _CUDA_VSTD::find(__first, __last, __value_); + if (__first != __last) + { + _ForwardIterator __i = __first; + while (++__i != __last) + { + if (!(*__i == __value_)) + { + *__first = _CUDA_VSTD::move(*__i); + ++__first; + } + } + } + return __first; +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___ALGORITHM_REMOVE_H diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/remove_copy.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/remove_copy.h new file mode 100644 index 00000000000..f3701e71b8e --- /dev/null +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/remove_copy.h @@ -0,0 +1,45 @@ +//===----------------------------------------------------------------------===// +// +// 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) 2023 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___ALGORITHM_REMOVE_COPY_H +#define _LIBCUDACXX___ALGORITHM_REMOVE_COPY_H + +#ifndef __cuda_std__ +# include <__config> +#endif // __cuda_std__ + +#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 + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 +_OutputIterator +remove_copy(_InputIterator __first, _InputIterator __last, _OutputIterator __result, const _Tp& __value_) +{ + for (; __first != __last; ++__first) + { + if (!(*__first == __value_)) + { + *__result = *__first; + ++__result; + } + } + return __result; +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___ALGORITHM_REMOVE_COPY_H diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/remove_copy_if.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/remove_copy_if.h new file mode 100644 index 00000000000..8443eee8a76 --- /dev/null +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/remove_copy_if.h @@ -0,0 +1,44 @@ +//===----------------------------------------------------------------------===// +// +// 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) 2023 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___ALGORITHM_REMOVE_COPY_IF_H +#define _LIBCUDACXX___ALGORITHM_REMOVE_COPY_IF_H + +#ifndef __cuda_std__ +# include <__config> +#endif // __cuda_std__ + +#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 + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _OutputIterator +remove_copy_if(_InputIterator __first, _InputIterator __last, _OutputIterator __result, _Predicate __pred) +{ + for (; __first != __last; ++__first) + { + if (!__pred(*__first)) + { + *__result = *__first; + ++__result; + } + } + return __result; +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___ALGORITHM_REMOVE_COPY_IF_H diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/remove_if.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/remove_if.h new file mode 100644 index 00000000000..7aca5396ac9 --- /dev/null +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/remove_if.h @@ -0,0 +1,54 @@ +//===----------------------------------------------------------------------===// +// +// 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) 2023 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___ALGORITHM_REMOVE_IF_H +#define _LIBCUDACXX___ALGORITHM_REMOVE_IF_H + +#ifndef __cuda_std__ +# include <__config> +#endif // __cuda_std__ + +#include "../__algorithm/find_if.h" +#include "../__type_traits/add_lvalue_reference.h" +#include "../__utility/move.h" + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +template +_LIBCUDACXX_NODISCARD_EXT inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 + _ForwardIterator + remove_if(_ForwardIterator __first, _ForwardIterator __last, _Predicate __pred) +{ + __first = _CUDA_VSTD::find_if<_ForwardIterator, __add_lvalue_reference_t<_Predicate>>(__first, __last, __pred); + if (__first != __last) + { + _ForwardIterator __i = __first; + while (++__i != __last) + { + if (!__pred(*__i)) + { + *__first = _CUDA_VSTD::move(*__i); + ++__first; + } + } + } + return __first; +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___ALGORITHM_REMOVE_IF_H diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/replace.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/replace.h new file mode 100644 index 00000000000..06964cb1102 --- /dev/null +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/replace.h @@ -0,0 +1,42 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___ALGORITHM_REPLACE_H +#define _LIBCUDACXX___ALGORITHM_REPLACE_H + +#ifndef __cuda_std__ +# include <__config> +#endif // __cuda_std__ + +#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 + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 void +replace(_ForwardIterator __first, _ForwardIterator __last, const _Tp& __old_value, const _Tp& __new_value) +{ + for (; __first != __last; ++__first) + { + if (*__first == __old_value) + { + *__first = __new_value; + } + } +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___ALGORITHM_REPLACE_H diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/replace_copy.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/replace_copy.h new file mode 100644 index 00000000000..519d524311e --- /dev/null +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/replace_copy.h @@ -0,0 +1,51 @@ +//===----------------------------------------------------------------------===// +// +// 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) 2023 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___ALGORITHM_REPLACE_COPY_H +#define _LIBCUDACXX___ALGORITHM_REPLACE_COPY_H + +#ifndef __cuda_std__ +# include <__config> +#endif // __cuda_std__ + +#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 + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _OutputIterator +replace_copy(_InputIterator __first, + _InputIterator __last, + _OutputIterator __result, + const _Tp& __old_value, + const _Tp& __new_value) +{ + for (; __first != __last; ++__first, (void) ++__result) + { + if (*__first == __old_value) + { + *__result = __new_value; + } + else + { + *__result = *__first; + } + } + return __result; +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___ALGORITHM_REPLACE_COPY_H diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/replace_copy_if.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/replace_copy_if.h new file mode 100644 index 00000000000..b3602a4997b --- /dev/null +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/replace_copy_if.h @@ -0,0 +1,48 @@ +//===----------------------------------------------------------------------===// +// +// 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) 2023 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___ALGORITHM_REPLACE_COPY_IF_H +#define _LIBCUDACXX___ALGORITHM_REPLACE_COPY_IF_H + +#ifndef __cuda_std__ +# include <__config> +#endif // __cuda_std__ + +#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 + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _OutputIterator +replace_copy_if( + _InputIterator __first, _InputIterator __last, _OutputIterator __result, _Predicate __pred, const _Tp& __new_value) +{ + for (; __first != __last; ++__first, (void) ++__result) + { + if (__pred(*__first)) + { + *__result = __new_value; + } + else + { + *__result = *__first; + } + } + return __result; +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___ALGORITHM_REPLACE_COPY_IF_H diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/replace_if.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/replace_if.h new file mode 100644 index 00000000000..0bf4bbb39af --- /dev/null +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/replace_if.h @@ -0,0 +1,42 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___ALGORITHM_REPLACE_IF_H +#define _LIBCUDACXX___ALGORITHM_REPLACE_IF_H + +#ifndef __cuda_std__ +# include <__config> +#endif // __cuda_std__ + +#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 + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 void +replace_if(_ForwardIterator __first, _ForwardIterator __last, _Predicate __pred, const _Tp& __new_value) +{ + for (; __first != __last; ++__first) + { + if (__pred(*__first)) + { + *__first = __new_value; + } + } +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___ALGORITHM_REPLACE_IF_H diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/transform.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/transform.h new file mode 100644 index 00000000000..801ab416502 --- /dev/null +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__algorithm/transform.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) 2023 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___ALGORITHM_TRANSFORM_H +#define _LIBCUDACXX___ALGORITHM_TRANSFORM_H + +#ifndef __cuda_std__ +# include <__config> +#endif // __cuda_std__ + +#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 + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _OutputIterator +transform(_InputIterator __first, _InputIterator __last, _OutputIterator __result, _UnaryOperation __op) +{ + for (; __first != __last; ++__first, (void) ++__result) + { + *__result = __op(*__first); + } + return __result; +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _OutputIterator +transform(_InputIterator1 __first1, + _InputIterator1 __last1, + _InputIterator2 __first2, + _OutputIterator __result, + _BinaryOperation __binary_op) +{ + for (; __first1 != __last1; ++__first1, (void) ++__first2, ++__result) + { + *__result = __binary_op(*__first1, *__first2); + } + return __result; +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___ALGORITHM_TRANSFORM_H diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/algorithm b/libcudacxx/include/cuda/std/detail/libcxx/include/algorithm index dc4b0402699..e8565a86c88 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/algorithm +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/algorithm @@ -645,6 +645,10 @@ template #include "__algorithm/any_of.h" #include "__algorithm/comp_ref_type.h" #include "__algorithm/comp.h" +#include "__algorithm/copy_backward.h" +#include "__algorithm/copy_if.h" +#include "__algorithm/copy_n.h" +#include "__algorithm/copy.h" #include "__algorithm/count_if.h" #include "__algorithm/count.h" #include "__algorithm/equal.h" @@ -657,14 +661,27 @@ template #include "__algorithm/find.h" #include "__algorithm/for_each_n.h" #include "__algorithm/for_each.h" +#include "__algorithm/generate_n.h" +#include "__algorithm/generate.h" #include "__algorithm/half_positive.h" #include "__algorithm/is_permutation.h" #include "__algorithm/lexicographical_compare.h" #include "__algorithm/mismatch.h" +#include "__algorithm/move_backward.h" +#include "__algorithm/move.h" #include "__algorithm/none_of.h" +#include "__algorithm/remove_copy_if.h" +#include "__algorithm/remove_copy.h" +#include "__algorithm/remove_if.h" +#include "__algorithm/remove.h" +#include "__algorithm/replace_copy_if.h" +#include "__algorithm/replace_copy.h" +#include "__algorithm/replace_if.h" +#include "__algorithm/replace.h" #include "__algorithm/search_n.h" #include "__algorithm/search.h" #include "__algorithm/swap_ranges.h" +#include "__algorithm/transform.h" #include "__assert" // all public C++ headers provide the assertion handler #include "__debug" #include "__iterator/distance.h" @@ -724,473 +741,6 @@ public: bool operator()(const _T1& __x, const _T2& __y) {return __p_(__y, __x);} }; -// copy -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_Iter -__unwrap_iter(_Iter __i) -{ - return __i; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -__enable_if_t -< - is_trivially_copy_assignable<_Tp>::value, - _Tp* -> -__unwrap_iter(move_iterator<_Tp*> __i) -{ - return __i.base(); -} - -#if _LIBCUDACXX_DEBUG_LEVEL < 2 - -template -inline _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_IF_NODEBUG -__enable_if_t -< - is_trivially_copy_assignable<_Tp>::value, - _Tp* -> -__unwrap_iter(__wrap_iter<_Tp*> __i) -{ - return __i.base(); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_IF_NODEBUG -__enable_if_t -< - is_trivially_copy_assignable<_Tp>::value, - const _Tp* -> -__unwrap_iter(__wrap_iter __i) -{ - return __i.base(); -} - -#else - -template -inline _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_IF_NODEBUG -__enable_if_t -< - is_trivially_copy_assignable<_Tp>::value, - __wrap_iter<_Tp*> -> -__unwrap_iter(__wrap_iter<_Tp*> __i) -{ - return __i; -} - -#endif // _LIBCUDACXX_DEBUG_LEVEL < 2 - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_OutputIterator -__copy(_InputIterator __first, _InputIterator __last, _OutputIterator __result) -{ - for (; __first != __last; ++__first, (void) ++__result) - *__result = *__first; - return __result; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -__enable_if_t -< - is_same<__remove_const_t<_Tp>, _Up>::value && - is_trivially_copy_assignable<_Up>::value, - _Up* -> -__copy(_Tp* __first, _Tp* __last, _Up* __result) -{ - const size_t __n = static_cast(__last - __first); - if (__n > 0) - _CUDA_VSTD::memmove(__result, __first, __n * sizeof(_Up)); - return __result + __n; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_OutputIterator -copy(_InputIterator __first, _InputIterator __last, _OutputIterator __result) -{ - return _CUDA_VSTD::__copy(__unwrap_iter(__first), __unwrap_iter(__last), __unwrap_iter(__result)); -} - -// copy_backward - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_OutputIterator -__copy_backward(_BidirectionalIterator __first, _BidirectionalIterator __last, _OutputIterator __result) -{ - while (__first != __last) - *--__result = *--__last; - return __result; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -__enable_if_t -< - is_same<__remove_const_t<_Tp>, _Up>::value && - is_trivially_copy_assignable<_Up>::value, - _Up* -> -__copy_backward(_Tp* __first, _Tp* __last, _Up* __result) -{ - const size_t __n = static_cast(__last - __first); - if (__n > 0) - { - __result -= __n; - _CUDA_VSTD::memmove(__result, __first, __n * sizeof(_Up)); - } - return __result; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_BidirectionalIterator2 -copy_backward(_BidirectionalIterator1 __first, _BidirectionalIterator1 __last, - _BidirectionalIterator2 __result) -{ - return _CUDA_VSTD::__copy_backward(__unwrap_iter(__first), - __unwrap_iter(__last), - __unwrap_iter(__result)); -} - -// copy_if - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_OutputIterator -copy_if(_InputIterator __first, _InputIterator __last, - _OutputIterator __result, _Predicate __pred) -{ - for (; __first != __last; ++__first) - { - if (__pred(*__first)) - { - *__result = *__first; - ++__result; - } - } - return __result; -} - -// copy_n - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -__enable_if_t -< - __is_cpp17_input_iterator<_InputIterator>::value && - !__is_cpp17_random_access_iterator<_InputIterator>::value, - _OutputIterator -> -copy_n(_InputIterator __first, _Size __orig_n, _OutputIterator __result) -{ - typedef decltype(__convert_to_integral(__orig_n)) _IntegralSize; - _IntegralSize __n = __orig_n; - if (__n > 0) - { - *__result = *__first; - ++__result; - for (--__n; __n > 0; --__n) - { - ++__first; - *__result = *__first; - ++__result; - } - } - return __result; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -__enable_if_t -< - __is_cpp17_random_access_iterator<_InputIterator>::value, - _OutputIterator -> -copy_n(_InputIterator __first, _Size __orig_n, _OutputIterator __result) -{ - typedef decltype(__convert_to_integral(__orig_n)) _IntegralSize; - _IntegralSize __n = __orig_n; - return _CUDA_VSTD::copy(__first, __first + __n, __result); -} - -// move - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_OutputIterator -__move(_InputIterator __first, _InputIterator __last, _OutputIterator __result) -{ - for (; __first != __last; ++__first, (void) ++__result) - *__result = _CUDA_VSTD::move(*__first); - return __result; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -__enable_if_t -< - is_same<__remove_const_t<_Tp>, _Up>::value && - is_trivially_copy_assignable<_Up>::value, - _Up* -> -__move(_Tp* __first, _Tp* __last, _Up* __result) -{ - const size_t __n = static_cast(__last - __first); - if (__n > 0) - _CUDA_VSTD::memmove(__result, __first, __n * sizeof(_Up)); - return __result + __n; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_OutputIterator -move(_InputIterator __first, _InputIterator __last, _OutputIterator __result) -{ - return _CUDA_VSTD::__move(__unwrap_iter(__first), __unwrap_iter(__last), __unwrap_iter(__result)); -} - -// move_backward - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_OutputIterator -__move_backward(_InputIterator __first, _InputIterator __last, _OutputIterator __result) -{ - while (__first != __last) - *--__result = _CUDA_VSTD::move(*--__last); - return __result; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -__enable_if_t -< - is_same<__remove_const_t<_Tp>, _Up>::value && - is_trivially_copy_assignable<_Up>::value, - _Up* -> -__move_backward(_Tp* __first, _Tp* __last, _Up* __result) -{ - const size_t __n = static_cast(__last - __first); - if (__n > 0) - { - __result -= __n; - _CUDA_VSTD::memmove(__result, __first, __n * sizeof(_Up)); - } - return __result; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_BidirectionalIterator2 -move_backward(_BidirectionalIterator1 __first, _BidirectionalIterator1 __last, - _BidirectionalIterator2 __result) -{ - return _CUDA_VSTD::__move_backward(__unwrap_iter(__first), __unwrap_iter(__last), __unwrap_iter(__result)); -} - -// iter_swap - -// moved to for better swap / noexcept support - -// transform - -template -inline _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 -_OutputIterator -transform(_InputIterator __first, _InputIterator __last, _OutputIterator __result, _UnaryOperation __op) -{ - for (; __first != __last; ++__first, (void) ++__result) - *__result = __op(*__first); - return __result; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 -_OutputIterator -transform(_InputIterator1 __first1, _InputIterator1 __last1, _InputIterator2 __first2, - _OutputIterator __result, _BinaryOperation __binary_op) -{ - for (; __first1 != __last1; ++__first1, (void) ++__first2, ++__result) - *__result = __binary_op(*__first1, *__first2); - return __result; -} - -// replace - -template -inline _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 -void -replace(_ForwardIterator __first, _ForwardIterator __last, const _Tp& __old_value, const _Tp& __new_value) -{ - for (; __first != __last; ++__first) - if (*__first == __old_value) - *__first = __new_value; -} - -// replace_if - -template -inline _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 -void -replace_if(_ForwardIterator __first, _ForwardIterator __last, _Predicate __pred, const _Tp& __new_value) -{ - for (; __first != __last; ++__first) - if (__pred(*__first)) - *__first = __new_value; -} - -// replace_copy - -template -inline _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 -_OutputIterator -replace_copy(_InputIterator __first, _InputIterator __last, _OutputIterator __result, - const _Tp& __old_value, const _Tp& __new_value) -{ - for (; __first != __last; ++__first, (void) ++__result) - if (*__first == __old_value) - *__result = __new_value; - else - *__result = *__first; - return __result; -} - -// replace_copy_if - -template -_LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 -inline _OutputIterator -replace_copy_if(_InputIterator __first, _InputIterator __last, _OutputIterator __result, - _Predicate __pred, const _Tp& __new_value) -{ - for (; __first != __last; ++__first, (void) ++__result) - if (__pred(*__first)) - *__result = __new_value; - else - *__result = *__first; - return __result; -} - -// generate - -template -inline _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 -void -generate(_ForwardIterator __first, _ForwardIterator __last, _Generator __gen) -{ - for (; __first != __last; ++__first) - *__first = __gen(); -} - -// generate_n - -template -inline _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 -_OutputIterator -generate_n(_OutputIterator __first, _Size __orig_n, _Generator __gen) -{ - typedef decltype(__convert_to_integral(__orig_n)) _IntegralSize; - _IntegralSize __n = __orig_n; - for (; __n > 0; ++__first, (void) --__n) - *__first = __gen(); - return __first; -} - -// remove - -template -_LIBCUDACXX_NODISCARD_EXT _LIBCUDACXX_INLINE_VISIBILITY -_LIBCUDACXX_CONSTEXPR_AFTER_CXX17 _ForwardIterator -remove(_ForwardIterator __first, _ForwardIterator __last, const _Tp& __value_) -{ - __first = _CUDA_VSTD::find(__first, __last, __value_); - if (__first != __last) - { - _ForwardIterator __i = __first; - while (++__i != __last) - { - if (!(*__i == __value_)) - { - *__first = _CUDA_VSTD::move(*__i); - ++__first; - } - } - } - return __first; -} - -// remove_if - -template -_LIBCUDACXX_NODISCARD_EXT _LIBCUDACXX_INLINE_VISIBILITY -_LIBCUDACXX_CONSTEXPR_AFTER_CXX17 _ForwardIterator -remove_if(_ForwardIterator __first, _ForwardIterator __last, _Predicate __pred) -{ - __first = _CUDA_VSTD::find_if<_ForwardIterator, __add_lvalue_reference_t<_Predicate>> - (__first, __last, __pred); - if (__first != __last) - { - _ForwardIterator __i = __first; - while (++__i != __last) - { - if (!__pred(*__i)) - { - *__first = _CUDA_VSTD::move(*__i); - ++__first; - } - } - } - return __first; -} - -// remove_copy - -template -inline _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 -_OutputIterator -remove_copy(_InputIterator __first, _InputIterator __last, _OutputIterator __result, const _Tp& __value_) -{ - for (; __first != __last; ++__first) - { - if (!(*__first == __value_)) - { - *__result = *__first; - ++__result; - } - } - return __result; -} - -// remove_copy_if - -template -inline _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX17 -_OutputIterator -remove_copy_if(_InputIterator __first, _InputIterator __last, _OutputIterator __result, _Predicate __pred) -{ - for (; __first != __last; ++__first) - { - if (!__pred(*__first)) - { - *__result = *__first; - ++__result; - } - } - return __result; -} - // unique template diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/cstring b/libcudacxx/include/cuda/std/detail/libcxx/include/cstring index 7bdca8cfff0..7eaf5349c98 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/cstring +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/cstring @@ -57,10 +57,14 @@ size_t strlen(const char* s); */ #ifndef __cuda_std__ -#include <__config> -#include -#include <__pragma_push> -#endif //__simet__ <- a tribute to ogiroux +# include <__config> +#endif // __cuda_std__ + +#if defined(_LIBCUDACXX_COMPILER_MSVC) +#include +#else // ^^^ _LIBCUDACXX_COMPILER_MSVC ^^^ / vvv !_LIBCUDACXX_COMPILER_MSVC vvv +#include "string.h" +#endif // !_LIBCUDACXX_COMPILER_MSVC #if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) # pragma GCC system_header @@ -72,36 +76,35 @@ size_t strlen(const char* s); _LIBCUDACXX_BEGIN_NAMESPACE_STD -using ::size_t; using ::memcpy; +using ::memset; +using ::size_t; + +#ifndef _LIBCUDACXX_COMPILER_NVRTC +using ::memchr; +using ::memcmp; using ::memmove; -using ::strcpy; -using ::strncpy; using ::strcat; -using ::strncat; -using ::memcmp; +using ::strchr; using ::strcmp; -using ::strncmp; using ::strcoll; -using ::strxfrm; -using ::memchr; -using ::strchr; +using ::strcpy; using ::strcspn; +using ::strncat; +using ::strncmp; +using ::strncpy; using ::strpbrk; using ::strrchr; using ::strspn; using ::strstr; +using ::strxfrm; #ifndef _LIBCUDACXX_HAS_NO_THREAD_UNSAFE_C_FUNCTIONS using ::strtok; #endif -using ::memset; using ::strerror; using ::strlen; +#endif // _LIBCUDACXX_COMPILER_NVRTC _LIBCUDACXX_END_NAMESPACE_STD -#ifndef __cuda_std__ -#include <__pragma_pop> -#endif //__cuda_std__ - -#endif // _LIBCUDACXX_CSTRING +#endif // _LIBCUDACXX_CSTRING diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/string.h b/libcudacxx/include/cuda/std/detail/libcxx/include/string.h index 90c4fe48bf9..ba410c42ab3 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/string.h +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/string.h @@ -51,7 +51,9 @@ size_t strlen(const char* s); */ -#include <__config> +#ifndef __cuda_std__ +# include <__config> +#endif // __cuda_std__ #if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) # pragma GCC system_header @@ -61,7 +63,9 @@ size_t strlen(const char* s); # pragma system_header #endif // no system header +#if !defined(_LIBCUDACXX_COMPILER_NVRTC) && !defined(_LIBCUDACXX_COMPILER_MSVC) #include_next +#endif // !_LIBCUDACXX_COMPILER_NVRTC && !_LIBCUDACXX_COMPILER_MSVC // MSVCRT, GNU libc and its derivates may already have the correct prototype in // . This macro can be defined by users if their C library provides diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.copy/copy.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.copy/copy.pass.cpp new file mode 100644 index 00000000000..1e34f0cc70b --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.copy/copy.pass.cpp @@ -0,0 +1,155 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// + +// template OutIter> +// constexpr OutIter // constexpr after C++17 +// copy(InIter first, InIter last, OutIter result); + +#include +#include + +#include "test_macros.h" +#include "test_iterators.h" + +struct NonTrivialCopy { + int data = 0; + bool copy_assigned_from = false; + + NonTrivialCopy() = default; + + __host__ __device__ TEST_CONSTEXPR_CXX20 + NonTrivialCopy(NonTrivialCopy&& other) noexcept : data(other.data), + copy_assigned_from(false) {} + __host__ __device__ TEST_CONSTEXPR_CXX20 + NonTrivialCopy(const NonTrivialCopy& other) noexcept + : data(other.data), + copy_assigned_from(false) {} + + __host__ __device__ TEST_CONSTEXPR_CXX20 NonTrivialCopy& + operator=(const NonTrivialCopy& other) noexcept { + data = other.data; + copy_assigned_from = true; + return *this; + } + + __host__ __device__ TEST_CONSTEXPR_CXX20 NonTrivialCopy& + operator=(NonTrivialCopy&& other) noexcept { + data = other.data; + copy_assigned_from = false; + return *this; + } + + __host__ __device__ TEST_CONSTEXPR_CXX20 + NonTrivialCopy(const int val) noexcept : data(val), + copy_assigned_from(false) {} + __host__ __device__ TEST_CONSTEXPR_CXX20 NonTrivialCopy& + operator=(const int val) noexcept { + data = val; + copy_assigned_from = false; + return *this; + } + + __host__ __device__ TEST_CONSTEXPR_CXX20 friend bool + operator==(const NonTrivialCopy& lhs, const NonTrivialCopy& rhs) noexcept { + // NOTE: This uses implicit knowledge that the right hand side has been copied from + return lhs.data == rhs.data && !lhs.copy_assigned_from && + rhs.copy_assigned_from; + } + __host__ __device__ TEST_CONSTEXPR_CXX20 bool + operator==(const int& other) const noexcept { + // NOTE: This uses implicit knowledge that the only elements we compare against were untouched + return data == other && !copy_assigned_from; + } +}; + +template +TEST_CONSTEXPR_CXX20 __host__ __device__ void test() { + using value_type = typename cuda::std::iterator_traits::value_type; + { + constexpr int N = 1000; + value_type ia[N] = {0}; + for (int i = 0; i < N; ++i) { + ia[i] = i; + } + value_type ib[N] = {0}; + + OutIter r = cuda::std::copy(InIter(ia), InIter(ia + N), OutIter(ib)); + assert(base(r) == ib + N); + for (int i = 0; i < N; ++i) { + assert(ia[i] == ib[i]); + } + } + { + constexpr int N = 5; + value_type ia[N] = {1, 2, 3, 4, 5}; + value_type ic[N + 2] = {6, 6, 6, 6, 6, 6, 6}; + + auto p = cuda::std::copy(ia, ia + 5, ic); + assert(p == (ic + N)); + for (int i = 0; i < N; ++i) { + assert(ia[i] == ic[i]); + } + + for (int i = N; i < N + 2; ++i) { + assert(ic[i] == 6); + } + } +} + +TEST_CONSTEXPR_CXX20 __host__ __device__ bool test() { + test, cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test >(); + test >(); + test >(); + test >(); + test >(); + test(); + + test(); + + return true; +} + +int main(int, char**) { + test(); + +#if TEST_STD_VER >= 20 + static_assert(test()); +#endif // TEST_STD_VER >= 20 + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.copy/copy_backward.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.copy/copy_backward.pass.cpp new file mode 100644 index 00000000000..233b97ccf90 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.copy/copy_backward.pass.cpp @@ -0,0 +1,143 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// + +// template +// requires OutputIterator +// constexpr OutIter // constexpr after C++17 +// copy_backward(InIter first, InIter last, OutIter result); + +#include +#include + +#include "test_macros.h" +#include "test_iterators.h" + +struct NonTrivialCopy { + int data = 0; + bool copy_assigned_from = false; + + NonTrivialCopy() = default; + + __host__ __device__ TEST_CONSTEXPR_CXX20 + NonTrivialCopy(NonTrivialCopy&& other) noexcept : data(other.data), + copy_assigned_from(false) {} + __host__ __device__ TEST_CONSTEXPR_CXX20 + NonTrivialCopy(const NonTrivialCopy& other) noexcept + : data(other.data), + copy_assigned_from(false) {} + + __host__ __device__ TEST_CONSTEXPR_CXX20 NonTrivialCopy& + operator=(const NonTrivialCopy& other) noexcept { + data = other.data; + copy_assigned_from = true; + return *this; + } + + __host__ __device__ TEST_CONSTEXPR_CXX20 NonTrivialCopy& + operator=(NonTrivialCopy&& other) noexcept { + data = other.data; + copy_assigned_from = false; + return *this; + } + + __host__ __device__ TEST_CONSTEXPR_CXX20 + NonTrivialCopy(const int val) noexcept : data(val), + copy_assigned_from(false) {} + __host__ __device__ TEST_CONSTEXPR_CXX20 NonTrivialCopy& + operator=(const int val) noexcept { + data = val; + copy_assigned_from = false; + return *this; + } + + __host__ __device__ TEST_CONSTEXPR_CXX20 friend bool + operator==(const NonTrivialCopy& lhs, const NonTrivialCopy& rhs) noexcept { + // NOTE: This uses implicit knowledge that the right hand side has been copied from + return lhs.data == rhs.data && !lhs.copy_assigned_from && + rhs.copy_assigned_from; + } + __host__ __device__ TEST_CONSTEXPR_CXX20 bool + operator==(const int& other) const noexcept { + return data == other; + } +}; + +template +TEST_CONSTEXPR_CXX20 __host__ __device__ void test() { + using value_type = typename cuda::std::iterator_traits::value_type; + { + constexpr int N = 1000; + value_type ia[N] = {0}; + for (int i = 0; i < N; ++i) { + ia[i] = i; + } + value_type ib[N] = {0}; + + OutIter r = cuda::std::copy_backward(InIter(ia), InIter(ia + N), OutIter(ib + N)); + assert(base(r) == ib); + for (int i = 0; i < N; ++i) { + assert(ia[i] == ib[i]); + } + } + { + constexpr int N = 5; + value_type ia[N] = {1, 2, 3, 4, 5}; + value_type ic[N + 2] = {6, 6, 6, 6, 6, 6, 6}; + + auto p = cuda::std::copy_backward(ia, ia + 5, ic + N); + assert(p == ic); + for (int i = 0; i < N; ++i) { + assert(ia[i] == ic[i]); + } + + for (int i = N; i < N + 2; ++i) { + assert(ic[i] == 6); + } + } + { // Ensure that we are properly preserving back elements when the ranges are overlapping + constexpr int N = 5; + value_type ia[N + 2] = {1, 2, 3, 4, 5, 6, 7}; + int expected[N + 2] = {1, 2, 1, 2, 3, 4, 5}; + + auto p = cuda::std::copy_backward(ia, ia + N, ia + N + 2); + assert(p == ia + 2); + for (int i = 0; i < N; ++i) { + assert(ia[i] == expected[i]); + } + } +} + +TEST_CONSTEXPR_CXX20 __host__ __device__ bool test() { + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test >(); + test >(); + test(); + + test(); + + return true; +} + +int main(int, char**) { + test(); + +#if TEST_STD_VER >= 20 + static_assert(test()); +#endif // TEST_STD_VER >= 20 + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.copy/copy_if.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.copy/copy_if.pass.cpp new file mode 100644 index 00000000000..f6ab1a3f37a --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.copy/copy_if.pass.cpp @@ -0,0 +1,117 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// + +// template OutIter, +// Predicate Pred> +// requires CopyConstructible +// constexpr OutIter // constexpr after C++17 +// copy_if(InIter first, InIter last, OutIter result, Pred pred); + +#include +#include + +#include "test_macros.h" +#include "test_iterators.h" + +struct PredMod3 { + __host__ __device__ constexpr bool operator()(int i) const noexcept { + return i % 3 == 0; + } +}; + +struct PredEqual6 { + __host__ __device__ constexpr bool operator()(int i) const noexcept { + return i == 6; + } +}; + +template +TEST_CONSTEXPR_CXX14 __host__ __device__ void test() { + { + constexpr unsigned N = 1000; + int ia[N] = {0}; + for (unsigned i = 0; i < N; ++i) { + ia[i] = i; + } + int ib[N] = {0}; + + OutIter r = + cuda::std::copy_if(InIter(ia), InIter(ia + N), OutIter(ib), PredMod3{}); + assert(base(r) == ib + N / 3 + 1); + for (unsigned i = 0; i < N / 3 + 1; ++i) { + assert(ib[i] % 3 == 0); + } + } + { + constexpr int N = 5; + constexpr int expected_copies = 2; + int ia[N] = {2, 4, 6, 8, 6}; + int ic[N + 2] = {0, 0, 0, 0, 0, 0}; + + auto p = cuda::std::copy_if(ia, ia + N, ic, PredEqual6{}); + assert(p == (ic + expected_copies)); + for (unsigned i = 0; i < expected_copies; ++i) { + assert(ic[i] == 6); + } + + for (unsigned i = expected_copies; i < N + 2; ++i) { + assert(ic[i] == 0); + } + } +} + +TEST_CONSTEXPR_CXX14 __host__ __device__ bool test() { + test, cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test >(); + test >(); + test >(); + test >(); + test >(); + test(); + + return true; +} + +int main(int, char**) { + test(); + +#if TEST_STD_VER >= 14 && !defined(TEST_COMPILER_MSVC_2017) + static_assert(test(), ""); +#endif // TEST_STD_VER >= 14 && !TEST_COMPILER_MSVC_2017 + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.copy/copy_n.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.copy/copy_n.pass.cpp new file mode 100644 index 00000000000..1e77fcfebf5 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.copy/copy_n.pass.cpp @@ -0,0 +1,104 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// + +// template OutIter> +// constexpr OutIter // constexpr after C++17 +// copy_n(InIter first, InIter::difference_type n, OutIter result); + +#include +#include + +#include "test_macros.h" +#include "test_iterators.h" +#include "user_defined_integral.h" + +using UDI = UserDefinedIntegral; + +template +TEST_CONSTEXPR_CXX20 __host__ __device__ void test() { + { + constexpr unsigned N = 1000; + int ia[N]; + for (unsigned i = 0; i < N; ++i) { + ia[i] = i; + } + int ib[N] = {0}; + + OutIter r = cuda::std::copy_n(InIter(ia), UDI(N / 2), OutIter(ib)); + assert(base(r) == ib + N / 2); + for (unsigned i = 0; i < N / 2; ++i) { + assert(ia[i] == ib[i]); + } + } + { + constexpr int N = 5; + int ia[N] = {1, 2, 3, 4, 5}; + int ic[N + 2] = {6, 6, 6, 6, 6, 6, 6}; + + auto p = cuda::std::copy_n(ia, N, ic); + assert(p == (ic + N)); + for (unsigned i = 0; i < N; ++i) { + assert(ia[i] == ic[i]); + } + + for (unsigned i = N; i < N + 2; ++i) { + assert(ic[i] == 6); + } + } +} + +TEST_CONSTEXPR_CXX20 __host__ __device__ bool test() { + test, cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test >(); + test >(); + test >(); + test >(); + test >(); + test(); + + return true; +} + +int main(int, char**) { + test(); + +#if TEST_STD_VER >= 20 + static_assert(test()); +#endif // TEST_STD_VER >= 20 + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.generate/generate.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.generate/generate.pass.cpp new file mode 100644 index 00000000000..4530488b22b --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.generate/generate.pass.cpp @@ -0,0 +1,61 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// + +// template +// requires OutputIterator +// && CopyConstructible +// constexpr void // constexpr after c++17 +// generate(Iter first, Iter last, Generator gen); + +#include +#include + +#include "test_macros.h" +#include "test_iterators.h" + +struct gen_test { + TEST_CONSTEXPR_CXX14 __host__ __device__ int operator()() const noexcept { + return 1; + } +}; + +template +TEST_CONSTEXPR_CXX14 __host__ __device__ void test() { + constexpr int N = 5; + int ia[N + 1] = {0}; + cuda::std::generate(Iter(ia), Iter(ia + N), gen_test()); + for (int i = 0; i < N; ++i) { + assert(ia[i] == 1); + } + + for (int i = N; i < N + 1; ++i) { + assert(ia[i] == 0); + } +} + +TEST_CONSTEXPR_CXX14 __host__ __device__ bool test() { + test >(); + test >(); + test >(); + test >(); + test(); + + return true; +} + +int main(int, char**) { + test(); + +#if TEST_STD_VER >= 14 + static_assert(test(), ""); +#endif // TEST_STD_VER >= 14 + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.generate/generate_n.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.generate/generate_n.pass.cpp new file mode 100644 index 00000000000..6e1e2a4077a --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.generate/generate_n.pass.cpp @@ -0,0 +1,76 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +// + +// template +// requires OutputIterator +// && CopyConstructible +// constexpr void // constexpr after c++17 +// generate_n(Iter first, Size n, Generator gen); + +#include +#include + +#include "test_macros.h" +#include "test_iterators.h" +#include "user_defined_integral.h" + +struct gen_test { + TEST_CONSTEXPR_CXX14 __host__ __device__ int operator()() const noexcept { + return 1; + } +}; + +template +TEST_CONSTEXPR_CXX14 __host__ __device__ void test() { + constexpr int N = 5; + int ia[N + 1] = {0}; + assert(cuda::std::generate_n(Iter(ia), Size(N), gen_test()) == Iter(ia + N)); + for (int i = 0; i < N; ++i) { + assert(ia[i] == 1); + } + + for (int i = N; i < N + 1; ++i) { + assert(ia[i] == 0); + } +} + +template +TEST_CONSTEXPR_CXX14 __host__ __device__ void test() { + test(); + test(); + test(); + test(); + test >(); + test(); + test(); +#ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE + test(); +#endif // _LIBCUDACXX_HAS_NO_LONG_DOUBLE +} + +TEST_CONSTEXPR_CXX14 __host__ __device__ bool test() { + test >(); + test >(); + test >(); + test >(); + test(); + + return true; +} + +int main(int, char**) { + test(); + +#if TEST_STD_VER >= 14 + static_assert(test(), ""); +#endif // TEST_STD_VER >= 14 + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.move/move.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.move/move.pass.cpp new file mode 100644 index 00000000000..b7b4e1c9da6 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.move/move.pass.cpp @@ -0,0 +1,242 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// + +// template +// requires OutputIterator::type> +// OutIter +// move(InIter first, InIter last, OutIter result); + +#include +#include +#if defined(_LIBCUDACXX_HAS_MEMORY) +#include +#endif // _LIBCUDACXX_HAS_MEMORY + +#include "test_macros.h" +#include "test_iterators.h" + +struct NonTrivialMove { + int data = 0; + bool move_assigned_from = false; + + NonTrivialMove() = default; + __host__ __device__ TEST_CONSTEXPR_CXX20 + NonTrivialMove(const NonTrivialMove& other) noexcept + : data(other.data), + move_assigned_from(false) {} + + __host__ __device__ TEST_CONSTEXPR_CXX20 + NonTrivialMove(NonTrivialMove&& other) noexcept : data(other.data), + move_assigned_from(false) {} + + __host__ __device__ TEST_CONSTEXPR_CXX20 NonTrivialMove& + operator=(const NonTrivialMove& other) noexcept { + data = other.data; + move_assigned_from = false; + return *this; + } + + __host__ __device__ TEST_CONSTEXPR_CXX20 NonTrivialMove& + operator=(NonTrivialMove&& other) noexcept { + data = other.data; + move_assigned_from = true; + return *this; + } + + __host__ __device__ TEST_CONSTEXPR_CXX20 + NonTrivialMove(const int val) noexcept : data(val), + move_assigned_from(false) {} + __host__ __device__ TEST_CONSTEXPR_CXX20 NonTrivialMove& + operator=(const int val) noexcept { + data = val; + move_assigned_from = false; + return *this; + } + + __host__ __device__ TEST_CONSTEXPR_CXX20 friend bool + operator==(const NonTrivialMove& lhs, const NonTrivialMove& rhs) noexcept { + // NOTE: This uses implicit knowledge that the right hand side has been moved from + return lhs.data == rhs.data && !lhs.move_assigned_from && + rhs.move_assigned_from; + } + __host__ __device__ TEST_CONSTEXPR_CXX20 bool + operator==(const int& other) const noexcept { + // NOTE: This uses implicit knowledge that the only elements we compare against were untouched + return data == other && !move_assigned_from; + } +}; + +template +TEST_CONSTEXPR_CXX20 __host__ __device__ void test() { + using value_type = typename cuda::std::iterator_traits::value_type; + { + constexpr int N = 1000; + value_type ia[N] = {0}; + for (int i = 0; i < N; ++i) { + ia[i] = i; + } + value_type ib[N] = {0}; + + OutIter r = cuda::std::move(InIter(ia), InIter(ia + N), OutIter(ib)); + assert(base(r) == ib + N); + for (int i = 0; i < N; ++i) { + assert(ia[i] == ib[i]); + } + } + + { + constexpr int N = 5; + value_type ia[N] = {1, 2, 3, 4, 5}; + value_type ic[N + 2] = {6, 6, 6, 6, 6, 6, 6}; + + auto p = cuda::std::move(ia, ia + 5, ic); + assert(p == (ic + N)); + for (int i = 0; i < N; ++i) { + assert(ia[i] == ic[i]); + } + + for (int i = N; i < N + 2; ++i) { + assert(ic[i] == 6); + } + } +} + +#if defined(_LIBCUDACXX_HAS_MEMORY) +template +TEST_CONSTEXPR_CXX20 __host__ __device__ void test() { + constexpr unsigned N = 100; + cuda::std::unique_ptr ia[N]; + for (unsigned i = 0; i < N; ++i) + ia[i].reset(new int(i)); + cuda::std::unique_ptr ib[N]; + + OutIter r = cuda::std::move(InIter(ia), InIter(ia + N), OutIter(ib)); + assert(base(r) == ib + N); + for (unsigned i = 0; i < N; ++i) + assert(*ib[i] == static_cast(i)); +} +#endif // _LIBCUDACXX_HAS_MEMORY + +TEST_CONSTEXPR_CXX20 __host__ __device__ bool test() { + test, cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test >(); + test >(); + test >(); + test >(); + test >(); + test(); + + test(); + +#if defined(_LIBCUDACXX_HAS_MEMORY) + test1*>, + cpp17_output_iterator*> >(); + test1*>, + cpp17_input_iterator*> >(); + test1*>, + forward_iterator*> >(); + test1*>, + bidirectional_iterator*> >(); + test1*>, + random_access_iterator*> >(); + test1*>, + cuda::std::unique_ptr*>(); + + test1*>, + cpp17_output_iterator*> >(); + test1*>, + cpp17_input_iterator*> >(); + test1*>, + forward_iterator*> >(); + test1*>, + bidirectional_iterator*> >(); + test1*>, + random_access_iterator*> >(); + test1*>, + cuda::std::unique_ptr*>(); + + test1*>, + cpp17_output_iterator*> >(); + test1*>, + cpp17_input_iterator*> >(); + test1*>, + forward_iterator*> >(); + test1*>, + bidirectional_iterator*> >(); + test1*>, + random_access_iterator*> >(); + test1*>, + cuda::std::unique_ptr*>(); + + test1*>, + cpp17_output_iterator*> >(); + test1*>, + cpp17_input_iterator*> >(); + test1*>, + forward_iterator*> >(); + test1*>, + bidirectional_iterator*> >(); + test1*>, + random_access_iterator*> >(); + test1*>, + cuda::std::unique_ptr*>(); + + test1*, + cpp17_output_iterator*> >(); + test1*, + cpp17_input_iterator*> >(); + test1*, + forward_iterator*> >(); + test1*, + bidirectional_iterator*> >(); + test1*, + random_access_iterator*> >(); + test1*, cuda::std::unique_ptr*>(); +#endif // _LIBCUDACXX_HAS_MEMORY + + return true; +} + +int main(int, char**) { + test(); + +#if TEST_STD_VER >= 20 + static_assert(test()); +#endif // TEST_STD_VER >= 20 + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.move/move_backward.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.move/move_backward.pass.cpp new file mode 100644 index 00000000000..70a3097fd67 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.move/move_backward.pass.cpp @@ -0,0 +1,186 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// + +// template +// requires OutputIterator::type> +// OutIter +// move_backward(InIter first, InIter last, OutIter result); + +#include +#include +#if defined(_LIBCUDACXX_HAS_MEMORY) +#include +#endif // _LIBCUDACXX_HAS_MEMORY + +#include "test_macros.h" +#include "test_iterators.h" + +struct NonTrivialMove { + int data = 0; + bool move_assigned_from = false; + + NonTrivialMove() = default; + __host__ __device__ TEST_CONSTEXPR_CXX20 + NonTrivialMove(const NonTrivialMove& other) noexcept + : data(other.data), + move_assigned_from(false) {} + + __host__ __device__ TEST_CONSTEXPR_CXX20 + NonTrivialMove(NonTrivialMove&& other) noexcept : data(other.data), + move_assigned_from(false) {} + + __host__ __device__ TEST_CONSTEXPR_CXX20 NonTrivialMove& + operator=(const NonTrivialMove& other) noexcept { + data = other.data; + move_assigned_from = false; + return *this; + } + + __host__ __device__ TEST_CONSTEXPR_CXX20 NonTrivialMove& + operator=(NonTrivialMove&& other) noexcept { + data = other.data; + move_assigned_from = true; + return *this; + } + + __host__ __device__ TEST_CONSTEXPR_CXX20 + NonTrivialMove(const int val) noexcept : data(val), + move_assigned_from(false) {} + __host__ __device__ TEST_CONSTEXPR_CXX20 NonTrivialMove& + operator=(const int val) noexcept { + data = val; + move_assigned_from = false; + return *this; + } + + __host__ __device__ TEST_CONSTEXPR_CXX20 friend bool + operator==(const NonTrivialMove& lhs, const NonTrivialMove& rhs) noexcept { + // NOTE: This uses implicit knowledge that the right hand side has been moved from + return lhs.data == rhs.data && !lhs.move_assigned_from && + rhs.move_assigned_from; + } + __host__ __device__ TEST_CONSTEXPR_CXX20 bool + operator==(const int& other) const noexcept { + return data == other; + } +}; + +template +TEST_CONSTEXPR_CXX20 __host__ __device__ void test() { + using value_type = typename cuda::std::iterator_traits::value_type; + { + constexpr int N = 1000; + value_type ia[N] = {0}; + for (int i = 0; i < N; ++i) { + ia[i] = i; + } + value_type ib[N] = {0}; + + OutIter r = + cuda::std::move_backward(InIter(ia), InIter(ia + N), OutIter(ib + N)); + assert(base(r) == ib); + for (int i = 0; i < N; ++i) { + assert(ia[i] == ib[i]); + } + } + { + constexpr int N = 5; + value_type ia[N] = {1, 2, 3, 4, 5}; + value_type ic[N + 2] = {6, 6, 6, 6, 6, 6, 6}; + + auto p = cuda::std::move_backward(ia, ia + 5, ic + N); + assert(p == ic); + for (int i = 0; i < N; ++i) { + assert(ia[i] == ic[i]); + } + + for (int i = N; i < N + 2; ++i) { + assert(ic[i] == 6); + } + } + { // Ensure that we are properly perserving back elements when the ranges are overlapping + constexpr int N = 5; + value_type ia[N + 2] = {1, 2, 3, 4, 5, 6, 7}; + int expected[N + 2] = {1, 2, 1, 2, 3, 4, 5}; + + auto p = cuda::std::move_backward(ia, ia + N, ia + N + 2); + assert(p == ia + 2); + for (int i = 0; i < N; ++i) { + assert(ia[i] == expected[i]); + } + } +} + +#if defined(_LIBCUDACXX_HAS_MEMORY) +template +TEST_CONSTEXPR_CXX20 __host__ __device__ void test1() { + const unsigned N = 100; + cuda::std::unique_ptr ia[N]; + for (unsigned i = 0; i < N; ++i) + ia[i].reset(new int(i)); + cuda::std::unique_ptr ib[N]; + + OutIter r = + cuda::std::move_backward(InIter(ia), InIter(ia + N), OutIter(ib + N)); + assert(base(r) == ib); + for (unsigned i = 0; i < N; ++i) + assert(*ib[i] == static_cast(i)); +} +#endif // _LIBCUDACXX_HAS_MEMORY + +TEST_CONSTEXPR_CXX20 __host__ __device__ bool test() { + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test >(); + test >(); + test(); + + test(); + +#if defined(_LIBCUDACXX_HAS_MEMORY) + test1*>, + bidirectional_iterator*> >(); + test1*>, + random_access_iterator*> >(); + test1*>, + cuda::std::unique_ptr*>(); + + test1*>, + bidirectional_iterator*> >(); + test1*>, + random_access_iterator*> >(); + test1*>, + cuda::std::unique_ptr*>(); + + test1*, + bidirectional_iterator*> >(); + test1*, + random_access_iterator*> >(); + test1*, cuda::std::unique_ptr*>(); +#endif // _LIBCUDACXX_HAS_MEMORY + + return true; +} + +int main(int, char**) { + test(); + +#if TEST_STD_VER >= 20 + static_assert(test()); +#endif // TEST_STD_VER >= 20 + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.remove/remove.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.remove/remove.pass.cpp new file mode 100644 index 00000000000..271b849ce72 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.remove/remove.pass.cpp @@ -0,0 +1,81 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// + +// template +// requires OutputIterator::type> +// && HasEqualTo +// constexpr Iter // constexpr after C++17 +// remove(Iter first, Iter last, const T& value); + +#include +#include + +#include "test_macros.h" +#include "test_iterators.h" + +struct MoveOnly { + int val_; + + __host__ __device__ constexpr MoveOnly(const int val) noexcept : val_(val) {} + + MoveOnly() = default; + MoveOnly(const MoveOnly&) = delete; + MoveOnly(MoveOnly&&) = default; + MoveOnly& operator=(const MoveOnly&) = delete; + MoveOnly& operator=(MoveOnly&&) = default; + + __host__ __device__ constexpr bool operator==(const int val) const noexcept { + return val_ == val; + } + + __host__ __device__ constexpr bool operator==(const MoveOnly& other) const noexcept { + return val_ == other.val_; + } +}; + +template +TEST_CONSTEXPR_CXX14 __host__ __device__ void test() { + using value_type = typename cuda::std::iterator_traits::value_type; + + constexpr int N = 9; + value_type ia[N] = {0, 1, 2, 3, 4, 2, 3, 4, 2}; + constexpr value_type expected[N - 3] = {0, 1, 3, 4, 3, 4}; + Iter r = cuda::std::remove(Iter(ia), Iter(ia + N), 2); + assert(base(r) == ia + N - 3); + for (int i = 0; i < N - 3; ++i) { + assert(ia[i] == expected[i]); + } +} + +TEST_CONSTEXPR_CXX14 __host__ __device__ bool test() { + test >(); + test >(); + test >(); + test >(); + test(); + + test >(); + test >(); + test >(); + test >(); + test(); + + return true; +} + +int main(int, char**) { + test(); + +#if TEST_STD_VER >= 14 + static_assert(test(), ""); +#endif // TEST_STD_VER >= 14 + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.remove/remove_copy.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.remove/remove_copy.pass.cpp new file mode 100644 index 00000000000..02838e61710 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.remove/remove_copy.pass.cpp @@ -0,0 +1,82 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// + +// template OutIter, class T> +// requires HasEqualTo +// constexpr OutIter // constexpr after C++17 +// remove_copy(InIter first, InIter last, OutIter result, const T& value); + +#include +#include + +#include "test_macros.h" +#include "test_iterators.h" + +template +TEST_CONSTEXPR_CXX14 __host__ __device__ void test() { + constexpr int N = 9; + int ia[N] = {0, 1, 2, 3, 4, 2, 3, 4, 2}; + constexpr int expected[N - 3] = {0, 1, 3, 4, 3, 4}; + int ib[N] = {0}; + OutIter r = + cuda::std::remove_copy(InIter(ia), InIter(ia + N), OutIter(ib), 2); + assert(base(r) == ib + N - 3); + for (int i = 0; i < N - 3; ++i) { + assert(ib[i] == expected[i]); + } + + for (int i = N - 3; i < N; ++i) { + assert(ib[i] == 0); + } +} + +TEST_CONSTEXPR_CXX14 __host__ __device__ bool test() { + test, cpp17_output_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test >(); + test >(); + test >(); + test >(); + test(); + + return true; +} + +int main(int, char**) { + test(); + +#if TEST_STD_VER >= 14 + static_assert(test(), ""); +#endif // TEST_STD_VER >= 14 + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.remove/remove_copy_if.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.remove/remove_copy_if.pass.cpp new file mode 100644 index 00000000000..286473e378c --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.remove/remove_copy_if.pass.cpp @@ -0,0 +1,86 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// + +// template OutIter, +// Predicate Pred> +// requires CopyConstructible +// constexpr OutIter // constexpr after C++17 +// remove_copy_if(InIter first, InIter last, OutIter result, Pred pred); + +#include +#include +#include + +#include "test_macros.h" +#include "test_iterators.h" + +__host__ __device__ TEST_CONSTEXPR bool equalToTwo(const int v) noexcept { return v == 2; } + +template +TEST_CONSTEXPR_CXX14 __host__ __device__ void test() { + constexpr int N = 9; + int ia[N] = {0, 1, 2, 3, 4, 2, 3, 4, 2}; + constexpr int expected[N - 3] = {0, 1, 3, 4, 3, 4}; + int ib[N] = {0}; + OutIter r = + cuda::std::remove_copy_if(InIter(ia), InIter(ia + N), OutIter(ib), equalToTwo); + assert(base(r) == ib + N - 3); + for (int i = 0; i < N - 3; ++i) { + assert(ib[i] == expected[i]); + } + + for (int i = N - 3; i < N; ++i) { + assert(ib[i] == 0); + } +} + +TEST_CONSTEXPR_CXX14 __host__ __device__ bool test() { + test, cpp17_output_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test >(); + test >(); + test >(); + test >(); + test(); + + return true; +} + +int main(int, char**) { + test(); + +#if TEST_STD_VER >= 14 + static_assert(test(), ""); +#endif // TEST_STD_VER >= 14 + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.remove/remove_if.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.remove/remove_if.pass.cpp new file mode 100644 index 00000000000..7de5c969631 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.remove/remove_if.pass.cpp @@ -0,0 +1,88 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// + +// template Pred> +// requires OutputIterator::type> +// && CopyConstructible +// constexpr Iter // constexpr after C++17 +// remove_if(Iter first, Iter last, Pred pred); + +#include +#include + +#include "test_macros.h" +#include "test_iterators.h" +#include "counting_predicates.h" + +struct MoveOnly { + int val_; + + __host__ __device__ constexpr MoveOnly(const int val) noexcept : val_(val) {} + + MoveOnly() = default; + MoveOnly(const MoveOnly&) = delete; + MoveOnly(MoveOnly&&) = default; + MoveOnly& operator=(const MoveOnly&) = delete; + MoveOnly& operator=(MoveOnly&&) = default; + + __host__ __device__ constexpr bool operator==(const int val) const noexcept { + return val_ == val; + } + + __host__ __device__ constexpr bool operator==(const MoveOnly& other) const noexcept { + return val_ == other.val_; + } +}; + +struct equal2 { + template + __host__ __device__ constexpr bool operator()(const T& i) const noexcept { + return i == 2; + } +}; + +template +TEST_CONSTEXPR_CXX14 __host__ __device__ void test() { + using value_type = typename cuda::std::iterator_traits::value_type; + constexpr int N = 9; + value_type ia[N] = {0, 1, 2, 3, 4, 2, 3, 4, 2}; + constexpr value_type expected[N - 3] = {0, 1, 3, 4, 3, 4}; + Iter r = cuda::std::remove_if(Iter(ia), Iter(ia + N), equal2{}); + assert(base(r) == ia + N - 3); + for (int i = 0; i < N - 3; ++i) { + assert(ia[i] == expected[i]); + } +} + +TEST_CONSTEXPR_CXX14 __host__ __device__ bool test() { + test >(); + test >(); + test >(); + test >(); + test(); + + test >(); + test >(); + test >(); + test >(); + test(); + + return true; +} + +int main(int, char**) { + test(); + +#if TEST_STD_VER >= 14 + static_assert(test(), ""); +#endif // TEST_STD_VER >= 14 + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.replace/replace.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.replace/replace.pass.cpp new file mode 100644 index 00000000000..a0d8a74f32c --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.replace/replace.pass.cpp @@ -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 +// +//===----------------------------------------------------------------------===// + +// + +// template +// requires OutputIterator +// && OutputIterator +// && HasEqualTo +// constexpr void // constexpr after C++17 +// replace(Iter first, Iter last, const T& old_value, const T& new_value); + +#include +#include + +#include "test_macros.h" +#include "test_iterators.h" + +template +TEST_CONSTEXPR_CXX14 __host__ __device__ void test() { + constexpr int N = 5; + int ia[N] = {0, 1, 2, 3, 4}; + constexpr int expected[N] = {0, 1, 5, 3, 4}; + cuda::std::replace(Iter(ia), Iter(ia + N), 2, 5); + + for (int i = 0; i < N; ++i) { + assert(ia[i] == expected[i]); + } +} + +TEST_CONSTEXPR_CXX14 __host__ __device__ bool test() { + test >(); + test >(); + test >(); + test >(); + test(); + + return true; +} + +int main(int, char**) +{ + test(); + +#if TEST_STD_VER >= 14 + static_assert(test(), ""); +#endif // TEST_STD_VER >= 14 + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.replace/replace_copy.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.replace/replace_copy.pass.cpp new file mode 100644 index 00000000000..b4eb79466ac --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.replace/replace_copy.pass.cpp @@ -0,0 +1,88 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// + +// template +// requires OutputIterator +// && OutputIterator +// && HasEqualTo +// constexpr OutIter // constexpr after C++17 +// replace_copy(InIter first, InIter last, OutIter result, const T& old_value, +// const T& new_value); + +#include +#include + +#include "test_macros.h" +#include "test_iterators.h" + +template +TEST_CONSTEXPR_CXX14 __host__ __device__ void test() { + { + constexpr int N = 5; + constexpr int ia[N] = {0, 1, 2, 3, 4}; + int ib[N + 1] = {0}; + constexpr int expected[N] = {0, 1, 5, 3, 4}; + OutIter r = + cuda::std::replace_copy(InIter(ia), InIter(ia + N), OutIter(ib), 2, 5); + + assert(base(r) == ib + N); + for (int i = 0; i < N; ++i) { + assert(ib[i] == expected[i]); + } + + for (int i = N; i < N + 1; ++i) { + assert(ib[i] == 0); + } + } +} + +TEST_CONSTEXPR_CXX14 __host__ __device__ bool test() { + test, cpp17_output_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test >(); + test >(); + test >(); + test >(); + test(); + + return true; +} + +int main(int, char**) { + test(); + +#if TEST_STD_VER >= 14 + static_assert(test(), ""); +#endif // TEST_STD_VER >= 14 + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.replace/replace_copy_if.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.replace/replace_copy_if.pass.cpp new file mode 100644 index 00000000000..78b56158ac0 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.replace/replace_copy_if.pass.cpp @@ -0,0 +1,89 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +// + +// template Pred, class T> +// requires OutputIterator +// && OutputIterator +// && CopyConstructible +// constexpr OutIter // constexpr after C++17 +// replace_copy_if(InIter first, InIter last, OutIter result, Pred pred, const T& new_value); + +#include +#include + +#include "test_macros.h" +#include "test_iterators.h" + +TEST_CONSTEXPR_CXX14 __host__ __device__ bool equalToTwo(const int v) noexcept { + return v == 2; +} + +template +TEST_CONSTEXPR_CXX14 __host__ __device__ void test() { + constexpr int N = 5; + constexpr int ia[N] = {0, 1, 2, 3, 4}; + int ib[N + 1] = {0}; + constexpr int expected[N] = {0, 1, 5, 3, 4}; + OutIter r = cuda::std::replace_copy_if(InIter(ia), InIter(ia + N), + OutIter(ib), equalToTwo, 5); + assert(base(r) == ib + N); + for (int i = 0; i < N; ++i) { + assert(ib[i] == expected[i]); + } + + for (int i = N; i < N + 1; ++i) { + assert(ib[i] == 0); + } +} + +TEST_CONSTEXPR_CXX14 __host__ __device__ bool test() { + test, cpp17_output_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test >(); + test >(); + test >(); + test >(); + test(); + + return true; +} + +int main(int, char**) { + test(); + +#if TEST_STD_VER >= 14 + static_assert(test(), ""); +#endif // TEST_STD_VER >= 14 + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.replace/replace_if.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.replace/replace_if.pass.cpp new file mode 100644 index 00000000000..e068dc6fd01 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.replace/replace_if.pass.cpp @@ -0,0 +1,57 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// + +// template Pred, class T> +// requires OutputIterator +// && OutputIterator +// && CopyConstructible +// constexpr void // constexpr after C++17 +// replace_if(Iter first, Iter last, Pred pred, const T& new_value); + +#include +#include + +#include "test_macros.h" +#include "test_iterators.h" + +TEST_CONSTEXPR_CXX14 __host__ __device__ bool equalToTwo(const int v) noexcept { return v == 2; } + +template +TEST_CONSTEXPR_CXX14 __host__ __device__ void test() { + constexpr int N = 5; + int ia[N] = {0, 1, 2, 3, 4}; + constexpr int expected[N] = {0, 1, 5, 3, 4}; + cuda::std::replace_if(Iter(ia), Iter(ia + N), equalToTwo, 5); + + for (int i = 0; i < N; ++i) { + assert(ia[i] == expected[i]); + } +} + +TEST_CONSTEXPR_CXX14 __host__ __device__ bool test() { + test >(); + test >(); + test >(); + test >(); + test(); + + return true; +} + +int main(int, char**) +{ + test(); + +#if TEST_STD_VER >= 14 + static_assert(test(), ""); +#endif // TEST_STD_VER >= 14 + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.transform/binary_transform.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.transform/binary_transform.pass.cpp new file mode 100644 index 00000000000..609f51b9017 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.transform/binary_transform.pass.cpp @@ -0,0 +1,360 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// + +// template BinaryOp> +// requires OutputIterator && CopyConstructible +// constexpr OutIter // constexpr after C++17 +// transform(InIter1 first1, InIter1 last1, InIter2 first2, OutIter result, BinaryOp binary_op); + +#include +#include +#include + +#include "test_macros.h" +#include "test_iterators.h" + +template +TEST_CONSTEXPR_CXX14 __host__ __device__ void test() { + { + constexpr int N = 5; + constexpr int ia[N] = {0, 1, 2, 3, 4}; + constexpr int ib[N] = {1, 2, 3, 4, 5}; + constexpr int expected[N] = {1, 3, 5, 7, 9}; + int ic[N + 1] = {0, 0, 0, 0, 0, 0}; + + OutIter r = cuda::std::transform(InIter1(ia), InIter1(ia + N), InIter2(ib), + OutIter(ic), cuda::std::plus{}); + assert(base(r) == ic + N); + for (int i = 0; i < N; ++i) { + assert(ic[i] == expected[i]); + } + + for (unsigned i = N; i < N + 1; ++i) { + assert(ic[i] == 0); + } + } +} + +TEST_CONSTEXPR_CXX14 __host__ __device__ bool test() { + test, cpp17_input_iterator, + cpp17_output_iterator >(); + test, cpp17_input_iterator, + cpp17_input_iterator >(); + test, cpp17_input_iterator, + forward_iterator >(); + test, cpp17_input_iterator, + bidirectional_iterator >(); + test, cpp17_input_iterator, + random_access_iterator >(); + test, cpp17_input_iterator, + int*>(); + + test, forward_iterator, + cpp17_output_iterator >(); + test, forward_iterator, + cpp17_input_iterator >(); + test, forward_iterator, + forward_iterator >(); + test, forward_iterator, + bidirectional_iterator >(); + test, forward_iterator, + random_access_iterator >(); + test, forward_iterator, int*>(); + + test, bidirectional_iterator, + cpp17_output_iterator >(); + test, bidirectional_iterator, + cpp17_input_iterator >(); + test, bidirectional_iterator, + forward_iterator >(); + test, bidirectional_iterator, + bidirectional_iterator >(); + test, bidirectional_iterator, + random_access_iterator >(); + test, bidirectional_iterator, + int*>(); + + test, random_access_iterator, + cpp17_output_iterator >(); + test, random_access_iterator, + cpp17_input_iterator >(); + test, random_access_iterator, + forward_iterator >(); + test, random_access_iterator, + bidirectional_iterator >(); + test, random_access_iterator, + random_access_iterator >(); + test, random_access_iterator, + int*>(); + + test, const int*, + cpp17_output_iterator >(); + test, const int*, + cpp17_input_iterator >(); + test, const int*, forward_iterator >(); + test, const int*, + bidirectional_iterator >(); + test, const int*, + random_access_iterator >(); + test, const int*, int*>(); + + test, cpp17_input_iterator, + cpp17_output_iterator >(); + test, cpp17_input_iterator, + cpp17_input_iterator >(); + test, cpp17_input_iterator, + forward_iterator >(); + test, cpp17_input_iterator, + bidirectional_iterator >(); + test, cpp17_input_iterator, + random_access_iterator >(); + test, cpp17_input_iterator, int*>(); + + test, forward_iterator, + cpp17_output_iterator >(); + test, forward_iterator, + cpp17_input_iterator >(); + test, forward_iterator, + forward_iterator >(); + test, forward_iterator, + bidirectional_iterator >(); + test, forward_iterator, + random_access_iterator >(); + test, forward_iterator, int*>(); + + test, bidirectional_iterator, + cpp17_output_iterator >(); + test, bidirectional_iterator, + cpp17_input_iterator >(); + test, bidirectional_iterator, + forward_iterator >(); + test, bidirectional_iterator, + bidirectional_iterator >(); + test, bidirectional_iterator, + random_access_iterator >(); + test, bidirectional_iterator, + int*>(); + + test, random_access_iterator, + cpp17_output_iterator >(); + test, random_access_iterator, + cpp17_input_iterator >(); + test, random_access_iterator, + forward_iterator >(); + test, random_access_iterator, + bidirectional_iterator >(); + test, random_access_iterator, + random_access_iterator >(); + test, random_access_iterator, + int*>(); + + test, const int*, + cpp17_output_iterator >(); + test, const int*, cpp17_input_iterator >(); + test, const int*, forward_iterator >(); + test, const int*, + bidirectional_iterator >(); + test, const int*, + random_access_iterator >(); + test, const int*, int*>(); + + test, cpp17_input_iterator, + cpp17_output_iterator >(); + test, cpp17_input_iterator, + cpp17_input_iterator >(); + test, cpp17_input_iterator, + forward_iterator >(); + test, cpp17_input_iterator, + bidirectional_iterator >(); + test, cpp17_input_iterator, + random_access_iterator >(); + test, cpp17_input_iterator, + int*>(); + + test, forward_iterator, + cpp17_output_iterator >(); + test, forward_iterator, + cpp17_input_iterator >(); + test, forward_iterator, + forward_iterator >(); + test, forward_iterator, + bidirectional_iterator >(); + test, forward_iterator, + random_access_iterator >(); + test, forward_iterator, + int*>(); + + test, bidirectional_iterator, + cpp17_output_iterator >(); + test, bidirectional_iterator, + cpp17_input_iterator >(); + test, bidirectional_iterator, + forward_iterator >(); + test, bidirectional_iterator, + bidirectional_iterator >(); + test, bidirectional_iterator, + random_access_iterator >(); + test, bidirectional_iterator, + int*>(); + + test, random_access_iterator, + cpp17_output_iterator >(); + test, random_access_iterator, + cpp17_input_iterator >(); + test, random_access_iterator, + forward_iterator >(); + test, random_access_iterator, + bidirectional_iterator >(); + test, random_access_iterator, + random_access_iterator >(); + test, random_access_iterator, + int*>(); + + test, const int*, + cpp17_output_iterator >(); + test, const int*, + cpp17_input_iterator >(); + test, const int*, + forward_iterator >(); + test, const int*, + bidirectional_iterator >(); + test, const int*, + random_access_iterator >(); + test, const int*, int*>(); + + test, cpp17_input_iterator, + cpp17_output_iterator >(); + test, cpp17_input_iterator, + cpp17_input_iterator >(); + test, cpp17_input_iterator, + forward_iterator >(); + test, cpp17_input_iterator, + bidirectional_iterator >(); + test, cpp17_input_iterator, + random_access_iterator >(); + test, cpp17_input_iterator, + int*>(); + + test, forward_iterator, + cpp17_output_iterator >(); + test, forward_iterator, + cpp17_input_iterator >(); + test, forward_iterator, + forward_iterator >(); + test, forward_iterator, + bidirectional_iterator >(); + test, forward_iterator, + random_access_iterator >(); + test, forward_iterator, + int*>(); + + test, bidirectional_iterator, + cpp17_output_iterator >(); + test, bidirectional_iterator, + cpp17_input_iterator >(); + test, bidirectional_iterator, + forward_iterator >(); + test, bidirectional_iterator, + bidirectional_iterator >(); + test, bidirectional_iterator, + random_access_iterator >(); + test, bidirectional_iterator, + int*>(); + + test, random_access_iterator, + cpp17_output_iterator >(); + test, random_access_iterator, + cpp17_input_iterator >(); + test, random_access_iterator, + forward_iterator >(); + test, random_access_iterator, + bidirectional_iterator >(); + test, random_access_iterator, + random_access_iterator >(); + test, random_access_iterator, + int*>(); + + test, const int*, + cpp17_output_iterator >(); + test, const int*, + cpp17_input_iterator >(); + test, const int*, + forward_iterator >(); + test, const int*, + bidirectional_iterator >(); + test, const int*, + random_access_iterator >(); + test, const int*, int*>(); + + test, + cpp17_output_iterator >(); + test, + cpp17_input_iterator >(); + test, forward_iterator >(); + test, + bidirectional_iterator >(); + test, + random_access_iterator >(); + test, int*>(); + + test, + cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, + bidirectional_iterator >(); + test, + random_access_iterator >(); + test, int*>(); + + test, + cpp17_output_iterator >(); + test, + cpp17_input_iterator >(); + test, + forward_iterator >(); + test, + bidirectional_iterator >(); + test, + random_access_iterator >(); + test, int*>(); + + test, + cpp17_output_iterator >(); + test, + cpp17_input_iterator >(); + test, + forward_iterator >(); + test, + bidirectional_iterator >(); + test, + random_access_iterator >(); + test, int*>(); + + test >(); + test >(); + test >(); + test >(); + test >(); + test(); + + return true; +} + +int main(int, char**) { + test(); + +#if TEST_STD_VER >= 14 + static_assert(test(), ""); +#endif // TEST_STD_VER >= 14 + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.transform/unary_transform.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.transform/unary_transform.pass.cpp new file mode 100644 index 00000000000..d7ec708dc57 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.transform/unary_transform.pass.cpp @@ -0,0 +1,92 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// + +// template Op> +// requires OutputIterator && CopyConstructible +// constexpr OutIter // constexpr after C++17 +// transform(InIter first, InIter last, OutIter result, Op op); + +#include +#include + +#include "test_macros.h" +#include "test_iterators.h" + +TEST_CONSTEXPR_CXX14 __host__ __device__ int plusOne(const int v) noexcept { return v + 1; } + +template +TEST_CONSTEXPR_CXX14 __host__ __device__ void test() { + { + constexpr int N = 4; + constexpr int ia[N] = {1, 3, 6, 7}; + constexpr int expected[N] = {2, 4, 7, 8}; + int ib[N + 1] = {0, 0, 0, 0, 0}; + + OutIter r = cuda::std::transform(InIter(ia), InIter(ia + N), OutIter(ib), plusOne); + assert(base(r) == ib + N); + for (int i = 0; i < N; ++i) { + assert(ib[i] == expected[i]); + } + + for (unsigned i = N; i < N + 1; ++i) { + assert(ib[i] == 0); + } + } +} + +TEST_CONSTEXPR_CXX14 __host__ __device__ bool test() { + test, cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test, cpp17_output_iterator >(); + test, cpp17_input_iterator >(); + test, forward_iterator >(); + test, bidirectional_iterator >(); + test, random_access_iterator >(); + test, int*>(); + + test >(); + test >(); + test >(); + test >(); + test >(); + test(); + + return true; +} + +int main(int, char**) { + test(); + +#if TEST_STD_VER >= 14 + static_assert(test(), ""); +#endif // TEST_STD_VER >= 14 + + return 0; +}