[CUB] Provide pointer traits to cache modified iterators#9623
Conversation
|
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (2)
🚧 Files skipped from review as they are similar to previous changes (2)
📝 WalkthroughSummary by CodeRabbit
Walkthrough
ChangesCacheModifiedInputIterator pointer_traits support
suggestion: important: The suggestion: The specialization omits 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: 2
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: a053f7a1-b012-455d-a2fc-d6aa17178267
📒 Files selected for processing (2)
cub/cub/iterator/cache_modified_input_iterator.cuhcub/test/catch2_test_iterator.cu
miscco
force-pushed
the
pointer_traits_cache_modifier
branch
from
391a2a5 to
bf6cb75
Compare
| [[nodiscard]] _CCCL_HOST_DEVICE_API static constexpr ValueType* to_address(const pointer iter) noexcept | ||
| { | ||
| return iter.__unwrap(); | ||
| } |
There was a problem hiding this comment.
@miscco should be we return a volatile ValueType* here in case the modifier is LOAD_VOLATILE?
| constexpr auto cache_modifier = c2h::get<1, TestType>::value; | ||
| constexpr int TEST_VALUES = 11000; | ||
|
|
||
| c2h::device_vector<T> d_data(TEST_VALUES); |
There was a problem hiding this comment.
| c2h::device_vector<T> d_data(TEST_VALUES); | |
| c2h::device_vector<T> d_data(TEST_VALUES, thrust::no_init); |
| cub::CacheModifiedInputIterator<cache_modifier, T>(const_cast<const T*>(thrust::raw_pointer_cast(d_data.data()))), | ||
| h_reference); | ||
|
|
||
| if constexpr (cache_modifier != cub::LOAD_CV && cache_modifier != cub::LOAD_VOLATILE) |
There was a problem hiding this comment.
Q: Why do we have to exclude volatile modifiers? Do they lead to errors?
|
|
||
| /// Structure dereference | ||
| _CCCL_HOST_DEVICE_API ValueType* __unwrap() const noexcept | ||
| { | ||
| return ptr; | ||
| } |
There was a problem hiding this comment.
Q: Do we need this? ptr is a public member :S
😬 CI Workflow Results🟥 Finished in 59m 13s: Pass: 1%/287 | Total: 1d 04h | Max: 59m 08s | Hits: 100%/6015See results here. |
2 participants
This provides
pointer_traitsto the cache modified iteratorsIts a bit dangerous, because it essentially drops the cache modifiers on the floor, but at the same time if we want to use a raw pointer anyway we might as well do the thing.
We could also just enable it for the default cases and leave .eg. volatile loads out