Drop more uses of thrust::pointer_traits

miscco · miscco · commit 3fbd348894fb · 2026-01-30T10:44:20.000+01:00
diff --git a/thrust/thrust/detail/allocator/allocator_system.h b/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>
diff --git a/thrust/thrust/detail/allocator/copy_construct_range.h b/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);
   }
diff --git a/thrust/thrust/detail/allocator/destroy_range.h b/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>)
diff --git a/thrust/thrust/detail/allocator/fill_construct_range.h b/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});
   }
diff --git a/thrust/thrust/detail/allocator/malloc_allocator.h b/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
 {
diff --git a/thrust/thrust/detail/allocator/tagged_allocator.h b/thrust/thrust/detail/allocator/tagged_allocator.h
@@ -27,7 +27,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
@@ -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
diff --git a/thrust/thrust/detail/execute_with_allocator.h b/thrust/thrust/detail/execute_with_allocator.h
@@ -28,7 +28,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
@@ -25,7 +25,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
@@ -26,8 +26,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
@@ -41,15 +39,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
@@ -30,6 +30,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>
@@ -75,7 +76,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
@@ -31,11 +31,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
@@ -43,33 +43,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;
 
@@ -174,8 +147,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
@@ -208,7 +181,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
@@ -237,7 +210,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
@@ -265,7 +238,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
@@ -293,12 +266,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
@@ -34,7 +34,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
@@ -35,6 +35,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>
@@ -46,7 +47,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
@@ -25,10 +25,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
@@ -32,6 +32,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
@@ -26,7 +26,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
@@ -28,7 +28,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
@@ -30,13 +30,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
`@@ -29,7 +29,6 @@`
`29`	`29`	`#include <thrust/detail/allocator/allocator_system.h>`
`30`	`30`	`#include <thrust/detail/copy.h>`
`31`	`31`	`#include <thrust/detail/execution_policy.h>`
`32`		`-#include <thrust/detail/type_traits/pointer_traits.h>`
`33`	`32`	`#include <thrust/for_each.h>`
`34`	`33`	`#include <thrust/iterator/iterator_traits.h>`
`35`	`34`	`#include <thrust/iterator/zip_iterator.h>`
`@@ -38,6 +37,7 @@`
`38`	`37`	`#include <cuda/std/__iterator/advance.h>`
`39`	`38`	`#include <cuda/std/__iterator/distance.h>`
`40`	`39`	`#include <cuda/std/__memory/allocator_traits.h>`
	`40`	`+#include <cuda/std/__memory/pointer_traits.h>`
`41`	`41`	`#include <cuda/std/__type_traits/is_convertible.h>`
`42`	`42`	`#include <cuda/std/__type_traits/is_trivially_copy_constructible.h>`
`43`	`43`	`#include <cuda/std/tuple>`
`@@ -164,7 +164,8 @@ template <typename System, typename Allocator, typename InputIterator, typename`
`164`	`164`	`_CCCL_HOST_DEVICE Pointer copy_construct_range(`
`165`	`165`	`thrust::execution_policy<System>& from_system, Allocator& a, InputIterator first, InputIterator last, Pointer result)`
`166`	`166`	`{`
`167`		`- if constexpr (needs_copy_construct_via_allocator<Allocator, typename pointer_element<Pointer>::type>)`
	`167`	`+ if constexpr (needs_copy_construct_via_allocator<Allocator,`
	`168`	`+ typename ::cuda::std::pointer_traits<Pointer>::element_type>)`
`168`	`169`	`{`
`169`	`170`	`return uninitialized_copy_with_allocator(a, from_system, allocator_system<Allocator>::get(a), first, last, result);`
`170`	`171`	`}`
`@@ -179,7 +180,8 @@ template <typename System, typename Allocator, typename InputIterator, typename`
`179`	`180`	`_CCCL_HOST_DEVICE Pointer copy_construct_range_n(`
`180`	`181`	`thrust::execution_policy<System>& from_system, Allocator& a, InputIterator first, Size n, Pointer result)`
`181`	`182`	`{`
`182`		`- if constexpr (needs_copy_construct_via_allocator<Allocator, typename pointer_element<Pointer>::type>)`
	`183`	`+ if constexpr (needs_copy_construct_via_allocator<Allocator,`
	`184`	`+ typename ::cuda::std::pointer_traits<Pointer>::element_type>)`
`183`	`185`	`{`
`184`	`186`	`return uninitialized_copy_with_allocator_n(a, from_system, allocator_system<Allocator>::get(a), first, n, result);`
`185`	`187`	`}`
Original file line number	Diff line number	Diff line change
`@@ -28,12 +28,12 @@`
`28`	`28`
`29`	`29`	`#include <thrust/detail/allocator/allocator_system.h>`
`30`	`30`	`#include <thrust/detail/type_traits.h>`
`31`		`-#include <thrust/detail/type_traits/pointer_traits.h>`
`32`	`31`	`#include <thrust/for_each.h>`
`33`	`32`	`#include <thrust/uninitialized_fill.h>`
`34`	`33`
`35`	`34`	`#include <cuda/std/__host_stdlib/memory>`
`36`	`35`	`#include <cuda/std/__memory/allocator_traits.h>`
	`36`	`+#include <cuda/std/__memory/pointer_traits.h>`
`37`	`37`
`38`	`38`	`THRUST_NAMESPACE_BEGIN`
`39`	`39`	`namespace detail`
`@@ -66,7 +66,7 @@ struct construct2_via_allocator`
`66`	`66`	`template <typename Allocator, typename Pointer, typename Size, typename T>`
`67`	`67`	`_CCCL_HOST_DEVICE void fill_construct_range(Allocator& a, Pointer p, Size n, const T& value)`
`68`	`68`	`{`
`69`		`- if constexpr (has_effectful_member_construct2<Allocator, typename pointer_element<Pointer>::type, T>)`
	`69`	`+ if constexpr (has_effectful_member_construct2<Allocator, typename ::cuda::std::pointer_traits<Pointer>::element_type, T>)`
`70`	`70`	`{`
`71`	`71`	`thrust::for_each_n(allocator_system<Allocator>::get(a), p, n, construct2_via_allocator<Allocator, T>{a, value});`
`72`	`72`	`}`
Original file line number	Diff line number	Diff line change
`@@ -28,11 +28,11 @@`
`28`	`28`
`29`	`29`	`#include <thrust/detail/allocator/allocator_system.h>`
`30`	`30`	`#include <thrust/detail/type_traits.h>`
`31`		`-#include <thrust/detail/type_traits/pointer_traits.h>`
`32`	`31`	`#include <thrust/for_each.h>`
`33`	`32`	`#include <thrust/uninitialized_fill.h>`
`34`	`33`
`35`	`34`	`#include <cuda/std/__memory/allocator_traits.h>`
	`35`	`+#include <cuda/std/__memory/pointer_traits.h>`
`36`	`36`
`37`	`37`	`THRUST_NAMESPACE_BEGIN`
`38`	`38`	`namespace detail`
`@@ -65,13 +65,15 @@ inline constexpr bool needs_default_construct_via_allocator<std::allocator<U>, T`
`65`	`65`	`template <typename Allocator, typename Pointer, typename Size>`
`66`	`66`	`_CCCL_HOST_DEVICE void value_initialize_range(Allocator& a, Pointer p, Size n)`
`67`	`67`	`{`
`68`		`- if constexpr (needs_default_construct_via_allocator<Allocator, typename pointer_element<Pointer>::type>)`
	`68`	`+ if constexpr (needs_default_construct_via_allocator<Allocator,`
	`69`	`+ typename ::cuda::std::pointer_traits<Pointer>::element_type>)`
`69`	`70`	`{`
`70`	`71`	`thrust::for_each_n(allocator_system<Allocator>::get(a), p, n, construct1_via_allocator<Allocator>{a});`
`71`	`72`	`}`
`72`	`73`	`else`
`73`	`74`	`{`
`74`		`- thrust::uninitialized_fill_n(allocator_system<Allocator>::get(a), p, n, typename pointer_element<Pointer>::type());`
	`75`	`+ thrust::uninitialized_fill_n(`
	`76`	`+ allocator_system<Allocator>::get(a), p, n, typename ::cuda::std::pointer_traits<Pointer>::element_type());`
`75`	`77`	`}`
`76`	`78`	`}`
`77`	`79`	`} // namespace detail`
Original file line number	Diff line number	Diff line change
`@@ -26,8 +26,6 @@`
`26`	`26`	`# pragma system_header`
`27`	`27`	`#endif // no system header`
`28`	`28`
`29`		`-#include <thrust/detail/type_traits/pointer_traits.h>`
`30`		`-`
`31`	`29`	`#include <cuda/std/__memory/pointer_traits.h>`
`32`	`30`
`33`	`31`	`THRUST_NAMESPACE_BEGIN`
`@@ -41,15 +39,15 @@ _CCCL_HOST_DEVICE auto raw_pointer_cast(Pointer ptr)`
`41`	`39`	`template <typename ToPointer, typename FromPointer>`
`42`	`40`	`_CCCL_HOST_DEVICE ToPointer reinterpret_pointer_cast(FromPointer ptr)`
`43`	`41`	`{`
`44`		`- using to_element = typename thrust::detail::pointer_element<ToPointer>::type;`
`45`		`- return ToPointer(reinterpret_cast<to_element*>(thrust::raw_pointer_cast(ptr)));`
	`42`	`+ using to_element = typename ::cuda::std::pointer_traits<ToPointer>::element_type;`
	`43`	`+ return ToPointer(reinterpret_cast<to_element*>(::cuda::std::to_address(ptr)));`
`46`	`44`	`}`
`47`	`45`
`48`	`46`	`template <typename ToPointer, typename FromPointer>`
`49`	`47`	`_CCCL_HOST_DEVICE ToPointer static_pointer_cast(FromPointer ptr)`
`50`	`48`	`{`
`51`		`- using to_element = typename thrust::detail::pointer_element<ToPointer>::type;`
`52`		`- return ToPointer(static_cast<to_element*>(thrust::raw_pointer_cast(ptr)));`
	`49`	`+ using to_element = typename ::cuda::std::pointer_traits<ToPointer>::element_type;`
	`50`	`+ return ToPointer(static_cast<to_element*>(::cuda::std::to_address(ptr)));`
`53`	`51`	`}`
`54`	`52`
`55`	`53`	`THRUST_NAMESPACE_END`