Fix scatter bug due to overlapping range in `pairwise_linestring_inte…

…rsection` (#1152)

This PR closes #1149 

`thrust::scatter` does not perform in place scatter. As the document says:

```
The iterator result + i shall not refer to any element referenced by any iterator j in the range [first,last) for all iterators i in the range [map,map + (last - first)).
```

The input and output range must not overlap. However, currently in intersection there is overlap. This may cause a bad scatter that only happens when input is large (device dependent).

This PR also fixes a bug in `find_duplicate_point` kernel where the `duplicate_flag` was incorrectly referenced.

Additionally, this PR includes a fix to the python API to handle sparse geoseries.

Authors:
  - Michael Wang (https://github.com/isVoid)

Approvers:
  - Mark Harris (https://github.com/harrism)
  - H. Thomson Comer (https://github.com/thomcom)

URL: #1152

cpp/include/cuspatial/detail/find/find_duplicate_points.cuh

@@ -45,10 +45,13 @@ void __global__ find_duplicate_points_kernel_simple(MultiPointRange multipoints,
       duplicate_flags[i + global_offset] = 0;
     }
 
-    for (auto i = 0; i < multipoint.size() && duplicate_flags[i] != 1; ++i)
+    // Loop over the point range to find duplicates, skipping if the point is already marked as
+    // duplicate.
+    for (auto i = 0; i < multipoint.size() && duplicate_flags[i + global_offset] != 1; ++i) {
       for (auto j = i + 1; j < multipoint.size(); ++j) {
         if (multipoint[i] == multipoint[j]) duplicate_flags[j + global_offset] = 1;
       }
+    }
   }
 }
 

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

@@ -36,6 +36,7 @@
 #include <rmm/mr/device/per_device_resource.hpp>
 
 #include <thrust/binary_search.h>
+#include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/reduce.h>
@@ -159,30 +160,33 @@ std::unique_ptr<rmm::device_uvector<types_t>> compute_types_buffer(
  * This is performing a group-by cumulative sum (pandas semantic) operation
  * to an "all 1s vector", using `types_buffer` as the key column.
  */
-template <typename index_t>
+template <typename index_t, typename types_t>
 std::unique_ptr<rmm::device_uvector<index_t>> compute_offset_buffer(
-  rmm::device_uvector<uint8_t> const& types_buffer,
+  rmm::device_uvector<types_t> const& types_buffer,
   rmm::mr::device_memory_resource* mr,
   rmm::cuda_stream_view stream)
 {
-  auto N            = types_buffer.size();
-  auto keys_copy    = rmm::device_uvector(types_buffer, stream);
-  auto indices_temp = rmm::device_uvector<index_t>(N, stream);
-  thrust::sequence(rmm::exec_policy(stream), indices_temp.begin(), indices_temp.end());
-  thrust::stable_sort_by_key(
-    rmm::exec_policy(stream), keys_copy.begin(), keys_copy.end(), indices_temp.begin());
+  auto N                                = types_buffer.size();
+  auto [offset_buffer_grouped, indices] = [&]() {
+    auto indices = rmm::device_uvector<index_t>(N, stream);
+    auto keys    = rmm::device_uvector<types_t>(types_buffer, stream);
+    thrust::sequence(rmm::exec_policy(stream), indices.begin(), indices.end());
+    thrust::stable_sort_by_key(rmm::exec_policy(stream), keys.begin(), keys.end(), indices.begin());
+
+    auto offset_buffer_grouped = std::make_unique<rmm::device_uvector<index_t>>(N, stream);
+
+    auto one_it = thrust::make_constant_iterator(1);
+    thrust::exclusive_scan_by_key(
+      rmm::exec_policy(stream), keys.begin(), keys.end(), one_it, offset_buffer_grouped->begin());
+
+    return std::pair{std::move(offset_buffer_grouped), std::move(indices)};
+  }();
 
   auto offset_buffer = std::make_unique<rmm::device_uvector<index_t>>(N, stream, mr);
-  thrust::uninitialized_fill_n(rmm::exec_policy(stream), offset_buffer->begin(), N, 1);
-  thrust::exclusive_scan_by_key(rmm::exec_policy(stream),
-                                keys_copy.begin(),
-                                keys_copy.end(),
-                                offset_buffer->begin(),
-                                offset_buffer->begin());
   thrust::scatter(rmm::exec_policy(stream),
-                  offset_buffer->begin(),
-                  offset_buffer->end(),
-                  indices_temp.begin(),
+                  offset_buffer_grouped->begin(),
+                  offset_buffer_grouped->end(),
+                  indices.begin(),
                   offset_buffer->begin());
   return offset_buffer;
 }
@@ -226,7 +230,7 @@ linestring_intersection_result<T, index_t> pairwise_linestring_intersection(
   // Phase 3: Remove duplicate points from intermediates
   // TODO: improve memory usage by using IIFE to
   // Remove the duplicate points
-  rmm::device_uvector<int32_t> point_flags(num_points, stream);
+  rmm::device_uvector<uint8_t> point_flags(num_points, stream);
   detail::find_duplicate_points(
     make_multipoint_range(points.offset_range(), points.geom_range()), point_flags.begin(), stream);
 
@@ -275,7 +279,7 @@ linestring_intersection_result<T, index_t> pairwise_linestring_intersection(
     stream,
     mr);
 
-  auto offsets_buffer = detail::compute_offset_buffer<index_t>(*types_buffer, mr, stream);
+  auto offsets_buffer = detail::compute_offset_buffer<index_t, types_t>(*types_buffer, mr, stream);
 
   // Assemble the look-back ids.
   auto lhs_linestring_id =

cpp/tests/CMakeLists.txt

@@ -219,7 +219,8 @@ ConfigureTest(LINESTRING_INTERSECTION_TEST_EXP
     intersection/linestring_intersection_count_test.cu
     intersection/linestring_intersection_intermediates_remove_if_test.cu
     intersection/linestring_intersection_with_duplicates_test.cu
-    intersection/linestring_intersection_test.cu)
+    intersection/linestring_intersection_test.cu
+    intersection/linestring_intersection_large_test.cu)
 
 # nearest points
 ConfigureTest(POINT_LINESTRING_NEAREST_POINT_TEST_EXP

cpp/tests/intersection/intersection_test_utils.cuh

@@ -16,9 +16,16 @@
 
 #pragma once
 
+#include <cuspatial_test/vector_factories.cuh>
+
 #include <cuspatial/cuda_utils.hpp>
 #include <cuspatial/geometry/segment.cuh>
+#include <cuspatial/intersection.cuh>
 
+#include <thrust/copy.h>
+#include <thrust/iterator/zip_iterator.h>
+#include <thrust/scatter.h>
+#include <thrust/sort.h>
 #include <thrust/tuple.h>
 
 namespace cuspatial {
@@ -56,5 +63,156 @@ struct order_key_value_pairs {
   }
 };
 
+/**
+ * @brief Perform sorting to the intersection result
+ *
+ * The result of intersection result is non-determinisitc. This algorithm sorts
+ * the geometries of the same types and the same list and makes the result deterministic.
+ *
+ * The example below contains 2 rows and 4 geometries. The order of the first
+ * and second point is non-deterministic.
+ * [
+ *  [Point(1.0, 1.5), Point(0.0, -0.3), Segment((0.0, 0.0), (1.0, 1.0))]
+ *   ^                ^
+ *  [Point(-3, -5)]
+ * ]
+ *
+ * After sorting, the result is deterministic:
+ * [
+ *  [Point(0.0, -0.3), Point(1.0, 1.5), Segment((0.0, 0.0), (1.0, 1.0))]
+ *   ^                 ^
+ *  [Point(-3, -5)]
+ * ]
+ *
+ * This function invalidates the input @p result and return a copy of sorted results.
+ */
+template <typename T, typename IndexType, typename type_t>
+linestring_intersection_result<T, IndexType> segment_sort_intersection_result(
+  linestring_intersection_result<T, IndexType>& result,
+  rmm::mr::device_memory_resource* mr,
+  rmm::cuda_stream_view stream)
+{
+  auto const num_points   = result.points_coords->size();
+  auto const num_segments = result.segments_coords->size();
+  auto const num_geoms    = num_points + num_segments;
+
+  rmm::device_uvector<IndexType> scatter_map(num_geoms, stream);
+  thrust::sequence(rmm::exec_policy(stream), scatter_map.begin(), scatter_map.end());
+
+  // Compute keys for each row in the union column. Rows of the same list
+  // are assigned the same label.
+  rmm::device_uvector<IndexType> geometry_collection_keys(num_geoms, stream);
+  auto geometry_collection_keys_begin = make_geometry_id_iterator<IndexType>(
+    result.geometry_collection_offset->begin(), result.geometry_collection_offset->end());
+
+  thrust::copy(rmm::exec_policy(stream),
+               geometry_collection_keys_begin,
+               geometry_collection_keys_begin + num_geoms,
+               geometry_collection_keys.begin());
+
+  // Perform "group-by" based on the list label and type of the row -
+  // This makes the geometry of the same type and of the same list neighbor.
+
+  // Make a copy of types buffer so that the sorting does not affect the original.
+  auto types_buffer = rmm::device_uvector<type_t>(*result.types_buffer, stream);
+  auto keys_begin =
+    thrust::make_zip_iterator(types_buffer.begin(), geometry_collection_keys.begin());
+  auto value_begin = thrust::make_zip_iterator(scatter_map.begin(),
+                                               result.lhs_linestring_id->begin(),
+                                               result.lhs_segment_id->begin(),
+                                               result.rhs_linestring_id->begin(),
+                                               result.rhs_segment_id->begin());
+
+  thrust::sort_by_key(rmm::exec_policy(stream), keys_begin, keys_begin + num_geoms, value_begin);
+
+  // Segment-sort the point array
+  auto keys_points_begin = thrust::make_zip_iterator(keys_begin, result.points_coords->begin());
+  thrust::sort_by_key(rmm::exec_policy(stream),
+                      keys_points_begin,
+                      keys_points_begin + num_points,
+                      scatter_map.begin(),
+                      order_key_value_pairs<thrust::tuple<IndexType, IndexType>, vec_2d<T>>{});
+
+  // Segment-sort the segment array
+  auto keys_segment_begin =
+    thrust::make_zip_iterator(keys_begin + num_points, result.segments_coords->begin());
+
+  thrust::sort_by_key(rmm::exec_policy(stream),
+                      keys_segment_begin,
+                      keys_segment_begin + num_segments,
+                      scatter_map.begin() + num_points,
+                      order_key_value_pairs<thrust::tuple<IndexType, IndexType>, segment<T>>{});
+
+  // Restore the order of indices
+  auto lhs_linestring_id = std::make_unique<rmm::device_uvector<IndexType>>(num_geoms, stream, mr);
+  auto lhs_segment_id    = std::make_unique<rmm::device_uvector<IndexType>>(num_geoms, stream, mr);
+  auto rhs_linestring_id = std::make_unique<rmm::device_uvector<IndexType>>(num_geoms, stream, mr);
+  auto rhs_segment_id    = std::make_unique<rmm::device_uvector<IndexType>>(num_geoms, stream, mr);
+
+  auto input_it = thrust::make_zip_iterator(result.lhs_linestring_id->begin(),
+                                            result.lhs_segment_id->begin(),
+                                            result.rhs_linestring_id->begin(),
+                                            result.rhs_segment_id->begin());
+
+  auto output_it = thrust::make_zip_iterator(lhs_linestring_id->begin(),
+                                             lhs_segment_id->begin(),
+                                             rhs_linestring_id->begin(),
+                                             rhs_segment_id->begin());
+
+  thrust::scatter(
+    rmm::exec_policy(stream), input_it, input_it + num_geoms, scatter_map.begin(), output_it);
+
+  return {std::move(result.geometry_collection_offset),
+          std::move(result.types_buffer),
+          std::move(result.offset_buffer),
+          std::move(result.points_coords),
+          std::move(result.segments_coords),
+          std::move(lhs_linestring_id),
+          std::move(lhs_segment_id),
+          std::move(rhs_linestring_id),
+          std::move(rhs_segment_id)};
+}
+
+template <typename T,
+          typename IndexType,
+          typename types_t,
+          typename point_t   = vec_2d<T>,
+          typename segment_t = segment<T>>
+auto make_linestring_intersection_result(
+  std::initializer_list<IndexType> geometry_collection_offset,
+  std::initializer_list<types_t> types_buffer,
+  std::initializer_list<IndexType> offset_buffer,
+  std::initializer_list<point_t> points_coords,
+  std::initializer_list<segment_t> segments_coords,
+  std::initializer_list<IndexType> lhs_linestring_ids,
+  std::initializer_list<IndexType> lhs_segment_ids,
+  std::initializer_list<IndexType> rhs_linestring_ids,
+  std::initializer_list<IndexType> rhs_segment_ids,
+  rmm::cuda_stream_view stream,
+  rmm::mr::device_memory_resource* mr)
+{
+  auto d_geometry_collection_offset =
+    make_device_uvector<IndexType>(geometry_collection_offset, stream, mr);
+  auto d_types_buffer       = make_device_uvector<types_t>(types_buffer, stream, mr);
+  auto d_offset_buffer      = make_device_uvector<IndexType>(offset_buffer, stream, mr);
+  auto d_points_coords      = make_device_uvector<point_t>(points_coords, stream, mr);
+  auto d_segments_coords    = make_device_uvector<segment_t>(segments_coords, stream, mr);
+  auto d_lhs_linestring_ids = make_device_uvector<IndexType>(lhs_linestring_ids, stream, mr);
+  auto d_lhs_segment_ids    = make_device_uvector<IndexType>(lhs_segment_ids, stream, mr);
+  auto d_rhs_linestring_ids = make_device_uvector<IndexType>(rhs_linestring_ids, stream, mr);
+  auto d_rhs_segment_ids    = make_device_uvector<IndexType>(rhs_segment_ids, stream, mr);
+
+  return linestring_intersection_result<T, IndexType>{
+    std::make_unique<rmm::device_uvector<IndexType>>(d_geometry_collection_offset, stream),
+    std::make_unique<rmm::device_uvector<types_t>>(d_types_buffer, stream),
+    std::make_unique<rmm::device_uvector<IndexType>>(d_offset_buffer, stream),
+    std::make_unique<rmm::device_uvector<point_t>>(d_points_coords, stream),
+    std::make_unique<rmm::device_uvector<segment_t>>(d_segments_coords, stream),
+    std::make_unique<rmm::device_uvector<IndexType>>(d_lhs_linestring_ids, stream),
+    std::make_unique<rmm::device_uvector<IndexType>>(d_lhs_segment_ids, stream),
+    std::make_unique<rmm::device_uvector<IndexType>>(d_rhs_linestring_ids, stream),
+    std::make_unique<rmm::device_uvector<IndexType>>(d_rhs_segment_ids, stream)};
+}
+
 }  // namespace test
 }  // namespace cuspatial