feat(hipccl): hipCUB's parity with CCCL 3.1.x by Beanavil · Pull Request #13083 · ROCm/rocm-libraries

Beanavil · 2026-10-05T12:25:20Z

Motivation

Add CCCL 3.1.x support to hipCUB and rocPRIM libraries.

The branch in this PR is the same as in #11292 but with the changes ported to the new directory structure.

JIRA ID : EXSWSTRHPC-320

Technical Details

Branch is based on the staging branch, which includes all the CCCL 3.0.x updates.

Most relevant changes:

Added support for large number of segments and large number of items to DeviceSegmentedRadixSort
Added WarpReduce overloadings
Added fixed size segmented reduce
Removed deprecated thread operators
Added support for large num_segments to DeviceSegmentedReduce with fixed segment size
Added support for large number of items to DeviceRunLengthEncode::Encode
Added env-based API for DeviceReduce
Bumped CCCL version to 3.1.4
Promoted libhipcxx to hard-dependency, also on Windows

Test Plan

Internal+AMD's CI runs

Test Result

Internal CI passes. AMD CI should be checked now

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

therock-pr-bot · 2026-10-09T22:20:53Z

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 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-09T22:35:56Z

🎉 All checks passed! This PR is ready for review.

NguyenNhuDi

I asked Claude to take a look at this PR and come up with potential issues. Here are some issues that Claude came up with that I think are relevant.

NguyenNhuDi · 2026-10-30T23:00:23Z

+        }
+    }
+}
+


Is it worth testing the two params overloads as well?

what params overloads do you refer to?

The guys with two params:
https://github.com/ROCm/rocm-libraries/pull/13083/changes#diff-ca4cd744224dcf3eca67fd28399df00c703d7d08c02e488479b347d53ec8cb5dR113

… device segmented radix sort

…mented_reduce

…ge number of items in device segmented radix sort

This commit works around one issue and fixes another. 1. The workaround: cmake probes the compiler to see what options it passes by default. One of them is a -I option that includes the path to the libhipcxx headers. For this reason, it removes that option in the compile commands it generates from our cmake files. However, it appears that the compiler is operating in a different mode/state when building our source than it is during the probe. This means that the -I option to include the path to the libhipcxx header is not passed at compile time. We can work around this with a temporary hack in the root level hipcub CMakeLists.txt that filters out any detected implicit include directories for libhipcxx. 2. libhipcxx 3.4 removed several macros hipCUB depended on. Modify the #defines so that everything compiler both with the current version as well as with older libhipcxx versions.

The workaround is no longer necessary.

Point projects/hipccl/hipccl3/libhipcxx at ROCm/libhipcxx amd-develop (87af709ba4349548a2bfaee73907635d889736dd), which is libhipcxx 3.4.0, required for CCCL 3.1.x/3.2.x compatibility. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>

…ead of _CCCL_* variants" This reverts commit 587574e.

This reverts commit 13b1344.

codecov · 2026-11-03T02:22:43Z

Codecov Report

❌ Patch coverage is 85.46012% with 237 lines in your changes missing coverage. Please review.

Files with missing lines	Patch %	Lines
...3/hipcub/test/hipcub/test_hipcub_device_reduce.cpp	67.21%	14 Missing and 26 partials ⚠️
...cl3/hipcub/test/hipcub/test_hipcub_warp_reduce.cpp	71.22%	20 Missing and 20 partials ⚠️
...hipcub/test_hipcub_device_segmented_radix_sort.hpp	82.87%	10 Missing and 21 partials ⚠️
...cl3/hipcub/test/hipcub/test_hipcub_device_scan.cpp	86.63%	4 Missing and 21 partials ⚠️
...st/hipcub/test_hipcub_device_run_length_encode.cpp	85.82%	3 Missing and 17 partials ⚠️
...est/hipcub/test_hipcub_device_segmented_reduce.cpp	95.96%	2 Missing and 16 partials ⚠️
...include/hipcub/backend/rocprim/detail/util_env.hpp	65.91%	11 Missing and 4 partials ⚠️
...backend/rocprim/device/device_segmented_reduce.hpp	93.60%	9 Missing and 4 partials ⚠️
...cl/hipccl3/hipcub/hipcub/include/hipcub/config.hpp	66.67%	4 Missing and 6 partials ⚠️
.../hipccl3/hipcub/test/hipcub/common_test_header.hpp	33.33%	8 Missing and 2 partials ⚠️
... and 2 more

Additional details and impacted files

@@             Coverage Diff              @@
##           develop   #13083       +/-   ##
============================================
+ Coverage    77.77%   80.83%    +3.06%     
============================================
  Files         2905      103     -2802     
  Lines       491621    26114   -465507     
  Branches     74545     5400    -69145     
============================================
- Hits        382355    21109   -361246     
+ Misses       84274     1627    -82647     
+ Partials     24992     3378    -21614

Flag	Coverage Δ
TensileLite-CPP	`?`
TensileLite-Unit	`?`
hipBLAS	`?`
hipBLASLt	`?`
hipCUB	`80.83% <85.46%> (-1.85%)`	⬇️
hipDNN	`?`
hipFFT	`?`
hipRAND	`?`
hipSOLVER	`?`
hipSPARSE	`?`
rocBLAS	`?`
rocFFT	`?`
rocRAND	`?`
rocSOLVER	`?`
rocSPARSE	`?`
rocThrust	`?`

Files with missing lines	Coverage Δ
...e/hipcub/backend/rocprim/block/block_histogram.hpp	`0.00% <ø> (ø)`
...nclude/hipcub/backend/rocprim/block/block_load.hpp	`0.00% <ø> (ø)`
.../hipcub/backend/rocprim/block/block_radix_sort.hpp	`0.00% <ø> (ø)`
...lude/hipcub/backend/rocprim/block/block_reduce.hpp	`0.00% <ø> (ø)`
...nclude/hipcub/backend/rocprim/block/block_scan.hpp	`0.00% <ø> (ø)`
...clude/hipcub/backend/rocprim/block/block_store.hpp	`0.00% <ø> (ø)`
...clude/hipcub/backend/rocprim/device/device_for.hpp	`100.00% <100.00%> (ø)`
...de/hipcub/backend/rocprim/device/device_reduce.hpp	`100.00% <100.00%> (ø)`
...ackend/rocprim/device/device_run_length_encode.hpp	`100.00% <ø> (ø)`
...end/rocprim/device/device_segmented_radix_sort.hpp	`100.00% <ø> (ø)`
... and 26 more

... and 2972 files with indirect coverage changes

🚀 New features to boost your workflow:

❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.
📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

…ity-upstream

stanleytsang-amd · 2026-11-04T15:32:21Z

@Beanavil Can you resolve the merge conflicts here please?

…ity-upstream

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

github-actions Bot added the documentation label Oct 5, 2026

Beanavil added project: hipcub project: rocprim project: rocthrust labels Oct 5, 2026

Beanavil changed the title ~~feat(hipcub,rocprim): Add CCCL 3.1.x support~~ feat(hipccl): hipCUB's parity with CCCL 3.1.x Oct 5, 2026

Naraenda force-pushed the users/naraenda/hipccl-3.x-staging branch from 2ec1db0 to c5ccc3d Compare October 5, 2026 14:31

Naraenda mentioned this pull request Oct 5, 2026

feat(hipccl): cccl 3.x parity #12952

Merged

1 task done

github-actions Bot added the project: hipccl label Oct 7, 2026

Beanavil force-pushed the users/beanavil/hipccl/hipcub-cccl-3-1-parity-upstream branch from 407c119 to a93bbfb Compare October 7, 2026 10:30

Base automatically changed from users/naraenda/hipccl-3.x-staging to develop October 9, 2026 22:16

Beanavil force-pushed the users/beanavil/hipccl/hipcub-cccl-3-1-parity-upstream branch from a93bbfb to cc9176f Compare October 12, 2026 09:05

github-actions Bot added the project: hipcub-deprecated label Oct 12, 2026

Beanavil force-pushed the users/beanavil/hipccl/hipcub-cccl-3-1-parity-upstream branch from cc9176f to 1690f2d Compare October 12, 2026 09:15

Beanavil marked this pull request as ready for review October 12, 2026 09:18

Beanavil requested review from a team as code owners October 12, 2026 09:18

Beanavil mentioned this pull request Oct 12, 2026

feat(hipcub,rocprim): Add CCCL 3.1.x support #11292

Closed

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

Beanavil force-pushed the users/beanavil/hipccl/hipcub-cccl-3-1-parity-upstream branch from 60bf9a3 to 5f5e253 Compare October 30, 2026 10:29

NguyenNhuDi requested changes Oct 30, 2026

View reviewed changes

Beanavil added 7 commits November 2, 2026 17:00

Add support for large number of segments and large number of items to…

550631c

… device segmented radix sort

Improve test times for device segmented radix sort

92d331f

Avoid instantiating discard_iterator while parsing

103679e

Remove _CCCL_INLINE_VAR

2fc15e0

Use [[maybe_unused]] for unused arguments

4c5a493

Remove HIPCUB_ARCH from template parameters

815addb

Add new WarpReduce overloadings

2b50416

NB4444 and others added 24 commits November 2, 2026 17:01

fix(hipcub): fix signature in benchmark for segmented_radix_sort

6ce5606

fix(hipcub): fix build issues test_hipcub_device_segmented_reduce

0dd0a2e

fix(hipcub): fix wrong init for the test

745c6c7

fix(hipcub): removed path for trigering static_assert

e5f3af1

fix(hipcub): add missing dependencies

e151d7f

fix(hipcub): build error on nvcc for segment

303a709

fix(rocprim,hipcub): fix support for large num_segments in device_seg…

56da7f0

…mented_reduce

fix(rocprim,hipcub): fix support for large number of segments and lar…

3bc6949

…ge number of items in device segmented radix sort

fix(hipcub): remove deprecated thread operators

90d40fe

fix(hipcub): patch segmented_reduce test for NVIDIA backend

554a2ca

fix(hipcub): fix host compiler complain when using vector.assign()

24762c4

Error if libhipcxx or libcudacxx is not found

123d21f

Promote libhipcxx to hard dependency on all platforms

da2160f

Add device transform test for large sizes

f89003a

Env-based API for hipcub::DeviceReduce::Reduce

6dc51fb

Env-based API for hipcub::DeviceReduce::Sum

a5a4d57

Fetch libhipcxx from system package

9416683

fix(hipcub): use libhipcxx's requires and templates

a45aa0b

Remove temporary workaround for libhipcxx include path

538e37b

The workaround is no longer necessary.

Revert "Use HIPCUB_PRAGMA_UNROLL_FULL and HIPCUB_PRAGMA_NOUNROLL inst…

784ab81

…ead of _CCCL_* variants" This reverts commit 587574e.

Revert "Use HIPCUB_SORT_MAYBE_UNROLL instead of _CCCL_* variant"

16f4e95

This reverts commit 13b1344.

fix(hipccl/hipcub): add missing libhipcxx dependency for hipCUB package

3f374c1

Beanavil force-pushed the users/beanavil/hipccl/hipcub-cccl-3-1-parity-upstream branch from 5f5e253 to 3f374c1 Compare November 2, 2026 17:02

NguyenNhuDi mentioned this pull request Nov 2, 2026

feat(hipccl): change BUILD_CODE_COVERAGE to <LIB>_ENABLE_COVERAGE #14895

Open

1 task

Merge branch 'develop' into users/beanavil/hipccl/hipcub-cccl-3-1-par…

da60eed

…ity-upstream

Merge branch 'develop' into users/beanavil/hipccl/hipcub-cccl-3-1-par…

b7cdc6a

…ity-upstream

Navigation Menu

Repository navigation

feat(hipccl): hipCUB's parity with CCCL 3.1.x - #13083

feat(hipccl): hipCUB's parity with CCCL 3.1.x#13083
Beanavil wants to merge 66 commits into
developfrom
users/beanavil/hipccl/hipcub-cccl-3-1-parity-upstream

Beanavil commented Oct 5, 2026

Uh oh!

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

edited

Loading

Uh oh!

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

edited

Loading

Uh oh!

NguyenNhuDi left a comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

NguyenNhuDi Oct 30, 2026

Uh oh!

Beanavil Nov 2, 2026

Uh oh!

NguyenNhuDi Nov 2, 2026 •

edited

Loading

Uh oh!

codecov Bot commented Nov 3, 2026 •

edited

Loading

Uh oh!

stanleytsang-amd commented Nov 4, 2026

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

Conversation

Beanavil commented Oct 5, 2026

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

Uh oh!

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

Uh oh!

✅ All Checks Passed — Ready for Review

Uh oh!

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

Uh oh!

Uh oh!

NguyenNhuDi left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

NguyenNhuDi Oct 30, 2026

Choose a reason for hiding this comment

Uh oh!

Beanavil Nov 2, 2026

Choose a reason for hiding this comment

Uh oh!

NguyenNhuDi Nov 2, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

codecov Bot commented Nov 3, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

stanleytsang-amd commented Nov 4, 2026

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

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

edited

Loading

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

edited

Loading

NguyenNhuDi Nov 2, 2026 •

edited

Loading

codecov Bot commented Nov 3, 2026 •

edited

Loading