Repository navigation
Refactor iterator usage in thrust/cuda find() - #6019
Merged
Merged
Conversation
Contributor
|
Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually. Contributors can view more details about this message here. |
bernhardmgruber
commented
Sep 25, 2025
Comment on lines
-215
to
-216
| // FIXME(bgruber): we should also be able to use transform_iterator here, but it makes nvc++ hang. See: | ||
| // https://github.com/NVIDIA/cccl/issues/3594. The problem does not occur with nvcc, so we could not add a test :/ |
Contributor
Author
There was a problem hiding this comment.
nvc++ is fine with this code
bernhardmgruber
commented
Sep 25, 2025
|
|
||
| result_type result = reduce( | ||
| const result_type result = reduce( | ||
| policy, interval_begin, interval_end, result_type(false, interval_end - begin), __find_if::functor<result_type>()); |
Contributor
Author
There was a problem hiding this comment.
I get an error here at interval_end - begin because zip_iterator::operator- is disabled because _Constraints::__all_sized_sentinel fails because the transform iterator above is not default constructible (the predicate isn't default constructible). The latter is fine, so I wonder why the constraint on zip_iterator::operator- is needed.
Contributor
There was a problem hiding this comment.
This is from ::std::ranges::zip_view
It kind of also boils down to the iterator concepts, because we want a forward iterator to be regular which requires default constructibility
Contributor
Author
There was a problem hiding this comment.
That's a problem now, because if we replace any thrust zip iterator by a cuda one, it can be a source breaking change for a user.
Contributor
Author
There was a problem hiding this comment.
Discussed with @miscco offline: he will relax the constraint on operator-
Contributor
There was a problem hiding this comment.
Put this in #6036 because we had to fix zip_iterator for NVHPC anyway
bernhardmgruber
force-pushed
the
ref_find
branch
from
September 25, 2025 17:04
56595e7 to
90bd55a
Compare
bernhardmgruber
marked this pull request as ready for review
September 25, 2025 17:04
3 tasks done
This comment has been minimized.
This comment has been minimized.
bernhardmgruber
force-pushed
the
ref_find
branch
from
September 26, 2025 09:11
90bd55a to
d6fdc19
Compare
This comment has been minimized.
This comment has been minimized.
bernhardmgruber
force-pushed
the
ref_find
branch
2 times, most recently
from
September 29, 2025 13:33
87c0680 to
dd654b8
Compare
bernhardmgruber
commented
Sep 29, 2025
This comment has been minimized.
This comment has been minimized.
bernhardmgruber
force-pushed
the
ref_find
branch
from
September 29, 2025 19:10
b79f5d3 to
02e4adb
Compare
This comment has been minimized.
This comment has been minimized.
bernhardmgruber
force-pushed
the
ref_find
branch
from
September 30, 2025 08:42
705995e to
7b9e329
Compare
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
bernhardmgruber
force-pushed
the
ref_find
branch
from
October 7, 2025 07:12
7b9e329 to
d285f7b
Compare
bernhardmgruber
commented
Oct 7, 2025
This comment has been minimized.
This comment has been minimized.
bernhardmgruber
force-pushed
the
ref_find
branch
from
October 14, 2025 09:50
55432b6 to
e41a347
Compare
Contributor
🥳 CI Workflow Results🟩 Finished in 4h 58m: Pass: 100%/118 | Total: 7d 17h | Max: 4h 57m | Hits: 33%/166765See results here. |
miscco
approved these changes
Oct 15, 2025
Naraenda
pushed a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 26, 2026
…hrust/cuda find() (NVIDIA/cccl#6019) Upstream: NVIDIA/cccl@1c3c7dd HIP: system/hip/detail/mismatch.h transform_pair_of_input_iterators_t got the iterator correctness fixes (defaulted default ctor, prefix operator++ returns self_t&, operator<); the find() zip/transform/counting_iterator rewrite is performance-only, so HIP find.h is unchanged. test/: no counterpart (no test change upstream beyond CUDA-only file). testing/: cuda/mismatch.cu unused-parameter cleanup taken. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
amd-jmahovsky
added a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 28, 2026
…ind/mismatch Upstream: NVIDIA/cccl@1c3c7ddee12 (NVIDIA/cccl#6019) Upstream: NVIDIA/cccl@8ca3dc3 (NVIDIA/cccl#6018) HIP previously took no change for these (performance-only). Bring system/hip closer to upstream: - find.h: find_if_n unwraps contiguous iterators and builds the zip/counting iterator with auto + variadic thrust::make_zip_iterator. transform_input_iterator_t now calls op(*input) directly instead of materializing the value_type (same semantics as ::cuda::transform_iterator; avoids NVIDIA/cccl#3591); prefix operator++ returns self_t&. The dedicated iterator is kept because thrust::transform_iterator still converts to the value_type and ::cuda::transform_iterator needs libhipcxx. - mismatch.h: drop transform_pair_of_input_iterators_t and enable the zip_iterator + zip_function implementation (thrust:: versions, so it also works without libhipcxx). Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
amd-jmahovsky
added a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 28, 2026
…tors in hip find/mismatch Related upstream commit: 8ca3dc31e90 Refactor thrust::mismatch to use CUDA iterators (NVIDIA/cccl#6018) Related upstream commit: 1c3c7ddee12 Refactor iterator usage in thrust/cuda find() (NVIDIA/cccl#6019) Related local port commit: e46dc36 e46dc36 brought the upstream find/mismatch rewrite to system/hip using thrust:: iterators, plus a local transform_input_iterator_t because thrust::transform_iterator materializes the value_type before calling the predicate (NVIDIA/cccl#3591). libhipcxx is now mandatory, so hip now uses the same libcu++ iterators as upstream's system/cuda: - find.h: find_if_n zips _THRUST_LIBCXX::make_transform_iterator( try_unwrap_contiguous_iterator(first), predicate) with _THRUST_LIBCXX::counting_iterator<Size>(0) and uses a _THRUST_STD::tuple<bool, Size> result type; __find_if:: transform_input_iterator_t and the dead !_THRUST_HAS_DEVICE_SYSTEM_STD include block are removed; upstream granular includes are taken (<thrust/functional.h> kept explicitly for the placeholders in find). - mismatch.h: _THRUST_LIBCXX::make_zip_iterator and _THRUST_LIBCXX::zip_function(binary_pred), with upstream's granular includes instead of thrust/{distance,pair,zip_function}.h and thrust/iterator/zip_iterator.h. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
Naraenda
pushed a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 29, 2026
…hrust/cuda find() (NVIDIA/cccl#6019) Upstream: NVIDIA/cccl@1c3c7dd HIP: system/hip/detail/mismatch.h transform_pair_of_input_iterators_t got the iterator correctness fixes (defaulted default ctor, prefix operator++ returns self_t&, operator<); the find() zip/transform/counting_iterator rewrite is performance-only, so HIP find.h is unchanged. test/: no counterpart (no test change upstream beyond CUDA-only file). testing/: cuda/mismatch.cu unused-parameter cleanup taken. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
Naraenda
pushed a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 29, 2026
…hrust/cuda find() (NVIDIA/cccl#6019) Upstream: NVIDIA/cccl@1c3c7dd HIP: system/hip/detail/mismatch.h transform_pair_of_input_iterators_t got the iterator correctness fixes (defaulted default ctor, prefix operator++ returns self_t&, operator<); the find() zip/transform/counting_iterator rewrite is performance-only, so HIP find.h is unchanged. test/: no counterpart (no test change upstream beyond CUDA-only file). testing/: cuda/mismatch.cu unused-parameter cleanup taken. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
Naraenda
pushed a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 29, 2026
…ind/mismatch Upstream: NVIDIA/cccl@1c3c7ddee12 (NVIDIA/cccl#6019) Upstream: NVIDIA/cccl@8ca3dc3 (NVIDIA/cccl#6018) HIP previously took no change for these (performance-only). Bring system/hip closer to upstream: - find.h: find_if_n unwraps contiguous iterators and builds the zip/counting iterator with auto + variadic thrust::make_zip_iterator. transform_input_iterator_t now calls op(*input) directly instead of materializing the value_type (same semantics as ::cuda::transform_iterator; avoids NVIDIA/cccl#3591); prefix operator++ returns self_t&. The dedicated iterator is kept because thrust::transform_iterator still converts to the value_type and ::cuda::transform_iterator needs libhipcxx. - mismatch.h: drop transform_pair_of_input_iterators_t and enable the zip_iterator + zip_function implementation (thrust:: versions, so it also works without libhipcxx). Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
Naraenda
pushed a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 29, 2026
…tors in hip find/mismatch Related upstream commit: 8ca3dc31e90 Refactor thrust::mismatch to use CUDA iterators (NVIDIA/cccl#6018) Related upstream commit: 1c3c7ddee12 Refactor iterator usage in thrust/cuda find() (NVIDIA/cccl#6019) Related local port commit: e46dc36 e46dc36 brought the upstream find/mismatch rewrite to system/hip using thrust:: iterators, plus a local transform_input_iterator_t because thrust::transform_iterator materializes the value_type before calling the predicate (NVIDIA/cccl#3591). libhipcxx is now mandatory, so hip now uses the same libcu++ iterators as upstream's system/cuda: - find.h: find_if_n zips _THRUST_LIBCXX::make_transform_iterator( try_unwrap_contiguous_iterator(first), predicate) with _THRUST_LIBCXX::counting_iterator<Size>(0) and uses a _THRUST_STD::tuple<bool, Size> result type; __find_if:: transform_input_iterator_t and the dead !_THRUST_HAS_DEVICE_SYSTEM_STD include block are removed; upstream granular includes are taken (<thrust/functional.h> kept explicitly for the placeholders in find). - mismatch.h: _THRUST_LIBCXX::make_zip_iterator and _THRUST_LIBCXX::zip_function(binary_pred), with upstream's granular includes instead of thrust/{distance,pair,zip_function}.h and thrust/iterator/zip_iterator.h. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
Naraenda
pushed a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 30, 2026
…hrust/cuda find() (NVIDIA/cccl#6019) Upstream: NVIDIA/cccl@1c3c7dd HIP: system/hip/detail/mismatch.h transform_pair_of_input_iterators_t got the iterator correctness fixes (defaulted default ctor, prefix operator++ returns self_t&, operator<); the find() zip/transform/counting_iterator rewrite is performance-only, so HIP find.h is unchanged. test/: no counterpart (no test change upstream beyond CUDA-only file). testing/: cuda/mismatch.cu unused-parameter cleanup taken. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
Beanavil
pushed a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 30, 2026
…hrust/cuda find() (NVIDIA/cccl#6019) Upstream: NVIDIA/cccl@1c3c7dd HIP: system/hip/detail/mismatch.h transform_pair_of_input_iterators_t got the iterator correctness fixes (defaulted default ctor, prefix operator++ returns self_t&, operator<); the find() zip/transform/counting_iterator rewrite is performance-only, so HIP find.h is unchanged. test/: no counterpart (no test change upstream beyond CUDA-only file). testing/: cuda/mismatch.cu unused-parameter cleanup taken. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
amd-jmahovsky
added a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 30, 2026
…ind/mismatch Upstream: NVIDIA/cccl@1c3c7ddee12 (NVIDIA/cccl#6019) Upstream: NVIDIA/cccl@8ca3dc3 (NVIDIA/cccl#6018) HIP previously took no change for these (performance-only). Bring system/hip closer to upstream: - find.h: find_if_n unwraps contiguous iterators and builds the zip/counting iterator with auto + variadic thrust::make_zip_iterator. transform_input_iterator_t now calls op(*input) directly instead of materializing the value_type (same semantics as ::cuda::transform_iterator; avoids NVIDIA/cccl#3591); prefix operator++ returns self_t&. The dedicated iterator is kept because thrust::transform_iterator still converts to the value_type and ::cuda::transform_iterator needs libhipcxx. - mismatch.h: drop transform_pair_of_input_iterators_t and enable the zip_iterator + zip_function implementation (thrust:: versions, so it also works without libhipcxx). Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
amd-jmahovsky
added a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 30, 2026
…tors in hip find/mismatch Related upstream commit: 8ca3dc31e90 Refactor thrust::mismatch to use CUDA iterators (NVIDIA/cccl#6018) Related upstream commit: 1c3c7ddee12 Refactor iterator usage in thrust/cuda find() (NVIDIA/cccl#6019) Related local port commit: e46dc36 e46dc36 brought the upstream find/mismatch rewrite to system/hip using thrust:: iterators, plus a local transform_input_iterator_t because thrust::transform_iterator materializes the value_type before calling the predicate (NVIDIA/cccl#3591). libhipcxx is now mandatory, so hip now uses the same libcu++ iterators as upstream's system/cuda: - find.h: find_if_n zips _THRUST_LIBCXX::make_transform_iterator( try_unwrap_contiguous_iterator(first), predicate) with _THRUST_LIBCXX::counting_iterator<Size>(0) and uses a _THRUST_STD::tuple<bool, Size> result type; __find_if:: transform_input_iterator_t and the dead !_THRUST_HAS_DEVICE_SYSTEM_STD include block are removed; upstream granular includes are taken (<thrust/functional.h> kept explicitly for the placeholders in find). - mismatch.h: _THRUST_LIBCXX::make_zip_iterator and _THRUST_LIBCXX::zip_function(binary_pred), with upstream's granular includes instead of thrust/{distance,pair,zip_function}.h and thrust/iterator/zip_iterator.h. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Merge first:
zip_iterator#6036