Skip to content

Improve thrust::cuda_cub::replace functor handling - #5949

Merged
bernhardmgruber merged 3 commits into
NVIDIA:mainfrom
bernhardmgruber:proclaim_replace
Sep 22, 2025
Merged

Improve thrust::cuda_cub::replace functor handling#5949
bernhardmgruber merged 3 commits into
NVIDIA:mainfrom
bernhardmgruber:proclaim_replace

Conversation

@bernhardmgruber

@bernhardmgruber bernhardmgruber commented Sep 19, 2025

Copy link
Copy Markdown
Contributor

Reuse functors which exist and proclaim they have copyable arguments when possible.

This is a utility PR in the context of #5179.

@bernhardmgruber
bernhardmgruber requested review from a team as code owners September 19, 2025 18:13
@github-project-automation github-project-automation Bot moved this to Todo in CCCL Sep 19, 2025
@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Sep 19, 2025
@bernhardmgruber bernhardmgruber changed the title Cleanup and improve thrust::cuda_cub::replace Refactor thrust::cuda_cub::replace functors Sep 19, 2025
@bernhardmgruber bernhardmgruber changed the title Refactor thrust::cuda_cub::replace functors Improve thrust::cuda_cub::replace functor handling Sep 19, 2025
fbusato
fbusato approved these changes Sep 19, 2025
View reviewed changes
Comment thread cub/cub/device/device_transform.cuh Outdated

THRUST_DEVICE_FUNCTION auto operator()() const -> T
template <typename... Args>
THRUST_DEVICE_FUNCTION auto operator()(Args&&...) const -> T

@fbusato fbusato Sep 19, 2025

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.

Confused by THRUST_DEVICE_FUNCTION in CUB

@bernhardmgruber bernhardmgruber Sep 19, 2025

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

yeah, that makes no sense! :D Fixed.

Comment thread cub/cub/device/device_transform.cuh

_CCCL_BEGIN_NAMESPACE_CUDA
template <typename T>
struct proclaims_copyable_arguments<CUB_NS_QUALIFIER::detail::__return_constant<T>> : ::cuda::std::true_type

@fbusato fbusato Sep 19, 2025

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.

can we use cub::detail:: instead?

@bernhardmgruber bernhardmgruber Sep 19, 2025

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

I don't think so. If the user uses a wrapped namespace for cub, then __return_constant would live in user_defined_ns::cub::detail. Since we are completely outside this namespace, lookup would not find it.

@bernhardmgruber
Cleanup and improve thrust::cuda_cub::replace
5aaae14 
Reuse functors which exist and proclaim they have copyable arguments when possible
@bernhardmgruber
attro
bf50f3c
@github-actions

This comment has been minimized.

Sign in to view
miscco
miscco approved these changes Sep 22, 2025
View reviewed changes
Comment thread thrust/thrust/detail/internal_functional.h Resolved
Comment thread thrust/thrust/detail/internal_functional.h Resolved
@bernhardmgruber
bernhardmgruber enabled auto-merge (squash) September 22, 2025 07:46
@github-actions

github-actions Bot commented Sep 22, 2025

Copy link
Copy Markdown
Contributor

🥳 CI Workflow Results

🟩 Finished in 3h 40m: Pass: 100%/185 | Total: 6d 09h | Max: 3h 31m | Hits: 61%/186335

See results here.

@bernhardmgruber
bernhardmgruber merged commit 1ceac37 into NVIDIA:main Sep 22, 2025
@github-project-automation github-project-automation Bot moved this from In Review to Done in CCCL Sep 22, 2025
@bernhardmgruber
bernhardmgruber deleted the proclaim_replace branch September 22, 2025 10:42
pciolkosz pushed a commit to pciolkosz/cccl that referenced this pull request Sep 24, 2025
@bernhardmgruber @miscco @pciolkosz
Improve thrust::cuda_cub::replace functor handling (NVIDIA#5949)
8085dce 
Reuse functors which exist and proclaim they have copyable arguments when possible

Co-authored-by: Michael Schellenberger Costa <miscco@nvidia.com>
Naraenda pushed a commit to ROCm/rocm-libraries that referenced this pull request Oct 26, 2026
@amd-jmahovsky @claude @Naraenda
feat(rocthrust): port CCCL 1ceac3731d5 - Improve `thrust::cuda_cub::r…
3adee1d 
…eplace` functor handling (NVIDIA/cccl#5949)

Upstream: NVIDIA/cccl@1ceac37

  - Generic `detail/internal_functional.h`: added `new_value_if_f` (`_CCCL_DEVICE_API` spelled `THRUST_DEVICE`) and `proclaims_copyable_arguments` specializations for `predicate_to_integral`, `tuple_binary_predicate`, `new_value_if_f`; specializations placed in `_THRUST_LIBCXX_NAMESPACE_BEGIN/END` or `namespace internal` (fallback), matching `zip_function.h`; address_stability pulled in via `thrust/detail/libcxx_wrapper/__functional/address_stability.h` instead of the cuda include.
  - cuda `replace.h` merged cleanly: uses `thrust::detail::new_value_if_f` with `cuda_cub::transform`, dropping its local functor.
  - HIP counterpart: no change: performance-only (copyable-argument proclamation only feeds CUDA DeviceTransform); `system/hip/detail/replace.h` keeps `__replace::new_value_if_f`. cub `device_transform.cuh` change belongs to hipCUB.
  - test/ counterpart: none; testing/ counterpart: none (no upstream test changes).
  - 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
@amd-jmahovsky @claude
refactor(rocthrust): CCCL 1ceac3731d5 parity - libcxx address_stabili…
93969ea 
…ty in internal_functional

Related upstream commit: 1ceac3731d5 Improve `thrust::cuda_cub::replace` functor handling (NVIDIA/cccl#5949)
Related local port commit: 3adee1d

libhipcxx is now mandatory, so internal_functional.h includes
_THRUST_LIBCXX_INCLUDE(__functional/address_stability.h) directly (upstream
<cuda/__functional/address_stability.h>) instead of the libcxx_wrapper header, and the
proclaims_copyable_arguments specializations for predicate_to_integral, tuple_binary_predicate and
new_value_if_f are only declared in the libcxx namespace (the `namespace internal` fallback is removed).

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
@amd-jmahovsky @claude
refactor(rocthrust): CCCL 1ceac3731d5 parity - hip replace uses share…
a5e6f00 
…d new_value_if_f

Related upstream commit: 1ceac3731d5 Improve `thrust::cuda_cub::replace` functor handling (NVIDIA/cccl#5949)
Related local port commit: 3d25701

The local port added the generic thrust::detail::new_value_if_f (with its
proclaims_copyable_arguments specialization) to detail/internal_functional.h
and moved system/cuda onto it, but left system/hip on its own copy as
"no change: performance-only". Bring system/hip in line:

- Drop hip's private __replace::new_value_if_f; replace_copy_if uses
  thrust::detail::new_value_if_f, as upstream's cuda backend does.
- replace() uses thrust::detail::equal_to_value<T> instead of the
  `_1 == old_value` placeholder expression, as upstream.
- __replace::constant_f is kept: upstream switched to CUB's
  detail::__return_constant, which has no rocThrust/libhipcxx equivalent.

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
@amd-jmahovsky @claude @Naraenda
feat(rocthrust): port CCCL 1ceac3731d5 - Improve `thrust::cuda_cub::r…
e752a98 
…eplace` functor handling (NVIDIA/cccl#5949)

Upstream: NVIDIA/cccl@1ceac37

  - Generic `detail/internal_functional.h`: added `new_value_if_f` (`_CCCL_DEVICE_API` spelled `THRUST_DEVICE`) and `proclaims_copyable_arguments` specializations for `predicate_to_integral`, `tuple_binary_predicate`, `new_value_if_f`; specializations placed in `_THRUST_LIBCXX_NAMESPACE_BEGIN/END` or `namespace internal` (fallback), matching `zip_function.h`; address_stability pulled in via `thrust/detail/libcxx_wrapper/__functional/address_stability.h` instead of the cuda include.
  - cuda `replace.h` merged cleanly: uses `thrust::detail::new_value_if_f` with `cuda_cub::transform`, dropping its local functor.
  - HIP counterpart: no change: performance-only (copyable-argument proclamation only feeds CUDA DeviceTransform); `system/hip/detail/replace.h` keeps `__replace::new_value_if_f`. cub `device_transform.cuh` change belongs to hipCUB.
  - test/ counterpart: none; testing/ counterpart: none (no upstream test changes).
  - 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
@amd-jmahovsky @claude @Naraenda
feat(rocthrust): port CCCL 1ceac3731d5 - Improve `thrust::cuda_cub::r…
4488ee3 
…eplace` functor handling (NVIDIA/cccl#5949)

Upstream: NVIDIA/cccl@1ceac37

  - Generic `detail/internal_functional.h`: added `new_value_if_f` (`_CCCL_DEVICE_API` spelled `THRUST_DEVICE`) and `proclaims_copyable_arguments` specializations for `predicate_to_integral`, `tuple_binary_predicate`, `new_value_if_f`; specializations placed in `_THRUST_LIBCXX_NAMESPACE_BEGIN/END` or `namespace internal` (fallback), matching `zip_function.h`; address_stability pulled in via `thrust/detail/libcxx_wrapper/__functional/address_stability.h` instead of the cuda include.
  - cuda `replace.h` merged cleanly: uses `thrust::detail::new_value_if_f` with `cuda_cub::transform`, dropping its local functor.
  - HIP counterpart: no change: performance-only (copyable-argument proclamation only feeds CUDA DeviceTransform); `system/hip/detail/replace.h` keeps `__replace::new_value_if_f`. cub `device_transform.cuh` change belongs to hipCUB.
  - test/ counterpart: none; testing/ counterpart: none (no upstream test changes).
  - 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
@amd-jmahovsky @claude @Naraenda
refactor(rocthrust): CCCL 1ceac3731d5 parity - libcxx address_stabili…
11eafe1 
…ty in internal_functional

Related upstream commit: 1ceac3731d5 Improve `thrust::cuda_cub::replace` functor handling (NVIDIA/cccl#5949)
Related local port commit: 3adee1d

libhipcxx is now mandatory, so internal_functional.h includes
_THRUST_LIBCXX_INCLUDE(__functional/address_stability.h) directly (upstream
<cuda/__functional/address_stability.h>) instead of the libcxx_wrapper header, and the
proclaims_copyable_arguments specializations for predicate_to_integral, tuple_binary_predicate and
new_value_if_f are only declared in the libcxx namespace (the `namespace internal` fallback is removed).

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
@amd-jmahovsky @claude @Naraenda
refactor(rocthrust): CCCL 1ceac3731d5 parity - hip replace uses share…
7ea19b1 
…d new_value_if_f

Related upstream commit: 1ceac3731d5 Improve `thrust::cuda_cub::replace` functor handling (NVIDIA/cccl#5949)
Related local port commit: 3d25701

The local port added the generic thrust::detail::new_value_if_f (with its
proclaims_copyable_arguments specialization) to detail/internal_functional.h
and moved system/cuda onto it, but left system/hip on its own copy as
"no change: performance-only". Bring system/hip in line:

- Drop hip's private __replace::new_value_if_f; replace_copy_if uses
  thrust::detail::new_value_if_f, as upstream's cuda backend does.
- replace() uses thrust::detail::equal_to_value<T> instead of the
  `_1 == old_value` placeholder expression, as upstream.
- __replace::constant_f is kept: upstream switched to CUB's
  detail::__return_constant, which has no rocThrust/libhipcxx equivalent.

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
@amd-jmahovsky @claude @Naraenda
feat(rocthrust): port CCCL 1ceac3731d5 - Improve `thrust::cuda_cub::r…
e58d058 
…eplace` functor handling (NVIDIA/cccl#5949)

Upstream: NVIDIA/cccl@1ceac37

  - Generic `detail/internal_functional.h`: added `new_value_if_f` (`_CCCL_DEVICE_API` spelled `THRUST_DEVICE`) and `proclaims_copyable_arguments` specializations for `predicate_to_integral`, `tuple_binary_predicate`, `new_value_if_f`; specializations placed in `_THRUST_LIBCXX_NAMESPACE_BEGIN/END` or `namespace internal` (fallback), matching `zip_function.h`; address_stability pulled in via `thrust/detail/libcxx_wrapper/__functional/address_stability.h` instead of the cuda include.
  - cuda `replace.h` merged cleanly: uses `thrust::detail::new_value_if_f` with `cuda_cub::transform`, dropping its local functor.
  - HIP counterpart: no change: performance-only (copyable-argument proclamation only feeds CUDA DeviceTransform); `system/hip/detail/replace.h` keeps `__replace::new_value_if_f`. cub `device_transform.cuh` change belongs to hipCUB.
  - test/ counterpart: none; testing/ counterpart: none (no upstream test changes).
  - 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
@amd-jmahovsky @claude @Beanavil
feat(rocthrust): port CCCL 1ceac3731d5 - Improve `thrust::cuda_cub::r…
636da1b 
…eplace` functor handling (NVIDIA/cccl#5949)

Upstream: NVIDIA/cccl@1ceac37

  - Generic `detail/internal_functional.h`: added `new_value_if_f` (`_CCCL_DEVICE_API` spelled `THRUST_DEVICE`) and `proclaims_copyable_arguments` specializations for `predicate_to_integral`, `tuple_binary_predicate`, `new_value_if_f`; specializations placed in `_THRUST_LIBCXX_NAMESPACE_BEGIN/END` or `namespace internal` (fallback), matching `zip_function.h`; address_stability pulled in via `thrust/detail/libcxx_wrapper/__functional/address_stability.h` instead of the cuda include.
  - cuda `replace.h` merged cleanly: uses `thrust::detail::new_value_if_f` with `cuda_cub::transform`, dropping its local functor.
  - HIP counterpart: no change: performance-only (copyable-argument proclamation only feeds CUDA DeviceTransform); `system/hip/detail/replace.h` keeps `__replace::new_value_if_f`. cub `device_transform.cuh` change belongs to hipCUB.
  - test/ counterpart: none; testing/ counterpart: none (no upstream test changes).
  - 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
@amd-jmahovsky @claude
refactor(rocthrust): CCCL 1ceac3731d5 parity - libcxx address_stabili…
abd50e5 
…ty in internal_functional

Related upstream commit: 1ceac3731d5 Improve `thrust::cuda_cub::replace` functor handling (NVIDIA/cccl#5949)
Related local port commit: 3adee1d

libhipcxx is now mandatory, so internal_functional.h includes
_THRUST_LIBCXX_INCLUDE(__functional/address_stability.h) directly (upstream
<cuda/__functional/address_stability.h>) instead of the libcxx_wrapper header, and the
proclaims_copyable_arguments specializations for predicate_to_integral, tuple_binary_predicate and
new_value_if_f are only declared in the libcxx namespace (the `namespace internal` fallback is removed).

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
@amd-jmahovsky @claude
refactor(rocthrust): CCCL 1ceac3731d5 parity - hip replace uses share…
0db5179 
…d new_value_if_f

Related upstream commit: 1ceac3731d5 Improve `thrust::cuda_cub::replace` functor handling (NVIDIA/cccl#5949)
Related local port commit: 3d25701

The local port added the generic thrust::detail::new_value_if_f (with its
proclaims_copyable_arguments specialization) to detail/internal_functional.h
and moved system/cuda onto it, but left system/hip on its own copy as
"no change: performance-only". Bring system/hip in line:

- Drop hip's private __replace::new_value_if_f; replace_copy_if uses
  thrust::detail::new_value_if_f, as upstream's cuda backend does.
- replace() uses thrust::detail::equal_to_value<T> instead of the
  `_1 == old_value` placeholder expression, as upstream.
- __replace::constant_f is kept: upstream switched to CUB's
  detail::__return_constant, which has no rocThrust/libhipcxx equivalent.

Not compiled, tests not run (per human)

Co-Authored-By: Claude <noreply@anthropic.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@miscco miscco miscco approved these changes
@fbusato fbusato fbusato approved these changes
@elstehle elstehle Awaiting requested review from elstehle elstehle is a code owner automatically assigned from NVIDIA/cccl-thrust-codeowners

Assignees

No one assigned

Labels

None yet

Projects

CCCL
Archived in project

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

3 participants

@bernhardmgruber @miscco @fbusato