Merge remote-tracking branch 'upstream' into worktree-fix-docs

Author: sleeepyjack

README.md

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023-2025, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2026, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -21,8 +21,7 @@
 
 #include <nvbench/nvbench.cuh>
 
-#include <thrust/iterator/iterator_traits.h>
-#include <thrust/iterator/tabulate_output_iterator.h>
+#include <cuda/iterator>
 
 #include <nv/target>
 
@@ -75,13 +74,13 @@ struct lazy_discard {
 };
 
 /**
- * @brief An output iterator similar to `thrust::discard_iterator` but prevents the write from being
+ * @brief An output iterator similar to `cuda::discard_iterator` but prevents the write from being
  * optimized out by the compiler.
  */
 template <class OutputIt>
 auto make_lazy_discard_iterator(OutputIt it)
 {
-  return thrust::tabulate_output_iterator(lazy_discard<OutputIt>{it});
+  return cuda::make_tabulate_output_iterator(lazy_discard<OutputIt>{it});
 }
 
 }  // namespace cuco::benchmark

benchmarks/benchmark_utils.hpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2024-2025, NVIDIA CORPORATION.
+ * Copyright (c) 2024-2026, 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,8 +24,8 @@
 
 #include <nvbench/nvbench.cuh>
 
+#include <cuda/iterator>
 #include <cuda/std/limits>
-#include <thrust/iterator/counting_iterator.h>
 
 #include <cstdint>
 #include <exception>
@@ -67,7 +67,7 @@ void bloom_filter_add(nvbench::state& state,
     state.skip("num_sub_filters too large for size_type");  // skip invalid configurations
   }
 
-  thrust::counting_iterator<Key> keys(0);
+  cuda::counting_iterator<Key> keys(0);
 
   state.add_element_count(num_keys);
 
@@ -108,7 +108,7 @@ void arrow_bloom_filter_add(nvbench::state& state, nvbench::type_list<Key, Dist>
                                                                                  // configurations
   }
 
-  thrust::counting_iterator<Key> keys(0);
+  cuda::counting_iterator<Key> keys(0);
 
   state.add_element_count(num_keys);
 

benchmarks/bloom_filter/add_bench.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2024-2025, NVIDIA CORPORATION.
+ * Copyright (c) 2024-2026, 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,9 +24,9 @@
 
 #include <nvbench/nvbench.cuh>
 
+#include <cuda/iterator>
 #include <cuda/std/limits>
 #include <thrust/device_vector.h>
-#include <thrust/iterator/counting_iterator.h>
 
 #include <exception>
 #include <limits>
@@ -73,7 +73,7 @@ void bloom_filter_contains(
     state.skip("num_sub_filters too large for size_type");  // skip invalid configurations
   }
 
-  thrust::counting_iterator<Key> keys(0);
+  cuda::counting_iterator<Key> keys(0);
   thrust::device_vector<bool> result(num_keys, false);
 
   state.add_element_count(num_keys);
@@ -119,7 +119,7 @@ void arrow_bloom_filter_contains(nvbench::state& state, nvbench::type_list<Key,
                                                                                  // configurations
   }
 
-  thrust::counting_iterator<Key> keys(0);
+  cuda::counting_iterator<Key> keys(0);
   thrust::device_vector<bool> result(num_keys, false);
 
   state.add_element_count(num_keys);

benchmarks/bloom_filter/contains_bench.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2024-2025, NVIDIA CORPORATION.
+ * Copyright (c) 2024-2026, 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,8 +24,8 @@
 #include <nvbench/nvbench.cuh>
 
 #include <cuda/functional>
+#include <cuda/iterator>
 #include <thrust/device_vector.h>
-#include <thrust/iterator/transform_iterator.h>
 
 #include <cmath>
 #include <cstddef>
@@ -40,7 +40,7 @@ template <typename InputIt>
   // Casting is valid since the keys generated are representable in int64_t.
   using T = std::int64_t;
 
-  auto cast_iter = thrust::make_transform_iterator(
+  auto cast_iter = cuda::make_transform_iterator(
     first, cuda::proclaim_return_type<T>([] __device__(auto i) { return static_cast<T>(i); }));
 
   auto set = cuco::static_set{n, 0.8, cuco::empty_key<T>{-1}};

benchmarks/hyperloglog/hyperloglog_bench.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2024-2025, NVIDIA CORPORATION.
+ * Copyright (c) 2024-2026, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -22,6 +22,7 @@
 
 #include <nvbench/nvbench.cuh>
 
+#include <cuda/iterator>
 #include <thrust/device_vector.h>
 #include <thrust/transform.h>
 
@@ -54,7 +55,7 @@ void static_multiset_retrieve(nvbench::state& state, nvbench::type_list<Key, Dis
 
   auto const output_size = set.count(keys.begin(), keys.end());
   thrust::device_vector<Key> output_match(output_size);
-  auto output_probe_begin = thrust::discard_iterator{};
+  auto output_probe_begin = cuda::discard_iterator{};
 
   state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
     set.retrieve(

benchmarks/static_multiset/retrieve_bench.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2024-2025, NVIDIA CORPORATION.
+ * Copyright (c) 2024-2026, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -22,6 +22,7 @@
 
 #include <nvbench/nvbench.cuh>
 
+#include <cuda/iterator>
 #include <thrust/device_vector.h>
 #include <thrust/transform.h>
 
@@ -54,7 +55,7 @@ void static_set_retrieve(nvbench::state& state, nvbench::type_list<Key, Dist>)
 
   auto const output_size = set.count(keys.begin(), keys.end());
   thrust::device_vector<Key> output_match(output_size);
-  auto output_probe_begin = thrust::discard_iterator{};
+  auto output_probe_begin = cuda::discard_iterator{};
 
   state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
     set.retrieve(

benchmarks/static_set/retrieve_bench.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023-2025, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2026, 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,6 +17,7 @@
 #include <cuco/static_map.cuh>
 
 #include <cub/block/block_reduce.cuh>
+#include <cuda/iterator>
 #include <cuda/std/atomic>
 #include <cuda/std/functional>
 #include <thrust/device_vector.h>
@@ -113,8 +114,8 @@ int main(void)
   thrust::device_vector<Key> insert_keys(num_keys);
   // Create a sequence of keys. Eeach distinct key has key_duplicates many matches.
   thrust::transform(
-    thrust::make_counting_iterator<Key>(0),
-    thrust::make_counting_iterator<Key>(insert_keys.size()),
+    cuda::make_counting_iterator<Key>(0),
+    cuda::make_counting_iterator<Key>(insert_keys.size()),
     insert_keys.begin(),
     [] __device__(auto i) -> Key { return static_cast<Key>(i % (num_keys / key_duplicates)); });
 

examples/static_map/count_by_key_example.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2026, 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,9 +17,8 @@
 #include <cuco/static_map.cuh>
 
 #include <cuda/functional>
+#include <cuda/iterator>
 #include <thrust/device_vector.h>
-#include <thrust/iterator/counting_iterator.h>
-#include <thrust/iterator/transform_iterator.h>
 #include <thrust/logical.h>
 #include <thrust/transform.h>
 
@@ -63,8 +62,8 @@ int main(void)
   auto const empty_value_sentinel = custom_value_type{-1};
 
   // Create an iterator of input key/value pairs
-  auto pairs_begin = thrust::make_transform_iterator(
-    thrust::make_counting_iterator<int32_t>(0),
+  auto pairs_begin = cuda::make_transform_iterator(
+    cuda::make_counting_iterator<int32_t>(0),
     cuda::proclaim_return_type<cuco::pair<custom_key_type, custom_value_type>>(
       [] __device__(auto i) { return cuco::pair{custom_key_type{i}, custom_value_type{i}}; }));
 
@@ -81,9 +80,9 @@ int main(void)
 
   // Reproduce inserted keys
   auto insert_keys =
-    thrust::make_transform_iterator(thrust::make_counting_iterator<int32_t>(0),
-                                    cuda::proclaim_return_type<custom_key_type>(
-                                      [] __device__(auto i) { return custom_key_type{i}; }));
+    cuda::make_transform_iterator(cuda::make_counting_iterator<int32_t>(0),
+                                  cuda::proclaim_return_type<custom_key_type>(
+                                    [] __device__(auto i) { return custom_key_type{i}; }));
 
   thrust::device_vector<bool> contained(num_pairs);
 

examples/static_map/custom_type_example.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2025, NVIDIA CORPORATION.
+ * Copyright (c) 2025-2026, 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 @@
 #include <cuco/static_map.cuh>
 
 #include <cuda/functional>
+#include <cuda/iterator>
 #include <cuda/std/tuple>
 #include <thrust/detail/raw_reference_cast.h>
 #include <thrust/device_vector.h>
 #include <thrust/host_vector.h>
-#include <thrust/iterator/counting_iterator.h>
-#include <thrust/iterator/transform_iterator.h>
 
 #include <iostream>
 
@@ -103,8 +102,8 @@ int main()
   };
   thrust::device_vector<value_type> d_values = {36.5f, 41.2f, 27.1f, 33.8f};
 
-  auto pairs_begin = thrust::make_transform_iterator(
-    thrust::counting_iterator{0},
+  auto pairs_begin = cuda::make_transform_iterator(
+    cuda::counting_iterator{0},
     cuda::proclaim_return_type<cuco::pair<stored_key, value_type>>(
       [keys = d_keys.begin(), values = d_values.begin()] __device__(int i) {
         return cuco::pair<stored_key, value_type>{keys[i], values[i]};