Replace Thrust iterator facilities with libcu++ ones

cpp/benchmarks/indexing/quadtree_on_points.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.
@@ -23,6 +23,7 @@
 #include <rmm/device_vector.hpp>
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/std/iterator>
 #include <thrust/host_vector.h>
 
 #include <nvbench/nvbench.cuh>
@@ -104,13 +105,13 @@ void quadtree_on_points_benchmark(nvbench::state& state, nvbench::type_list<T>)
     thrust::device, d_points.begin(), d_points.end(), [] __device__(auto const& a, auto const& b) {
       return a.x <= b.x && a.y <= b.y;
     });
-  auto const vertex_1 = h_points[thrust::distance(d_points.begin(), vertex_1_itr)];
+  auto const vertex_1 = h_points[cuda::std::distance(d_points.begin(), vertex_1_itr)];
 
   auto const vertex_2_itr = thrust::max_element(
     thrust::device, d_points.begin(), d_points.end(), [] __device__(auto const& a, auto const& b) {
       return a.x >= b.x && a.y >= b.y;
     });
-  auto const vertex_2 = h_points[thrust::distance(d_points.begin(), vertex_2_itr)];
+  auto const vertex_2 = h_points[cuda::std::distance(d_points.begin(), vertex_2_itr)];
 
   // TODO: to be replaced by nvbench fixture once it's ready
   cuspatial::rmm_pool_raii rmm_pool;

cpp/include/cuspatial/detail/distance/hausdorff.cuh

@@ -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.
@@ -25,9 +25,8 @@
 #include <rmm/exec_policy.hpp>
 
 #include <cuda/atomic>
-#include <thrust/advance.h>
+#include <cuda/std/iterator>
 #include <thrust/binary_search.h>
-#include <thrust/distance.h>
 #include <thrust/execution_policy.h>
 #include <thrust/fill.h>
 #include <thrust/memory.h>
@@ -92,7 +91,7 @@ CUSPATIAL_KERNEL void kernel_hausdorff(
     auto const lhs_space_iter =
       thrust::upper_bound(thrust::seq, space_offsets, space_offsets + num_spaces, lhs_p_idx);
     // determine the LHS space this point belongs to.
-    Index const lhs_space_idx = thrust::distance(space_offsets, thrust::prev(lhs_space_iter));
+    Index const lhs_space_idx = cuda::std::distance(space_offsets, cuda::std::prev(lhs_space_iter));
 
     // get the coordinates of this LHS point.
     Point const lhs_p = points[lhs_p_idx];

cpp/include/cuspatial/detail/geometry/linestring_ref.cuh

@@ -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.
@@ -22,6 +22,7 @@
 #include <cuspatial/iterator_factory.cuh>
 #include <cuspatial/traits.hpp>
 
+#include <cuda/std/iterator>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/tuple.h>
 
@@ -58,8 +59,8 @@ template <typename VecIterator>
 CUSPATIAL_HOST_DEVICE auto linestring_ref<VecIterator>::num_segments() const
 {
   // The number of segment equals the number of points minus 1. And the number of points
-  // is thrust::distance(_point_begin, _point_end).
-  return thrust::distance(_point_begin, _point_end) - 1;
+  // is cuda::std::distance(_point_begin, _point_end).
+  return cuda::std::distance(_point_begin, _point_end) - 1;
 }
 
 template <typename VecIterator>

cpp/include/cuspatial/detail/geometry/polygon_ref.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.
@@ -23,6 +23,7 @@
 #include <cuspatial/iterator_factory.cuh>
 #include <cuspatial/traits.hpp>
 
+#include <cuda/std/iterator>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/tuple.h>
 
@@ -44,7 +45,7 @@ CUSPATIAL_HOST_DEVICE polygon_ref<RingIterator, VecIterator>::polygon_ref(RingIt
 template <typename RingIterator, typename VecIterator>
 CUSPATIAL_HOST_DEVICE auto polygon_ref<RingIterator, VecIterator>::num_rings() const
 {
-  return thrust::distance(_ring_begin, _ring_end) - 1;
+  return cuda::std::distance(_ring_begin, _ring_end) - 1;
 }
 
 template <typename RingIterator, typename VecIterator>
@@ -63,13 +64,13 @@ CUSPATIAL_HOST_DEVICE auto polygon_ref<RingIterator, VecIterator>::ring_end() co
 template <typename RingIterator, typename VecIterator>
 CUSPATIAL_HOST_DEVICE auto polygon_ref<RingIterator, VecIterator>::point_begin() const
 {
-  return thrust::next(_point_begin, *_ring_begin);
+  return cuda::std::next(_point_begin, *_ring_begin);
 }
 
 template <typename RingIterator, typename VecIterator>
 CUSPATIAL_HOST_DEVICE auto polygon_ref<RingIterator, VecIterator>::point_end() const
 {
-  return thrust::next(_point_begin, *thrust::prev(_ring_end));
+  return cuda::std::next(_point_begin, *cuda::std::prev(_ring_end));
 }
 
 template <typename RingIterator, typename VecIterator>

cpp/include/cuspatial/detail/geometry_collection/multilinestring_ref.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.
@@ -20,6 +20,7 @@
 #include <cuspatial/geometry/linestring_ref.cuh>
 #include <cuspatial/iterator_factory.cuh>
 
+#include <cuda/std/iterator>
 #include <thrust/iterator/transform_iterator.h>
 
 namespace cuspatial {
@@ -55,7 +56,7 @@ CUSPATIAL_HOST_DEVICE multilinestring_ref<PartIterator, VecIterator>::multilines
 template <typename PartIterator, typename VecIterator>
 CUSPATIAL_HOST_DEVICE auto multilinestring_ref<PartIterator, VecIterator>::num_linestrings() const
 {
-  return thrust::distance(_part_begin, _part_end) - 1;
+  return cuda::std::distance(_part_begin, _part_end) - 1;
 }
 
 template <typename PartIterator, typename VecIterator>
@@ -74,15 +75,15 @@ CUSPATIAL_HOST_DEVICE auto multilinestring_ref<PartIterator, VecIterator>::part_
 template <typename PartIterator, typename VecIterator>
 CUSPATIAL_HOST_DEVICE auto multilinestring_ref<PartIterator, VecIterator>::point_begin() const
 {
-  return thrust::next(_point_begin, *_part_begin);
+  return cuda::std::next(_point_begin, *_part_begin);
 }
 
 template <typename PartIterator, typename VecIterator>
 CUSPATIAL_HOST_DEVICE auto multilinestring_ref<PartIterator, VecIterator>::point_end() const
 {
   // _part_end refers to the one past the last part index to the points of this multilinestring.
   // So prior to computing the end point index, we need to decrement _part_end.
-  return thrust::next(_point_begin, *thrust::prev(_part_end));
+  return cuda::std::next(_point_begin, *cuda::std::prev(_part_end));
 }
 
 template <typename PartIterator, typename VecIterator>

cpp/include/cuspatial/detail/geometry_collection/multipoint_ref.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022-2023, 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.
@@ -19,7 +19,7 @@
 #include <cuspatial/cuda_utils.hpp>
 #include <cuspatial/iterator_factory.cuh>
 
-#include <thrust/distance.h>
+#include <cuda/std/iterator>
 
 namespace cuspatial {
 
@@ -45,7 +45,7 @@ CUSPATIAL_HOST_DEVICE auto multipoint_ref<VecIterator>::point_end() const
 template <typename VecIterator>
 CUSPATIAL_HOST_DEVICE auto multipoint_ref<VecIterator>::num_points() const
 {
-  return thrust::distance(_points_begin, _points_end);
+  return cuda::std::distance(_points_begin, _points_end);
 }
 
 template <typename VecIterator>

cpp/include/cuspatial/detail/geometry_collection/multipolygon_ref.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.
@@ -21,6 +21,7 @@
 #include <cuspatial/geometry/polygon_ref.cuh>
 #include <cuspatial/iterator_factory.cuh>
 
+#include <cuda/std/iterator>
 #include <thrust/iterator/transform_iterator.h>
 
 namespace cuspatial {
@@ -45,7 +46,7 @@ struct to_polygon_functor {
   CUSPATIAL_HOST_DEVICE auto operator()(difference_type i)
   {
     return polygon_ref{ring_begin + part_begin[i],
-                       thrust::next(ring_begin + part_begin[i + 1]),
+                       cuda::std::next(ring_begin + part_begin[i + 1]),
                        point_begin,
                        point_end};
   }
@@ -75,7 +76,7 @@ template <typename PartIterator, typename RingIterator, typename VecIterator>
 CUSPATIAL_HOST_DEVICE auto multipolygon_ref<PartIterator, RingIterator, VecIterator>::num_polygons()
   const
 {
-  return thrust::distance(_part_begin, _part_end) - 1;
+  return cuda::std::distance(_part_begin, _part_end) - 1;
 }
 
 template <typename PartIterator, typename RingIterator, typename VecIterator>
@@ -112,14 +113,14 @@ template <typename PartIterator, typename RingIterator, typename VecIterator>
 CUSPATIAL_HOST_DEVICE auto multipolygon_ref<PartIterator, RingIterator, VecIterator>::point_begin()
   const
 {
-  return thrust::next(_point_begin, *thrust::next(_ring_begin, *_part_begin));
+  return cuda::std::next(_point_begin, *cuda::std::next(_ring_begin, *_part_begin));
 }
 
 template <typename PartIterator, typename RingIterator, typename VecIterator>
 CUSPATIAL_HOST_DEVICE auto multipolygon_ref<PartIterator, RingIterator, VecIterator>::point_end()
   const
 {
-  return thrust::next(_point_begin, *thrust::next(_ring_begin, *thrust::prev(_part_end)));
+  return cuda::std::next(_point_begin, *cuda::std::next(_ring_begin, *cuda::std::prev(_part_end)));
 }
 
 template <typename PartIterator, typename RingIterator, typename VecIterator>

cpp/include/cuspatial/detail/index/construction/phase_1.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.
@@ -27,8 +27,8 @@
 #include <rmm/resource_ref.hpp>
 
 #include <cuda/functional>
+#include <cuda/std/iterator>
 #include <thrust/copy.h>
-#include <thrust/distance.h>
 #include <thrust/fill.h>
 #include <thrust/functional.h>
 #include <thrust/iterator/constant_iterator.h>
@@ -68,7 +68,7 @@ compute_point_keys_and_sorted_indices(PointIt points_first,
                                       rmm::cuda_stream_view stream,
                                       rmm::device_async_resource_ref mr)
 {
-  auto num_points = thrust::distance(points_first, points_last);
+  auto num_points = cuda::std::distance(points_first, points_last);
   rmm::device_uvector<uint32_t> keys(num_points, stream);
   thrust::transform(
     rmm::exec_policy(stream),
@@ -121,7 +121,7 @@ inline IndexT build_tree_level(KeyInputIterator keys_begin,
                                       vals_out,
                                       thrust::equal_to<uint32_t>(),
                                       binary_op);
-  return thrust::distance(keys_out, result.first);
+  return cuda::std::distance(keys_out, result.first);
 }
 
 /**
@@ -262,7 +262,7 @@ inline auto make_full_levels(PointIt points_first,
                              rmm::cuda_stream_view stream,
                              rmm::device_async_resource_ref mr)
 {
-  auto num_points = thrust::distance(points_first, points_last);
+  auto num_points = cuda::std::distance(points_first, points_last);
   // Compute point keys and sort into bottom-level quadrants
   // (i.e. quads at level `max_depth - 1`)
 

cpp/include/cuspatial/detail/index/construction/phase_2.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.
@@ -23,9 +23,9 @@
 #include <rmm/exec_policy.hpp>
 #include <rmm/resource_ref.hpp>
 
+#include <cuda/std/iterator>
 #include <thrust/copy.h>
 #include <thrust/count.h>
-#include <thrust/distance.h>
 #include <thrust/fill.h>
 #include <thrust/functional.h>
 #include <thrust/iterator/counting_iterator.h>
@@ -64,7 +64,7 @@ inline rmm::device_uvector<uint32_t> compute_leaf_positions(
                                 leaf_pos.begin(),
                                 !thrust::placeholders::_1);
   // Shrink leaf_pos's underlying device allocation
-  leaf_pos.resize(thrust::distance(leaf_pos.begin(), result), stream);
+  leaf_pos.resize(cuda::std::distance(leaf_pos.begin(), result), stream);
   leaf_pos.shrink_to_fit(stream);
   return leaf_pos;
 }
@@ -281,7 +281,7 @@ inline std::pair<uint32_t, uint32_t> remove_unqualified_quads(
                       [max_size] __device__(auto const n) { return n <= max_size; });
 
   // add the number of level 1 nodes back in to num_valid_nodes
-  auto num_valid_nodes = thrust::distance(tree, last_valid) + level_1_size;
+  auto num_valid_nodes = cuda::std::distance(tree, last_valid) + level_1_size;
 
   quad_keys.resize(num_valid_nodes, stream);
   quad_keys.shrink_to_fit(stream);

cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh

@@ -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.
@@ -37,6 +37,7 @@
 #include <rmm/resource_ref.hpp>
 
 #include <cuda/atomic>
+#include <cuda/std/iterator>
 #include <thrust/binary_search.h>
 #include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/discard_iterator.h>
@@ -116,9 +117,9 @@ struct types_buffer_functor {
   template <typename index_t>
   type_t __device__ operator()(index_t i)
   {
-    auto geometry_idx = thrust::distance(
+    auto geometry_idx = cuda::std::distance(
       geometric_column_offset.begin(),
-      thrust::prev(thrust::upper_bound(
+      cuda::std::prev(thrust::upper_bound(
         thrust::seq, geometric_column_offset.begin(), geometric_column_offset.end(), i)));
 
     auto num_points   = points_offset[geometry_idx + 1] - points_offset[geometry_idx];