Replace the deprecated thrust::identity with cuda::std::identity

include/cuco/detail/bloom_filter/bloom_filter_impl.cuh

@@ -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.
@@ -29,10 +29,10 @@
 #include <cuda/std/__algorithm/min.h>  // TODO #include <cuda/std/algorithm> once available
 #include <cuda/std/array>
 #include <cuda/std/bit>
+#include <cuda/std/functional>
 #include <cuda/std/tuple>
 #include <cuda/std/type_traits>
 #include <cuda/stream_ref>
-#include <thrust/functional.h>
 #include <thrust/iterator/constant_iterator.h>
 
 #include <cstdint>
@@ -182,7 +182,7 @@ class bloom_filter_impl {
         stream.get()));
     } else {
       auto const always_true = thrust::constant_iterator<bool>{true};
-      this->add_if_async(first, last, always_true, thrust::identity{}, stream);
+      this->add_if_async(first, last, always_true, cuda::std::identity{}, stream);
     }
   }
 
@@ -283,7 +283,7 @@ class bloom_filter_impl {
                                          cuda::stream_ref stream) const noexcept
   {
     auto const always_true = thrust::constant_iterator<bool>{true};
-    this->contains_if_async(first, last, always_true, thrust::identity{}, output_begin, stream);
+    this->contains_if_async(first, last, always_true, cuda::std::identity{}, output_begin, stream);
   }
 
   template <class InputIt, class StencilIt, class Predicate, class OutputIt>

include/cuco/detail/hash_functions/identity_hash.cuh

@@ -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.
@@ -16,8 +16,8 @@
 
 #pragma once
 
+#include <cuda/std/functional>
 #include <cuda/std/type_traits>
-#include <thrust/functional.h>
 
 namespace cuco::detail {
 
@@ -33,7 +33,7 @@ namespace cuco::detail {
  * @tparam Key The type of the values to hash
  */
 template <typename Key>
-struct identity_hash : private thrust::identity<Key> {
+struct identity_hash : private cuda::std::identity {
   using argument_type = Key;  ///< The type of the values taken as argument
   /// The type of the hash values produced
   using result_type = cuda::std::conditional_t<sizeof(Key) <= 4, uint32_t, uint64_t>;
@@ -49,7 +49,7 @@ struct identity_hash : private thrust::identity<Key> {
    */
   __host__ __device__ result_type operator()(Key const& x) const
   {
-    return static_cast<result_type>(thrust::identity<Key>::operator()(x));
+    return static_cast<result_type>(cuda::std::identity::operator()(x));
   }
 };  // identity_hash
 

include/cuco/detail/open_addressing/open_addressing_impl.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.
@@ -31,6 +31,7 @@
 #include <cub/device/device_for.cuh>
 #include <cub/device/device_select.cuh>
 #include <cuda/atomic>
+#include <cuda/std/functional>
 #include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
@@ -273,7 +274,7 @@ class open_addressing_impl {
   size_type insert(InputIt first, InputIt last, Ref container_ref, cuda::stream_ref stream)
   {
     auto const always_true = thrust::constant_iterator<bool>{true};
-    return this->insert_if(first, last, always_true, thrust::identity{}, container_ref, stream);
+    return this->insert_if(first, last, always_true, cuda::std::identity{}, container_ref, stream);
   }
 
   /**
@@ -296,7 +297,7 @@ class open_addressing_impl {
                     cuda::stream_ref stream) noexcept
   {
     auto const always_true = thrust::constant_iterator<bool>{true};
-    this->insert_if_async(first, last, always_true, thrust::identity{}, container_ref, stream);
+    this->insert_if_async(first, last, always_true, cuda::std::identity{}, container_ref, stream);
   }
 
   /**
@@ -494,7 +495,7 @@ class open_addressing_impl {
   {
     auto const always_true = thrust::constant_iterator<bool>{true};
     this->contains_if_async(
-      first, last, always_true, thrust::identity{}, output_begin, container_ref, stream);
+      first, last, always_true, cuda::std::identity{}, output_begin, container_ref, stream);
   }
 
   /**
@@ -569,7 +570,7 @@ class open_addressing_impl {
     auto const always_true = thrust::constant_iterator<bool>{true};
 
     this->find_if_async(
-      first, last, always_true, thrust::identity{}, output_begin, container_ref, stream);
+      first, last, always_true, cuda::std::identity{}, output_begin, container_ref, stream);
   }
 
   /**

include/cuco/detail/open_addressing/open_addressing_ref_impl.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.
@@ -24,6 +24,7 @@
 #include <cuco/probing_scheme.cuh>
 
 #include <cuda/atomic>
+#include <cuda/std/functional>
 #include <cuda/std/type_traits>
 #include <thrust/distance.h>
 #include <thrust/execution_policy.h>
@@ -1246,7 +1247,7 @@ class open_addressing_ref_impl {
 
             // Fill the buffer if any matching keys are found
             auto const lane_id = probing_tile.thread_rank();
-            if (thrust::any_of(thrust::seq, exists, exists + bucket_size, thrust::identity{})) {
+            if (thrust::any_of(thrust::seq, exists, exists + bucket_size, cuda::std::identity{})) {
               if constexpr (IsOuter) { found_match = true; }
 
               int32_t num_matches[bucket_size];