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

amd-jmahovsky · 2026-10-29T16:53:21Z

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.

…ind/mismatch Upstream: NVIDIA/cccl@1c3c7ddee12 (NVIDIA/cccl#6019) Upstream: NVIDIA/cccl@8ca3dc3 (NVIDIA/cccl#6018) HIP previously took no change for these (performance-only). Bring system/hip closer to upstream: - find.h: find_if_n unwraps contiguous iterators and builds the zip/counting iterator with auto + variadic thrust::make_zip_iterator. transform_input_iterator_t now calls op(*input) directly instead of materializing the value_type (same semantics as ::cuda::transform_iterator; avoids NVIDIA/cccl#3591); prefix operator++ returns self_t&. The dedicated iterator is kept because thrust::transform_iterator still converts to the value_type and ::cuda::transform_iterator needs libhipcxx. - mismatch.h: drop transform_pair_of_input_iterators_t and enable the zip_iterator + zip_function implementation (thrust:: versions, so it also works without libhipcxx). Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…function from Thrust (NVIDIA/cccl#6188) Upstream: NVIDIA/cccl@5303f65 Completes the earlier partial port. libhipcxx is now mandatory (libcxx.h requires >= 3.4), so the compatibility reason for keeping thrust::detail integer helpers no longer applies: - delete thrust/detail/integer_math.h. - execute_with_allocator.h: divide_ri -> _THRUST_LIBCXX::ceil_div. - mr/{pool,disjoint_pool}.h: log2_ri -> _THRUST_LIBCXX::ceil_ilog2, is_power_of_2 -> _THRUST_LIBCXX::is_power_of_two. - mr/pool_options.h: validate() uses is_power_of_two with the upstream explicit != 0 checks. - system/cuda/detail/dispatch.h: take upstream (defines its own thrust::detail::is_negative). - system/hip/detail/dispatch.h: same is_negative, spelled with _THRUST_STD/THRUST_HOST_DEVICE; hip scan.h/sort.h drop the integer_math.h include. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…r includes Related upstream commit: 6d3ffc900fb Port `thrust::transform_input_output_iterator` to `cuda` (NVIDIA/cccl#5113) Related local port commit: 251fd49 libhipcxx is now mandatory, so internal_functional.h includes ::cuda discard_iterator, tabulate_output_iterator and transform_input_output_iterator unconditionally, as upstream does. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…e_cvref_t in pointer Related upstream commit: 30c02fd29dd Refactor thrust::pointer (NVIDIA/cccl#5478) Related local port commit: a3de492 libhipcxx is now mandatory, so thrust::pointer uses _THRUST_STD::remove_cvref_t (upstream ::cuda::std::remove_cvref_t) instead of the ::internal fallback alias. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

… type_traits checks Related upstream commit: ca57b79824c Inherit `thrust::transform_iterator` traversal from base iterator traversal (NVIDIA/cccl#5883) Related local port commit: 308ae2c libhipcxx is now mandatory, so drop the _THRUST_HAS_DEVICE_SYSTEM_STD guards in test/test_type_traits.cpp and testing/type_traits.cu: the libcxx_make_{counting,transform,zip}_iterator propagation combinations and the thrust::pair/tuple relocatability checks always compile, as upstream. The !_THRUST_HAS_DEVICE_SYSTEM_STD <type_traits> fallback include is removed. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…n iterator_facade(_category) Related upstream commit: abcb70514b7 Refactor thrust::iterator_facade_category (NVIDIA/cccl#5877) Related local port commit: 42453aa libhipcxx is now mandatory, so iterator_facade_category.h and iterator_facade.h use _THRUST_STD::_If (upstream ::cuda::std::_If) instead of the ::internal::If fallback from libcxx_wrapper/std/__type_traits/conditional.h, and iterator_facade.h drops its !_THRUST_HAS_DEVICE_SYSTEM_STD <cstddef>/<type_traits> fallback includes. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…hipcxx include fallbacks Related upstream commit: da893945e45 Modernize Thrust examples (NVIDIA/cccl#5670) Related local port commit: ec82559 libhipcxx is now mandatory, so remove the `#if !_THRUST_HAS_DEVICE_SYSTEM_STD` <iterator> fallback includes from examples/{repeated,strided,tiled}_range.cu and the <functional> fallback from examples/minimal_custom_backend.cu (pre-existing, same pattern); the include lists now match upstream apart from the local include/host_device.h. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…des in xor_combine_engine_max Related upstream commit: bb718514624 Drop Thrust mpl math (NVIDIA/cccl#5897) Related local port commit: 1a7094e Include _THRUST_STD_INCLUDE(__algorithm/max.h) and __algorithm/min.h instead of the <algorithm> and <type_traits> umbrella headers, as upstream (both granular headers exist in libhipcxx, which is now mandatory). The host <cstddef> include is kept for the unqualified size_t. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…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>

…ocator <type_traits> fallbacks Related upstream commit: 700f2a867a7 Refactor Thrust `destroy_range` and `device_[new|delete|free]` (NVIDIA/cccl#6134) Related upstream commit: fbb9479172c Refactor Thrust allocator internals (NVIDIA/cccl#6136) Related local port commits: 63becda, 17ee233 libhipcxx is now mandatory, so remove the `#if !_THRUST_HAS_DEVICE_SYSTEM_STD` <type_traits> fallbacks from detail/allocator/{destroy_range,value_initialize_range,fill_construct_range, copy_construct_range}.h, detail/temporary_array.h and detail/execute_with_allocator.h, and add the upstream granular _THRUST_STD_INCLUDE headers (__iterator/advance.h, __iterator/distance.h, __type_traits/is_convertible.h, __type_traits/is_trivially_copy_constructible.h, __type_traits/type_identity.h, __utility/pair.h). thrust/detail/memory_wrapper.h stays: <cuda/std/__cccl/memory_wrapper.h> is not in libhipcxx. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…ceil_div in omp/tbb scan Related upstream commit: 1238f287f4f add parallel scan support for TBB and OMP (NVIDIA/cccl#6178) Related upstream commit: e73cc82956a Move CPP, OMP and TBB exec policies to detail (NVIDIA/cccl#6361) Related upstream commit: c22298179e0 Refactor Thrust OMP system headers (NVIDIA/cccl#6372) Related local port commits: 6cb9dc4, 43961ea, e6c2a3e libhipcxx is now mandatory, so: - system/omp/detail/scan.h: replace the <algorithm>/<iterator>/<numeric>/<type_traits> umbrella includes and thrust/{advance,distance}.h with the upstream granular headers (__cmath/ceil_div.h, __functional/invoke.h, __iterator/{advance,distance}.h, __numeric/{exclusive_scan,inclusive_scan,reduce}.h, __type_traits/{conditional,is_same}.h; plus __algorithm/{max,min}.h for the max/min calls), and compute block_size with _THRUST_LIBCXX::ceil_div(n, num_threads) as upstream. - system/tbb/detail/scan.h: thrust/{advance,distance}.h -> _THRUST_STD __iterator/{advance,distance}.h. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

… scan Related upstream commit: 26c79caf60e fix omp scan bug (#6560) (NVIDIA/cccl#6704) Related local port commit: ffd4daf The port replaced `_CCCL_ASSERT(num_threads > 1, ...)` with a comment because the macro was assumed unavailable. libhipcxx (now mandatory) provides _CCCL_ASSERT (cuda/std/__cccl/assert.h, reached via thrust/detail/config.h), so restore the upstream assertion. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…unt in hip set_operations Related upstream commit: 638c03d54f7 Replace __popc with cude::std::popcounter (NVIDIA/cccl#6213) Related local port commit: baac406 The port changed only system/cuda. libhipcxx is now mandatory, so system/hip/detail/set_operations.h follows upstream: thread_output_count uses _THRUST_STD::popcount(static_cast<unsigned>(active_mask)) instead of the __popc intrinsic, includes _THRUST_STD_INCLUDE(__bit/popcount.h), and the max/min includes are no longer guarded (the !_THRUST_HAS_DEVICE_SYSTEM_STD algorithm_wrapper.h fallback is removed). Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…hs in detail/type_traits.h Related upstream commit: c859e948453 Improve `invoke` machinery (NVIDIA/cccl#6227) Related local port commit: 33e87b6 (pre-rebase SHA; folded into the rebased history) The port switched only the libhipcxx branch to _THRUST_STD::invoke_result and left the "TODO(libhipcxx): remove this path" rocprim/std fallbacks. libhipcxx is now mandatory, so thrust::detail::invoke_result is just _THRUST_STD::invoke_result as upstream, and the dead rocprim/std fallbacks for the ::internal aliases (is_arithmetic, is_integral, make_unsigned, decay_t, remove_cvref, ...), ::internal::detail::common_type and the <rocprim/type_traits_functions.hpp> include are removed. The ::internal aliases themselves stay for now (they still have users). Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

… in is_contiguous_iterator Related upstream commit: 13886f72f16 Add host standard library detection (NVIDIA/cccl#6244) Related local port commit: 1a0de38 The port was N/A because _CCCL_HOST_STD_LIB was unavailable without libhipcxx. libhipcxx is now mandatory and provides it (cuda/std/__cccl/host_std_lib.h, via the libcxx config headers), so is_contiguous_iterator.h detects libc++, libstdc++ and the MSVC STL with _CCCL_HOST_STD_LIB(LIBCXX/LIBSTDCXX/STL) instead of defined(_LIBCPP_VERSION)/defined(__GLIBCXX__)/defined(_MSC_VER), as upstream. Note: the old "Use defined(_MSC_VER) due to Windows CI failures" workaround is dropped; please check Windows CI. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…evice_new.h Related upstream commit: f1b4a17019b Always include `<new>` when we need operator new for clang-cuda (NVIDIA/cccl#6310) Related local port commit: 5c1c329 libhipcxx is now mandatory, so: - detail/internal_functional.h and system/detail/generic/uninitialized_copy.inl include _THRUST_STD_INCLUDE(__new/device_new.h) unconditionally (the generic .inl also drops its !_THRUST_HAS_DEVICE_SYSTEM_STD <type_traits> fallback). - system/hip/detail/uninitialized_{copy,fill}.h now include __new/device_new.h like their system/cuda counterparts, since they also use placement ::new in device code (clang-based HIP has the same need as clang-cuda). Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…with _THRUST_STD::apply Related upstream commit: 52eb28883c4 Simplify thrust::zip_function (NVIDIA/cccl#6321) Related local port commit: 0d3216d Also related: 7b52284299d Rename our namespace macros (NVIDIA/cccl#5538) (the namespace fallback removed below) The port kept the local zip_detail::apply helper and the non-libhipcxx `namespace internal` path for compatibility. libhipcxx is now mandatory, so zip_function.h now matches upstream: - zip_function::operator() returns _THRUST_STD::apply(func, _THRUST_STD::forward<Tuple>(args)); thrust::detail::zip_detail::{apply,apply_impl} are removed. - include _THRUST_LIBCXX_INCLUDE(functional) and _THRUST_STD_INCLUDE(tuple) instead of thrust/tuple.h, thrust/type_traits/integer_sequence.h and the libcxx_wrapper fallback. - the proclaims_copyable_arguments<zip_function<F>> specialization is only declared in the libcxx namespace (_THRUST_LIBCXX_NAMESPACE_BEGIN/END). The leftover "C++11 decltype(auto) workaround" comments are removed as upstream did. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…xx thrust::tuple/pair implementation Related upstream commit: 81b9aa5f1cc Move the implementation of `tuple` into its own file (NVIDIA/cccl#6336) Related local port commit: fe7ee5f Upstream thrust::tuple/thrust::pair are only aliases of ::cuda::std::tuple/pair. rocThrust still carried the old standalone implementation behind `#if !_THRUST_HAS_DEVICE_SYSTEM_STD` ("TODO(libhipcxx): remove this path"). libhipcxx is now mandatory, so: - thrust/tuple.h and thrust/pair.h keep only the _THRUST_STD alias branch, as upstream. - delete thrust/detail/tuple.inl and thrust/detail/pair.inl (only included by the removed branch). - drop the fallback <cstddef>/<type_traits>/<utility> includes and the fallback explanation comment. The local THRUST_DOXYGEN_INVOKED spelling and the _THRUST_STD_INCLUDE(utility) include in pair.h are kept. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…onal> in generic scan Related upstream commit: 85bf1b00c4d Refactor generic thrust scan dispatch (NVIDIA/cccl#6360) Related local port commit: fc66966 system/detail/generic/scan.h uses _THRUST_STD::plus<>; include _THRUST_LIBCXX_INCLUDE(functional) (upstream <cuda/functional>) instead of <thrust/functional.h>, now that libhipcxx is mandatory. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…allbacks, granular includes Related upstream commit: 865e3d6429f Inline remaining `*.inl` files in tbb and seq backends (NVIDIA/cccl#6437) Related local port commit: cd4d303 libhipcxx is now mandatory, so the inlined seq/tbb backend headers follow upstream's includes: - sequential/sort.h: drop the unused libcxx_wrapper conjunction.h include and the !_THRUST_HAS_DEVICE_SYSTEM_STD <type_traits> fallback. - sequential/stable_merge_sort.h: include only _THRUST_STD_INCLUDE(__algorithm/min.h), unguarded (algorithm_wrapper.h fallback and the <iterator> umbrella removed). - sequential/stable_primitive_sort.h: drop the <type_traits> fallback. - tbb/detail/reduce_by_key.h: unguarded __algorithm/{max,min}.h and __type_traits/void_t.h, plus upstream's __iterator/reverse_iterator.h and __utility/pair.h instead of the <iterator> umbrella; algorithm_wrapper.h/<type_traits> fallback removed. (thrust::pair -> _THRUST_STD::pair is done separately for CCCL 91434d76ab3.) The local _THRUST_IF_TARGET/THRUST_COMPILER spellings in sequential/sort.h are kept. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…th includes Related upstream commit: 24deb145aec Cleanup `thrust::complex` math includes and functions (NVIDIA/cccl#6546) Related upstream commit: 369f2f72eca Fix merge conflicts from dropping headers (NVIDIA/cccl#6563) Related local port commits: f58e510, 556bd2a libhipcxx is now mandatory, so the thrust::complex headers can use the same granular libhipcxx includes as upstream instead of the <cmath> umbrella: - thrust/detail/complex/*.h include the upstream __cmath/* headers (plus limits / __type_traits/common_type.h where upstream does). arithmetic.h keeps cfloat for FLT_MIN/DBL_MIN. - thrust/complex.h includes __type_traits/{common_type,remove_reference}.h as upstream and drops the dead !_THRUST_HAS_DEVICE_SYSTEM_STD <type_traits> fallback. - ::internal::common_type_t is replaced by _THRUST_STD::common_type_t (as upstream) in complex.h, arithmetic.h, cpow.h and test_complex.cpp, and the ::internal::common_type_t alias (TODO(libhipcxx)) is removed from thrust/detail/type_traits.h. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…cludes Related upstream commit: 4b10fd26255 Cleanup includes in thrust (NVIDIA/cccl#6547) Related local port commit: ed71224 The port kept the ::internal::is_integral / ::internal::If wrappers, the local umbrella include blocks and the dual less/greater list in cuda sort.h because of non-libhipcxx builds. libhipcxx is now mandatory, so: - ::internal::is_integral(_v) -> _THRUST_STD::is_integral(_v) in detail/type_traits.h, detail/vector_base.inl and counting_iterator.h (is_integral_v in distance_to, as upstream); the now-unused ::internal::is_integral/is_integral_v aliases are removed. - ::internal::If -> _THRUST_STD::_If (as upstream) in counting_iterator.h, iterator_adaptor_base.h, iterator_category_to_traversal.h, generic/reduce_by_key.inl and hip/detail/reduce_by_key.h; the libcxx_wrapper conditional.h includes are replaced by _THRUST_STD_INCLUDE(__type_traits/conditional.h). (The wrapper file itself stays because the dead fallback branch of the construct_at wrapper still includes it.) - counting_iterator.h and iterator_adaptor_base.h take the upstream granular __type_traits includes (drops the rocprim/type_traits.hpp fallback); minimum_system.h and iterator_category_to_traversal.h drop the type_traits includes as upstream. - cuda sort.h: can_use_primitive_sort uses the upstream is_same_v list without the duplicated _THRUST_STD:: comparators (_THRUST_STD is ::cuda::std there). thrust/detail/type_traits.h keeps its _THRUST_STD_INCLUDE(type_traits) umbrella; switching it to upstream's granular list would require the upstream include fixups in ~100 other headers. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…jection Related upstream commit: b85d41e36d4 Fixes for thrust::shuffle (#7172) (NVIDIA/cccl#7189) Related local port commit: bd2d9c0 The port replaced _THRUST_STD::max/_THRUST_STD::bit_width in feistel_bijection with a hand-written bit-count loop because bit_width is C++20 and unavailable without libhipcxx. libhipcxx is now mandatory, so random_bijection.h follows upstream: max_index/total_bits use _THRUST_STD::max and _THRUST_STD::bit_width, the raw <cuda/std/...> includes and ::cuda::std:: spellings become _THRUST_STD_INCLUDE / _THRUST_STD::, and the upstream granular includes are taken (__type_traits/{is_convertible,is_integral}.h, __utility/forward.h, bit, cstdint). __algorithm/max.h is included explicitly (upstream gets it transitively). Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…tors in hip find/mismatch Related upstream commit: 8ca3dc31e90 Refactor thrust::mismatch to use CUDA iterators (NVIDIA/cccl#6018) Related upstream commit: 1c3c7ddee12 Refactor iterator usage in thrust/cuda find() (NVIDIA/cccl#6019) Related local port commit: e46dc36 e46dc36 brought the upstream find/mismatch rewrite to system/hip using thrust:: iterators, plus a local transform_input_iterator_t because thrust::transform_iterator materializes the value_type before calling the predicate (NVIDIA/cccl#3591). libhipcxx is now mandatory, so hip now uses the same libcu++ iterators as upstream's system/cuda: - find.h: find_if_n zips _THRUST_LIBCXX::make_transform_iterator( try_unwrap_contiguous_iterator(first), predicate) with _THRUST_LIBCXX::counting_iterator<Size>(0) and uses a _THRUST_STD::tuple<bool, Size> result type; __find_if:: transform_input_iterator_t and the dead !_THRUST_HAS_DEVICE_SYSTEM_STD include block are removed; upstream granular includes are taken (<thrust/functional.h> kept explicitly for the placeholders in find). - mismatch.h: _THRUST_LIBCXX::make_zip_iterator and _THRUST_LIBCXX::zip_function(binary_pred), with upstream's granular includes instead of thrust/{distance,pair,zip_function}.h and thrust/iterator/zip_iterator.h. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…llbacks in zip_iterator and iterator tests Related upstream commit: 7b52284299d Rename our namespace macros (NVIDIA/cccl#5538) Related upstream commit: 8b68d958473 Drop internal uses of `thrust::reverse_iterator` (NVIDIA/cccl#5616) Related local port commit: 148b6e8 libhipcxx is now mandatory, so every !_THRUST_HAS_DEVICE_SYSTEM_STD branch in the files these upstream commits touched is dead code. Upstream has no such fallbacks; remove them to match: - thrust/iterator/zip_iterator.h: remove the fallback zip_iterator machinery (advance_iterator, tuple_impl_specific, tuple_of_value_types, tuple_of_iterator_references helpers, ...) and the associated TODO; includes are now unconditional _THRUST_STD_INCLUDE(...). - thrust/iterator/detail/tuple_of_iterator_references.h: drop fallback branches. - thrust/detail/vector_base.h: drop fallback <type_traits> include. - testing/{copy,scan}.cu, testing/unittest/testframework.h: unconditionally specialize _THRUST_STD traits and drop fallback includes. - testing/{reverse,transform,transform_input_output,transform_output, tabulate_output,zip}_iterator.cu: iterator concept static_asserts are now unconditional, as upstream. The namespace macro rename from 7b52284299d needs no further change; the local _THRUST_STD_NAMESPACE_BEGIN/_THRUST_LIBCXX_NAMESPACE_BEGIN spellings are already the equivalents. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…ialized_copy delegates to copy_n Related upstream commit: 4b2201184c4 Use `thrust::copy` in `thrust::uninitialized_copy[_n]` in CUDA system when possible (NVIDIA/cccl#5181) Related upstream commit: ee141c6b175 Fix uninitialized read in uninitialized_copy_n (NVIDIA/cccl#5811) Related local port commit: 9d7099b Related local port commit: c91e8ff Both upstream commits were ported to system/cuda only, with the HIP counterpart marked "no change: performance-only". Bring system/hip in line: - hip_rocprim::uninitialized_copy_n now delegates to hip_rocprim::copy_n when the output value type is trivially constructible from, and trivially assignable from, the input reference (the #5811 condition). For trivially relocatable device-to-device copies this turns the placement-new parallel_for kernel into a hipMemcpyAsync. - copy_n is forward-declared rather than included, as upstream does, to avoid the circular include via internal/copy_cross_system.h. - Includes follow upstream's granular libhipcxx headers (__iterator/distance.h, __iterator/incrementable_traits.h, and the is_trivially_{constructible,assignable} traits actually used). Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…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>

…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>

…seq.h into execution_policy.h Related upstream commit: 2f39a342b93 Refactor Thrust execution policies (NVIDIA/cccl#5821) Related local port commit: db142f7 The local port applied this refactor to the cuda backend only. Do the same for hip: - thrust/system/hip/detail/par.h and par_to_seq.h are removed; their contents (execute_on_stream*, par_t, par_nosync_t and the hip-specific par_det_t/par_det_nosync_t, and the par/par_nosync/par_det/par_det_nosync constants) now live in hip/detail/execution_policy.h. All existing aliases in thrust::hip, thrust::system::hip and thrust::system::hip::detail are kept. - cvt_to_seq takes upstream's simple form (returns thrust::detail::seq_t), dropping the unused has_par trait and the #if 0 block. - hip/detail/util.h forward-declares execution_policy instead of including hip/detail/execution_policy.h at the top, as cuda/detail/util.h does. It includes it at the bottom so headers that relied on util.h for the policy definitions keep working. - Algorithm headers, temporary_buffer.h, hip/execution_policy.h and the allocator_aware_policies tests include hip/detail/execution_policy.h instead of the removed headers. Compatibility: code that includes thrust/system/hip/detail/par.h or par_to_seq.h directly must switch to thrust/system/hip/execution_policy.h. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

… benchmark counterparts Related upstream commit: 106665c64f4 Use SPDX license identifiers in CUB (NVIDIA/cccl#6441) Related local port commit: 6339d01 The local port marked these upstream files NOLOCAL, but rocThrust has counterparts under different names, which kept the long-form BSD header: - set_operations/base.cuh -> benchmark/bench/set_operations/base.hpp - set_operations/by_key.cuh -> benchmark/bench/set_operations/by_key.hpp - vectorized_search/base.cu -> benchmark/bench/vectorized_search/basic.cu - transform/{babelstream,fib}.cu -> benchmark/bench/transform/basic.cu Take upstream's SPDX header, keeping the AMD Modifications line as an SPDX-FileCopyrightText line, as the local port did for the other benchmarks. The AMD-only MIT-licensed set_operations/*.cu files are unchanged. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…ementation Related upstream commit: 5197e4b8a58 Refactor and condense thrust::copy implementation (NVIDIA/cccl#5491) Related local port commit: f0a16da The local port only condensed the CUDA backend. Apply the same refactor to the HIP backend: fold hip/detail/internal/copy_cross_system.h and hip/detail/internal/copy_device_to_device.h into hip/detail/copy.h, in a single __copy namespace mirroring upstream: - trivial_cross_system_copy_n (H->D and D->H overloads) - non_trivial_cross_system_copy_n (H->D, and D->H under the HIP compiler) - cross_system_copy_n, which returns early for n == 0 and selects the trivial or non-trivial path with if constexpr on is_indirectly_trivially_relocate_to_v instead of true_type/false_type tag dispatch - device_to_device, likewise using if constexpr cross_system_copy is removed; the cross-system copy() overload now calls cross_system_copy_n with the computed distance, as upstream does. The HIP host/device workaround structs in copy()/copy_n() are kept. Not ported: upstream's device_to_device unwraps cuda::zip_transform_iterator into cub_transform_many. hip_rocprim::transform has no multi-input transform entry point to forward to, so non-trivial D->D copies still go through hip_rocprim::transform with identity. Upstream's proclaim_copyable_arguments wrapper is also omitted because the HIP transform does not use address-stability information. The removed detail headers were only included by hip/detail/copy.h. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

… cleanup Related upstream commit: 4b10fd26255 Cleanup includes in thrust (NVIDIA/cccl#6547) Related local port commit: ee821be Bring the include structure of thrust/ in line with upstream v3.2.1 now that libhipcxx is mandatory: - Apply the upstream granular libcu++ include regions to the files touched by #6547 (umbrella <cuda/std/X> headers replaced by __X/*.h headers). - Remove the ::internal:: type-trait aliases from detail/type_traits.h and use _THRUST_STD:: directly, matching upstream. - Drop dead !_THRUST_HAS_DEVICE_SYSTEM_STD fallback code (thrust::swap fallback, aligned_storage_size fallback, actor argument variant, raw_reference_cast apply branch, iter_swap/swap_ranges fallbacks). - Replace umbrella includes in the hip backend with granular headers and add direct includes for libcu++ symbols used by hip backend files. Local adaptations kept: thrust/detail/memory_wrapper.h, __host_stdlib/ stdexcept, hip hints in #if 0 blocks, THRUST_COMPILER(NVRTC), _CCCL_CUDA_COMPILATION(), the [NON-CCCL PARITY] cstddef include in type_traits.h, erfcinv includes, and the find.h tuple umbrella. The libcxx_wrapper directory and tests are left for a separate effort. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…pers to libhipcxx Related upstream commit: ab99d87e866 Move nvhpc header wrappers to libcudacxx. (NVIDIA/cccl#6167) Related local port commit: 48c541b Apply the thrust/thrust portion of upstream CCCL commit ab99d87e866 ("Move nvhpc header wrappers to libcudacxx (#6167)"): delete the local thrust::detail::{algorithm,memory,numeric}_wrapper.h nvhpc include-cycle guards and pull the equivalents in from libhipcxx instead. Since rocThrust treats libhipcxx as a hard dependency, the upstream include of <cuda/std/__cccl/{algorithm,memory}_wrapper.h> is wrapped in the _THRUST_STD_INCLUDE macro, and further updated to the paths CCCL uses as of v3.4.0: <cuda/std/__host_stdlib/{algorithm,memory}> (the __cccl/*_wrapper.h headers were themselves renamed upstream). No numeric_wrapper.h include existed in any of the touched files. thrust/detail/allocator/allocator_traits.h was deleted entirely further upstream in favor of using cuda::std::allocator_traits directly; that is a separate, unrelated refactor and is not part of this change - the file and its callers are kept as-is here, with only its own wrapper include updated. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

…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>

…through parallel_for Related upstream commit: e93c9c33537 Port thrust::generate[_n] to cub::DeviceTransform::Fill (NVIDIA/cccl#5807) Upstream: NVIDIA/cccl@e93c9c3 Related local port commit: 575e1bf The local port only changed the CUDA backend; the HIP counterpart was left behind as performance-only. Upstream stops routing generate_n through for_each_n with a generate_f functor that const_casts each element, and writes the generated values straight to the output instead (cub::DeviceTransform with no inputs). rocPRIM has no zero-input transform, so hip/detail/generate.h does the same with parallel_for and a __generate::functor that assigns result[idx] = generator(). This drops the const_cast and the per-element reference dereference through for_each_f. parallel_for still synchronizes when the policy requires it, which covers the stream synchronization upstream added in d91b32f1bf6 (NVIDIA/cccl#5889). Also as upstream: THRUST_EXEC_CHECK_DISABLE on generate_n, the generate_f functor and the "for_each" comments are removed, and the includes switch from thrust/distance.h and for_each.h to execution_policy.h, parallel_for.h, util.h and _THRUST_STD_INCLUDE(iterator). Not compiled, tests not run Co-Authored-By: Claude <noreply@anthropic.com>

…_n delegates to fill_n for trivial types Related upstream commit: 50929671faf Use DeviceTransform::Fill in uninitialized_fill[_n] when possible (NVIDIA/cccl#5813) Upstream: NVIDIA/cccl@5092967 Related local port commit: e592f35 The local port took the CUDA side of this commit but left the HIP backend always constructing elements with placement new through parallel_for. Mirror upstream in thrust/system/hip/detail/uninitialized_fill.h: when the value type is trivially constructible and trivially assignable from T, there is no constructor side effect to preserve, so delegate to hip_rocprim::fill_n. Otherwise keep the existing placement-new functor. Not compiled, tests not run Co-Authored-By: Claude <noreply@anthropic.com>

… a custom kernel Related upstream commit: 49b642e53dd Use a custom `iter_swap` kernel in Thrust (NVIDIA/cccl#5843) Upstream: NVIDIA/cccl@49b642e Related local port commit: eedd8e8 The local port updated the CUDA backend only. Mirror upstream in thrust/system/hip/detail/iter_swap.h: the host path no longer goes through thrust::swap_ranges (and drops the thrust/swap.h include) but launches a single-thread detail::iter_swap_kernel on the policy's stream, checks the launch and calls synchronize_optional. The device path swaps in place with _THRUST_STD::swap. The war_nvbugs_881631 struct becomes host_path with upstream's comment explaining why the host path sits outside the target switch. The triple_chevron launch maps to hipLaunchKernelGGL, as in parallel_for.h. Not compiled, tests not run Co-Authored-By: Claude <noreply@anthropic.com>

… policies like cuda Related upstream commit: 2f39a342b93 Refactor Thrust execution policies (NVIDIA/cccl#5821) Upstream: NVIDIA/cccl@2f39a34 Related local port commit: 7544a6b Related local parity commit: 2b563b3 2b563b3 already folded par.h/par_to_seq.h into thrust/system/hip/detail/execution_policy.h. This commit applies the rest of upstream's cuda execution_policy.h refactor to the hip file: - tag derives from thrust::detail::allocator_aware_execution_policy<execution_policy>, and execution_policy<Derived>::operator tag() returns {}; - execute_on_stream{,_nosync,_deterministic,_nosync_deterministic} use `= default` default constructors and construct the injected base class name instead of a base_t alias; - par_t, par_nosync_t, par_det_t and par_det_nosync_t drop their user-written constexpr default constructors and base_t aliases; - the public types sit in the execution_policies Doxygen group, with upstream's comments adapted to HIP and a usage example on par; - `namespace system::hip`, and system::hip::detail no longer re-exports execution_policy and tag, matching upstream's system::cuda::detail. The AMD-only deterministic stream bases are unchanged. Not compiled, tests not run Co-Authored-By: Claude <noreply@anthropic.com>

…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>

…r_if proclaim copyable arguments Related upstream commit: fe943f23ced Ensure thrust::scatter_if uses UBLKCP (NVIDIA/cccl#5952) Upstream: NVIDIA/cccl@fe943f2 Related local port commit: cd93370 The local port updated the CUDA backend only. Mirror upstream in thrust/system/hip/detail/scatter.h: scatter, the predicated scatter_if and the stencil-only scatter_if overload now wrap _THRUST_STD::identity{} in _THRUST_LIBCXX::proclaim_copyable_arguments, so the transform may copy its arguments instead of reading them in place. The hip stencil-only overload never had the stray Predicate template parameter that upstream removes. The file now includes __functional/address_stability.h explicitly instead of relying on transform.h to bring it in. Not compiled, tests not run Co-Authored-By: Claude <noreply@anthropic.com>

…sform_n with copyable arguments Related upstream commit: ad70011e990 Support fancy iterators in vectorized transform and port thrust::tabulate to it (NVIDIA/cccl#6012) Upstream: NVIDIA/cccl@ad70011 Related local port commit: 39a06d7 Related local parity commit: 4698506 4698506 turned hip tabulate into a transform over a counting_iterator, but called transform(begin, begin + count, ...) with the bare tabulate_op. Finish matching upstream in thrust/system/hip/detail/tabulate.h: call hip_rocprim::transform_n (added by the 76fc4572271 parity commit) on _THRUST_LIBCXX::counting_iterator<size_type>{} and count, and wrap tabulate_op in _THRUST_LIBCXX::proclaim_copyable_arguments. Includes follow upstream: __functional/address_stability.h and the _THRUST_STD iterator header replace __iterator/distance.h and __iterator/incrementable_traits.h. The counting_iterator include stays, since upstream only gets it through transform.h. Not compiled, tests not run Co-Authored-By: Claude <noreply@anthropic.com>

…ity in counting_iterator

Naraenda

I have added 6b1e529 to ensure it builds with passing tests (locally at least).

There are also some missing includes that could be added, but weren't needed to make the tests build/pass. I have attached those to the review. I'm not sure if those should be addressed here or not.

Naraenda · 2026-11-02T14:04:11Z

+#  include _THRUST_STD_INCLUDE(__utility/swap.h)

 THRUST_NAMESPACE_BEGIN
 namespace hip_rocprim
 {
+namespace detail
+{
+template <typename Pointer1, typename Pointer2>
+ROCPRIM_KERNEL void iter_swap_kernel(Pointer1 a, Pointer2 b)


Use of ROCPRIM_KERNEL without including rocprim/config.hpp.

Suggested change

# include _THRUST_STD_INCLUDE(__utility/swap.h)

THRUST_NAMESPACE_BEGIN

namespace hip_rocprim

{

namespace detail

{

template <typename Pointer1, typename Pointer2>

ROCPRIM_KERNEL void iter_swap_kernel(Pointer1 a, Pointer2 b)

# include _THRUST_STD_INCLUDE(__utility/swap.h)

# include <rocprim/config.hpp>

THRUST_NAMESPACE_BEGIN

namespace hip_rocprim

{

namespace detail

{

template <typename Pointer1, typename Pointer2>

ROCPRIM_KERNEL void iter_swap_kernel(Pointer1 a, Pointer2 b)

Naraenda · 2026-11-02T14:07:59Z

+  // copy constructor which could have a side effect and we can delegate to copy_n.
+  using input_ref_t  = thrust::detail::raw_reference_t<_THRUST_STD::iter_reference_t<InputIt>>;
+  using output_ref_t = thrust::detail::raw_reference_t<_THRUST_STD::iter_reference_t<OutputIt>>;
+  using output_t     = thrust::detail::it_value_t<OutputIt>;


This is missing #include <thrust/iterator/iterator_traits.h>

Naraenda · 2026-11-02T14:08:27Z

+  // if the output type is trivially constructible from the input, it has no side effect, and we can skip placement new
+  // and calling a constructor. Furthermore, if assigning the input value to an output element is also trivial, there is
+  // no copy constructor which could have a side effect and we can delegate to fill_n.
+  using value_t = thrust::detail::it_value_t<Iterator>;


Missing #include <thrust/iterator/iterator_traits.h>

todo-rocthrust.md and cccl-investigation-v3.2.1.md are working files of the CCCL sync skills and were not meant to be checked in (added in 787f309). Co-Authored-By: Claude <noreply@anthropic.com>

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

…/amd-jmahovsky/hipccl/rocthrust_cccl32_additional

…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 a review from a team as a code owner October 29, 2026 16:53

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

amd-jmahovsky changed the title ~~Users/amd jmahovsky/hipccl/rocthrust cccl32 additional~~ feat(hipccl): rocThrust's parity with CCCL 3.2 ADDITIONAL Oct 29, 2026

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

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

Naraenda requested a review from a team as a code owner October 30, 2026 09:41

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

amd-jmahovsky and others added 22 commits October 30, 2026 12:48

amd-jmahovsky and others added 11 commits October 30, 2026 12:48

amd-jmahovsky force-pushed the users/amd-jmahovsky/hipccl/rocthrust_cccl32_additional branch from cc299b7 to 464aa44 Compare October 30, 2026 16:49

amd-jmahovsky and others added 9 commits October 30, 2026 18:54

fix(rocthrust): break hip include cycles, use _THRUST_STD::type_ident…

6b1e529

…ity in counting_iterator

Naraenda self-requested a review November 2, 2026 13:25

Naraenda reviewed Nov 2, 2026

View reviewed changes

amd-jmahovsky and others added 4 commits November 3, 2026 16:33

chore(rocthrust): remove CCCL 3.2 sync scratch files

b43af49

todo-rocthrust.md and cccl-investigation-v3.2.1.md are working files of the CCCL sync skills and were not meant to be checked in (added in 787f309). Co-Authored-By: Claude <noreply@anthropic.com>

chore(libhipcxx): bump submodule to 87af709ba43

1ecc3a9

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

chore(libhipcxx): bump submodule to 5cf4deb633f

91fc71c

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

Merge branch 'users/amd-jmahovsky/hipccl/rocthrust_cccl32' into users…

0644c30

…/amd-jmahovsky/hipccl/rocthrust_cccl32_additional

amd-jmahovsky merged commit 1f19360 into users/amd-jmahovsky/hipccl/rocthrust_cccl32 Nov 4, 2026

amd-jmahovsky deleted the users/amd-jmahovsky/hipccl/rocthrust_cccl32_additional branch November 4, 2026 17:59

Navigation Menu

Repository navigation

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

feat(hipccl): rocThrust's parity with CCCL 3.2 ADDITIONAL#14670
amd-jmahovsky merged 49 commits into
users/amd-jmahovsky/hipccl/rocthrust_cccl32from
users/amd-jmahovsky/hipccl/rocthrust_cccl32_additional

amd-jmahovsky commented Oct 29, 2026

Uh oh!

Naraenda left a comment

Uh oh!

Naraenda Nov 2, 2026

Uh oh!

Naraenda Nov 2, 2026

Uh oh!

Naraenda Nov 2, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

amd-jmahovsky commented Oct 29, 2026

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

Uh oh!

Naraenda left a comment

Choose a reason for hiding this comment

Uh oh!

Naraenda Nov 2, 2026

Choose a reason for hiding this comment

Uh oh!

Naraenda Nov 2, 2026

Choose a reason for hiding this comment

Uh oh!

Naraenda Nov 2, 2026

Choose a reason for hiding this comment

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants