Get rid of the deprecated thrust::identity (#17942)

This PR removes the usage of the deprecated `thrust::identity` and replaces it with `cuda::std::identity` where applicable. In cases where `thrust::identity` was used as a type conversion callable and `cuda::std::identity` is not suitable, it uses a new utility called `cast_fn` in the `detail` namespace. Authors: - Yunsong Wang (https://github.com/PointKernel) Approvers: - Michael Schellenberger Costa (https://github.com/miscco) - David Wendt (https://github.com/davidwendt) - MithunR (https://github.com/mythrocks) URL: #17942
rapidsai · Feb 10, 2025 · 2d0c85c · 2d0c85c
1 parent 10cd7c4
commit 2d0c85c
Show file tree

Hide file tree

Showing 24 changed files with 104 additions and 79 deletions.
diff --git a/cpp/benchmarks/common/generate_input.cu b/cpp/benchmarks/common/generate_input.cu
@@ -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.
@@ -46,7 +46,6 @@
 #include <thrust/execution_policy.h>
 #include <thrust/fill.h>
 #include <thrust/for_each.h>
-#include <thrust/functional.h>
 #include <thrust/gather.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
@@ -511,7 +510,7 @@ std::unique_ptr<cudf::column> create_random_column(data_profile const& profile,
   auto [result_bitmask, null_count] =
     cudf::detail::valid_if(null_mask.begin(),
                            null_mask.end(),
-                           thrust::identity<bool>{},
+                           cuda::std::identity{},
                            cudf::get_default_stream(),
                            cudf::get_current_device_resource_ref());
 
@@ -600,7 +599,7 @@ std::unique_ptr<cudf::column> create_random_utf8_string_column(data_profile cons
   auto [result_bitmask, null_count] =
     profile.get_null_probability().has_value()
       ? cudf::detail::valid_if(
-          null_mask.begin(), null_mask.end() - 1, thrust::identity<bool>{}, stream, mr)
+          null_mask.begin(), null_mask.end() - 1, cuda::std::identity{}, stream, mr)
       : std::pair{rmm::device_buffer{}, 0};
 
   return cudf::make_strings_column(
@@ -693,7 +692,7 @@ std::unique_ptr<cudf::column> create_random_column<cudf::struct_view>(data_profi
           auto valids = valid_dist(engine, num_rows);
           return cudf::detail::valid_if(valids.begin(),
                                         valids.end(),
-                                        thrust::identity<bool>{},
+                                        cuda::std::identity{},
                                         cudf::get_default_stream(),
                                         cudf::get_current_device_resource_ref());
         }
@@ -787,7 +786,7 @@ std::unique_ptr<cudf::column> create_random_column<cudf::list_view>(data_profile
 
     auto [null_mask, null_count] = cudf::detail::valid_if(valids.begin(),
                                                           valids.end(),
-                                                          thrust::identity<bool>{},
+                                                          cuda::std::identity{},
                                                           cudf::get_default_stream(),
                                                           cudf::get_current_device_resource_ref());
     list_column                  = cudf::make_lists_column(

diff --git a/cpp/benchmarks/join/join_common.hpp b/cpp/benchmarks/join/join_common.hpp
@@ -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.
@@ -31,7 +31,7 @@
 #include <cudf/utilities/error.hpp>
 #include <cudf/utilities/memory_resource.hpp>
 
-#include <thrust/functional.h>
+#include <cuda/std/functional>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/random/linear_congruential_engine.h>
@@ -85,7 +85,7 @@ void BM_join(state_type& state, Join JoinFunc)
       thrust::make_transform_iterator(thrust::make_counting_iterator(0), null75_generator{});
     return cudf::detail::valid_if(validity,
                                   validity + size,
-                                  thrust::identity<bool>{},
+                                  cuda::std::identity{},
                                   cudf::get_default_stream(),
                                   cudf::get_current_device_resource_ref());
   };

diff --git a/cpp/include/cudf/detail/utilities/cast_functor.cuh b/cpp/include/cudf/detail/utilities/cast_functor.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023, 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.
@@ -31,6 +31,23 @@
 namespace cudf {
 namespace detail {
 
+/**
+ * @brief Functor that casts a primitive input value to a specified type
+ */
+template <typename T>
+struct cast_fn {
+  template <typename U>
+  CUDF_HOST_DEVICE constexpr T operator()(U&& val) const
+  {
+    return static_cast<T>(cuda::std::forward<U>(val));
+  }
+
+  CUDF_HOST_DEVICE constexpr T&& operator()(T&& val) const noexcept
+  {
+    return cuda::std::forward<T>(val);
+  }
+};
+
 /**
  * @brief Functor that casts another functor's result to a specified type.
  *

diff --git a/cpp/include/cudf/reduction/detail/reduction_operators.cuh b/cpp/include/cudf/reduction/detail/reduction_operators.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.
@@ -17,12 +17,11 @@
 #pragma once
 
 #include <cudf/detail/iterator.cuh>
+#include <cudf/detail/utilities/cast_functor.cuh>
 #include <cudf/detail/utilities/device_operators.cuh>
 #include <cudf/detail/utilities/transform_unary_functions.cuh>
 #include <cudf/types.hpp>  //for CUDF_HOST_DEVICE
 
-#include <thrust/functional.h>
-
 #include <cmath>
 
 namespace cudf {
@@ -156,15 +155,15 @@ struct sum : public simple_op<sum> {
   using op = cudf::DeviceSum;
 
   template <typename ResultType>
-  using transformer = thrust::identity<ResultType>;
+  using transformer = cudf::detail::cast_fn<ResultType>;
 };
 
 // operator for `product`
 struct product : public simple_op<product> {
   using op = cudf::DeviceProduct;
 
   template <typename ResultType>
-  using transformer = thrust::identity<ResultType>;
+  using transformer = cudf::detail::cast_fn<ResultType>;
 };
 
 // operator for `sum_of_squares`
@@ -180,15 +179,15 @@ struct min : public simple_op<min> {
   using op = cudf::DeviceMin;
 
   template <typename ResultType>
-  using transformer = thrust::identity<ResultType>;
+  using transformer = cudf::detail::cast_fn<ResultType>;
 };
 
 // operator for `max`
 struct max : public simple_op<max> {
   using op = cudf::DeviceMax;
 
   template <typename ResultType>
-  using transformer = thrust::identity<ResultType>;
+  using transformer = cudf::detail::cast_fn<ResultType>;
 };
 
 /**
@@ -246,7 +245,7 @@ struct mean : public compound_op<mean> {
   using op = cudf::DeviceSum;
 
   template <typename ResultType>
-  using transformer = thrust::identity<ResultType>;
+  using transformer = cudf::detail::cast_fn<ResultType>;
 
   template <typename ResultType>
   struct intermediate {

diff --git a/cpp/include/cudf_test/column_wrapper.hpp b/cpp/include/cudf_test/column_wrapper.hpp
@@ -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.
@@ -39,8 +39,8 @@
 
 #include <rmm/device_buffer.hpp>
 
+#include <cuda/std/functional>
 #include <thrust/copy.h>
-#include <thrust/functional.h>
 #include <thrust/host_vector.h>
 #include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/counting_iterator.h>
@@ -1645,8 +1645,11 @@ class lists_column_wrapper : public detail::column_wrapper {
 
     // concatenate them together, skipping children that are null.
     std::vector<column_view> children;
-    thrust::copy_if(
-      std::cbegin(cols), std::cend(cols), valids, std::back_inserter(children), thrust::identity{});
+    thrust::copy_if(std::cbegin(cols),
+                    std::cend(cols),
+                    valids,
+                    std::back_inserter(children),
+                    cuda::std::identity{});
 
     auto data = children.empty() ? cudf::empty_like(expected_hierarchy)
                                  : cudf::concatenate(children,

diff --git a/cpp/src/groupby/sort/group_merge_m2.cu b/cpp/src/groupby/sort/group_merge_m2.cu
@@ -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.
@@ -27,7 +27,7 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
 
-#include <thrust/functional.h>
+#include <cuda/std/functional>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
@@ -180,7 +180,7 @@ std::unique_ptr<column> group_merge_m2(column_view const& values,
   // Generate bitmask for the output.
   // Only mean and M2 values can be nullable. Count column must be non-nullable.
   auto [null_mask, null_count] =
-    cudf::detail::valid_if(validities.begin(), validities.end(), thrust::identity{}, stream, mr);
+    cudf::detail::valid_if(validities.begin(), validities.end(), cuda::std::identity{}, stream, mr);
   if (null_count > 0) {
     result_means->set_null_mask(null_mask, null_count, stream);   // copy null_mask
     result_M2s->set_null_mask(std::move(null_mask), null_count);  // take over null_mask

diff --git a/cpp/src/groupby/sort/group_scan_util.cuh b/cpp/src/groupby/sort/group_scan_util.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.
@@ -27,6 +27,7 @@
 #include <cudf/detail/iterator.cuh>
 #include <cudf/detail/null_mask.hpp>
 #include <cudf/detail/structs/utilities.hpp>
+#include <cudf/detail/utilities/cast_functor.cuh>
 #include <cudf/table/table_device_view.cuh>
 #include <cudf/types.hpp>
 #include <cudf/utilities/memory_resource.hpp>
@@ -36,7 +37,6 @@
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 
-#include <thrust/functional.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/scan.h>
@@ -129,12 +129,12 @@ struct group_scan_functor<K, T, std::enable_if_t<is_group_scan_supported<K, T>()
     if (values.has_nulls()) {
       auto input = thrust::make_transform_iterator(
         make_null_replacement_iterator(*values_view, OpType::template identity<DeviceType>()),
-        thrust::identity<ResultDeviceType>{});
+        cudf::detail::cast_fn<ResultDeviceType>{});
       do_scan(input, result_view->begin<ResultDeviceType>(), OpType{});
       result->set_null_mask(cudf::detail::copy_bitmask(values, stream, mr), values.null_count());
     } else {
       auto input = thrust::make_transform_iterator(values_view->begin<DeviceType>(),
-                                                   thrust::identity<ResultDeviceType>{});
+                                                   cudf::detail::cast_fn<ResultDeviceType>{});
       do_scan(input, result_view->begin<ResultDeviceType>(), OpType{});
     }
     return result;

diff --git a/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh b/cpp/src/groupby/sort/group_single_pass_reduction_util.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.
@@ -33,7 +33,7 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
 
-#include <thrust/functional.h>
+#include <cuda/std/functional>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/discard_iterator.h>
 #include <thrust/reduce.h>
@@ -204,7 +204,7 @@ struct group_reduction_functor<
                    thrust::logical_or{});
 
       auto [null_mask, null_count] =
-        cudf::detail::valid_if(validity.begin(), validity.end(), thrust::identity{}, stream, mr);
+        cudf::detail::valid_if(validity.begin(), validity.end(), cuda::std::identity{}, stream, mr);
       result->set_null_mask(std::move(null_mask), null_count);
     }
     return result;
@@ -257,7 +257,7 @@ struct group_reduction_functor<
                    thrust::logical_or{});
 
       auto [null_mask, null_count] =
-        cudf::detail::valid_if(validity.begin(), validity.end(), thrust::identity{}, stream, mr);
+        cudf::detail::valid_if(validity.begin(), validity.end(), cuda::std::identity{}, stream, mr);
       result->set_null_mask(std::move(null_mask), null_count);
     }
 

diff --git a/cpp/src/groupby/sort/sort_helper.cu b/cpp/src/groupby/sort/sort_helper.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.
@@ -165,7 +165,7 @@ sort_groupby_helper::index_vector const& sort_groupby_helper::group_offsets(
                                  itr + size,
                                  result.begin(),
                                  group_offsets->begin(),
-                                 thrust::identity<bool>{});
+                                 cuda::std::identity{});
   } else {
     auto const d_key_equal = comparator.equal_to<false>(
       cudf::nullate::DYNAMIC{cudf::has_nested_nulls(_keys)}, null_equality::EQUAL);

diff --git a/cpp/src/io/json/json_normalization.cu b/cpp/src/io/json/json_normalization.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2024, NVIDIA CORPORATION.
+ * Copyright (c) 2024-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 <cub/device/device_copy.cuh>
 #include <cuda/atomic>
+#include <cuda/std/functional>
 #include <thrust/binary_search.h>
 #include <thrust/distance.h>
 #include <thrust/gather.h>
@@ -442,7 +443,7 @@ std::
                     inbuf.begin(),
                     inbuf.end(),
                     stencil.begin(),
-                    thrust::identity<int>());
+                    cuda::std::identity{});
   inbuf.resize(inbuf_size - num_deletions, stream);
 
   thrust::exclusive_scan(rmm::exec_policy_nosync(stream),

diff --git a/cpp/src/join/hash_join.cu b/cpp/src/join/hash_join.cu
@@ -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.
@@ -31,8 +31,8 @@
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/std/functional>
 #include <thrust/count.h>
-#include <thrust/functional.h>
 #include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
@@ -357,7 +357,7 @@ std::size_t get_full_join_size(
     left_join_complement_size = thrust::count_if(rmm::exec_policy(stream),
                                                  invalid_index_map->begin(),
                                                  invalid_index_map->end(),
-                                                 thrust::identity());
+                                                 cuda::std::identity());
   }
   return join_size + left_join_complement_size;
 }

diff --git a/cpp/src/join/join_utils.cu b/cpp/src/join/join_utils.cu
@@ -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.
@@ -20,8 +20,8 @@
 
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/std/functional>
 #include <thrust/copy.h>
-#include <thrust/functional.h>
 #include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/scatter.h>
@@ -141,7 +141,7 @@ get_left_join_indices_complement(std::unique_ptr<rmm::device_uvector<size_type>>
                                               thrust::make_counting_iterator(end_counter),
                                               invalid_index_map->begin(),
                                               right_indices_complement->begin(),
-                                              thrust::identity{}) -
+                                              cuda::std::identity{}) -
                               right_indices_complement->begin();
     right_indices_complement->resize(indices_count, stream);
   }