Try to ensure that thrust::tuple and thrust::pair work with CTAD

libcudacxx/include/cuda/std/__utility/pair.h

@@ -283,18 +283,21 @@ struct _LIBCUDACXX_TEMPLATE_VIS pair : public __pair_base<_T1, _T2>
       : __base(__t1, __t2)
   {}
 
-  template <class _U1          = _T1,
-            class _U2          = _T2,
-            class _Constraints = typename __pair_constraints<_T1, _T2>::template __constructible<_U1, _U2>,
+  template <class _U1, class _U2>
+  using __pair_constructible = typename __pair_constraints<_T1, _T2>::template __constructible<_U1, _U2>;
+
+  template <class _U1                                                  = _T1,
+            class _U2                                                  = _T2,
+            class _Constraints                                         = __pair_constructible<_U1, _U2>,
             __enable_if_t<_Constraints::__explicit_constructible, int> = 0>
   _LIBCUDACXX_INLINE_VISIBILITY explicit constexpr pair(_U1&& __u1, _U2&& __u2) noexcept(
     _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T1, _U1) && _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T2, _U2))
       : __base(_CUDA_VSTD::forward<_U1>(__u1), _CUDA_VSTD::forward<_U2>(__u2))
   {}
 
-  template <class _U1          = _T1,
-            class _U2          = _T2,
-            class _Constraints = typename __pair_constraints<_T1, _T2>::template __constructible<_U1, _U2>,
+  template <class _U1                                                  = _T1,
+            class _U2                                                  = _T2,
+            class _Constraints                                         = __pair_constructible<_U1, _U2>,
             __enable_if_t<_Constraints::__implicit_constructible, int> = 0>
   _LIBCUDACXX_INLINE_VISIBILITY constexpr pair(_U1&& __u1, _U2&& __u2) noexcept(
     _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T1, _U1) && _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T2, _U2))
@@ -316,19 +319,19 @@ struct _LIBCUDACXX_TEMPLATE_VIS pair : public __pair_base<_T1, _T2>
   pair(pair const&) = default;
   pair(pair&&)      = default;
 
-  template <class _U1 = _T1,
-            class _U2 = _T2,
-            class _Constraints = typename __pair_constraints<_T1, _T2>::template __constructible<const _U1&, const _U2&>,
+  template <class _U1                                                  = _T1,
+            class _U2                                                  = _T2,
+            class _Constraints                                         = __pair_constructible<const _U1&, const _U2&>,
             __enable_if_t<_Constraints::__explicit_constructible, int> = 0>
   _LIBCUDACXX_INLINE_VISIBILITY explicit _CCCL_CONSTEXPR_CXX14 pair(const pair<_U1, _U2>& __p) noexcept(
     _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T1, const _U1&)
     && _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T2, const _U2&))
       : __base(__p.first, __p.second)
   {}
 
-  template <class _U1 = _T1,
-            class _U2 = _T2,
-            class _Constraints = typename __pair_constraints<_T1, _T2>::template __constructible<const _U1&, const _U2&>,
+  template <class _U1                                                  = _T1,
+            class _U2                                                  = _T2,
+            class _Constraints                                         = __pair_constructible<const _U1&, const _U2&>,
             __enable_if_t<_Constraints::__implicit_constructible, int> = 0>
   _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 pair(const pair<_U1, _U2>& __p) noexcept(
     _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T1, const _U1&)
@@ -337,18 +340,18 @@ struct _LIBCUDACXX_TEMPLATE_VIS pair : public __pair_base<_T1, _T2>
   {}
 
   // move constructors
-  template <class _U1          = _T1,
-            class _U2          = _T2,
-            class _Constraints = typename __pair_constraints<_T1, _T2>::template __constructible<_U1, _U2>,
+  template <class _U1                                                  = _T1,
+            class _U2                                                  = _T2,
+            class _Constraints                                         = __pair_constructible<_U1, _U2>,
             __enable_if_t<_Constraints::__explicit_constructible, int> = 0>
   _LIBCUDACXX_INLINE_VISIBILITY explicit _CCCL_CONSTEXPR_CXX14 pair(pair<_U1, _U2>&& __p) noexcept(
     _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T1, _U1) && _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T2, _U2))
       : __base(_CUDA_VSTD::forward<_U1>(__p.first), _CUDA_VSTD::forward<_U2>(__p.second))
   {}
 
-  template <class _U1          = _T1,
-            class _U2          = _T2,
-            class _Constraints = typename __pair_constraints<_T1, _T2>::template __constructible<_U1, _U2>,
+  template <class _U1                                                  = _T1,
+            class _U2                                                  = _T2,
+            class _Constraints                                         = __pair_constructible<_U1, _U2>,
             __enable_if_t<_Constraints::__implicit_constructible, int> = 0>
   _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 pair(pair<_U1, _U2>&& __p) noexcept(
     _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T1, _U1) && _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T2, _U2))
@@ -359,7 +362,7 @@ struct _LIBCUDACXX_TEMPLATE_VIS pair : public __pair_base<_T1, _T2>
 #if defined(__cuda_std__) && !defined(_CCCL_COMPILER_NVRTC)
   template <class _U1,
             class _U2,
-            class _Constraints = typename __pair_constraints<_T1, _T2>::template __constructible<const _U1&, const _U2&>,
+            class _Constraints                                         = __pair_constructible<const _U1&, const _U2&>,
             __enable_if_t<_Constraints::__explicit_constructible, int> = 0>
   _CCCL_HOST _LIBCUDACXX_HIDE_FROM_ABI explicit _CCCL_CONSTEXPR_CXX14 pair(const ::std::pair<_U1, _U2>& __p) noexcept(
     _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T1, const _U1&)
@@ -369,7 +372,7 @@ struct _LIBCUDACXX_TEMPLATE_VIS pair : public __pair_base<_T1, _T2>
 
   template <class _U1,
             class _U2,
-            class _Constraints = typename __pair_constraints<_T1, _T2>::template __constructible<const _U1&, const _U2&>,
+            class _Constraints                                         = __pair_constructible<const _U1&, const _U2&>,
             __enable_if_t<_Constraints::__implicit_constructible, int> = 0>
   _CCCL_HOST _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 pair(const ::std::pair<_U1, _U2>& __p) noexcept(
     _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T1, const _U1&)
@@ -379,7 +382,7 @@ struct _LIBCUDACXX_TEMPLATE_VIS pair : public __pair_base<_T1, _T2>
 
   template <class _U1,
             class _U2,
-            class _Constraints = typename __pair_constraints<_T1, _T2>::template __constructible<_U1, _U2>,
+            class _Constraints                                         = __pair_constructible<_U1, _U2>,
             __enable_if_t<_Constraints::__explicit_constructible, int> = 0>
   _CCCL_HOST _LIBCUDACXX_HIDE_FROM_ABI explicit _CCCL_CONSTEXPR_CXX14 pair(::std::pair<_U1, _U2>&& __p) noexcept(
     _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T1, _U1) && _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T2, _U2))
@@ -388,7 +391,7 @@ struct _LIBCUDACXX_TEMPLATE_VIS pair : public __pair_base<_T1, _T2>
 
   template <class _U1,
             class _U2,
-            class _Constraints = typename __pair_constraints<_T1, _T2>::template __constructible<_U1, _U2>,
+            class _Constraints                                         = __pair_constructible<_U1, _U2>,
             __enable_if_t<_Constraints::__implicit_constructible, int> = 0>
   _CCCL_HOST _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 pair(::std::pair<_U1, _U2>&& __p) noexcept(
     _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T1, _U1) && _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T2, _U2))

libcudacxx/include/cuda/std/detail/libcxx/include/tuple

@@ -194,10 +194,6 @@ template <class... Types>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
-template <class>
-struct __is_tuple_of_iterator_references : false_type
-{};
-
 // __tuple_leaf
 struct __tuple_leaf_default_constructor_tag
 {};
@@ -808,15 +804,6 @@ public:
         typename __tuple_constraints<_Tp...>::template __tuple_like_constraints<_Tuple>,
         __invalid_tuple_constraints>;
 
-  // Horrible hack to make tuple_of_iterator_references work
-  template <class _TupleOfIteratorReferences,
-            __enable_if_t<__is_tuple_of_iterator_references<_TupleOfIteratorReferences>::value, int> = 0,
-            __enable_if_t<(tuple_size<_TupleOfIteratorReferences>::value == sizeof...(_Tp)), int>    = 0>
-  _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 tuple(_TupleOfIteratorReferences&& __t)
-      : tuple(_CUDA_VSTD::forward<_TupleOfIteratorReferences>(__t).template __to_tuple<_Tp...>(
-        __make_tuple_indices_t<sizeof...(_Tp)>()))
-  {}
-
   template <class _Tuple,
             class _Constraints                                                  = __tuple_like_constraints<_Tuple>,
             __enable_if_t<!_PackExpandsToThisTuple<_Tuple>::value, int>         = 0,

thrust/testing/pair.cu

@@ -322,4 +322,16 @@ void TestPairStructuredBindings(void)
   ASSERT_EQUAL(b, b2);
 }
 DECLARE_UNITTEST(TestPairStructuredBindings);
-#endif
+
+void TestPairCTAD(void)
+{
+  const int a = 42;
+  const int b = 1337;
+  thrust::pair p(a, b);
+
+  auto [a2, b2] = p;
+  ASSERT_EQUAL(a, a2);
+  ASSERT_EQUAL(b, b2);
+}
+DECLARE_UNITTEST(TestPairCTAD);
+#endif // _CCCL_STD_VER >= 2017

thrust/testing/tuple.cu

@@ -511,7 +511,21 @@ void TestTupleStructuredBindings(void)
   ASSERT_EQUAL(c, c2);
 }
 DECLARE_UNITTEST(TestTupleStructuredBindings);
-#endif
+
+void TestTupleCTAD(void)
+{
+  const int a   = 0;
+  const char b  = 42;
+  const short c = 1337;
+  thrust::tuple t(a, b, c);
+
+  auto [a2, b2, c2] = t;
+  ASSERT_EQUAL(a, a2);
+  ASSERT_EQUAL(b, b2);
+  ASSERT_EQUAL(c, c2);
+}
+DECLARE_UNITTEST(TestTupleCTAD);
+#endif // _CCCL_STD_VER >= 2017
 
 // Ensure that we are backwards compatible with the old thrust::tuple implementation
 static_assert(

thrust/thrust/detail/functional/actor.h

@@ -65,7 +65,7 @@ struct actor : Eval
 {
   typedef Eval eval_type;
 
-  _CCCL_HOST_DEVICE constexpr actor();
+  constexpr actor() = default;
 
   _CCCL_HOST_DEVICE actor(const Eval& base);
 

thrust/thrust/detail/functional/actor.inl

@@ -48,11 +48,6 @@ namespace detail
 namespace functional
 {
 
-template <typename Eval>
-_CCCL_HOST_DEVICE constexpr actor<Eval>::actor()
-    : eval_type()
-{}
-
 template <typename Eval>
 _CCCL_HOST_DEVICE actor<Eval>::actor(const Eval& base)
     : eval_type(base)

thrust/thrust/iterator/detail/tuple_of_iterator_references.h

@@ -67,9 +67,7 @@ class tuple_of_iterator_references : public thrust::tuple<Ts...>
   using super_t = thrust::tuple<Ts...>;
   using super_t::super_t;
 
-  inline _CCCL_HOST_DEVICE tuple_of_iterator_references()
-      : super_t()
-  {}
+  tuple_of_iterator_references() = default;
 
   // allow implicit construction from tuple<refs>
   inline _CCCL_HOST_DEVICE tuple_of_iterator_references(const super_t& other)
@@ -138,15 +136,15 @@ class tuple_of_iterator_references : public thrust::tuple<Ts...>
 
 } // namespace detail
 
-THRUST_NAMESPACE_END
-
-_LIBCUDACXX_BEGIN_NAMESPACE_STD
-
 template <class... Ts>
 struct __is_tuple_of_iterator_references<THRUST_NS_QUALIFIER::detail::tuple_of_iterator_references<Ts...>>
-    : integral_constant<bool, true>
+    : _CUDA_VSTD::true_type
 {};
 
+THRUST_NAMESPACE_END
+
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
+
 // define tuple_size, tuple_element, etc.
 template <class... Ts>
 struct tuple_size<THRUST_NS_QUALIFIER::detail::tuple_of_iterator_references<Ts...>>

thrust/thrust/optional.h

@@ -838,7 +838,12 @@ struct nullopt_t
 /// void foo (thrust::optional<int>);
 /// foo(thrust::nullopt); //pass an empty optional
 /// ```
-static constexpr nullopt_t nullopt{nullopt_t::do_not_use{}, nullopt_t::do_not_use{}};
+#ifdef __CUDA_ARCH__
+__device__ static _LIBCUDACXX_CONSTEXPR_GLOBAL
+#else
+static constexpr
+#endif // __CUDA_ARCH__
+  nullopt_t nullopt{nullopt_t::do_not_use{}, nullopt_t::do_not_use{}};
 
 class bad_optional_access : public std::exception
 {

thrust/thrust/pair.h

@@ -50,7 +50,7 @@ THRUST_NAMESPACE_BEGIN
  *  \tparam T A \c pair type of interest.
  */
 template <size_t N, class T>
-using tuple_element = ::cuda::std::tuple_element<N, T>;
+using tuple_element = _CUDA_VSTD::tuple_element<N, T>;
 
 /*! This convenience metafunction is included for compatibility with
  *  \p tuple. It returns \c 2, the number of elements of a \p pair,
@@ -59,7 +59,7 @@ using tuple_element = ::cuda::std::tuple_element<N, T>;
  *  \tparam Pair A \c pair type of interest.
  */
 template <class T>
-using tuple_size = ::cuda::std::tuple_size<T>;
+using tuple_size = _CUDA_VSTD::tuple_size<T>;
 
 /*! \p pair is a generic data structure encapsulating a heterogeneous
  *  pair of values.
@@ -73,10 +73,49 @@ using tuple_size = ::cuda::std::tuple_size<T>;
  *          provided by <tt>pair::second_type</tt>.
  */
 template <class T, class U>
-using pair = ::cuda::std::pair<T, U>;
-
-using ::cuda::std::get;
-using ::cuda::std::make_pair;
+struct pair : public _CUDA_VSTD::pair<T, U>
+{
+  using super_t = _CUDA_VSTD::pair<T, U>;
+  using super_t::super_t;
+
+#if (defined(_CCCL_COMPILER_GCC) && __GNUC__ < 9) || (defined(_CCCL_COMPILER_CLANG) && __clang_major__ < 12)
+  // For whatever reason nvcc complains about that constructor being used before being defined in a constexpr variable
+  constexpr pair() = default;
+
+  template <class _U1          = T,
+            class _U2          = U,
+            class _Constraints = typename _CUDA_VSTD::__pair_constraints<T, U>::template __constructible<_U1, _U2>,
+            _CUDA_VSTD::__enable_if_t<_Constraints::__implicit_constructible, int> = 0>
+  _CCCL_HOST_DEVICE constexpr pair(_U1&& __u1, _U2&& __u2)
+      : super_t(_CUDA_VSTD::forward<_U1>(__u1), _CUDA_VSTD::forward<_U2>(__u2))
+  {}
+#endif // _CCCL_COMPILER_GCC < 9 || _CCCL_COMPILER_CLANG < 12
+};
+
+#if _CCCL_STD_VER >= 2017
+template <class _T1, class _T2>
+_CCCL_HOST_DEVICE pair(_T1, _T2) -> pair<_T1, _T2>;
+#endif // _CCCL_STD_VER >= 2017
+
+template <class T1, class T2>
+inline _CCCL_HOST_DEVICE
+  _CUDA_VSTD::__enable_if_t<_CUDA_VSTD::__is_swappable<T1>::value && _CUDA_VSTD::__is_swappable<T2>::value, void>
+  swap(pair<T1, T2>& lhs, pair<T1, T2>& rhs) noexcept(
+    (_CUDA_VSTD::__is_nothrow_swappable<T1>::value && _CUDA_VSTD::__is_nothrow_swappable<T2>::value))
+{
+  lhs.swap(rhs);
+}
+
+template <class T1, class T2>
+inline _CCCL_HOST_DEVICE
+  pair<typename _CUDA_VSTD::__unwrap_ref_decay<T1>::type, typename _CUDA_VSTD::__unwrap_ref_decay<T2>::type>
+  make_pair(T1&& t1, T2&& t2)
+{
+  return pair<typename _CUDA_VSTD::__unwrap_ref_decay<T1>::type, typename _CUDA_VSTD::__unwrap_ref_decay<T2>::type>(
+    _CUDA_VSTD::forward<T1>(t1), _CUDA_VSTD::forward<T2>(t2));
+}
+
+using _CUDA_VSTD::get;
 
 /*! \endcond
  */
@@ -88,3 +127,38 @@ using ::cuda::std::make_pair;
  */
 
 THRUST_NAMESPACE_END
+
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
+
+template <class T1, class T2>
+struct tuple_size<THRUST_NS_QUALIFIER::pair<T1, T2>> : tuple_size<pair<T1, T2>>
+{};
+
+template <size_t Id, class T1, class T2>
+struct tuple_element<Id, THRUST_NS_QUALIFIER::pair<T1, T2>> : tuple_element<Id, pair<T1, T2>>
+{};
+
+template <class T1, class T2>
+struct __tuple_like_ext<THRUST_NS_QUALIFIER::pair<T1, T2>> : true_type
+{};
+
+_LIBCUDACXX_END_NAMESPACE_STD
+
+// This is a workaround for the fact that structured bindings require that the specializations of
+// `tuple_size` and `tuple_element` reside in namespace std (https://eel.is/c++draft/dcl.struct.bind#4).
+// See https://github.com/NVIDIA/libcudacxx/issues/316 for a short discussion
+#if _CCCL_STD_VER >= 2017
+
+#  include <utility>
+
+namespace std
+{
+template <class T1, class T2>
+struct tuple_size<THRUST_NS_QUALIFIER::pair<T1, T2>> : tuple_size<pair<T1, T2>>
+{};
+
+template <size_t Id, class T1, class T2>
+struct tuple_element<Id, THRUST_NS_QUALIFIER::pair<T1, T2>> : tuple_element<Id, pair<T1, T2>>
+{};
+} // namespace std
+#endif // _CCCL_STD_VER >= 2017