From 03ccb9958a0c079c9947eed34b4ffe2ba728a66f Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Thu, 16 Jul 2026 10:34:07 +0200
Subject: [PATCH 1/3] Use `thrust::is_contiguous_iterator_v` over
 `cuda::std::contiguous_iterator`

The former trait includes more types that we can turn into pointers
---
 cub/cub/agent/agent_find.cuh                       | 6 ++++--
 cub/cub/block/block_load.cuh                       | 4 +++-
 cub/cub/block/block_store.cuh                      | 5 ++++-
 cub/cub/iterator/cache_modified_input_iterator.cuh | 4 +++-
 cub/cub/thread/thread_load.cuh                     | 4 +++-
 cub/cub/thread/thread_store.cuh                    | 4 +++-
 6 files changed, 20 insertions(+), 7 deletions(-)

diff --git a/cub/cub/agent/agent_find.cuh b/cub/cub/agent/agent_find.cuh
index 89229ee21cd7..080e922d2047 100644
--- a/cub/cub/agent/agent_find.cuh
+++ b/cub/cub/agent/agent_find.cuh
@@ -10,13 +10,15 @@
 #include <cub/util_type.cuh>
 
 #include <thrust/detail/raw_reference_cast.h>
+#include <thrust/type_traits/is_contiguous_iterator.h>
 
 #include <cuda/__memory/is_aligned.h>
 #include <cuda/__type_traits/is_trivially_copyable.h>
+#include <cuda/std/__type_traits/integral_constant.h>
+
 #if !_CCCL_HAS_NV_ATOMIC_BUILTINS()
 #  include <cuda/atomic>
 #endif // !_CCCL_HAS_NV_ATOMIC_BUILTINS()
-#include <cuda/std/__type_traits/integral_constant.h>
 
 CUB_NAMESPACE_BEGIN
 namespace detail::find
@@ -40,7 +42,7 @@ struct agent_t
 
   // Can vectorize according to the policy if the input iterator is a native pointer to a primitive type
   static constexpr bool attempt_vectorization =
-    (VecSize > 1) && (ItemsPerThread % VecSize == 0) && (::cuda::std::contiguous_iterator<InputIteratorT>)
+    (VecSize > 1) && (ItemsPerThread % VecSize == 0) && (THRUST_NS_QUALIFIER::is_contiguous_iterator_v<InputIteratorT>)
     && ::cuda::is_trivially_copyable_v<InputT>;
 
   static constexpr CacheLoadModifier load_modifier = LoadModifier;
diff --git a/cub/cub/block/block_load.cuh b/cub/cub/block/block_load.cuh
index 077abb17ebc6..d5db8789071d 100644
--- a/cub/cub/block/block_load.cuh
+++ b/cub/cub/block/block_load.cuh
@@ -9,6 +9,8 @@
 
 #include <cub/config.cuh>
 
+#include <thrust/type_traits/is_contiguous_iterator.h>
+
 #if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
 #  pragma GCC system_header
 #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
@@ -1009,7 +1011,7 @@ public:
       {
         InternalLoadDirectBlockedVectorized<RandomAccessIterator::__modifier>(linear_tid, block_src_it.ptr, dst_items);
       }
-      else if constexpr (::cuda::std::contiguous_iterator<RandomAccessIterator>
+      else if constexpr (THRUST_NS_QUALIFIER::is_contiguous_iterator_v<RandomAccessIterator>
                          && ::cuda::std::__can_to_address<RandomAccessIterator>)
       {
         InternalLoadDirectBlockedVectorized<LOAD_DEFAULT>(linear_tid, ::cuda::std::to_address(block_src_it), dst_items);
diff --git a/cub/cub/block/block_store.cuh b/cub/cub/block/block_store.cuh
index 59113bfc4ab6..b2ccaebcd1bf 100644
--- a/cub/cub/block/block_store.cuh
+++ b/cub/cub/block/block_store.cuh
@@ -21,6 +21,8 @@
 #include <cub/util_ptx.cuh>
 #include <cub/util_type.cuh>
 
+#include <thrust/type_traits/is_contiguous_iterator.h>
+
 #include <cuda/std/__concepts/same_as.h>
 #include <cuda/std/__fwd/format.h>
 #include <cuda/std/__host_stdlib/ostream>
@@ -838,7 +840,8 @@ public:
     }
     else if constexpr (Algorithm == BLOCK_STORE_VECTORIZE)
     {
-      if constexpr (::cuda::std::contiguous_iterator<OutputIteratorT> && ::cuda::std::__can_to_address<OutputIteratorT>)
+      if constexpr (THRUST_NS_QUALIFIER::is_contiguous_iterator_v<OutputIteratorT>
+                    && ::cuda::std::__can_to_address<OutputIteratorT>)
       {
         StoreDirectBlockedVectorized(linear_tid, ::cuda::std::to_address(block_itr), items);
       }
diff --git a/cub/cub/iterator/cache_modified_input_iterator.cuh b/cub/cub/iterator/cache_modified_input_iterator.cuh
index e8055488c1a7..ff69f68a6974 100644
--- a/cub/cub/iterator/cache_modified_input_iterator.cuh
+++ b/cub/cub/iterator/cache_modified_input_iterator.cuh
@@ -24,6 +24,7 @@
 
 #include <thrust/detail/raw_pointer_cast.h>
 #include <thrust/iterator/iterator_facade.h>
+#include <thrust/type_traits/is_contiguous_iterator.h>
 
 #include <cuda/std/__host_stdlib/ostream>
 #include <cuda/std/__iterator/iterator_traits.h>
@@ -226,9 +227,10 @@ inline constexpr bool is_CacheModifiedInputIterator<CacheModifiedInputIterator<M
 template <CacheLoadModifier LoadModifier, typename Iterator>
 _CCCL_HOST_DEVICE _CCCL_FORCEINLINE auto try_make_cache_modified_iterator(Iterator it)
 {
-  if constexpr (::cuda::std::contiguous_iterator<Iterator>)
+  if constexpr (THRUST_NS_QUALIFIER::is_contiguous_iterator_v<Iterator>)
   {
     return CacheModifiedInputIterator<LoadModifier, it_value_t<Iterator>, it_difference_t<Iterator>>{
+      // FIXME(bgruber): this should be THRUST_NS_QUALIFIER::unwrap_contiguous_iterator
       THRUST_NS_QUALIFIER::raw_pointer_cast(&*it)};
   }
   else
diff --git a/cub/cub/thread/thread_load.cuh b/cub/cub/thread/thread_load.cuh
index d98eb37e27a6..c98d1b311b53 100644
--- a/cub/cub/thread/thread_load.cuh
+++ b/cub/cub/thread/thread_load.cuh
@@ -22,6 +22,8 @@
 #include <cub/util_ptx.cuh>
 #include <cub/util_type.cuh>
 
+#include <thrust/type_traits/is_contiguous_iterator.h>
+
 #include <cuda/std/__concepts/same_as.h>
 #include <cuda/std/__fwd/format.h>
 #include <cuda/std/__host_stdlib/ostream>
@@ -348,7 +350,7 @@ template <CacheLoadModifier MODIFIER, typename RandomAccessIterator>
 _CCCL_DEVICE _CCCL_FORCEINLINE detail::it_value_t<RandomAccessIterator> ThreadLoad(RandomAccessIterator itr)
 {
   using T = detail::it_value_t<RandomAccessIterator>;
-  if constexpr (!::cuda::std::__4::contiguous_iterator<RandomAccessIterator> || MODIFIER == LOAD_DEFAULT)
+  if constexpr (!THRUST_NS_QUALIFIER::is_contiguous_iterator_v<RandomAccessIterator> || MODIFIER == LOAD_DEFAULT)
   {
     return *itr;
   }
diff --git a/cub/cub/thread/thread_store.cuh b/cub/cub/thread/thread_store.cuh
index c4cb9307d182..2624d3e2361b 100644
--- a/cub/cub/thread/thread_store.cuh
+++ b/cub/cub/thread/thread_store.cuh
@@ -20,6 +20,8 @@
 #include <cub/util_ptx.cuh>
 #include <cub/util_type.cuh>
 
+#include <thrust/type_traits/is_contiguous_iterator.h>
+
 #include <cuda/std/__iterator/concepts.h>
 #include <cuda/std/__memory/pointer_traits.h>
 #include <cuda/std/__type_traits/integral_constant.h>
@@ -316,7 +318,7 @@ ThreadStore(T* ptr, T val, detail::constant_t<MODIFIER> /*modifier*/, ::cuda::st
 template <CacheStoreModifier MODIFIER, typename OutputIteratorT, typename T>
 _CCCL_DEVICE _CCCL_FORCEINLINE void ThreadStore(OutputIteratorT itr, T val)
 {
-  if constexpr (!::cuda::std::contiguous_iterator<OutputIteratorT> || MODIFIER == STORE_DEFAULT)
+  if constexpr (!THRUST_NS_QUALIFIER::is_contiguous_iterator_v<OutputIteratorT> || MODIFIER == STORE_DEFAULT)
   {
     *itr = val;
   }

From fc45b0c91c3988ec031a461f56d9d3121570ef94 Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Thu, 23 Jul 2026 10:47:42 +0200
Subject: [PATCH 2/3] Fix cache-modified iterator construction for broader
 contiguous iterators

With is_contiguous_iterator_v now matching thrust fancy iterators (e.g.
normal_iterator<device_ptr<T>>), AgentDifference constructed its LoadIt
directly from the raw input iterator, which no longer converts. Build it
through try_make_cache_modified_iterator instead, and unwrap contiguous
iterators via unwrap_contiguous_iterator (to_address) rather than
raw_pointer_cast(&*it).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
---
 cub/cub/iterator/cache_modified_input_iterator.cuh | 5 ++---
 1 file changed, 2 insertions(+), 3 deletions(-)

diff --git a/cub/cub/iterator/cache_modified_input_iterator.cuh b/cub/cub/iterator/cache_modified_input_iterator.cuh
index ff69f68a6974..fd80b8a83636 100644
--- a/cub/cub/iterator/cache_modified_input_iterator.cuh
+++ b/cub/cub/iterator/cache_modified_input_iterator.cuh
@@ -22,9 +22,9 @@
 #include <cub/thread/thread_load.cuh>
 #include <cub/thread/thread_store.cuh>
 
-#include <thrust/detail/raw_pointer_cast.h>
 #include <thrust/iterator/iterator_facade.h>
 #include <thrust/type_traits/is_contiguous_iterator.h>
+#include <thrust/type_traits/unwrap_contiguous_iterator.h>
 
 #include <cuda/std/__host_stdlib/ostream>
 #include <cuda/std/__iterator/iterator_traits.h>
@@ -230,8 +230,7 @@ _CCCL_HOST_DEVICE _CCCL_FORCEINLINE auto try_make_cache_modified_iterator(Iterat
   if constexpr (THRUST_NS_QUALIFIER::is_contiguous_iterator_v<Iterator>)
   {
     return CacheModifiedInputIterator<LoadModifier, it_value_t<Iterator>, it_difference_t<Iterator>>{
-      // FIXME(bgruber): this should be THRUST_NS_QUALIFIER::unwrap_contiguous_iterator
-      THRUST_NS_QUALIFIER::raw_pointer_cast(&*it)};
+      THRUST_NS_QUALIFIER::unwrap_contiguous_iterator(it)};
   }
   else
   {

From 6370f24801bf5890764dfab92215c8594420a39a Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Mon, 17 Aug 2026 00:40:16 +0200
Subject: [PATCH 3/3] Apply feedback

---
 cub/cub/block/block_load.cuh                       | 3 +--
 cub/cub/block/block_store.cuh                      | 3 +--
 cub/cub/iterator/cache_modified_input_iterator.cuh | 2 +-
 3 files changed, 3 insertions(+), 5 deletions(-)

diff --git a/cub/cub/block/block_load.cuh b/cub/cub/block/block_load.cuh
index d5db8789071d..59c154879b74 100644
--- a/cub/cub/block/block_load.cuh
+++ b/cub/cub/block/block_load.cuh
@@ -1011,8 +1011,7 @@ public:
       {
         InternalLoadDirectBlockedVectorized<RandomAccessIterator::__modifier>(linear_tid, block_src_it.ptr, dst_items);
       }
-      else if constexpr (THRUST_NS_QUALIFIER::is_contiguous_iterator_v<RandomAccessIterator>
-                         && ::cuda::std::__can_to_address<RandomAccessIterator>)
+      else if constexpr (THRUST_NS_QUALIFIER::is_contiguous_iterator_v<RandomAccessIterator>)
       {
         InternalLoadDirectBlockedVectorized<LOAD_DEFAULT>(linear_tid, ::cuda::std::to_address(block_src_it), dst_items);
       }
diff --git a/cub/cub/block/block_store.cuh b/cub/cub/block/block_store.cuh
index b2ccaebcd1bf..b8366f38cfa2 100644
--- a/cub/cub/block/block_store.cuh
+++ b/cub/cub/block/block_store.cuh
@@ -840,8 +840,7 @@ public:
     }
     else if constexpr (Algorithm == BLOCK_STORE_VECTORIZE)
     {
-      if constexpr (THRUST_NS_QUALIFIER::is_contiguous_iterator_v<OutputIteratorT>
-                    && ::cuda::std::__can_to_address<OutputIteratorT>)
+      if constexpr (THRUST_NS_QUALIFIER::is_contiguous_iterator_v<OutputIteratorT>)
       {
         StoreDirectBlockedVectorized(linear_tid, ::cuda::std::to_address(block_itr), items);
       }
diff --git a/cub/cub/iterator/cache_modified_input_iterator.cuh b/cub/cub/iterator/cache_modified_input_iterator.cuh
index fd80b8a83636..01e1f3ddec93 100644
--- a/cub/cub/iterator/cache_modified_input_iterator.cuh
+++ b/cub/cub/iterator/cache_modified_input_iterator.cuh
@@ -230,7 +230,7 @@ _CCCL_HOST_DEVICE _CCCL_FORCEINLINE auto try_make_cache_modified_iterator(Iterat
   if constexpr (THRUST_NS_QUALIFIER::is_contiguous_iterator_v<Iterator>)
   {
     return CacheModifiedInputIterator<LoadModifier, it_value_t<Iterator>, it_difference_t<Iterator>>{
-      THRUST_NS_QUALIFIER::unwrap_contiguous_iterator(it)};
+      ::cuda::std::to_address(it)};
   }
   else
   {