Skip to content

Ensure non-trivial thrust::copy from D2D uses UBLKCP - #5954

Merged
bernhardmgruber merged 2 commits into
NVIDIA:mainfrom
bernhardmgruber:copy_BULK
Sep 23, 2025
Merged

bernhardmgruber merged 2 commits into
NVIDIA:mainfrom
bernhardmgruber:copy_BULK

Conversation

@bernhardmgruber

@bernhardmgruber bernhardmgruber commented Sep 19, 2025 •

Copy link
Copy Markdown
Contributor

Fixes: #5955

@bernhardmgruber
bernhardmgruber requested a review from a team as a code owner September 19, 2025 21:37
@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
@github-actions

This comment has been minimized.

@github-actions

This comment has been minimized.

@github-actions

Copy link
Copy Markdown
Contributor

🥳 CI Workflow Results

🟩 Finished in 3h 31m: Pass: 100%/185 | Total: 6d 10h | Max: 3h 29m | Hits: 75%/186504

See results here.

@bernhardmgruber
bernhardmgruber merged commit 36ec00b into NVIDIA:main Sep 23, 2025
1 check passed
@github-project-automation github-project-automation Bot moved this from In Review to Done in CCCL Sep 23, 2025
@bernhardmgruber
bernhardmgruber deleted the copy_BULK branch September 24, 2025 12:46
pciolkosz pushed a commit to pciolkosz/cccl that referenced this pull request Sep 24, 2025
Naraenda pushed a commit to ROCm/rocm-libraries that referenced this pull request Oct 26, 2026
…opy from D2D uses UBLKCP (NVIDIA/cccl#5954)

Upstream: NVIDIA/cccl@36ec00b

  - cuda `copy.h`: non-trivial device-to-device copy wraps `::cuda::std::identity` in `::cuda::proclaim_copyable_arguments` so DeviceTransform can use UBLKCP (merged cleanly). cub `kernels/transform.cuh` belongs to hipCUB.
  - HIP counterpart: no change: performance-only.
  - test/ counterpart: none; testing/ counterpart: none.
  - 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
…opy from D2D uses UBLKCP (NVIDIA/cccl#5954)

Upstream: NVIDIA/cccl@36ec00b

  - cuda `copy.h`: non-trivial device-to-device copy wraps `::cuda::std::identity` in `::cuda::proclaim_copyable_arguments` so DeviceTransform can use UBLKCP (merged cleanly). cub `kernels/transform.cuh` belongs to hipCUB.
  - HIP counterpart: no change: performance-only.
  - test/ counterpart: none; testing/ counterpart: none.
  - 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
…opy from D2D uses UBLKCP (NVIDIA/cccl#5954)

Upstream: NVIDIA/cccl@36ec00b

  - cuda `copy.h`: non-trivial device-to-device copy wraps `::cuda::std::identity` in `::cuda::proclaim_copyable_arguments` so DeviceTransform can use UBLKCP (merged cleanly). cub `kernels/transform.cuh` belongs to hipCUB.
  - HIP counterpart: no change: performance-only.
  - test/ counterpart: none; testing/ counterpart: none.
  - 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
…opy from D2D uses UBLKCP (NVIDIA/cccl#5954)

Upstream: NVIDIA/cccl@36ec00b

  - cuda `copy.h`: non-trivial device-to-device copy wraps `::cuda::std::identity` in `::cuda::proclaim_copyable_arguments` so DeviceTransform can use UBLKCP (merged cleanly). cub `kernels/transform.cuh` belongs to hipCUB.
  - HIP counterpart: no change: performance-only.
  - test/ counterpart: none; testing/ counterpart: none.
  - 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
…opy from D2D uses UBLKCP (NVIDIA/cccl#5954)

Upstream: NVIDIA/cccl@36ec00b

  - cuda `copy.h`: non-trivial device-to-device copy wraps `::cuda::std::identity` in `::cuda::proclaim_copyable_arguments` so DeviceTransform can use UBLKCP (merged cleanly). cub `kernels/transform.cuh` belongs to hipCUB.
  - HIP counterpart: no change: performance-only.
  - test/ counterpart: none; testing/ counterpart: none.
  - 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
…roclaims copyable arguments

Related upstream commit: 36ec00be03f Ensure non-trivial thrust::copy from D2D uses UBLKCP (NVIDIA/cccl#5954)
Upstream: NVIDIA/cccl@36ec00b
Related local port commit: 31847ac
Related local parity commit: 713e047

713e047 left the proclaim_copyable_arguments wrapper out of the HIP
backend. Add it, as upstream does for CUDA: the non-trivial branch of
__copy::device_to_device in thrust/system/hip/detail/copy.h now passes
_THRUST_LIBCXX::proclaim_copyable_arguments(_THRUST_STD::identity{}) to
hip_rocprim::transform, so the transform may copy its arguments instead of
reading them in place.

Upstream needs no new include for this. Here copy.h and transform.h include
each other, so copy.h includes __functional/address_stability.h itself.

Not compiled, tests not run

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

Labels

None yet

Projects

Archived in project

Development

Successfully merging this pull request may close these issues.

Ensure thrust::copy uses the UBLKCP path of cub::DeviceTransform

3 participants