matyas-streamhpc · Jun 2, 2026 · Jun 2, 2026 · Jun 2, 2026 · Jun 2, 2026 · Jun 2, 2026
@@ -28,12 +28,12 @@ AttributeMacros: [
                   'THRUST_FORCEINLINE',
                   'THRUST_HOST_DEVICE',
                   'THRUST_HOST',
-                  'THRUST_NODISCARD',
                   'THRUST_DEPRECATED',
                   'THRUST_GLOBAL_CONSTANT',
                   # thrust
                   '_CCCL_ALIGNAS_TYPE',
                   '_CCCL_ALIGNAS',
+                  '_CCCL_API',
                   '_CCCL_CONSTEXPR_CXX14',
                   '_CCCL_CONSTEXPR_CXX17',
                   '_CCCL_CONSTEXPR_CXX20',
@@ -45,7 +45,6 @@ AttributeMacros: [
                   '_CCCL_HOST',
                   '_CCCL_NO_UNIQUE_ADDRESS',
                   '_CCCL_NODISCARD_FRIEND',
-                  '_CCCL_NODISCARD',
                   '_CCCL_NORETURN',
                   '_CCCL_VISIBILITY_HIDDEN',
                   'CUB_RUNTIME_FUNCTION',
@@ -66,8 +65,6 @@ AttributeMacros: [
                   '_LIBCUDACXX_EXPORTED_FROM_ABI',
                   '_LIBCUDACXX_EXTERN_TEMPLATE_TYPE_VIS',
                   '_LIBCUDACXX_HIDDEN',
-                  '_LIBCUDACXX_HIDE_FROM_ABI_AFTER_V1',
-                  '_LIBCUDACXX_HIDE_FROM_ABI',
                   '_LIBCUDACXX_INLINE_VISIBILITY',
                   '_LIBCUDACXX_INTERNAL_LINKAGE',
                   '_LIBCUDACXX_METHOD_TEMPLATE_IMPLICIT_INSTANTIATION_VIS',

@@ -3,8 +3,31 @@
 Documentation for rocThrust available at
 [https://rocm.docs.amd.com/projects/rocThrust/en/latest/](https://rocm.docs.amd.com/projects/rocThrust/en/latest/).
 
+## rocThrust *.*.* for ROCm *.*.* (CCCL 3.1 parity)
+
+### Added
+
+* Add `THRUST_DISABLE_BROKEN_GCC_VECTORIZER`
+* Add `thrust::reduce_into`.
+* Add `thrust::shuffle_iterator`.
+* Add `thrust::transform_n` and `thrust::transform_if_n`.
+
+### Changed
+
+* Replaced bad_alloc, cassert, cfloat, cmath, complex, cstddef, cstdint, cstring, initializer_list, limits, move and type_traits operators with `hip::std` operators.
+* Replaced `hip::std::__identity` with `hip::std::identity`.
+* Replaced ``_LIBCUDACXX_HIDE_FROM_ABI`` with ``_CCCL_API``.
+* Defined `THRUST_COMP_DEPR_IMPL` with `_CCCL_WARNING`.
+* Replaced `thrust::constant_iterator`, `thrust::make_discard_iterator`, `thrust::permutation_iterator`, `thrust::strided_iterator`, `thrust::tabulate_output_iterator` and `thrust::transform_output_iterator` with `hip::constant_iterator`, `hip::make_discard_iterator`, `hip::permutation_iterator`, `hip::strided_iterator`, `hip::tabulate_output_iterator` and `hip::transform_output_iterator` instead.
+
+### Removed
+
+* Deprecated `thrust::advance`, `thrust::bit_and`, `thrust::bit_or`, `thrust::bit_xor`, `thrust::distance`, `thrust::equal_to`, `thrust::greater`, `thrust::greater_equal`, `thrust::less`, `thrust::less_equal`, `thrust::logical_and`, `thrust::logical_not`, `thrust::logical_or`, `thrust::maximum`, `thrust::minimum`, `thrust::modulus`, `thrust::multiplies`, `thrust::negate`, `thrust::next`, `thrust::not_equal_to`, `thrust::not_fn`, `thrust::operator_name`, and `thrust::plus`. Use `hip::std::advance`, `hip::std::bit_and`, `hip::std::bit_or`, `hip::std::bit_xor`, `hip::std::distance`, `hip::std::equal_to`, `hip::std::greater`, `hip::std::greater_equal`, `hip::std::less`, `hip::std::less_equal`, `hip::std::logical_and`, `hip::std::logical_not`, `hip::std::logical_or`, `hip::maximum`, `hip::minimum`, `hip::modulus`, `hip::std::multiplies`, `hip::std::negate`, `hip::std::next`, `hip::std::not_equal_to`, `hip::std::not_fn`, `hip::std::operator_name` and `hip::std::plus` instead.
+* Deprecated `thrust::detail::::identity_`. Use `hip::std::type_identity` instead.
+* Removed `THRUST_NODISCARD`.
+
 ## rocThrust 5.0.0 for ROCm 10.2.0
- 
+
 ### Added
 
 * Largely in feature parity with CCCL/thrust v3.0.3.

@@ -38,17 +38,6 @@
 
 PRIMBENCH_REGISTER_TYPE(bench_utils::large_data, "large_data")
 
-template <class T>
-struct less_then_t
-{
-  T m_val;
-
-  __host__ __device__ bool operator()(const T& val) const
-  {
-    return val < m_val;
-  }
-};
-
 template <typename T>
 struct copy_benchmark : public primbench::benchmark_interface
 {
@@ -69,8 +58,8 @@ struct copy_benchmark : public primbench::benchmark_interface
 
   void run(primbench::state& state) override
   {
-    T val = bench_utils::value_from_entropy<T>(bench_utils::get_entropy_percentage(entropy));
-    less_then_t<T> select_op{val};
+    const T val = bench_utils::lerp_min_max<T>(bench_utils::get_entropy_percentage(entropy));
+    bench_utils::less_then_t<T> select_op{val};
 
     bench_utils::caching_allocator_t alloc{};
     thrust::detail::device_t policy{};

@@ -34,17 +34,6 @@
 
 #include "bench_utils.hpp"
 
-template <class T>
-struct less_then_t
-{
-  T m_val;
-
-  __host__ __device__ bool operator()(const T& val) const
-  {
-    return val < m_val;
-  }
-};
-
 template <typename T>
 struct partition_benchmark : public primbench::benchmark_interface
 {
@@ -68,9 +57,9 @@ struct partition_benchmark : public primbench::benchmark_interface
     bench_utils::caching_allocator_t alloc{};
     thrust::detail::device_t policy{};
 
-    T val = bench_utils::value_from_entropy<T>(bench_utils::get_entropy_percentage(entropy_reduction));
+    const T val = bench_utils::lerp_min_max<T>(bench_utils::get_entropy_percentage(entropy_reduction));
 
-    less_then_t<T> select_op{val};
+    bench_utils::less_then_t<T> select_op{val};
 
     thrust::device_vector<T> in = bench_utils::generate(m_items, state.seed);
     thrust::device_vector<T> out(m_items);

@@ -71,7 +71,8 @@ struct reduce_benchmark : public primbench::benchmark_interface
     out_keys = in_keys;
     in_vals  = thrust::device_vector<T>(m_items);
 
-    const size_t unique_keys = thrust::distance(out_keys.begin(), thrust::unique(out_keys.begin(), out_keys.end()));
+    const size_t unique_keys =
+      _THRUST_STD::distance(out_keys.begin(), thrust::unique(out_keys.begin(), out_keys.end()));
     thrust::device_vector<T> out_vals;
     out_vals = thrust::device_vector<T>(unique_keys);
 

@@ -83,7 +83,7 @@ struct base_set_benchmark : public primbench::benchmark_interface
          in.cend(),
          out.begin());
 
-    const size_t items_in_AB = thrust::distance(out.begin(), result_ends);
+    const size_t items_in_AB = _THRUST_STD::distance(out.begin(), result_ends);
 
     state.set_items(m_items);
     state.add_reads<T>(m_items);

@@ -90,7 +90,7 @@ struct by_key_set_benchmark : public primbench::benchmark_interface
       out_keys.begin(),
       out_vals.begin());
 
-    const size_t item_in_AB  = thrust::distance(out_keys.begin(), result_ends.first);
+    const size_t item_in_AB  = _THRUST_STD::distance(out_keys.begin(), result_ends.first);
     const size_t values_read = OpT::read_all_values ? m_items : items_in_A;
 
     state.set_items(m_items);

@@ -33,7 +33,6 @@
 #include <thrust/copy.h>
 #include <thrust/count.h>
 #include <thrust/detail/config/namespace.h>
-#include <thrust/detail/libcxx_wrapper/__functional/address_stability.h>
 #include <thrust/device_vector.h>
 #include <thrust/execution_policy.h>
 #include <thrust/fill.h>
@@ -174,7 +173,7 @@ struct mul
   {
     const T scalar = startScalar;
     thrust::transform(
-      policy, c.begin(), c.end(), b.begin(), ::internal::proclaim_copyable_arguments([=] THRUST_DEVICE(const T& ci) {
+      policy, c.begin(), c.end(), b.begin(), _THRUST_LIBCXX::proclaim_copyable_arguments([=] THRUST_DEVICE(const T& ci) {
         return ci * scalar;
       }));
   }
@@ -195,7 +194,7 @@ struct add
       a.end(),
       b.begin(),
       c.begin(),
-      ::internal::proclaim_copyable_arguments([] THRUST_DEVICE(const T& ai, const T& bi) -> T {
+      _THRUST_LIBCXX::proclaim_copyable_arguments([] THRUST_DEVICE(const T& ai, const T& bi) -> T {
         return ai + bi;
       }));
   }
@@ -217,7 +216,7 @@ struct triad
       b.end(),
       c.begin(),
       a.begin(),
-      ::internal::proclaim_copyable_arguments([=] THRUST_DEVICE(const T& bi, const T& ci) {
+      _THRUST_LIBCXX::proclaim_copyable_arguments([=] THRUST_DEVICE(const T& bi, const T& ci) {
         return bi + scalar * ci;
       }));
   }
@@ -239,7 +238,7 @@ struct nstream
       thrust::make_zip_iterator(a.end(), b.end(), c.end()),
       a.begin(),
       thrust::make_zip_function(
-        ::internal::proclaim_copyable_arguments([=] THRUST_DEVICE(const T& ai, const T& bi, const T& ci) {
+        _THRUST_LIBCXX::proclaim_copyable_arguments([=] THRUST_DEVICE(const T& ai, const T& bi, const T& ci) {
           return ai + bi + scalar * ci;
         })));
   }

@@ -72,7 +72,7 @@ struct transform_reduce_benchmark : public primbench::benchmark_interface
 
     state.run([&] {
       thrust::transform_reduce(
-        policy(alloc).on(state.stream), in.begin(), in.end(), square_t<T>{}, T{}, thrust::plus<T>{});
+        policy(alloc).on(state.stream), in.begin(), in.end(), square_t<T>{}, T{}, _THRUST_STD::plus<T>{});
     });
   }
 

@@ -65,7 +65,7 @@ struct unique_benchmark : public primbench::benchmark_interface
 
     // not a warm-up run, we need to run once to determine the size of the output
     const auto new_end = thrust::unique_copy(policy(alloc).on(state.stream), in.cbegin(), in.cend(), out.begin());
-    const size_t unique_items = thrust::distance(out.begin(), new_end);
+    const size_t unique_items = _THRUST_STD::distance(out.begin(), new_end);
 
     state.set_items(m_items);
     state.add_reads<T>(m_items);

@@ -93,7 +93,7 @@ struct unique_benchmark : public primbench::benchmark_interface
       out_keys.begin(),
       out_vals.begin());
 
-    const size_t unique_elements = thrust::distance(out_keys.begin(), new_key_end);
+    const size_t unique_elements = _THRUST_STD::distance(out_keys.begin(), new_key_end);
 
     state.set_items(m_items);
     state.add_reads<T>(m_items);

@@ -45,6 +45,10 @@
 // rocRAND
 #include <rocrand/rocrand.h>
 
+// LIBHIPCXX
+#include _THRUST_STD_INCLUDE(cmath)
+#include _THRUST_STD_INCLUDE(functional)
+
 // STL
 #include <algorithm>
 #include <cstdint>
@@ -85,25 +89,29 @@ float get_entropy_percentage(int entropy_reduction)
   }
 }
 
+// Creates an interpolated value of type T between min (at = 0.0) and max (at = 1.0).
 template <typename T>
-T value_from_entropy(double percentage)
+[[nodiscard]] T lerp_min_max(double at) noexcept
 {
-  if (percentage == 100.0)
+  if (at == 1.0)
   {
-    return std::numeric_limits<T>::max();
+    return _THRUST_STD::numeric_limits<T>::max();
   }
-
-  percentage /= 100; // convert percentage to per one
-
-  // Select value from the line between the lowest and the highest representable
-  // values of type T based on the entropy value.
-  const auto max_val = static_cast<double>(std::numeric_limits<T>::max());
-  const auto min_val = static_cast<double>(std::numeric_limits<T>::lowest());
-  const auto result  = min_val + percentage * (max_val - min_val);
-  return static_cast<T>(result);
+  const auto min_val = static_cast<double>(_THRUST_STD::numeric_limits<T>::lowest());
+  const auto max_val = static_cast<double>(_THRUST_STD::numeric_limits<T>::max());
+  return static_cast<T>(_THRUST_STD::lerp(min_val, max_val, at));
 }
 
-const int entropy_reductions[] = {0, 2, 4, 6};
+template <class T>
+struct less_then_t
+{
+  T m_val;
+
+  [[nodiscard]] __device__ bool operator()(const T& val) const noexcept
+  {
+    return val < m_val;
+  }
+};
 
 namespace detail
 {
@@ -353,15 +361,15 @@ std::size_t gen_uniform_offsets(
     policy,
     segment_offsets.data(),
     segment_offsets.data() + segment_offsets.size(),
-    segment_offsets.data() /*, thrust::plus<>{}*/);
+    segment_offsets.data() /*, _THRUST_STD::plus<>{}*/);
 
   // Find first sum of offsets greater than 'elements', we are sure that there is
   // going to be one because we added elements + 1 at the end of the segment_offsets.
   auto iter = thrust::find_if(
     policy, segment_offsets.data(), segment_offsets.data() + segment_offsets.size(), geq_t<T>{elements});
 
   // Compute the element's index.
-  auto dist = thrust::distance(segment_offsets.data(), iter);
+  auto dist = _THRUST_STD::distance(segment_offsets.data(), iter);
   // Fill next item with 'elements'.
   thrust::fill_n(policy, segment_offsets.data() + dist, 1, elements);
   // Return next item's index.

@@ -1605,8 +1605,8 @@ Finally, a new TBB reduce_by_key implementation provides 80% faster
   - `thrust::stable_partition_copy with stencil`
   - `thrust::tabulate`
 - Memory Allocation:
-	- `thrust::malloc`
-	- `thrust::free`
+  - `thrust::malloc`
+  - `thrust::free`
   - `thrust::get_temporary_buffer`
   - `thrust::return_temporary_buffer`
 

@@ -2396,7 +2396,6 @@ PREDEFINED             = THRUST_DOXYGEN_INVOKED \
                          THRUST_DEVICE_SYSTEM=5 \
                          THRUST_STD_VER=2017 \
                          THRUST_GLOBAL_CONSTANT="inline constexpr" \
-                         THRUST_NODISCARD=[[nodiscard]] \
                          THRUST_MR_DEFAULT_ALIGNMENT=alignof(std::max_align_t) \
                          __cpp_lib_remove_cvref=201711 \
                          THRUST_PREVENT_MACRO_SUBSTITUTION= \

@@ -1,4 +1,4 @@
-// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2020-2026 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -36,7 +36,8 @@ int main(void)
   data[3] = 5;
 
   // add 10 to all values in data
-  thrust::transform(data.begin(), data.end(), thrust::constant_iterator<int>(10), data.begin(), thrust::plus<int>());
+  thrust::transform(
+    data.begin(), data.end(), thrust::constant_iterator<int>(10), data.begin(), _THRUST_STD::plus<int>());
 
   // data is now [13, 17, 12, 15]
 

@@ -19,7 +19,6 @@
 // SOFTWARE.
 
 #include <thrust/copy.h>
-#include <thrust/detail/libcxx_wrapper/std/__functional/identity.h>
 #include <thrust/device_vector.h>
 #include <thrust/functional.h>
 #include <thrust/iterator/counting_iterator.h>
@@ -52,7 +51,7 @@ int main()
   // compute indices of nonzero elements
   using IndexIterator = thrust::device_vector<int>::iterator;
 
-  IndexIterator indices_end = thrust::copy_if(first, last, stencil.begin(), indices.begin(), ::internal::identity{});
+  IndexIterator indices_end = thrust::copy_if(first, last, stencil.begin(), indices.begin(), _THRUST_STD::identity{});
   // indices now contains [1,2,5,7]
 
   // print result

@@ -1,4 +1,4 @@
-// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2020-2026 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -59,7 +59,7 @@ int main()
   cudaStreamCreate(&s);
 
   // launch a CUDA kernel with only 1 thread on our stream
-  reduce_kernel<<<1, 1, 0, s>>>(data.begin(), data.end(), 0, thrust::plus<int>(), result.data());
+  reduce_kernel<<<1, 1, 0, s>>>(data.begin(), data.end(), 0, _THRUST_STD::plus<int>(), result.data());
 
   // wait for the stream to finish
   cudaStreamSynchronize(s);
@@ -77,7 +77,7 @@ int main()
   auto begin        = data.begin();
   auto end          = data.end();
   unsigned int init = 0;
-  auto binary_op    = thrust::plus<unsigned int>();
+  auto binary_op    = _THRUST_STD::plus<unsigned int>();
 
   // std::async captures the algorithm parameters by value
   // use std::launch::async to ensure the creation of a new thread

@@ -60,7 +60,7 @@ public:
 
   __host__ __device__ difference_type size() const
   {
-    return thrust::distance(first, last);
+    return _THRUST_STD::distance(first, last);
   }
 
   __host__ __device__ reference operator[](difference_type n)
@@ -216,7 +216,7 @@ int main()
     make_range_view(thrust::make_transform_iterator(X.cbegin(), f1()), thrust::make_transform_iterator(X.cend(), f1())),
 
     // range view of normal_iterators
-    make_range_view(Y.begin(), thrust::distance(Y.begin(), Y.end())),
+    make_range_view(Y.begin(), _THRUST_STD::distance(Y.begin(), Y.end())),
 
     // range view of naked pointers
     make_range_view(Z.data().get(), 4));