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

cub/test/catch2_segmented_sort_helper.cuh

@@ -42,6 +42,7 @@
 #include <cuda/std/limits>
 #include <cuda/std/tuple>
 #include <cuda/std/type_traits>
+#include <cuda/std/utility>
 
 #include <cstdio>
 
@@ -290,7 +291,7 @@ private:
           {
             if (current_value == d_values[probe_unchecked_idx])
             {
-              using thrust::swap;
+              using ::cuda::std::swap;
               swap(d_values[probe_unchecked_idx], d_values[unchecked_values_for_current_dup_key_begin]);
               unchecked_values_for_current_dup_key_begin++;
               break;
@@ -520,7 +521,7 @@ struct unstable_segmented_value_checker
           if (current_value == test_values[probe_unchecked_idx])
           {
             // Swap the found value out of the unchecked region to reduce the search space in future iterations:
-            using thrust::swap;
+            using ::cuda::std::swap;
             swap(test_values[probe_unchecked_idx], test_values[unchecked_values_for_current_dup_key_begin]);
             unchecked_values_for_current_dup_key_begin++;
             break;

docs/thrust/api_docs/utility.rst

@@ -9,5 +9,4 @@ Utility
 
    utility/pair
    utility/tuple
-   utility/swap
    utility/type_traits

libcudacxx/include/cuda/std/__type_traits/is_swappable.h

@@ -32,6 +32,7 @@
 #include <cuda/std/__type_traits/is_same.h>
 #include <cuda/std/__type_traits/is_void.h>
 #include <cuda/std/__type_traits/nat.h>
+#include <cuda/std/__type_traits/type_identity.h>
 #include <cuda/std/__utility/declval.h>
 #include <cuda/std/cstddef>
 
@@ -100,8 +101,11 @@ using __swap_result_t _CCCL_NODEBUG_ALIAS =
   enable_if_t<__detect_adl_swap::__can_define_swap<_Tp>::value && _CCCL_TRAIT(is_move_constructible, _Tp)
               && _CCCL_TRAIT(is_move_assignable, _Tp)>;
 
+// we use type_identity_t<_Tp> as second parameter, to avoid ambiguity with std::swap, which will thus be preferred by
+// overload resolution (which is ok since std::swap is only considered when explicitly called, or found by ADL for types
+// from std::)
 template <class _Tp>
-_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 __swap_result_t<_Tp> swap(_Tp& __x, _Tp& __y) noexcept(
+_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 __swap_result_t<_Tp> swap(_Tp& __x, type_identity_t<_Tp>& __y) noexcept(
   _CCCL_TRAIT(is_nothrow_move_constructible, _Tp) && _CCCL_TRAIT(is_nothrow_move_assignable, _Tp));
 
 template <class _Tp, size_t _Np>

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

@@ -31,9 +31,12 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
+// we use type_identity_t<_Tp> as second parameter, to avoid ambiguity with std::swap, which will thus be preferred by
+// overload resolution (which is ok since std::swap is only considered when explicitly called, or found by ADL for types
+// from std::)
 _CCCL_EXEC_CHECK_DISABLE
 template <class _Tp>
-_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 __swap_result_t<_Tp> swap(_Tp& __x, _Tp& __y) noexcept(
+_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 __swap_result_t<_Tp> swap(_Tp& __x, type_identity_t<_Tp>& __y) noexcept(
   _CCCL_TRAIT(is_nothrow_move_constructible, _Tp) && _CCCL_TRAIT(is_nothrow_move_assignable, _Tp))
 {
   _Tp __t(_CUDA_VSTD::move(__x));

libcudacxx/test/libcudacxx/std/utilities/utility/utility.swap/swap.pass.cpp

@@ -21,6 +21,7 @@
 #include "test_macros.h"
 
 #if !defined(TEST_COMPILER_NVRTC)
+#  include <memory>
 #  include <utility>
 #endif // !TEST_COMPILER_NVRTC
 
@@ -102,26 +103,58 @@ struct swap_with_friend
   __host__ __device__ friend void swap(swap_with_friend&, swap_with_friend&) {}
 };
 
+template <typename T>
 __host__ __device__ void test_ambiguous_std()
 {
-#if !defined(TEST_COMPILER_NVRTC)
   // clang-format off
   NV_IF_TARGET(NV_IS_HOST, (
+    // fully qualified calls
+    {
+      T i = {};
+      T j = {};
+      cuda::std::swap(i,j);
+    }
+  ))
+#if !defined(TEST_COMPILER_NVRTC)
+  NV_IF_TARGET(NV_IS_HOST, (
+    {
+      T i = {};
+      T j = {};
+      std::swap(i,j);
+    }
+  ))
+#endif // !TEST_COMPILER_NVRTC
+  NV_IF_TARGET(NV_IS_HOST, (
+    // ADL calls
+    {
+      T i = {};
+      T j = {};
+      swap(i,j);
+    }
+  ))
+#if !defined(TEST_COMPILER_NVRTC)
+  NV_IF_TARGET(NV_IS_HOST, (
+    {
+      T i = {};
+      T j = {};
+      using cuda::std::swap;
+      swap(i,j);
+    }
     {
-      cuda::std::pair<::std::pair<int, int>, int> i = {};
-      cuda::std::pair<::std::pair<int, int>, int> j = {};
+      T i = {};
+      T j = {};
+      using std::swap;
       swap(i,j);
     }
-    { // Ensure that we do not SFINAE swap out if there is a free function as that will take precedent
-      swap_with_friend<::std::pair<int, int>> with_friend;
-      cuda::std::swap(with_friend, with_friend);
+    {
+      T i = {};
+      T j = {};
+      using std::swap;
+      using cuda::std::swap;
+      swap(i,j);
     }
   ))
   // clang-format on
-#  if TEST_STD_VER >= 2014
-  static_assert(cuda::std::is_swappable<cuda::std::pair<::std::pair<int, int>, int>>::value, "");
-  static_assert(cuda::std::is_swappable<swap_with_friend<::std::pair<int, int>>>::value, "");
-#  endif // TEST_STD_VER >= 2014
 #endif // !TEST_COMPILER_NVRTC
 }
 
@@ -162,7 +195,20 @@ int main(int, char**)
   static_assert(test_swap_constexpr(), "");
 #endif // TEST_STD_VER >= 2014
 
-  test_ambiguous_std();
+  test_ambiguous_std<cuda::std::pair<int, int>>(); // has cuda::std::swap overload
+#if !defined(TEST_COMPILER_NVRTC)
+  test_ambiguous_std<::std::pair<int, int>>(); // has std::swap overload
+  test_ambiguous_std<cuda::std::pair<::std::pair<int, int>, int>>(); // has std:: and cuda::std as associated namespaces
+  test_ambiguous_std<::std::allocator<char>>(); // no std::swap overload
+
+  // Ensure that we do not SFINAE swap out if there is a free function as that will take precedent
+  test_ambiguous_std<swap_with_friend<::std::pair<int, int>>>();
+#endif // !TEST_COMPILER_NVRTC
+
+#if !defined(TEST_COMPILER_NVRTC) && TEST_STD_VER >= 2014
+  static_assert(cuda::std::is_swappable<cuda::std::pair<::std::pair<int, int>, int>>::value, "");
+  static_assert(cuda::std::is_swappable<swap_with_friend<::std::pair<int, int>>>::value, "");
+#endif // !defined(TEST_COMPILER_NVRTC) && TEST_STD_VER >= 2014
 
   return 0;
 }

thrust/testing/device_reference.cu

@@ -218,8 +218,9 @@ void TestDeviceReferenceSwap()
   ref1 = 7;
   ref2 = 13;
 
-  // test thrust::swap()
-  thrust::swap(ref1, ref2);
+  // test ADL two-step swap
+  using ::cuda::std::swap;
+  swap(ref1, ref2);
   ASSERT_EQUAL(13, ref1);
   ASSERT_EQUAL(7, ref2);
 

thrust/testing/pair.cu

@@ -286,7 +286,8 @@ void TestPairSwap()
   thrust::pair<int, int> a(x, y);
   thrust::pair<int, int> b(z, w);
 
-  thrust::swap(a, b);
+  using ::cuda::std::swap;
+  swap(a, b);
 
   ASSERT_EQUAL(z, a.first);
   ASSERT_EQUAL(w, a.second);

thrust/testing/swap_ranges.cu

@@ -133,7 +133,8 @@ struct type_with_swap
 
 inline _CCCL_HOST_DEVICE void swap(type_with_swap& a, type_with_swap& b)
 {
-  thrust::swap(a.m_x, b.m_x);
+  using ::cuda::std::swap;
+  swap(a.m_x, b.m_x);
   a.m_swapped = true;
   b.m_swapped = true;
 }

thrust/testing/tuple.cu

@@ -471,7 +471,8 @@ void TestTupleSwap()
   thrust::tuple<int, int, int> t1(a, b, c);
   thrust::tuple<int, int, int> t2(x, y, z);
 
-  thrust::swap(t1, t2);
+  using ::cuda::std::swap;
+  swap(t1, t2);
 
   ASSERT_EQUAL(x, thrust::get<0>(t1));
   ASSERT_EQUAL(y, thrust::get<1>(t1));

thrust/testing/vector_allocators.cu

@@ -252,14 +252,15 @@ void TestVectorAllocatorPropagateOnSwap()
 
   Vector v1(10, alloc1);
   Vector v2(17, alloc1);
-  thrust::swap(v1, v2);
+  using ::cuda::std::swap;
+  swap(v1, v2);
 
   ASSERT_EQUAL(v1.size(), 17u);
   ASSERT_EQUAL(v2.size(), 10u);
 
   Vector v3(15, alloc1);
   Vector v4(31, alloc2);
-  ASSERT_THROWS(thrust::swap(v3, v4), thrust::detail::allocator_mismatch_on_swap);
+  ASSERT_THROWS(swap(v3, v4), thrust::detail::allocator_mismatch_on_swap);
 }
 
 void TestVectorAllocatorPropagateOnSwapHost()

thrust/testing/vector_insert.cu

@@ -113,7 +113,8 @@ struct TestVectorRangeInsert
     size_t end   = n > 0 ? (size_t) h_src[n + 1] % n : 0;
     if (end < begin)
     {
-      thrust::swap(begin, end);
+      using ::cuda::std::swap;
+      swap(begin, end);
     }
 
     // choose insertion position at random

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,15 @@ 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 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
@@ -190,16 +191,18 @@ _CCCL_HOST_DEVICE void contiguous_storage<T, Alloc>::deallocate() noexcept
   } // end if
 } // end contiguous_storage::deallocate()
 
+_CCCL_EXEC_CHECK_DISABLE
 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);
+  swap(m_allocator, x.m_allocator);
 } // end contiguous_storage::swap()
 
 template <typename T, typename Alloc>
@@ -335,16 +338,19 @@ template <typename T, typename Alloc>
 _CCCL_HOST_DEVICE void contiguous_storage<T, Alloc>::swap_allocators(true_type, const Alloc&)
 {} // end contiguous_storage::swap_allocators()
 
+_CCCL_EXEC_CHECK_DISABLE
 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);
+  using ::cuda::std::swap;
+  swap(m_allocator, other);
 } // end contiguous_storage::swap_allocators()
 
 template <typename T, typename Alloc>
@@ -419,10 +425,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);
 }