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

@@ -2023,9 +2023,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 inline
 #endif
 
 #endif // __cplusplus

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

@@ -69,14 +69,20 @@ __indexed<_Idx, _Tp> __at_index(__indexed<_Idx, _Tp> const&);
 
 } // namespace __indexer_detail
 
+// MSVC 14.11 loses its mind over the below if the decltype occurs at global scope
 template <size_t _Idx, class ..._Types>
-using __type_pack_element _LIBCUDACXX_NODEBUG_TYPE = typename decltype(
-    __indexer_detail::__at_index<_Idx>(
-        __indexer_detail::__indexer<
-            __tuple_types<_Types...>,
-            typename __make_tuple_indices<sizeof...(_Types)>::type
-        >{})
-  )::type;
+struct __type_pack_element_impl {
+    using __at = decltype(
+        __indexer_detail::__at_index<_Idx>(
+            __indexer_detail::__indexer<
+                __tuple_types<_Types...>,
+                typename __make_tuple_indices<sizeof...(_Types)>::type
+            >{}));
+};
+
+template <size_t _Idx, class ..._Types>
+using __type_pack_element _LIBCUDACXX_NODEBUG_TYPE =
+    typename __type_pack_element_impl<_Idx, _Types...>::__at::type;
 #endif
 
 template <size_t _Ip, class ..._Types>

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/tuple

@@ -1247,10 +1247,9 @@ struct __ignore_t
 
 namespace {
 #ifdef __CUDA_ARCH__
-  _LIBCUDACXX_INLINE_VAR __constant__ __ignore_t<unsigned char> ignore = __ignore_t<unsigned char>();
-#else
-  _LIBCUDACXX_INLINE_VAR constexpr __ignore_t<unsigned char> ignore = __ignore_t<unsigned char>();
+    static
 #endif
+    _LIBCUDACXX_INLINE_VAR constexpr __ignore_t<unsigned char> ignore = __ignore_t<unsigned char>();
 }
 
 template <class... _Tp>