Drop implementation of thrust::pair and thrust::tuple (#2395)

* Drop implementation of `thrust::pair` and `thrust::tuple` We previously moved them back to proper class definitions, as using alias declarations broke CTAD. Thanks to @bernhardmgruber who realized that instead of making them an alias we can just pull them in and be done with it. Co-authored-by: Bernhard Manfred Gruber <[email protected]>
NVIDIA · Sep 11, 2024 · e7ade77 · e7ade77
1 parent d5492d5
commit e7ade77
Show file tree

Hide file tree

Showing 6 changed files with 60 additions and 215 deletions.
diff --git a/docs/thrust/api_docs/utility/pair.rst b/docs/thrust/api_docs/utility/pair.rst
@@ -3,11 +3,12 @@
 Pair
 ----
 
-  - :cpp:struct:`thrust::pair <thrust::pair>`
+Thrust provides an implementations of `std::pair`, that is pulled in from `libcu++`.
+
+It is recommended to replace `thrust::pair` and `thrust::make_pair` with `cuda::std::pair` and `cuda::std::make_pair`.
 
 .. toctree::
    :glob:
    :maxdepth: 1
 
-   ${repo_docs_api_path}/*struct*pair*
-   ${repo_docs_api_path}/*function_group__pair*
+   ${repo_docs_api_path}/*typedef*pair*
diff --git a/docs/thrust/api_docs/utility/tuple.rst b/docs/thrust/api_docs/utility/tuple.rst
@@ -3,10 +3,12 @@
 Tuple
 ------
 
-  - :cpp:struct:`thrust::tuple <thrust::tuple>`
+Thrust provides an implementation of `std::tuple`, that is pulled in from `libcu++`.
+
+It is recommended to replace `thrust::tuple` with `cuda::std::tuple`.
 
 .. toctree::
    :glob:
    :maxdepth: 1
 
-   ${repo_docs_api_path}/*function_group__tuple*
+   ${repo_docs_api_path}/*typedef*tuple*
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/tuple b/libcudacxx/include/cuda/std/detail/libcxx/include/tuple
@@ -194,6 +194,10 @@ _CCCL_PUSH_MACROS
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
+template <class>
+struct __is_tuple_of_iterator_references : _CUDA_VSTD::false_type
+{};
+
 // __tuple_leaf
 struct __tuple_leaf_default_constructor_tag
 {};
@@ -854,6 +858,15 @@ 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_HIDE_FROM_ABI _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,

diff --git a/thrust/thrust/iterator/detail/tuple_of_iterator_references.h b/thrust/thrust/iterator/detail/tuple_of_iterator_references.h
@@ -67,7 +67,9 @@ class tuple_of_iterator_references : public thrust::tuple<Ts...>
   using super_t = thrust::tuple<Ts...>;
   using super_t::super_t;
 
-  tuple_of_iterator_references() = default;
+  inline _CCCL_HOST_DEVICE tuple_of_iterator_references()
+      : super_t()
+  {}
 
   // allow implicit construction from tuple<refs>
   inline _CCCL_HOST_DEVICE tuple_of_iterator_references(const super_t& other)
@@ -136,15 +138,14 @@ class tuple_of_iterator_references : public thrust::tuple<Ts...>
 
 } // namespace detail
 
-template <class... Ts>
-struct __is_tuple_of_iterator_references<THRUST_NS_QUALIFIER::detail::tuple_of_iterator_references<Ts...>>
-    : _CUDA_VSTD::true_type
-{};
-
 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...>> : true_type
+{};
+
 // define tuple_size, tuple_element, etc.
 template <class... Ts>
 struct tuple_size<THRUST_NS_QUALIFIER::detail::tuple_of_iterator_references<Ts...>>

diff --git a/thrust/thrust/pair.h b/thrust/thrust/pair.h
@@ -30,9 +30,6 @@
 #  pragma system_header
 #endif // no system header
 
-#include <thrust/type_traits/is_trivially_relocatable.h>
-
-#include <cuda/std/__type_traits/conjunction.h>
 #include <cuda/std/utility>
 
 THRUST_NAMESPACE_BEGIN
@@ -52,17 +49,25 @@ THRUST_NAMESPACE_BEGIN
  *  \tparam N This parameter selects the member of interest.
  *  \tparam T A \c pair type of interest.
  */
+#ifdef DOXYGEN_SHOULD_SKIP_THIS // Provide a fake alias for doxygen
 template <size_t N, class T>
 using tuple_element = _CUDA_VSTD::tuple_element<N, T>;
+#else // ^^^ DOXYGEN_SHOULD_SKIP_THIS ^^^ / vvv !DOXYGEN_SHOULD_SKIP_THIS vvv
+using _CUDA_VSTD::tuple_element;
+#endif // DOXYGEN_SHOULD_SKIP_THIS
 
 /*! This convenience metafunction is included for compatibility with
  *  \p tuple. It returns \c 2, the number of elements of a \p pair,
  *  in its nested data member, \c value.
  *
  *  \tparam Pair A \c pair type of interest.
  */
+#ifdef DOXYGEN_SHOULD_SKIP_THIS // Provide a fake alias for doxygen
 template <class T>
 using tuple_size = _CUDA_VSTD::tuple_size<T>;
+#else // ^^^ DOXYGEN_SHOULD_SKIP_THIS ^^^ / vvv !DOXYGEN_SHOULD_SKIP_THIS vvv
+using _CUDA_VSTD::tuple_size;
+#endif // DOXYGEN_SHOULD_SKIP_THIS
 
 /*! \p pair is a generic data structure encapsulating a heterogeneous
  *  pair of values.
@@ -75,55 +80,15 @@ using tuple_size = _CUDA_VSTD::tuple_size<T>;
  *          requirements on the type of \p T2. <tt>T2</tt>'s type is
  *          provided by <tt>pair::second_type</tt>.
  */
+#ifdef DOXYGEN_SHOULD_SKIP_THIS // Provide a fake alias for doxygen
 template <class T, class U>
-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 pair = _CUDA_VSTD::pair<T, U>;
+#else // ^^^ DOXYGEN_SHOULD_SKIP_THIS ^^^ / vvv !DOXYGEN_SHOULD_SKIP_THIS vvv
+using _CUDA_VSTD::pair;
+#endif // DOXYGEN_SHOULD_SKIP_THIS
 
 using _CUDA_VSTD::get;
-
-template <typename T, typename U>
-struct proclaim_trivially_relocatable<pair<T, U>>
-    : ::cuda::std::conjunction<is_trivially_relocatable<T>, is_trivially_relocatable<U>>
-{};
+using _CUDA_VSTD::make_pair;
 
 /*! \endcond
  */
@@ -135,38 +100,3 @@ struct proclaim_trivially_relocatable<pair<T, U>>
  */
 
 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