Use thrust::is_contiguous_iterator_v over cuda::std::contiguous_iterator#9894
Use thrust::is_contiguous_iterator_v over cuda::std::contiguous_iterator#9894bernhardmgruber wants to merge 1 commit into
thrust::is_contiguous_iterator_v over cuda::std::contiguous_iterator#9894Conversation
…erator` The former trait includes more types that we can turn into pointers
📝 WalkthroughSummary by CodeRabbit
WalkthroughChangesCUB’s load, store, and cache-modified iterator paths now use Contiguous load paths
Contiguous store paths
Cache-modified iterator path
Possibly related issues
Possibly related PRs
Suggested reviewers: Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 1
🧹 Nitpick comments (2)
cub/cub/block/block_load.cuh (1)
1014-1017: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: 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 liftsuggestion: 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
📒 Files selected for processing (6)
cub/cub/agent/agent_find.cuhcub/cub/block/block_load.cuhcub/cub/block/block_store.cuhcub/cub/iterator/cache_modified_input_iterator.cuhcub/cub/thread/thread_load.cuhcub/cub/thread/thread_store.cuh
| (VecSize > 1) && (ItemsPerThread % VecSize == 0) && (THRUST_NS_QUALIFIER::is_contiguous_iterator_v<InputIteratorT>) | ||
| && THRUST_NS_QUALIFIER::is_trivially_relocatable_v<InputT>; |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟠 Major | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
#!/bin/bash
set -euo pipefail
printf '\n== agent_find ==\n'
sed -n '1,120p' cub/cub/agent/agent_find.cuh
printf '\n== thread_load ==\n'
sed -n '320,390p' cub/cub/thread/thread_load.cuh
printf '\n== thread_store ==\n'
sed -n '300,360p' cub/cub/thread/thread_store.cuh
printf '\n== contiguity trait defs/usages ==\n'
rg -n "is_contiguous_iterator_v|__can_to_address|to_address\\(" thrust cub -g '!**/build/**' -g '!**/dist/**'
printf '\n== trait definition file(s) ==\n'
fd -a "is_contiguous_iterator.h" thrust cubRepository: NVIDIA/cccl
Length of output: 26490
🏁 Script executed:
#!/bin/bash
set -euo pipefail
printf '\n== locate trait definition ==\n'
fd -a "is_contiguous_iterator.h" .
printf '\n== search minimal trait-related symbols ==\n'
rg -n "is_contiguous_iterator|__can_to_address|contiguous_iterator" . -g '!**/build/**' -g '!**/dist/**' -g '!**/.git/**'Repository: NVIDIA/cccl
Length of output: 50367
🌐 Web query:
thrust type_traits is_contiguous_iterator.h definition contiguous_iterator to_address
💡 Result:
Thrust provides several tools within its type traits and iterator utilities to handle contiguous iterators and address extraction [1][2]. is_contiguous_iterator thrust::is_contiguous_iterator is a type trait used to determine if an iterator points to elements that are contiguous in memory [3]. It is defined in the header thrust/type_traits/is_contiguous_iterator.h [1]. The trait returns a true_type if the given Iterator satisfies the ContiguousIterator concept, and false_type otherwise [3]. A corresponding convenience variable template, thrust::is_contiguous_iterator_v, is also available for ease of use [2][4]. to_address and Unwrapping Thrust does not define a custom to_address function; instead, it relies on the C++ standard library's std::to_address, typically accessed via the CUDA Core Compute Libraries (CCCL) integration (often via cuda::std::to_address) [5]. To bridge the gap between Thrust iterators and raw pointers, Thrust provides specific utilities: - unwrap_contiguous_iterator: This function takes a contiguous iterator and converts it to its underlying raw pointer [6]. It includes a static_assert to ensure the iterator is indeed contiguous [6]. - try_unwrap_contiguous_iterator: This utility conditionally unwraps an iterator to a raw pointer if it is contiguous; otherwise, it returns the iterator unmodified [6]. These utilities are located in thrust/type_traits/unwrap_contiguous_iterator.h and are designed to facilitate high-performance operations by allowing direct memory access when safe [6].
Citations:
- 1: https://github.com/NVIDIA/thrust/blob/main/thrust/type_traits/is_contiguous_iterator.h
- 2: https://gevtushenko.github.io/cccl/thrust/api_docs/utility/type_traits.html
- 3: https://gevtushenko.github.io/cccl/thrust/api/group__type__traits_1gafcd0e9ee2ab2f7123caf84f6c85a691b.html
- 4: https://gevtushenko.github.io/cccl/thrust/api/group__type__traits.html
- 5: https://cppreference.com/cpp/memory/to_address
- 6: https://github.com/NVIDIA/cccl/blob/0d9331dc/thrust/thrust/type_traits/unwrap_contiguous_iterator.h
🏁 Script executed:
#!/bin/bash
set -euo pipefail
sed -n '1,120p' thrust/thrust/type_traits/unwrap_contiguous_iterator.h
printf '\n== block_load branch ==\n'
sed -n '1000,1035p' cub/cub/block/block_load.cuh
printf '\n== block_store branch ==\n'
sed -n '830,860p' cub/cub/block/block_store.cuh
printf '\n== contig trait definition ==\n'
sed -n '1,180p' thrust/thrust/type_traits/is_contiguous_iterator.hRepository: NVIDIA/cccl
Length of output: 9635
🏁 Script executed:
#!/bin/bash
set -euo pipefail
printf '== proclaim_contiguous_iterator uses ==\n'
rg -n "THRUST_PROCLAIM_CONTIGUOUS_ITERATOR|proclaim_contiguous_iterator<" . -g '!**/build/**' -g '!**/dist/**' -g '!**/.git/**'
printf '\n== thread_load/thread_store callers ==\n'
rg -n "ThreadLoad<|ThreadLoad\\(|ThreadStore<|ThreadStore\\(" cub thrust -g '!**/build/**' -g '!**/dist/**' -g '!**/.git/**' | sed -n '1,120p'Repository: NVIDIA/cccl
Length of output: 12658
🏁 Script executed:
#!/bin/bash
set -euo pipefail
printf '== normal_iterator ==\n'
sed -n '1,140p' thrust/thrust/iterator/detail/normal_iterator.h
printf '\n== tagged_iterator ==\n'
sed -n '1,120p' thrust/thrust/iterator/detail/tagged_iterator.h
printf '\n== to_address support search ==\n'
rg -n "to_address\\(|__can_to_address|pointer_traits.*to_address|struct pointer_traits" thrust/thrust/iterator thrust/thrust/detail libcudacxx/include cuda -g '!**/build/**' -g '!**/dist/**' -g '!**/.git/**'Repository: NVIDIA/cccl
Length of output: 16545
🏁 Script executed:
#!/bin/bash
set -euo pipefail
printf '== iterator_adaptor operator-> / base access ==\n'
rg -n "operator->|to_address|base\\(" thrust/thrust/iterator/iterator_adaptor.h thrust/thrust/iterator/detail/normal_iterator.h thrust/thrust/iterator/detail/tagged_iterator.h
printf '\n== iterator_adaptor excerpt ==\n'
sed -n '1,260p' thrust/thrust/iterator/iterator_adaptor.hRepository: NVIDIA/cccl
Length of output: 8566
🏁 Script executed:
#!/bin/bash
set -euo pipefail
nl -ba cub/cub/agent/agent_find.cuh | sed -n '1,260p'Repository: NVIDIA/cccl
Length of output: 189
🏁 Script executed:
#!/bin/bash
set -euo pipefail
printf '== agent_find slice ==\n'
awk 'NR>=1 && NR<=260 { printf "%4d %s\n", NR, $0 }' cub/cub/agent/agent_find.cuh | sed -n '1,260p'Repository: NVIDIA/cccl
Length of output: 8416
🏁 Script executed:
#!/bin/bash
set -euo pipefail
printf '== agent_find vectorization control flow ==\n'
rg -n "attempt_vectorization|ConsumeTile\\(|is_aligned_and_full_tile\\(" cub/cub/agent/agent_find.cuh
printf '\n== surrounding lines ==\n'
awk 'NR>=120 && NR<=260 { printf "%4d %s\n", NR, $0 }' cub/cub/agent/agent_find.cuhRepository: NVIDIA/cccl
Length of output: 4198
🏁 Script executed:
#!/bin/bash
set -euo pipefail
rg -n "operator->|to_address" thrust/thrust/iterator/iterator_facade.h thrust/thrust/iterator/detail/iterator_facade* thrust/thrust/iterator/iterator_adaptor.h thrust/thrust/iterator/detail/normal_iterator.h thrust/thrust/iterator/detail/tagged_iterator.hRepository: NVIDIA/cccl
Length of output: 685
🏁 Script executed:
#!/bin/bash
set -euo pipefail
rg -n "operator->|arrow|to_address" thrust/thrust/iterator/iterator_facade.h thrust/thrust/iterator/detail/iterator_facade*.h thrust/thrust/iterator/detail/iterator_facade*.cuh thrust/thrust/iterator/iterator_adaptor.hRepository: NVIDIA/cccl
Length of output: 524
🏁 Script executed:
#!/bin/bash
set -euo pipefail
rg -n "operator->" thrust/thrust/iterator/iterator_facade.h thrust/thrust/iterator/detail -g '!**/build/**' -g '!**/dist/**'Repository: NVIDIA/cccl
Length of output: 433
🏁 Script executed:
#!/bin/bash
set -euo pipefail
rg -n "attempt_vectorization|ConsumeTile\\(|is_aligned_and_full_tile\\(" cub/cub/agent/agent_find.cuh
printf '\n== agent_find around vectorization ==\n'
awk 'NR>=120 && NR<=260 { printf "%4d %s\n", NR, $0 }' cub/cub/agent/agent_find.cuhRepository: NVIDIA/cccl
Length of output: 4168
important: is_contiguous_iterator_v is broader than the raw-pointer path these overloads need.
cub/cub/agent/agent_find.cuh: keep vectorization pointer-only; the vectorized tile path still does rawreinterpret_castloads, andis_aligned_and_full_tile()already requiresstd::is_pointer_v.cub/cub/thread/thread_load.cuh/cub/cub/thread/thread_store.cuh: keep the::cuda::std::__can_to_addressguard so contiguous iterators that can't be unwrapped still take the dereference fallback.
📍 Affects 3 files
cub/cub/agent/agent_find.cuh#L45-L46(this comment)cub/cub/thread/thread_load.cuh#L353-L359cub/cub/thread/thread_store.cuh#L321-L327
Source: Learnings
😬 CI Workflow Results🟥 Finished in 1h 48m: Pass: 14%/287 | Total: 3d 05h | Max: 1h 25m | Hits: 19%/155820See results here. |
1 participant
The former trait includes more types that we can turn into pointers