Cleanup some of thrust pointer traits

libcudacxx/include/cuda/std/__memory/pointer_traits.h

@@ -159,12 +159,6 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT pointer_traits<_Tp*>
   }
 };
 
-template <class _From, class _To>
-struct __rebind_pointer
-{
-  using type = typename pointer_traits<_From>::template rebind<_To>;
-};
-
 // to_address
 
 template <class _Pointer, class = void>

thrust/testing/memory.cu

@@ -71,8 +71,8 @@ get_temporary_buffer(my_old_temporary_allocation_system, std::ptrdiff_t)
 template <typename Pointer>
 void return_temporary_buffer(my_old_temporary_allocation_system, Pointer p)
 {
-  using RP = typename thrust::detail::pointer_traits<Pointer>::raw_pointer;
-  ASSERT_EQUAL(p.get(), reinterpret_cast<RP>(4217));
+  using RP = typename cuda::std::pointer_traits<Pointer>::raw_pointer;
+  ASSERT_EQUAL(::cuda::std::to_address(p), reinterpret_cast<RP>(4217));
 }
 } // namespace my_old_namespace
 
@@ -101,8 +101,8 @@ void return_temporary_buffer(my_new_temporary_allocation_system, Pointer)
 template <typename Pointer>
 void return_temporary_buffer(my_new_temporary_allocation_system, Pointer p, std::ptrdiff_t n)
 {
-  using RP = typename thrust::detail::pointer_traits<Pointer>::raw_pointer;
-  ASSERT_EQUAL(p.get(), reinterpret_cast<RP>(1742));
+  using RP = typename cuda::std::pointer_traits<Pointer>::raw_pointer;
+  ASSERT_EQUAL(::cuda::std::to_address(p), reinterpret_cast<RP>(1742));
   ASSERT_EQUAL(n, 413);
 }
 } // namespace my_new_namespace

thrust/testing/mr_pool.cu

@@ -4,6 +4,8 @@
 #include <thrust/mr/pool.h>
 #include <thrust/mr/sync_pool.h>
 
+#include <cuda/std/memory>
+
 #include <unittest/unittest.h>
 
 template <typename T>
@@ -100,6 +102,25 @@ struct tracked_pointer
   }
 };
 
+template <typename T>
+struct cuda::std::pointer_traits<tracked_pointer<T>>
+{
+  using pointer         = tracked_pointer<T>;
+  using element_type    = T;
+  using difference_type = cuda::std::ptrdiff_t;
+
+  template <typename U>
+  struct rebind
+  {
+    using other = typename THRUST_NS_QUALIFIER::detail::rebind_pointer<pointer, U>::type;
+  };
+
+  [[nodiscard]] _CCCL_API static constexpr element_type* to_address(const pointer iter) noexcept
+  {
+    return iter.get();
+  }
+};
+
 class tracked_resource final : public thrust::mr::memory_resource<tracked_pointer<void>>
 {
 public:

thrust/thrust/allocate_unique.h

@@ -69,7 +69,7 @@ struct allocator_delete final
     using traits = ::cuda::std::allocator_traits<::cuda::std::remove_cvref_t<Allocator>>;
     typename traits::allocator_type alloc_T(alloc_);
 
-    if (nullptr != detail::pointer_traits<pointer>::get(p))
+    if (nullptr != ::cuda::std::to_address(p))
     {
       if constexpr (!Uninitialized)
       {
@@ -144,7 +144,7 @@ struct array_allocator_delete final
   {
     using traits = ::cuda::std::allocator_traits<::cuda::std::remove_cvref_t<Allocator>>;
     typename traits::allocator_type alloc_T(get_allocator());
-    if (nullptr != detail::pointer_traits<pointer>::get(p))
+    if (nullptr != ::cuda::std::to_address(p))
     {
       if constexpr (!Uninitialized)
       {

thrust/thrust/detail/allocator/allocator_system.h

@@ -31,11 +31,11 @@
 #endif // no system header
 
 #include <thrust/detail/type_traits.h>
-#include <thrust/detail/type_traits/pointer_traits.h>
 #include <thrust/iterator/iterator_traits.h>
 
 #include <cuda/std/__iterator/iterator_traits.h>
 #include <cuda/std/__memory/allocator_traits.h>
+#include <cuda/std/__memory/pointer_traits.h>
 #include <cuda/std/__type_traits/add_lvalue_reference.h>
 #include <cuda/std/__type_traits/conditional.h>
 #include <cuda/std/__type_traits/integral_constant.h>

thrust/thrust/detail/allocator/copy_construct_range.h

@@ -29,7 +29,6 @@
 #include <thrust/detail/allocator/allocator_system.h>
 #include <thrust/detail/copy.h>
 #include <thrust/detail/execution_policy.h>
-#include <thrust/detail/type_traits/pointer_traits.h>
 #include <thrust/for_each.h>
 #include <thrust/iterator/iterator_traits.h>
 #include <thrust/iterator/zip_iterator.h>
@@ -38,6 +37,7 @@
 #include <cuda/std/__iterator/advance.h>
 #include <cuda/std/__iterator/distance.h>
 #include <cuda/std/__memory/allocator_traits.h>
+#include <cuda/std/__memory/pointer_traits.h>
 #include <cuda/std/__type_traits/is_convertible.h>
 #include <cuda/std/__type_traits/is_trivially_copy_constructible.h>
 #include <cuda/std/tuple>
@@ -164,7 +164,8 @@ template <typename System, typename Allocator, typename InputIterator, typename
 _CCCL_HOST_DEVICE Pointer copy_construct_range(
   thrust::execution_policy<System>& from_system, Allocator& a, InputIterator first, InputIterator last, Pointer result)
 {
-  if constexpr (needs_copy_construct_via_allocator<Allocator, typename pointer_element<Pointer>::type>)
+  if constexpr (needs_copy_construct_via_allocator<Allocator,
+                                                   typename ::cuda::std::pointer_traits<Pointer>::element_type>)
   {
     return uninitialized_copy_with_allocator(a, from_system, allocator_system<Allocator>::get(a), first, last, result);
   }
@@ -179,7 +180,8 @@ template <typename System, typename Allocator, typename InputIterator, typename
 _CCCL_HOST_DEVICE Pointer copy_construct_range_n(
   thrust::execution_policy<System>& from_system, Allocator& a, InputIterator first, Size n, Pointer result)
 {
-  if constexpr (needs_copy_construct_via_allocator<Allocator, typename pointer_element<Pointer>::type>)
+  if constexpr (needs_copy_construct_via_allocator<Allocator,
+                                                   typename ::cuda::std::pointer_traits<Pointer>::element_type>)
   {
     return uninitialized_copy_with_allocator_n(a, from_system, allocator_system<Allocator>::get(a), first, n, result);
   }

thrust/thrust/detail/allocator/destroy_range.h

@@ -28,11 +28,11 @@
 
 #include <thrust/detail/allocator/allocator_system.h>
 #include <thrust/detail/allocator/destroy_range.h>
-#include <thrust/detail/type_traits/pointer_traits.h>
 #include <thrust/for_each.h>
 
 #include <cuda/std/__host_stdlib/memory>
 #include <cuda/std/__memory/allocator_traits.h>
+#include <cuda/std/__memory/pointer_traits.h>
 
 THRUST_NAMESPACE_BEGIN
 namespace detail
@@ -78,7 +78,7 @@ template <typename Allocator, typename Pointer, typename Size>
 _CCCL_HOST_DEVICE void
 destroy_range([[maybe_unused]] Allocator& a, [[maybe_unused]] Pointer p, [[maybe_unused]] Size n) noexcept
 {
-  using pe_t = typename pointer_element<Pointer>::type;
+  using pe_t = typename ::cuda::std::pointer_traits<Pointer>::element_type;
 
   // case 1: destroy via allocator
   if constexpr (has_effectful_member_destroy<Allocator, pe_t>)

thrust/thrust/detail/allocator/fill_construct_range.h

@@ -28,12 +28,12 @@
 
 #include <thrust/detail/allocator/allocator_system.h>
 #include <thrust/detail/type_traits.h>
-#include <thrust/detail/type_traits/pointer_traits.h>
 #include <thrust/for_each.h>
 #include <thrust/uninitialized_fill.h>
 
 #include <cuda/std/__host_stdlib/memory>
 #include <cuda/std/__memory/allocator_traits.h>
+#include <cuda/std/__memory/pointer_traits.h>
 
 THRUST_NAMESPACE_BEGIN
 namespace detail
@@ -66,7 +66,7 @@ struct construct2_via_allocator
 template <typename Allocator, typename Pointer, typename Size, typename T>
 _CCCL_HOST_DEVICE void fill_construct_range(Allocator& a, Pointer p, Size n, const T& value)
 {
-  if constexpr (has_effectful_member_construct2<Allocator, typename pointer_element<Pointer>::type, T>)
+  if constexpr (has_effectful_member_construct2<Allocator, typename ::cuda::std::pointer_traits<Pointer>::element_type, T>)
   {
     thrust::for_each_n(allocator_system<Allocator>::get(a), p, n, construct2_via_allocator<Allocator, T>{a, value});
   }

thrust/thrust/detail/allocator/malloc_allocator.h

@@ -30,10 +30,11 @@
 #include <thrust/detail/allocator/tagged_allocator.h>
 #include <thrust/detail/malloc_and_free.h>
 #include <thrust/detail/raw_pointer_cast.h>
-#include <thrust/detail/type_traits/pointer_traits.h>
 #include <thrust/system/detail/bad_alloc.h>
 #include <thrust/system/detail/generic/select_system.h>
 
+#include <cuda/std/__memory/pointer_traits.h>
+
 THRUST_NAMESPACE_BEGIN
 namespace detail
 {

thrust/thrust/detail/allocator/tagged_allocator.h

@@ -25,10 +25,11 @@
 #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
 #  pragma system_header
 #endif // no system header
+
 #include <thrust/detail/allocator/tagged_allocator.h>
-#include <thrust/detail/type_traits/pointer_traits.h>
 #include <thrust/iterator/iterator_traits.h>
 
+#include <cuda/std/__memory/pointer_traits.h>
 #include <cuda/std/limits>
 
 THRUST_NAMESPACE_BEGIN
@@ -42,10 +43,10 @@ class tagged_allocator<void, Tag, Pointer>
 {
 public:
   using value_type      = void;
-  using pointer         = typename thrust::detail::pointer_traits<Pointer>::template rebind<void>::other;
-  using const_pointer   = typename thrust::detail::pointer_traits<Pointer>::template rebind<const void>::other;
+  using pointer         = typename ::cuda::std::pointer_traits<Pointer>::template rebind<void>::other;
+  using const_pointer   = typename ::cuda::std::pointer_traits<Pointer>::template rebind<const void>::other;
   using size_type       = std::size_t;
-  using difference_type = typename thrust::detail::pointer_traits<Pointer>::difference_type;
+  using difference_type = typename ::cuda::std::pointer_traits<Pointer>::difference_type;
   using system_type     = Tag;
 
   template <typename U>
@@ -60,12 +61,12 @@ class tagged_allocator
 {
 public:
   using value_type      = T;
-  using pointer         = typename thrust::detail::pointer_traits<Pointer>::template rebind<T>::other;
-  using const_pointer   = typename thrust::detail::pointer_traits<Pointer>::template rebind<const T>::other;
+  using pointer         = typename ::cuda::std::pointer_traits<Pointer>::template rebind<T>::other;
+  using const_pointer   = typename ::cuda::std::pointer_traits<Pointer>::template rebind<const T>::other;
   using reference       = thrust::detail::it_reference_t<pointer>;
   using const_reference = thrust::detail::it_reference_t<const_pointer>;
   using size_type       = std::size_t;
-  using difference_type = typename thrust::detail::pointer_traits<pointer>::difference_type;
+  using difference_type = typename ::cuda::std::pointer_traits<pointer>::difference_type;
   using system_type     = Tag;
 
   template <typename U>

thrust/thrust/detail/allocator/value_initialize_range.h

@@ -28,11 +28,11 @@
 
 #include <thrust/detail/allocator/allocator_system.h>
 #include <thrust/detail/type_traits.h>
-#include <thrust/detail/type_traits/pointer_traits.h>
 #include <thrust/for_each.h>
 #include <thrust/uninitialized_fill.h>
 
 #include <cuda/std/__memory/allocator_traits.h>
+#include <cuda/std/__memory/pointer_traits.h>
 
 THRUST_NAMESPACE_BEGIN
 namespace detail
@@ -65,13 +65,15 @@ inline constexpr bool needs_default_construct_via_allocator<std::allocator<U>, T
 template <typename Allocator, typename Pointer, typename Size>
 _CCCL_HOST_DEVICE void value_initialize_range(Allocator& a, Pointer p, Size n)
 {
-  if constexpr (needs_default_construct_via_allocator<Allocator, typename pointer_element<Pointer>::type>)
+  if constexpr (needs_default_construct_via_allocator<Allocator,
+                                                      typename ::cuda::std::pointer_traits<Pointer>::element_type>)
   {
     thrust::for_each_n(allocator_system<Allocator>::get(a), p, n, construct1_via_allocator<Allocator>{a});
   }
   else
   {
-    thrust::uninitialized_fill_n(allocator_system<Allocator>::get(a), p, n, typename pointer_element<Pointer>::type());
+    thrust::uninitialized_fill_n(
+      allocator_system<Allocator>::get(a), p, n, typename ::cuda::std::pointer_traits<Pointer>::element_type());
   }
 }
 } // namespace detail

thrust/thrust/detail/execute_with_allocator.h

@@ -28,10 +28,10 @@
 
 #include <thrust/detail/execute_with_allocator_fwd.h>
 #include <thrust/detail/raw_pointer_cast.h>
-#include <thrust/detail/type_traits/pointer_traits.h>
 
 #include <cuda/__cmath/ceil_div.h>
 #include <cuda/std/__memory/allocator_traits.h>
+#include <cuda/std/__memory/pointer_traits.h>
 #include <cuda/std/__utility/pair.h>
 
 THRUST_NAMESPACE_BEGIN
@@ -67,7 +67,7 @@ _CCCL_HOST void return_temporary_buffer(
   using pointer         = typename alloc_traits::pointer;
   using size_type       = typename alloc_traits::size_type;
   using value_type      = typename alloc_traits::value_type;
-  using T               = typename thrust::detail::pointer_traits<Pointer>::element_type;
+  using T               = typename ::cuda::std::pointer_traits<Pointer>::element_type;
 
   size_type num_elements = ::cuda::ceil_div(sizeof(T) * n, sizeof(value_type));
 

thrust/thrust/detail/get_iterator_value.h

@@ -25,11 +25,12 @@
 #  pragma system_header
 #endif // no system header
 
-#include <thrust/detail/type_traits/pointer_traits.h>
 #include <thrust/execution_policy.h>
 #include <thrust/iterator/iterator_traits.h>
 #include <thrust/system/detail/generic/memory.h> // for get_value()
 
+#include <cuda/std/__memory/pointer_traits.h>
+
 THRUST_NAMESPACE_BEGIN
 
 namespace detail
@@ -50,7 +51,7 @@ _CCCL_HOST_DEVICE it_value_t<Iterator> get_iterator_value(thrust::execution_poli
 // we use get_value(exec,pointer*) function
 // to perform a dereferencing consistent with the execution policy
 template <typename DerivedPolicy, typename Pointer>
-_CCCL_HOST_DEVICE typename thrust::detail::pointer_traits<Pointer*>::element_type
+_CCCL_HOST_DEVICE typename ::cuda::std::pointer_traits<Pointer*>::element_type
 get_iterator_value(thrust::execution_policy<DerivedPolicy>& exec, Pointer* ptr)
 {
   return get_value(derived_cast(exec), ptr);