Fix various issues in CPOs and tuple

.upstream-tests/test/cuda/memory_resource/get_property/forward_property.pass.cpp

@@ -8,7 +8,7 @@
 //===----------------------------------------------------------------------===//
 
 // UNSUPPORTED: c++03, c++11
-
+// XFAIL: clang-7, clang-8
 // cuda::forward_property
 
 #define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE

.upstream-tests/test/std/utilities/tuple/tuple.general/ignore.pass.cpp

@@ -49,10 +49,12 @@ constexpr bool test_ignore_constexpr()
 }
 
 int main(int, char**) {
+#ifndef __CUDA_ARCH__
     {
         constexpr auto& ignore_v = cuda::std::ignore;
         ((void)ignore_v);
     }
+#endif
     {
         static_assert(test_ignore_constexpr(), "");
     }

.upstream-tests/test/std/utilities/tuple/tuple.tuple/tuple.creation/tie.pass.cpp

@@ -48,6 +48,7 @@ __host__ __device__ constexpr bool test_tie_constexpr() {
 int main(int, char**)
 {
     {
+#ifndef __CUDA_ARCH__
         int i = 0;
         const char *_s = "C++";
         // cuda::std::string not supported
@@ -56,6 +57,7 @@ int main(int, char**)
         cuda::std::tie(i, cuda::std::ignore, s) = cuda::std::make_tuple(42, 3.14, _s);
         assert(i == 42);
         assert(s == _s);
+#endif
     }
 #if TEST_STD_VER > 11
     {

.upstream-tests/test/std/utilities/utility/forward/move.pass.cpp

@@ -37,8 +37,8 @@ __device__ const int& cx = x;
 template <class QualInt>
 __host__ __device__ QualInt get() TEST_NOEXCEPT { return static_cast<QualInt>(x); }
 
-_LIBCUDACXX_CPO_ACCESSIBILITY int copy_ctor = 0;
-_LIBCUDACXX_CPO_ACCESSIBILITY int move_ctor = 0;
+STATIC_TEST_GLOBAL_VAR int copy_ctor = 0;
+STATIC_TEST_GLOBAL_VAR int move_ctor = 0;
 
 struct A {
     __host__ __device__ A() {}

.upstream-tests/test/support/test_macros.h

@@ -399,6 +399,13 @@ inline void DoNotOptimize(Tp const& value) {
 # define _STATIC_MEMBER_IMPL(type) static type v;
 #endif
 
+#if defined(__CUDA_ARCH__)
+#  define STATIC_TEST_GLOBAL_VAR static __device__
+#else
+#  define STATIC_TEST_GLOBAL_VAR
+#endif
+
+
 #define STATIC_MEMBER_VAR(name, type) \
   __host__ __device__ static type& name() { \
     _STATIC_MEMBER_IMPL(type); \

include/cuda/memory_resource

@@ -195,7 +195,7 @@ struct __fn {
 _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo {
-  _LIBCUDACXX_CPO_ACCESSIBILITY _LIBCUDACXX_INLINE_VAR constexpr auto get_property = __get_property::__fn{};
+  _LIBCUDACXX_CPO_ACCESSIBILITY constexpr auto get_property = __get_property::__fn{};
 } // namespace __cpo
 
 namespace mr

include/cuda/std/detail/__config

@@ -221,10 +221,4 @@
 #undef _LIBCUDACXX_END_NAMESPACE_CPO
 #define _LIBCUDACXX_END_NAMESPACE_CPO } }
 
-#if defined(__CUDA_ARCH__)
-#  define _LIBCUDACXX_CPO_ACCESSIBILITY static __device__
-#else
-#  define _LIBCUDACXX_CPO_ACCESSIBILITY
-#endif
-
 #endif //__cuda_std__

include/cuda/std/detail/libcxx/include/__concepts/swappable.h

@@ -84,7 +84,7 @@ _LIBCUDACXX_BEGIN_NAMESPACE_CPO(__swap)
     move_constructible<_Tp> &&
     assignable_from<_Tp&, _Tp>;
 
-#else // ^^^ CXX20 ^^^ / vvv CXX17 vvv 
+#else // ^^^ CXX20 ^^^ / vvv CXX17 vvv
 
   template<class _Tp, class _Up>
   _LIBCUDACXX_CONCEPT_FRAGMENT(
@@ -95,10 +95,10 @@ _LIBCUDACXX_BEGIN_NAMESPACE_CPO(__swap)
 
   template<class _Tp, class _Up>
   _LIBCUDACXX_CONCEPT __unqualified_swappable_with = _LIBCUDACXX_FRAGMENT(__unqualified_swappable_with_, _Tp, _Up);
-    
+
   template<class _Tp, class _Up, size_t _Size, class = void>
   _LIBCUDACXX_INLINE_VAR constexpr bool __swappable_arrays = false;
-  
+
   template<class _Tp>
   _LIBCUDACXX_CONCEPT_FRAGMENT(
     __exchangeable_,
@@ -107,7 +107,7 @@ _LIBCUDACXX_BEGIN_NAMESPACE_CPO(__swap)
       requires(move_constructible<_Tp>),
       requires(assignable_from<_Tp&, _Tp>)
     ));
-     
+
   template<class _Tp>
   _LIBCUDACXX_CONCEPT __exchangeable = _LIBCUDACXX_FRAGMENT(__exchangeable_, _Tp);
 #endif // _LIBCUDACXX_STD_VER < 20
@@ -128,7 +128,7 @@ _LIBCUDACXX_BEGIN_NAMESPACE_CPO(__swap)
 
     // 2.2   Otherwise, if `E1` and `E2` are lvalues of array types with equal extent and...
     _LIBCUDACXX_TEMPLATE(class _Tp, class _Up, size_t _Size)
-      (requires __swappable_arrays<_Tp, _Up, _Size>)      
+      (requires __swappable_arrays<_Tp, _Up, _Size>)
     _LIBCUDACXX_INLINE_VISIBILITY constexpr void operator()(_Tp(& __t)[_Size], _Up(& __u)[_Size]) const
       noexcept(__noexcept_swappable_arrays<_Tp, _Up>)
     {
@@ -140,7 +140,7 @@ _LIBCUDACXX_BEGIN_NAMESPACE_CPO(__swap)
 
     // 2.3   Otherwise, if `E1` and `E2` are lvalues of the same type `T` that models...
     _LIBCUDACXX_TEMPLATE(class _Tp)
-      (requires __exchangeable<_Tp>)      
+      (requires __exchangeable<_Tp>)
     _LIBCUDACXX_INLINE_VISIBILITY constexpr void operator()(_Tp& __x, _Tp& __y) const
       noexcept(is_nothrow_move_constructible_v<_Tp> && is_nothrow_move_assignable_v<_Tp>)
     {
@@ -157,7 +157,7 @@ template<class _Tp, class _Up, class _Size>
       requires(extent_v<_Tp> == extent_v<_Up>),
       typename(decltype(__swap(__t[0], __u[0])))
     ));
-    
+
   template<class _Tp, class _Up, size_t _Size>
   _LIBCUDACXX_INLINE_VAR constexpr bool __swappable_arrays<_Tp, _Up, _Size, void_t<type_identity_t<_Tp>>> =
     _LIBCUDACXX_FRAGMENT(__swappable_arrays_, _Tp, _Up, _CUDA_VSTD::integral_constant<size_t, _Size>);
@@ -170,7 +170,7 @@ template<class _Tp, class _Up, class _Size>
 _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo {
-  _LIBCUDACXX_CPO_ACCESSIBILITY _LIBCUDACXX_INLINE_VAR constexpr auto swap = __swap::__fn{};
+  _LIBCUDACXX_CPO_ACCESSIBILITY constexpr auto swap = __swap::__fn{};
 } // namespace __cpo
 _LIBCUDACXX_END_NAMESPACE_RANGES
 

include/cuda/std/detail/libcxx/include/__config

@@ -1324,15 +1324,21 @@ typedef __char32_t char32_t;
 #define _LIBCUDACXX_BEGIN_NAMESPACE_STD namespace std { inline namespace _LIBCUDACXX_ABI_NAMESPACE {
 #define _LIBCUDACXX_END_NAMESPACE_STD  } }
 #define _CUDA_VSTD std::_LIBCUDACXX_ABI_NAMESPACE
+
+#if !defined(__cuda_std__)
 _LIBCUDACXX_BEGIN_NAMESPACE_STD _LIBCUDACXX_END_NAMESPACE_STD
+#endif
 
 #define _LIBCUDACXX_BEGIN_NAMESPACE_STD_NOVERSION namespace std {
 #define _LIBCUDACXX_END_NAMESPACE_STD_NOVERSION }
 
 #define _LIBCUDACXX_BEGIN_NAMESPACE_RANGES \
-   namespace std { namespace ranges { inline namespace _LIBCUDACXX_CUDA_ABI_NAMESPACE {
+    namespace std { namespace ranges { inline namespace _LIBCUDACXX_CUDA_ABI_NAMESPACE {
 #define _LIBCUDACXX_END_NAMESPACE_RANGES } } }
+
+#if !defined(__cuda_std__)
 _LIBCUDACXX_BEGIN_NAMESPACE_RANGES _LIBCUDACXX_END_NAMESPACE_RANGES
+#endif
 
 #define _CUDA_VRANGES std::ranges::_LIBCUDACXX_CUDA_ABI_NAMESPACE
 
@@ -2023,9 +2029,9 @@ extern "C" _LIBCUDACXX_FUNC_VIS void __sanitizer_annotate_contiguous_container(
 #endif
 
 #if defined(__CUDA_ARCH__)
-#  define _LIBCUDACXX_CPO_ACCESSIBILITY static __device__
+#  define _LIBCUDACXX_CPO_ACCESSIBILITY __device__
 #else
-#  define _LIBCUDACXX_CPO_ACCESSIBILITY
+#  define _LIBCUDACXX_CPO_ACCESSIBILITY _LIBCUDACXX_INLINE_VAR
 #endif
 
 #endif // __cplusplus

include/cuda/std/detail/libcxx/include/__debug

@@ -215,7 +215,7 @@ _LIBCUDACXX_END_NAMESPACE_STD
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
 template <class _Tp>
-_LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY
+_LIBCUDACXX_INLINE_VISIBILITY
 _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 inline void __debug_db_insert_c(_Tp* __c) {
 #ifdef _LIBCUDACXX_ENABLE_DEBUG_MODE
     if (!__libcpp_is_constant_evaluated())
@@ -226,7 +226,7 @@ _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 inline void __debug_db_insert_c(_Tp* __c) {
 }
 
 template <class _Tp>
-_LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY
+_LIBCUDACXX_INLINE_VISIBILITY
 _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 inline void __debug_db_insert_i(_Tp* __i) {
 #ifdef _LIBCUDACXX_ENABLE_DEBUG_MODE
     if (!__libcpp_is_constant_evaluated())
@@ -237,7 +237,7 @@ _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 inline void __debug_db_insert_i(_Tp* __i) {
 }
 
 template <class _Tp>
-_LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY
+_LIBCUDACXX_INLINE_VISIBILITY
 _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 inline void __debug_db_erase_c(_Tp* __c) {
 #ifdef _LIBCUDACXX_ENABLE_DEBUG_MODE
     if (!__libcpp_is_constant_evaluated())
@@ -248,7 +248,7 @@ _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 inline void __debug_db_erase_c(_Tp* __c) {
 }
 
 template <class _Tp>
-_LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY
+_LIBCUDACXX_INLINE_VISIBILITY
 _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 inline void __debug_db_swap(_Tp* __lhs, _Tp* __rhs) {
 #ifdef _LIBCUDACXX_ENABLE_DEBUG_MODE
     if (!__libcpp_is_constant_evaluated())
@@ -260,7 +260,7 @@ _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 inline void __debug_db_swap(_Tp* __lhs, _Tp* _
 }
 
 template <class _Tp>
-_LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY
+_LIBCUDACXX_INLINE_VISIBILITY
 _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 inline void __debug_db_invalidate_all(_Tp* __c) {
 #ifdef _LIBCUDACXX_ENABLE_DEBUG_MODE
     if (!__libcpp_is_constant_evaluated())

include/cuda/std/detail/libcxx/include/__tuple_dir/tuple_element.h

@@ -49,7 +49,36 @@ struct _LIBCUDACXX_TEMPLATE_VIS tuple_element<_Ip, const volatile _Tp>
 
 #ifndef _LIBCUDACXX_CXX03_LANG
 
-#if !__has_builtin(__type_pack_element)
+#ifdef _LIBCUDACXX_COMPILER_MSVC
+
+namespace __indexer_detail {
+
+template <size_t _Idx, class ..._Types>
+struct _nth_of;
+
+template <class _Head, class ..._Tail>
+struct _nth_of<0, _Head, _Tail...> {
+    using type = _Head;
+};
+
+template <size_t _Idx, class _Head, class ..._Tail>
+struct _nth_of<_Idx, _Head, _Tail...> {
+    using type = typename _nth_of<_Idx-1, _Tail...>::type;
+};
+
+template <size_t _Idx, class ..._Types>
+struct nth_of {
+    static_assert(_Idx < sizeof...(_Types), "");
+    using _impl = _nth_of<_Idx, _Types...>;
+    using type = typename _impl::type;
+};
+
+} // namespace __indexer_detail
+
+template <size_t _Idx, class ..._Types>
+using __type_pack_element _LIBCUDACXX_NODEBUG_TYPE = typename __indexer_detail::nth_of<_Idx, _Types...>::type;
+
+#elif !__has_builtin(__type_pack_element)
 
 namespace __indexer_detail {
 

include/cuda/std/detail/libcxx/include/__type_traits/common_reference.h

@@ -43,16 +43,18 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD
 // Let COND_RES(X, Y) be:
 #ifdef _LIBCUDACXX_COMPILER_MSVC // Workaround for DevCom-1627396
 template <class _Xp, class _Yp>
-using __cond_res_if_right =
-    decltype(false ? _CUDA_VSTD::declval<_Xp(&)()>()() : _CUDA_VSTD::declval<_Yp(&)()>()());
-
+struct __cond_res_if_right_wrapper {
+    using __type = decltype(false ? _CUDA_VSTD::declval<_Xp(&)()>()() : _CUDA_VSTD::declval<_Yp(&)()>()());
+};
+template <class _Xp, class _Yp>
+using __cond_res_if_right = typename __cond_res_if_right_wrapper<_Xp, _Yp>::__type;
 template <class _Tp, class _Up, class = void>
 struct __cond_res_workaround {};
 
 template <class _Tp, class _Up>
 struct __cond_res_workaround<_Tp, _Up, void_t<__cond_res_if_right<_Tp, _Up>>> {
     using _RTp = remove_cvref_t<_Tp>;
-    using type = conditional_t<is_same_v<_RTp, remove_cvref_t<_Up>> && 
+    using type = conditional_t<is_same_v<_RTp, remove_cvref_t<_Up>> &&
                                (is_scalar_v<_RTp> || is_array_v<_RTp>) &&
                                ((is_lvalue_reference_v<_Tp> && is_rvalue_reference_v<_Up>) || (is_rvalue_reference_v<_Tp> && is_lvalue_reference_v<_Up>)),
                  decay_t<__copy_cv_t<remove_reference_t<_Tp>, remove_reference_t<_Up>>>, __cond_res_if_right<_Tp, _Up>>;
@@ -108,7 +110,7 @@ struct __common_ref_rr {};
 
 template<class _Ap, class _Bp>
 struct __common_ref_rr<_Ap&&, _Bp&&, enable_if_t<
-                            is_convertible_v<_Ap&&, __common_ref_C<_Ap, _Bp>> 
+                            is_convertible_v<_Ap&&, __common_ref_C<_Ap, _Bp>>
                          && is_convertible_v<_Bp&&, __common_ref_C<_Ap, _Bp>>>>
 {
     using __type = __common_ref_C<_Ap, _Bp>;
@@ -165,17 +167,17 @@ struct common_reference<_Tp>
 
 // bullet 3 - sizeof...(T) == 2
 template <class _Tp, class _Up, class = void> struct __common_reference_sub_bullet3;
-template <class _Tp, class _Up, class = void> struct __common_reference_sub_bullet2 
+template <class _Tp, class _Up, class = void> struct __common_reference_sub_bullet2
     : __common_reference_sub_bullet3<_Tp, _Up> {};
-template <class _Tp, class _Up, class = void> struct __common_reference_sub_bullet1 
+template <class _Tp, class _Up, class = void> struct __common_reference_sub_bullet1
     : __common_reference_sub_bullet2<_Tp, _Up> {};
 
 // sub-bullet 1 - If T1 and T2 are reference types and COMMON-REF(T1, T2) is well-formed, then
 // the member typedef `type` denotes that type.
 template <class _Tp, class _Up> struct common_reference<_Tp, _Up> : __common_reference_sub_bullet1<_Tp, _Up> {};
 
 template <class _Tp, class _Up>
-struct __common_reference_sub_bullet1<_Tp, _Up, void_t<__common_ref_t<_Tp, _Up>, 
+struct __common_reference_sub_bullet1<_Tp, _Up, void_t<__common_ref_t<_Tp, _Up>,
     enable_if_t<is_reference_v<_Tp> && is_reference_v<_Up>>>>
 {
     using type = __common_ref_t<_Tp, _Up>;
@@ -224,4 +226,4 @@ template <class...> struct common_reference {};
 
 _LIBCUDACXX_END_NAMESPACE_STD
 
-#endif // _LIBCUDACXX___TYPE_TRAITS_COMMON_REFERENCE_H
+#endif // _LIBCUDACXX___TYPE_TRAITS_COMMON_REFERENCE_H

include/cuda/std/detail/libcxx/include/__type_traits/is_convertible.h

@@ -53,16 +53,16 @@ template <class _Ty>
 struct is_convertible<volatile _Ty&, const volatile _Ty&> : true_type {};
 
 template <class _Ty>
-_INLINE_VAR constexpr bool is_convertible_v<_Ty&, volatile _Ty&> = true;
+_LIBCUDACXX_INLINE_VAR constexpr bool is_convertible_v<_Ty&, volatile _Ty&> = true;
 
 template <class _Ty>
-_INLINE_VAR constexpr bool is_convertible_v<volatile _Ty&, volatile _Ty&> = true;
+_LIBCUDACXX_INLINE_VAR constexpr bool is_convertible_v<volatile _Ty&, volatile _Ty&> = true;
 
 template <class _Ty>
-_INLINE_VAR constexpr bool is_convertible_v<_Ty&, const volatile _Ty&> = true;
+_LIBCUDACXX_INLINE_VAR constexpr bool is_convertible_v<_Ty&, const volatile _Ty&> = true;
 
 template <class _Ty>
-_INLINE_VAR constexpr bool is_convertible_v<volatile _Ty&, const volatile _Ty&> = true;
+_LIBCUDACXX_INLINE_VAR constexpr bool is_convertible_v<volatile _Ty&, const volatile _Ty&> = true;
 #endif // _LIBCUDACXX_COMPILER_MSVC
 
 #else  // __has_builtin(__is_convertible_to) && !defined(_LIBCUDACXX_USE_IS_CONVERTIBLE_FALLBACK)

include/cuda/std/detail/libcxx/include/__utility/in_place.h

@@ -28,21 +28,21 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD
 struct _LIBCUDACXX_TYPE_VIS in_place_t {
     explicit in_place_t() = default;
 };
-_LIBCUDACXX_CPO_ACCESSIBILITY inline constexpr in_place_t in_place{};
+_LIBCUDACXX_CPO_ACCESSIBILITY constexpr in_place_t in_place{};
 
 template <class _Tp>
 struct _LIBCUDACXX_TEMPLATE_VIS in_place_type_t {
     explicit in_place_type_t() = default;
 };
 template <class _Tp>
-_LIBCUDACXX_CPO_ACCESSIBILITY inline constexpr in_place_type_t<_Tp> in_place_type{};
+_LIBCUDACXX_CPO_ACCESSIBILITY constexpr in_place_type_t<_Tp> in_place_type{};
 
 template <size_t _Idx>
 struct _LIBCUDACXX_TEMPLATE_VIS in_place_index_t {
     explicit in_place_index_t() = default;
 };
 template <size_t _Idx>
-_LIBCUDACXX_CPO_ACCESSIBILITY inline constexpr in_place_index_t<_Idx> in_place_index{};
+_LIBCUDACXX_CPO_ACCESSIBILITY constexpr in_place_index_t<_Idx> in_place_index{};
 
 template <class _Tp> struct __is_inplace_type_imp : false_type {};
 template <class _Tp> struct __is_inplace_type_imp<in_place_type_t<_Tp>> : true_type {};

include/cuda/std/detail/libcxx/include/__utility/unreachable.h

@@ -22,15 +22,15 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
-_LIBCUDACXX_NORETURN _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY
+_LIBCUDACXX_NORETURN _LIBCUDACXX_INLINE_VISIBILITY
 inline void __libcpp_unreachable()
 {
   _LIBCUDACXX_UNREACHABLE();
 }
 
 #if _LIBCUDACXX_STD_VER > 20
 
-[[noreturn]] _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY
+[[noreturn]] _LIBCUDACXX_INLINE_VISIBILITY
 inline void unreachable() { _LIBCUDACXX_UNREACHABLE(); }
 
 #endif // _LIBCUDACXX_STD_VER > 20