Use `thrust::is_contiguous_iterator_v` over `cuda::std::contiguous_iterator` by bernhardmgruber · Pull Request #9894 · NVIDIA/cccl

bernhardmgruber · 2026-07-16T08:34:35Z

The former trait includes more types that we can turn into pointers

coderabbitai · 2026-07-16T08:45:26Z

📝 Walkthrough

Summary by CodeRabbit

Performance
- Improved detection of contiguous iterators across loading, storing, caching, and vectorization operations.
- Enables more reliable optimized paths for contiguous data while preserving fallback behavior for other iterator types.
Compatibility
- Improved interoperability with Thrust iterator types.
- No changes to public APIs or expected operation.

Walkthrough

CUB load and store paths now use Thrust’s contiguous-iterator trait. Cache-modified iterator construction uses the trait and cuda::std::to_address. Non-contiguous iterator fallbacks remain unchanged.

Changes

Contiguous iterator trait migration

Layer / File(s)	Summary
Load-path contiguity checks `cub/cub/agent/agent_find.cuh`, `cub/cub/block/block_load.cuh`, `cub/cub/thread/thread_load.cuh`	Vectorization and load selection use `THRUST_NS_QUALIFIER::is_contiguous_iterator_v`. Existing fallback and modifier conditions remain unchanged.
Store-path contiguity checks `cub/cub/block/block_store.cuh`, `cub/cub/thread/thread_store.cuh`	Store selection uses Thrust’s contiguous-iterator trait. Existing modifier conditions remain unchanged.
Cache-modified iterator adaptation `cub/cub/iterator/cache_modified_input_iterator.cuh`	Contiguous iterators are detected with Thrust’s trait and converted with `cuda::std::to_address`. Non-contiguous iterators remain unchanged.

Suggested reviewers: elstehle, griwes, naderalawar

Merge Risk: 🔵 Low · up to 6370f

The PR changes contiguous-iterator handling and has a bounded include-hygiene follow-up to make the dependency on address conversion explicit and remove an unused include. It is otherwise mergeable with owner awareness.

_{Comment @coderabbitai help to get the list of available commands.}

coderabbitai

Actionable comments posted: 1

🧹 Nitpick comments (2)

cub/cub/block/block_load.cuh (1)

1014-1017: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

suggestion: Align both vectorization docs with the new selection rule. The implementation now uses Thrust’s contiguous-iterator trait plus __can_to_address, so the existing “simple pointer” restrictions are stale.

cub/cub/block/block_load.cuh#L1014-L1017: update the BlockLoad vectorization requirements.

cub/cub/block/block_store.cuh#L843-L846: update the BlockStore vectorization requirements. (raw.githubusercontent.com)

cub/cub/iterator/cache_modified_input_iterator.cuh (1)

230-234: 🎯 Functional Correctness | 🔵 Trivial | 🏗️ Heavy lift

suggestion: Validate the broadened iterator-unwrapping contract. This branch now admits additional Thrust pointer-like iterators, but still derives the base pointer through raw_pointer_cast(&*it) and leaves a FIXME for a dedicated unwrap facility. Add compile coverage for the newly accepted iterator types and either adopt the supported unwrap operation or document why this expression is valid for every trait-positive type. (raw.githubusercontent.com)

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 9ae78556-04bf-4cfb-8ab5-2b87557b17f4

📥 Commits

Reviewing files that changed from the base of the PR and between 0d919ed and cf41e69.

📒 Files selected for processing (6)

cub/cub/agent/agent_find.cuh
cub/cub/block/block_load.cuh
cub/cub/block/block_store.cuh
cub/cub/iterator/cache_modified_input_iterator.cuh
cub/cub/thread/thread_load.cuh
cub/cub/thread/thread_store.cuh

Jacobfaib

Can we not just extend cuda::std::contiguous_iterator to cover these cases? Looking at the definition of thrust::is_contiguous_iterator_v:

template <typename Iterator>
inline constexpr bool is_contiguous_iterator_impl_v =
  ::cuda::std::contiguous_iterator<Iterator> || is_thrust_pointer_v<Iterator> || is_libcxx_wrap_iter_v<Iterator>
  || is_libstdcxx_normal_iterator_v<Iterator> || is_msvc_contiguous_iterator_v<Iterator>
  || proclaim_contiguous_iterator<Iterator>::value;

The stdlib ones are OK (and correct even as a transparent extension) while is_thrust_pointer_v should be similarly doable since we can change the thrust pointer definitions to themselves proclaim contiguousness via the normal ways.

I think the only ones we wouldn't be able to do is proclaim_contiguous_iterator<Iterator>, but maybe we should provide a cuda::contiguous_iterator concept that covers this in addition?

cc @miscco

miscco · 2026-07-23T13:04:50Z

Can we not just extend cuda::std::contiguous_iterator to cover these cases? Looking at the definition of thrust::is_contiguous_iterator_v:

We cannot do that under any circumstances

The only thing we can and should do is introduce an extended version of it under a different name.

I believe we should rather use the standard to_address which is only specialized for pointer-like things.

miscco · 2026-07-24T07:23:20Z

  // 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>)


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

miscco · 2026-07-24T07:24:11Z

+      else if constexpr (THRUST_NS_QUALIFIER::is_contiguous_iterator_v<RandomAccessIterator>
                         && ::cuda::std::__can_to_address<RandomAccessIterator>)


I believe this is actually too much ^^

contiguous_iterator already requires to_address, so I believe the solution here is to just check for to_address. We can be a bit saver and also require random_access operations though

Suggested change

else if constexpr (THRUST_NS_QUALIFIER::is_contiguous_iterator_v<RandomAccessIterator>

&& ::cuda::std::__can_to_address<RandomAccessIterator>)

else if constexpr (::cuda::std::__has_random_access_traversal<RandomAccessIterator>

&& ::cuda::std::__can_to_address<RandomAccessIterator>)

Wait, why don't we just require cuda::std::__has_contiguous_traversal here? Wouldn't that solve our problem?

@miscco I extended the Thrust test for contiguous iterators to also check for __can_to_address and it seems the latter works for any contiguous Thrust iterator: #10734

While working on #10734 I realized that to_address works for too many types. For example, we __can_to_address<reverse_iterator<...>> but turning that into a pointer to read from the underlying memory does not preserve the right semantics. So I think we need to check for is_contiguous_iterator_v. But checking __can_to_address additionally is too much.

miscco · 2026-07-24T07:24:56Z

+      if constexpr (THRUST_NS_QUALIFIER::is_contiguous_iterator_v<OutputIteratorT>
+                    && ::cuda::std::__can_to_address<OutputIteratorT>)


Ditto this should probably just be

Suggested change

if constexpr (THRUST_NS_QUALIFIER::is_contiguous_iterator_v<OutputIteratorT>

&& ::cuda::std::__can_to_address<OutputIteratorT>)

if constexpr (::cuda::std::__has_random_access_traversal<OutputIteratorT>

&& ::cuda::std::__can_to_address<OutputIteratorT>)

…erator` The former trait includes more types that we can turn into pointers

…tors 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>

coderabbitai · 2026-08-16T22:29:36Z

Note

GitHub couldn't provide a complete incremental comparison for this pull request, so CodeRabbit is performing a full review instead. This review may take a little longer.

coderabbitai

Caution

Some comments are outside the diff and can’t be posted inline due to platform limitations.

⚠️ Outside diff range comments (1)

cub/cub/iterator/cache_modified_input_iterator.cuh (1)

25-27: 📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win

important: Include <cuda/std/__memory/pointer_traits.h> directly for cuda::std::to_address, and remove the unused <thrust/type_traits/unwrap_contiguous_iterator.h> include.

Source: Coding guidelines

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 6831412c-c2a4-4ddb-9e46-9cb10d534f11

📥 Commits

Reviewing files that changed from the base of the PR and between fc45b0c and 6370f24.

📒 Files selected for processing (3)

cub/cub/block/block_load.cuh
cub/cub/block/block_store.cuh
cub/cub/iterator/cache_modified_input_iterator.cuh

🚧 Files skipped from review as they are similar to previous changes (2)

cub/cub/block/block_store.cuh
cub/cub/block/block_load.cuh

Included review availability: Your plan includes up to 12 reviews per rolling hour; 8 remain after this review.

github-actions · 2026-08-16T23:22:55Z

⏱️ CCCL compile-time benchmark comparison: Public headers compile-time bench

Result: 0 regression row(s), 2 improvement row(s) above threshold.

Run	Value
Config	`public-headers-gcc13`
Baseline	`origin/main`
Preset	`all-dev`
Targets	`cub.headers.base, thrust.cpp.cuda.headers.base, libcudacxx.test.public_headers`
GPU / launch args	`rtx2080` / `--cuda 13.3 --host gcc13`

Artifacts: reports and traces

Direct file processing

-f file-processing exclusive --sort total

🟢 Direct file processing — Improvements

Rank	Improvement impact	Selected Δ	Baseline	Current	Event	Matched traces
1	`0.660744`	`-0.660744`	`5.414224`	`4.753480`	Processing Header File: `libcudacxx/include/cuda/std/__cccl/prologue.h`	552
2	`0.212704`	`-0.212704`	`1.669849`	`1.457145`	Processing Header File: `libcudacxx/include/cuda/std/__cccl/epilogue.h`	552

github-actions · 2026-08-17T01:38:51Z

🥳 CI Workflow Results

🟩 Finished in 2h 54m: Pass: 100%/272 | Total: 13d 00h | Max: 2h 54m | Hits: 13%/1251260

See results here.

bernhardmgruber requested a review from a team as a code owner July 16, 2026 08:34

github-project-automation Bot added this to CCCL Jul 16, 2026

bernhardmgruber requested a review from srinivasyadav18 July 16, 2026 08:34

github-project-automation Bot moved this to Todo in CCCL Jul 16, 2026

cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Jul 16, 2026

coderabbitai Bot reviewed Jul 16, 2026

View reviewed changes

Comment thread cub/cub/agent/agent_find.cuh Outdated

This comment has been minimized.

Sign in to view

coderabbitai Bot mentioned this pull request Jul 16, 2026

[CUB] Fix DeviceAdjacentDifference for cuda::device_buffer iterators #9861

Merged

2 tasks

bernhardmgruber requested review from a team as code owners July 23, 2026 08:53

bernhardmgruber requested review from NaderAlAwar, elstehle and griwes July 23, 2026 08:53

bernhardmgruber force-pushed the fix_contig_iter branch from f84862e to 1fadbdf Compare July 23, 2026 08:56

Jacobfaib reviewed Jul 23, 2026

View reviewed changes

Comment thread cub/cub/iterator/cache_modified_input_iterator.cuh Outdated

This comment has been minimized.

Sign in to view

miscco reviewed Jul 24, 2026

View reviewed changes

elstehle mentioned this pull request Jul 29, 2026

Implement Segmented TopK using Thread-Block-Clusters #9224

Merged

3 tasks

bernhardmgruber and others added 2 commits August 17, 2026 00:29

Use thrust::is_contiguous_iterator_v over `cuda::std::contiguous_it…

03ccb99

…erator` The former trait includes more types that we can turn into pointers

bernhardmgruber force-pushed the fix_contig_iter branch from 1fadbdf to fc45b0c Compare August 16, 2026 22:29

Apply feedback

6370f24

coderabbitai Bot reviewed Aug 16, 2026

View reviewed changes

		else if constexpr (THRUST_NS_QUALIFIER::is_contiguous_iterator_v<RandomAccessIterator>
		&& ::cuda::std::__can_to_address<RandomAccessIterator>)

		if constexpr (THRUST_NS_QUALIFIER::is_contiguous_iterator_v<OutputIteratorT>
		&& ::cuda::std::__can_to_address<OutputIteratorT>)

Uh oh!

Conversation

bernhardmgruber commented Jul 16, 2026

Uh oh!

coderabbitai Bot commented Jul 16, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary by CodeRabbit

Walkthrough

Changes

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

This comment has been minimized.

Jacobfaib left a comment • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

miscco commented Jul 23, 2026

Uh oh!

This comment has been minimized.

miscco Jul 24, 2026

Choose a reason for hiding this comment

Uh oh!

miscco Jul 24, 2026

Choose a reason for hiding this comment

Uh oh!

Jacobfaib Jul 24, 2026

Choose a reason for hiding this comment

Uh oh!

bernhardmgruber Aug 10, 2026

Choose a reason for hiding this comment

Uh oh!

bernhardmgruber Aug 16, 2026

Choose a reason for hiding this comment

Uh oh!

miscco Jul 24, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

coderabbitai Bot commented Aug 16, 2026

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

github-actions Bot commented Aug 16, 2026

⏱️ CCCL compile-time benchmark comparison: Public headers compile-time bench

Direct file processing

Uh oh!

github-actions Bot commented Aug 17, 2026

🥳 CI Workflow Results

🟩 Finished in 2h 54m: Pass: 100%/272 | Total: 13d 00h | Max: 2h 54m | Hits: 13%/1251260

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

coderabbitai Bot commented Jul 16, 2026 •

edited

Loading

Jacobfaib left a comment •

edited

Loading