diff --git a/thrust/thrust/allocate_unique.h b/thrust/thrust/allocate_unique.h
index 100d6541118..1ec032acdda 100644
--- a/thrust/thrust/allocate_unique.h
+++ b/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_);
   }
diff --git a/thrust/thrust/detail/contiguous_storage.h b/thrust/thrust/detail/contiguous_storage.h
index 41f3f5f0265..0092f0a1f61 100644
--- a/thrust/thrust/detail/contiguous_storage.h
+++ b/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>
diff --git a/thrust/thrust/detail/contiguous_storage.inl b/thrust/thrust/detail/contiguous_storage.inl
index d886f05efbe..a5b67bca8f8 100644
--- a/thrust/thrust/detail/contiguous_storage.inl
+++ b/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
diff --git a/thrust/thrust/detail/reference.h b/thrust/thrust/detail/reference.h
index 88fdf3d87d6..9685f141e40 100644
--- a/thrust/thrust/detail/reference.h
+++ b/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);
 }
diff --git a/thrust/thrust/detail/swap.h b/thrust/thrust/detail/swap.h
deleted file mode 100644
index 14fdd04c095..00000000000
--- a/thrust/thrust/detail/swap.h
+++ /dev/null
@@ -1,40 +0,0 @@
-/*
- *  Copyright 2008-2013 NVIDIA Corporation
- *
- *  Licensed under the Apache License, Version 2.0 (the "License");
- *  you may not use this file except in compliance with the License.
- *  You may obtain a copy of the License at
- *
- *      http://www.apache.org/licenses/LICENSE-2.0
- *
- *  Unless required by applicable law or agreed to in writing, software
- *  distributed under the License is distributed on an "AS IS" BASIS,
- *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- *  See the License for the specific language governing permissions and
- *  limitations under the License.
- */
-
-#pragma once
-
-#include <thrust/detail/config.h>
-
-#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
-#  pragma GCC system_header
-#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
-#  pragma clang system_header
-#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
-#  pragma system_header
-#endif // no system header
-
-THRUST_NAMESPACE_BEGIN
-
-_CCCL_EXEC_CHECK_DISABLE
-template <typename Assignable1, typename Assignable2>
-_CCCL_HOST_DEVICE inline void swap(Assignable1& a, Assignable2& b)
-{
-  Assignable1 temp = a;
-  a                = b;
-  b                = temp;
-} // end swap()
-
-THRUST_NAMESPACE_END
diff --git a/thrust/thrust/detail/swap.inl b/thrust/thrust/detail/swap.inl
deleted file mode 100644
index 38b4cd92f82..00000000000
--- a/thrust/thrust/detail/swap.inl
+++ /dev/null
@@ -1,31 +0,0 @@
-/*
- *  Copyright 2008-2013 NVIDIA Corporation
- *
- *  Licensed under the Apache License, Version 2.0 (the "License");
- *  you may not use this file except in compliance with the License.
- *  You may obtain a copy of the License at
- *
- *      http://www.apache.org/licenses/LICENSE-2.0
- *
- *  Unless required by applicable law or agreed to in writing, software
- *  distributed under the License is distributed on an "AS IS" BASIS,
- *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- *  See the License for the specific language governing permissions and
- *  limitations under the License.
- */
-
-#pragma once
-
-#include <thrust/detail/config.h>
-
-#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
-#  pragma GCC system_header
-#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
-#  pragma clang system_header
-#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
-#  pragma system_header
-#endif // no system header
-
-#include <thrust/detail/swap.h>
-#include <thrust/detail/swap_ranges.inl>
-#include <thrust/swap.h>
diff --git a/thrust/thrust/detail/vector_base.h b/thrust/thrust/detail/vector_base.h
index 3c0dabac418..0adfac1a6f2 100644
--- a/thrust/thrust/detail/vector_base.h
+++ b/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);
+  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 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.
diff --git a/thrust/thrust/detail/vector_base.inl b/thrust/thrust/detail/vector_base.inl
index ee06ccb6f09..2c1e6edfdb6 100644
--- a/thrust/thrust/detail/vector_base.inl
+++ b/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)
diff --git a/thrust/thrust/device_reference.h b/thrust/thrust/device_reference.h
index 545d5449bee..cfbcbdc3314 100644
--- a/thrust/thrust/device_reference.h
+++ b/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 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
diff --git a/thrust/thrust/device_vector.h b/thrust/thrust/device_vector.h
index a63def5ba89..e009eb98567 100644
--- a/thrust/thrust/device_vector.h
+++ b/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 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
  */
 
diff --git a/thrust/thrust/host_vector.h b/thrust/thrust/host_vector.h
index 6ede3b41e29..ad4ddd90016 100644
--- a/thrust/thrust/host_vector.h
+++ b/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);
+  }
+};
 
 /*! \}
  */
diff --git a/thrust/thrust/optional.h b/thrust/thrust/optional.h
index dbee5ebda24..788f6597953 100644
--- a/thrust/thrust/optional.h
+++ 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
diff --git a/thrust/thrust/swap.h b/thrust/thrust/swap.h
index 420be8f526c..89e0c618ce5 100644
--- a/thrust/thrust/swap.h
+++ b/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>
diff --git a/thrust/thrust/system/cuda/detail/iter_swap.h b/thrust/thrust/system/cuda/detail/iter_swap.h
index 2aee2a31f87..6cd6d6a19fd 100644
--- a/thrust/thrust/system/cuda/detail/iter_swap.h
+++ b/thrust/thrust/system/cuda/detail/iter_swap.h
@@ -34,6 +34,8 @@
 #  include <thrust/swap.h>
 #  include <thrust/system/cuda/detail/execution_policy.h>
 
+#  include <cuda/std/utility>
+
 #  include <nv/target>
 
 THRUST_NAMESPACE_BEGIN
@@ -53,7 +55,7 @@ inline _CCCL_HOST_DEVICE void iter_swap(thrust::cuda::execution_policy<DerivedPo
 
     _CCCL_DEVICE inline static void device_path(Pointer1 a, Pointer2 b)
     {
-      using thrust::swap;
+      using ::cuda::std::swap;
       swap(*thrust::raw_pointer_cast(a), *thrust::raw_pointer_cast(b));
     }
   };
diff --git a/thrust/thrust/system/cuda/detail/swap_ranges.h b/thrust/thrust/system/cuda/detail/swap_ranges.h
index e1d642574c0..fe63d16a0ff 100644
--- a/thrust/thrust/system/cuda/detail/swap_ranges.h
+++ b/thrust/thrust/system/cuda/detail/swap_ranges.h
@@ -43,6 +43,8 @@
 #  include <thrust/system/cuda/detail/parallel_for.h>
 #  include <thrust/system/cuda/detail/transform.h>
 
+#  include <cuda/std/utility>
+
 #  include <iterator>
 
 THRUST_NAMESPACE_BEGIN
@@ -71,14 +73,10 @@ struct swap_f
   template <class Size>
   void THRUST_DEVICE_FUNCTION operator()(Size idx)
   {
+    // TODO(bgruber): this should probably use ::cuda::std::iter_swap(items1 + idx, items2 + idx);
     value1_type item1 = items1[idx];
     value2_type item2 = items2[idx];
-    // XXX thrust::swap is buggy
-    // if reference_type of ItemIt1/ItemsIt2
-    // is a proxy reference, then KABOOM!
-    // to avoid this, just copy the value first before swap
-    // *todo* specialize on real & proxy references
-    using thrust::swap;
+    using ::cuda::std::swap;
     swap(item1, item2);
     items1[idx] = item1;
     items2[idx] = item2;
diff --git a/thrust/thrust/system/detail/generic/swap_ranges.inl b/thrust/thrust/system/detail/generic/swap_ranges.inl
index 1a626a38acf..a120bd6209e 100644
--- a/thrust/thrust/system/detail/generic/swap_ranges.inl
+++ b/thrust/thrust/system/detail/generic/swap_ranges.inl
@@ -31,6 +31,7 @@
 #include <thrust/system/detail/generic/swap_ranges.h>
 #include <thrust/tuple.h>
 
+#include <cuda/std/utility>
 THRUST_NAMESPACE_BEGIN
 namespace system
 {
@@ -49,7 +50,7 @@ struct swap_pair_elements
   _CCCL_HOST_DEVICE void operator()(Tuple t)
   {
     // use unqualified swap to allow ADL to catch any user-defined swap
-    using thrust::swap;
+    using ::cuda::std::swap;
     swap(thrust::get<0>(t), thrust::get<1>(t));
   }
 }; // end swap_pair_elements
diff --git a/thrust/thrust/system/detail/sequential/iter_swap.h b/thrust/thrust/system/detail/sequential/iter_swap.h
index 1d2e33efb42..6789ee9d935 100644
--- a/thrust/thrust/system/detail/sequential/iter_swap.h
+++ b/thrust/thrust/system/detail/sequential/iter_swap.h
@@ -26,9 +26,10 @@
 #  pragma system_header
 #endif // no system header
 #include <thrust/detail/raw_pointer_cast.h>
-#include <thrust/detail/swap.h>
 #include <thrust/system/detail/sequential/execution_policy.h>
 
+#include <cuda/std/utility>
+
 THRUST_NAMESPACE_BEGIN
 namespace system
 {
@@ -40,7 +41,7 @@ namespace sequential
 template <typename DerivedPolicy, typename Pointer1, typename Pointer2>
 _CCCL_HOST_DEVICE void iter_swap(sequential::execution_policy<DerivedPolicy>&, Pointer1 a, Pointer2 b)
 {
-  using thrust::swap;
+  using ::cuda::std::swap;
   swap(*thrust::raw_pointer_cast(a), *thrust::raw_pointer_cast(b));
 } // end iter_swap()
 
diff --git a/thrust/thrust/system/detail/sequential/partition.h b/thrust/thrust/system/detail/sequential/partition.h
index 3a869a8eead..cc744bb8866 100644
--- a/thrust/thrust/system/detail/sequential/partition.h
+++ b/thrust/thrust/system/detail/sequential/partition.h
@@ -51,18 +51,17 @@ namespace detail
 namespace sequential
 {
 
+// TODO(bgruber): we should make this an alias of cuda::std::iter_swap
 _CCCL_EXEC_CHECK_DISABLE
 template <typename ForwardIterator1, typename ForwardIterator2>
 _CCCL_HOST_DEVICE void iter_swap(ForwardIterator1 iter1, ForwardIterator2 iter2)
 {
-  // XXX this isn't correct because it doesn't use thrust::swap
-  using namespace thrust::detail;
-
+  // note: we cannot use swap(*iter1, *iter2) here, because the reference_type's could be proxy references, for which
+  // swap() is not guaranteed to work
   using T = typename thrust::iterator_value<ForwardIterator1>::type;
-
-  T temp = *iter1;
-  *iter1 = *iter2;
-  *iter2 = temp;
+  T temp  = ::cuda::std::move(*iter1);
+  *iter1  = ::cuda::std::move(*iter2);
+  *iter2  = ::cuda::std::move(temp);
 }
 
 _CCCL_EXEC_CHECK_DISABLE