feat(hipccl): rocThrust's parity with CCCL 3.2 by amd-jmahovsky · Pull Request #13963 · ROCm/rocm-libraries

amd-jmahovsky · 2026-10-19T21:35:00Z

Motivation

Add CCCL 3.2.x support to rocThrust library.

Each NVIDIA commit has a corresponding commit in this PR with an explanation from the AI of how it ported the commit.

See the files in the root rocThrust dir cccl-investigation-v3.2.1.md and todo-rocthrust.md. The investigation file is a summary of the overall changes and the TODO has detailed information about each commit.

Technical Details

This is an AI-generated but human-verified port of NVIDIA Thrust 3.2 to rocThrust, based on rocThrust 3.1.

This PR is stacked on top of:

feat(hipccl): rocThrust's parity with CCCL 3.1.x #13086

This PR required the following PRs to be merged:

fix(rocthrust): resolve test failures in rocthrust hipccl 3.2 #14467

Test Plan

TheRock and Math CI tests should pass.

Test Result

n/a

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.

Naraenda · 2026-10-20T16:21:15Z

I got it building locally with some extra required patches. There's also a possible hang on device reduce by key that I'm investigating. Not sure if it's related to this though.

amd-jmahovsky · 2026-10-20T16:44:45Z

I got it building locally with some extra required patches. There's also a possible hang on device reduce by key that I'm investigating. Not sure if it's related to this though.

Very nice! That's a good sign if it took less than a day to get the AI port running on its own.

Naraenda · 2026-10-21T14:51:19Z

Created a drafted PR (#14113) with WIP changes to make it build and pass tests.

The following tests still fail with a single case failure:

pair.hip
reduce_by_key.hip
shuffle.hip

Naraenda · 2026-10-23T14:43:40Z

Retargeted the branch so the merge stack is a bit more clearer.

Naraenda · 2026-10-26T14:09:55Z

Rebased

Naraenda · 2026-10-26T14:11:17Z

99% tests passed, 2 tests failed out of 320

Label Time Summary:
hip         = 295.59 sec*proc (163 tests)
upstream    =  87.23 sec*proc (157 tests)

Total Test time (real) = 382.90 sec

The following tests FAILED:
	240 - reduce_by_key.hip (Failed)                        hip
	272 - shuffle.hip (Failed)                              hip

We already have fixes for these pending internally.

therock-pr-bot · 2026-10-26T14:15:42Z

❌ PR Check — Action Required

Check	Status	Details
📝 PR Description	❌ Fail	Error: PR description must reference a JIRA ID, ISSUE ID, or a GitHub closing keyword. Expected: include a `JIRA ID` / `ISSUE ID` line (separator `:` or `-`, or omitted; value may be a JIRA key, a number with/without `#`, or a link), OR a closing keyword + issue reference. Accepted examples: • `JIRA ID : TESTAUTO-6039` • `JIRA ID - #330` • `JIRA ID #330` • `JIRA ID` (on separate line) `ROCM-25757` • `ISSUE ID : TESTUTO-3334` • `ISSUE ID #3334` • `ISSUE ID - TESTAUTO-3433` • `ISSUE ID` (on separate line) `AIRUNTIME-2352` • `ISSUE ID : https://github.com/<org_name>/<repo_name>/issues/1234` • `Closes #10` • `Fixes octo-org/octo-repo#100` • `Resolves: #123` • `#123` • `https://github.com/<org_name>/<repo_name>/issues/123` Current: no valid JIRA/ISSUE/closing-keyword reference found
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	⏳ Pending	⏳ Still running…
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—

⚠️ 1 policy check(s) failed. Please address the issues above before this PR can be Reviewed.

🚫 Please fix the failed policies

❌ PR Description

The Not ready to Review label was added to this PR. Once all policies pass, the label is removed automatically.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

therock-pr-bot · 2026-10-26T14:15:44Z

🚫 Please fix the failed policies before requesting reviews.

The following policy checks failed:

❌ PR Description

The Not ready to Review label has been added to this PR.
Once all policies pass, the label will be removed automatically.

…VIDIA/cccl#6557) Upstream: NVIDIA/cccl@2845628 HIP counterpart (system/hip): none. N/A: gersemi formatting-only reflow of CMake files; local CMakeLists.txt/*.cmake are rocThrust-owned (HIP build) and the monorepo does not use gersemi, so no reformat ported; no functional change. test/ counterpart: none. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

… "sum columns" examples (NVIDIA/cccl#4462) Upstream: NVIDIA/cccl@3a4892e HIP counterpart (system/hip): none. Ported rewritten examples/sum_rows.cu and new examples/sum_columns.cu (registered in examples/CMakeLists.txt), adapted for non-libhipcxx compatibility: cuda::std::mdspan replaced by row-major raw pointer indexing, cuda::counting_iterator by thrust::counting_iterator<int>; added missing permutation_iterator include; AMD MIT header kept on sum_rows. CMake sidecars (sum_*.cmake, GCC7/nvcc-only), sum_columns filecheck and gersemi-only address_stability.cmake reflow not ported (local build has no per-example sidecar/FileCheck support). test/ counterpart: none. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…expr in CUB/Thrust (NVIDIA/cccl#6480) Upstream: NVIDIA/cccl@06f2025 HIP counterpart (system/hip): none; system/hip has no anonymous-enum constants of this form. Ported enum -> static constexpr in system/cuda (agent_launcher.h, core/util.h, reduce.h, reduce_by_key.h, set_operations.h); util.h conflict resolved to static constexpr bool value while keeping local ::cuda::std::true_type/false_type. test/ counterpart: none. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…e resilient against user-defined macros (NVIDIA/cccl#6564) Upstream: NVIDIA/cccl@8b26307 HIP counterpart (system/hip): mirrored in namespace.h HIP branch: ABI namespace now THRUST_PP_CAT(THRUST_PP_CAT(_V_, THRUST_VERSION), THRUST_PP_SPLICE_WITH(_, _R, ROCTHRUST_VERSION)) so user macros named THRUST/NS cannot break it (namespace name changes from THRUST_<ver>_<rocver>_NS to _V_<ver>_R_<rocver>; ABI-only, inline namespace). CUDA branches take upstream. cub/libcudacxx parts not applicable. test/ counterpart: none. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…r` with `cuda::std::pair` (NVIDIA/cccl#6616) Upstream: NVIDIA/cccl@91434d7 HIP counterpart (system/hip): none; system/hip keeps thrust::pair. Ported only system/cuda (extrema, merge, mismatch, partition, reduce_by_key, set_operations, temporary_buffer, unique_by_key) and CUDA-only testing/cuda tests, where thrust::pair is an alias of cuda::std::pair. Generic/public headers, sequential/omp/tbb backends, top-level tests, examples and benchmarks keep thrust::pair: in the non-libhipcxx fallback thrust::pair is a distinct struct (pair.h), so switching signatures to _THRUST_STD::pair (std::pair) would break the fallback and diverge from system/hip return types (compatibility over upstream parity). Remaining upstream changes in those files were formatting-only. test/ counterpart: none (test/ uses thrust::pair; unchanged). Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…lgorithm>` (NVIDIA/cccl#3741) Upstream: NVIDIA/cccl@fba9f83 HIP counterpart (system/hip): none. testing/counting_iterator.cu now includes _THRUST_STD_INCLUDE(algorithm) instead of the private __algorithm_ header; testing/zip_iterator_sort.cu qualifies thrust::stable_sort (avoids ADL ambiguity with now-public cuda::std algorithms). test/ counterpart: mirrored in test/test_counting_iterator.cpp (algorithm include) and test/test_zip_iterator_sort.cpp (thrust::stable_sort). Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…emporaries (NVIDIA/cccl#6617) Upstream: NVIDIA/cccl@b934754 HIP counterpart (system/hip): none; system/hip uses rocPRIM temp storage, not cub AliasTemporaries. Ported call-site updates in system/cuda (copy_if, core/util, extrema, partition, reduce_by_key, unique, unique_by_key). test/ counterpart: none. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…s for `thrust::seq` (NVIDIA/cccl#6415) Upstream: NVIDIA/cccl@99976dc HIP counterpart (system/hip): none. N/A: adds nvtx_policy.h and _CCCL_NVTX_RANGE_SCOPE_IF to algorithm entry points to skip NVTX for thrust::seq; rocThrust carries no NVTX ranges in generic entry points (kept ours, nvtx_policy.h not added). test/ counterpart: none. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

… to catch2_test_complex.cu (NVIDIA/cccl#6625) Upstream: NVIDIA/cccl@2f82c4d HIP counterpart (system/hip): none. N/A: upstream moves testing/complex.cu to Catch2 testing/catch2_test_complex.cu; rocThrust has no Catch2 harness, so the catch2 file is not added and legacy testing/complex.cu is kept (same coverage). test/ counterpart: test/test_complex.cpp already covers these cases; unchanged. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…l#6560) (NVIDIA/cccl#6704) Upstream: NVIDIA/cccl@26c79ca HIP counterpart (system/hip): none (rocPRIM scan unaffected). Ported OMP scan fix in system/omp/detail/scan.h: blocks reduce with their first element as init (no accum_t{} identity assumption), no-init inclusive path seeds from block_sums[0] and thread 0 scans without prefix, serial fallback when n < max(threshold, num_threads); _CCCL_ASSERT dropped (not available locally; replaced by a comment). testing/scan.cu and testing/transform_scan.cu gain the new tests (THRUST_HOST_DEVICE spelling; kept local include guards, added <numeric>). test/ counterpart: mirrored in test/test_scan.cpp (Int with friend ops, TestScanWithUserDefinedTypeAndInit replacing TestInclusiveScanWithUserDefinedType, TestScanBug6317 using get_random_data, TestScanEdgeCases) and test/test_transform_scan.cpp (TransformInt, TestTransformScanEdgeCases). Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

… about device lambdas that need to query the return type (NVIDIA/cccl#6765) (NVIDIA/cccl#6782) Upstream: NVIDIA/cccl@2f529ed HIP counterpart (system/hip): none. transform_iterator_reference now uses _THRUST_STD::invoke_result_t<UnaryFunc&, ...> (also for functional::actor) so the device-lambda return-type diagnostic fires; kept local include block (type_traits umbrella provides invoke_result_t; upstream's granular __functional/__type_traits includes not added). libcudacxx invoke.h part not applicable. test/ counterpart: none. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…DIA/cccl#6850) Upstream: NVIDIA/cccl@783ebe6 HIP counterpart (system/hip): none. N/A: upstream CCCL CMake backports (Thrust*.cmake, DetectSupportedStandards, ThrustRunTest, gersemi-formatted CMakeLists, async_reduce.cmake removal); rocThrust's CMake (top-level, benchmark, examples, testing) is independent of upstream's and these files are absent or AMD-owned locally. test/ counterpart: none. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…rors in thrust memory resource (NVIDIA/cccl#7002) (NVIDIA/cccl#7033) Upstream: NVIDIA/cccl@2c7c41e HIP counterpart (system/hip): mirrored in system/hip/memory_resource.h: do_deallocate no longer throws on hipFree errors (status discarded with [[maybe_unused]]), matching system/cuda/memory_resource.h. test/ counterpart: none. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…ivial copies (NVIDIA/cccl#7006) (NVIDIA/cccl#7134) Upstream: NVIDIA/cccl@7407ca1 HIP counterpart (system/hip): mirrored in system/hip/detail/util.h: trivial_copy_from_device/trivial_copy_to_device use hipMemcpyDefault (both hipMemcpyWithStream and legacy hipMemcpyAsync paths) so the runtime picks the direction when a device-accessible pointer is host memory. system/cuda/detail/util.h and CUDA-only testing/cuda/trivial_copy_memcpy_default.cu taken as upstream. test/ counterpart: added TestTrivialCopyFromDeviceHostSource/TestTrivialCopyToDeviceHostDest to test/test_copy.cpp. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…rrays in vec transform (NVIDIA/cccl#7162) (NVIDIA/cccl#7163) Upstream: NVIDIA/cccl@5281dca HIP counterpart (system/hip): none; the Alignment parameter only serves CUB's vectorized transform (cub part not applicable), and system/hip uses rocPRIM transform. Ported the uninitialized_array Alignment template parameter in system/cuda/detail/core/util.h. test/ counterpart: none. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…VIDIA/cccl#7172) (NVIDIA/cccl#7189) Upstream: NVIDIA/cccl@b85d41e HIP counterpart (system/hip): none (generic shuffle/random_bijection shared by all backends). Took upstream's new Feistel bijection (L/R split, size() = 1 << (L_bits+R_bits), min 256, nearest_power_of_two/mulhilo/get_cipher_bits/round_state removed) and generic/shuffle.inl now uses bijection.size(). Replaced cuda::std::bit_width with a portable bit-count loop (C++20 std::bit_width unavailable without libhipcxx); kept local umbrella includes instead of granular __type_traits/__bit includes. testing/shuffle.cu taken as upstream (Cephes and EvenSpacing/EvenDistribution tests removed, z-test key-position test, chi-squared critical value 145.461) with cuda::std:: -> _THRUST_STD:: and _THRUST_STD_INCLUDE(algorithm/cmath). test/ counterpart: mirrored in test/test_shuffle.cpp (size_t(256) bound, FeistelBijectionLength 345->512/256->256/1->256, inverse_erf z-test, critical value 145.461, Cephes/EvenSpacing/EvenDistribution removed, <cmath> added). Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…VIDIA/cccl#7329) Upstream: NVIDIA/cccl@64bf707 HIP counterpart (system/hip): none. N/A: upstream version bump only touches thrust/version.h / CCCL version files; rocThrust versioning (thrust/version.h, CHANGELOG.md) is rocThrust-owned and not modified by the sync. test/ counterpart: none. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…LIBCXX_NAMESPACE_{BEGIN,END}

Co-Authored-By: Claude <noreply@anthropic.com>

Beanavil · 2026-11-04T12:31:35Z

I think we still need to bump rocthrust's version

- #define THRUST_VERSION 300100 // macro expansion with ## requires this to be a single value
+ #define THRUST_VERSION 300201 // macro expansion with ## requires this to be a single value

Co-Authored-By: Claude <noreply@anthropic.com>

## Motivation Fixes two test found after #13963. ## Technical Details * a6eccea updates a test that was missed in the preceding PR. * 2736c6c resolves an issue that highlighted an issue where expected behavior in reduce by key is different between hipCUB and rocThrust. ## Test Plan Build and test with libhipcxx 3.4. ## Test Result Works on my machine. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. --------- Co-authored-by: Nick Breed <nick@streamhpc.com>

amd-jmahovsky · 2026-11-04T16:33:11Z

I think we still need to bump rocthrust's version

- #define THRUST_VERSION 300100 // macro expansion with ## requires this to be a single value
+ #define THRUST_VERSION 300201 // macro expansion with ## requires this to be a single value

Fixed

## Motivation Additional parity fixes on top of the original rocThrust CCCL 3.2 PR. ## Technical Details These fix omissions in the original branch/PR. ## Test Plan Code builds and passes unit tests. ## Test Result  ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. --------- Co-authored-by: Claude <noreply@anthropic.com> Co-authored-by: Wayne Franz <wayfranz@amd.com> Co-authored-by: Nara Prasetya <nara@streamhpc.com>

…py.h par_to_seq.h was folded into hip/detail/execution_policy.h by the CCCL 3.2 ADDITIONAL changes (#14670), while #14467 added an include of it to copy.h. copy.h already includes hip/detail/execution_policy.h, which provides cvt_to_seq. Co-Authored-By: Claude <noreply@anthropic.com>

amd-jmahovsky requested review from a team as code owners October 19, 2026 21:35

github-actions Bot added documentation project: rocthrust project: hipccl labels Oct 19, 2026

assistant-librarian Bot added the organization: ROCm label Oct 19, 2026

Naraenda changed the base branch from users/matyas-streamhpc/hipccl/rocthrust_cccl31 to users/naraenda/fix-rocthrust-cccl31 October 23, 2026 14:42

Naraenda added this pull request to stack #14284 October 23, 2026 14:42

NguyenNhuDi mentioned this pull request Oct 23, 2026

fix(rocthrust): fix rocthrust compile errors with libhipcxx 3.4 #14319

Closed

1 task

Naraenda force-pushed the users/naraenda/fix-rocthrust-cccl31 branch from a07b0d7 to 9472de8 Compare October 26, 2026 12:54

Naraenda force-pushed the users/amd-jmahovsky/hipccl/rocthrust_cccl32 branch from 80cd0f1 to b4d7c62 Compare October 26, 2026 14:09

Naraenda assigned Naraenda and amd-jmahovsky Oct 26, 2026

Naraenda removed this pull request from stack #14284 October 26, 2026 14:15

therock-pr-bot Bot added the Not ready to Review label Oct 26, 2026

Naraenda added this pull request to stack #14381 October 26, 2026 14:16

Naraenda mentioned this pull request Oct 27, 2026

fix(rocthrust): resolve test failures in rocthrust hipccl 3.2 #14467

Merged

1 task done

Beanavil mentioned this pull request Oct 28, 2026

feat(hipccl): hipCUB's parity with CCCL 3.2.x #14589

Open

1 task

stanleytsang-amd force-pushed the users/naraenda/fix-rocthrust-cccl31 branch from 03cbde5 to 5e8b1aa Compare October 28, 2026 19:23

Naraenda removed this pull request from stack #14381 October 29, 2026 08:52

Base automatically changed from users/naraenda/fix-rocthrust-cccl31 to users/matyas-streamhpc/hipccl/rocthrust_cccl31 October 29, 2026 09:04

Naraenda force-pushed the users/amd-jmahovsky/hipccl/rocthrust_cccl32 branch from 8834f3f to 7e110ef Compare October 29, 2026 13:18

amd-jmahovsky and others added 20 commits October 30, 2026 09:09

add upgrade md files

787f309

fix(rocthrust): capture by copy implicitely instead of explicitely

5fc5953

fix(rocthrust): fix incorrectly resolved merge conflict with _THRUST_…

7bd4dfc

…LIBCXX_NAMESPACE_{BEGIN,END}

Naraenda force-pushed the users/amd-jmahovsky/hipccl/rocthrust_cccl32 branch from 6263bbe to 7bd4dfc Compare October 30, 2026 09:41

chore(libhipcxx): bump submodule to 87af709ba43

4a8b41f

Co-Authored-By: Claude <noreply@anthropic.com>

amd-jmahovsky and others added 3 commits November 4, 2026 11:09

Update version.h

1e16ec9

chore(libhipcxx): bump submodule to 5cf4deb633f

494abab

Co-Authored-By: Claude <noreply@anthropic.com>

github-actions Bot added the project: rocprim label Nov 4, 2026

amd-jmahovsky and others added 2 commits November 4, 2026 09:59

Navigation Menu

Repository navigation

feat(hipccl): rocThrust's parity with CCCL 3.2 - #13963

feat(hipccl): rocThrust's parity with CCCL 3.2#13963
amd-jmahovsky wants to merge 142 commits into
users/matyas-streamhpc/hipccl/rocthrust_cccl31from
users/amd-jmahovsky/hipccl/rocthrust_cccl32

amd-jmahovsky commented Oct 19, 2026 •

edited by Naraenda

Loading

Uh oh!

Naraenda commented Oct 20, 2026

Uh oh!

amd-jmahovsky commented Oct 20, 2026

Uh oh!

Naraenda commented Oct 21, 2026 •

edited

Loading

Uh oh!

Naraenda commented Oct 23, 2026

Uh oh!

Naraenda commented Oct 26, 2026

Uh oh!

Naraenda commented Oct 26, 2026

Uh oh!

therock-pr-bot Bot commented Oct 26, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Oct 26, 2026

Uh oh!

Beanavil commented Nov 4, 2026

Uh oh!

amd-jmahovsky commented Nov 4, 2026

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

amd-jmahovsky commented Oct 19, 2026 • edited by Naraenda Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

Uh oh!

Naraenda commented Oct 20, 2026

Uh oh!

amd-jmahovsky commented Oct 20, 2026

Uh oh!

Naraenda commented Oct 21, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Naraenda commented Oct 23, 2026

Uh oh!

Naraenda commented Oct 26, 2026

Uh oh!

Naraenda commented Oct 26, 2026

Uh oh!

therock-pr-bot Bot commented Oct 26, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

❌ PR Check — Action Required

Uh oh!

therock-pr-bot Bot commented Oct 26, 2026

Uh oh!

Beanavil commented Nov 4, 2026

Uh oh!

amd-jmahovsky commented Nov 4, 2026

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

amd-jmahovsky commented Oct 19, 2026 •

edited by Naraenda

Loading

Naraenda commented Oct 21, 2026 •

edited

Loading

therock-pr-bot Bot commented Oct 26, 2026 •

edited

Loading