Replace thrust::swap by cuda::std::swap

Fixes: NVIDIA#2948

thrust/thrust/allocate_unique.h

@@ -20,7 +20,7 @@
 #include <thrust/detail/raw_pointer_cast.h>
 #include <thrust/detail/type_deduction.h>
 
-#include <utility>
+#include <cuda/std/utility>
 
 THRUST_NAMESPACE_BEGIN
 
@@ -114,7 +114,7 @@ struct allocator_delete final
 
   void swap(allocator_delete& other) noexcept
   {
-    using std::swap;
+    using ::cuda::std::swap;
     swap(alloc_, other.alloc_);
   }
 
@@ -216,7 +216,7 @@ struct array_allocator_delete final
 
   void swap(array_allocator_delete& other) noexcept
   {
-    using std::swap;
+    using ::cuda::std::swap;
     swap(alloc_, other.alloc_);
     swap(count_, other.count_);
   }

thrust/thrust/detail/contiguous_storage.h

@@ -179,13 +179,16 @@ class contiguous_storage
   _CCCL_HOST_DEVICE void propagate_allocator_dispatch(true_type, contiguous_storage& other);
 
   _CCCL_HOST_DEVICE void propagate_allocator_dispatch(false_type, contiguous_storage& other);
+
+  friend _CCCL_HOST_DEVICE constexpr void
+  swap(contiguous_storage& lhs, contiguous_storage& rhs) noexcept(noexcept(lhs.swap(rhs)))
+  {
+    lhs.swap(rhs);
+  }
 }; // end contiguous_storage
 
 } // namespace detail
 
-template <typename T, typename Alloc>
-_CCCL_HOST_DEVICE void swap(detail::contiguous_storage<T, Alloc>& lhs, detail::contiguous_storage<T, Alloc>& rhs);
-
 THRUST_NAMESPACE_END
 
 #include <thrust/detail/contiguous_storage.inl>

thrust/thrust/detail/contiguous_storage.inl

@@ -31,7 +31,8 @@
 #include <thrust/detail/allocator/fill_construct_range.h>
 #include <thrust/detail/allocator/value_initialize_range.h>
 #include <thrust/detail/contiguous_storage.h>
-#include <thrust/detail/swap.h>
+
+#include <cuda/std/utility>
 
 #include <stdexcept> // for std::runtime_error
 #include <utility> // for use of std::swap in the WAR below
@@ -193,13 +194,17 @@ _CCCL_HOST_DEVICE void contiguous_storage<T, Alloc>::deallocate() noexcept
 template <typename T, typename Alloc>
 _CCCL_HOST_DEVICE void contiguous_storage<T, Alloc>::swap(contiguous_storage& x)
 {
-  thrust::swap(m_begin, x.m_begin);
-  thrust::swap(m_size, x.m_size);
+  using ::cuda::std::swap;
+  swap(m_begin, x.m_begin);
+  swap(m_size, x.m_size);
 
+  // FIXME(bgruber): swap_allocators already swaps m_allocator, so we are swapping twice here !!
   swap_allocators(integral_constant<bool, allocator_traits<Alloc>::propagate_on_container_swap::value>(),
                   x.m_allocator);
-
-  thrust::swap(m_allocator, x.m_allocator);
+  // FIXME(bgruber): this should use ADL-two-step swap, but this creates an ambiguity with std::swap until
+  // https://github.com/NVIDIA/cccl/issues/2984 is resolved.
+  // swap(m_allocator, x.m_allocator);
+  ::cuda::std::swap(m_allocator, x.m_allocator);
 } // end contiguous_storage::swap()
 
 template <typename T, typename Alloc>
@@ -338,13 +343,18 @@ _CCCL_HOST_DEVICE void contiguous_storage<T, Alloc>::swap_allocators(true_type,
 template <typename T, typename Alloc>
 _CCCL_HOST_DEVICE void contiguous_storage<T, Alloc>::swap_allocators(false_type, Alloc& other)
 {
+  // FIXME(bgruber): it is really concerning, that swapping an allocator can throw. swap() should be noexcept in
+  // general.
   NV_IF_TARGET(NV_IS_DEVICE,
                (
                  // allocators must be equal when swapping containers with allocators that propagate on swap
                  assert(!is_allocator_not_equal(other));),
                (if (is_allocator_not_equal(other)) { throw allocator_mismatch_on_swap(); }));
-
-  thrust::swap(m_allocator, other);
+  // FIXME(bgruber): this should use ADL-two-step swap, but this creates an ambiguity with std::swap until
+  // https://github.com/NVIDIA/cccl/issues/2984 is resolved.
+  // using ::cuda::std::swap;
+  // swap(m_allocator, other);
+  ::cuda::std::swap(m_allocator, other);
 } // end contiguous_storage::swap_allocators()
 
 template <typename T, typename Alloc>
@@ -419,10 +429,4 @@ _CCCL_HOST_DEVICE void contiguous_storage<T, Alloc>::propagate_allocator_dispatc
 
 } // namespace detail
 
-template <typename T, typename Alloc>
-_CCCL_HOST_DEVICE void swap(detail::contiguous_storage<T, Alloc>& lhs, detail::contiguous_storage<T, Alloc>& rhs)
-{
-  lhs.swap(rhs);
-} // end swap()
-
 THRUST_NAMESPACE_END

thrust/thrust/detail/reference.h

@@ -491,8 +491,10 @@ class tagged_reference<void const, Tag>
  *  \param x The first \p tagged_reference of interest.
  *  \param y The second \p tagged_reference of interest.
  */
+// note: this is not a hidden friend, because we have template specializations of tagged_reference
 template <typename Element, typename Tag>
-_CCCL_HOST_DEVICE void swap(tagged_reference<Element, Tag>& x, tagged_reference<Element, Tag>& y)
+_CCCL_HOST_DEVICE void
+swap(tagged_reference<Element, Tag>& x, tagged_reference<Element, Tag>& y) noexcept(noexcept(x.swap(y)))
 {
   x.swap(y);
 }

thrust/thrust/detail/vector_base.h

@@ -38,6 +38,7 @@
 #include <thrust/iterator/reverse_iterator.h>
 
 #include <cuda/std/__iterator/iterator_traits.h>
+#include <cuda/std/utility>
 
 #include <initializer_list>
 #include <vector>
@@ -406,7 +407,12 @@ class vector_base
   /*! This method swaps the contents of this vector_base with another vector_base.
    *  \param v The vector_base with which to swap.
    */
-  void swap(vector_base& v);
+  constexpr void swap(vector_base& v)
+  {
+    using ::cuda::std::swap;
+    swap(m_storage, v.m_storage);
+    swap(m_size, v.m_size);
+  }
 
   /*! This method removes the element at position pos.
    *  \param pos The position of the element of interest.
@@ -551,18 +557,20 @@ class vector_base
   template <typename ForwardIterator>
   void
   allocate_and_copy(size_type requested_size, ForwardIterator first, ForwardIterator last, storage_type& new_storage);
-}; // end vector_base
 
-/*! This function assigns the contents of vector a to vector b and the
- *  contents of vector b to vector a.
- *
- *  \param a The first vector of interest. After completion, the contents
- *           of b will be returned here.
- *  \param b The second vector of interest. After completion, the contents
- *           of a will be returned here.
- */
-template <typename T, typename Alloc>
-void swap(vector_base<T, Alloc>& a, vector_base<T, Alloc>& b);
+  /*! This function assigns the contents of vector a to vector b and the
+   *  contents of vector b to vector a.
+   *
+   *  \param a The first vector of interest. After completion, the contents
+   *           of b will be returned here.
+   *  \param b The second vector of interest. After completion, the contents
+   *           of a will be returned here.
+   */
+  friend constexpr void swap(vector_base& a, vector_base& b) noexcept(noexcept(a.swap(b)))
+  {
+    a.swap(b);
+  }
+}; // end vector_base
 
 /*! This operator allows comparison between two vectors.
  *  \param lhs The first \p vector to compare.

thrust/thrust/detail/vector_base.inl

@@ -585,13 +585,6 @@ typename vector_base<T, Alloc>::iterator vector_base<T, Alloc>::erase(iterator f
   return first;
 } // end vector_base::erase()
 
-template <typename T, typename Alloc>
-void vector_base<T, Alloc>::swap(vector_base& v)
-{
-  thrust::swap(m_storage, v.m_storage);
-  thrust::swap(m_size, v.m_size);
-} // end vector_base::swap()
-
 template <typename T, typename Alloc>
 void vector_base<T, Alloc>::assign(size_type n, const T& x)
 {
@@ -1110,12 +1103,6 @@ void vector_base<T, Alloc>::allocate_and_copy(
   } // end catch
 } // end vector_base::allocate_and_copy()
 
-template <typename T, typename Alloc>
-void swap(vector_base<T, Alloc>& a, vector_base<T, Alloc>& b)
-{
-  a.swap(b);
-} // end swap()
-
 // iterator tags match
 template <typename InputIterator1, typename InputIterator2>
 bool vector_equal(InputIterator1 first1, InputIterator1 last1, InputIterator2 first2, thrust::detail::true_type)

thrust/thrust/device_reference.h

@@ -947,17 +947,16 @@ class device_reference : public thrust::reference<T, thrust::device_ptr<T>, thru
      */
     device_reference &operator^=(const T &rhs);
 #endif // end doxygen-only members
-}; // end device_reference
 
-/*! swaps the value of one \p device_reference with another.
- *  \p x The first \p device_reference of interest.
- *  \p y The second \p device_reference of interest.
- */
-template <typename T>
-_CCCL_HOST_DEVICE void swap(device_reference<T>& x, device_reference<T>& y)
-{
-  x.swap(y);
-}
+  /*! swaps the value of one \p device_reference with another.
+   *  \p x The first \p device_reference of interest.
+   *  \p y The second \p device_reference of interest.
+   */
+  friend _CCCL_HOST_DEVICE constexpr void swap(device_reference& x, device_reference& y) noexcept(noexcept(x.swap(y)))
+  {
+    x.swap(y);
+  }
+}; // end device_reference
 
 // declare these methods for the purpose of Doxygenating them
 // they actually are defined for a base class

thrust/thrust/device_vector.h

@@ -271,6 +271,15 @@ class device_vector : public detail::vector_base<T, Alloc>
       : Parent(first, last, alloc)
   {}
 
+  /*! Exchanges the values of two vectors.
+   *  \p x The first \p device_vector of interest.
+   *  \p y The second \p device_vector of interest.
+   */
+  friend constexpr void swap(device_vector& a, device_vector& b) noexcept(noexcept(a.swap(b)))
+  {
+    a.swap(b);
+  }
+
 // declare these members for the purpose of Doxygenating them
 // they actually exist in a base class
 #if 0
@@ -532,16 +541,6 @@ class device_vector : public detail::vector_base<T, Alloc>
 #endif // end doxygen-only members
 };
 
-/*! Exchanges the values of two vectors.
- *  \p x The first \p device_vector of interest.
- *  \p y The second \p device_vector of interest.
- */
-template <typename T, typename Alloc>
-void swap(device_vector<T, Alloc>& a, device_vector<T, Alloc>& b)
-{
-  a.swap(b);
-}
-
 /*! \} // containres
  */
 

thrust/thrust/host_vector.h

@@ -532,17 +532,16 @@ class host_vector : public detail::vector_base<T, Alloc>
      */
     allocator_type get_allocator() const;
 #endif // end doxygen-only members
-};
 
-/*! Exchanges the values of two vectors.
- *  \p x The first \p host_vector of interest.
- *  \p y The second \p host_vector of interest.
- */
-template <typename T, typename Alloc>
-void swap(host_vector<T, Alloc>& a, host_vector<T, Alloc>& b)
-{
-  a.swap(b);
-}
+  /*! Exchanges the values of two vectors.
+   *  \p x The first \p host_vector of interest.
+   *  \p y The second \p host_vector of interest.
+   */
+  friend void swap(host_vector& a, host_vector& b) noexcept(noexcept(a.swap(b)))
+  {
+    a.swap(b);
+  }
+};
 
 /*! \}
  */

thrust/thrust/optional.h

@@ -1548,7 +1548,7 @@ class optional
     {
       if (rhs.has_value())
       {
-        using thrust::swap;
+        using ::cuda::std::swap;
         swap(**this, *rhs);
       }
       else

thrust/thrust/swap.h

@@ -31,6 +31,8 @@
 #endif // no system header
 #include <thrust/detail/execution_policy.h>
 
+#include <cuda/std/utility>
+
 THRUST_NAMESPACE_BEGIN
 
 /*! \addtogroup utility
@@ -66,8 +68,7 @@ THRUST_NAMESPACE_BEGIN
  *  // x == 2, y == 1
  *  \endcode
  */
-template <typename Assignable1, typename Assignable2>
-_CCCL_HOST_DEVICE inline void swap(Assignable1& a, Assignable2& b);
+using ::cuda::std::swap;
 
 /*! \} // swap
  */
@@ -182,4 +183,4 @@ ForwardIterator2 swap_ranges(ForwardIterator1 first1, ForwardIterator1 last1, Fo
 
 THRUST_NAMESPACE_END
 
-#include <thrust/detail/swap.inl>
+#include <thrust/detail/swap_ranges.inl>