diff --git a/libcudacxx/include/cuda/std/__algorithm/unwrap_range.h b/libcudacxx/include/cuda/std/__algorithm/unwrap_range.h new file mode 100644 index 00000000000..04cd9abc63d --- /dev/null +++ b/libcudacxx/include/cuda/std/__algorithm/unwrap_range.h @@ -0,0 +1,120 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___ALGORITHM_UNWRAP_RANGE_H +#define _LIBCUDACXX___ALGORITHM_UNWRAP_RANGE_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include +#include +#include +#include +#include + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +// __unwrap_range and __rewrap_range are used to unwrap ranges which may have different iterator and sentinel types. +// __unwrap_iter and __rewrap_iter don't work for this, because they assume that the iterator and sentinel have +// the same type. __unwrap_range tries to get two iterators and then forward to __unwrap_iter. + +#if _CCCL_STD_VER >= 2020 +template +struct __unwrap_range_impl +{ + inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY static constexpr auto + __unwrap(_Iter __first, _Sent __sent) + requires random_access_iterator<_Iter> && sized_sentinel_for<_Sent, _Iter> + { + auto __last = ranges::next(__first, __sent); + return pair{_CUDA_VSTD::__unwrap_iter(_CUDA_VSTD::move(__first)), + _CUDA_VSTD::__unwrap_iter(_CUDA_VSTD::move(__last))}; + } + + inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY static constexpr auto + __unwrap(_Iter __first, _Sent __last) + { + return pair{_CUDA_VSTD::move(__first), _CUDA_VSTD::move(__last)}; + } + + inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY static constexpr auto + __rewrap(_Iter __orig_iter, decltype(_CUDA_VSTD::__unwrap_iter(_CUDA_VSTD::move(__orig_iter))) __iter) + requires random_access_iterator<_Iter> && sized_sentinel_for<_Sent, _Iter> + { + return _CUDA_VSTD::__rewrap_iter(_CUDA_VSTD::move(__orig_iter), _CUDA_VSTD::move(__iter)); + } + + inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY static constexpr auto + __rewrap(const _Iter&, _Iter __iter) + requires(!(random_access_iterator<_Iter> && sized_sentinel_for<_Sent, _Iter>) ) + { + return __iter; + } +}; + +template +struct __unwrap_range_impl<_Iter, _Iter> +{ + inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY static constexpr auto + __unwrap(_Iter __first, _Iter __last) + { + return pair{_CUDA_VSTD::__unwrap_iter(_CUDA_VSTD::move(__first)), + _CUDA_VSTD::__unwrap_iter(_CUDA_VSTD::move(__last))}; + } + + inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY static constexpr auto + __rewrap(_Iter __orig_iter, decltype(_CUDA_VSTD::__unwrap_iter(__orig_iter)) __iter) + { + return _CUDA_VSTD::__rewrap_iter(_CUDA_VSTD::move(__orig_iter), _CUDA_VSTD::move(__iter)); + } +}; + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr auto __unwrap_range(_Iter __first, _Sent __last) +{ + return __unwrap_range_impl<_Iter, _Sent>::__unwrap(_CUDA_VSTD::move(__first), _CUDA_VSTD::move(__last)); +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr _Iter +__rewrap_range(_Iter __orig_iter, _Unwrapped __iter) +{ + return __unwrap_range_impl<_Iter, _Sent>::__rewrap(_CUDA_VSTD::move(__orig_iter), _CUDA_VSTD::move(__iter)); +} +#else // ^^^ C++20 ^^^ / vvv C++17 vvv +template ()))> +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 pair<_Unwrapped, _Unwrapped> +__unwrap_range(_Iter __first, _Iter __last) +{ + return _CUDA_VSTD::make_pair( + _CUDA_VSTD::__unwrap_iter(_CUDA_VSTD::move(__first)), _CUDA_VSTD::__unwrap_iter(_CUDA_VSTD::move(__last))); +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 _Iter +__rewrap_range(_Iter __orig_iter, _Unwrapped __iter) +{ + return _CUDA_VSTD::__rewrap_iter(_CUDA_VSTD::move(__orig_iter), _CUDA_VSTD::move(__iter)); +} +#endif // _CCCL_STD_VER <= 2017 + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___ALGORITHM_UNWRAP_RANGE_H diff --git a/libcudacxx/include/cuda/std/__functional/function.h b/libcudacxx/include/cuda/std/__functional/function.h index ef1dfdd91a1..0995996ed5e 100644 --- a/libcudacxx/include/cuda/std/__functional/function.h +++ b/libcudacxx/include/cuda/std/__functional/function.h @@ -25,6 +25,10 @@ #include #include #include +#include +#include +#include +#include #include #include #include diff --git a/libcudacxx/include/cuda/std/__memory/allocator_arg_t.h b/libcudacxx/include/cuda/std/__memory/allocator_arg_t.h index bba91a6d621..33b03c45d15 100644 --- a/libcudacxx/include/cuda/std/__memory/allocator_arg_t.h +++ b/libcudacxx/include/cuda/std/__memory/allocator_arg_t.h @@ -42,7 +42,7 @@ _LIBCUDACXX_INLINE_VAR constexpr allocator_arg_t allocator_arg = allocator_arg_t template struct __uses_alloc_ctor_imp { - typedef _LIBCUDACXX_NODEBUG __remove_cvref_t<_Alloc> _RawAlloc; + typedef _LIBCUDACXX_NODEBUG_TYPE __remove_cvref_t<_Alloc> _RawAlloc; static const bool __ua = uses_allocator<_Tp, _RawAlloc>::value; static const bool __ic = is_constructible<_Tp, allocator_arg_t, _Alloc, _Args...>::value; static const int value = __ua ? 2 - __ic : 0; diff --git a/libcudacxx/include/cuda/std/__memory/allocator_traits.h b/libcudacxx/include/cuda/std/__memory/allocator_traits.h index a54a80bdcb4..dbd053eec0c 100644 --- a/libcudacxx/include/cuda/std/__memory/allocator_traits.h +++ b/libcudacxx/include/cuda/std/__memory/allocator_traits.h @@ -562,7 +562,7 @@ struct _LIBCUDACXX_TEMPLATE_VIS allocator_traits }; template -using __rebind_alloc _LIBCUDACXX_NODEBUG = typename _Traits::template rebind_alloc<_Tp>; +using __rebind_alloc _LIBCUDACXX_NODEBUG_TYPE = typename _Traits::template rebind_alloc<_Tp>; template struct __rebind_alloc_helper diff --git a/libcudacxx/include/cuda/std/__memory/builtin_new_allocator.h b/libcudacxx/include/cuda/std/__memory/builtin_new_allocator.h new file mode 100644 index 00000000000..a72b3a63148 --- /dev/null +++ b/libcudacxx/include/cuda/std/__memory/builtin_new_allocator.h @@ -0,0 +1,87 @@ +// -*- C++ -*- +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___MEMORY_BUILTIN_NEW_ALLOCATOR_H +#define _LIBCUDACXX___MEMORY_BUILTIN_NEW_ALLOCATOR_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +// __builtin_new_allocator -- A non-templated helper for allocating and +// deallocating memory using __builtin_operator_new and +// __builtin_operator_delete. It should be used in preference to +// `std::allocator` to avoid additional instantiations. +struct __builtin_new_allocator +{ + struct __builtin_new_deleter + { + typedef void* pointer_type; + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr explicit __builtin_new_deleter( + size_t __size, size_t __align) noexcept + : __size_(__size) + , __align_(__align) + {} + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY void operator()(void* __p) const noexcept + { + _CUDA_VSTD::__libcpp_deallocate(__p, __size_, __align_); + } + + private: + size_t __size_; + size_t __align_; + }; + + typedef unique_ptr __holder_t; + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY static __holder_t __allocate_bytes(size_t __s, size_t __align) + { + return __holder_t(_CUDA_VSTD::__libcpp_allocate(__s, __align), __builtin_new_deleter(__s, __align)); + } + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY static void + __deallocate_bytes(void* __p, size_t __s, size_t __align) noexcept + { + _CUDA_VSTD::__libcpp_deallocate(__p, __s, __align); + } + + template + _LIBCUDACXX_NODEBUG_TYPE _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY static __holder_t + __allocate_type(size_t __n) + { + return __allocate_bytes(__n * sizeof(_Tp), _LIBCUDACXX_ALIGNOF(_Tp)); + } + + template + _LIBCUDACXX_NODEBUG_TYPE _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY static void + __deallocate_type(void* __p, size_t __n) noexcept + { + __deallocate_bytes(__p, __n * sizeof(_Tp), _LIBCUDACXX_ALIGNOF(_Tp)); + } +}; + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___MEMORY_BUILTIN_NEW_ALLOCATOR_H diff --git a/libcudacxx/include/cuda/std/__memory/compressed_pair.h b/libcudacxx/include/cuda/std/__memory/compressed_pair.h new file mode 100644 index 00000000000..7c00a60f0c8 --- /dev/null +++ b/libcudacxx/include/cuda/std/__memory/compressed_pair.h @@ -0,0 +1,235 @@ +// -*- C++ -*- +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___MEMORY_COMPRESSED_PAIR_H +#define _LIBCUDACXX___MEMORY_COMPRESSED_PAIR_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +// Tag used to default initialize one or both of the pair's elements. +struct __default_init_tag +{}; +struct __value_init_tag +{}; + +template ::value> +struct __compressed_pair_elem +{ + using _ParamT = _Tp; + using reference = _Tp&; + using const_reference = const _Tp&; + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr explicit __compressed_pair_elem( + __default_init_tag) noexcept(_LIBCUDACXX_TRAIT(is_nothrow_default_constructible, _Tp)) + {} + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr explicit __compressed_pair_elem( + __value_init_tag) noexcept(_LIBCUDACXX_TRAIT(is_nothrow_default_constructible, _Tp)) + : __value_() + {} + + template ), int> = 0> + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr explicit __compressed_pair_elem(_Up&& __u) noexcept( + _LIBCUDACXX_TRAIT(is_nothrow_constructible, _Tp, _Up)) + : __value_(_CUDA_VSTD::forward<_Up>(__u)) + {} + + template + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX17 explicit __compressed_pair_elem( + piecewise_construct_t, + tuple<_Args...> __args, + __tuple_indices<_Indices...>) noexcept(_LIBCUDACXX_TRAIT(is_nothrow_constructible, _Tp, _Args...)) + : __value_(_CUDA_VSTD::forward<_Args>(_CUDA_VSTD::get<_Indices>(__args))...) + {} + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 reference __get() noexcept + { + return __value_; + } + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr const_reference __get() const noexcept + { + return __value_; + } + +private: + _Tp __value_; +}; + +template +struct __compressed_pair_elem<_Tp, _Idx, true> : private _Tp +{ + using _ParamT = _Tp; + using reference = _Tp&; + using const_reference = const _Tp&; + using __value_type = _Tp; + + _LIBCUDACXX_HIDE_FROM_ABI constexpr explicit __compressed_pair_elem() = default; + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr explicit __compressed_pair_elem( + __default_init_tag) noexcept(_LIBCUDACXX_TRAIT(is_nothrow_default_constructible, _Tp)) + {} + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr explicit __compressed_pair_elem( + __value_init_tag) noexcept(_LIBCUDACXX_TRAIT(is_nothrow_default_constructible, _Tp)) + : __value_type() + {} + + template ), int> = 0> + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr explicit __compressed_pair_elem(_Up&& __u) noexcept( + _LIBCUDACXX_TRAIT(is_nothrow_constructible, _Tp, _Up)) + : __value_type(_CUDA_VSTD::forward<_Up>(__u)) + {} + + template + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX17 __compressed_pair_elem( + piecewise_construct_t, + tuple<_Args...> __args, + __tuple_indices<_Indices...>) noexcept(_LIBCUDACXX_TRAIT(is_nothrow_constructible, _Tp, _Args...)) + : __value_type(_CUDA_VSTD::forward<_Args>(_CUDA_VSTD::get<_Indices>(__args))...) + {} + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 reference __get() noexcept + { + return *this; + } + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr const_reference __get() const noexcept + { + return *this; + } +}; + +template +class __compressed_pair + : private __compressed_pair_elem<_T1, 0> + , private __compressed_pair_elem<_T2, 1> +{ +public: + // NOTE: This static assert should never fire because __compressed_pair + // is *almost never* used in a scenario where it's possible for T1 == T2. + // (The exception is std::function where it is possible that the function + // object and the allocator have the same type). + static_assert((!_LIBCUDACXX_TRAIT(is_same, _T1, _T2)), + "__compressed_pair cannot be instantiated when T1 and T2 are the same type; " + "The current implementation is NOT ABI-compatible with the previous implementation for this " + "configuration"); + + using _Base1 _LIBCUDACXX_NODEBUG_TYPE = __compressed_pair_elem<_T1, 0>; + using _Base2 _LIBCUDACXX_NODEBUG_TYPE = __compressed_pair_elem<_T2, 1>; + + template , _Dummy>::value + && __dependent_type, _Dummy>::value>> + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr explicit __compressed_pair() noexcept( + _LIBCUDACXX_TRAIT(is_nothrow_default_constructible, _T1) + && _LIBCUDACXX_TRAIT(is_nothrow_default_constructible, _T2)) + : _Base1(__value_init_tag()) + , _Base2(__value_init_tag()) + {} + + template + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr explicit __compressed_pair( + _U1&& __t1, + _U2&& __t2) noexcept(_LIBCUDACXX_TRAIT(is_constructible, _T1, _U1) && _LIBCUDACXX_TRAIT(is_constructible, _T2, _U2)) + : _Base1(_CUDA_VSTD::forward<_U1>(__t1)) + , _Base2(_CUDA_VSTD::forward<_U2>(__t2)) + {} + + template + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX17 explicit __compressed_pair( + piecewise_construct_t __pc, + tuple<_Args1...> __first_args, + tuple<_Args2...> __second_args) noexcept(_LIBCUDACXX_TRAIT(is_constructible, _T1, _Args1...) + && _LIBCUDACXX_TRAIT(is_constructible, _T2, _Args2...)) + : _Base1(__pc, _CUDA_VSTD::move(__first_args), typename __make_tuple_indices::type()) + , _Base2(__pc, _CUDA_VSTD::move(__second_args), typename __make_tuple_indices::type()) + {} + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 typename _Base1::reference + first() noexcept + { + return static_cast<_Base1&>(*this).__get(); + } + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr typename _Base1::const_reference + first() const noexcept + { + return static_cast<_Base1 const&>(*this).__get(); + } + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 typename _Base2::reference + second() noexcept + { + return static_cast<_Base2&>(*this).__get(); + } + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr typename _Base2::const_reference + second() const noexcept + { + return static_cast<_Base2 const&>(*this).__get(); + } + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr static _Base1* + __get_first_base(__compressed_pair* __pair) noexcept + { + return static_cast<_Base1*>(__pair); + } + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr static _Base2* + __get_second_base(__compressed_pair* __pair) noexcept + { + return static_cast<_Base2*>(__pair); + } + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 void + swap(__compressed_pair& __x) noexcept(__is_nothrow_swappable<_T1>::value && __is_nothrow_swappable<_T2>::value) + { + using _CUDA_VSTD::swap; + swap(first(), __x.first()); + swap(second(), __x.second()); + } +}; + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 void swap( + __compressed_pair<_T1, _T2>& __x, + __compressed_pair<_T1, _T2>& __y) noexcept(__is_nothrow_swappable<_T1>::value && __is_nothrow_swappable<_T2>::value) +{ + __x.swap(__y); +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___MEMORY_COMPRESSED_PAIR_H diff --git a/libcudacxx/include/cuda/std/__memory/destruct_n.h b/libcudacxx/include/cuda/std/__memory/destruct_n.h new file mode 100644 index 00000000000..93a4354a068 --- /dev/null +++ b/libcudacxx/include/cuda/std/__memory/destruct_n.h @@ -0,0 +1,87 @@ +// -*- C++ -*- +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___MEMORY_DESTRUCT_N_H +#define _LIBCUDACXX___MEMORY_DESTRUCT_N_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +struct __destruct_n +{ +private: + size_t __size_; + + template + _LIBCUDACXX_INLINE_VISIBILITY void __process(_Tp* __p, false_type) noexcept + { + for (size_t __i = 0; __i < __size_; ++__i, ++__p) + { + __p->~_Tp(); + } + } + + template + _LIBCUDACXX_INLINE_VISIBILITY void __process(_Tp*, true_type) noexcept + {} + + _LIBCUDACXX_INLINE_VISIBILITY void __incr(false_type) noexcept + { + ++__size_; + } + _LIBCUDACXX_INLINE_VISIBILITY void __incr(true_type) noexcept {} + + _LIBCUDACXX_INLINE_VISIBILITY void __set(size_t __s, false_type) noexcept + { + __size_ = __s; + } + _LIBCUDACXX_INLINE_VISIBILITY void __set(size_t, true_type) noexcept {} + +public: + _LIBCUDACXX_INLINE_VISIBILITY explicit __destruct_n(size_t __s) noexcept + : __size_(__s) + {} + + template + _LIBCUDACXX_INLINE_VISIBILITY void __incr() noexcept + { + __incr(integral_constant::value>()); + } + + template + _LIBCUDACXX_INLINE_VISIBILITY void __set(size_t __s, _Tp*) noexcept + { + __set(__s, integral_constant::value>()); + } + + template + _LIBCUDACXX_INLINE_VISIBILITY void operator()(_Tp* __p) noexcept + { + __process(__p, integral_constant::value>()); + } +}; + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___MEMORY_DESTRUCT_N_H diff --git a/libcudacxx/include/cuda/std/__memory/temporary_buffer.h b/libcudacxx/include/cuda/std/__memory/temporary_buffer.h new file mode 100644 index 00000000000..869b3851fd5 --- /dev/null +++ b/libcudacxx/include/cuda/std/__memory/temporary_buffer.h @@ -0,0 +1,88 @@ +// -*- C++ -*- +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___MEMORY_TEMPORARY_BUFFER_H +#define _LIBCUDACXX___MEMORY_TEMPORARY_BUFFER_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include +#include +#include +#include +#include +#include + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +template +_CCCL_NODISCARD _LIBCUDACXX_NO_CFI _LIBCUDACXX_INLINE_VISIBILITY pair<_Tp*, ptrdiff_t> +get_temporary_buffer(ptrdiff_t __n) noexcept +{ + pair<_Tp*, ptrdiff_t> __r(0, 0); + const ptrdiff_t __m = + (~ptrdiff_t(0) ^ ptrdiff_t(ptrdiff_t(1) << (sizeof(ptrdiff_t) * __CHAR_BIT__ - 1))) / sizeof(_Tp); + if (__n > __m) + { + __n = __m; + } + while (__n > 0) + { +#if !defined(_LIBCUDACXX_HAS_NO_ALIGNED_ALLOCATION) + if (__is_overaligned_for_new(_LIBCUDACXX_ALIGNOF(_Tp))) + { + _CUDA_VSTD::align_val_t __al = _CUDA_VSTD::align_val_t(_CUDA_VSTD::alignment_of<_Tp>::value); + __r.first = static_cast<_Tp*>(::operator new(__n * sizeof(_Tp), __al)); + } + else + { + __r.first = static_cast<_Tp*>(::operator new(__n * sizeof(_Tp))); + } +#else // ^^^ !_LIBCUDACXX_HAS_NO_ALIGNED_ALLOCATION ^^^ / vvv _LIBCUDACXX_HAS_NO_ALIGNED_ALLOCATION vvv + if (__is_overaligned_for_new(_LIBCUDACXX_ALIGNOF(_Tp))) + { + // Since aligned operator new is unavailable, return an empty + // buffer rather than one with invalid alignment. + return __r; + } + + __r.first = static_cast<_Tp*>(::operator new(__n * sizeof(_Tp))); +#endif // _LIBCUDACXX_HAS_NO_ALIGNED_ALLOCATION + + if (__r.first) + { + __r.second = __n; + break; + } + __n /= 2; + } + return __r; +} + +template +inline _LIBCUDACXX_INLINE_VISIBILITY void return_temporary_buffer(_Tp* __p) noexcept +{ + _CUDA_VSTD::__libcpp_deallocate_unsized((void*) __p, _LIBCUDACXX_ALIGNOF(_Tp)); +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___MEMORY_TEMPORARY_BUFFER_H diff --git a/libcudacxx/include/cuda/std/__memory/uninitialized_algorithms.h b/libcudacxx/include/cuda/std/__memory/uninitialized_algorithms.h new file mode 100644 index 00000000000..02b04a077b9 --- /dev/null +++ b/libcudacxx/include/cuda/std/__memory/uninitialized_algorithms.h @@ -0,0 +1,686 @@ +// -*- C++ -*- +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___MEMORY_UNINITIALIZED_ALGORITHMS_H +#define _LIBCUDACXX___MEMORY_UNINITIALIZED_ALGORITHMS_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +struct __always_false +{ + template + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr bool operator()(_Args&&...) const noexcept + { + return false; + } +}; + +template +struct __simple_rollback +{ + _ForwardIterator& __first_; + _ForwardIterator& __current_; + + _LIBCUDACXX_INLINE_VISIBILITY __simple_rollback(_ForwardIterator& __first, _ForwardIterator& __current) + : __first_(__first) + , __current_(__current) + {} + + _LIBCUDACXX_INLINE_VISIBILITY void operator()() const noexcept + { + _CUDA_VSTD::__destroy(__first_, __current_); + } +}; + +// uninitialized_copy + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY pair<_InputIterator, _ForwardIterator> +__uninitialized_copy( + _InputIterator __ifirst, _Sentinel1 __ilast, _ForwardIterator __ofirst, _EndPredicate __stop_copying) +{ + _ForwardIterator __idx = __ofirst; + auto __guard = __make_exception_guard(__simple_rollback<_ForwardIterator>{__ofirst, __idx}); + for (; __ifirst != __ilast && !__stop_copying(__idx); ++__ifirst, (void) ++__idx) + { + ::new (_CUDA_VSTD::__voidify(*__idx)) _ValueType(*__ifirst); + } + __guard.__complete(); + + return pair<_InputIterator, _ForwardIterator>(_CUDA_VSTD::move(__ifirst), _CUDA_VSTD::move(__idx)); +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _ForwardIterator +uninitialized_copy(_InputIterator __ifirst, _InputIterator __ilast, _ForwardIterator __ofirst) +{ + typedef typename iterator_traits<_ForwardIterator>::value_type _ValueType; + auto __result = _CUDA_VSTD::__uninitialized_copy<_ValueType>( + _CUDA_VSTD::move(__ifirst), _CUDA_VSTD::move(__ilast), _CUDA_VSTD::move(__ofirst), __always_false{}); + return _CUDA_VSTD::move(__result.second); +} + +// uninitialized_copy_n + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY pair<_InputIterator, _ForwardIterator> +__uninitialized_copy_n(_InputIterator __ifirst, _Size __n, _ForwardIterator __ofirst, _EndPredicate __stop_copying) +{ + _ForwardIterator __idx = __ofirst; + auto __guard = __make_exception_guard(__simple_rollback<_ForwardIterator>{__ofirst, __idx}); + for (; __n > 0 && !__stop_copying(__idx); ++__ifirst, (void) ++__idx, (void) --__n) + { + ::new (_CUDA_VSTD::__voidify(*__idx)) _ValueType(*__ifirst); + } + __guard.__complete(); + + return pair<_InputIterator, _ForwardIterator>(_CUDA_VSTD::move(__ifirst), _CUDA_VSTD::move(__idx)); +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _ForwardIterator +uninitialized_copy_n(_InputIterator __ifirst, _Size __n, _ForwardIterator __ofirst) +{ + typedef typename iterator_traits<_ForwardIterator>::value_type _ValueType; + auto __result = _CUDA_VSTD::__uninitialized_copy_n<_ValueType>( + _CUDA_VSTD::move(__ifirst), __n, _CUDA_VSTD::move(__ofirst), __always_false{}); + return _CUDA_VSTD::move(__result.second); +} + +// uninitialized_fill + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _ForwardIterator +__uninitialized_fill(_ForwardIterator __first, _Sentinel __last, const _Tp& __x) +{ + _ForwardIterator __idx = __first; + auto __guard = __make_exception_guard(__simple_rollback<_ForwardIterator>{__first, __idx}); + for (; __idx != __last; ++__idx) + { + ::new (_CUDA_VSTD::__voidify(*__idx)) _ValueType(__x); + } + __guard.__complete(); + + return __idx; +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY void +uninitialized_fill(_ForwardIterator __first, _ForwardIterator __last, const _Tp& __x) +{ + typedef typename iterator_traits<_ForwardIterator>::value_type _ValueType; + (void) _CUDA_VSTD::__uninitialized_fill<_ValueType>(__first, __last, __x); +} + +// uninitialized_fill_n + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _ForwardIterator +__uninitialized_fill_n(_ForwardIterator __first, _Size __n, const _Tp& __x) +{ + _ForwardIterator __idx = __first; + auto __guard = __make_exception_guard(__simple_rollback<_ForwardIterator>{__first, __idx}); + for (; __n > 0; ++__idx, (void) --__n) + { + ::new (_CUDA_VSTD::__voidify(*__idx)) _ValueType(__x); + } + __guard.__complete(); + + return __idx; +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _ForwardIterator +uninitialized_fill_n(_ForwardIterator __first, _Size __n, const _Tp& __x) +{ + typedef typename iterator_traits<_ForwardIterator>::value_type _ValueType; + return _CUDA_VSTD::__uninitialized_fill_n<_ValueType>(__first, __n, __x); +} + +// uninitialized_default_construct + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _ForwardIterator +__uninitialized_default_construct(_ForwardIterator __first, _Sentinel __last) +{ + auto __idx = __first; + auto __guard = __make_exception_guard(__simple_rollback<_ForwardIterator>{__first, __idx}); + for (; __idx != __last; ++__idx) + { + ::new (_CUDA_VSTD::__voidify(*__idx)) _ValueType; + } + __guard.__complete(); + + return __idx; +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY void +uninitialized_default_construct(_ForwardIterator __first, _ForwardIterator __last) +{ + using _ValueType = typename iterator_traits<_ForwardIterator>::value_type; + (void) _CUDA_VSTD::__uninitialized_default_construct<_ValueType>(_CUDA_VSTD::move(__first), _CUDA_VSTD::move(__last)); +} + +// uninitialized_default_construct_n + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _ForwardIterator +__uninitialized_default_construct_n(_ForwardIterator __first, _Size __n) +{ + auto __idx = __first; + auto __guard = __make_exception_guard(__simple_rollback<_ForwardIterator>{__first, __idx}); + for (; __n > 0; ++__idx, (void) --__n) + { + ::new (_CUDA_VSTD::__voidify(*__idx)) _ValueType; + } + __guard.__complete(); + + return __idx; +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _ForwardIterator +uninitialized_default_construct_n(_ForwardIterator __first, _Size __n) +{ + using _ValueType = typename iterator_traits<_ForwardIterator>::value_type; + return _CUDA_VSTD::__uninitialized_default_construct_n<_ValueType>(_CUDA_VSTD::move(__first), __n); +} + +// uninitialized_value_construct + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _ForwardIterator +__uninitialized_value_construct(_ForwardIterator __first, _Sentinel __last) +{ + auto __idx = __first; + auto __guard = __make_exception_guard(__simple_rollback<_ForwardIterator>{__first, __idx}); + for (; __idx != __last; ++__idx) + { + ::new (_CUDA_VSTD::__voidify(*__idx)) _ValueType(); + } + __guard.__complete(); + + return __idx; +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY void +uninitialized_value_construct(_ForwardIterator __first, _ForwardIterator __last) +{ + using _ValueType = typename iterator_traits<_ForwardIterator>::value_type; + (void) _CUDA_VSTD::__uninitialized_value_construct<_ValueType>(_CUDA_VSTD::move(__first), _CUDA_VSTD::move(__last)); +} + +// uninitialized_value_construct_n + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _ForwardIterator +__uninitialized_value_construct_n(_ForwardIterator __first, _Size __n) +{ + auto __idx = __first; + auto __guard = __make_exception_guard(__simple_rollback<_ForwardIterator>{__first, __idx}); + for (; __n > 0; ++__idx, (void) --__n) + { + ::new (_CUDA_VSTD::__voidify(*__idx)) _ValueType(); + } + __guard.__complete(); + + return __idx; +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _ForwardIterator +uninitialized_value_construct_n(_ForwardIterator __first, _Size __n) +{ + using _ValueType = typename iterator_traits<_ForwardIterator>::value_type; + return _CUDA_VSTD::__uninitialized_value_construct_n<_ValueType>(_CUDA_VSTD::move(__first), __n); +} + +// uninitialized_move + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY pair<_InputIterator, _ForwardIterator> +__uninitialized_move(_InputIterator __ifirst, _Sentinel1 __ilast, _ForwardIterator __ofirst, _EndPredicate __stop_moving) +{ + auto __idx = __ofirst; + auto __guard = __make_exception_guard(__simple_rollback<_ForwardIterator>{__ofirst, __idx}); + for (; __ifirst != __ilast && !__stop_moving(__idx); ++__idx, (void) ++__ifirst) + { + ::new (_CUDA_VSTD::__voidify(*__idx)) _ValueType(_IterOps::__iter_move(__ifirst)); + } + __guard.__complete(); + + return {_CUDA_VSTD::move(__ifirst), _CUDA_VSTD::move(__idx)}; +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _ForwardIterator +uninitialized_move(_InputIterator __ifirst, _InputIterator __ilast, _ForwardIterator __ofirst) +{ + using _ValueType = typename iterator_traits<_ForwardIterator>::value_type; + auto __result = _CUDA_VSTD::__uninitialized_move<_ValueType, _IterOps<_ClassicAlgPolicy>>( + _CUDA_VSTD::move(__ifirst), _CUDA_VSTD::move(__ilast), _CUDA_VSTD::move(__ofirst), __always_false{}); + return _CUDA_VSTD::move(__result.second); +} + +// uninitialized_move_n + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY pair<_InputIterator, _ForwardIterator> +__uninitialized_move_n(_InputIterator __ifirst, _Size __n, _ForwardIterator __ofirst, _EndPredicate __stop_moving) +{ + auto __idx = __ofirst; + auto __guard = __make_exception_guard(__simple_rollback<_ForwardIterator>{__ofirst, __idx}); + for (; __n > 0 && !__stop_moving(__idx); ++__idx, (void) ++__ifirst, --__n) + { + ::new (_CUDA_VSTD::__voidify(*__idx)) _ValueType(_IterOps::__iter_move(__ifirst)); + } + __guard.__complete(); + + return {_CUDA_VSTD::move(__ifirst), _CUDA_VSTD::move(__idx)}; +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY pair<_InputIterator, _ForwardIterator> +uninitialized_move_n(_InputIterator __ifirst, _Size __n, _ForwardIterator __ofirst) +{ + using _ValueType = typename iterator_traits<_ForwardIterator>::value_type; + return _CUDA_VSTD::__uninitialized_move_n<_ValueType, _IterOps<_ClassicAlgPolicy>>( + _CUDA_VSTD::move(__ifirst), __n, _CUDA_VSTD::move(__ofirst), __always_false{}); +} + +// TODO: Rewrite this to iterate left to right and use reverse_iterators when calling +// Destroys every element in the range [first, last) FROM RIGHT TO LEFT using allocator +// destruction. If elements are themselves C-style arrays, they are recursively destroyed +// in the same manner. +// +// This function assumes that destructors do not throw, and that the allocator is bound to +// the correct type. +template ::value, int> = 0> +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 void +__allocator_destroy_multidimensional(_Alloc& __alloc, _BidirIter __first, _BidirIter __last) noexcept +{ + using _ValueType = typename iterator_traits<_BidirIter>::value_type; + static_assert(_LIBCUDACXX_TRAIT(is_same, typename allocator_traits<_Alloc>::value_type, _ValueType), + "The allocator should already be rebound to the correct type"); + + if (__first == __last) + { + return; + } + + _CCCL_IF_CONSTEXPR (_LIBCUDACXX_TRAIT(is_array, _ValueType)) + { + static_assert(!__libcpp_is_unbounded_array<_ValueType>::value, + "arrays of unbounded arrays don't exist, but if they did we would mess up here"); + + using _Element = __remove_extent_t<_ValueType>; + __allocator_traits_rebind_t<_Alloc, _Element> __elem_alloc(__alloc); + do + { + --__last; + auto&& __array = *__last; + _CUDA_VSTD::__allocator_destroy_multidimensional( + __elem_alloc, __array, __array + _LIBCUDACXX_TRAIT(extent, _ValueType)); + } while (__last != __first); + } + else + { + do + { + --__last; + allocator_traits<_Alloc>::destroy(__alloc, _CUDA_VSTD::addressof(*__last)); + } while (__last != __first); + } +} + +// Constructs the object at the given location using the allocator's construct method. +// +// If the object being constructed is an array, each element of the array is allocator-constructed, +// recursively. If an exception is thrown during the construction of an array, the initialized +// elements are destroyed in reverse order of initialization using allocator destruction. +// +// This function assumes that the allocator is bound to the correct type. +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 void +__allocator_construct_at_multidimensional(_Alloc& __alloc, _Tp* __loc) +{ + static_assert(_LIBCUDACXX_TRAIT(is_same, typename allocator_traits<_Alloc>::value_type, _Tp), + "The allocator should already be rebound to the correct type"); + + _CCCL_IF_CONSTEXPR (_LIBCUDACXX_TRAIT(is_array, _Tp)) + { + using _Element = __remove_extent_t<_Tp>; + __allocator_traits_rebind_t<_Alloc, _Element> __elem_alloc(__alloc); + size_t __i = 0; + _Tp& __array = *__loc; + + // If an exception is thrown, destroy what we have constructed so far in reverse order. + auto __guard = _CUDA_VSTD::__make_exception_guard([&]() { + _CUDA_VSTD::__allocator_destroy_multidimensional(__elem_alloc, __array, __array + __i); + }); + + for (; __i != _LIBCUDACXX_TRAIT(extent, _Tp); ++__i) + { + _CUDA_VSTD::__allocator_construct_at_multidimensional(__elem_alloc, _CUDA_VSTD::addressof(__array[__i])); + } + __guard.__complete(); + } + else + { + allocator_traits<_Alloc>::construct(__alloc, __loc); + } +} + +// Constructs the object at the given location using the allocator's construct method, passing along +// the provided argument. +// +// If the object being constructed is an array, the argument is also assumed to be an array. Each +// each element of the array being constructed is allocator-constructed from the corresponding +// element of the argument array. If an exception is thrown during the construction of an array, +// the initialized elements are destroyed in reverse order of initialization using allocator +// destruction. +// +// This function assumes that the allocator is bound to the correct type. +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 void +__allocator_construct_at_multidimensional(_Alloc& __alloc, _Tp* __loc, _Arg const& __arg) +{ + static_assert(_LIBCUDACXX_TRAIT(is_same, typename allocator_traits<_Alloc>::value_type, _Tp), + "The allocator should already be rebound to the correct type"); + + _CCCL_IF_CONSTEXPR (_LIBCUDACXX_TRAIT(is_array, _Tp)) + { + static_assert(_LIBCUDACXX_TRAIT(is_array, _Arg), + "Provided non-array initialization argument to __allocator_construct_at_multidimensional when " + "trying to construct an array."); + + using _Element = __remove_extent_t<_Tp>; + __allocator_traits_rebind_t<_Alloc, _Element> __elem_alloc(__alloc); + size_t __i = 0; + _Tp& __array = *__loc; + + // If an exception is thrown, destroy what we have constructed so far in reverse order. + auto __guard = _CUDA_VSTD::__make_exception_guard([&]() { + _CUDA_VSTD::__allocator_destroy_multidimensional(__elem_alloc, __array, __array + __i); + }); + for (; __i != _LIBCUDACXX_TRAIT(extent, _Tp); ++__i) + { + _CUDA_VSTD::__allocator_construct_at_multidimensional( + __elem_alloc, _CUDA_VSTD::addressof(__array[__i]), __arg[__i]); + } + __guard.__complete(); + } + else + { + allocator_traits<_Alloc>::construct(__alloc, __loc, __arg); + } +} + +// Given a range starting at it and containing n elements, initializes each element in the +// range from left to right using the construct method of the allocator (rebound to the +// correct type). +// +// If an exception is thrown, the initialized elements are destroyed in reverse order of +// initialization using allocator_traits destruction. If the elements in the range are C-style +// arrays, they are initialized element-wise using allocator construction, and recursively so. +template ::difference_type> +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 void +__uninitialized_allocator_fill_n_multidimensional(_Alloc& __alloc, _BidirIter __it, _Size __n, _Tp const& __value) +{ + using _ValueType = typename iterator_traits<_BidirIter>::value_type; + __allocator_traits_rebind_t<_Alloc, _ValueType> __value_alloc(__alloc); + _BidirIter __begin = __it; + + // If an exception is thrown, destroy what we have constructed so far in reverse order. + auto __guard = _CUDA_VSTD::__make_exception_guard([&]() { + _CUDA_VSTD::__allocator_destroy_multidimensional(__value_alloc, __begin, __it); + }); + for (; __n != 0; --__n, ++__it) + { + _CUDA_VSTD::__allocator_construct_at_multidimensional(__value_alloc, _CUDA_VSTD::addressof(*__it), __value); + } + __guard.__complete(); +} + +// Same as __uninitialized_allocator_fill_n_multidimensional, but doesn't pass any initialization argument +// to the allocator's construct method, which results in value initialization. +template ::difference_type> +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 void +__uninitialized_allocator_value_construct_n_multidimensional(_Alloc& __alloc, _BidirIter __it, _Size __n) +{ + using _ValueType = typename iterator_traits<_BidirIter>::value_type; + __allocator_traits_rebind_t<_Alloc, _ValueType> __value_alloc(__alloc); + _BidirIter __begin = __it; + + // If an exception is thrown, destroy what we have constructed so far in reverse order. + auto __guard = _CUDA_VSTD::__make_exception_guard([&]() { + _CUDA_VSTD::__allocator_destroy_multidimensional(__value_alloc, __begin, __it); + }); + for (; __n != 0; --__n, ++__it) + { + _CUDA_VSTD::__allocator_construct_at_multidimensional(__value_alloc, _CUDA_VSTD::addressof(*__it)); + } + __guard.__complete(); +} + +// Destroy all elements in [__first, __last) from left to right using allocator destruction. +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 void +__allocator_destroy(_Alloc& __alloc, _Iter __first, _Sent __last) +{ + for (; __first != __last; ++__first) + { + allocator_traits<_Alloc>::destroy(__alloc, _CUDA_VSTD::__to_address(__first)); + } +} + +template +class _AllocatorDestroyRangeReverse +{ +public: + inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 + _AllocatorDestroyRangeReverse(_Alloc& __alloc, _Iter& __first, _Iter& __last) + : __alloc_(__alloc) + , __first_(__first) + , __last_(__last) + {} + + inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 void operator()() const + { + _CUDA_VSTD::__allocator_destroy( + __alloc_, _CUDA_VSTD::reverse_iterator<_Iter>(__last_), _CUDA_VSTD::reverse_iterator<_Iter>(__first_)); + } + +private: + _Alloc& __alloc_; + _Iter& __first_; + _Iter& __last_; +}; + +// Copy-construct [__first1, __last1) in [__first2, __first2 + N), where N is distance(__first1, __last1). +// +// The caller has to ensure that __first2 can hold at least N uninitialized elements. If an exception is thrown the +// already copied elements are destroyed in reverse order of their construction. +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 _Iter2 +__uninitialized_allocator_copy_impl(_Alloc& __alloc, _Iter1 __first1, _Sent1 __last1, _Iter2 __first2) +{ + auto __destruct_first = __first2; + auto __guard = _CUDA_VSTD::__make_exception_guard( + _AllocatorDestroyRangeReverse<_Alloc, _Iter2>(__alloc, __destruct_first, __first2)); + while (__first1 != __last1) + { + allocator_traits<_Alloc>::construct(__alloc, _CUDA_VSTD::__to_address(__first2), *__first1); + ++__first1; + ++__first2; + } + __guard.__complete(); + return __first2; +} + +template +struct __allocator_has_trivial_copy_construct : _Not<__has_construct<_Alloc, _Type*, const _Type&>> +{}; + +template +struct __allocator_has_trivial_copy_construct, _Type> : true_type +{}; + +template , + class _Out, + __enable_if_t< + // using _RawTypeIn because of the allocator extension + _LIBCUDACXX_TRAIT(is_trivially_copy_constructible, _RawTypeIn) + && _LIBCUDACXX_TRAIT(is_trivially_copy_assignable, _RawTypeIn) + && _LIBCUDACXX_TRAIT(is_same, __remove_const_t<_In>, __remove_const_t<_Out>) + && __allocator_has_trivial_copy_construct<_Alloc, _RawTypeIn>::value>* = nullptr> +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 _Out* +__uninitialized_allocator_copy_impl(_Alloc&, _In* __first1, _In* __last1, _Out* __first2) +{ + if (__libcpp_is_constant_evaluated()) + { + while (__first1 != __last1) + { + _CUDA_VSTD::__construct_at(_CUDA_VSTD::__to_address(__first2), *__first1); + ++__first1; + ++__first2; + } + return __first2; + } + else + { + return _CUDA_VSTD::copy(__first1, __last1, __first2); + } +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 _Iter2 +__uninitialized_allocator_copy(_Alloc& __alloc, _Iter1 __first1, _Sent1 __last1, _Iter2 __first2) +{ + auto __unwrapped_range = _CUDA_VSTD::__unwrap_range(__first1, __last1); + auto __result = _CUDA_VSTD::__uninitialized_allocator_copy_impl( + __alloc, __unwrapped_range.first, __unwrapped_range.second, _CUDA_VSTD::__unwrap_iter(__first2)); + return _CUDA_VSTD::__rewrap_iter(__first2, __result); +} + +// Move-construct the elements [__first1, __last1) into [__first2, __first2 + N) +// if the move constructor is noexcept, where N is distance(__first1, __last1). +// +// Otherwise try to copy all elements. If an exception is thrown the already copied +// elements are destroyed in reverse order of their construction. +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 _Iter2 +__uninitialized_allocator_move_if_noexcept(_Alloc& __alloc, _Iter1 __first1, _Sent1 __last1, _Iter2 __first2) +{ + static_assert(__is_cpp17_move_insertable<_Alloc>::value, + "The specified type does not meet the requirements of Cpp17MoveInsertable"); + auto __destruct_first = __first2; + auto __guard = _CUDA_VSTD::__make_exception_guard( + _AllocatorDestroyRangeReverse<_Alloc, _Iter2>(__alloc, __destruct_first, __first2)); + while (__first1 != __last1) + { +#ifndef _LIBCUDACXX_NO_EXCEPTIONS + allocator_traits<_Alloc>::construct( + __alloc, _CUDA_VSTD::__to_address(__first2), _CUDA_VSTD::move_if_noexcept(*__first1)); +#else + allocator_traits<_Alloc>::construct(__alloc, _CUDA_VSTD::__to_address(__first2), _CUDA_VSTD::move(*__first1)); +#endif + ++__first1; + ++__first2; + } + __guard.__complete(); + return __first2; +} + +template +struct __allocator_has_trivial_move_construct : _Not<__has_construct<_Alloc, _Type*, _Type&&>> +{}; + +template +struct __allocator_has_trivial_move_construct, _Type> : true_type +{}; + +#ifndef _LIBCUDACXX_COMPILER_GCC +template ::value_type, + class = __enable_if_t<_LIBCUDACXX_TRAIT(is_trivially_move_constructible, _Type) + && _LIBCUDACXX_TRAIT(is_trivially_move_assignable, _Type) + && __allocator_has_trivial_move_construct<_Alloc, _Type>::value>> +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 _Iter2 +__uninitialized_allocator_move_if_noexcept(_Alloc&, _Iter1 __first1, _Iter1 __last1, _Iter2 __first2) +{ + if (__libcpp_is_constant_evaluated()) + { + while (__first1 != __last1) + { + _CUDA_VSTD::__construct_at(_CUDA_VSTD::__to_address(__first2), _CUDA_VSTD::move(*__first1)); + ++__first1; + ++__first2; + } + return __first2; + } + else + { + return _CUDA_VSTD::move(__first1, __last1, __first2); + } +} +#endif // _LIBCUDACXX_COMPILER_GCC + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___MEMORY_UNINITIALIZED_ALGORITHMS_H diff --git a/libcudacxx/include/cuda/std/__memory/unique_ptr.h b/libcudacxx/include/cuda/std/__memory/unique_ptr.h new file mode 100644 index 00000000000..c46cc02bed4 --- /dev/null +++ b/libcudacxx/include/cuda/std/__memory/unique_ptr.h @@ -0,0 +1,801 @@ +// -*- C++ -*- +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___MEMORY_UNIQUE_PTR_H +#define _LIBCUDACXX___MEMORY_UNIQUE_PTR_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#ifndef _LIBCUDACXX_HAS_NO_SPACESHIP_OPERATOR +# include +# include +# include +#endif // _LIBCUDACXX_HAS_NO_SPACESHIP_OPERATOR +#include +#include +#include // __pointer +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +template +struct _LIBCUDACXX_TEMPLATE_VIS default_delete +{ + static_assert(!_LIBCUDACXX_TRAIT(is_function, _Tp), "default_delete cannot be instantiated for function types"); + + _LIBCUDACXX_HIDE_FROM_ABI constexpr default_delete() noexcept = default; + + template + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 default_delete( + const default_delete<_Up>&, __enable_if_t<_LIBCUDACXX_TRAIT(is_convertible, _Up*, _Tp*), int> = 0) noexcept + {} + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 void + operator()(_Tp* __ptr) const noexcept + { + static_assert(sizeof(_Tp) >= 0, "cannot delete an incomplete type"); + static_assert(!is_void<_Tp>::value, "cannot delete an incomplete type"); + delete __ptr; + } +}; + +template +struct _LIBCUDACXX_TEMPLATE_VIS default_delete<_Tp[]> +{ + _LIBCUDACXX_HIDE_FROM_ABI constexpr default_delete() noexcept = default; + + template + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 + default_delete(const default_delete<_Up[]>&, + __enable_if_t<_LIBCUDACXX_TRAIT(is_convertible, _Up (*)[], _Tp (*)[]), int> = 0) noexcept + {} + + template + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 + __enable_if_t<_LIBCUDACXX_TRAIT(is_convertible, _Up (*)[], _Tp (*)[]), void> + operator()(_Up* __ptr) const noexcept + { + static_assert(sizeof(_Up) >= 0, "cannot delete an incomplete type"); + delete[] __ptr; + } +}; + +template +struct __unique_ptr_deleter_sfinae +{ + static_assert(!_LIBCUDACXX_TRAIT(is_reference, _Deleter), "incorrect specialization"); + typedef const _Deleter& __lval_ref_type; + typedef _Deleter&& __good_rval_ref_type; + typedef true_type __enable_rval_overload; +}; + +template +struct __unique_ptr_deleter_sfinae<_Deleter const&> +{ + typedef const _Deleter& __lval_ref_type; + typedef const _Deleter&& __bad_rval_ref_type; + typedef false_type __enable_rval_overload; +}; + +template +struct __unique_ptr_deleter_sfinae<_Deleter&> +{ + typedef _Deleter& __lval_ref_type; + typedef _Deleter&& __bad_rval_ref_type; + typedef false_type __enable_rval_overload; +}; + +#if defined(_LIBCUDACXX_ABI_ENABLE_UNIQUE_PTR_TRIVIAL_ABI) +# define _LIBCUDACXX_UNIQUE_PTR_TRIVIAL_ABI __attribute__((__trivial_abi__)) +#else +# define _LIBCUDACXX_UNIQUE_PTR_TRIVIAL_ABI +#endif + +template > +class _LIBCUDACXX_UNIQUE_PTR_TRIVIAL_ABI _LIBCUDACXX_TEMPLATE_VIS unique_ptr +{ +public: + typedef _Tp element_type; + typedef _Dp deleter_type; + typedef _LIBCUDACXX_NODEBUG_TYPE typename __pointer<_Tp, deleter_type>::type pointer; + + static_assert(!_LIBCUDACXX_TRAIT(is_rvalue_reference, deleter_type), + "the specified deleter type cannot be an rvalue reference"); + +private: + __compressed_pair __ptr_; + + struct __nat + { + int __for_bool_; + }; + + typedef _LIBCUDACXX_NODEBUG_TYPE __unique_ptr_deleter_sfinae<_Dp> _DeleterSFINAE; + + template + using _LValRefType _LIBCUDACXX_NODEBUG_TYPE = typename __dependent_type<_DeleterSFINAE, _Dummy>::__lval_ref_type; + + template + using _GoodRValRefType _LIBCUDACXX_NODEBUG_TYPE = + typename __dependent_type<_DeleterSFINAE, _Dummy>::__good_rval_ref_type; + + template + using _BadRValRefType _LIBCUDACXX_NODEBUG_TYPE = + typename __dependent_type<_DeleterSFINAE, _Dummy>::__bad_rval_ref_type; + + template , _Dummy>::type> + using _EnableIfDeleterDefaultConstructible _LIBCUDACXX_NODEBUG_TYPE = + typename enable_if::value && !is_pointer<_Deleter>::value>::type; + + template + using _EnableIfDeleterConstructible _LIBCUDACXX_NODEBUG_TYPE = + typename enable_if::value>::type; + + template + using _EnableIfMoveConvertible _LIBCUDACXX_NODEBUG_TYPE = + typename enable_if::value && !is_array<_Up>::value>::type; + + template + using _EnableIfDeleterConvertible _LIBCUDACXX_NODEBUG_TYPE = + typename enable_if<(is_reference<_Dp>::value && is_same<_Dp, _UDel>::value) + || (!is_reference<_Dp>::value && is_convertible<_UDel, _Dp>::value)>::type; + + template + using _EnableIfDeleterAssignable = typename enable_if::value>::type; + +public: + template > + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr unique_ptr() noexcept + : __ptr_(__value_init_tag(), __value_init_tag()) + {} + + template > + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr unique_ptr(nullptr_t) noexcept + : __ptr_(__value_init_tag(), __value_init_tag()) + {} + + template > + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 explicit unique_ptr(pointer __p) noexcept + : __ptr_(__p, __value_init_tag()) + {} + + template >> + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 + unique_ptr(pointer __p, _LValRefType<_Dummy> __d) noexcept + : __ptr_(__p, __d) + {} + + template >> + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 + unique_ptr(pointer __p, _GoodRValRefType<_Dummy> __d) noexcept + : __ptr_(__p, _CUDA_VSTD::move(__d)) + { + static_assert(!is_reference::value, "rvalue deleter bound to reference"); + } + + template >> + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY unique_ptr(pointer __p, _BadRValRefType<_Dummy> __d) = delete; + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 unique_ptr(unique_ptr&& __u) noexcept + : __ptr_(__u.release(), _CUDA_VSTD::forward(__u.get_deleter())) + {} + + template , _Up>, + class = _EnableIfDeleterConvertible<_Ep>> + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 + unique_ptr(unique_ptr<_Up, _Ep>&& __u) noexcept + : __ptr_(__u.release(), _CUDA_VSTD::forward<_Ep>(__u.get_deleter())) + {} + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 unique_ptr& + operator=(unique_ptr&& __u) noexcept + { + reset(__u.release()); + __ptr_.second() = _CUDA_VSTD::forward(__u.get_deleter()); + return *this; + } + + template , _Up>, + class = _EnableIfDeleterAssignable<_Ep>> + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 unique_ptr& + operator=(unique_ptr<_Up, _Ep>&& __u) noexcept + { + reset(__u.release()); + __ptr_.second() = _CUDA_VSTD::forward<_Ep>(__u.get_deleter()); + return *this; + } + + unique_ptr(unique_ptr const&) = delete; + unique_ptr& operator=(unique_ptr const&) = delete; + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 ~unique_ptr() + { + reset(); + } + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 unique_ptr& operator=(nullptr_t) noexcept + { + reset(); + return *this; + } + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 __add_lvalue_reference_t<_Tp> + operator*() const + { + return *__ptr_.first(); + } + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 pointer operator->() const noexcept + { + return __ptr_.first(); + } + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 pointer get() const noexcept + { + return __ptr_.first(); + } + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 deleter_type& get_deleter() noexcept + { + return __ptr_.second(); + } + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 const deleter_type& + get_deleter() const noexcept + { + return __ptr_.second(); + } + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 explicit operator bool() const noexcept + { + return __ptr_.first() != nullptr; + } + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 pointer release() noexcept + { + pointer __t = __ptr_.first(); + __ptr_.first() = pointer(); + return __t; + } + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 void + reset(pointer __p = pointer()) noexcept + { + pointer __tmp = __ptr_.first(); + __ptr_.first() = __p; + if (__tmp) + { + __ptr_.second()(__tmp); + } + } + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 void swap(unique_ptr& __u) noexcept + { + __ptr_.swap(__u.__ptr_); + } +}; + +template +class _LIBCUDACXX_UNIQUE_PTR_TRIVIAL_ABI _LIBCUDACXX_TEMPLATE_VIS unique_ptr<_Tp[], _Dp> +{ +public: + typedef _Tp element_type; + typedef _Dp deleter_type; + typedef typename __pointer<_Tp, deleter_type>::type pointer; + +private: + __compressed_pair __ptr_; + + template + struct _CheckArrayPointerConversion : is_same<_From, pointer> + {}; + + template + struct _CheckArrayPointerConversion<_FromElem*> + : integral_constant< + bool, + is_same<_FromElem*, pointer>::value + || (is_same::value && is_convertible<_FromElem (*)[], element_type (*)[]>::value)> + {}; + + typedef __unique_ptr_deleter_sfinae<_Dp> _DeleterSFINAE; + + template + using _LValRefType _LIBCUDACXX_NODEBUG_TYPE = typename __dependent_type<_DeleterSFINAE, _Dummy>::__lval_ref_type; + + template + using _GoodRValRefType _LIBCUDACXX_NODEBUG_TYPE = + typename __dependent_type<_DeleterSFINAE, _Dummy>::__good_rval_ref_type; + + template + using _BadRValRefType _LIBCUDACXX_NODEBUG_TYPE = + typename __dependent_type<_DeleterSFINAE, _Dummy>::__bad_rval_ref_type; + + template , _Dummy>::type> + using _EnableIfDeleterDefaultConstructible _LIBCUDACXX_NODEBUG_TYPE = + typename enable_if::value && !is_pointer<_Deleter>::value>::type; + + template + using _EnableIfDeleterConstructible _LIBCUDACXX_NODEBUG_TYPE = + typename enable_if::value>::type; + + template + using _EnableIfPointerConvertible _LIBCUDACXX_NODEBUG_TYPE = + typename enable_if<_CheckArrayPointerConversion<_Pp>::value>::type; + + template + using _EnableIfMoveConvertible _LIBCUDACXX_NODEBUG_TYPE = typename enable_if< + is_array<_Up>::value && is_same::value && is_same::value + && is_convertible<_ElemT (*)[], element_type (*)[]>::value>::type; + + template + using _EnableIfDeleterConvertible _LIBCUDACXX_NODEBUG_TYPE = + typename enable_if<(is_reference<_Dp>::value && is_same<_Dp, _UDel>::value) + || (!is_reference<_Dp>::value && is_convertible<_UDel, _Dp>::value)>::type; + + template + using _EnableIfDeleterAssignable _LIBCUDACXX_NODEBUG_TYPE = + typename enable_if::value>::type; + +public: + template > + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr unique_ptr() noexcept + : __ptr_(__value_init_tag(), __value_init_tag()) + {} + + template > + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr unique_ptr(nullptr_t) noexcept + : __ptr_(__value_init_tag(), __value_init_tag()) + {} + + template , + class = _EnableIfPointerConvertible<_Pp>> + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 explicit unique_ptr(_Pp __p) noexcept + : __ptr_(__p, __value_init_tag()) + {} + + template >, + class = _EnableIfPointerConvertible<_Pp>> + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 + unique_ptr(_Pp __p, _LValRefType<_Dummy> __d) noexcept + : __ptr_(__p, __d) + {} + + template >> + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 + unique_ptr(nullptr_t, _LValRefType<_Dummy> __d) noexcept + : __ptr_(nullptr, __d) + {} + + template >, + class = _EnableIfPointerConvertible<_Pp>> + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 + unique_ptr(_Pp __p, _GoodRValRefType<_Dummy> __d) noexcept + : __ptr_(__p, _CUDA_VSTD::move(__d)) + { + static_assert(!is_reference::value, "rvalue deleter bound to reference"); + } + + template >> + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 + unique_ptr(nullptr_t, _GoodRValRefType<_Dummy> __d) noexcept + : __ptr_(nullptr, _CUDA_VSTD::move(__d)) + { + static_assert(!is_reference::value, "rvalue deleter bound to reference"); + } + + template >, + class = _EnableIfPointerConvertible<_Pp>> + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY unique_ptr(_Pp __p, _BadRValRefType<_Dummy> __d) = delete; + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 unique_ptr(unique_ptr&& __u) noexcept + : __ptr_(__u.release(), _CUDA_VSTD::forward(__u.get_deleter())) + {} + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 unique_ptr& + operator=(unique_ptr&& __u) noexcept + { + reset(__u.release()); + __ptr_.second() = _CUDA_VSTD::forward(__u.get_deleter()); + return *this; + } + + template , _Up>, + class = _EnableIfDeleterConvertible<_Ep>> + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 + unique_ptr(unique_ptr<_Up, _Ep>&& __u) noexcept + : __ptr_(__u.release(), _CUDA_VSTD::forward<_Ep>(__u.get_deleter())) + {} + + template , _Up>, + class = _EnableIfDeleterAssignable<_Ep>> + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 unique_ptr& + operator=(unique_ptr<_Up, _Ep>&& __u) noexcept + { + reset(__u.release()); + __ptr_.second() = _CUDA_VSTD::forward<_Ep>(__u.get_deleter()); + return *this; + } + +public: + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 ~unique_ptr() + { + reset(); + } + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 unique_ptr& operator=(nullptr_t) noexcept + { + reset(); + return *this; + } + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 __add_lvalue_reference_t<_Tp> + operator[](size_t __i) const + { + return __ptr_.first()[__i]; + } + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 pointer get() const noexcept + { + return __ptr_.first(); + } + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 deleter_type& get_deleter() noexcept + { + return __ptr_.second(); + } + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 const deleter_type& + get_deleter() const noexcept + { + return __ptr_.second(); + } + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 explicit operator bool() const noexcept + { + return __ptr_.first() != nullptr; + } + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 pointer release() noexcept + { + pointer __t = __ptr_.first(); + __ptr_.first() = pointer(); + return __t; + } + + template ::value, int> = 0> + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 void reset(_Pp __p) noexcept + { + pointer __tmp = __ptr_.first(); + __ptr_.first() = __p; + if (__tmp) + { + __ptr_.second()(__tmp); + } + } + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 void reset(nullptr_t = nullptr) noexcept + { + pointer __tmp = __ptr_.first(); + __ptr_.first() = nullptr; + if (__tmp) + { + __ptr_.second()(__tmp); + } + } + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 void swap(unique_ptr& __u) noexcept + { + __ptr_.swap(__u.__ptr_); + } +}; + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 + __enable_if_t<__is_swappable<_Dp>::value, void> + swap(unique_ptr<_Tp, _Dp>& __x, unique_ptr<_Tp, _Dp>& __y) noexcept +{ + __x.swap(__y); +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 bool +operator==(const unique_ptr<_T1, _D1>& __x, const unique_ptr<_T2, _D2>& __y) +{ + return __x.get() == __y.get(); +} + +#if _CCCL_STD_VER <= 2017 +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY + + bool + operator!=(const unique_ptr<_T1, _D1>& __x, const unique_ptr<_T2, _D2>& __y) +{ + return !(__x == __y); +} +#endif + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY + + bool + operator<(const unique_ptr<_T1, _D1>& __x, const unique_ptr<_T2, _D2>& __y) +{ + typedef typename unique_ptr<_T1, _D1>::pointer _P1; + typedef typename unique_ptr<_T2, _D2>::pointer _P2; + typedef typename common_type<_P1, _P2>::type _Vp; + return less<_Vp>()(__x.get(), __y.get()); +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY + + bool + operator>(const unique_ptr<_T1, _D1>& __x, const unique_ptr<_T2, _D2>& __y) +{ + return __y < __x; +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY + + bool + operator<=(const unique_ptr<_T1, _D1>& __x, const unique_ptr<_T2, _D2>& __y) +{ + return !(__y < __x); +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY + + bool + operator>=(const unique_ptr<_T1, _D1>& __x, const unique_ptr<_T2, _D2>& __y) +{ + return !(__x < __y); +} + +#ifndef _LIBCUDACXX_HAS_NO_SPACESHIP_OPERATOR +# if _CCCL_STD_VER >= 2020 +template + requires three_way_comparable_with::pointer, typename unique_ptr<_T2, _D2>::pointer> +_LIBCUDACXX_HIDE_FROM_ABI + compare_three_way_result_t::pointer, typename unique_ptr<_T2, _D2>::pointer> + operator<=>(const unique_ptr<_T1, _D1>& __x, const unique_ptr<_T2, _D2>& __y) +{ + return compare_three_way()(__x.get(), __y.get()); +} +# endif // _CCCL_STD_VER >= 2020 +#endif // _LIBCUDACXX_HAS_NO_SPACESHIP_OPERATOR + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 bool +operator==(const unique_ptr<_T1, _D1>& __x, nullptr_t) noexcept +{ + return !__x; +} + +#if _CCCL_STD_VER <= 2017 +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY + + bool + operator==(nullptr_t, const unique_ptr<_T1, _D1>& __x) noexcept +{ + return !__x; +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY + + bool + operator!=(const unique_ptr<_T1, _D1>& __x, nullptr_t) noexcept +{ + return static_cast(__x); +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY + + bool + operator!=(nullptr_t, const unique_ptr<_T1, _D1>& __x) noexcept +{ + return static_cast(__x); +} +#endif // _CCCL_STD_VER <= 2017 + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 bool +operator<(const unique_ptr<_T1, _D1>& __x, nullptr_t) +{ + typedef typename unique_ptr<_T1, _D1>::pointer _P1; + return less<_P1>()(__x.get(), nullptr); +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 bool +operator<(nullptr_t, const unique_ptr<_T1, _D1>& __x) +{ + typedef typename unique_ptr<_T1, _D1>::pointer _P1; + return less<_P1>()(nullptr, __x.get()); +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 bool +operator>(const unique_ptr<_T1, _D1>& __x, nullptr_t) +{ + return nullptr < __x; +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 bool +operator>(nullptr_t, const unique_ptr<_T1, _D1>& __x) +{ + return __x < nullptr; +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 bool +operator<=(const unique_ptr<_T1, _D1>& __x, nullptr_t) +{ + return !(nullptr < __x); +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 bool +operator<=(nullptr_t, const unique_ptr<_T1, _D1>& __x) +{ + return !(__x < nullptr); +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 bool +operator>=(const unique_ptr<_T1, _D1>& __x, nullptr_t) +{ + return !(__x < nullptr); +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 bool +operator>=(nullptr_t, const unique_ptr<_T1, _D1>& __x) +{ + return !(nullptr < __x); +} + +#ifndef _LIBCUDACXX_HAS_NO_SPACESHIP_OPERATOR +# if _CCCL_STD_VER >= 2020 +template + requires three_way_comparable::pointer> +_LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 + compare_three_way_result_t::pointer> + operator<=>(const unique_ptr<_T1, _D1>& __x, nullptr_t) +{ + return compare_three_way()(__x.get(), static_cast::pointer>(nullptr)); +} +# endif // _CCCL_STD_VER >= 2020 +#endif // _LIBCUDACXX_HAS_NO_SPACESHIP_OPERATOR + +template +struct __unique_if +{ + typedef unique_ptr<_Tp> __unique_single; +}; + +template +struct __unique_if<_Tp[]> +{ + typedef unique_ptr<_Tp[]> __unique_array_unknown_bound; +}; + +template +struct __unique_if<_Tp[_Np]> +{ + typedef void __unique_array_known_bound; +}; + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 + typename __unique_if<_Tp>::__unique_single + make_unique(_Args&&... __args) +{ + return unique_ptr<_Tp>(new _Tp(_CUDA_VSTD::forward<_Args>(__args)...)); +} + +template +inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 + typename __unique_if<_Tp>::__unique_array_unknown_bound + make_unique(size_t __n) +{ + typedef __remove_extent_t<_Tp> _Up; + return unique_ptr<_Tp>(new _Up[__n]()); +} + +template +_LIBCUDACXX_INLINE_VISIBILITY typename __unique_if<_Tp>::__unique_array_known_bound make_unique(_Args&&...) = delete; + +template +_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 typename __unique_if<_Tp>::__unique_single +make_unique_for_overwrite() +{ + return unique_ptr<_Tp>(new _Tp); +} + +template +_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 typename __unique_if<_Tp>::__unique_array_unknown_bound +make_unique_for_overwrite(size_t __n) +{ + return unique_ptr<_Tp>(new __remove_extent_t<_Tp>[__n]); +} + +template +_LIBCUDACXX_INLINE_VISIBILITY typename __unique_if<_Tp>::__unique_array_known_bound +make_unique_for_overwrite(_Args&&...) = delete; + +template +struct _LIBCUDACXX_TEMPLATE_VIS hash; + +#ifndef __cuda_std__ +template +struct _LIBCUDACXX_TEMPLATE_VIS hash> +{ +# if _CCCL_STD_VER <= 2017 || defined(_LIBCUDACXX_ENABLE_CXX20_REMOVED_BINDER_TYPEDEFS) + _LIBCUDACXX_DEPRECATED_IN_CXX17 typedef unique_ptr<_Tp, _Dp> argument_type; + _LIBCUDACXX_DEPRECATED_IN_CXX17 typedef size_t result_type; +# endif + + _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY size_t operator()(const unique_ptr<_Tp, _Dp>& __ptr) const + { + typedef typename unique_ptr<_Tp, _Dp>::pointer pointer; + return hash()(__ptr.get()); + } +}; +#endif // __cuda_std__ + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___MEMORY_UNIQUE_PTR_H diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/algorithm b/libcudacxx/include/cuda/std/detail/libcxx/include/algorithm index bc48a73211d..78c168c8bfd 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/algorithm +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/algorithm @@ -1,6 +1,7 @@ //===----------------------------------------------------------------------===// // -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. @@ -644,6 +645,8 @@ template # pragma system_header #endif // no system header +#include // all public C++ headers provide the assertion handler +#include #include #include #include @@ -732,8 +735,6 @@ template #include #include #include -#include // all public C++ headers provide the assertion handler -#include #include #include #include @@ -741,6 +742,8 @@ template #include #include #include +#include +#include #include #include #include @@ -1206,12 +1209,6 @@ __second_half_done: // | } -struct __return_temporary_buffer -{ - template - _LIBCUDACXX_INLINE_VISIBILITY void operator()(_Tp* __p) const {_CUDA_VSTD::return_temporary_buffer(__p);} -}; - template _CCCL_HOST_DEVICE _ForwardIterator diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/memory b/libcudacxx/include/cuda/std/detail/libcxx/include/memory index 9a863a60a4d..c974532a025 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/memory +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/memory @@ -4,4039 +4,47 @@ // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. // //===----------------------------------------------------------------------===// #ifndef _LIBCUDACXX_MEMORY #define _LIBCUDACXX_MEMORY -/* - memory synopsis +#include -namespace 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 -struct allocator_arg_t { }; -inline constexpr allocator_arg_t allocator_arg = allocator_arg_t(); - -template struct uses_allocator; - -template -struct pointer_traits -{ - typedef Ptr pointer; - typedef
element_type; - typedef
difference_type; - - template using rebind =
; - - static pointer pointer_to(
); -}; - -template -struct pointer_traits -{ - typedef T* pointer; - typedef T element_type; - typedef ptrdiff_t difference_type; - - template using rebind = U*; - - static pointer pointer_to(
) noexcept; // constexpr in C++20 -}; - -template constexpr T* to_address(T* p) noexcept; // C++20 -template auto to_address(const Ptr& p) noexcept; // C++20 - -template -struct allocator_traits -{ - typedef Alloc allocator_type; - typedef typename allocator_type::value_type - value_type; - - typedef Alloc::pointer | value_type* pointer; - typedef Alloc::const_pointer - | pointer_traits::rebind - const_pointer; - typedef Alloc::void_pointer - | pointer_traits::rebind - void_pointer; - typedef Alloc::const_void_pointer - | pointer_traits::rebind - const_void_pointer; - typedef Alloc::difference_type - | pointer_traits::difference_type - difference_type; - typedef Alloc::size_type - | make_unsigned::type - size_type; - typedef Alloc::propagate_on_container_copy_assignment - | false_type propagate_on_container_copy_assignment; - typedef Alloc::propagate_on_container_move_assignment - | false_type propagate_on_container_move_assignment; - typedef Alloc::propagate_on_container_swap - | false_type propagate_on_container_swap; - typedef Alloc::is_always_equal - | is_empty is_always_equal; - - template using rebind_alloc = Alloc::rebind::other | Alloc; - template using rebind_traits = allocator_traits>; - - static pointer allocate(allocator_type& a, size_type n); // [[nodiscard]] in C++20 - static pointer allocate(allocator_type& a, size_type n, const_void_pointer hint); // [[nodiscard]] in C++20 - - static void deallocate(allocator_type& a, pointer p, size_type n) noexcept; - - template - static void construct(allocator_type& a, T* p, Args&&... args); - - template - static void destroy(allocator_type& a, T* p); - - static size_type max_size(const allocator_type& a); // noexcept in C++14 - - static allocator_type - select_on_container_copy_construction(const allocator_type& a); -}; - -template <> -class allocator -{ -public: - typedef void* pointer; - typedef const void* const_pointer; - typedef void value_type; - - template struct rebind {typedef allocator<_Up> other;}; -}; - -template -class allocator -{ -public: - typedef size_t size_type; - typedef ptrdiff_t difference_type; - typedef T* pointer; - typedef const T* const_pointer; - typedef typename add_lvalue_reference::type reference; - typedef typename add_lvalue_reference::type const_reference; - typedef T value_type; - - template struct rebind {typedef allocator other;}; - - constexpr allocator() noexcept; // constexpr in C++20 - constexpr allocator(const allocator&) noexcept; // constexpr in C++20 - template - constexpr allocator(const allocator&) noexcept; // constexpr in C++20 - ~allocator(); - pointer address(reference x) const noexcept; - const_pointer address(const_reference x) const noexcept; - pointer allocate(size_type, allocator::const_pointer hint = 0); - void deallocate(pointer p, size_type n) noexcept; - size_type max_size() const noexcept; - template - void construct(U* p, Args&&... args); - template - void destroy(U* p); -}; - -template -bool operator==(const allocator&, const allocator&) noexcept; - -template -bool operator!=(const allocator&, const allocator&) noexcept; - -template -class raw_storage_iterator - : public iterator // purposefully not C++03 -{ -public: - explicit raw_storage_iterator(OutputIterator x); - raw_storage_iterator& operator*(); - raw_storage_iterator& operator=(const T& element); - raw_storage_iterator& operator++(); - raw_storage_iterator operator++(int); -}; - -template pair get_temporary_buffer(ptrdiff_t n) noexcept; -template void return_temporary_buffer(T* p) noexcept; - -template T* addressof(T& r) noexcept; -template T* addressof(const T&& r) noexcept = delete; - -template -ForwardIterator -uninitialized_copy(InputIterator first, InputIterator last, ForwardIterator result); - -template -ForwardIterator -uninitialized_copy_n(InputIterator first, Size n, ForwardIterator result); - -template -void uninitialized_fill(ForwardIterator first, ForwardIterator last, const T& x); - -template -ForwardIterator -uninitialized_fill_n(ForwardIterator first, Size n, const T& x); - -template -void destroy_at(T* location); - -template - void destroy(ForwardIterator first, ForwardIterator last); - -template - ForwardIterator destroy_n(ForwardIterator first, Size n); - -template - ForwardIterator uninitialized_move(InputIterator first, InputIterator last, ForwardIterator result); - -template - pair uninitialized_move_n(InputIterator first, Size n, ForwardIterator result); - -template - void uninitialized_value_construct(ForwardIterator first, ForwardIterator last); - -template - ForwardIterator uninitialized_value_construct_n(ForwardIterator first, Size n); - -template - void uninitialized_default_construct(ForwardIterator first, ForwardIterator last); - -template - ForwardIterator uninitialized_default_construct_n(ForwardIterator first, Size n); - -template struct auto_ptr_ref {}; // deprecated in C++11, removed in C++17 - -template -class auto_ptr // deprecated in C++11, removed in C++17 -{ -public: - typedef X element_type; - - explicit auto_ptr(X* p =0) throw(); - auto_ptr(auto_ptr&) throw(); - template auto_ptr(auto_ptr&) throw(); - auto_ptr& operator=(auto_ptr&) throw(); - template auto_ptr& operator=(auto_ptr&) throw(); - auto_ptr& operator=(auto_ptr_ref r) throw(); - ~auto_ptr() throw(); - - typename add_lvalue_reference::type operator*() const throw(); - X* operator->() const throw(); - X* get() const throw(); - X* release() throw(); - void reset(X* p =0) throw(); - - auto_ptr(auto_ptr_ref) throw(); - template operator auto_ptr_ref() throw(); - template operator auto_ptr() throw(); -}; - -template -struct default_delete -{ - constexpr default_delete() noexcept = default; - template default_delete(const default_delete&) noexcept; - - void operator()(T*) const noexcept; -}; - -template -struct default_delete -{ - constexpr default_delete() noexcept = default; - void operator()(T*) const noexcept; - template void operator()(U*) const = delete; -}; - -template > -class unique_ptr -{ -public: - typedef see below pointer; - typedef T element_type; - typedef D deleter_type; - - // constructors - constexpr unique_ptr() noexcept; - explicit unique_ptr(pointer p) noexcept; - unique_ptr(pointer p, see below d1) noexcept; - unique_ptr(pointer p, see below d2) noexcept; - unique_ptr(unique_ptr&& u) noexcept; - unique_ptr(nullptr_t) noexcept : unique_ptr() { } - template - unique_ptr(unique_ptr&& u) noexcept; - template - unique_ptr(auto_ptr&& u) noexcept; // removed in C++17 - - // destructor - ~unique_ptr(); - - // assignment - unique_ptr& operator=(unique_ptr&& u) noexcept; - template unique_ptr& operator=(unique_ptr&& u) noexcept; - unique_ptr& operator=(nullptr_t) noexcept; - - // observers - typename add_lvalue_reference::type operator*() const; - pointer operator->() const noexcept; - pointer get() const noexcept; - deleter_type& get_deleter() noexcept; - const deleter_type& get_deleter() const noexcept; - explicit operator bool() const noexcept; - - // modifiers - pointer release() noexcept; - void reset(pointer p = pointer()) noexcept; - void swap(unique_ptr& u) noexcept; -}; - -template -class unique_ptr -{ -public: - typedef implementation-defined pointer; - typedef T element_type; - typedef D deleter_type; - - // constructors - constexpr unique_ptr() noexcept; - explicit unique_ptr(pointer p) noexcept; - unique_ptr(pointer p, see below d) noexcept; - unique_ptr(pointer p, see below d) noexcept; - unique_ptr(unique_ptr&& u) noexcept; - unique_ptr(nullptr_t) noexcept : unique_ptr() { } - - // destructor - ~unique_ptr(); - - // assignment - unique_ptr& operator=(unique_ptr&& u) noexcept; - unique_ptr& operator=(nullptr_t) noexcept; - - // observers - T& operator[](size_t i) const; - pointer get() const noexcept; - deleter_type& get_deleter() noexcept; - const deleter_type& get_deleter() const noexcept; - explicit operator bool() const noexcept; - - // modifiers - pointer release() noexcept; - void reset(pointer p = pointer()) noexcept; - void reset(nullptr_t) noexcept; - template void reset(U) = delete; - void swap(unique_ptr& u) noexcept; -}; - -template - void swap(unique_ptr& x, unique_ptr& y) noexcept; - -template - bool operator==(const unique_ptr& x, const unique_ptr& y); -template - bool operator!=(const unique_ptr& x, const unique_ptr& y); -template - bool operator<(const unique_ptr& x, const unique_ptr& y); -template - bool operator<=(const unique_ptr& x, const unique_ptr& y); -template - bool operator>(const unique_ptr& x, const unique_ptr& y); -template - bool operator>=(const unique_ptr& x, const unique_ptr& y); - -template - bool operator==(const unique_ptr& x, nullptr_t) noexcept; -template - bool operator==(nullptr_t, const unique_ptr& y) noexcept; -template - bool operator!=(const unique_ptr& x, nullptr_t) noexcept; -template - bool operator!=(nullptr_t, const unique_ptr& y) noexcept; - -template - bool operator<(const unique_ptr& x, nullptr_t); -template - bool operator<(nullptr_t, const unique_ptr& y); -template - bool operator<=(const unique_ptr& x, nullptr_t); -template - bool operator<=(nullptr_t, const unique_ptr& y); -template - bool operator>(const unique_ptr& x, nullptr_t); -template - bool operator>(nullptr_t, const unique_ptr& y); -template - bool operator>=(const unique_ptr& x, nullptr_t); -template - bool operator>=(nullptr_t, const unique_ptr& y); - -class bad_weak_ptr - : public std::exception -{ - bad_weak_ptr() noexcept; -}; - -template unique_ptr make_unique(Args&&... args); // C++14 -template unique_ptr make_unique(size_t n); // C++14 -template unspecified make_unique(Args&&...) = delete; // C++14, T == U[N] - -template - basic_ostream& operator<< (basic_ostream& os, unique_ptr const& p); - -template -class shared_ptr -{ -public: - typedef T element_type; - typedef weak_ptr weak_type; // C++17 - - // constructors: - constexpr shared_ptr() noexcept; - template explicit shared_ptr(Y* p); - template shared_ptr(Y* p, D d); - template shared_ptr(Y* p, D d, A a); - template shared_ptr(nullptr_t p, D d); - template shared_ptr(nullptr_t p, D d, A a); - template shared_ptr(const shared_ptr& r, T *p) noexcept; - shared_ptr(const shared_ptr& r) noexcept; - template shared_ptr(const shared_ptr& r) noexcept; - shared_ptr(shared_ptr&& r) noexcept; - template shared_ptr(shared_ptr&& r) noexcept; - template explicit shared_ptr(const weak_ptr& r); - template shared_ptr(auto_ptr&& r); // removed in C++17 - template shared_ptr(unique_ptr&& r); - shared_ptr(nullptr_t) : shared_ptr() { } - - // destructor: - ~shared_ptr(); - - // assignment: - shared_ptr& operator=(const shared_ptr& r) noexcept; - template shared_ptr& operator=(const shared_ptr& r) noexcept; - shared_ptr& operator=(shared_ptr&& r) noexcept; - template shared_ptr& operator=(shared_ptr&& r); - template shared_ptr& operator=(auto_ptr&& r); // removed in C++17 - template shared_ptr& operator=(unique_ptr&& r); - - // modifiers: - void swap(shared_ptr& r) noexcept; - void reset() noexcept; - template void reset(Y* p); - template void reset(Y* p, D d); - template void reset(Y* p, D d, A a); - - // observers: - T* get() const noexcept; - T& operator*() const noexcept; - T* operator->() const noexcept; - long use_count() const noexcept; - bool unique() const noexcept; - explicit operator bool() const noexcept; - template bool owner_before(shared_ptr const& b) const noexcept; - template bool owner_before(weak_ptr const& b) const noexcept; -}; - -// shared_ptr comparisons: -template - bool operator==(shared_ptr const& a, shared_ptr const& b) noexcept; -template - bool operator!=(shared_ptr const& a, shared_ptr const& b) noexcept; -template - bool operator<(shared_ptr const& a, shared_ptr const& b) noexcept; -template - bool operator>(shared_ptr const& a, shared_ptr const& b) noexcept; -template - bool operator<=(shared_ptr const& a, shared_ptr const& b) noexcept; -template - bool operator>=(shared_ptr const& a, shared_ptr const& b) noexcept; - -template - bool operator==(const shared_ptr& x, nullptr_t) noexcept; -template - bool operator==(nullptr_t, const shared_ptr& y) noexcept; -template - bool operator!=(const shared_ptr& x, nullptr_t) noexcept; -template - bool operator!=(nullptr_t, const shared_ptr& y) noexcept; -template - bool operator<(const shared_ptr& x, nullptr_t) noexcept; -template -bool operator<(nullptr_t, const shared_ptr& y) noexcept; -template - bool operator<=(const shared_ptr& x, nullptr_t) noexcept; -template - bool operator<=(nullptr_t, const shared_ptr& y) noexcept; -template - bool operator>(const shared_ptr& x, nullptr_t) noexcept; -template - bool operator>(nullptr_t, const shared_ptr& y) noexcept; -template - bool operator>=(const shared_ptr& x, nullptr_t) noexcept; -template - bool operator>=(nullptr_t, const shared_ptr& y) noexcept; - -// shared_ptr specialized algorithms: -template void swap(shared_ptr& a, shared_ptr& b) noexcept; - -// shared_ptr casts: -template - shared_ptr static_pointer_cast(shared_ptr const& r) noexcept; -template - shared_ptr dynamic_pointer_cast(shared_ptr const& r) noexcept; -template - shared_ptr const_pointer_cast(shared_ptr const& r) noexcept; - -// shared_ptr I/O: -template - basic_ostream& operator<< (basic_ostream& os, shared_ptr const& p); - -// shared_ptr get_deleter: -template D* get_deleter(shared_ptr const& p) noexcept; - -template - shared_ptr make_shared(Args&&... args); -template - shared_ptr allocate_shared(const A& a, Args&&... args); - -template -class weak_ptr -{ -public: - typedef T element_type; - - // constructors - constexpr weak_ptr() noexcept; - template weak_ptr(shared_ptr const& r) noexcept; - weak_ptr(weak_ptr const& r) noexcept; - template weak_ptr(weak_ptr const& r) noexcept; - weak_ptr(weak_ptr&& r) noexcept; // C++14 - template weak_ptr(weak_ptr&& r) noexcept; // C++14 - - // destructor - ~weak_ptr(); - - // assignment - weak_ptr& operator=(weak_ptr const& r) noexcept; - template weak_ptr& operator=(weak_ptr const& r) noexcept; - template weak_ptr& operator=(shared_ptr const& r) noexcept; - weak_ptr& operator=(weak_ptr&& r) noexcept; // C++14 - template weak_ptr& operator=(weak_ptr&& r) noexcept; // C++14 - - // modifiers - void swap(weak_ptr& r) noexcept; - void reset() noexcept; - - // observers - long use_count() const noexcept; - bool expired() const noexcept; - shared_ptr lock() const noexcept; - template bool owner_before(shared_ptr const& b) const noexcept; - template bool owner_before(weak_ptr const& b) const noexcept; -}; - -// weak_ptr specialized algorithms: -template void swap(weak_ptr& a, weak_ptr& b) noexcept; - -// class owner_less: -template struct owner_less; - -template -struct owner_less > - : __binary_function, shared_ptr, bool> -{ - typedef bool result_type; - bool operator()(shared_ptr const&, shared_ptr const&) const noexcept; - bool operator()(shared_ptr const&, weak_ptr const&) const noexcept; - bool operator()(weak_ptr const&, shared_ptr const&) const noexcept; -}; - -template -struct owner_less > - : __binary_function, weak_ptr, bool> -{ - typedef bool result_type; - bool operator()(weak_ptr const&, weak_ptr const&) const noexcept; - bool operator()(shared_ptr const&, weak_ptr const&) const noexcept; - bool operator()(weak_ptr const&, shared_ptr const&) const noexcept; -}; - -template <> // Added in C++14 -struct owner_less -{ - template - bool operator()( shared_ptr<_Tp> const& __x, shared_ptr<_Up> const& __y) const noexcept; - template - bool operator()( shared_ptr<_Tp> const& __x, weak_ptr<_Up> const& __y) const noexcept; - template - bool operator()( weak_ptr<_Tp> const& __x, shared_ptr<_Up> const& __y) const noexcept; - template - bool operator()( weak_ptr<_Tp> const& __x, weak_ptr<_Up> const& __y) const noexcept; - - typedef void is_transparent; -}; - -template -class enable_shared_from_this -{ -protected: - constexpr enable_shared_from_this() noexcept; - enable_shared_from_this(enable_shared_from_this const&) noexcept; - enable_shared_from_this& operator=(enable_shared_from_this const&) noexcept; - ~enable_shared_from_this(); -public: - shared_ptr shared_from_this(); - shared_ptr shared_from_this() const; -}; - -template - bool atomic_is_lock_free(const shared_ptr* p); -template - shared_ptr atomic_load(const shared_ptr* p); -template - shared_ptr atomic_load_explicit(const shared_ptr* p, memory_order mo); -template - void atomic_store(shared_ptr* p, shared_ptr r); -template - void atomic_store_explicit(shared_ptr* p, shared_ptr r, memory_order mo); -template - shared_ptr atomic_exchange(shared_ptr* p, shared_ptr r); -template - shared_ptr - atomic_exchange_explicit(shared_ptr* p, shared_ptr r, memory_order mo); -template - bool - atomic_compare_exchange_weak(shared_ptr* p, shared_ptr* v, shared_ptr w); -template - bool - atomic_compare_exchange_strong( shared_ptr* p, shared_ptr* v, shared_ptr w); -template - bool - atomic_compare_exchange_weak_explicit(shared_ptr* p, shared_ptr* v, - shared_ptr w, memory_order success, - memory_order failure); -template - bool - atomic_compare_exchange_strong_explicit(shared_ptr* p, shared_ptr* v, - shared_ptr w, memory_order success, - memory_order failure); -// Hash support -template struct hash; -template struct hash >; -template struct hash >; - -template - inline constexpr bool uses_allocator_v = uses_allocator::value; - -// Pointer safety -enum class pointer_safety { relaxed, preferred, strict }; -void declare_reachable(void *p); -template T *undeclare_reachable(T *p); -void declare_no_pointers(char *p, size_t n); -void undeclare_no_pointers(char *p, size_t n); -pointer_safety get_pointer_safety() noexcept; - -void* align(size_t alignment, size_t size, void*& ptr, size_t& space); - -} // std - -*/ -#ifndef __cuda_std__ -#include <__config> -#include -#include -#if !defined(_LIBCUDACXX_HAS_NO_ATOMIC_HEADER) -# include -#endif -#endif //__cuda_std__ - -#include // all public C++ headers provide the assertion handler -#include -#include -#include #include #include #include #include +#include #include -#include #include -#include #include #include +#include +#include #include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include +#include // all public C++ headers provide the assertion handler // standard-mandated includes #include -#include - -#ifndef __cuda_std__ - -_LIBCUDACXX_BEGIN_NAMESPACE_STD - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_ValueType __libcpp_relaxed_load(_ValueType const* __value) { -#if !defined(_LIBCUDACXX_HAS_NO_THREADS) && \ - defined(__ATOMIC_RELAXED) && \ - (__has_builtin(__atomic_load_n) || defined(_CCCL_COMPILER_GCC)) - return __atomic_load_n(__value, __ATOMIC_RELAXED); -#else - return *__value; -#endif -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_ValueType __libcpp_acquire_load(_ValueType const* __value) { -#if !defined(_LIBCUDACXX_HAS_NO_THREADS) && \ - defined(__ATOMIC_ACQUIRE) && \ - (__has_builtin(__atomic_load_n) || defined(_CCCL_COMPILER_GCC)) - return __atomic_load_n(__value, __ATOMIC_ACQUIRE); -#else - return *__value; -#endif -} - -template -class _LIBCUDACXX_TEMPLATE_VIS raw_storage_iterator - : public iterator&> // purposefully not C++03 -{ -private: - _OutputIterator __x_; -public: - _LIBCUDACXX_INLINE_VISIBILITY explicit raw_storage_iterator(_OutputIterator __x) : __x_(__x) {} - _LIBCUDACXX_INLINE_VISIBILITY raw_storage_iterator& operator*() {return *this;} - _LIBCUDACXX_INLINE_VISIBILITY raw_storage_iterator& operator=(const _Tp& __element) - {::new(_CUDA_VSTD::addressof(*__x_)) _Tp(__element); return *this;} -#if _CCCL_STD_VER >= 2014 - _LIBCUDACXX_INLINE_VISIBILITY raw_storage_iterator& operator=(_Tp&& __element) - {::new(_CUDA_VSTD::addressof(*__x_)) _Tp(_CUDA_VSTD::move(__element)); return *this;} -#endif - _LIBCUDACXX_INLINE_VISIBILITY raw_storage_iterator& operator++() {++__x_; return *this;} - _LIBCUDACXX_INLINE_VISIBILITY raw_storage_iterator operator++(int) - {raw_storage_iterator __t(*this); ++__x_; return __t;} -#if _CCCL_STD_VER >= 2014 - _LIBCUDACXX_INLINE_VISIBILITY _OutputIterator base() const { return __x_; } -#endif -}; - -template -_CCCL_NODISCARD _LIBCUDACXX_NO_CFI -pair<_Tp*, ptrdiff_t> -get_temporary_buffer(ptrdiff_t __n) noexcept -{ - pair<_Tp*, ptrdiff_t> __r(0, 0); - const ptrdiff_t __m = (~ptrdiff_t(0) ^ - ptrdiff_t(ptrdiff_t(1) << (sizeof(ptrdiff_t) * __CHAR_BIT__ - 1))) - / sizeof(_Tp); - if (__n > __m) - __n = __m; - while (__n > 0) - { -#if !defined(_LIBCUDACXX_HAS_NO_ALIGNED_ALLOCATION) - if (__is_overaligned_for_new(_LIBCUDACXX_ALIGNOF(_Tp))) - { - std::align_val_t __al = - std::align_val_t(std::alignment_of<_Tp>::value); - __r.first = static_cast<_Tp*>(::operator new( - __n * sizeof(_Tp), __al, nothrow)); - } else { - __r.first = static_cast<_Tp*>(::operator new( - __n * sizeof(_Tp), nothrow)); - } -#else - if (__is_overaligned_for_new(_LIBCUDACXX_ALIGNOF(_Tp))) - { - // Since aligned operator new is unavailable, return an empty - // buffer rather than one with invalid alignment. - return __r; - } - - __r.first = static_cast<_Tp*>(::operator new(__n * sizeof(_Tp), nothrow)); -#endif - - if (__r.first) - { - __r.second = __n; - break; - } - __n /= 2; - } - return __r; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -void return_temporary_buffer(_Tp* __p) noexcept -{ - _CUDA_VSTD::__libcpp_deallocate_unsized((void*)__p, _LIBCUDACXX_ALIGNOF(_Tp)); -} - -#if _CCCL_STD_VER <= 2014 || defined(_LIBCUDACXX_ENABLE_CXX17_REMOVED_AUTO_PTR) -template -struct _LIBCUDACXX_DEPRECATED_IN_CXX11 auto_ptr_ref -{ - _Tp* __ptr_; -}; - -template -class _LIBCUDACXX_TEMPLATE_VIS _LIBCUDACXX_DEPRECATED_IN_CXX11 auto_ptr -{ -private: - _Tp* __ptr_; -public: - typedef _Tp element_type; - - _LIBCUDACXX_INLINE_VISIBILITY explicit auto_ptr(_Tp* __p = 0) throw() : __ptr_(__p) {} - _LIBCUDACXX_INLINE_VISIBILITY auto_ptr(auto_ptr& __p) throw() : __ptr_(__p.release()) {} - template _LIBCUDACXX_INLINE_VISIBILITY auto_ptr(auto_ptr<_Up>& __p) throw() - : __ptr_(__p.release()) {} - _LIBCUDACXX_INLINE_VISIBILITY auto_ptr& operator=(auto_ptr& __p) throw() - {reset(__p.release()); return *this;} - template _LIBCUDACXX_INLINE_VISIBILITY auto_ptr& operator=(auto_ptr<_Up>& __p) throw() - {reset(__p.release()); return *this;} - _LIBCUDACXX_INLINE_VISIBILITY auto_ptr& operator=(auto_ptr_ref<_Tp> __p) throw() - {reset(__p.__ptr_); return *this;} - _LIBCUDACXX_INLINE_VISIBILITY ~auto_ptr() throw() {delete __ptr_;} - - _LIBCUDACXX_INLINE_VISIBILITY _Tp& operator*() const throw() - {return *__ptr_;} - _LIBCUDACXX_INLINE_VISIBILITY _Tp* operator->() const throw() {return __ptr_;} - _LIBCUDACXX_INLINE_VISIBILITY _Tp* get() const throw() {return __ptr_;} - _LIBCUDACXX_INLINE_VISIBILITY _Tp* release() throw() - { - _Tp* __t = __ptr_; - __ptr_ = 0; - return __t; - } - _LIBCUDACXX_INLINE_VISIBILITY void reset(_Tp* __p = 0) throw() - { - if (__ptr_ != __p) - delete __ptr_; - __ptr_ = __p; - } - - _LIBCUDACXX_INLINE_VISIBILITY auto_ptr(auto_ptr_ref<_Tp> __p) throw() : __ptr_(__p.__ptr_) {} - template _LIBCUDACXX_INLINE_VISIBILITY operator auto_ptr_ref<_Up>() throw() - {auto_ptr_ref<_Up> __t; __t.__ptr_ = release(); return __t;} - template _LIBCUDACXX_INLINE_VISIBILITY operator auto_ptr<_Up>() throw() - {return auto_ptr<_Up>(release());} -}; - -template <> -class _LIBCUDACXX_TEMPLATE_VIS _LIBCUDACXX_DEPRECATED_IN_CXX11 auto_ptr -{ -public: - typedef void element_type; -}; -#endif - -template ::value && !__libcpp_is_final<_Tp>::value> -struct __compressed_pair_elem { - typedef _Tp _ParamT; - typedef _Tp& reference; - typedef const _Tp& const_reference; - - _LIBCUDACXX_INLINE_VISIBILITY constexpr __compressed_pair_elem() : __value_() {} - - template >::value - >::type> - _LIBCUDACXX_INLINE_VISIBILITY - constexpr explicit - __compressed_pair_elem(_Up&& __u) - : __value_(_CUDA_VSTD::forward<_Up>(__u)) - { - } - - template - _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX17 - __compressed_pair_elem(piecewise_construct_t, tuple<_Args...> __args, - __tuple_indices<_Indexes...>) - : __value_(_CUDA_VSTD::forward<_Args>(_CUDA_VSTD::get<_Indexes>(__args))...) {} - - _LIBCUDACXX_INLINE_VISIBILITY reference __get() noexcept { return __value_; } - _LIBCUDACXX_INLINE_VISIBILITY - const_reference __get() const noexcept { return __value_; } - -private: - _Tp __value_; -}; - -template -struct __compressed_pair_elem<_Tp, _Idx, true> : private _Tp { - typedef _Tp _ParamT; - typedef _Tp& reference; - typedef const _Tp& const_reference; - typedef _Tp __value_type; - - _LIBCUDACXX_INLINE_VISIBILITY constexpr __compressed_pair_elem() = default; - - template >::value - >::type> - _LIBCUDACXX_INLINE_VISIBILITY - constexpr explicit - __compressed_pair_elem(_Up&& __u) - : __value_type(_CUDA_VSTD::forward<_Up>(__u)) - {} - - template - _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX17 - __compressed_pair_elem(piecewise_construct_t, tuple<_Args...> __args, - __tuple_indices<_Indexes...>) - : __value_type(_CUDA_VSTD::forward<_Args>(_CUDA_VSTD::get<_Indexes>(__args))...) {} - - _LIBCUDACXX_INLINE_VISIBILITY reference __get() noexcept { return *this; } - _LIBCUDACXX_INLINE_VISIBILITY - const_reference __get() const noexcept { return *this; } -}; - -// Tag used to construct the second element of the compressed pair. -struct __second_tag {}; - -template -class __compressed_pair : private __compressed_pair_elem<_T1, 0>, - private __compressed_pair_elem<_T2, 1> { - typedef _LIBCUDACXX_NODEBUG_TYPE __compressed_pair_elem<_T1, 0> _Base1; - typedef _LIBCUDACXX_NODEBUG_TYPE __compressed_pair_elem<_T2, 1> _Base2; - - // NOTE: This static assert should never fire because __compressed_pair - // is *almost never* used in a scenario where it's possible for T1 == T2. - // (The exception is std::function where it is possible that the function - // object and the allocator have the same type). - static_assert((!is_same<_T1, _T2>::value), - "__compressed_pair cannot be instantated when T1 and T2 are the same type; " - "The current implementation is NOT ABI-compatible with the previous " - "implementation for this configuration"); - -public: - template , _Dummy>::value && - __dependent_type, _Dummy>::value - >::type - > - _LIBCUDACXX_INLINE_VISIBILITY - constexpr __compressed_pair() {} - - template , - __compressed_pair>::value, - bool>::type = true> - _LIBCUDACXX_INLINE_VISIBILITY constexpr explicit - __compressed_pair(_Tp&& __t) - : _Base1(std::forward<_Tp>(__t)), _Base2() {} - - template - _LIBCUDACXX_INLINE_VISIBILITY constexpr - __compressed_pair(__second_tag, _Tp&& __t) - : _Base1(), _Base2(std::forward<_Tp>(__t)) {} - - template - _LIBCUDACXX_INLINE_VISIBILITY constexpr - __compressed_pair(_U1&& __t1, _U2&& __t2) - : _Base1(std::forward<_U1>(__t1)), _Base2(std::forward<_U2>(__t2)) {} - - template - _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX17 - __compressed_pair(piecewise_construct_t __pc, tuple<_Args1...> __first_args, - tuple<_Args2...> __second_args) - : _Base1(__pc, _CUDA_VSTD::move(__first_args), - typename __make_tuple_indices::type()), - _Base2(__pc, _CUDA_VSTD::move(__second_args), - typename __make_tuple_indices::type()) {} - - _LIBCUDACXX_INLINE_VISIBILITY - typename _Base1::reference first() noexcept { - return static_cast<_Base1&>(*this).__get(); - } - - _LIBCUDACXX_INLINE_VISIBILITY - typename _Base1::const_reference first() const noexcept { - return static_cast<_Base1 const&>(*this).__get(); - } - - _LIBCUDACXX_INLINE_VISIBILITY - typename _Base2::reference second() noexcept { - return static_cast<_Base2&>(*this).__get(); - } - - _LIBCUDACXX_INLINE_VISIBILITY - typename _Base2::const_reference second() const noexcept { - return static_cast<_Base2 const&>(*this).__get(); - } - - _LIBCUDACXX_INLINE_VISIBILITY - void swap(__compressed_pair& __x) - noexcept(__is_nothrow_swappable<_T1>::value && - __is_nothrow_swappable<_T2>::value) - { - using std::swap; - swap(first(), __x.first()); - swap(second(), __x.second()); - } -}; - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -void swap(__compressed_pair<_T1, _T2>& __x, __compressed_pair<_T1, _T2>& __y) - noexcept(__is_nothrow_swappable<_T1>::value && - __is_nothrow_swappable<_T2>::value) { - __x.swap(__y); -} - -// default_delete - -template -struct _LIBCUDACXX_TEMPLATE_VIS default_delete { - static_assert(!is_function<_Tp>::value, - "default_delete cannot be instantiated for function types"); - _LIBCUDACXX_INLINE_VISIBILITY constexpr default_delete() noexcept = default; - - template - _LIBCUDACXX_INLINE_VISIBILITY - default_delete(const default_delete<_Up>&, - typename enable_if::value>::type* = - 0) noexcept {} - - _LIBCUDACXX_INLINE_VISIBILITY void operator()(_Tp* __ptr) const noexcept { - static_assert(sizeof(_Tp) > 0, - "default_delete can not delete incomplete type"); - static_assert(!is_void<_Tp>::value, - "default_delete can not delete incomplete type"); - delete __ptr; - } -}; - -template -struct _LIBCUDACXX_TEMPLATE_VIS default_delete<_Tp[]> { -private: - template - struct _EnableIfConvertible - : enable_if::value> {}; - -public: - _LIBCUDACXX_INLINE_VISIBILITY constexpr default_delete() noexcept = default; - - template - _LIBCUDACXX_INLINE_VISIBILITY - default_delete(const default_delete<_Up[]>&, - typename _EnableIfConvertible<_Up>::type* = 0) noexcept {} - - template - _LIBCUDACXX_INLINE_VISIBILITY - typename _EnableIfConvertible<_Up>::type - operator()(_Up* __ptr) const noexcept { - static_assert(sizeof(_Tp) > 0, - "default_delete can not delete incomplete type"); - static_assert(!is_void<_Tp>::value, - "default_delete can not delete void type"); - delete[] __ptr; - } -}; - -template -struct __unique_ptr_deleter_sfinae { - static_assert(!is_reference<_Deleter>::value, "incorrect specialization"); - typedef const _Deleter& __lval_ref_type; - typedef _Deleter&& __good_rval_ref_type; - typedef true_type __enable_rval_overload; -}; - -template -struct __unique_ptr_deleter_sfinae<_Deleter const&> { - typedef const _Deleter& __lval_ref_type; - typedef const _Deleter&& __bad_rval_ref_type; - typedef false_type __enable_rval_overload; -}; - -template -struct __unique_ptr_deleter_sfinae<_Deleter&> { - typedef _Deleter& __lval_ref_type; - typedef _Deleter&& __bad_rval_ref_type; - typedef false_type __enable_rval_overload; -}; - -template > -class _LIBCUDACXX_TEMPLATE_VIS unique_ptr { -public: - typedef _Tp element_type; - typedef _Dp deleter_type; - typedef _LIBCUDACXX_NODEBUG_TYPE typename __pointer<_Tp, deleter_type>::type pointer; - - static_assert(!is_rvalue_reference::value, - "the specified deleter type cannot be an rvalue reference"); - -private: - __compressed_pair __ptr_; - - struct __nat { int __for_bool_; }; - - typedef _LIBCUDACXX_NODEBUG_TYPE __unique_ptr_deleter_sfinae<_Dp> _DeleterSFINAE; - - template - using _LValRefType _LIBCUDACXX_NODEBUG_TYPE = - typename __dependent_type<_DeleterSFINAE, _Dummy>::__lval_ref_type; - - template - using _GoodRValRefType _LIBCUDACXX_NODEBUG_TYPE = - typename __dependent_type<_DeleterSFINAE, _Dummy>::__good_rval_ref_type; - - template - using _BadRValRefType _LIBCUDACXX_NODEBUG_TYPE = - typename __dependent_type<_DeleterSFINAE, _Dummy>::__bad_rval_ref_type; - - template , _Dummy>::type> - using _EnableIfDeleterDefaultConstructible _LIBCUDACXX_NODEBUG_TYPE = - typename enable_if::value && - !is_pointer<_Deleter>::value>::type; - - template - using _EnableIfDeleterConstructible _LIBCUDACXX_NODEBUG_TYPE = - typename enable_if::value>::type; - - template - using _EnableIfMoveConvertible _LIBCUDACXX_NODEBUG_TYPE = typename enable_if< - is_convertible::value && - !is_array<_Up>::value - >::type; - - template - using _EnableIfDeleterConvertible _LIBCUDACXX_NODEBUG_TYPE = typename enable_if< - (is_reference<_Dp>::value && is_same<_Dp, _UDel>::value) || - (!is_reference<_Dp>::value && is_convertible<_UDel, _Dp>::value) - >::type; - - template - using _EnableIfDeleterAssignable = typename enable_if< - is_assignable<_Dp&, _UDel&&>::value - >::type; - -public: - template > - _LIBCUDACXX_INLINE_VISIBILITY - constexpr unique_ptr() noexcept : __ptr_(pointer()) {} - - template > - _LIBCUDACXX_INLINE_VISIBILITY - constexpr unique_ptr(nullptr_t) noexcept : __ptr_(pointer()) {} - - template > - _LIBCUDACXX_INLINE_VISIBILITY - explicit unique_ptr(pointer __p) noexcept : __ptr_(__p) {} - - template > > - _LIBCUDACXX_INLINE_VISIBILITY - unique_ptr(pointer __p, _LValRefType<_Dummy> __d) noexcept - : __ptr_(__p, __d) {} - - template > > - _LIBCUDACXX_INLINE_VISIBILITY - unique_ptr(pointer __p, _GoodRValRefType<_Dummy> __d) noexcept - : __ptr_(__p, _CUDA_VSTD::move(__d)) { - static_assert(!is_reference::value, - "rvalue deleter bound to reference"); - } - - template > > - _LIBCUDACXX_INLINE_VISIBILITY - unique_ptr(pointer __p, _BadRValRefType<_Dummy> __d) = delete; - - _LIBCUDACXX_INLINE_VISIBILITY - unique_ptr(unique_ptr&& __u) noexcept - : __ptr_(__u.release(), _CUDA_VSTD::forward(__u.get_deleter())) { - } - - template , _Up>, - class = _EnableIfDeleterConvertible<_Ep> - > - _LIBCUDACXX_INLINE_VISIBILITY - unique_ptr(unique_ptr<_Up, _Ep>&& __u) noexcept - : __ptr_(__u.release(), _CUDA_VSTD::forward<_Ep>(__u.get_deleter())) {} - -#if _CCCL_STD_VER <= 2014 || defined(_LIBCUDACXX_ENABLE_CXX17_REMOVED_AUTO_PTR) - template - _LIBCUDACXX_INLINE_VISIBILITY - unique_ptr(auto_ptr<_Up>&& __p, - typename enable_if::value && - is_same<_Dp, default_delete<_Tp> >::value, - __nat>::type = __nat()) noexcept - : __ptr_(__p.release()) {} -#endif - - _LIBCUDACXX_INLINE_VISIBILITY - unique_ptr& operator=(unique_ptr&& __u) noexcept { - reset(__u.release()); - __ptr_.second() = _CUDA_VSTD::forward(__u.get_deleter()); - return *this; - } - - template , _Up>, - class = _EnableIfDeleterAssignable<_Ep> - > - _LIBCUDACXX_INLINE_VISIBILITY - unique_ptr& operator=(unique_ptr<_Up, _Ep>&& __u) noexcept { - reset(__u.release()); - __ptr_.second() = _CUDA_VSTD::forward<_Ep>(__u.get_deleter()); - return *this; - } - -#if _CCCL_STD_VER <= 2014 || defined(_LIBCUDACXX_ENABLE_CXX17_REMOVED_AUTO_PTR) - template - _LIBCUDACXX_INLINE_VISIBILITY - typename enable_if::value && - is_same<_Dp, default_delete<_Tp> >::value, - unique_ptr&>::type - operator=(auto_ptr<_Up> __p) { - reset(__p.release()); - return *this; - } -#endif - - _LIBCUDACXX_INLINE_VISIBILITY - ~unique_ptr() { reset(); } - - _LIBCUDACXX_INLINE_VISIBILITY - unique_ptr& operator=(nullptr_t) noexcept { - reset(); - return *this; - } - - _LIBCUDACXX_INLINE_VISIBILITY - __add_lvalue_reference_t<_Tp> - operator*() const { - return *__ptr_.first(); - } - _LIBCUDACXX_INLINE_VISIBILITY - pointer operator->() const noexcept { - return __ptr_.first(); - } - _LIBCUDACXX_INLINE_VISIBILITY - pointer get() const noexcept { - return __ptr_.first(); - } - _LIBCUDACXX_INLINE_VISIBILITY - deleter_type& get_deleter() noexcept { - return __ptr_.second(); - } - _LIBCUDACXX_INLINE_VISIBILITY - const deleter_type& get_deleter() const noexcept { - return __ptr_.second(); - } - _LIBCUDACXX_INLINE_VISIBILITY - _LIBCUDACXX_EXPLICIT operator bool() const noexcept { - return __ptr_.first() != nullptr; - } - - _LIBCUDACXX_INLINE_VISIBILITY - pointer release() noexcept { - pointer __t = __ptr_.first(); - __ptr_.first() = pointer(); - return __t; - } - - _LIBCUDACXX_INLINE_VISIBILITY - void reset(pointer __p = pointer()) noexcept { - pointer __tmp = __ptr_.first(); - __ptr_.first() = __p; - if (__tmp) - __ptr_.second()(__tmp); - } - - _LIBCUDACXX_INLINE_VISIBILITY - void swap(unique_ptr& __u) noexcept { - __ptr_.swap(__u.__ptr_); - } -}; - - -template -class _LIBCUDACXX_TEMPLATE_VIS unique_ptr<_Tp[], _Dp> { -public: - typedef _Tp element_type; - typedef _Dp deleter_type; - typedef typename __pointer<_Tp, deleter_type>::type pointer; - -private: - __compressed_pair __ptr_; - - template - struct _CheckArrayPointerConversion : is_same<_From, pointer> {}; - - template - struct _CheckArrayPointerConversion<_FromElem*> - : integral_constant::value || - (is_same::value && - is_convertible<_FromElem(*)[], element_type(*)[]>::value) - > - {}; - - typedef __unique_ptr_deleter_sfinae<_Dp> _DeleterSFINAE; - - template - using _LValRefType _LIBCUDACXX_NODEBUG_TYPE = - typename __dependent_type<_DeleterSFINAE, _Dummy>::__lval_ref_type; - - template - using _GoodRValRefType _LIBCUDACXX_NODEBUG_TYPE = - typename __dependent_type<_DeleterSFINAE, _Dummy>::__good_rval_ref_type; - - template - using _BadRValRefType _LIBCUDACXX_NODEBUG_TYPE = - typename __dependent_type<_DeleterSFINAE, _Dummy>::__bad_rval_ref_type; - - template , _Dummy>::type> - using _EnableIfDeleterDefaultConstructible _LIBCUDACXX_NODEBUG_TYPE = - typename enable_if::value && - !is_pointer<_Deleter>::value>::type; - - template - using _EnableIfDeleterConstructible _LIBCUDACXX_NODEBUG_TYPE = - typename enable_if::value>::type; - - template - using _EnableIfPointerConvertible _LIBCUDACXX_NODEBUG_TYPE = typename enable_if< - _CheckArrayPointerConversion<_Pp>::value - >::type; - - template - using _EnableIfMoveConvertible _LIBCUDACXX_NODEBUG_TYPE = typename enable_if< - is_array<_Up>::value && - is_same::value && - is_same::value && - is_convertible<_ElemT(*)[], element_type(*)[]>::value - >::type; - - template - using _EnableIfDeleterConvertible _LIBCUDACXX_NODEBUG_TYPE = typename enable_if< - (is_reference<_Dp>::value && is_same<_Dp, _UDel>::value) || - (!is_reference<_Dp>::value && is_convertible<_UDel, _Dp>::value) - >::type; - - template - using _EnableIfDeleterAssignable _LIBCUDACXX_NODEBUG_TYPE = typename enable_if< - is_assignable<_Dp&, _UDel&&>::value - >::type; - -public: - template > - _LIBCUDACXX_INLINE_VISIBILITY - constexpr unique_ptr() noexcept : __ptr_(pointer()) {} - - template > - _LIBCUDACXX_INLINE_VISIBILITY - constexpr unique_ptr(nullptr_t) noexcept : __ptr_(pointer()) {} - - template , - class = _EnableIfPointerConvertible<_Pp> > - _LIBCUDACXX_INLINE_VISIBILITY - explicit unique_ptr(_Pp __p) noexcept - : __ptr_(__p) {} - - template >, - class = _EnableIfPointerConvertible<_Pp> > - _LIBCUDACXX_INLINE_VISIBILITY - unique_ptr(_Pp __p, _LValRefType<_Dummy> __d) noexcept - : __ptr_(__p, __d) {} - - template > > - _LIBCUDACXX_INLINE_VISIBILITY - unique_ptr(nullptr_t, _LValRefType<_Dummy> __d) noexcept - : __ptr_(nullptr, __d) {} - - template >, - class = _EnableIfPointerConvertible<_Pp> > - _LIBCUDACXX_INLINE_VISIBILITY - unique_ptr(_Pp __p, _GoodRValRefType<_Dummy> __d) noexcept - : __ptr_(__p, _CUDA_VSTD::move(__d)) { - static_assert(!is_reference::value, - "rvalue deleter bound to reference"); - } - - template > > - _LIBCUDACXX_INLINE_VISIBILITY - unique_ptr(nullptr_t, _GoodRValRefType<_Dummy> __d) noexcept - : __ptr_(nullptr, _CUDA_VSTD::move(__d)) { - static_assert(!is_reference::value, - "rvalue deleter bound to reference"); - } - - template >, - class = _EnableIfPointerConvertible<_Pp> > - _LIBCUDACXX_INLINE_VISIBILITY - unique_ptr(_Pp __p, _BadRValRefType<_Dummy> __d) = delete; - - _LIBCUDACXX_INLINE_VISIBILITY - unique_ptr(unique_ptr&& __u) noexcept - : __ptr_(__u.release(), _CUDA_VSTD::forward(__u.get_deleter())) { - } - - _LIBCUDACXX_INLINE_VISIBILITY - unique_ptr& operator=(unique_ptr&& __u) noexcept { - reset(__u.release()); - __ptr_.second() = _CUDA_VSTD::forward(__u.get_deleter()); - return *this; - } - - template , _Up>, - class = _EnableIfDeleterConvertible<_Ep> - > - _LIBCUDACXX_INLINE_VISIBILITY - unique_ptr(unique_ptr<_Up, _Ep>&& __u) noexcept - : __ptr_(__u.release(), _CUDA_VSTD::forward<_Ep>(__u.get_deleter())) { - } - - template , _Up>, - class = _EnableIfDeleterAssignable<_Ep> - > - _LIBCUDACXX_INLINE_VISIBILITY - unique_ptr& - operator=(unique_ptr<_Up, _Ep>&& __u) noexcept { - reset(__u.release()); - __ptr_.second() = _CUDA_VSTD::forward<_Ep>(__u.get_deleter()); - return *this; - } - -public: - _LIBCUDACXX_INLINE_VISIBILITY - ~unique_ptr() { reset(); } - - _LIBCUDACXX_INLINE_VISIBILITY - unique_ptr& operator=(nullptr_t) noexcept { - reset(); - return *this; - } - - _LIBCUDACXX_INLINE_VISIBILITY - __add_lvalue_reference_t<_Tp> - operator[](size_t __i) const { - return __ptr_.first()[__i]; - } - _LIBCUDACXX_INLINE_VISIBILITY - pointer get() const noexcept { - return __ptr_.first(); - } - - _LIBCUDACXX_INLINE_VISIBILITY - deleter_type& get_deleter() noexcept { - return __ptr_.second(); - } - - _LIBCUDACXX_INLINE_VISIBILITY - const deleter_type& get_deleter() const noexcept { - return __ptr_.second(); - } - _LIBCUDACXX_INLINE_VISIBILITY - _LIBCUDACXX_EXPLICIT operator bool() const noexcept { - return __ptr_.first() != nullptr; - } - - _LIBCUDACXX_INLINE_VISIBILITY - pointer release() noexcept { - pointer __t = __ptr_.first(); - __ptr_.first() = pointer(); - return __t; - } - - template - _LIBCUDACXX_INLINE_VISIBILITY - typename enable_if< - _CheckArrayPointerConversion<_Pp>::value - >::type - reset(_Pp __p) noexcept { - pointer __tmp = __ptr_.first(); - __ptr_.first() = __p; - if (__tmp) - __ptr_.second()(__tmp); - } - - _LIBCUDACXX_INLINE_VISIBILITY - void reset(nullptr_t = nullptr) noexcept { - pointer __tmp = __ptr_.first(); - __ptr_.first() = nullptr; - if (__tmp) - __ptr_.second()(__tmp); - } - - _LIBCUDACXX_INLINE_VISIBILITY - void swap(unique_ptr& __u) noexcept { - __ptr_.swap(__u.__ptr_); - } - -}; - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -typename enable_if< - __is_swappable<_Dp>::value, - void ->::type -swap(unique_ptr<_Tp, _Dp>& __x, unique_ptr<_Tp, _Dp>& __y) noexcept {__x.swap(__y);} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator==(const unique_ptr<_T1, _Dest1>& __x, const unique_ptr<_T2, _Dest2>& __y) {return __x.get() == __y.get();} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator!=(const unique_ptr<_T1, _Dest1>& __x, const unique_ptr<_T2, _Dest2>& __y) {return !(__x == __y);} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator< (const unique_ptr<_T1, _Dest1>& __x, const unique_ptr<_T2, _Dest2>& __y) -{ - typedef typename unique_ptr<_T1, _Dest1>::pointer _P1; - typedef typename unique_ptr<_T2, _Dest2>::pointer _P2; - typedef typename common_type<_P1, _P2>::type _Vp; - return less<_Vp>()(__x.get(), __y.get()); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator> (const unique_ptr<_T1, _Dest1>& __x, const unique_ptr<_T2, _Dest2>& __y) {return __y < __x;} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator<=(const unique_ptr<_T1, _Dest1>& __x, const unique_ptr<_T2, _Dest2>& __y) {return !(__y < __x);} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator>=(const unique_ptr<_T1, _Dest1>& __x, const unique_ptr<_T2, _Dest2>& __y) {return !(__x < __y);} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator==(const unique_ptr<_T1, _Dest1>& __x, nullptr_t) noexcept -{ - return !__x; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator==(nullptr_t, const unique_ptr<_T1, _Dest1>& __x) noexcept -{ - return !__x; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator!=(const unique_ptr<_T1, _Dest1>& __x, nullptr_t) noexcept -{ - return static_cast(__x); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator!=(nullptr_t, const unique_ptr<_T1, _Dest1>& __x) noexcept -{ - return static_cast(__x); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator<(const unique_ptr<_T1, _Dest1>& __x, nullptr_t) -{ - typedef typename unique_ptr<_T1, _Dest1>::pointer _P1; - return less<_P1>()(__x.get(), nullptr); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator<(nullptr_t, const unique_ptr<_T1, _Dest1>& __x) -{ - typedef typename unique_ptr<_T1, _Dest1>::pointer _P1; - return less<_P1>()(nullptr, __x.get()); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator>(const unique_ptr<_T1, _Dest1>& __x, nullptr_t) -{ - return nullptr < __x; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator>(nullptr_t, const unique_ptr<_T1, _Dest1>& __x) -{ - return __x < nullptr; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator<=(const unique_ptr<_T1, _Dest1>& __x, nullptr_t) -{ - return !(nullptr < __x); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator<=(nullptr_t, const unique_ptr<_T1, _Dest1>& __x) -{ - return !(__x < nullptr); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator>=(const unique_ptr<_T1, _Dest1>& __x, nullptr_t) -{ - return !(__x < nullptr); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator>=(nullptr_t, const unique_ptr<_T1, _Dest1>& __x) -{ - return !(nullptr < __x); -} - -#if _CCCL_STD_VER > 2011 - -template -struct __unique_if -{ - typedef unique_ptr<_Tp> __unique_single; -}; - -template -struct __unique_if<_Tp[]> -{ - typedef unique_ptr<_Tp[]> __unique_array_unknown_bound; -}; - -template -struct __unique_if<_Tp[_Np]> -{ - typedef void __unique_array_known_bound; -}; - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -typename __unique_if<_Tp>::__unique_single -make_unique(_Args&&... __args) -{ - return unique_ptr<_Tp>(new _Tp(_CUDA_VSTD::forward<_Args>(__args)...)); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -typename __unique_if<_Tp>::__unique_array_unknown_bound -make_unique(size_t __n) -{ - typedef typename remove_extent<_Tp>::type _Up; - return unique_ptr<_Tp>(new _Up[__n]()); -} - -template - typename __unique_if<_Tp>::__unique_array_known_bound - make_unique(_Args&&...) = delete; - -#endif // _CCCL_STD_VER > 2011 - -template -struct _LIBCUDACXX_TEMPLATE_VIS hash<__enable_hash_helper< - unique_ptr<_Tp, _Dp>, typename unique_ptr<_Tp, _Dp>::pointer> > -{ - typedef unique_ptr<_Tp, _Dp> argument_type; - typedef size_t result_type; - _LIBCUDACXX_INLINE_VISIBILITY - result_type operator()(const argument_type& __ptr) const - { - typedef typename argument_type::pointer pointer; - return hash()(__ptr.get()); - } -}; - -struct __destruct_n -{ -private: - size_t __size_; - - template - _LIBCUDACXX_INLINE_VISIBILITY void __process(_Tp* __p, false_type) noexcept - {for (size_t __i = 0; __i < __size_; ++__i, ++__p) __p->~_Tp();} - - template - _LIBCUDACXX_INLINE_VISIBILITY void __process(_Tp*, true_type) noexcept - {} - - _LIBCUDACXX_INLINE_VISIBILITY void __incr(false_type) noexcept - {++__size_;} - _LIBCUDACXX_INLINE_VISIBILITY void __incr(true_type) noexcept - {} - - _LIBCUDACXX_INLINE_VISIBILITY void __set(size_t __s, false_type) noexcept - {__size_ = __s;} - _LIBCUDACXX_INLINE_VISIBILITY void __set(size_t, true_type) noexcept - {} -public: - _LIBCUDACXX_INLINE_VISIBILITY explicit __destruct_n(size_t __s) noexcept - : __size_(__s) {} - - template - _LIBCUDACXX_INLINE_VISIBILITY void __incr(_Tp*) noexcept - {__incr(integral_constant::value>());} - - template - _LIBCUDACXX_INLINE_VISIBILITY void __set(size_t __s, _Tp*) noexcept - {__set(__s, integral_constant::value>());} - - template - _LIBCUDACXX_INLINE_VISIBILITY void operator()(_Tp* __p) noexcept - {__process(__p, integral_constant::value>());} -}; - -template -_ForwardIterator -uninitialized_copy(_InputIterator __f, _InputIterator __l, _ForwardIterator __r) -{ - typedef typename iterator_traits<_ForwardIterator>::value_type value_type; -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - _ForwardIterator __s = __r; - try - { -#endif - for (; __f != __l; ++__f, (void) ++__r) - ::new (static_cast(_CUDA_VSTD::addressof(*__r))) value_type(*__f); -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - } - catch (...) - { - for (; __s != __r; ++__s) - __s->~value_type(); - throw; - } -#endif - return __r; -} - -template -_ForwardIterator -uninitialized_copy_n(_InputIterator __f, _Size __n, _ForwardIterator __r) -{ - typedef typename iterator_traits<_ForwardIterator>::value_type value_type; -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - _ForwardIterator __s = __r; - try - { -#endif - for (; __n > 0; ++__f, (void) ++__r, (void) --__n) - ::new (static_cast(_CUDA_VSTD::addressof(*__r))) value_type(*__f); -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - } - catch (...) - { - for (; __s != __r; ++__s) - __s->~value_type(); - throw; - } -#endif - return __r; -} - -template -void -uninitialized_fill(_ForwardIterator __f, _ForwardIterator __l, const _Tp& __x) -{ - typedef typename iterator_traits<_ForwardIterator>::value_type value_type; -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - _ForwardIterator __s = __f; - try - { -#endif - for (; __f != __l; ++__f) - ::new (static_cast(_CUDA_VSTD::addressof(*__f))) value_type(__x); -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - } - catch (...) - { - for (; __s != __f; ++__s) - __s->~value_type(); - throw; - } -#endif -} - -template -_ForwardIterator -uninitialized_fill_n(_ForwardIterator __f, _Size __n, const _Tp& __x) -{ - typedef typename iterator_traits<_ForwardIterator>::value_type value_type; -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - _ForwardIterator __s = __f; - try - { -#endif - for (; __n > 0; ++__f, (void) --__n) - ::new (static_cast(_CUDA_VSTD::addressof(*__f))) value_type(__x); -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - } - catch (...) - { - for (; __s != __f; ++__s) - __s->~value_type(); - throw; - } -#endif - return __f; -} - -#if _CCCL_STD_VER > 2014 - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -void uninitialized_default_construct(_ForwardIterator __first, _ForwardIterator __last) { - using _Vt = typename iterator_traits<_ForwardIterator>::value_type; - auto __idx = __first; -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - try { -#endif - for (; __idx != __last; ++__idx) - ::new((void*)_CUDA_VSTD::addressof(*__idx)) _Vt; -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - } catch (...) { - _CUDA_VSTD::destroy(__first, __idx); - throw; - } -#endif -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_ForwardIterator uninitialized_default_construct_n(_ForwardIterator __first, _Size __n) { - using _Vt = typename iterator_traits<_ForwardIterator>::value_type; - auto __idx = __first; -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - try { -#endif - for (; __n > 0; (void)++__idx, --__n) - ::new((void*)_CUDA_VSTD::addressof(*__idx)) _Vt; - return __idx; -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - } catch (...) { - _CUDA_VSTD::destroy(__first, __idx); - throw; - } -#endif -} - - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -void uninitialized_value_construct(_ForwardIterator __first, _ForwardIterator __last) { - using _Vt = typename iterator_traits<_ForwardIterator>::value_type; - auto __idx = __first; -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - try { -#endif - for (; __idx != __last; ++__idx) - ::new((void*)_CUDA_VSTD::addressof(*__idx)) _Vt(); -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - } catch (...) { - _CUDA_VSTD::destroy(__first, __idx); - throw; - } -#endif -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_ForwardIterator uninitialized_value_construct_n(_ForwardIterator __first, _Size __n) { - using _Vt = typename iterator_traits<_ForwardIterator>::value_type; - auto __idx = __first; -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - try { -#endif - for (; __n > 0; (void)++__idx, --__n) - ::new((void*)_CUDA_VSTD::addressof(*__idx)) _Vt(); - return __idx; -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - } catch (...) { - _CUDA_VSTD::destroy(__first, __idx); - throw; - } -#endif -} - - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_ForwardIt uninitialized_move(_InputIt __first, _InputIt __last, _ForwardIt __first_res) { - using _Vt = typename iterator_traits<_ForwardIt>::value_type; - auto __idx = __first_res; -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - try { -#endif - for (; __first != __last; (void)++__idx, ++__first) - ::new((void*)_CUDA_VSTD::addressof(*__idx)) _Vt(std::move(*__first)); - return __idx; -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - } catch (...) { - _CUDA_VSTD::destroy(__first_res, __idx); - throw; - } -#endif -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -pair<_InputIt, _ForwardIt> -uninitialized_move_n(_InputIt __first, _Size __n, _ForwardIt __first_res) { - using _Vt = typename iterator_traits<_ForwardIt>::value_type; - auto __idx = __first_res; -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - try { -#endif - for (; __n > 0; ++__idx, (void)++__first, --__n) - ::new((void*)_CUDA_VSTD::addressof(*__idx)) _Vt(std::move(*__first)); - return {__first, __idx}; -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - } catch (...) { - _CUDA_VSTD::destroy(__first_res, __idx); - throw; - } -#endif -} - - -#endif // _CCCL_STD_VER > 2014 - -// NOTE: Relaxed and acq/rel atomics (for increment and decrement respectively) -// should be sufficient for thread safety. -// See https://bugs.llvm.org/show_bug.cgi?id=22803 -#if defined(_CCCL_COMPILER_CLANG) \ - && __has_builtin(__atomic_add_fetch) \ - && defined(__ATOMIC_RELAXED) \ - && defined(__ATOMIC_ACQ_REL) -# define _LIBCUDACXX_HAS_BUILTIN_ATOMIC_SUPPORT -#elif defined(_CCCL_COMPILER_GCC) -# define _LIBCUDACXX_HAS_BUILTIN_ATOMIC_SUPPORT -#endif - -template -inline _LIBCUDACXX_INLINE_VISIBILITY _Tp -__libcpp_atomic_refcount_increment(_Tp& __t) noexcept -{ -#if defined(_LIBCUDACXX_HAS_BUILTIN_ATOMIC_SUPPORT) && !defined(_LIBCUDACXX_HAS_NO_THREADS) - return __atomic_add_fetch(&__t, 1, __ATOMIC_RELAXED); -#else - return __t += 1; -#endif -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY _Tp -__libcpp_atomic_refcount_decrement(_Tp& __t) noexcept -{ -#if defined(_LIBCUDACXX_HAS_BUILTIN_ATOMIC_SUPPORT) && !defined(_LIBCUDACXX_HAS_NO_THREADS) - return __atomic_add_fetch(&__t, -1, __ATOMIC_ACQ_REL); -#else - return __t -= 1; -#endif -} - -class _LIBCUDACXX_EXCEPTION_ABI bad_weak_ptr - : public std::exception -{ -public: - virtual ~bad_weak_ptr() noexcept; - virtual const char* what() const noexcept; -}; - -_CCCL_NORETURN inline _LIBCUDACXX_INLINE_VISIBILITY -void __throw_bad_weak_ptr() -{ -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - throw bad_weak_ptr(); -#else - _CUDA_VSTD::abort(); -#endif -} - -template class _LIBCUDACXX_TEMPLATE_VIS weak_ptr; - -class _LIBCUDACXX_TYPE_VIS __shared_count -{ - __shared_count(const __shared_count&); - __shared_count& operator=(const __shared_count&); - -protected: - long __shared_owners_; - virtual ~__shared_count(); -private: - virtual void __on_zero_shared() noexcept = 0; - -public: - _LIBCUDACXX_INLINE_VISIBILITY - explicit __shared_count(long __refs = 0) noexcept - : __shared_owners_(__refs) {} - -#if defined(_LIBCUDACXX_BUILDING_LIBRARY) && \ - defined(_LIBCUDACXX_DEPRECATED_ABI_LEGACY_LIBRARY_DEFINITIONS_FOR_INLINE_FUNCTIONS) - void __add_shared() noexcept; - bool __release_shared() noexcept; -#else - _LIBCUDACXX_INLINE_VISIBILITY - void __add_shared() noexcept { - __libcpp_atomic_refcount_increment(__shared_owners_); - } - _LIBCUDACXX_INLINE_VISIBILITY - bool __release_shared() noexcept { - if (__libcpp_atomic_refcount_decrement(__shared_owners_) == -1) { - __on_zero_shared(); - return true; - } - return false; - } -#endif - _LIBCUDACXX_INLINE_VISIBILITY - long use_count() const noexcept { - return __libcpp_relaxed_load(&__shared_owners_) + 1; - } -}; - -class _LIBCUDACXX_TYPE_VIS __shared_weak_count - : private __shared_count -{ - long __shared_weak_owners_; - -public: - _LIBCUDACXX_INLINE_VISIBILITY - explicit __shared_weak_count(long __refs = 0) noexcept - : __shared_count(__refs), - __shared_weak_owners_(__refs) {} -protected: - virtual ~__shared_weak_count(); - -public: -#if defined(_LIBCUDACXX_BUILDING_LIBRARY) && \ - defined(_LIBCUDACXX_DEPRECATED_ABI_LEGACY_LIBRARY_DEFINITIONS_FOR_INLINE_FUNCTIONS) - void __add_shared() noexcept; - void __add_weak() noexcept; - void __release_shared() noexcept; -#else - _LIBCUDACXX_INLINE_VISIBILITY - void __add_shared() noexcept { - __shared_count::__add_shared(); - } - _LIBCUDACXX_INLINE_VISIBILITY - void __add_weak() noexcept { - __libcpp_atomic_refcount_increment(__shared_weak_owners_); - } - _LIBCUDACXX_INLINE_VISIBILITY - void __release_shared() noexcept { - if (__shared_count::__release_shared()) - __release_weak(); - } -#endif - void __release_weak() noexcept; - _LIBCUDACXX_INLINE_VISIBILITY - long use_count() const noexcept {return __shared_count::use_count();} - __shared_weak_count* lock() noexcept; - - // Define the function out only if we build static libc++ without RTTI. - // Otherwise we may break clients who need to compile their projects with - // -fno-rtti and yet link against a libc++.dylib compiled - // without -fno-rtti. -#if !defined(_LIBCUDACXX_NO_RTTI) || !defined(_LIBCUDACXX_BUILD_STATIC) - virtual const void* __get_deleter(const type_info&) const noexcept; -#endif -private: - virtual void __on_zero_shared_weak() noexcept = 0; -}; - -template -class __shared_ptr_pointer - : public __shared_weak_count -{ - __compressed_pair<__compressed_pair<_Tp, _Dp>, _Alloc> __data_; -public: - _LIBCUDACXX_INLINE_VISIBILITY - __shared_ptr_pointer(_Tp __p, _Dp __d, _Alloc __a) - : __data_(__compressed_pair<_Tp, _Dp>(__p, _CUDA_VSTD::move(__d)), _CUDA_VSTD::move(__a)) {} - -#ifndef _LIBCUDACXX_NO_RTTI - virtual const void* __get_deleter(const type_info&) const noexcept; -#endif - -private: - virtual void __on_zero_shared() noexcept; - virtual void __on_zero_shared_weak() noexcept; -}; - -#ifndef _LIBCUDACXX_NO_RTTI - -template -const void* -__shared_ptr_pointer<_Tp, _Dp, _Alloc>::__get_deleter(const type_info& __t) const noexcept -{ - return __t == typeid(_Dp) ? _CUDA_VSTD::addressof(__data_.first().second()) : nullptr; -} - -#endif // _LIBCUDACXX_NO_RTTI - -template -void -__shared_ptr_pointer<_Tp, _Dp, _Alloc>::__on_zero_shared() noexcept -{ - __data_.first().second()(__data_.first().first()); - __data_.first().second().~_Dp(); -} - -template -void -__shared_ptr_pointer<_Tp, _Dp, _Alloc>::__on_zero_shared_weak() noexcept -{ - typedef typename __allocator_traits_rebind<_Alloc, __shared_ptr_pointer>::type _Al; - typedef allocator_traits<_Al> _ATraits; - typedef pointer_traits _PTraits; - - _Al __a(__data_.second()); - __data_.second().~_Alloc(); - __a.deallocate(_PTraits::pointer_to(*this), 1); -} - -template -class __shared_ptr_emplace - : public __shared_weak_count -{ - __compressed_pair<_Alloc, _Tp> __data_; -public: -#ifndef _LIBCUDACXX_HAS_NO_VARIADICS - - _LIBCUDACXX_INLINE_VISIBILITY - __shared_ptr_emplace(_Alloc __a) - : __data_(_CUDA_VSTD::move(__a)) {} - - template - _LIBCUDACXX_INLINE_VISIBILITY - __shared_ptr_emplace(_Alloc __a, _Args&& ...__args) - : __data_(piecewise_construct, _CUDA_VSTD::forward_as_tuple(__a), - _CUDA_VSTD::forward_as_tuple(_CUDA_VSTD::forward<_Args>(__args)...)) {} - -#else // _LIBCUDACXX_HAS_NO_VARIADICS - - _LIBCUDACXX_INLINE_VISIBILITY - __shared_ptr_emplace(_Alloc __a) - : __data_(__a) {} - - template - _LIBCUDACXX_INLINE_VISIBILITY - __shared_ptr_emplace(_Alloc __a, _A0& __a0) - : __data_(__a, _Tp(__a0)) {} - - template - _LIBCUDACXX_INLINE_VISIBILITY - __shared_ptr_emplace(_Alloc __a, _A0& __a0, _A1& __a1) - : __data_(__a, _Tp(__a0, __a1)) {} - - template - _LIBCUDACXX_INLINE_VISIBILITY - __shared_ptr_emplace(_Alloc __a, _A0& __a0, _A1& __a1, _A2& __a2) - : __data_(__a, _Tp(__a0, __a1, __a2)) {} - -#endif // _LIBCUDACXX_HAS_NO_VARIADICS - -private: - virtual void __on_zero_shared() noexcept; - virtual void __on_zero_shared_weak() noexcept; -public: - _LIBCUDACXX_INLINE_VISIBILITY - _Tp* get() noexcept {return _CUDA_VSTD::addressof(__data_.second());} -}; - -template -void -__shared_ptr_emplace<_Tp, _Alloc>::__on_zero_shared() noexcept -{ - __data_.second().~_Tp(); -} - -template -void -__shared_ptr_emplace<_Tp, _Alloc>::__on_zero_shared_weak() noexcept -{ - typedef typename __allocator_traits_rebind<_Alloc, __shared_ptr_emplace>::type _Al; - typedef allocator_traits<_Al> _ATraits; - typedef pointer_traits _PTraits; - _Al __a(__data_.first()); - __data_.first().~_Alloc(); - __a.deallocate(_PTraits::pointer_to(*this), 1); -} - -struct __shared_ptr_dummy_rebind_allocator_type; -template <> -class _LIBCUDACXX_TEMPLATE_VIS allocator<__shared_ptr_dummy_rebind_allocator_type> -{ -public: - template - struct rebind - { - typedef allocator<_Other> other; - }; -}; - -template class _LIBCUDACXX_TEMPLATE_VIS enable_shared_from_this; - -template -class _LIBCUDACXX_TEMPLATE_VIS shared_ptr -{ -public: - typedef _Tp element_type; - -#if _CCCL_STD_VER > 2014 - typedef weak_ptr<_Tp> weak_type; -#endif -private: - element_type* __ptr_; - __shared_weak_count* __cntrl_; - - struct __nat {int __for_bool_;}; -public: - _LIBCUDACXX_INLINE_VISIBILITY - constexpr shared_ptr() noexcept; - _LIBCUDACXX_INLINE_VISIBILITY - constexpr shared_ptr(nullptr_t) noexcept; - template - explicit shared_ptr(_Yp* __p, - typename enable_if::value, __nat>::type = __nat()); - template - shared_ptr(_Yp* __p, _Dp __d, - typename enable_if::value, __nat>::type = __nat()); - template - shared_ptr(_Yp* __p, _Dp __d, _Alloc __a, - typename enable_if::value, __nat>::type = __nat()); - template shared_ptr(nullptr_t __p, _Dp __d); - template shared_ptr(nullptr_t __p, _Dp __d, _Alloc __a); - template _LIBCUDACXX_INLINE_VISIBILITY shared_ptr(const shared_ptr<_Yp>& __r, element_type* __p) noexcept; - _LIBCUDACXX_INLINE_VISIBILITY - shared_ptr(const shared_ptr& __r) noexcept; - template - _LIBCUDACXX_INLINE_VISIBILITY - shared_ptr(const shared_ptr<_Yp>& __r, - typename enable_if::value, __nat>::type = __nat()) - noexcept; -#ifndef _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - _LIBCUDACXX_INLINE_VISIBILITY - shared_ptr(shared_ptr&& __r) noexcept; - template _LIBCUDACXX_INLINE_VISIBILITY shared_ptr(shared_ptr<_Yp>&& __r, - typename enable_if::value, __nat>::type = __nat()) - noexcept; -#endif // _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - template explicit shared_ptr(const weak_ptr<_Yp>& __r, - typename enable_if::value, __nat>::type= __nat()); -#if _CCCL_STD_VER <= 2014 || defined(_LIBCUDACXX_ENABLE_CXX17_REMOVED_AUTO_PTR) -#ifndef _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - template - shared_ptr(auto_ptr<_Yp>&& __r, - typename enable_if::value, __nat>::type = __nat()); -#else - template - shared_ptr(auto_ptr<_Yp> __r, - typename enable_if::value, __nat>::type = __nat()); -#endif -#endif -#ifndef _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - template - shared_ptr(unique_ptr<_Yp, _Dp>&&, - typename enable_if - < - !is_lvalue_reference<_Dp>::value && - !is_array<_Yp>::value && - is_convertible::pointer, element_type*>::value, - __nat - >::type = __nat()); - template - shared_ptr(unique_ptr<_Yp, _Dp>&&, - typename enable_if - < - is_lvalue_reference<_Dp>::value && - !is_array<_Yp>::value && - is_convertible::pointer, element_type*>::value, - __nat - >::type = __nat()); -#else // _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - template - shared_ptr(unique_ptr<_Yp, _Dp>, - typename enable_if - < - !is_lvalue_reference<_Dp>::value && - !is_array<_Yp>::value && - is_convertible::pointer, element_type*>::value, - __nat - >::type = __nat()); - template - shared_ptr(unique_ptr<_Yp, _Dp>, - typename enable_if - < - is_lvalue_reference<_Dp>::value && - !is_array<_Yp>::value && - is_convertible::pointer, element_type*>::value, - __nat - >::type = __nat()); -#endif // _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - - ~shared_ptr(); - - _LIBCUDACXX_INLINE_VISIBILITY - shared_ptr& operator=(const shared_ptr& __r) noexcept; - template - typename enable_if - < - is_convertible<_Yp*, element_type*>::value, - shared_ptr& - >::type - _LIBCUDACXX_INLINE_VISIBILITY - operator=(const shared_ptr<_Yp>& __r) noexcept; -#ifndef _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - _LIBCUDACXX_INLINE_VISIBILITY - shared_ptr& operator=(shared_ptr&& __r) noexcept; - template - typename enable_if - < - is_convertible<_Yp*, element_type*>::value, - shared_ptr<_Tp>& - >::type - _LIBCUDACXX_INLINE_VISIBILITY - operator=(shared_ptr<_Yp>&& __r); -#if _CCCL_STD_VER <= 2014 || defined(_LIBCUDACXX_ENABLE_CXX17_REMOVED_AUTO_PTR) - template - _LIBCUDACXX_INLINE_VISIBILITY - typename enable_if - < - !is_array<_Yp>::value && - is_convertible<_Yp*, element_type*>::value, - shared_ptr - >::type& - operator=(auto_ptr<_Yp>&& __r); -#endif -#else // _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES -#if _CCCL_STD_VER <= 2014 || defined(_LIBCUDACXX_ENABLE_CXX17_REMOVED_AUTO_PTR) - template - _LIBCUDACXX_INLINE_VISIBILITY - typename enable_if - < - !is_array<_Yp>::value && - is_convertible<_Yp*, element_type*>::value, - shared_ptr& - >::type - operator=(auto_ptr<_Yp> __r); -#endif -#endif - template - typename enable_if - < - !is_array<_Yp>::value && - is_convertible::pointer, element_type*>::value, - shared_ptr& - >::type -#ifndef _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - _LIBCUDACXX_INLINE_VISIBILITY - operator=(unique_ptr<_Yp, _Dp>&& __r); -#else // _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - _LIBCUDACXX_INLINE_VISIBILITY - operator=(unique_ptr<_Yp, _Dp> __r); -#endif - - _LIBCUDACXX_INLINE_VISIBILITY - void swap(shared_ptr& __r) noexcept; - _LIBCUDACXX_INLINE_VISIBILITY - void reset() noexcept; - template - typename enable_if - < - is_convertible<_Yp*, element_type*>::value, - void - >::type - _LIBCUDACXX_INLINE_VISIBILITY - reset(_Yp* __p); - template - typename enable_if - < - is_convertible<_Yp*, element_type*>::value, - void - >::type - _LIBCUDACXX_INLINE_VISIBILITY - reset(_Yp* __p, _Dp __d); - template - typename enable_if - < - is_convertible<_Yp*, element_type*>::value, - void - >::type - _LIBCUDACXX_INLINE_VISIBILITY - reset(_Yp* __p, _Dp __d, _Alloc __a); - - _LIBCUDACXX_INLINE_VISIBILITY - element_type* get() const noexcept {return __ptr_;} - _LIBCUDACXX_INLINE_VISIBILITY - __add_lvalue_reference_t operator*() const noexcept - {return *__ptr_;} - _LIBCUDACXX_INLINE_VISIBILITY - element_type* operator->() const noexcept {return __ptr_;} - _LIBCUDACXX_INLINE_VISIBILITY - long use_count() const noexcept {return __cntrl_ ? __cntrl_->use_count() : 0;} - _LIBCUDACXX_INLINE_VISIBILITY - bool unique() const noexcept {return use_count() == 1;} - _LIBCUDACXX_INLINE_VISIBILITY - _LIBCUDACXX_EXPLICIT operator bool() const noexcept {return get() != 0;} - template - _LIBCUDACXX_INLINE_VISIBILITY - bool owner_before(shared_ptr<_Up> const& __p) const noexcept - {return __cntrl_ < __p.__cntrl_;} - template - _LIBCUDACXX_INLINE_VISIBILITY - bool owner_before(weak_ptr<_Up> const& __p) const noexcept - {return __cntrl_ < __p.__cntrl_;} - _LIBCUDACXX_INLINE_VISIBILITY - bool - __owner_equivalent(const shared_ptr& __p) const - {return __cntrl_ == __p.__cntrl_;} - -#ifndef _LIBCUDACXX_NO_RTTI - template - _LIBCUDACXX_INLINE_VISIBILITY - _Dp* __get_deleter() const noexcept - {return static_cast<_Dp*>(__cntrl_ - ? const_cast(__cntrl_->__get_deleter(typeid(_Dp))) - : nullptr);} -#endif // _LIBCUDACXX_NO_RTTI - - template - static - shared_ptr<_Tp> - make_shared(_Args&& ...__args); - - template - static - shared_ptr<_Tp> - allocate_shared(const _Alloc& __a, _Args&& ...__args); - -private: - template ::value> - struct __shared_ptr_default_allocator - { - typedef allocator<_Yp> type; - }; - - template - struct __shared_ptr_default_allocator<_Yp, true> - { - typedef allocator<__shared_ptr_dummy_rebind_allocator_type> type; - }; - - template - _LIBCUDACXX_INLINE_VISIBILITY - typename enable_if* - >::value, - void>::type - __enable_weak_this(const enable_shared_from_this<_Yp>* __e, - _OrigPtr* __ptr) noexcept - { - typedef __remove_cv_t<_Yp> _RawYp; - if (__e && __e->__weak_this_.expired()) - { - __e->__weak_this_ = shared_ptr<_RawYp>(*this, - const_cast<_RawYp*>(static_cast(__ptr))); - } - } - - _LIBCUDACXX_INLINE_VISIBILITY void __enable_weak_this(...) noexcept {} - - template friend class _LIBCUDACXX_TEMPLATE_VIS shared_ptr; - template friend class _LIBCUDACXX_TEMPLATE_VIS weak_ptr; -}; - - -template -inline constexpr -shared_ptr<_Tp>::shared_ptr() noexcept - : __ptr_(0), - __cntrl_(0) -{ -} - -template -inline constexpr -shared_ptr<_Tp>::shared_ptr(nullptr_t) noexcept - : __ptr_(0), - __cntrl_(0) -{ -} - -template -template -shared_ptr<_Tp>::shared_ptr(_Yp* __p, - typename enable_if::value, __nat>::type) - : __ptr_(__p) -{ - unique_ptr<_Yp> __hold(__p); - typedef typename __shared_ptr_default_allocator<_Yp>::type _AllocT; - typedef __shared_ptr_pointer<_Yp*, default_delete<_Yp>, _AllocT > _CntrlBlk; - __cntrl_ = new _CntrlBlk(__p, default_delete<_Yp>(), _AllocT()); - __hold.release(); - __enable_weak_this(__p, __p); -} - -template -template -shared_ptr<_Tp>::shared_ptr(_Yp* __p, _Dp __d, - typename enable_if::value, __nat>::type) - : __ptr_(__p) -{ -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - try - { -#endif // _LIBCUDACXX_NO_EXCEPTIONS - typedef typename __shared_ptr_default_allocator<_Yp>::type _AllocT; - typedef __shared_ptr_pointer<_Yp*, _Dp, _AllocT > _CntrlBlk; - __cntrl_ = new _CntrlBlk(__p, __d, _AllocT()); - __enable_weak_this(__p, __p); -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - } - catch (...) - { - __d(__p); - throw; - } -#endif // _LIBCUDACXX_NO_EXCEPTIONS -} - -template -template -shared_ptr<_Tp>::shared_ptr(nullptr_t __p, _Dp __d) - : __ptr_(0) -{ -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - try - { -#endif // _LIBCUDACXX_NO_EXCEPTIONS - typedef typename __shared_ptr_default_allocator<_Tp>::type _AllocT; - typedef __shared_ptr_pointer _CntrlBlk; - __cntrl_ = new _CntrlBlk(__p, __d, _AllocT()); -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - } - catch (...) - { - __d(__p); - throw; - } -#endif // _LIBCUDACXX_NO_EXCEPTIONS -} - -template -template -shared_ptr<_Tp>::shared_ptr(_Yp* __p, _Dp __d, _Alloc __a, - typename enable_if::value, __nat>::type) - : __ptr_(__p) -{ -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - try - { -#endif // _LIBCUDACXX_NO_EXCEPTIONS - typedef __shared_ptr_pointer<_Yp*, _Dp, _Alloc> _CntrlBlk; - typedef typename __allocator_traits_rebind<_Alloc, _CntrlBlk>::type _A2; - typedef __allocator_destructor<_A2> _Dest2; - _A2 __a2(__a); - unique_ptr<_CntrlBlk, _Dest2> __hold2(__a2.allocate(1), _Dest2(__a2, 1)); - ::new(static_cast(_CUDA_VSTD::addressof(*__hold2.get()))) - _CntrlBlk(__p, __d, __a); - __cntrl_ = _CUDA_VSTD::addressof(*__hold2.release()); - __enable_weak_this(__p, __p); -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - } - catch (...) - { - __d(__p); - throw; - } -#endif // _LIBCUDACXX_NO_EXCEPTIONS -} - -template -template -shared_ptr<_Tp>::shared_ptr(nullptr_t __p, _Dp __d, _Alloc __a) - : __ptr_(0) -{ -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - try - { -#endif // _LIBCUDACXX_NO_EXCEPTIONS - typedef __shared_ptr_pointer _CntrlBlk; - typedef typename __allocator_traits_rebind<_Alloc, _CntrlBlk>::type _A2; - typedef __allocator_destructor<_A2> _Dest2; - _A2 __a2(__a); - unique_ptr<_CntrlBlk, _Dest2> __hold2(__a2.allocate(1), _Dest2(__a2, 1)); - ::new(static_cast(_CUDA_VSTD::addressof(*__hold2.get()))) - _CntrlBlk(__p, __d, __a); - __cntrl_ = _CUDA_VSTD::addressof(*__hold2.release()); -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - } - catch (...) - { - __d(__p); - throw; - } -#endif // _LIBCUDACXX_NO_EXCEPTIONS -} - -template -template -inline -shared_ptr<_Tp>::shared_ptr(const shared_ptr<_Yp>& __r, element_type *__p) noexcept - : __ptr_(__p), - __cntrl_(__r.__cntrl_) -{ - if (__cntrl_) - __cntrl_->__add_shared(); -} - -template -inline -shared_ptr<_Tp>::shared_ptr(const shared_ptr& __r) noexcept - : __ptr_(__r.__ptr_), - __cntrl_(__r.__cntrl_) -{ - if (__cntrl_) - __cntrl_->__add_shared(); -} - -template -template -inline -shared_ptr<_Tp>::shared_ptr(const shared_ptr<_Yp>& __r, - typename enable_if::value, __nat>::type) - noexcept - : __ptr_(__r.__ptr_), - __cntrl_(__r.__cntrl_) -{ - if (__cntrl_) - __cntrl_->__add_shared(); -} - -#ifndef _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - -template -inline -shared_ptr<_Tp>::shared_ptr(shared_ptr&& __r) noexcept - : __ptr_(__r.__ptr_), - __cntrl_(__r.__cntrl_) -{ - __r.__ptr_ = 0; - __r.__cntrl_ = 0; -} - -template -template -inline -shared_ptr<_Tp>::shared_ptr(shared_ptr<_Yp>&& __r, - typename enable_if::value, __nat>::type) - noexcept - : __ptr_(__r.__ptr_), - __cntrl_(__r.__cntrl_) -{ - __r.__ptr_ = 0; - __r.__cntrl_ = 0; -} - -#endif // _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - -#if _CCCL_STD_VER <= 2014 || defined(_LIBCUDACXX_ENABLE_CXX17_REMOVED_AUTO_PTR) -template -template -#ifndef _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES -shared_ptr<_Tp>::shared_ptr(auto_ptr<_Yp>&& __r, -#else -shared_ptr<_Tp>::shared_ptr(auto_ptr<_Yp> __r, -#endif - typename enable_if::value, __nat>::type) - : __ptr_(__r.get()) -{ - typedef __shared_ptr_pointer<_Yp*, default_delete<_Yp>, allocator<_Yp> > _CntrlBlk; - __cntrl_ = new _CntrlBlk(__r.get(), default_delete<_Yp>(), allocator<_Yp>()); - __enable_weak_this(__r.get(), __r.get()); - __r.release(); -} -#endif - -template -template -#ifndef _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES -shared_ptr<_Tp>::shared_ptr(unique_ptr<_Yp, _Dp>&& __r, -#else -shared_ptr<_Tp>::shared_ptr(unique_ptr<_Yp, _Dp> __r, -#endif - typename enable_if - < - !is_lvalue_reference<_Dp>::value && - !is_array<_Yp>::value && - is_convertible::pointer, element_type*>::value, - __nat - >::type) - : __ptr_(__r.get()) -{ -#if _CCCL_STD_VER > 2011 - if (__ptr_ == nullptr) - __cntrl_ = nullptr; - else -#endif - { - typedef typename __shared_ptr_default_allocator<_Yp>::type _AllocT; - typedef __shared_ptr_pointer<_Yp*, _Dp, _AllocT > _CntrlBlk; - __cntrl_ = new _CntrlBlk(__r.get(), __r.get_deleter(), _AllocT()); - __enable_weak_this(__r.get(), __r.get()); - } - __r.release(); -} - -template -template -#ifndef _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES -shared_ptr<_Tp>::shared_ptr(unique_ptr<_Yp, _Dp>&& __r, -#else -shared_ptr<_Tp>::shared_ptr(unique_ptr<_Yp, _Dp> __r, -#endif - typename enable_if - < - is_lvalue_reference<_Dp>::value && - !is_array<_Yp>::value && - is_convertible::pointer, element_type*>::value, - __nat - >::type) - : __ptr_(__r.get()) -{ -#if _CCCL_STD_VER > 2011 - if (__ptr_ == nullptr) - __cntrl_ = nullptr; - else -#endif - { - typedef typename __shared_ptr_default_allocator<_Yp>::type _AllocT; - typedef __shared_ptr_pointer<_Yp*, - reference_wrapper<__libcpp_remove_reference_t<_Dp>>, - _AllocT > _CntrlBlk; - __cntrl_ = new _CntrlBlk(__r.get(), ref(__r.get_deleter()), _AllocT()); - __enable_weak_this(__r.get(), __r.get()); - } - __r.release(); -} - -template -template -shared_ptr<_Tp> -shared_ptr<_Tp>::make_shared(_Args&& ...__args) -{ - static_assert( is_constructible<_Tp, _Args...>::value, "Can't construct object in make_shared" ); - typedef __shared_ptr_emplace<_Tp, allocator<_Tp> > _CntrlBlk; - typedef allocator<_CntrlBlk> _A2; - typedef __allocator_destructor<_A2> _Dest2; - _A2 __a2; - unique_ptr<_CntrlBlk, _Dest2> __hold2(__a2.allocate(1), _Dest2(__a2, 1)); - ::new(__hold2.get()) _CntrlBlk(__a2, _CUDA_VSTD::forward<_Args>(__args)...); - shared_ptr<_Tp> __r; - __r.__ptr_ = __hold2.get()->get(); - __r.__cntrl_ = __hold2.release(); - __r.__enable_weak_this(__r.__ptr_, __r.__ptr_); - return __r; -} - -template -template -shared_ptr<_Tp> -shared_ptr<_Tp>::allocate_shared(const _Alloc& __a, _Args&& ...__args) -{ - static_assert( is_constructible<_Tp, _Args...>::value, "Can't construct object in allocate_shared" ); - typedef __shared_ptr_emplace<_Tp, _Alloc> _CntrlBlk; - typedef typename __allocator_traits_rebind<_Alloc, _CntrlBlk>::type _A2; - typedef __allocator_destructor<_A2> _Dest2; - _A2 __a2(__a); - unique_ptr<_CntrlBlk, _Dest2> __hold2(__a2.allocate(1), _Dest2(__a2, 1)); - ::new(static_cast(_CUDA_VSTD::addressof(*__hold2.get()))) - _CntrlBlk(__a, _CUDA_VSTD::forward<_Args>(__args)...); - shared_ptr<_Tp> __r; - __r.__ptr_ = __hold2.get()->get(); - __r.__cntrl_ = _CUDA_VSTD::addressof(*__hold2.release()); - __r.__enable_weak_this(__r.__ptr_, __r.__ptr_); - return __r; -} - -template -shared_ptr<_Tp>::~shared_ptr() -{ - if (__cntrl_) - __cntrl_->__release_shared(); -} - -template -inline -shared_ptr<_Tp>& -shared_ptr<_Tp>::operator=(const shared_ptr& __r) noexcept -{ - shared_ptr(__r).swap(*this); - return *this; -} - -template -template -inline -typename enable_if -< - is_convertible<_Yp*, typename shared_ptr<_Tp>::element_type*>::value, - shared_ptr<_Tp>& ->::type -shared_ptr<_Tp>::operator=(const shared_ptr<_Yp>& __r) noexcept -{ - shared_ptr(__r).swap(*this); - return *this; -} - -#ifndef _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - -template -inline -shared_ptr<_Tp>& -shared_ptr<_Tp>::operator=(shared_ptr&& __r) noexcept -{ - shared_ptr(_CUDA_VSTD::move(__r)).swap(*this); - return *this; -} - -template -template -inline -typename enable_if -< - is_convertible<_Yp*, typename shared_ptr<_Tp>::element_type*>::value, - shared_ptr<_Tp>& ->::type -shared_ptr<_Tp>::operator=(shared_ptr<_Yp>&& __r) -{ - shared_ptr(_CUDA_VSTD::move(__r)).swap(*this); - return *this; -} - -#if _CCCL_STD_VER <= 2014 || defined(_LIBCUDACXX_ENABLE_CXX17_REMOVED_AUTO_PTR) -template -template -inline -typename enable_if -< - !is_array<_Yp>::value && - is_convertible<_Yp*, typename shared_ptr<_Tp>::element_type*>::value, - shared_ptr<_Tp> ->::type& -shared_ptr<_Tp>::operator=(auto_ptr<_Yp>&& __r) -{ - shared_ptr(_CUDA_VSTD::move(__r)).swap(*this); - return *this; -} -#endif - -template -template -inline -typename enable_if -< - !is_array<_Yp>::value && - is_convertible::pointer, - typename shared_ptr<_Tp>::element_type*>::value, - shared_ptr<_Tp>& ->::type -shared_ptr<_Tp>::operator=(unique_ptr<_Yp, _Dp>&& __r) -{ - shared_ptr(_CUDA_VSTD::move(__r)).swap(*this); - return *this; -} - -#else // _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - -#if _CCCL_STD_VER <= 2014 || defined(_LIBCUDACXX_ENABLE_CXX17_REMOVED_AUTO_PTR) -template -template -inline _LIBCUDACXX_INLINE_VISIBILITY -typename enable_if -< - !is_array<_Yp>::value && - is_convertible<_Yp*, typename shared_ptr<_Tp>::element_type*>::value, - shared_ptr<_Tp>& ->::type -shared_ptr<_Tp>::operator=(auto_ptr<_Yp> __r) -{ - shared_ptr(__r).swap(*this); - return *this; -} -#endif - -template -template -inline _LIBCUDACXX_INLINE_VISIBILITY -typename enable_if -< - !is_array<_Yp>::value && - is_convertible::pointer, - typename shared_ptr<_Tp>::element_type*>::value, - shared_ptr<_Tp>& ->::type -shared_ptr<_Tp>::operator=(unique_ptr<_Yp, _Dp> __r) -{ - shared_ptr(_CUDA_VSTD::move(__r)).swap(*this); - return *this; -} - -#endif // _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - -template -inline -void -shared_ptr<_Tp>::swap(shared_ptr& __r) noexcept -{ - _CUDA_VSTD::swap(__ptr_, __r.__ptr_); - _CUDA_VSTD::swap(__cntrl_, __r.__cntrl_); -} - -template -inline -void -shared_ptr<_Tp>::reset() noexcept -{ - shared_ptr().swap(*this); -} - -template -template -inline -typename enable_if -< - is_convertible<_Yp*, typename shared_ptr<_Tp>::element_type*>::value, - void ->::type -shared_ptr<_Tp>::reset(_Yp* __p) -{ - shared_ptr(__p).swap(*this); -} - -template -template -inline -typename enable_if -< - is_convertible<_Yp*, typename shared_ptr<_Tp>::element_type*>::value, - void ->::type -shared_ptr<_Tp>::reset(_Yp* __p, _Dp __d) -{ - shared_ptr(__p, __d).swap(*this); -} - -template -template -inline -typename enable_if -< - is_convertible<_Yp*, typename shared_ptr<_Tp>::element_type*>::value, - void ->::type -shared_ptr<_Tp>::reset(_Yp* __p, _Dp __d, _Alloc __a) -{ - shared_ptr(__p, __d, __a).swap(*this); -} - -#ifndef _LIBCUDACXX_HAS_NO_VARIADICS - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -typename enable_if -< - !is_array<_Tp>::value, - shared_ptr<_Tp> ->::type -make_shared(_Args&& ...__args) -{ - return shared_ptr<_Tp>::make_shared(_CUDA_VSTD::forward<_Args>(__args)...); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -typename enable_if -< - !is_array<_Tp>::value, - shared_ptr<_Tp> ->::type -allocate_shared(const _Alloc& __a, _Args&& ...__args) -{ - return shared_ptr<_Tp>::allocate_shared(__a, _CUDA_VSTD::forward<_Args>(__args)...); -} - -#else // _LIBCUDACXX_HAS_NO_VARIADICS - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -shared_ptr<_Tp> -make_shared() -{ - return shared_ptr<_Tp>::make_shared(); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -shared_ptr<_Tp> -make_shared(_A0& __a0) -{ - return shared_ptr<_Tp>::make_shared(__a0); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -shared_ptr<_Tp> -make_shared(_A0& __a0, _A1& __a1) -{ - return shared_ptr<_Tp>::make_shared(__a0, __a1); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -shared_ptr<_Tp> -make_shared(_A0& __a0, _A1& __a1, _A2& __a2) -{ - return shared_ptr<_Tp>::make_shared(__a0, __a1, __a2); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -shared_ptr<_Tp> -allocate_shared(const _Alloc& __a) -{ - return shared_ptr<_Tp>::allocate_shared(__a); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -shared_ptr<_Tp> -allocate_shared(const _Alloc& __a, _A0& __a0) -{ - return shared_ptr<_Tp>::allocate_shared(__a, __a0); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -shared_ptr<_Tp> -allocate_shared(const _Alloc& __a, _A0& __a0, _A1& __a1) -{ - return shared_ptr<_Tp>::allocate_shared(__a, __a0, __a1); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -shared_ptr<_Tp> -allocate_shared(const _Alloc& __a, _A0& __a0, _A1& __a1, _A2& __a2) -{ - return shared_ptr<_Tp>::allocate_shared(__a, __a0, __a1, __a2); -} - -#endif // _LIBCUDACXX_HAS_NO_VARIADICS - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator==(const shared_ptr<_Tp>& __x, const shared_ptr<_Up>& __y) noexcept -{ - return __x.get() == __y.get(); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator!=(const shared_ptr<_Tp>& __x, const shared_ptr<_Up>& __y) noexcept -{ - return !(__x == __y); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator<(const shared_ptr<_Tp>& __x, const shared_ptr<_Up>& __y) noexcept -{ -#if _CCCL_STD_VER <= 2011 - typedef typename common_type<_Tp*, _Up*>::type _Vp; - return less<_Vp>()(__x.get(), __y.get()); -#else - return less<>()(__x.get(), __y.get()); -#endif - -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator>(const shared_ptr<_Tp>& __x, const shared_ptr<_Up>& __y) noexcept -{ - return __y < __x; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator<=(const shared_ptr<_Tp>& __x, const shared_ptr<_Up>& __y) noexcept -{ - return !(__y < __x); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator>=(const shared_ptr<_Tp>& __x, const shared_ptr<_Up>& __y) noexcept -{ - return !(__x < __y); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator==(const shared_ptr<_Tp>& __x, nullptr_t) noexcept -{ - return !__x; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator==(nullptr_t, const shared_ptr<_Tp>& __x) noexcept -{ - return !__x; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator!=(const shared_ptr<_Tp>& __x, nullptr_t) noexcept -{ - return static_cast(__x); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator!=(nullptr_t, const shared_ptr<_Tp>& __x) noexcept -{ - return static_cast(__x); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator<(const shared_ptr<_Tp>& __x, nullptr_t) noexcept -{ - return less<_Tp*>()(__x.get(), nullptr); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator<(nullptr_t, const shared_ptr<_Tp>& __x) noexcept -{ - return less<_Tp*>()(nullptr, __x.get()); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator>(const shared_ptr<_Tp>& __x, nullptr_t) noexcept -{ - return nullptr < __x; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator>(nullptr_t, const shared_ptr<_Tp>& __x) noexcept -{ - return __x < nullptr; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator<=(const shared_ptr<_Tp>& __x, nullptr_t) noexcept -{ - return !(nullptr < __x); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator<=(nullptr_t, const shared_ptr<_Tp>& __x) noexcept -{ - return !(__x < nullptr); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator>=(const shared_ptr<_Tp>& __x, nullptr_t) noexcept -{ - return !(__x < nullptr); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -operator>=(nullptr_t, const shared_ptr<_Tp>& __x) noexcept -{ - return !(nullptr < __x); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -void -swap(shared_ptr<_Tp>& __x, shared_ptr<_Tp>& __y) noexcept -{ - __x.swap(__y); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -typename enable_if -< - !is_array<_Tp>::value && !is_array<_Up>::value, - shared_ptr<_Tp> ->::type -static_pointer_cast(const shared_ptr<_Up>& __r) noexcept -{ - return shared_ptr<_Tp>(__r, static_cast<_Tp*>(__r.get())); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -typename enable_if -< - !is_array<_Tp>::value && !is_array<_Up>::value, - shared_ptr<_Tp> ->::type -dynamic_pointer_cast(const shared_ptr<_Up>& __r) noexcept -{ - _Tp* __p = dynamic_cast<_Tp*>(__r.get()); - return __p ? shared_ptr<_Tp>(__r, __p) : shared_ptr<_Tp>(); -} - -template -typename enable_if -< - is_array<_Tp>::value == is_array<_Up>::value, - shared_ptr<_Tp> ->::type -const_pointer_cast(const shared_ptr<_Up>& __r) noexcept -{ - typedef typename remove_extent<_Tp>::type _RTp; - return shared_ptr<_Tp>(__r, const_cast<_RTp*>(__r.get())); -} - -#ifndef _LIBCUDACXX_NO_RTTI - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_Dp* -get_deleter(const shared_ptr<_Tp>& __p) noexcept -{ - return __p.template __get_deleter<_Dp>(); -} - -#endif // _LIBCUDACXX_NO_RTTI - -template -class _LIBCUDACXX_TEMPLATE_VIS weak_ptr -{ -public: - typedef _Tp element_type; -private: - element_type* __ptr_; - __shared_weak_count* __cntrl_; - -public: - _LIBCUDACXX_INLINE_VISIBILITY - constexpr weak_ptr() noexcept; - template _LIBCUDACXX_INLINE_VISIBILITY weak_ptr(shared_ptr<_Yp> const& __r, - typename enable_if::value, __nat*>::type = 0) - noexcept; - _LIBCUDACXX_INLINE_VISIBILITY - weak_ptr(weak_ptr const& __r) noexcept; - template _LIBCUDACXX_INLINE_VISIBILITY weak_ptr(weak_ptr<_Yp> const& __r, - typename enable_if::value, __nat*>::type = 0) - noexcept; - -#ifndef _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - _LIBCUDACXX_INLINE_VISIBILITY - weak_ptr(weak_ptr&& __r) noexcept; - template _LIBCUDACXX_INLINE_VISIBILITY weak_ptr(weak_ptr<_Yp>&& __r, - typename enable_if::value, __nat*>::type = 0) - noexcept; -#endif // _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - ~weak_ptr(); - - _LIBCUDACXX_INLINE_VISIBILITY - weak_ptr& operator=(weak_ptr const& __r) noexcept; - template - typename enable_if - < - is_convertible<_Yp*, element_type*>::value, - weak_ptr& - >::type - _LIBCUDACXX_INLINE_VISIBILITY - operator=(weak_ptr<_Yp> const& __r) noexcept; - -#ifndef _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - - _LIBCUDACXX_INLINE_VISIBILITY - weak_ptr& operator=(weak_ptr&& __r) noexcept; - template - typename enable_if - < - is_convertible<_Yp*, element_type*>::value, - weak_ptr& - >::type - _LIBCUDACXX_INLINE_VISIBILITY - operator=(weak_ptr<_Yp>&& __r) noexcept; - -#endif // _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - - template - typename enable_if - < - is_convertible<_Yp*, element_type*>::value, - weak_ptr& - >::type - _LIBCUDACXX_INLINE_VISIBILITY - operator=(shared_ptr<_Yp> const& __r) noexcept; - - _LIBCUDACXX_INLINE_VISIBILITY - void swap(weak_ptr& __r) noexcept; - _LIBCUDACXX_INLINE_VISIBILITY - void reset() noexcept; - - _LIBCUDACXX_INLINE_VISIBILITY - long use_count() const noexcept - {return __cntrl_ ? __cntrl_->use_count() : 0;} - _LIBCUDACXX_INLINE_VISIBILITY - bool expired() const noexcept - {return __cntrl_ == 0 || __cntrl_->use_count() == 0;} - shared_ptr<_Tp> lock() const noexcept; - template - _LIBCUDACXX_INLINE_VISIBILITY - bool owner_before(const shared_ptr<_Up>& __r) const noexcept - {return __cntrl_ < __r.__cntrl_;} - template - _LIBCUDACXX_INLINE_VISIBILITY - bool owner_before(const weak_ptr<_Up>& __r) const noexcept - {return __cntrl_ < __r.__cntrl_;} - - template friend class _LIBCUDACXX_TEMPLATE_VIS weak_ptr; - template friend class _LIBCUDACXX_TEMPLATE_VIS shared_ptr; -}; - -template -inline constexpr -weak_ptr<_Tp>::weak_ptr() noexcept - : __ptr_(0), - __cntrl_(0) -{ -} - -template -inline -weak_ptr<_Tp>::weak_ptr(weak_ptr const& __r) noexcept - : __ptr_(__r.__ptr_), - __cntrl_(__r.__cntrl_) -{ - if (__cntrl_) - __cntrl_->__add_weak(); -} - -template -template -inline -weak_ptr<_Tp>::weak_ptr(shared_ptr<_Yp> const& __r, - typename enable_if::value, __nat*>::type) - noexcept - : __ptr_(__r.__ptr_), - __cntrl_(__r.__cntrl_) -{ - if (__cntrl_) - __cntrl_->__add_weak(); -} - -template -template -inline -weak_ptr<_Tp>::weak_ptr(weak_ptr<_Yp> const& __r, - typename enable_if::value, __nat*>::type) - noexcept - : __ptr_(__r.__ptr_), - __cntrl_(__r.__cntrl_) -{ - if (__cntrl_) - __cntrl_->__add_weak(); -} - -#ifndef _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - -template -inline -weak_ptr<_Tp>::weak_ptr(weak_ptr&& __r) noexcept - : __ptr_(__r.__ptr_), - __cntrl_(__r.__cntrl_) -{ - __r.__ptr_ = 0; - __r.__cntrl_ = 0; -} - -template -template -inline -weak_ptr<_Tp>::weak_ptr(weak_ptr<_Yp>&& __r, - typename enable_if::value, __nat*>::type) - noexcept - : __ptr_(__r.__ptr_), - __cntrl_(__r.__cntrl_) -{ - __r.__ptr_ = 0; - __r.__cntrl_ = 0; -} - -#endif // _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - -template -weak_ptr<_Tp>::~weak_ptr() -{ - if (__cntrl_) - __cntrl_->__release_weak(); -} - -template -inline -weak_ptr<_Tp>& -weak_ptr<_Tp>::operator=(weak_ptr const& __r) noexcept -{ - weak_ptr(__r).swap(*this); - return *this; -} - -template -template -inline -typename enable_if -< - is_convertible<_Yp*, _Tp*>::value, - weak_ptr<_Tp>& ->::type -weak_ptr<_Tp>::operator=(weak_ptr<_Yp> const& __r) noexcept -{ - weak_ptr(__r).swap(*this); - return *this; -} - -#ifndef _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - -template -inline -weak_ptr<_Tp>& -weak_ptr<_Tp>::operator=(weak_ptr&& __r) noexcept -{ - weak_ptr(_CUDA_VSTD::move(__r)).swap(*this); - return *this; -} - -template -template -inline -typename enable_if -< - is_convertible<_Yp*, _Tp*>::value, - weak_ptr<_Tp>& ->::type -weak_ptr<_Tp>::operator=(weak_ptr<_Yp>&& __r) noexcept -{ - weak_ptr(_CUDA_VSTD::move(__r)).swap(*this); - return *this; -} - -#endif // _LIBCUDACXX_HAS_NO_RVALUE_REFERENCES - -template -template -inline -typename enable_if -< - is_convertible<_Yp*, _Tp*>::value, - weak_ptr<_Tp>& ->::type -weak_ptr<_Tp>::operator=(shared_ptr<_Yp> const& __r) noexcept -{ - weak_ptr(__r).swap(*this); - return *this; -} - -template -inline -void -weak_ptr<_Tp>::swap(weak_ptr& __r) noexcept -{ - _CUDA_VSTD::swap(__ptr_, __r.__ptr_); - _CUDA_VSTD::swap(__cntrl_, __r.__cntrl_); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -void -swap(weak_ptr<_Tp>& __x, weak_ptr<_Tp>& __y) noexcept -{ - __x.swap(__y); -} - -template -inline -void -weak_ptr<_Tp>::reset() noexcept -{ - weak_ptr().swap(*this); -} - -template -template -shared_ptr<_Tp>::shared_ptr(const weak_ptr<_Yp>& __r, - typename enable_if::value, __nat>::type) - : __ptr_(__r.__ptr_), - __cntrl_(__r.__cntrl_ ? __r.__cntrl_->lock() : __r.__cntrl_) -{ - if (__cntrl_ == 0) - __throw_bad_weak_ptr(); -} - -template -shared_ptr<_Tp> -weak_ptr<_Tp>::lock() const noexcept -{ - shared_ptr<_Tp> __r; - __r.__cntrl_ = __cntrl_ ? __cntrl_->lock() : __cntrl_; - if (__r.__cntrl_) - __r.__ptr_ = __ptr_; - return __r; -} - -#if _CCCL_STD_VER > 2014 -template struct owner_less; -#else -template struct owner_less; -#endif - -template -struct _LIBCUDACXX_TEMPLATE_VIS owner_less > - : __binary_function, shared_ptr<_Tp>, bool> -{ - typedef bool result_type; - _LIBCUDACXX_INLINE_VISIBILITY - bool operator()(shared_ptr<_Tp> const& __x, shared_ptr<_Tp> const& __y) const noexcept - {return __x.owner_before(__y);} - _LIBCUDACXX_INLINE_VISIBILITY - bool operator()(shared_ptr<_Tp> const& __x, weak_ptr<_Tp> const& __y) const noexcept - {return __x.owner_before(__y);} - _LIBCUDACXX_INLINE_VISIBILITY - bool operator()( weak_ptr<_Tp> const& __x, shared_ptr<_Tp> const& __y) const noexcept - {return __x.owner_before(__y);} -}; - -template -struct _LIBCUDACXX_TEMPLATE_VIS owner_less > - : __binary_function, weak_ptr<_Tp>, bool> -{ - typedef bool result_type; - _LIBCUDACXX_INLINE_VISIBILITY - bool operator()( weak_ptr<_Tp> const& __x, weak_ptr<_Tp> const& __y) const noexcept - {return __x.owner_before(__y);} - _LIBCUDACXX_INLINE_VISIBILITY - bool operator()(shared_ptr<_Tp> const& __x, weak_ptr<_Tp> const& __y) const noexcept - {return __x.owner_before(__y);} - _LIBCUDACXX_INLINE_VISIBILITY - bool operator()( weak_ptr<_Tp> const& __x, shared_ptr<_Tp> const& __y) const noexcept - {return __x.owner_before(__y);} -}; - -#if _CCCL_STD_VER > 2014 -template <> -struct _LIBCUDACXX_TEMPLATE_VIS owner_less -{ - template - _LIBCUDACXX_INLINE_VISIBILITY - bool operator()( shared_ptr<_Tp> const& __x, shared_ptr<_Up> const& __y) const noexcept - {return __x.owner_before(__y);} - template - _LIBCUDACXX_INLINE_VISIBILITY - bool operator()( shared_ptr<_Tp> const& __x, weak_ptr<_Up> const& __y) const noexcept - {return __x.owner_before(__y);} - template - _LIBCUDACXX_INLINE_VISIBILITY - bool operator()( weak_ptr<_Tp> const& __x, shared_ptr<_Up> const& __y) const noexcept - {return __x.owner_before(__y);} - template - _LIBCUDACXX_INLINE_VISIBILITY - bool operator()( weak_ptr<_Tp> const& __x, weak_ptr<_Up> const& __y) const noexcept - {return __x.owner_before(__y);} - typedef void is_transparent; -}; -#endif - -template -class _LIBCUDACXX_TEMPLATE_VIS enable_shared_from_this -{ - mutable weak_ptr<_Tp> __weak_this_; -protected: - _LIBCUDACXX_INLINE_VISIBILITY constexpr - enable_shared_from_this() noexcept {} - _LIBCUDACXX_INLINE_VISIBILITY - enable_shared_from_this(enable_shared_from_this const&) noexcept {} - _LIBCUDACXX_INLINE_VISIBILITY - enable_shared_from_this& operator=(enable_shared_from_this const&) noexcept - {return *this;} - _LIBCUDACXX_INLINE_VISIBILITY - ~enable_shared_from_this() {} -public: - _LIBCUDACXX_INLINE_VISIBILITY - shared_ptr<_Tp> shared_from_this() - {return shared_ptr<_Tp>(__weak_this_);} - _LIBCUDACXX_INLINE_VISIBILITY - shared_ptr<_Tp const> shared_from_this() const - {return shared_ptr(__weak_this_);} - -#if _CCCL_STD_VER > 2014 - _LIBCUDACXX_INLINE_VISIBILITY - weak_ptr<_Tp> weak_from_this() noexcept - { return __weak_this_; } - - _LIBCUDACXX_INLINE_VISIBILITY - weak_ptr weak_from_this() const noexcept - { return __weak_this_; } -#endif // _CCCL_STD_VER > 2014 - - template friend class shared_ptr; -}; - -template -struct _LIBCUDACXX_TEMPLATE_VIS hash > -{ - typedef shared_ptr<_Tp> argument_type; - typedef size_t result_type; - - _LIBCUDACXX_INLINE_VISIBILITY - result_type operator()(const argument_type& __ptr) const noexcept - { - return hash<_Tp*>()(__ptr.get()); - } -}; - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -basic_ostream<_CharT, _Traits>& -operator<<(basic_ostream<_CharT, _Traits>& __os, shared_ptr<_Yp> const& __p); - - -#if !defined(_LIBCUDACXX_HAS_NO_ATOMIC_HEADER) - -class _LIBCUDACXX_TYPE_VIS __sp_mut -{ - void* __lx; -public: - void lock() noexcept; - void unlock() noexcept; - -private: - constexpr __sp_mut(void*) noexcept; - __sp_mut(const __sp_mut&); - __sp_mut& operator=(const __sp_mut&); - - friend _LIBCUDACXX_FUNC_VIS __sp_mut& __get_sp_mut(const void*); -}; - -_LIBCUDACXX_FUNC_VIS _LIBCUDACXX_AVAILABILITY_ATOMIC_SHARED_PTR -__sp_mut& __get_sp_mut(const void*); - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -bool -atomic_is_lock_free(const shared_ptr<_Tp>*) -{ - return false; -} - -template -_LIBCUDACXX_AVAILABILITY_ATOMIC_SHARED_PTR -shared_ptr<_Tp> -atomic_load(const shared_ptr<_Tp>* __p) -{ - __sp_mut& __m = __get_sp_mut(__p); - __m.lock(); - shared_ptr<_Tp> __q = *__p; - __m.unlock(); - return __q; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_LIBCUDACXX_AVAILABILITY_ATOMIC_SHARED_PTR -shared_ptr<_Tp> -atomic_load_explicit(const shared_ptr<_Tp>* __p, memory_order) -{ - return atomic_load(__p); -} - -template -_LIBCUDACXX_AVAILABILITY_ATOMIC_SHARED_PTR -void -atomic_store(shared_ptr<_Tp>* __p, shared_ptr<_Tp> __r) -{ - __sp_mut& __m = __get_sp_mut(__p); - __m.lock(); - __p->swap(__r); - __m.unlock(); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_LIBCUDACXX_AVAILABILITY_ATOMIC_SHARED_PTR -void -atomic_store_explicit(shared_ptr<_Tp>* __p, shared_ptr<_Tp> __r, memory_order) -{ - atomic_store(__p, __r); -} - -template -_LIBCUDACXX_AVAILABILITY_ATOMIC_SHARED_PTR -shared_ptr<_Tp> -atomic_exchange(shared_ptr<_Tp>* __p, shared_ptr<_Tp> __r) -{ - __sp_mut& __m = __get_sp_mut(__p); - __m.lock(); - __p->swap(__r); - __m.unlock(); - return __r; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_LIBCUDACXX_AVAILABILITY_ATOMIC_SHARED_PTR -shared_ptr<_Tp> -atomic_exchange_explicit(shared_ptr<_Tp>* __p, shared_ptr<_Tp> __r, memory_order) -{ - return atomic_exchange(__p, __r); -} - -template -_LIBCUDACXX_AVAILABILITY_ATOMIC_SHARED_PTR -bool -atomic_compare_exchange_strong(shared_ptr<_Tp>* __p, shared_ptr<_Tp>* __v, shared_ptr<_Tp> __w) -{ - shared_ptr<_Tp> __temp; - __sp_mut& __m = __get_sp_mut(__p); - __m.lock(); - if (__p->__owner_equivalent(*__v)) - { - _CUDA_VSTD::swap(__temp, *__p); - *__p = __w; - __m.unlock(); - return true; - } - _CUDA_VSTD::swap(__temp, *__v); - *__v = *__p; - __m.unlock(); - return false; -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_LIBCUDACXX_AVAILABILITY_ATOMIC_SHARED_PTR -bool -atomic_compare_exchange_weak(shared_ptr<_Tp>* __p, shared_ptr<_Tp>* __v, shared_ptr<_Tp> __w) -{ - return atomic_compare_exchange_strong(__p, __v, __w); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_LIBCUDACXX_AVAILABILITY_ATOMIC_SHARED_PTR -bool -atomic_compare_exchange_strong_explicit(shared_ptr<_Tp>* __p, shared_ptr<_Tp>* __v, - shared_ptr<_Tp> __w, memory_order, memory_order) -{ - return atomic_compare_exchange_strong(__p, __v, __w); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_LIBCUDACXX_AVAILABILITY_ATOMIC_SHARED_PTR -bool -atomic_compare_exchange_weak_explicit(shared_ptr<_Tp>* __p, shared_ptr<_Tp>* __v, - shared_ptr<_Tp> __w, memory_order, memory_order) -{ - return atomic_compare_exchange_weak(__p, __v, __w); -} - -#endif // !defined(_LIBCUDACXX_HAS_NO_ATOMIC_HEADER) - -//enum class -#if defined(_LIBCUDACXX_ABI_POINTER_SAFETY_ENUM_TYPE) -enum class pointer_safety : unsigned char { - relaxed, - preferred, - strict -}; -#else -struct _LIBCUDACXX_TYPE_VIS pointer_safety -{ - enum __lx - { - relaxed, - preferred, - strict - }; - - __lx __v_; - - _LIBCUDACXX_INLINE_VISIBILITY - pointer_safety() : __v_() {} - - _LIBCUDACXX_INLINE_VISIBILITY - pointer_safety(__lx __v) : __v_(__v) {} - _LIBCUDACXX_INLINE_VISIBILITY - operator int() const {return __v_;} -}; -#endif - -#if !defined(_LIBCUDACXX_ABI_POINTER_SAFETY_ENUM_TYPE) && \ - defined(_LIBCUDACXX_BUILDING_LIBRARY) -_LIBCUDACXX_FUNC_VIS pointer_safety get_pointer_safety() noexcept; -#else -// This function is only offered in C++03 under ABI v1. -inline _LIBCUDACXX_INLINE_VISIBILITY -pointer_safety get_pointer_safety() noexcept { - return pointer_safety::relaxed; -} -#endif - - -_LIBCUDACXX_FUNC_VIS void declare_reachable(void* __p); -_LIBCUDACXX_FUNC_VIS void declare_no_pointers(char* __p, size_t __n); -_LIBCUDACXX_FUNC_VIS void undeclare_no_pointers(char* __p, size_t __n); -_LIBCUDACXX_FUNC_VIS void* __undeclare_reachable(void* __p); - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -_Tp* -undeclare_reachable(_Tp* __p) -{ - return static_cast<_Tp*>(__undeclare_reachable(__p)); -} - -// --- Helper for container swap -- -template -inline _LIBCUDACXX_INLINE_VISIBILITY -void __swap_allocator(_Alloc & __a1, _Alloc & __a2) -#if _CCCL_STD_VER >= 2014 - noexcept -#else - noexcept(__is_nothrow_swappable<_Alloc>::value) -#endif -{ - __swap_allocator(__a1, __a2, - integral_constant::propagate_on_container_swap::value>()); -} - -template -_LIBCUDACXX_INLINE_VISIBILITY -void __swap_allocator(_Alloc & __a1, _Alloc & __a2, true_type) -#if _CCCL_STD_VER >= 2014 - noexcept -#else - noexcept(__is_nothrow_swappable<_Alloc>::value) -#endif -{ - using _CUDA_VSTD::swap; - swap(__a1, __a2); -} - -template -inline _LIBCUDACXX_INLINE_VISIBILITY -void __swap_allocator(_Alloc &, _Alloc &, false_type) noexcept {} - -template > -struct __noexcept_move_assign_container : public integral_constant 2014 - || _Traits::is_always_equal::value -#else - && is_nothrow_move_assignable<_Alloc>::value -#endif - > {}; - - -#ifndef _LIBCUDACXX_HAS_NO_VARIADICS -template -struct __temp_value { - typedef allocator_traits<_Alloc> _Traits; - - typename aligned_storage::type __v; - _Alloc &__a; - - _Tp *__addr() { return reinterpret_cast<_Tp *>(addressof(__v)); } - _Tp & get() { return *__addr(); } - - template - _LIBCUDACXX_NO_CFI - __temp_value(_Alloc &__alloc, _Args&& ... __args) : __a(__alloc) { - _Traits::construct(__a, reinterpret_cast<_Tp*>(addressof(__v)), - _CUDA_VSTD::forward<_Args>(__args)...); - } - - ~__temp_value() { _Traits::destroy(__a, __addr()); } - }; -#endif - -// __builtin_new_allocator -- A non-templated helper for allocating and -// deallocating memory using __builtin_operator_new and -// __builtin_operator_delete. It should be used in preference to -// `std::allocator` to avoid additional instantiations. -struct __builtin_new_allocator { - struct __builtin_new_deleter { - typedef void* pointer_type; - - constexpr explicit __builtin_new_deleter(size_t __size, size_t __align) - : __size_(__size), __align_(__align) {} - - void operator()(void* p) const noexcept { - std::__libcpp_deallocate(p, __size_, __align_); - } - - private: - size_t __size_; - size_t __align_; - }; - - typedef unique_ptr __holder_t; - - static __holder_t __allocate_bytes(size_t __s, size_t __align) { - return __holder_t(std::__libcpp_allocate(__s, __align), - __builtin_new_deleter(__s, __align)); - } - - static void __deallocate_bytes(void* __p, size_t __s, - size_t __align) noexcept { - std::__libcpp_deallocate(__p, __s, __align); - } - - template - _LIBCUDACXX_NODEBUG_TYPE _LIBCUDACXX_ALWAYS_INLINE - static __holder_t __allocate_type(size_t __n) { - return __allocate_bytes(__n * sizeof(_Tp), _LIBCUDACXX_ALIGNOF(_Tp)); - } - - template - _LIBCUDACXX_NODEBUG_TYPE _LIBCUDACXX_ALWAYS_INLINE - static void __deallocate_type(void* __p, size_t __n) noexcept { - __deallocate_bytes(__p, __n * sizeof(_Tp), _LIBCUDACXX_ALIGNOF(_Tp)); - } -}; - - -_LIBCUDACXX_END_NAMESPACE_STD - -#endif // __cuda_std__ - -#include //__cuda_std__ +// [memory.syn] +#ifndef _LIBCUDACXX_HAS_NO_SPACESHIP_OPERATOR +# include +#endif // !_LIBCUDACXX_HAS_NO_SPACESHIP_OPERATOR #if defined(_LIBCUDACXX_HAS_PARALLEL_ALGORITHMS) && _CCCL_STD_VER >= 2017 -# include <__pstl_memory> +# include <__pstl_memory> #endif -#endif // _LIBCUDACXX_MEMORY +#endif // _LIBCUDACXX_MEMORY diff --git a/libcudacxx/test/libcudacxx/std/containers/sequences/array/array.cons/implicit_copy.pass.cpp b/libcudacxx/test/libcudacxx/std/containers/sequences/array/array.cons/implicit_copy.pass.cpp index 23a3d3edd8b..3854fb2e04f 100644 --- a/libcudacxx/test/libcudacxx/std/containers/sequences/array/array.cons/implicit_copy.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/containers/sequences/array/array.cons/implicit_copy.pass.cpp @@ -106,7 +106,7 @@ __host__ __device__ TEST_CONSTEXPR_CXX14_NOT_MSVC_2017 bool tests() // NVCC believes `copy = array` accesses uninitialized memory #if defined(TEST_COMPILER_NVCC) || defined(TEST_COMPILER_NVRTC) - if (!cuda::std::__libcpp_is_constant_evaluated()) + if (!TEST_IS_CONSTANT_EVALUATED()) #endif // TEST_COMPILER_NVCC { typedef cuda::std::array Array; @@ -118,7 +118,7 @@ __host__ __device__ TEST_CONSTEXPR_CXX14_NOT_MSVC_2017 bool tests() } // NVCC believes `copy = array` accesses uninitialized memory #if defined(TEST_COMPILER_NVCC) || defined(TEST_COMPILER_NVRTC) - if (!cuda::std::__libcpp_is_constant_evaluated()) + if (!TEST_IS_CONSTANT_EVALUATED()) #endif // TEST_COMPILER_NVCC { typedef cuda::std::array Array; diff --git a/libcudacxx/test/libcudacxx/std/containers/sequences/array/iterators.pass.cpp b/libcudacxx/test/libcudacxx/std/containers/sequences/array/iterators.pass.cpp index 18af754392b..1a205d21cf4 100644 --- a/libcudacxx/test/libcudacxx/std/containers/sequences/array/iterators.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/containers/sequences/array/iterators.pass.cpp @@ -166,7 +166,7 @@ __host__ __device__ assert(c.begin() == cuda::std::begin(c)); assert(c.cbegin() == cuda::std::cbegin(c)); # if TEST_STD_VER < 2017 - if (!cuda::std::__libcpp_is_constant_evaluated()) + if (!TEST_IS_CONSTANT_EVALUATED()) # endif // TEST_STD_VER < 2017 { assert(c.rbegin() == cuda::std::rbegin(c)); @@ -175,7 +175,7 @@ __host__ __device__ assert(c.end() == cuda::std::end(c)); assert(c.cend() == cuda::std::cend(c)); # if TEST_STD_VER < 2017 - if (!cuda::std::__libcpp_is_constant_evaluated()) + if (!TEST_IS_CONSTANT_EVALUATED()) # endif // TEST_STD_VER < 2017 { assert(c.rend() == cuda::std::rend(c)); @@ -185,7 +185,7 @@ __host__ __device__ assert(cuda::std::begin(c) != cuda::std::end(c)); assert(cuda::std::cbegin(c) != cuda::std::cend(c)); # if TEST_STD_VER < 2017 - if (!cuda::std::__libcpp_is_constant_evaluated()) + if (!TEST_IS_CONSTANT_EVALUATED()) # endif // TEST_STD_VER < 2017 { assert(cuda::std::rbegin(c) != cuda::std::rend(c)); @@ -212,7 +212,7 @@ __host__ __device__ assert(cii == nullptr); # endif // TEST_COMPILER_CUDACC_BELOW_11_3 // This breaks NVCCs constexpr evaluator - if (!cuda::std::__libcpp_is_constant_evaluated()) + if (!TEST_IS_CONSTANT_EVALUATED()) { assert(!(ii1 < cii)); assert(!(cii < ii1)); @@ -231,7 +231,7 @@ __host__ __device__ assert(c.begin() == cuda::std::begin(c)); assert(c.cbegin() == cuda::std::cbegin(c)); # if TEST_STD_VER < 2017 - if (!cuda::std::__libcpp_is_constant_evaluated()) + if (!TEST_IS_CONSTANT_EVALUATED()) # endif // TEST_STD_VER < 2017 { assert(c.rbegin() == cuda::std::rbegin(c)); @@ -240,7 +240,7 @@ __host__ __device__ assert(c.end() == cuda::std::end(c)); assert(c.cend() == cuda::std::cend(c)); # if TEST_STD_VER < 2017 - if (!cuda::std::__libcpp_is_constant_evaluated()) + if (!TEST_IS_CONSTANT_EVALUATED()) # endif // TEST_STD_VER < 2017 { assert(c.rend() == cuda::std::rend(c)); @@ -250,7 +250,7 @@ __host__ __device__ assert(cuda::std::begin(c) == cuda::std::end(c)); assert(cuda::std::cbegin(c) == cuda::std::cend(c)); # if TEST_STD_VER < 2017 - if (!cuda::std::__libcpp_is_constant_evaluated()) + if (!TEST_IS_CONSTANT_EVALUATED()) # endif // TEST_STD_VER < 2017 { assert(cuda::std::rbegin(c) == cuda::std::rend(c)); diff --git a/libcudacxx/test/libcudacxx/std/utilities/function.objects/func.bind_front/bind_front.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/function.objects/func.bind_front/bind_front.pass.cpp index 3e0832595df..96151050747 100644 --- a/libcudacxx/test/libcudacxx/std/utilities/function.objects/func.bind_front/bind_front.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/utilities/function.objects/func.bind_front/bind_front.pass.cpp @@ -207,7 +207,7 @@ __host__ __device__ constexpr bool test() // Make sure we don't treat cuda::std::reference_wrapper specially. #if TEST_STD_VER > 2017 # if defined(TEST_COMPILER_NVRTC) // reference_wrapper requires `addressof` which is currently not supported with nvrtc - if (!cuda::std::__libcpp_is_constant_evaluated()) + if (!TEST_IS_CONSTANT_EVALUATED()) # endif // TEST_COMPILER_NVRTC { auto add = [](cuda::std::reference_wrapper a, cuda::std::reference_wrapper b) { diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/destroy.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/destroy.pass.cpp index 9417053999b..87d40b345c8 100644 --- a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/destroy.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/destroy.pass.cpp @@ -99,7 +99,9 @@ struct CountDestructor __host__ __device__ TEST_CONSTEXPR_CXX20 bool test() { - if (!cuda::std::__libcpp_is_constant_evaluated()) +#if TEST_STD_VER >= 2020 + if (!TEST_IS_CONSTANT_EVALUATED()) +#endif // TEST_STD_VER >= 2020 { using Alloc = NoDestroy; int destructors = 0; @@ -115,7 +117,7 @@ __host__ __device__ TEST_CONSTEXPR_CXX20 bool test() cuda::std::allocator_traits::deallocate(alloc, pool, 1); } #if !defined(TEST_COMPILER_MSVC) && TEST_STD_VER >= 2020 // incomplete type not allowed - if (!cuda::std::__libcpp_is_constant_evaluated()) + if (!TEST_IS_CONSTANT_EVALUATED()) { typedef IncompleteHolder* T; typedef NoDestroy Alloc; diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/smartptr/unique.ptr/README.TXT b/libcudacxx/test/libcudacxx/std/utilities/memory/smartptr/unique.ptr/README.TXT new file mode 100644 index 00000000000..20f77f61827 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/utilities/memory/smartptr/unique.ptr/README.TXT @@ -0,0 +1,16 @@ +Test Naming and Directory Structure +=================================== + +The directory structure for the unique_ptr class templates differs from the +normal test directory naming conventions (e.g. matching the stable name in the standard). + +Instead of having a [unique.ptr.single] and [unique.ptr.runtime] directory, +each containing their own tests, a single directory, "unique.ptr.class", +contains both sets of tests. + +This allows the common behavior of the two unique_ptr specializations to be +tested in the same place without duplication. + +Tests specific to [unique.ptr.single] have the suffix ".single.pass.cpp" +and those specific to [unique.ptr.runtime] are named "*.runtime.pass.cpp". +Tests for both specializations are named normally. diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/smartptr/unique.ptr/unique.ptr.class/pointer_type.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/smartptr/unique.ptr/unique.ptr.class/pointer_type.pass.cpp new file mode 100644 index 00000000000..f5ab1de0e2c --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/utilities/memory/smartptr/unique.ptr/unique.ptr.class/pointer_type.pass.cpp @@ -0,0 +1,84 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// +// + +// unique_ptr + +// Test unique_ptr::pointer type + +#include +#include + +#include "test_macros.h" + +struct Deleter +{ + struct pointer + {}; +}; + +#if !defined(TEST_COMPILER_GCC) && !defined(TEST_COMPILER_MSVC) +struct D2 +{ +private: + typedef void pointer; +}; +#endif // !TEST_COMPILER_GCC && !TEST_COMPILER_MSVC + +#ifndef TEST_COMPILER_NVRTC // A class static data member with non-const type is considered a host variable +struct D3 +{ + static long pointer; +}; +#endif // !TEST_COMPILER_NVRTC + +template +__host__ __device__ TEST_CONSTEXPR_CXX23 void test_basic() +{ + typedef typename cuda::std::conditional::type VT; + { + typedef cuda::std::unique_ptr P; + static_assert((cuda::std::is_same::value), ""); + } + { + typedef cuda::std::unique_ptr P; + static_assert((cuda::std::is_same::value), ""); + } +#if !defined(TEST_COMPILER_GCC) && !defined(TEST_COMPILER_MSVC) + { + typedef cuda::std::unique_ptr P; + static_assert(cuda::std::is_same::value, ""); + } +#endif // !TEST_COMPILER_GCC && !TEST_COMPILER_MSVC +#ifndef TEST_COMPILER_NVRTC + { + typedef cuda::std::unique_ptr P; + static_assert(cuda::std::is_same::value, ""); + } +#endif // !TEST_COMPILER_NVRTC +} + +__host__ __device__ TEST_CONSTEXPR_CXX23 bool test() +{ + test_basic(); + test_basic(); + + return true; +} + +int main(int, char**) +{ + test(); +#if TEST_STD_VER >= 2023 + static_assert(test()); +#endif // TEST_STD_VER >= 2023 + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/smartptr/unique.ptr/unique.ptr.class/unique.ptr.asgn/move.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/smartptr/unique.ptr/unique.ptr.class/unique.ptr.asgn/move.pass.cpp new file mode 100644 index 00000000000..2729f512b93 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/utilities/memory/smartptr/unique.ptr/unique.ptr.class/unique.ptr.asgn/move.pass.cpp @@ -0,0 +1,176 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// +// UNSUPPORTED: c++03 + +// Self assignement post-conditions are tested. +// ADDITIONAL_COMPILE_FLAGS: -Wno-self-move + +// + +// unique_ptr + +// Test unique_ptr move assignment + +// test move assignment. Should only require a MoveConstructible deleter, or if +// deleter is a reference, not even that. + +#include +#include +#include + +#include "deleter_types.h" +#include "test_macros.h" +#include "unique_ptr_test_helper.h" + +struct GenericDeleter +{ + __host__ __device__ void operator()(void*) const; +}; + +template +__host__ __device__ TEST_CONSTEXPR_CXX23 void test_basic() +{ + typedef typename cuda::std::conditional::type VT; + const int expect_alive = IsArray ? 5 : 1; + { + cuda::std::unique_ptr s1(newValue(expect_alive)); + A* p = s1.get(); + cuda::std::unique_ptr s2(newValue(expect_alive)); + if (!TEST_IS_CONSTANT_EVALUATED_CXX23()) + { + assert(A_count == (expect_alive * 2)); + } + s2 = cuda::std::move(s1); + if (!TEST_IS_CONSTANT_EVALUATED_CXX23()) + { + assert(A_count == expect_alive); + } + assert(s2.get() == p); + assert(s1.get() == 0); + } + if (!TEST_IS_CONSTANT_EVALUATED_CXX23()) + { + assert(A_count == 0); + } + { + cuda::std::unique_ptr> s1(newValue(expect_alive), Deleter(5)); + A* p = s1.get(); + cuda::std::unique_ptr> s2(newValue(expect_alive)); + if (!TEST_IS_CONSTANT_EVALUATED_CXX23()) + { + assert(A_count == (expect_alive * 2)); + } + s2 = cuda::std::move(s1); + assert(s2.get() == p); + assert(s1.get() == 0); + if (!TEST_IS_CONSTANT_EVALUATED_CXX23()) + { + assert(A_count == expect_alive); + } + assert(s2.get_deleter().state() == 5); + assert(s1.get_deleter().state() == 0); + } + if (!TEST_IS_CONSTANT_EVALUATED_CXX23()) + { + assert(A_count == 0); + } + { + CDeleter d1(5); + cuda::std::unique_ptr&> s1(newValue(expect_alive), d1); + A* p = s1.get(); + CDeleter d2(6); + cuda::std::unique_ptr&> s2(newValue(expect_alive), d2); + s2 = cuda::std::move(s1); + assert(s2.get() == p); + assert(s1.get() == 0); + if (!TEST_IS_CONSTANT_EVALUATED_CXX23()) + { + assert(A_count == expect_alive); + } + assert(d1.state() == 5); + assert(d2.state() == 5); + } + if (!TEST_IS_CONSTANT_EVALUATED_CXX23()) + { + assert(A_count == 0); + } + { + cuda::std::unique_ptr s(newValue(expect_alive)); + A* p = s.get(); + s = cuda::std::move(s); + if (!TEST_IS_CONSTANT_EVALUATED_CXX23()) + { + assert(A_count == expect_alive); + } + assert(s.get() == p); + } + if (!TEST_IS_CONSTANT_EVALUATED_CXX23()) + { + assert(A_count == 0); + } +} + +template +__host__ __device__ TEST_CONSTEXPR_CXX23 void test_sfinae() +{ + typedef typename cuda::std::conditional::type VT; + { + typedef cuda::std::unique_ptr U; + static_assert(!cuda::std::is_assignable::value, ""); + static_assert(!cuda::std::is_assignable::value, ""); + static_assert(!cuda::std::is_assignable::value, ""); + static_assert(cuda::std::is_nothrow_assignable::value, ""); + } + { + typedef cuda::std::unique_ptr U; + static_assert(!cuda::std::is_assignable::value, ""); + static_assert(!cuda::std::is_assignable::value, ""); + static_assert(!cuda::std::is_assignable::value, ""); + static_assert(cuda::std::is_nothrow_assignable::value, ""); + } + { + typedef cuda::std::unique_ptr&> U; + static_assert(!cuda::std::is_assignable::value, ""); + static_assert(!cuda::std::is_assignable::value, ""); + static_assert(!cuda::std::is_assignable::value, ""); + static_assert(cuda::std::is_nothrow_assignable::value, ""); + } + { + typedef cuda::std::unique_ptr&> U; + static_assert(!cuda::std::is_assignable::value, ""); + static_assert(!cuda::std::is_assignable::value, ""); + static_assert(!cuda::std::is_assignable::value, ""); + static_assert(cuda::std::is_nothrow_assignable::value, ""); + } +} + +__host__ __device__ TEST_CONSTEXPR_CXX23 bool test() +{ + { + test_basic(); + test_sfinae(); + } + { + test_basic(); + test_sfinae(); + } + + return true; +} + +int main(int, char**) +{ + test(); +#if TEST_STD_VER >= 2023 + static_assert(test()); +#endif // TEST_STD_VER >= 2023 + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/smartptr/unique.ptr/unique.ptr.class/unique.ptr.asgn/move_convert.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/smartptr/unique.ptr/unique.ptr.class/unique.ptr.asgn/move_convert.pass.cpp new file mode 100644 index 00000000000..9abc631c5e6 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/utilities/memory/smartptr/unique.ptr/unique.ptr.class/unique.ptr.asgn/move_convert.pass.cpp @@ -0,0 +1,478 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// +// UNSUPPORTED: c++03 +// UNSUPPORTED: nvrtc + +// + +// unique_ptr + +// Test unique_ptr converting move ctor + +#include +#include + +#include "test_macros.h" +#include "type_id.h" +#include "unique_ptr_test_helper.h" + +template +struct GenericDeleter +{ + __host__ __device__ TEST_CONSTEXPR_CXX23 void operator()(void*) const {} +}; + +template +struct GenericConvertingDeleter +{ + template + __host__ __device__ TEST_CONSTEXPR_CXX23 GenericConvertingDeleter(GenericConvertingDeleter) + {} + + template + __host__ __device__ TEST_CONSTEXPR_CXX23 GenericConvertingDeleter& operator=(GenericConvertingDeleter const&) + { + return *this; + } + + __host__ __device__ TEST_CONSTEXPR_CXX23 void operator()(void*) const {} +}; + +template +using EnableIfNotSame = typename cuda::std::enable_if< + !cuda::std::is_same::type, typename cuda::std::decay::type>::value>::type; + +template +struct is_specialization; + +template