Repository navigation
Extend and refactor transform overloads in CUDA system - #5238
Merged
bernhardmgruber merged 6 commits intoJul 17, 2025
Merged
Conversation
3 tasks done
miscco
reviewed
Jul 14, 2025
bernhardmgruber
force-pushed
the
transform_cuda_overloads
branch
from
July 14, 2025 12:30
529b2b1 to
e347898
Compare
Contributor
🟨 CI finished in 6h 40m: Pass: 28%/139 | Total: 1d 08h | Avg: 14m 01s | Max: 1h 39m | Hits: 95%/16741
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| CCCL Packaging | |
| libcu++ | |
| CUB | |
| +/- | Thrust |
| CUDA Experimental | |
| stdpar | |
| python | |
| CCCL C Parallel Library | |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | CCCL Packaging |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| +/- | CUDA Experimental |
| +/- | stdpar |
| python | |
| +/- | CCCL C Parallel Library |
| +/- | Catch2Helper |
🏃 Runner counts (total jobs: 139)
| # | Runner |
|---|---|
| 91 | linux-amd64-cpu16 |
| 17 | windows-amd64-cpu16 |
| 10 | linux-arm64-cpu16 |
| 7 | linux-amd64-gpu-rtx2080-latest-1 |
| 6 | linux-amd64-gpu-rtxa6000-latest-1 |
| 5 | linux-amd64-gpu-h100-latest-1 |
| 3 | linux-amd64-gpu-rtx4090-latest-1 |
bernhardmgruber
force-pushed
the
transform_cuda_overloads
branch
from
July 17, 2025 10:29
e347898 to
9c3bcc3
Compare
bernhardmgruber
commented
Jul 17, 2025
bernhardmgruber
commented
Jul 17, 2025
miscco
reviewed
Jul 17, 2025
| */ | ||
| template <typename InputIterator, typename ForwardIterator, typename UnaryFunction, typename Predicate> | ||
| ForwardIterator | ||
| _CCCL_HOST_DEVICE ForwardIterator |
Contributor
There was a problem hiding this comment.
Should those and the other rather use _CCCL_API ?
Contributor
Author
There was a problem hiding this comment.
I frankly don't know. Looking into a few files, it looks like the Thrust public API has either _CCCL_HOST_DEVICE or no macro.
It also seems the overloads without execution policy have no macro. I guess I have to update this PR then.
miscco
approved these changes
Jul 17, 2025
bernhardmgruber
force-pushed
the
transform_cuda_overloads
branch
from
July 17, 2025 11:32
0a76a3f to
b43fbb4
Compare
Contributor
🟩 CI finished in 2h 09m: Pass: 100%/139 | Total: 3d 13h | Avg: 36m 54s | Max: 1h 45m | Hits: 78%/173788
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| CCCL Packaging | |
| libcu++ | |
| CUB | |
| +/- | Thrust |
| CUDA Experimental | |
| stdpar | |
| python | |
| CCCL C Parallel Library | |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | CCCL Packaging |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| +/- | CUDA Experimental |
| +/- | stdpar |
| python | |
| +/- | CCCL C Parallel Library |
| +/- | Catch2Helper |
🏃 Runner counts (total jobs: 139)
| # | Runner |
|---|---|
| 91 | linux-amd64-cpu16 |
| 17 | windows-amd64-cpu16 |
| 10 | linux-arm64-cpu16 |
| 7 | linux-amd64-gpu-rtx2080-latest-1 |
| 6 | linux-amd64-gpu-rtxa6000-latest-1 |
| 5 | linux-amd64-gpu-h100-latest-1 |
| 3 | linux-amd64-gpu-rtx4090-latest-1 |
davebayer
pushed a commit
to davebayer/cccl
that referenced
this pull request
Sep 23, 2025
Naraenda
pushed a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 26, 2026
…m overloads in CUDA system (NVIDIA/cccl#5238) Upstream: NVIDIA/cccl@76fc457 - clean port — applied to projects/hipccl/hipccl3/rocthrust/thrust/system/cuda/detail/transform.h (now identical to upstream; the local `_THRUST_STD::distance` spellings were all on lines upstream rewrote). - HIP counterpart: no change — `thrust::transform_n`/`transform_if_n` already reach the HIP backend through the generic implementations in `system/detail/generic/transform.h`; the new CUDA overloads only shortcut onto the CUDA fast path (performance-only, per sync decision). - test/ counterpart: none (no testing/*.cu paths touched). - 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
…sform over counting_iterator Related upstream commit: ad70011e990 Support fancy iterators in vectorized transform and port thrust::tabulate to it (NVIDIA/cccl#6012) Related local port commit: 291f437 The local port took upstream's cuda tabulate.h but left system/hip on its own parallel_for functor as "no change: performance-only". Bring system/hip in line with upstream's structure: - Drop __tabulate::functor; tabulate is a transform of _THRUST_LIBCXX::counting_iterator<iter_difference_t<Iterator>> into the output range, with granular libhipcxx includes as upstream. - hip has no transform_n overload (NVIDIA/cccl#5238 was cuda-only), so hip_rocprim::transform is called with [begin, begin + count). - tabulate_op is not wrapped in _THRUST_LIBCXX::proclaim_copyable_arguments: the proclamation only feeds cub::DeviceTransform, which hip's transform does not use, and the wrapper inherits from the functor, so it would reject function pointers and final classes for no benefit. 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
…m overloads in CUDA system (NVIDIA/cccl#5238) Upstream: NVIDIA/cccl@76fc457 - clean port — applied to projects/hipccl/hipccl3/rocthrust/thrust/system/cuda/detail/transform.h (now identical to upstream; the local `_THRUST_STD::distance` spellings were all on lines upstream rewrote). - HIP counterpart: no change — `thrust::transform_n`/`transform_if_n` already reach the HIP backend through the generic implementations in `system/detail/generic/transform.h`; the new CUDA overloads only shortcut onto the CUDA fast path (performance-only, per sync decision). - test/ counterpart: none (no testing/*.cu paths touched). - 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
…m overloads in CUDA system (NVIDIA/cccl#5238) Upstream: NVIDIA/cccl@76fc457 - clean port — applied to projects/hipccl/hipccl3/rocthrust/thrust/system/cuda/detail/transform.h (now identical to upstream; the local `_THRUST_STD::distance` spellings were all on lines upstream rewrote). - HIP counterpart: no change — `thrust::transform_n`/`transform_if_n` already reach the HIP backend through the generic implementations in `system/detail/generic/transform.h`; the new CUDA overloads only shortcut onto the CUDA fast path (performance-only, per sync decision). - test/ counterpart: none (no testing/*.cu paths touched). - 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
…sform over counting_iterator Related upstream commit: ad70011e990 Support fancy iterators in vectorized transform and port thrust::tabulate to it (NVIDIA/cccl#6012) Related local port commit: 291f437 The local port took upstream's cuda tabulate.h but left system/hip on its own parallel_for functor as "no change: performance-only". Bring system/hip in line with upstream's structure: - Drop __tabulate::functor; tabulate is a transform of _THRUST_LIBCXX::counting_iterator<iter_difference_t<Iterator>> into the output range, with granular libhipcxx includes as upstream. - hip has no transform_n overload (NVIDIA/cccl#5238 was cuda-only), so hip_rocprim::transform is called with [begin, begin + count). - tabulate_op is not wrapped in _THRUST_LIBCXX::proclaim_copyable_arguments: the proclamation only feeds cub::DeviceTransform, which hip's transform does not use, and the wrapper inherits from the functor, so it would reject function pointers and final classes for no benefit. 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
…m overloads in CUDA system (NVIDIA/cccl#5238) Upstream: NVIDIA/cccl@76fc457 - clean port — applied to projects/hipccl/hipccl3/rocthrust/thrust/system/cuda/detail/transform.h (now identical to upstream; the local `_THRUST_STD::distance` spellings were all on lines upstream rewrote). - HIP counterpart: no change — `thrust::transform_n`/`transform_if_n` already reach the HIP backend through the generic implementations in `system/detail/generic/transform.h`; the new CUDA overloads only shortcut onto the CUDA fast path (performance-only, per sync decision). - test/ counterpart: none (no testing/*.cu paths touched). - 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
…m overloads in CUDA system (NVIDIA/cccl#5238) Upstream: NVIDIA/cccl@76fc457 - clean port — applied to projects/hipccl/hipccl3/rocthrust/thrust/system/cuda/detail/transform.h (now identical to upstream; the local `_THRUST_STD::distance` spellings were all on lines upstream rewrote). - HIP counterpart: no change — `thrust::transform_n`/`transform_if_n` already reach the HIP backend through the generic implementations in `system/detail/generic/transform.h`; the new CUDA overloads only shortcut onto the CUDA fast path (performance-only, per sync decision). - test/ counterpart: none (no testing/*.cu paths touched). - 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
…sform over counting_iterator Related upstream commit: ad70011e990 Support fancy iterators in vectorized transform and port thrust::tabulate to it (NVIDIA/cccl#6012) Related local port commit: 291f437 The local port took upstream's cuda tabulate.h but left system/hip on its own parallel_for functor as "no change: performance-only". Bring system/hip in line with upstream's structure: - Drop __tabulate::functor; tabulate is a transform of _THRUST_LIBCXX::counting_iterator<iter_difference_t<Iterator>> into the output range, with granular libhipcxx includes as upstream. - hip has no transform_n overload (NVIDIA/cccl#5238 was cuda-only), so hip_rocprim::transform is called with [begin, begin + count). - tabulate_op is not wrapped in _THRUST_LIBCXX::proclaim_copyable_arguments: the proclamation only feeds cub::DeviceTransform, which hip's transform does not use, and the wrapper inherits from the functor, so it would reject function pointers and final classes for no benefit. 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
…p transform overloads Related upstream commit: 76fc4572271 Extend and refactor transform overloads in CUDA system (NVIDIA/cccl#5238) Upstream: NVIDIA/cccl@76fc457 Related local port commit: 208d3c3 The local port only changed the CUDA backend; the HIP counterpart was left behind as performance-only. Apply the same overload set to hip/detail/transform.h: - add transform_n (unary and binary), transform_if_n (unary, unary with stencil, binary with stencil), taking _THRUST_STD::iter_difference_t<InputIt> counts as upstream does. HIP calls to thrust::transform_n/transform_if_n no longer fall back to the generic implementation. - transform/transform_if call __transform::unary/binary directly with _THRUST_STD::distance instead of bouncing through the stencil-less transform_if overload, and order the overloads like upstream. - construct the functors with braces, drop the trailing "// struct" comments and the section banners, and replace thrust/distance.h with the libhipcxx __iterator includes. The HIP backend keeps parallel_for instead of rocprim::transform, because thrust::transform allows the functor to modify its input. Not compiled, tests not run 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.
This PR mostly adds a few more
transform_ifandtransform[_if]_noverloads in the Thrust CUDA system, so we can put more custimized algorithms there in the (near) future. I also reordered the overloads more logically and switched the CDP fallback path to usecuda::std::transform.