ROCm 7 thrust namespace change adaptation

src/common/cuda_context.cuh

@@ -18,19 +18,19 @@ struct CUDAContext {
    */
   auto CTP() const {
 #if THRUST_MAJOR_VERSION >= 2
-    return thrust::cuda::par_nosync(caching_alloc_).on(dh::DefaultStream());
+    return thrust::system::hip::par_nosync(caching_alloc_).on(dh::DefaultStream());
 #else
-    return thrust::cuda::par(caching_alloc_).on(dh::DefaultStream());
+    return thrust::system::hip::par(caching_alloc_).on(dh::DefaultStream());
 #endif  // THRUST_MAJOR_VERSION >= 2
   }
   /**
    * \brief Thrust policy without caching allocator.
    */
   auto TP() const {
 #if THRUST_MAJOR_VERSION >= 2
-    return thrust::cuda::par_nosync(alloc_).on(dh::DefaultStream());
+    return thrust::system::hip::par_nosync(alloc_).on(dh::DefaultStream());
 #else
-    return thrust::cuda::par(alloc_).on(dh::DefaultStream());
+    return thrust::system::hip::par(alloc_).on(dh::DefaultStream());
 #endif  // THRUST_MAJOR_VERSION >= 2
   }
   auto Stream() const { return dh::DefaultStream(); }

src/common/device_helpers.hip.h

@@ -888,7 +888,7 @@ template <typename... Inputs,
               * = nullptr>
 size_t SegmentedUnique(Inputs &&...inputs) {
   dh::XGBCachingDeviceAllocator<char> alloc;
-  return SegmentedUnique(thrust::hip::par(alloc),
+  return SegmentedUnique(thrust::system::hip::par(alloc),
                          std::forward<Inputs &&>(inputs)...,
                          thrust::equal_to<size_t>{});
 }
@@ -988,7 +988,7 @@ void CopyIf(InIt in_first, InIt in_second, OutIt out_first, Predicate pred) {
   for (size_t offset = 0; offset < length; offset += kMaxCopySize) {
     auto begin_input = in_first + offset;
     auto end_input = in_first + std::min(offset + kMaxCopySize, length);
-    out_first = thrust::copy_if(thrust::hip::par(alloc), begin_input,
+    out_first = thrust::copy_if(thrust::system::hip::par(alloc), begin_input,
                                 end_input, out_first, pred);
   }
 }

src/common/stats.cuh

@@ -217,7 +217,7 @@ void SegmentedWeightedQuantile(Context const* ctx, AlphaIt alpha_it, SegIt seg_b
   auto scan_val = dh::MakeTransformIterator<float>(thrust::make_counting_iterator(0ul),
                                                    detail::WeightOp<WIter>{w_begin, d_sorted_idx});
 
-  thrust::inclusive_scan_by_key(thrust::cuda::par(caching), scan_key, scan_key + n_weights,
+  thrust::inclusive_scan_by_key(thrust::system::hip::par(caching), scan_key, scan_key + n_weights,
                                 scan_val, weights_cdf.begin());
 
   auto n_segments = std::distance(seg_beg, seg_end) - 1;

src/metric/elementwise_metric.cu

@@ -52,7 +52,7 @@ PackedReduceResult Reduce(Context const* ctx, MetaInfo const& info, Fn&& loss) {
     thrust::counting_iterator<size_t> begin(0);
     thrust::counting_iterator<size_t> end = begin + labels.Size();
     result = thrust::transform_reduce(
-        thrust::cuda::par(alloc), begin, end,
+        thrust::system::hip::par(alloc), begin, end,
         [=] XGBOOST_DEVICE(size_t i) {
           auto idx = linalg::UnravelIndex(i, labels.Shape());
           auto sample_id = std::get<0>(idx);