Use cuda::std::addressof in Thrust

NVIDIA · Jan 13, 2025 · 6e31a64 · 6e31a64
1 parent cc7c1bb
commit 6e31a64
Show file tree

Hide file tree

Showing 4 changed files with 28 additions and 38 deletions.
diff --git a/thrust/thrust/addressof.h b/thrust/thrust/addressof.h
@@ -15,20 +15,8 @@
 #  pragma system_header
 #endif // no system header
 
-#include <thrust/detail/memory_wrapper.h>
+#include <cuda/std/__memory/addressof.h>
 
 THRUST_NAMESPACE_BEGIN
-
-///////////////////////////////////////////////////////////////////////////////
-
-/*! Obtains the actual address of the object or function arg, even in presence of overloaded operator&.
- */
-template <typename T>
-_CCCL_HOST_DEVICE T* addressof(T& arg)
-{
-  return reinterpret_cast<T*>(&const_cast<char&>(reinterpret_cast<const volatile char&>(arg)));
-}
-
-///////////////////////////////////////////////////////////////////////////////
-
+using ::cuda::std::addressof;
 THRUST_NAMESPACE_END
diff --git a/thrust/thrust/detail/memory_algorithms.h b/thrust/thrust/detail/memory_algorithms.h
@@ -17,12 +17,13 @@
 #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
 #  pragma system_header
 #endif // no system header
-#include <thrust/addressof.h>
 #include <thrust/detail/allocator/allocator_traits.h>
 #include <thrust/detail/memory_wrapper.h>
 #include <thrust/detail/type_traits.h>
 #include <thrust/iterator/iterator_traits.h>
 
+#include <cuda/std/__memory/addressof.h>
+
 #include <new>
 #include <utility>
 
@@ -54,7 +55,7 @@ _CCCL_HOST_DEVICE ForwardIt destroy(ForwardIt first, ForwardIt last) noexcept
 {
   for (; first != last; ++first)
   {
-    destroy_at(addressof(*first));
+    destroy_at(::cuda::std::addressof(*first));
   }
 
   return first;
@@ -71,7 +72,7 @@ _CCCL_HOST_DEVICE ForwardIt destroy(Allocator const& alloc, ForwardIt first, For
 
   for (; first != last; ++first)
   {
-    destroy_at(alloc_T, addressof(*first));
+    destroy_at(alloc_T, ::cuda::std::addressof(*first));
   }
 
   return first;
@@ -82,7 +83,7 @@ _CCCL_HOST_DEVICE ForwardIt destroy_n(ForwardIt first, Size n) noexcept
 {
   for (; n > 0; (void) ++first, --n)
   {
-    destroy_at(addressof(*first));
+    destroy_at(::cuda::std::addressof(*first));
   }
 
   return first;
@@ -99,7 +100,7 @@ _CCCL_HOST_DEVICE ForwardIt destroy_n(Allocator const& alloc, ForwardIt first, S
 
   for (; n > 0; (void) ++first, --n)
   {
-    destroy_at(alloc_T, addressof(*first));
+    destroy_at(alloc_T, ::cuda::std::addressof(*first));
   }
 
   return first;
@@ -119,13 +120,13 @@ _CCCL_HOST_DEVICE void uninitialized_construct(ForwardIt first, ForwardIt last,
       try {
         for (; current != last; ++current)
         {
-          ::new (static_cast<void*>(addressof(*current))) T(args...);
+          ::new (static_cast<void*>(::cuda::std::addressof(*current))) T(args...);
         }
       } catch (...) {
         destroy(first, current);
         throw;
       }),
-    (for (; current != last; ++current) { ::new (static_cast<void*>(addressof(*current))) T(args...); }));
+    (for (; current != last; ++current) { ::new (static_cast<void*>(::cuda::std::addressof(*current))) T(args...); }));
 }
 
 template <typename Allocator, typename ForwardIt, typename... Args>
@@ -146,13 +147,13 @@ void uninitialized_construct_with_allocator(Allocator const& alloc, ForwardIt fi
       try {
         for (; current != last; ++current)
         {
-          traits::construct(alloc_T, addressof(*current), args...);
+          traits::construct(alloc_T, ::cuda::std::addressof(*current), args...);
         }
       } catch (...) {
         destroy(alloc_T, first, current);
         throw;
       }),
-    (for (; current != last; ++current) { traits::construct(alloc_T, addressof(*current), args...); }));
+    (for (; current != last; ++current) { traits::construct(alloc_T, ::cuda::std::addressof(*current), args...); }));
 }
 
 template <typename ForwardIt, typename Size, typename... Args>
@@ -169,13 +170,13 @@ void uninitialized_construct_n(ForwardIt first, Size n, Args const&... args)
       try {
         for (; n > 0; ++current, --n)
         {
-          ::new (static_cast<void*>(addressof(*current))) T(args...);
+          ::new (static_cast<void*>(::cuda::std::addressof(*current))) T(args...);
         }
       } catch (...) {
         destroy(first, current);
         throw;
       }),
-    (for (; n > 0; ++current, --n) { ::new (static_cast<void*>(addressof(*current))) T(args...); }));
+    (for (; n > 0; ++current, --n) { ::new (static_cast<void*>(::cuda::std::addressof(*current))) T(args...); }));
 }
 
 template <typename Allocator, typename ForwardIt, typename Size, typename... Args>
@@ -196,13 +197,13 @@ void uninitialized_construct_n_with_allocator(Allocator const& alloc, ForwardIt
       try {
         for (; n > 0; (void) ++current, --n)
         {
-          traits::construct(alloc_T, addressof(*current), args...);
+          traits::construct(alloc_T, ::cuda::std::addressof(*current), args...);
         }
       } catch (...) {
         destroy(alloc_T, first, current);
         throw;
       }),
-    (for (; n > 0; (void) ++current, --n) { traits::construct(alloc_T, addressof(*current), args...); }));
+    (for (; n > 0; (void) ++current, --n) { traits::construct(alloc_T, ::cuda::std::addressof(*current), args...); }));
 }
 
 ///////////////////////////////////////////////////////////////////////////////

diff --git a/thrust/thrust/optional.h b/thrust/thrust/optional.h
@@ -22,10 +22,10 @@
 #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
 #  pragma system_header
 #endif // no system header
-#include <thrust/addressof.h>
 #include <thrust/detail/type_traits.h>
 #include <thrust/swap.h>
 
+#include <cuda/std/__memory/addressof.h>
 #include <cuda/std/__type_traits/void_t.h>
 
 #define THRUST_OPTIONAL_VERSION_MAJOR 0
@@ -398,7 +398,7 @@ struct optional_operations_base : optional_storage_base<T>
   template <class... Args>
   _CCCL_HOST_DEVICE void construct(Args&&... args) noexcept
   {
-    new (thrust::addressof(this->m_value)) T(std::forward<Args>(args)...);
+    new (::cuda::std::addressof(this->m_value)) T(std::forward<Args>(args)...);
     this->m_has_value = true;
   }
 
@@ -1555,13 +1555,13 @@ class optional
       }
       else
       {
-        new (thrust::addressof(rhs.m_value)) T(std::move(this->m_value));
+        new (::cuda::std::addressof(rhs.m_value)) T(std::move(this->m_value));
         this->m_value.T::~T();
       }
     }
     else if (rhs.has_value())
     {
-      new (thrust::addressof(this->m_value)) T(std::move(rhs.m_value));
+      new (::cuda::std::addressof(this->m_value)) T(std::move(rhs.m_value));
       rhs.m_value.T::~T();
     }
   }
@@ -1573,15 +1573,15 @@ class optional
   _CCCL_EXEC_CHECK_DISABLE
   _CCCL_HOST_DEVICE constexpr const T* operator->() const
   {
-    return thrust::addressof(this->m_value);
+    return ::cuda::std::addressof(this->m_value);
   }
 
   /// \group pointer
   /// \synopsis constexpr T *operator->();
   _CCCL_EXEC_CHECK_DISABLE
   _CCCL_HOST_DEVICE THRUST_OPTIONAL_CPP11_CONSTEXPR T* operator->()
   {
-    return thrust::addressof(this->m_value);
+    return ::cuda::std::addressof(this->m_value);
   }
 
   /// \return the stored value
@@ -2614,7 +2614,7 @@ class optional<T&>
   _CCCL_EXEC_CHECK_DISABLE
   template <class U = T, detail::enable_if_t<!detail::is_optional<detail::decay_t<U>>::value>* = nullptr>
   _CCCL_HOST_DEVICE constexpr optional(U&& u)
-      : m_value(thrust::addressof(u))
+      : m_value(::cuda::std::addressof(u))
   {
     static_assert(std::is_lvalue_reference<U>::value, "U must be an lvalue");
   }
@@ -2656,7 +2656,7 @@ class optional<T&>
   _CCCL_HOST_DEVICE optional& operator=(U&& u)
   {
     static_assert(std::is_lvalue_reference<U>::value, "U must be an lvalue");
-    m_value = thrust::addressof(u);
+    m_value = ::cuda::std::addressof(u);
     return *this;
   }
 
@@ -2668,7 +2668,7 @@ class optional<T&>
   template <class U>
   _CCCL_HOST_DEVICE optional& operator=(const optional<U>& rhs)
   {
-    m_value = thrust::addressof(rhs.value());
+    m_value = ::cuda::std::addressof(rhs.value());
     return *this;
   }
 
@@ -2680,7 +2680,7 @@ class optional<T&>
   template <class U>
   _CCCL_HOST_DEVICE T& emplace(U& u) noexcept
   {
-    m_value = thrust::addressof(u);
+    m_value = ::cuda::std::addressof(u);
     return *m_value;
   }
 

diff --git a/thrust/thrust/system/cuda/detail/future.inl b/thrust/thrust/system/cuda/detail/future.inl
@@ -36,6 +36,7 @@
 #  include <thrust/type_traits/integer_sequence.h>
 #  include <thrust/type_traits/remove_cvref.h>
 
+#  include <cuda/std/__memory/addressof.h>
 #  include <cuda/std/__memory/unique_ptr.h>
 
 #  include <type_traits>
@@ -639,7 +640,7 @@ public:
   // For testing only.
   _CCCL_HOST_DEVICE raw_const_pointer data() const
   {
-    return addressof(value_);
+    return ::cuda::std::addressof(value_);
   }
 #  endif
 };