Drop more uses of thrust::pointer_traits

miscco · miscco · commit 138f219a8430 · 2026-03-06T07:51:00.000+01:00
diff --git a/thrust/thrust/detail/allocator/allocator_system.h b/thrust/thrust/detail/allocator/allocator_system.h
@@ -18,11 +18,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>
diff --git a/thrust/thrust/detail/allocator/copy_construct_range.h b/thrust/thrust/detail/allocator/copy_construct_range.h
@@ -16,7 +16,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>
@@ -25,6 +24,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>
@@ -151,7 +151,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);
   }
@@ -166,7 +167,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);
   }
diff --git a/thrust/thrust/detail/allocator/destroy_range.h b/thrust/thrust/detail/allocator/destroy_range.h
@@ -15,11 +15,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
@@ -65,7 +65,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>)
diff --git a/thrust/thrust/detail/allocator/fill_construct_range.h b/thrust/thrust/detail/allocator/fill_construct_range.h
@@ -15,12 +15,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
@@ -53,7 +53,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});
   }
diff --git a/thrust/thrust/detail/allocator/malloc_allocator.h b/thrust/thrust/detail/allocator/malloc_allocator.h
@@ -17,10 +17,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
 {
diff --git a/thrust/thrust/detail/allocator/tagged_allocator.h b/thrust/thrust/detail/allocator/tagged_allocator.h
@@ -14,7 +14,6 @@
 #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>
diff --git a/thrust/thrust/detail/allocator/value_initialize_range.h b/thrust/thrust/detail/allocator/value_initialize_range.h
@@ -15,11 +15,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
@@ -52,13 +52,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
diff --git a/thrust/thrust/detail/execute_with_allocator.h b/thrust/thrust/detail/execute_with_allocator.h
@@ -15,7 +15,6 @@
 
 #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>
diff --git a/thrust/thrust/detail/get_iterator_value.h b/thrust/thrust/detail/get_iterator_value.h
@@ -13,7 +13,6 @@
 #  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()
diff --git a/thrust/thrust/detail/raw_pointer_cast.h b/thrust/thrust/detail/raw_pointer_cast.h
@@ -13,8 +13,6 @@
 #  pragma system_header
 #endif // no system header
 
-#include <thrust/detail/type_traits/pointer_traits.h>
-
 #include <cuda/std/__memory/pointer_traits.h>
 
 THRUST_NAMESPACE_BEGIN
@@ -28,15 +26,15 @@ _CCCL_HOST_DEVICE auto raw_pointer_cast(Pointer ptr)
 template <typename ToPointer, typename FromPointer>
 _CCCL_HOST_DEVICE ToPointer reinterpret_pointer_cast(FromPointer ptr)
 {
-  using to_element = typename thrust::detail::pointer_element<ToPointer>::type;
-  return ToPointer(reinterpret_cast<to_element*>(thrust::raw_pointer_cast(ptr)));
+  using to_element = typename ::cuda::std::pointer_traits<ToPointer>::element_type;
+  return ToPointer(reinterpret_cast<to_element*>(::cuda::std::to_address(ptr)));
 }
 
 template <typename ToPointer, typename FromPointer>
 _CCCL_HOST_DEVICE ToPointer static_pointer_cast(FromPointer ptr)
 {
-  using to_element = typename thrust::detail::pointer_element<ToPointer>::type;
-  return ToPointer(static_cast<to_element*>(thrust::raw_pointer_cast(ptr)));
+  using to_element = typename ::cuda::std::pointer_traits<ToPointer>::element_type;
+  return ToPointer(static_cast<to_element*>(::cuda::std::to_address(ptr)));
 }
 
 THRUST_NAMESPACE_END
diff --git a/thrust/thrust/detail/raw_reference_cast.h b/thrust/thrust/detail/raw_reference_cast.h
@@ -17,6 +17,7 @@
 #include <thrust/detail/type_traits.h>
 #include <thrust/detail/type_traits/has_nested_type.h>
 
+#include <cuda/std/__memory/pointer_traits.h>
 #include <cuda/std/__tuple_dir/apply.h>
 #include <cuda/std/__type_traits/add_lvalue_reference.h>
 #include <cuda/std/__type_traits/enable_if.h>
@@ -62,7 +63,7 @@ struct raw_reference_impl : ::cuda::std::add_lvalue_reference<T>
 
 template <typename T>
 struct raw_reference_impl<T, ::cuda::std::enable_if_t<is_wrapped_reference<::cuda::std::remove_cv_t<T>>::value>>
-    : ::cuda::std::add_lvalue_reference<typename pointer_element<typename T::pointer>::type>
+    : ::cuda::std::add_lvalue_reference<typename ::cuda::std::pointer_traits<typename T::pointer>::element_type>
 {};
 
 template <typename T>
diff --git a/thrust/thrust/detail/reference.h b/thrust/thrust/detail/reference.h
@@ -18,11 +18,11 @@
 #  pragma system_header
 #endif // no system header
 #include <thrust/detail/reference_forward_declaration.h>
-#include <thrust/detail/type_traits/pointer_traits.h>
 #include <thrust/iterator/iterator_traits.h>
 #include <thrust/system/detail/generic/memory.h>
 #include <thrust/system/detail/generic/select_system.h>
 
+#include <cuda/std/__memory/pointer_traits.h>
 #include <cuda/std/__type_traits/enable_if.h>
 #include <cuda/std/__type_traits/is_comparable.h>
 #include <cuda/std/__type_traits/is_convertible.h>
diff --git a/thrust/thrust/detail/type_traits/pointer_traits.h b/thrust/thrust/detail/type_traits/pointer_traits.h
@@ -30,33 +30,6 @@
 THRUST_NAMESPACE_BEGIN
 namespace detail
 {
-template <typename Ptr>
-struct pointer_element;
-
-template <template <typename...> class Ptr, typename FirstArg, typename... Args>
-struct pointer_element<Ptr<FirstArg, Args...>>
-{
-  using type = FirstArg;
-};
-
-template <typename T>
-struct pointer_element<T*>
-{
-  using type = T;
-};
-
-template <typename Ptr>
-struct pointer_difference
-{
-  using type = typename Ptr::difference_type;
-};
-
-template <typename T>
-struct pointer_difference<T*>
-{
-  using type = ::cuda::std::ptrdiff_t;
-};
-
 template <typename Ptr, typename T>
 struct rebind_pointer;
 
@@ -161,8 +134,8 @@ struct pointer_traits
 {
   using pointer         = Ptr;
   using reference       = typename Ptr::reference;
-  using element_type    = typename pointer_element<Ptr>::type;
-  using difference_type = typename pointer_difference<Ptr>::type;
+  using element_type    = typename ::cuda::std::pointer_traits<pointer>::element_type;
+  using difference_type = typename ::cuda::std::pointer_traits<pointer>::difference_type;
 
   template <typename U>
   struct rebind
@@ -195,7 +168,7 @@ struct pointer_traits<T*>
   using pointer         = T*;
   using reference       = T&;
   using element_type    = T;
-  using difference_type = typename pointer_difference<T*>::type;
+  using difference_type = typename ::cuda::std::pointer_traits<pointer>::difference_type;
 
   template <typename U>
   struct rebind
@@ -224,7 +197,7 @@ struct pointer_traits<void*>
   using pointer         = void*;
   using reference       = void;
   using element_type    = void;
-  using difference_type = pointer_difference<void*>::type;
+  using difference_type = typename ::cuda::std::pointer_traits<pointer>::difference_type;
 
   template <typename U>
   struct rebind
@@ -252,7 +225,7 @@ struct pointer_traits<const void*>
   using pointer         = const void*;
   using reference       = const void;
   using element_type    = const void;
-  using difference_type = pointer_difference<const void*>::type;
+  using difference_type = typename ::cuda::std::pointer_traits<pointer>::difference_type;
 
   template <typename U>
   struct rebind
@@ -280,12 +253,13 @@ inline constexpr bool is_pointer_system_convertible_v =
 
 template <typename FromPtr, typename ToPtr>
 inline constexpr bool is_pointer_convertible_v =
-  ::cuda::std::is_convertible_v<typename pointer_element<FromPtr>::type*, typename pointer_element<ToPtr>::type*>
+  ::cuda::std::is_convertible_v<typename ::cuda::std::pointer_traits<FromPtr>::element_type*,
+                                typename ::cuda::std::pointer_traits<ToPtr>::element_type*>
   && is_pointer_system_convertible_v<FromPtr, ToPtr>;
 
 template <typename FromPtr, typename ToPtr>
 inline constexpr bool is_void_pointer_system_convertible_v =
-  ::cuda::std::is_void_v<typename pointer_element<FromPtr>::type> //
+  ::cuda::std::is_void_v<typename ::cuda::std::pointer_traits<FromPtr>::element_type> //
   && is_pointer_system_convertible_v<FromPtr, ToPtr>;
 
 // avoid inspecting traits of the arguments if they aren't known to be pointers
diff --git a/thrust/thrust/memory.h b/thrust/thrust/memory.h
@@ -21,7 +21,6 @@
 #include <thrust/detail/raw_reference_cast.h>
 #include <thrust/detail/reference.h>
 #include <thrust/detail/temporary_buffer.h>
-#include <thrust/detail/type_traits/pointer_traits.h>
 
 #include <cuda/std/__memory/pointer_traits.h>
 
diff --git a/thrust/thrust/mr/disjoint_pool.h b/thrust/thrust/mr/disjoint_pool.h
@@ -22,6 +22,7 @@
 
 #include <thrust/binary_search.h>
 #include <thrust/detail/seq.h>
+#include <thrust/detail/type_traits/pointer_traits.h>
 #include <thrust/find.h>
 #include <thrust/host_vector.h>
 #include <thrust/mr/allocator.h>
@@ -33,7 +34,6 @@
 #include <cuda/std/__algorithm/max.h>
 #include <cuda/std/__algorithm/min.h>
 #include <cuda/std/__host_stdlib/algorithm>
-#include <cuda/std/__memory/pointer_traits.h>
 #include <cuda/std/cassert>
 #include <cuda/std/cstdint>
 
diff --git a/thrust/thrust/mr/fancy_pointer_resource.h b/thrust/thrust/mr/fancy_pointer_resource.h
@@ -12,10 +12,12 @@
 #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
 #  pragma system_header
 #endif // no system header
-#include <thrust/detail/type_traits/pointer_traits.h>
+
 #include <thrust/mr/memory_resource.h>
 #include <thrust/mr/validator.h>
 
+#include <cuda/std/__memory/pointer_traits.h>
+
 THRUST_NAMESPACE_BEGIN
 namespace mr
 {
diff --git a/thrust/thrust/mr/pool.h b/thrust/thrust/mr/pool.h
@@ -19,6 +19,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <thrust/detail/type_traits/pointer_traits.h>
 #include <thrust/host_vector.h>
 #include <thrust/mr/allocator.h>
 #include <thrust/mr/memory_resource.h>
diff --git a/thrust/thrust/system/cuda/detail/temporary_buffer.h b/thrust/thrust/system/cuda/detail/temporary_buffer.h
@@ -13,7 +13,6 @@
 #  pragma system_header
 #endif // no system header
 
-#include <thrust/detail/type_traits/pointer_traits.h>
 #include <thrust/system/cuda/detail/execution_policy.h>
 #include <thrust/system/detail/bad_alloc.h>
 
diff --git a/thrust/thrust/system/detail/generic/memory.inl b/thrust/thrust/system/detail/generic/memory.inl
@@ -15,7 +15,6 @@
 
 #include <thrust/detail/malloc_and_free_fwd.h>
 #include <thrust/detail/static_assert.h>
-#include <thrust/detail/type_traits/pointer_traits.h>
 
 #include <cuda/std/__type_traits/is_void.h>
 
diff --git a/thrust/thrust/system/detail/sequential/copy.h b/thrust/thrust/system/detail/sequential/copy.h
@@ -17,13 +17,14 @@
 #  pragma system_header
 #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 <thrust/system/detail/sequential/execution_policy.h>
 #include <thrust/system/detail/sequential/general_copy.h>
 #include <thrust/system/detail/sequential/trivial_copy.h>
 #include <thrust/type_traits/is_trivially_relocatable.h>
 
+#include <cuda/std/__memory/pointer_traits.h>
+
 THRUST_NAMESPACE_BEGIN
 namespace system::detail::sequential
 {

Original file line number	Diff line number	Diff line change
`@@ -16,7 +16,6 @@`
`16`	`16`	`#include <thrust/detail/allocator/allocator_system.h>`
`17`	`17`	`#include <thrust/detail/copy.h>`
`18`	`18`	`#include <thrust/detail/execution_policy.h>`
`19`		`-#include <thrust/detail/type_traits/pointer_traits.h>`
`20`	`19`	`#include <thrust/for_each.h>`
`21`	`20`	`#include <thrust/iterator/iterator_traits.h>`
`22`	`21`	`#include <thrust/iterator/zip_iterator.h>`
`@@ -25,6 +24,7 @@`
`25`	`24`	`#include <cuda/std/__iterator/advance.h>`
`26`	`25`	`#include <cuda/std/__iterator/distance.h>`
`27`	`26`	`#include <cuda/std/__memory/allocator_traits.h>`
	`27`	`+#include <cuda/std/__memory/pointer_traits.h>`
`28`	`28`	`#include <cuda/std/__type_traits/is_convertible.h>`
`29`	`29`	`#include <cuda/std/__type_traits/is_trivially_copy_constructible.h>`
`30`	`30`	`#include <cuda/std/tuple>`
`@@ -151,7 +151,8 @@ template <typename System, typename Allocator, typename InputIterator, typename`
`151`	`151`	`_CCCL_HOST_DEVICE Pointer copy_construct_range(`
`152`	`152`	`thrust::execution_policy<System>& from_system, Allocator& a, InputIterator first, InputIterator last, Pointer result)`
`153`	`153`	`{`
`154`		`- if constexpr (needs_copy_construct_via_allocator<Allocator, typename pointer_element<Pointer>::type>)`
	`154`	`+ if constexpr (needs_copy_construct_via_allocator<Allocator,`
	`155`	`+ typename ::cuda::std::pointer_traits<Pointer>::element_type>)`
`155`	`156`	`{`
`156`	`157`	`return uninitialized_copy_with_allocator(a, from_system, allocator_system<Allocator>::get(a), first, last, result);`
`157`	`158`	`}`
`@@ -166,7 +167,8 @@ template <typename System, typename Allocator, typename InputIterator, typename`
`166`	`167`	`_CCCL_HOST_DEVICE Pointer copy_construct_range_n(`
`167`	`168`	`thrust::execution_policy<System>& from_system, Allocator& a, InputIterator first, Size n, Pointer result)`
`168`	`169`	`{`
`169`		`- if constexpr (needs_copy_construct_via_allocator<Allocator, typename pointer_element<Pointer>::type>)`
	`170`	`+ if constexpr (needs_copy_construct_via_allocator<Allocator,`
	`171`	`+ typename ::cuda::std::pointer_traits<Pointer>::element_type>)`
`170`	`172`	`{`
`171`	`173`	`return uninitialized_copy_with_allocator_n(a, from_system, allocator_system<Allocator>::get(a), first, n, result);`
`172`	`174`	`}`
Original file line number	Diff line number	Diff line change
`@@ -15,12 +15,12 @@`
`15`	`15`
`16`	`16`	`#include <thrust/detail/allocator/allocator_system.h>`
`17`	`17`	`#include <thrust/detail/type_traits.h>`
`18`		`-#include <thrust/detail/type_traits/pointer_traits.h>`
`19`	`18`	`#include <thrust/for_each.h>`
`20`	`19`	`#include <thrust/uninitialized_fill.h>`
`21`	`20`
`22`	`21`	`#include <cuda/std/__host_stdlib/memory>`
`23`	`22`	`#include <cuda/std/__memory/allocator_traits.h>`
	`23`	`+#include <cuda/std/__memory/pointer_traits.h>`
`24`	`24`
`25`	`25`	`THRUST_NAMESPACE_BEGIN`
`26`	`26`	`namespace detail`
`@@ -53,7 +53,7 @@ struct construct2_via_allocator`
`53`	`53`	`template <typename Allocator, typename Pointer, typename Size, typename T>`
`54`	`54`	`_CCCL_HOST_DEVICE void fill_construct_range(Allocator& a, Pointer p, Size n, const T& value)`
`55`	`55`	`{`
`56`		`- if constexpr (has_effectful_member_construct2<Allocator, typename pointer_element<Pointer>::type, T>)`
	`56`	`+ if constexpr (has_effectful_member_construct2<Allocator, typename ::cuda::std::pointer_traits<Pointer>::element_type, T>)`
`57`	`57`	`{`
`58`	`58`	`thrust::for_each_n(allocator_system<Allocator>::get(a), p, n, construct2_via_allocator<Allocator, T>{a, value});`
`59`	`59`	`}`
Original file line number	Diff line number	Diff line change
`@@ -15,11 +15,11 @@`
`15`	`15`
`16`	`16`	`#include <thrust/detail/allocator/allocator_system.h>`
`17`	`17`	`#include <thrust/detail/type_traits.h>`
`18`		`-#include <thrust/detail/type_traits/pointer_traits.h>`
`19`	`18`	`#include <thrust/for_each.h>`
`20`	`19`	`#include <thrust/uninitialized_fill.h>`
`21`	`20`
`22`	`21`	`#include <cuda/std/__memory/allocator_traits.h>`
	`22`	`+#include <cuda/std/__memory/pointer_traits.h>`
`23`	`23`
`24`	`24`	`THRUST_NAMESPACE_BEGIN`
`25`	`25`	`namespace detail`
`@@ -52,13 +52,15 @@ inline constexpr bool needs_default_construct_via_allocator<std::allocator<U>, T`
`52`	`52`	`template <typename Allocator, typename Pointer, typename Size>`
`53`	`53`	`_CCCL_HOST_DEVICE void value_initialize_range(Allocator& a, Pointer p, Size n)`
`54`	`54`	`{`
`55`		`- if constexpr (needs_default_construct_via_allocator<Allocator, typename pointer_element<Pointer>::type>)`
	`55`	`+ if constexpr (needs_default_construct_via_allocator<Allocator,`
	`56`	`+ typename ::cuda::std::pointer_traits<Pointer>::element_type>)`
`56`	`57`	`{`
`57`	`58`	`thrust::for_each_n(allocator_system<Allocator>::get(a), p, n, construct1_via_allocator<Allocator>{a});`
`58`	`59`	`}`
`59`	`60`	`else`
`60`	`61`	`{`
`61`		`- thrust::uninitialized_fill_n(allocator_system<Allocator>::get(a), p, n, typename pointer_element<Pointer>::type());`
	`62`	`+ thrust::uninitialized_fill_n(`
	`63`	`+ allocator_system<Allocator>::get(a), p, n, typename ::cuda::std::pointer_traits<Pointer>::element_type());`
`62`	`64`	`}`
`63`	`65`	`}`
`64`	`66`	`} // namespace detail`
Original file line number	Diff line number	Diff line change
`@@ -13,8 +13,6 @@`
`13`	`13`	`# pragma system_header`
`14`	`14`	`#endif // no system header`
`15`	`15`
`16`		`-#include <thrust/detail/type_traits/pointer_traits.h>`
`17`		`-`
`18`	`16`	`#include <cuda/std/__memory/pointer_traits.h>`
`19`	`17`
`20`	`18`	`THRUST_NAMESPACE_BEGIN`
`@@ -28,15 +26,15 @@ _CCCL_HOST_DEVICE auto raw_pointer_cast(Pointer ptr)`
`28`	`26`	`template <typename ToPointer, typename FromPointer>`
`29`	`27`	`_CCCL_HOST_DEVICE ToPointer reinterpret_pointer_cast(FromPointer ptr)`
`30`	`28`	`{`
`31`		`- using to_element = typename thrust::detail::pointer_element<ToPointer>::type;`
`32`		`- return ToPointer(reinterpret_cast<to_element*>(thrust::raw_pointer_cast(ptr)));`
	`29`	`+ using to_element = typename ::cuda::std::pointer_traits<ToPointer>::element_type;`
	`30`	`+ return ToPointer(reinterpret_cast<to_element*>(::cuda::std::to_address(ptr)));`
`33`	`31`	`}`
`34`	`32`
`35`	`33`	`template <typename ToPointer, typename FromPointer>`
`36`	`34`	`_CCCL_HOST_DEVICE ToPointer static_pointer_cast(FromPointer ptr)`
`37`	`35`	`{`
`38`		`- using to_element = typename thrust::detail::pointer_element<ToPointer>::type;`
`39`		`- return ToPointer(static_cast<to_element*>(thrust::raw_pointer_cast(ptr)));`
	`36`	`+ using to_element = typename ::cuda::std::pointer_traits<ToPointer>::element_type;`
	`37`	`+ return ToPointer(static_cast<to_element*>(::cuda::std::to_address(ptr)));`
`40`	`38`	`}`
`41`	`39`
`42`	`40`	`THRUST_NAMESPACE_END`