diff --git a/libcudacxx/include/cuda/std/__type_traits/is_swappable.h b/libcudacxx/include/cuda/std/__type_traits/is_swappable.h
index 0764139cd39..87c129e67c0 100644
--- a/libcudacxx/include/cuda/std/__type_traits/is_swappable.h
+++ b/libcudacxx/include/cuda/std/__type_traits/is_swappable.h
@@ -22,6 +22,7 @@
 
 #include <cuda/std/__type_traits/add_lvalue_reference.h>
 #include <cuda/std/__type_traits/conditional.h>
+#include <cuda/std/__type_traits/disjunction.h>
 #include <cuda/std/__type_traits/enable_if.h>
 #include <cuda/std/__type_traits/is_move_assignable.h>
 #include <cuda/std/__type_traits/is_move_constructible.h>
@@ -36,6 +37,56 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
+// We need to detect whether there is already a free function swap that would end up being ambiguous.
+// This can happen when a type pulls in both namespace std and namespace cuda::std via ADL.
+// In that case we are always safe to just not do anything because that type must be host only.
+// However, we must be carefull to ensure that we still create the overload if there is just a hidden friend swap
+namespace __detect_hidden_friend_swap
+{
+// This will intentionally create an ambiguity with std::swap if that is find-able by ADL. But it will not interfere
+// with hidden friend swap
+template <class _Tp>
+_CCCL_HOST_DEVICE void swap(_Tp&, _Tp&);
+
+struct __hidden_friend_swap_found
+{};
+
+template <class _Tp>
+_LIBCUDACXX_INLINE_VISIBILITY auto __swap(_Tp& __lhs, _Tp& __rhs) -> decltype(swap(__lhs, __rhs));
+_LIBCUDACXX_INLINE_VISIBILITY auto __swap(...) -> __hidden_friend_swap_found;
+template <class _Tp>
+struct __has_hidden_friend_swap
+    : is_same<decltype(__detect_hidden_friend_swap::__swap(_CUDA_VSTD::declval<_Tp&>(), _CUDA_VSTD::declval<_Tp&>())),
+              void>
+{};
+} // namespace __detect_hidden_friend_swap
+
+namespace __detect_adl_swap
+{
+struct __no_adl_swap_found
+{};
+template <class _Tp>
+_LIBCUDACXX_INLINE_VISIBILITY auto __swap(_Tp& __lhs, _Tp& __rhs) -> decltype(swap(__lhs, __rhs));
+_LIBCUDACXX_INLINE_VISIBILITY auto __swap(...) -> __no_adl_swap_found;
+template <class _Tp>
+struct __has_no_adl_swap
+    : is_same<decltype(__detect_adl_swap::__swap(_CUDA_VSTD::declval<_Tp&>(), _CUDA_VSTD::declval<_Tp&>())),
+              __no_adl_swap_found>
+{};
+template <class _Tp, size_t _Np>
+struct __has_no_adl_swap_array
+    : is_same<
+        decltype(__detect_adl_swap::__swap(_CUDA_VSTD::declval<_Tp (&)[_Np]>(), _CUDA_VSTD::declval<_Tp (&)[_Np]>())),
+        __no_adl_swap_found>
+{};
+
+// We should only define swap if there is no ADL found function or it is a hidden friend
+template <class _Tp>
+struct __can_define_swap
+    : _Or<__has_no_adl_swap<_Tp>, __detect_hidden_friend_swap::__has_hidden_friend_swap<_Tp>>
+{};
+} // namespace __detect_adl_swap
+
 template <class _Tp>
 struct __is_swappable;
 template <class _Tp>
@@ -43,15 +94,17 @@ struct __is_nothrow_swappable;
 
 template <class _Tp>
 using __swap_result_t =
-  __enable_if_t<_LIBCUDACXX_TRAIT(is_move_constructible, _Tp) && _LIBCUDACXX_TRAIT(is_move_assignable, _Tp)>;
+  __enable_if_t<__detect_adl_swap::__can_define_swap<_Tp>::value && _LIBCUDACXX_TRAIT(is_move_constructible, _Tp)
+                && _LIBCUDACXX_TRAIT(is_move_assignable, _Tp)>;
 
 template <class _Tp>
 inline _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 __swap_result_t<_Tp> swap(_Tp& __x, _Tp& __y) noexcept(
   _LIBCUDACXX_TRAIT(is_nothrow_move_constructible, _Tp) && _LIBCUDACXX_TRAIT(is_nothrow_move_assignable, _Tp));
 
 template <class _Tp, size_t _Np>
-inline _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 __enable_if_t<__is_swappable<_Tp>::value>
-  swap(_Tp (&__a)[_Np], _Tp (&__b)[_Np]) noexcept(__is_nothrow_swappable<_Tp>::value);
+inline _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14
+  __enable_if_t<__detect_adl_swap::__has_no_adl_swap_array<_Tp, _Np>::value && __is_swappable<_Tp>::value>
+    swap(_Tp (&__a)[_Np], _Tp (&__b)[_Np]) noexcept(__is_nothrow_swappable<_Tp>::value);
 
 namespace __detail
 {
diff --git a/libcudacxx/include/cuda/std/__utility/swap.h b/libcudacxx/include/cuda/std/__utility/swap.h
index fbf5a7593b9..5c2143da524 100644
--- a/libcudacxx/include/cuda/std/__utility/swap.h
+++ b/libcudacxx/include/cuda/std/__utility/swap.h
@@ -41,8 +41,9 @@ inline _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 __swap_result_t<_Tp>
 }
 
 template <class _Tp, size_t _Np>
-inline _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 __enable_if_t<__is_swappable<_Tp>::value>
-swap(_Tp (&__a)[_Np], _Tp (&__b)[_Np]) noexcept(__is_nothrow_swappable<_Tp>::value)
+inline _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14
+  __enable_if_t<__detect_adl_swap::__has_no_adl_swap_array<_Tp, _Np>::value && __is_swappable<_Tp>::value>
+  swap(_Tp (&__a)[_Np], _Tp (&__b)[_Np]) noexcept(__is_nothrow_swappable<_Tp>::value)
 {
   for (size_t __i = 0; __i != _Np; ++__i)
   {
diff --git a/libcudacxx/test/libcudacxx/std/utilities/utility/utility.swap/swap.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/utility/utility.swap/swap.pass.cpp
new file mode 100644
index 00000000000..2ba70fd4b3c
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/utility/utility.swap/swap.pass.cpp
@@ -0,0 +1,152 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// <utility>
+
+// template<class T>
+//   requires MoveAssignable<T> && MoveConstructible<T>
+//   void
+//   swap(T& a, T& b);
+
+#include <cuda/std/__memory_>
+#include <cuda/std/cassert>
+#include <cuda/std/type_traits>
+#include <cuda/std/utility>
+
+#include "test_macros.h"
+
+#if !defined(TEST_COMPILER_NVRTC)
+#  include <utility>
+#endif // !TEST_COMPILER_NVRTC
+
+struct CopyOnly
+{
+  __host__ __device__ CopyOnly() {}
+  __host__ __device__ CopyOnly(CopyOnly const&) noexcept {}
+  __host__ __device__ CopyOnly& operator=(CopyOnly const&)
+  {
+    return *this;
+  }
+};
+
+struct MoveOnly
+{
+  __host__ __device__ MoveOnly() {}
+  __host__ __device__ MoveOnly(MoveOnly&&) {}
+  __host__ __device__ MoveOnly& operator=(MoveOnly&&) noexcept
+  {
+    return *this;
+  }
+};
+
+struct NoexceptMoveOnly
+{
+  __host__ __device__ NoexceptMoveOnly() {}
+  __host__ __device__ NoexceptMoveOnly(NoexceptMoveOnly&&) noexcept {}
+  __host__ __device__ NoexceptMoveOnly& operator=(NoexceptMoveOnly&&) noexcept
+  {
+    return *this;
+  }
+};
+
+struct NotMoveConstructible
+{
+  __host__ __device__ NotMoveConstructible& operator=(NotMoveConstructible&&)
+  {
+    return *this;
+  }
+
+private:
+  __host__ __device__ NotMoveConstructible(NotMoveConstructible&&);
+};
+
+struct NotMoveAssignable
+{
+  __host__ __device__ NotMoveAssignable(NotMoveAssignable&&);
+
+private:
+  __host__ __device__ NotMoveAssignable& operator=(NotMoveAssignable&&);
+};
+
+template <class Tp>
+__host__ __device__ auto can_swap_test(int)
+  -> decltype(cuda::std::swap(cuda::std::declval<Tp>(), cuda::std::declval<Tp>()));
+
+template <class Tp>
+__host__ __device__ auto can_swap_test(...) -> cuda::std::false_type;
+
+template <class Tp>
+__host__ __device__ constexpr bool can_swap()
+{
+  return cuda::std::is_same<decltype(can_swap_test<Tp>(0)), void>::value;
+}
+
+#if TEST_STD_VER >= 2014
+__host__ __device__ constexpr bool test_swap_constexpr()
+{
+  int i = 1;
+  int j = 2;
+  cuda::std::swap(i, j);
+  return i == 2 && j == 1;
+}
+#endif // TEST_STD_VER >= 2014
+
+__host__ __device__ void test_ambiguous_std()
+{
+#if !defined(TEST_COMPILER_NVRTC)
+  // clang-format off
+  NV_IF_TARGET(NV_IS_HOST, (
+    cuda::std::pair<::std::pair<int, int>, int> i = {};
+    cuda::std::pair<::std::pair<int, int>, int> j = {};
+    swap(i,j);
+  ))
+  // clang-format on
+#endif // !TEST_COMPILER_NVRTC
+}
+
+int main(int, char**)
+{
+  {
+    int i = 1;
+    int j = 2;
+    cuda::std::swap(i, j);
+    assert(i == 2);
+    assert(j == 1);
+  }
+  {
+    cuda::std::unique_ptr<int> i(new int(1));
+    cuda::std::unique_ptr<int> j(new int(2));
+    cuda::std::swap(i, j);
+    assert(*i == 2);
+    assert(*j == 1);
+  }
+  {
+    // test that the swap
+    static_assert(can_swap<CopyOnly&>(), "");
+    static_assert(can_swap<MoveOnly&>(), "");
+    static_assert(can_swap<NoexceptMoveOnly&>(), "");
+
+    static_assert(!can_swap<NotMoveConstructible&>(), "");
+    static_assert(!can_swap<NotMoveAssignable&>(), "");
+
+    CopyOnly c;
+    MoveOnly m;
+    NoexceptMoveOnly nm;
+    static_assert(!noexcept(cuda::std::swap(c, c)), "");
+    static_assert(!noexcept(cuda::std::swap(m, m)), "");
+    static_assert(noexcept(cuda::std::swap(nm, nm)), "");
+  }
+
+#if TEST_STD_VER >= 2014
+  static_assert(test_swap_constexpr());
+#endif // TEST_STD_VER >= 2014
+
+  test_ambiguous_std();
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/utility/utility.swap/swap_array.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/utility/utility.swap/swap_array.pass.cpp
new file mode 100644
index 00000000000..b34144f265d
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/utility/utility.swap/swap_array.pass.cpp
@@ -0,0 +1,150 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// <utility>
+
+// template<ValueType T, size_t N>
+//   requires Swappable<T>
+//   void
+//   swap(T (&a)[N], T (&b)[N]);
+
+#include <cuda/std/__algorithm_>
+#include <cuda/std/__memory_>
+#include <cuda/std/array>
+#include <cuda/std/cassert>
+#include <cuda/std/type_traits>
+#include <cuda/std/utility>
+
+#include "test_macros.h"
+
+#if !defined(TEST_COMPILER_NVRTC)
+#  include <utility>
+#endif // !TEST_COMPILER_NVRTC
+
+struct CopyOnly
+{
+  __host__ __device__ CopyOnly() {}
+  __host__ __device__ CopyOnly(CopyOnly const&) noexcept {}
+  __host__ __device__ CopyOnly& operator=(CopyOnly const&)
+  {
+    return *this;
+  }
+};
+
+struct NoexceptMoveOnly
+{
+  __host__ __device__ NoexceptMoveOnly() {}
+  __host__ __device__ NoexceptMoveOnly(NoexceptMoveOnly&&) noexcept {}
+  __host__ __device__ NoexceptMoveOnly& operator=(NoexceptMoveOnly&&) noexcept
+  {
+    return *this;
+  }
+};
+
+struct NotMoveConstructible
+{
+  __host__ __device__ NotMoveConstructible() {}
+  __host__ __device__ NotMoveConstructible& operator=(NotMoveConstructible&&)
+  {
+    return *this;
+  }
+
+private:
+  __host__ __device__ NotMoveConstructible(NotMoveConstructible&&);
+};
+
+template <class Tp>
+__host__ __device__ auto can_swap_test(int)
+  -> decltype(cuda::std::swap(cuda::std::declval<Tp>(), cuda::std::declval<Tp>()));
+
+template <class Tp>
+__host__ __device__ auto can_swap_test(...) -> cuda::std::false_type;
+
+template <class Tp>
+__host__ __device__ constexpr bool can_swap()
+{
+  return cuda::std::is_same<decltype(can_swap_test<Tp>(0)), void>::value;
+}
+
+#if TEST_STD_VER >= 2014
+__host__ __device__ constexpr bool test_swap_constexpr()
+{
+  int i[3] = {1, 2, 3};
+  int j[3] = {4, 5, 6};
+  cuda::std::swap(i, j);
+  return i[0] == 4 && i[1] == 5 && i[2] == 6 && j[0] == 1 && j[1] == 2 && j[2] == 3;
+}
+#endif // TEST_STD_VER >= 2014
+
+__host__ __device__ void test_ambiguous_std()
+{
+#if !defined(TEST_COMPILER_NVRTC)
+  // clang-format off
+  NV_IF_TARGET(NV_IS_HOST, (
+    cuda::std::pair<::std::pair<int, int>, int> i[3] = {};
+    cuda::std::pair<::std::pair<int, int>, int> j[3] = {};
+    swap(i,j);
+  ))
+  // clang-format on
+#endif // !TEST_COMPILER_NVRTC
+}
+
+int main(int, char**)
+{
+  {
+    int i[3] = {1, 2, 3};
+    int j[3] = {4, 5, 6};
+    cuda::std::swap(i, j);
+    assert(i[0] == 4);
+    assert(i[1] == 5);
+    assert(i[2] == 6);
+    assert(j[0] == 1);
+    assert(j[1] == 2);
+    assert(j[2] == 3);
+  }
+  {
+    cuda::std::unique_ptr<int> i[3];
+    for (int k = 0; k < 3; ++k)
+    {
+      i[k].reset(new int(k + 1));
+    }
+    cuda::std::unique_ptr<int> j[3];
+    for (int k = 0; k < 3; ++k)
+    {
+      j[k].reset(new int(k + 4));
+    }
+    cuda::std::swap(i, j);
+    assert(*i[0] == 4);
+    assert(*i[1] == 5);
+    assert(*i[2] == 6);
+    assert(*j[0] == 1);
+    assert(*j[1] == 2);
+    assert(*j[2] == 3);
+  }
+  {
+    using CA = CopyOnly[42];
+    using MA = NoexceptMoveOnly[42];
+    using NA = NotMoveConstructible[42];
+    static_assert(can_swap<CA&>(), "");
+    static_assert(can_swap<MA&>(), "");
+    static_assert(!can_swap<NA&>(), "");
+
+    CA ca;
+    MA ma;
+    static_assert(!noexcept(cuda::std::swap(ca, ca)), "");
+    static_assert(noexcept(cuda::std::swap(ma, ma)), "");
+  }
+
+#if TEST_STD_VER >= 2014
+  static_assert(test_swap_constexpr());
+#endif // TEST_STD_VER >= 2014
+
+  test_ambiguous_std();
+
+  return 0;
+}