Skip to content

Use thrust::is_contiguous_iterator_v over cuda::std::contiguous_iterator #9894

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
bernhardmgruber wants to merge 3 commits into
base: main
Choose a base branch
from
Open

Use thrust::is_contiguous_iterator_v over cuda::std::contiguous_iterator #9894

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
03ccb99
Use `thrust::is_contiguous_iterator_v` over `cuda::std::contiguous_it…
bernhardmgruber Jul 16, 2026
fc45b0c
Fix cache-modified iterator construction for broader contiguous itera…
bernhardmgruber Jul 23, 2026
6370f24
Apply feedback
bernhardmgruber Aug 16, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 4 additions & 2 deletions cub/cub/agent/agent_find.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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>)

@miscco miscco Jul 24, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I am not sure whether this is something we really want in the long run. We should rather make sure that to_address works properly

At the same time we do need to worry about cache modified iterators

&& ::cuda::is_trivially_copyable_v<InputT>;

static constexpr CacheLoadModifier load_modifier = LoadModifier;
Expand Down
5 changes: 3 additions & 2 deletions cub/cub/block/block_load.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down Expand Up @@ -1009,8 +1011,7 @@ public:
{
InternalLoadDirectBlockedVectorized<RandomAccessIterator::__modifier>(linear_tid, block_src_it.ptr, dst_items);
}
else if constexpr (::cuda::std::contiguous_iterator<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);
}
Expand Down
4 changes: 3 additions & 1 deletion cub/cub/block/block_store.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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>
Expand Down Expand Up @@ -838,7 +840,7 @@ 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>)
{
StoreDirectBlockedVectorized(linear_tid, ::cuda::std::to_address(block_itr), items);
}
Expand Down
7 changes: 4 additions & 3 deletions cub/cub/iterator/cache_modified_input_iterator.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,8 +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>
Expand Down Expand Up @@ -226,10 +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>>{
THRUST_NS_QUALIFIER::raw_pointer_cast(&*it)};
::cuda::std::to_address(it)};
}
else
{
Expand Down
4 changes: 3 additions & 1 deletion cub/cub/thread/thread_load.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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>
Expand Down Expand Up @@ -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;
}
Expand Down
4 changes: 3 additions & 1 deletion cub/cub/thread/thread_store.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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>
Expand Down Expand Up @@ -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;
}
Expand Down
Loading