Use cuda::proclaim_return_type on device lambdas.

cpp/include/cuspatial/column/geometry_column_view.hpp

@@ -16,7 +16,6 @@
 
 #pragma once
 
-#include <cuspatial/range/range.cuh>
 #include <cuspatial/types.hpp>
 
 #include <cudf/lists/lists_column_view.hpp>

cpp/include/cuspatial/detail/index/construction/phase_1.cuh

@@ -38,6 +38,8 @@
 #include <thrust/transform.h>
 #include <thrust/tuple.h>
 
+#include <cuda/functional>
+
 #include <memory>
 #include <tuple>
 #include <utility>
@@ -73,14 +75,14 @@ compute_point_keys_and_sorted_indices(PointIt points_first,
     points_first,
     points_last,
     keys.begin(),
-    [=] __device__(vec_2d<T> const& point) {
+    cuda::proclaim_return_type<std::uint32_t>([=] __device__(vec_2d<T> const& point) {
       if (point.x < min.x || point.x > max.x || point.y < min.y || point.y > max.y) {
         // If the point is outside the bbox, return a max_level key
         return static_cast<uint32_t>((1 << (2 * max_depth)) - 1);
       }
       return cuspatial::detail::utility::z_order(static_cast<uint16_t>((point.x - min.x) / scale),
                                                  static_cast<uint16_t>((point.y - min.y) / scale));
-    });
+    }));
 
   rmm::device_uvector<uint32_t> indices(keys.size(), stream, mr);
 
@@ -145,7 +147,9 @@ inline std::tuple<IndexT, IndexT, std::vector<IndexT>, std::vector<IndexT>> buil
 
   // iterator for the parent level's quad node keys
   auto parent_keys = thrust::make_transform_iterator(
-    keys_begin, [] __device__(uint32_t const child_key) { return (child_key >> 2); });
+    keys_begin, cuda::proclaim_return_type<uint32_t>([] __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 = thrust::make_zip_iterator(quad_point_count_begin, quad_child_count_begin);

cpp/include/cuspatial/detail/join/quadtree_bbox_filtering.cuh

@@ -24,6 +24,8 @@
 
 #include <thrust/iterator/discard_iterator.h>
 
+#include <cuda/functional>
+
 #include <iterator>
 #include <utility>
 
@@ -47,10 +49,10 @@ join_quadtree_and_bounding_boxes(point_quadtree_ref quadtree,
 
   // Count the number of top-level nodes to start.
   // This could be provided explicitly, but count_if should be fast enough.
-  auto num_top_level_leaves = thrust::count_if(rmm::exec_policy(stream),
-                                               quadtree.level_begin(),
-                                               quadtree.level_end(),
-                                               thrust::placeholders::_1 == 0);
+  int32_t num_top_level_leaves = thrust::count_if(rmm::exec_policy(stream),
+                                                  quadtree.level_begin(),
+                                                  quadtree.level_end(),
+                                                  thrust::placeholders::_1 == 0);
 
   auto num_pairs = num_top_level_leaves * num_boxes;
 
@@ -89,10 +91,14 @@ join_quadtree_and_bounding_boxes(point_quadtree_ref quadtree,
                                bounding_boxes_first,
                                // The top-level node indices
                                detail::make_counting_transform_iterator(
-                                 0, [=] __device__(auto i) { return i % num_top_level_leaves; }),
+                                 0, cuda::proclaim_return_type<int32_t>([=] __device__(auto i) {
+                                   return i % num_top_level_leaves;
+                                 })),
                                // The top-level bbox indices
                                detail::make_counting_transform_iterator(
-                                 0, [=] __device__(auto i) { return i / num_top_level_leaves; }),
+                                 0, cuda::proclaim_return_type<int32_t>([=] __device__(auto i) {
+                                   return i / num_top_level_leaves;
+                                 })),
                                make_current_level_iter(),  // intermediate intersections or parent
                                                            // quadrants found during traversal
                                // found intersecting quadrant and bbox indices for output

cpp/include/cuspatial/detail/join/quadtree_point_to_nearest_linestring.cuh

@@ -25,6 +25,7 @@
 #include <cuspatial/range/multilinestring_range.cuh>
 #include <cuspatial/traits.hpp>
 
+#include <functional>
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 
@@ -247,7 +248,8 @@ quadtree_point_to_nearest_linestring(LinestringIndexIterator linestring_indices_
 
   auto all_point_indices =
     thrust::make_transform_iterator(all_point_linestring_indices_and_distances,
-                                    [] __device__(auto const& x) { return thrust::get<0>(x); });
+                                    cuda::proclaim_return_type<uint32_t>(
+                                      [] __device__(auto const& x) { return thrust::get<0>(x); }));
 
   // Allocate vectors for the distances min reduction
   auto num_points = std::distance(point_indices_first, point_indices_last);
@@ -272,22 +274,23 @@ quadtree_point_to_nearest_linestring(LinestringIndexIterator linestring_indices_
       thrust::make_discard_iterator(), output_linestring_idxs.begin(), output_distances.begin()),
     thrust::equal_to<uint32_t>(),  // comparator
     // binop to select the point/linestring pair with the smallest distance
-    [] __device__(auto const& lhs, auto const& rhs) {
-      T const& d_lhs = thrust::get<2>(lhs);
-      T const& d_rhs = thrust::get<2>(rhs);
-      // If lhs distance is 0, choose rhs
-      if (d_lhs == T{0}) { return rhs; }
-      // if rhs distance is 0, choose lhs
-      if (d_rhs == T{0}) { return lhs; }
-      // If distances to lhs/rhs are the same, choose linestring with smallest id
-      if (d_lhs == d_rhs) {
-        auto const& i_lhs = thrust::get<1>(lhs);
-        auto const& i_rhs = thrust::get<1>(rhs);
-        return i_lhs < i_rhs ? lhs : rhs;
-      }
-      // Otherwise choose linestring with smallest distance
-      return d_lhs < d_rhs ? lhs : rhs;
-    });
+    cuda::proclaim_return_type<thrust::tuple<uint32_t, uint32_t, T>>(
+      [] __device__(auto const& lhs, auto const& rhs) {
+        T const& d_lhs = thrust::get<2>(lhs);
+        T const& d_rhs = thrust::get<2>(rhs);
+        // If lhs distance is 0, choose rhs
+        if (d_lhs == T{0}) { return rhs; }
+        // if rhs distance is 0, choose lhs
+        if (d_rhs == T{0}) { return lhs; }
+        // If distances to lhs/rhs are the same, choose linestring with smallest id
+        if (d_lhs == d_rhs) {
+          auto const& i_lhs = thrust::get<1>(lhs);
+          auto const& i_rhs = thrust::get<1>(rhs);
+          return i_lhs < i_rhs ? lhs : rhs;
+        }
+        // Otherwise choose linestring with smallest distance
+        return d_lhs < d_rhs ? lhs : rhs;
+      }));
 
   auto const num_distances = thrust::distance(point_idxs.begin(), point_idxs_end.first);
 

cpp/include/cuspatial/detail/point_quadtree.cuh

@@ -30,6 +30,8 @@
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/tuple.h>
 
+#include <cuda/functional>
+
 #include <tuple>
 
 namespace cuspatial {
@@ -108,9 +110,9 @@ inline point_quadtree make_quad_tree(rmm::device_uvector<uint32_t>& keys,
                       offsets_iter + num_valid_nodes,
                       offsets.begin(),
                       // return is_internal_node ? lhs : rhs
-                      [] __device__(auto const& t) {
+                      cuda::proclaim_return_type<uint32_t>([] __device__(auto const& t) {
                         return thrust::get<0>(t) ? thrust::get<1>(t) : thrust::get<2>(t);
-                      });
+                      }));
 
     return std::move(offsets);
   }();
@@ -126,9 +128,9 @@ inline point_quadtree make_quad_tree(rmm::device_uvector<uint32_t>& keys,
                     lengths_iter + num_valid_nodes,
                     lengths.begin(),
                     // return bool ? lhs : rhs
-                    [] __device__(auto const& t) {
+                    cuda::proclaim_return_type<uint32_t>([] __device__(auto const& t) {
                       return thrust::get<0>(t) ? thrust::get<1>(t) : thrust::get<2>(t);
-                    });
+                    }));
 
   // Shrink keys to the number of valid nodes
   keys.resize(num_valid_nodes, stream);

cpp/include/cuspatial_test/vector_factories.cuh

@@ -68,10 +68,12 @@ auto make_device_uvector(std::initializer_list<T> inl,
   return res;
 }
 
+// TODO: this can be eliminated when Thrust 2.1.0 is the minimum because
+// thrust::host_vector has a constructor that takes an initializer_list
 template <typename T>
 auto make_host_vector(std::initializer_list<T> inl)
 {
-  return thrust::host_vector<T>{inl.begin(), inl.end()};
+  return thrust::host_vector<T>(inl.begin(), inl.end());
 }
 
 /**

cpp/tests/distance/hausdorff_test.cu

@@ -30,6 +30,8 @@
 #include <gmock/gmock.h>
 #include <gtest/gtest.h>
 
+#include <cuda/functional>
+
 #include <iterator>
 
 template <typename T>
@@ -43,7 +45,7 @@ struct HausdorffTest : public ::testing::Test {
     auto const d_space_offsets = rmm::device_vector<Index>{space_offsets};
 
     auto const num_distances = space_offsets.size() * space_offsets.size();
-    auto distances           = rmm::device_vector<T>{num_distances};
+    auto distances           = rmm::device_vector<T>(num_distances);
 
     auto const distances_end = cuspatial::directed_hausdorff_distance(d_points.begin(),
                                                                       d_points.end(),
@@ -150,10 +152,12 @@ void generic_hausdorff_test()
   auto zero_iter         = thrust::make_constant_iterator<vec_2d>({0, 0});
   auto counting_iter     = thrust::make_counting_iterator<uint32_t>(0);
   auto space_offset_iter = thrust::make_transform_iterator(
-    counting_iter, [] __device__(auto idx) { return idx * elements_per_space; });
+    counting_iter, cuda::proclaim_return_type<uint32_t>([] __device__(auto idx) {
+      return idx * elements_per_space;
+    }));
 
-  auto distances = rmm::device_vector<T>{num_distances};
-  auto expected  = rmm::device_vector<T>{num_distances, 0};
+  auto distances = rmm::device_vector<T>(num_distances);
+  auto expected  = rmm::device_vector<T>(num_distances, 0);
 
   auto distances_end = cuspatial::directed_hausdorff_distance(zero_iter,
                                                               zero_iter + num_points,

cpp/tests/distance/haversine_test.cu

@@ -59,8 +59,8 @@ TYPED_TEST(HaversineTest, Zero)
   auto a_lonlat = rmm::device_vector<Location>(1, Location{0, 0});
   auto b_lonlat = rmm::device_vector<Location>(1, Location{0, 0});
 
-  auto distance = rmm::device_vector<T>{1, -1};
-  auto expected = rmm::device_vector<T>{1, 0};
+  auto distance = rmm::device_vector<T>(1, -1);
+  auto expected = rmm::device_vector<T>(1, 0);
 
   auto distance_end = cuspatial::haversine_distance(
     a_lonlat.begin(), a_lonlat.end(), b_lonlat.begin(), distance.begin());

cpp/tests/distance/linestring_distance_test.cu

@@ -209,8 +209,8 @@ TYPED_TEST(PairwiseLinestringDistanceTest, FromLongInputs)
   auto offset =
     rmm::device_vector<int32_t>{std::vector<int32_t>{0, 100, 200, 300, 400, num_points}};
 
-  auto got      = rmm::device_vector<T>{num_pairs};
-  auto expected = rmm::device_vector<T>{std::vector<T>{42.0, 42.0, 42.0, 42.0, 42.0}};
+  auto got      = rmm::device_vector<T>(num_pairs);
+  auto expected = rmm::device_vector<T>{{42.0, 42.0, 42.0, 42.0, 42.0}};
 
   auto mlinestrings1 = make_multilinestring_range(num_pairs,
                                                   thrust::make_counting_iterator(0),

cpp/tests/trajectory/trajectory_distances_and_speeds_test.cu

@@ -29,6 +29,8 @@
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/reduce.h>
 
+#include <cuda/functional>
+
 #include <gtest/gtest.h>
 
 #include <limits>
@@ -66,7 +68,8 @@ struct TrajectoryDistancesAndSpeedsTest : public ::testing::Test {
       thrust::reduce(expected_speeds.begin(),
                      expected_speeds.end(),
                      std::numeric_limits<T>::lowest(),
-                     [] __device__(T const& a, T const& b) { return max(abs(a), abs(b)); });
+                     cuda::proclaim_return_type<T>(
+                       [] __device__(T const& a, T const& b) { return max(abs(a), abs(b)); }));
 
     // We expect the floating point error (in ulps) due to be proportional to the  number of
     // operations to compute the relevant quantity. For distance, this computation is

cpp/tests/trajectory/trajectory_test_utils.cuh

@@ -296,7 +296,7 @@ struct trajectory_test_data {
 
     auto id_and_position = thrust::make_zip_iterator(ids_sorted.begin(), points_sorted.begin());
 
-    auto distance_per_step = rmm::device_vector<T>{points.size()};
+    auto distance_per_step = rmm::device_vector<T>(points.size());
 
     thrust::transform(rmm::exec_policy(),
                       id_and_position,