Update to new Thrust 1.12.0 APIs

cpp/src/indexing/construction/detail/phase_1.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -65,8 +65,8 @@ compute_point_keys_and_sorted_indices(cudf::column_view const &x,
 {
   rmm::device_uvector<uint32_t> keys(x.size(), stream);
   thrust::transform(rmm::exec_policy(stream),
-                    make_zip_iterator(x.begin<T>(), y.begin<T>()),
-                    make_zip_iterator(x.begin<T>(), y.begin<T>()) + x.size(),
+                    thrust::make_zip_iterator(x.begin<T>(), y.begin<T>()),
+                    thrust::make_zip_iterator(x.begin<T>(), y.begin<T>()) + x.size(),
                     keys.begin(),
                     [=] __device__(auto const &point) {
                       T x, y;
@@ -147,11 +147,11 @@ build_tree_levels(int8_t max_depth,
     keys_begin, [] __device__(uint32_t const child_key) { return (child_key >> 2); });
 
   // iterator for the current level's quad node point and child counts
-  auto child_nodes = make_zip_iterator(quad_point_count_begin, quad_child_count_begin);
+  auto child_nodes = thrust::make_zip_iterator(quad_point_count_begin, quad_child_count_begin);
 
   // iterator for the current level's initial values
   auto child_values =
-    make_zip_iterator(quad_point_count_begin, thrust::make_constant_iterator<uint32_t>(1));
+    thrust::make_zip_iterator(quad_point_count_begin, thrust::make_constant_iterator<uint32_t>(1));
 
   for (cudf::size_type level = max_depth - 1; level >= 0; --level) {
     auto num_full_quads = build_tree_level(parent_keys + begin,

cpp/src/indexing/construction/detail/phase_2.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -70,7 +70,7 @@ inline rmm::device_uvector<uint32_t> flatten_point_keys(
 {
   rmm::device_uvector<uint32_t> flattened_keys(num_valid_nodes, stream);
   auto keys_and_levels =
-    make_zip_iterator(quad_keys.begin(), quad_level.begin(), indicator.begin<bool>());
+    thrust::make_zip_iterator(quad_keys.begin(), quad_level.begin(), indicator.begin<bool>());
   thrust::transform(rmm::exec_policy(stream),
                     keys_and_levels,
                     keys_and_levels + num_valid_nodes,
@@ -122,7 +122,7 @@ inline rmm::device_uvector<uint32_t> compute_flattened_first_point_positions(
       rmm::exec_policy(stream),
       flattened_keys.begin(),
       flattened_keys.end(),
-      make_zip_iterator(initial_sort_indices.begin(), quad_point_count_tmp.begin()));
+      thrust::make_zip_iterator(initial_sort_indices.begin(), quad_point_count_tmp.begin()));
 
     thrust::remove_if(rmm::exec_policy(stream),
                       quad_point_count_tmp.begin(),
@@ -157,7 +157,7 @@ inline rmm::device_uvector<uint32_t> compute_flattened_first_point_positions(
                          quad_point_offsets_tmp.begin());
 
   auto counts_and_offsets =
-    make_zip_iterator(quad_point_count_tmp.begin(), quad_point_offsets_tmp.begin());
+    thrust::make_zip_iterator(quad_point_count_tmp.begin(), quad_point_offsets_tmp.begin());
 
   thrust::stable_sort_by_key(rmm::exec_policy(stream),
                              initial_sort_indices.begin(),
@@ -170,7 +170,7 @@ inline rmm::device_uvector<uint32_t> compute_flattened_first_point_positions(
                   counts_and_offsets,
                   counts_and_offsets + leaf_offsets.size(),
                   leaf_offsets.begin(),
-                  make_zip_iterator(quad_point_count.begin(), quad_point_offsets.begin()));
+                  thrust::make_zip_iterator(quad_point_count.begin(), quad_point_offsets.begin()));
 
   quad_point_offsets.shrink_to_fit(stream);
 
@@ -261,10 +261,10 @@ inline std::pair<uint32_t, uint32_t> remove_unqualified_quads(
   // Start counting nodes at level 2, since children of the root node should not
   // be discarded.
   // line 5 of algorithm in Fig. 5 in ref.
-  auto tree = make_zip_iterator(quad_keys.begin() + level_1_size,
-                                quad_point_count.begin() + level_1_size,
-                                quad_child_count.begin() + level_1_size,
-                                quad_levels.begin() + level_1_size);
+  auto tree = thrust::make_zip_iterator(quad_keys.begin() + level_1_size,
+                                        quad_point_count.begin() + level_1_size,
+                                        quad_child_count.begin() + level_1_size,
+                                        quad_levels.begin() + level_1_size);
 
   auto last_valid =
     thrust::remove_if(rmm::exec_policy(stream),

cpp/src/indexing/construction/detail/utilities.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -18,26 +18,14 @@
 
 #include <cudf/column/column_factories.hpp>
 #include <cudf/types.hpp>
-#include <cudf/utilities/type_dispatcher.hpp>
 
-#include <rmm/thrust_rmm_allocator.h>
 #include <rmm/cuda_stream_view.hpp>
 
-#include <thrust/iterator/zip_iterator.h>
 #include <thrust/tuple.h>
 
 namespace cuspatial {
 namespace detail {
 
-/**
- * @brief Helper function to reduce verbosity creating thrust zip iterators
- */
-template <typename... Ts>
-inline auto make_zip_iterator(Ts... its)
-{
-  return thrust::make_zip_iterator(thrust::make_tuple(std::forward<Ts>(its)...));
-}
-
 template <typename T>
 struct tuple_sum {
   inline __device__ thrust::tuple<T, T> operator()(thrust::tuple<T, T> const& a,

cpp/src/indexing/construction/point_quadtree.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -107,9 +107,9 @@ inline std::unique_ptr<cudf::table> make_quad_tree(rmm::device_uvector<uint32_t>
                            level_1_size);
 
     auto &offsets     = quad_child_pos;
-    auto offsets_iter = make_zip_iterator(is_quad->view().begin<bool>(),
-                                          quad_child_pos->view().template begin<uint32_t>(),
-                                          quad_point_pos.begin());
+    auto offsets_iter = thrust::make_zip_iterator(is_quad->view().begin<bool>(),
+                                                  quad_child_pos->view().template begin<uint32_t>(),
+                                                  quad_point_pos.begin());
 
     // for each value in `is_quad` copy from `quad_child_pos` if true, else
     // `quad_point_pos`
@@ -129,9 +129,9 @@ inline std::unique_ptr<cudf::table> make_quad_tree(rmm::device_uvector<uint32_t>
   auto lengths = make_fixed_width_column<uint32_t>(num_valid_nodes, stream, mr);
   // for each value in `is_quad` copy from `quad_child_count` if true, else
   // `quad_point_count`
-  auto lengths_iter = make_zip_iterator(is_quad->view().begin<bool>(),  //
-                                        quad_child_count.begin(),
-                                        quad_point_count.begin());
+  auto lengths_iter = thrust::make_zip_iterator(is_quad->view().begin<bool>(),  //
+                                                quad_child_count.begin(),
+                                                quad_point_count.begin());
   thrust::transform(rmm::exec_policy(stream),
                     lengths_iter,
                     lengths_iter + num_valid_nodes,

cpp/src/join/detail/intersection.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, 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,6 +29,7 @@
 
 #include <thrust/copy.h>
 #include <thrust/iterator/transform_iterator.h>
+#include <thrust/iterator/zip_iterator.h>
 #include <thrust/remove.h>
 #include <thrust/tuple.h>
 
@@ -97,8 +98,8 @@ inline std::pair<cudf::size_type, cudf::size_type> find_intersections(
   auto d_poly_y_max = cudf::column_device_view::create(poly_bbox.column(3), stream);
 
   thrust::transform(rmm::exec_policy(stream),
-                    make_zip_iterator(node_indices, poly_indices),
-                    make_zip_iterator(node_indices, poly_indices) + num_pairs,
+                    thrust::make_zip_iterator(node_indices, poly_indices),
+                    thrust::make_zip_iterator(node_indices, poly_indices) + num_pairs,
                     node_pairs,
                     [x_min,
                      y_min,

cpp/src/join/detail/traversal.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, 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 <rmm/exec_policy.hpp>
 
 #include <thrust/gather.h>
+#include <thrust/iterator/zip_iterator.h>
 #include <thrust/scan.h>
 #include <thrust/scatter.h>
 #include <thrust/tuple.h>
@@ -106,15 +107,15 @@ descend_quadtree(LengthsIter counts,
                  parent_indices.begin(),
                  parent_indices.begin() + num_children,
                  // curr level iterator
-                 make_zip_iterator(parent_types.begin(),
-                                   parent_levels.begin(),
-                                   parent_node_indices.begin(),
-                                   parent_poly_indices.begin()),
+                 thrust::make_zip_iterator(parent_types.begin(),
+                                           parent_levels.begin(),
+                                           parent_node_indices.begin(),
+                                           parent_poly_indices.begin()),
                  // next level iterator
-                 make_zip_iterator(child_types.begin(),
-                                   child_levels.begin(),
-                                   child_quad_indices.begin(),
-                                   child_poly_indices.begin()));
+                 thrust::make_zip_iterator(child_types.begin(),
+                                           child_levels.begin(),
+                                           child_quad_indices.begin(),
+                                           child_poly_indices.begin()));
 
   rmm::device_uvector<uint32_t> relative_child_offsets(num_children, stream);
   thrust::exclusive_scan_by_key(rmm::exec_policy(stream),

cpp/src/join/quadtree_point_in_polygon.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -34,6 +34,7 @@
 #include <thrust/copy.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/permutation_iterator.h>
+#include <thrust/iterator/zip_iterator.h>
 #include <thrust/remove.h>
 #include <thrust/transform.h>
 
@@ -144,7 +145,7 @@ struct compute_quadtree_point_in_polygon {
     rmm::device_uvector<uint32_t> poly_idxs(num_total_points, stream);
     rmm::device_uvector<uint32_t> point_idxs(num_total_points, stream);
 
-    auto poly_and_point_indices = make_zip_iterator(poly_idxs.begin(), point_idxs.begin());
+    auto poly_and_point_indices = thrust::make_zip_iterator(poly_idxs.begin(), point_idxs.begin());
 
     // Compute the combination of polygon and point index pairs. For each polygon/quadrant pair,
     // enumerate pairs of (poly_index, point_index) for each point in each quadrant.
@@ -169,7 +170,8 @@ struct compute_quadtree_point_in_polygon {
 
     // Enumerate the point X/Ys using the sorted `point_indices` (from quadtree construction)
     auto point_xys_iter = thrust::make_permutation_iterator(
-      make_zip_iterator(point_x.begin<T>(), point_y.begin<T>()), point_indices.begin<uint32_t>());
+      thrust::make_zip_iterator(point_x.begin<T>(), point_y.begin<T>()),
+      point_indices.begin<uint32_t>());
 
     // Compute the number of intersections by removing (poly, point) pairs that don't intersect
     auto num_intersections = thrust::distance(

cpp/src/join/quadtree_point_to_nearest_polyline.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -34,6 +34,7 @@
 #include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
+#include <thrust/iterator/zip_iterator.h>
 #include <thrust/reduce.h>
 
 #include <vector>

cpp/src/join/quadtree_poly_filtering.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, 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,8 @@
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 
+#include <thrust/iterator/zip_iterator.h>
+
 #include <tuple>
 
 namespace cuspatial {
@@ -89,13 +91,13 @@ inline std::unique_ptr<cudf::table> join_quadtree_and_bboxes(cudf::table_view co
   cudf::size_type num_parents{0};
 
   auto make_current_level_iter = [&]() {
-    return make_zip_iterator(
+    return thrust::make_zip_iterator(
       cur_types.begin(), cur_levels.begin(), cur_node_idxs.begin(), cur_poly_idxs.begin());
   };
 
   auto make_output_values_iter = [&]() {
     return num_results +
-           make_zip_iterator(
+           thrust::make_zip_iterator(
              out_types.begin(), out_levels.begin(), out_node_idxs.begin(), out_poly_idxs.begin());
   };
 

cpp/tests/join/point_in_polygon_test_large.cpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -96,7 +96,7 @@ inline auto make_polygons_geometry(thrust::host_vector<uint32_t> const &poly_off
     }
     polygons.push_back(polygon);
   }
-  return std::move(polygons);
+  return polygons;
 }
 
 template <typename T>