Drop deprecated thrust features and replace with libcu++ ones (#6248)

CCCL is deprecating a lot of legacy thrust features, so replace them with the standard ones from libc++ or just drop them altogether Authors: - Michael Schellenberger Costa (https://github.com/miscco) Approvers: - Dante Gama Dessavre (https://github.com/dantegd) URL: #6248
rapidsai · Jan 27, 2025 · 66abfc5 · 66abfc5
1 parent 3a32228
commit 66abfc5
Show file tree

Hide file tree

Showing 14 changed files with 38 additions and 44 deletions.
diff --git a/cpp/src/dbscan/vertexdeg/algo.cuh b/cpp/src/dbscan/vertexdeg/algo.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2018-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2018-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -46,7 +46,7 @@ namespace VertexDeg {
 namespace Algo {
 
 template <typename index_t = int>
-struct column_counter : public thrust::unary_function<index_t, index_t> {
+struct column_counter {
   index_t* ia_;
   index_t n_;
 

diff --git a/cpp/src/glm/qn/glm_base.cuh b/cpp/src/glm/qn/glm_base.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2018-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2018-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -26,6 +26,7 @@
 #include <raft/util/cuda_utils.cuh>
 #include <raft/util/cudart_utils.hpp>
 
+#include <cuda/std/functional>
 #include <thrust/execution_policy.h>
 #include <thrust/functional.h>
 #include <thrust/reduce.h>
@@ -124,7 +125,7 @@ struct GLMBase : GLMDims {
                                        sample_weights,
                                        sample_weights + n_samples,
                                        (T)0,
-                                       thrust::plus<T>());
+                                       cuda::std::plus<T>());
   }
 
   /*

diff --git a/cpp/src/hdbscan/condensed_hierarchy.cu b/cpp/src/hdbscan/condensed_hierarchy.cu
@@ -27,7 +27,6 @@
 
 #include <cub/cub.cuh>
 #include <cuda/functional>
-#include <cuda/std/functional>
 #include <thrust/copy.h>
 #include <thrust/device_ptr.h>
 #include <thrust/execution_policy.h>
@@ -161,7 +160,7 @@ void CondensedHierarchy<value_idx, value_t>::condense(value_idx* full_parents,
     cuda::proclaim_return_type<value_idx>(
       [=] __device__(value_idx a) -> value_idx { return static_cast<value_idx>(a != -1); }),
     static_cast<value_idx>(0),
-    thrust::plus<value_idx>());
+    cuda::std::plus<value_idx>());
 
   parents.resize(n_edges, stream);
   children.resize(n_edges, stream);

diff --git a/cpp/src/hdbscan/detail/membership.cuh b/cpp/src/hdbscan/detail/membership.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -80,8 +80,7 @@ void get_probabilities(const raft::handle_t& handle,
                              int,
                              const value_idx*,
                              const value_idx*,
-                             cudaStream_t,
-                             bool) =
+                             cudaStream_t) =
     cub::DeviceSegmentedReduce::Max<const value_t*, value_t*, const value_idx*, const value_idx*>;
   Utils::cub_segmented_reduce(
     lambdas, deaths.data(), n_clusters, sorted_parents_offsets.data(), stream, reduce_func);

diff --git a/cpp/src/hdbscan/detail/select.cuh b/cpp/src/hdbscan/detail/select.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -30,6 +30,7 @@
 #include <rmm/exec_policy.hpp>
 
 #include <cub/cub.cuh>
+#include <cuda/std/functional>
 #include <thrust/copy.h>
 #include <thrust/execution_policy.h>
 #include <thrust/fill.h>
@@ -223,7 +224,7 @@ void excess_of_mass(const raft::handle_t& handle,
                                  cuda::proclaim_return_type<value_t>(
                                    [=] __device__(value_idx a) -> value_t { return stability[a]; }),
                                  0.0,
-                                 thrust::plus<value_t>());
+                                 cuda::std::plus<value_t>());
     }
 
     if (subtree_stability > node_stability || cluster_sizes_h[node] > max_cluster_size) {

diff --git a/cpp/src/hdbscan/detail/stabilities.cuh b/cpp/src/hdbscan/detail/stabilities.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -104,8 +104,7 @@ void compute_stabilities(const raft::handle_t& handle,
                              int,
                              const value_idx*,
                              const value_idx*,
-                             cudaStream_t,
-                             bool) =
+                             cudaStream_t) =
     cub::DeviceSegmentedReduce::Min<const value_t*, value_t*, const value_idx*, const value_idx*>;
   Utils::cub_segmented_reduce(lambdas,
                               births_parent_min.data() + 1,

diff --git a/cpp/src/hdbscan/detail/utils.h b/cpp/src/hdbscan/detail/utils.h
@@ -76,19 +76,11 @@ void cub_segmented_reduce(const value_t* in,
 {
   rmm::device_uvector<char> d_temp_storage(0, stream);
   size_t temp_storage_bytes = 0;
-  cub_reduce_func(
-    nullptr, temp_storage_bytes, in, out, n_segments, offsets, offsets + 1, stream, false);
+  cub_reduce_func(nullptr, temp_storage_bytes, in, out, n_segments, offsets, offsets + 1, stream);
   d_temp_storage.resize(temp_storage_bytes, stream);
 
-  cub_reduce_func(d_temp_storage.data(),
-                  temp_storage_bytes,
-                  in,
-                  out,
-                  n_segments,
-                  offsets,
-                  offsets + 1,
-                  stream,
-                  false);
+  cub_reduce_func(
+    d_temp_storage.data(), temp_storage_bytes, in, out, n_segments, offsets, offsets + 1, stream);
 }
 
 /**
@@ -118,7 +110,7 @@ Common::CondensedHierarchy<value_idx, value_t> make_cluster_tree(
     cuda::proclaim_return_type<value_idx>(
       [=] __device__(value_idx a) -> value_idx { return static_cast<value_idx>(a > 1); }),
     static_cast<value_idx>(0),
-    thrust::plus<value_idx>());
+    cuda::std::plus<value_idx>());
 
   // remove leaves from condensed tree
   rmm::device_uvector<value_idx> cluster_parents(cluster_tree_edges, stream);

diff --git a/cpp/src/hdbscan/prediction_data.cu b/cpp/src/hdbscan/prediction_data.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -130,7 +130,7 @@ void generate_prediction_data(const raft::handle_t& handle,
 
   // this is to find maximum lambdas of all children under a parent
   cudaError_t (*reduce_func)(
-    void*, size_t&, const float*, float*, int, const int*, const int*, cudaStream_t, bool) =
+    void*, size_t&, const float*, float*, int, const int*, const int*, cudaStream_t) =
     cub::DeviceSegmentedReduce::Max<const float*, float*, const int*, const int*>;
   detail::Utils::cub_segmented_reduce(lambdas,
                                       prediction_data.get_deaths(),

diff --git a/cpp/src/svm/kernelcache.cuh b/cpp/src/svm/kernelcache.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -81,7 +81,7 @@ CUML_KERNEL void mapColumnIndicesToSVRSpace(
 }
 
 template <typename math_t>
-struct select_at_index : public thrust::unary_function<int, math_t> {
+struct select_at_index {
   const math_t* dot_;
   select_at_index(const math_t* dot) : dot_(dot) {}
 

diff --git a/cpp/src/svm/sparse_util.cuh b/cpp/src/svm/sparse_util.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -26,6 +26,7 @@
 
 #include <rmm/device_uvector.hpp>
 
+#include <cuda/std/functional>
 #include <thrust/device_ptr.h>
 #include <thrust/execution_policy.h>
 #include <thrust/iterator/constant_iterator.h>
@@ -325,7 +326,7 @@ raft::device_csr_matrix_view<math_t, int, int, int> getMatrixBatch(
                         inptr_src + batch_size + 1,
                         thrust::make_constant_iterator(nnz_offset),
                         inptr_tgt,
-                        thrust::minus<int>());
+                        cuda::std::minus<int>());
     }
 
     auto csr_struct_out = raft::make_device_compressed_structure_view<int, int, int>(
@@ -496,7 +497,7 @@ static void copySparseRowsToDense(const int* indptr,
   RAFT_CUDA_TRY(cudaPeekAtLastError());
 }
 
-struct rowsize : public thrust::unary_function<int, int> {
+struct rowsize {
   const int* indptr_;
   rowsize(const int* indptr) : indptr_(indptr) {}
 
@@ -610,7 +611,7 @@ int computeIndptrForSubset(
                                    row_new_indices_ptr + num_indices,
                                    row_sizes_ptr + 1,
                                    rowsize(indptr_in),
-                                   thrust::plus<int>());
+                                   cuda::std::plus<int>());
 
   // retrieve nnz from indptr_in[num_indices]
   int nnz;

diff --git a/cpp/src/tsa/auto_arima.cuh b/cpp/src/tsa/auto_arima.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -181,7 +181,7 @@ inline void divide_by_mask_execute(const DataT* d_in,
  * a matrix from its index. This makes possible a 2d scan with thrust.
  * Found in thrust/examples/scan_matrix_by_rows.cu
  */
-struct which_col : thrust::unary_function<int, int> {
+struct which_col {
   MLCommon::FastIntDiv divisor;
   __host__ which_col(int col_length) : divisor(col_length) {}
   __host__ __device__ int operator()(int idx) const { return idx / divisor; }

diff --git a/cpp/src/tsne/fft_tsne.cuh b/cpp/src/tsne/fft_tsne.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -35,6 +35,7 @@
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
 
+#include <cuda/std/functional>
 #include <thrust/device_ptr.h>
 #include <thrust/fill.h>
 #include <thrust/functional.h>
@@ -512,7 +513,7 @@ value_t FFT_TSNE(value_t* VAL,
                                     norm_vec_thrust,
                                     norm_vec_thrust + normalization_vec_device.size(),
                                     0.0f,
-                                    thrust::plus<value_t>());
+                                    cuda::std::plus<value_t>());
       normalization = sumQ - n;
     }
 
@@ -565,7 +566,7 @@ value_t FFT_TSNE(value_t* VAL,
                       att_forces_thrust + n,
                       att_forces_thrust + n,
                       att_forces_thrust,
-                      thrust::plus<value_t>());
+                      cuda::std::plus<value_t>());
 
     thrust::transform(thrust_policy,
                       att_forces_thrust,
@@ -577,7 +578,7 @@ value_t FFT_TSNE(value_t* VAL,
                                        att_forces_thrust,
                                        att_forces_thrust + attractive_forces_device.size(),
                                        0.0f,
-                                       thrust::plus<value_t>()) /
+                                       cuda::std::plus<value_t>()) /
                         attractive_forces_device.size();
 
     if (grad_norm <= params.min_grad_norm) {

diff --git a/cpp/src_prims/timeSeries/stationarity.cuh b/cpp/src_prims/timeSeries/stationarity.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -174,7 +174,7 @@ CUML_KERNEL void kpss_stationarity_check_kernel(bool* results,
  * Found in thrust/examples/scan_matrix_by_rows.cu
  */
 template <typename IdxT>
-struct which_col : thrust::unary_function<IdxT, IdxT> {
+struct which_col {
   IdxT col_length;
   __host__ __device__ which_col(IdxT col_length_) : col_length(col_length_) {}
   __host__ __device__ IdxT operator()(IdxT idx) const { return idx / col_length; }

diff --git a/cpp/tests/sg/rf_test.cu b/cpp/tests/sg/rf_test.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -25,6 +25,7 @@
 #include <raft/util/cuda_utils.cuh>
 #include <raft/util/cudart_utils.hpp>
 
+#include <cuda/std/functional>
 #include <thrust/binary_search.h>
 #include <thrust/copy.h>
 #include <thrust/device_vector.h>
@@ -305,7 +306,7 @@ class RfSpecialisedTest {
       thrust::device_vector<double> normal(params.n_rows);
       r.normal(normal.data().get(), normal.size(), 0.0, 2.0, nullptr);
       thrust::transform(
-        normal.begin(), normal.end(), y_temp.begin(), y.begin(), thrust::plus<LabelT>());
+        normal.begin(), normal.end(), y_temp.begin(), y.begin(), cuda::std::plus<LabelT>());
     }
     raft::linalg::transpose(
       handle, X.data().get(), X_transpose.data().get(), params.n_rows, params.n_cols, nullptr);