From 545344b2f08eaff78ca833e36a7f63e1da7c95a1 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 28 Jun 2023 17:11:07 -0700 Subject: [PATCH 01/18] add long input test case --- .../linestring_intersection_large_test.cu | 30 +++++++++++++++++++ 1 file changed, 30 insertions(+) diff --git a/cpp/tests/intersection/linestring_intersection_large_test.cu b/cpp/tests/intersection/linestring_intersection_large_test.cu index aed4c00f3..3f7ba4fcd 100644 --- a/cpp/tests/intersection/linestring_intersection_large_test.cu +++ b/cpp/tests/intersection/linestring_intersection_large_test.cu @@ -27,6 +27,7 @@ #include #include +#include #include template @@ -2025,3 +2026,32 @@ TYPED_TEST(LinestringIntersectionLargeTest, LongInput) CUSPATIAL_RUN_TEST( this->template verify_legal_result, multilinestrings1.range(), multilinestrings2.range()); } + +template +struct coordinate_functor { + cuspatial::vec_2d __device__ operator()(std::size_t i) + { + return cuspatial::vec_2d{i / T{2.0}, i / T{2.0}}; + } +}; + +TYPED_TEST(LinestringIntersectionLargeTest, LongInput_2) +{ + using P = cuspatial::vec_2d; + auto geometry_offset = cuspatial::test::make_device_vector({0, 1}); + auto part_offset = cuspatial::test::make_device_vector({0, 130}); + auto coordinates = rmm::device_uvector

(260, this->stream()); + + thrust::tabulate(rmm::exec_policy(this->stream()), + coordinates.begin(), + thrust::next(coordinates.begin(), 128), + coordinate_functor{}); + + coordinates.set_element(128, P{127.0, 0.0}, this->stream()); + coordinates.set_element(129, P{0.0, 0.0}, this->stream()); + + auto rng = cuspatial::make_multilinestring_range( + 1, geometry_offset.begin(), 1, part_offset.begin(), 130, coordinates.begin()); + + CUSPATIAL_RUN_TEST(this->template verify_legal_result, rng, rng); +} From cbd803fa64b20bacf511a5846b19e181e6ea6cb6 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 28 Jun 2023 21:16:51 -0700 Subject: [PATCH 02/18] sort the segments before merging --- .../detail/find/find_and_combine_segment.cuh | 44 ++++++++++++++++++- .../intersection/linestring_intersection.cuh | 12 +++++ ...inestring_intersection_with_duplicates.cuh | 12 +++++ cpp/include/cuspatial/geometry/segment.cuh | 11 +++++ .../linestring_intersection_large_test.cu | 2 +- 5 files changed, 78 insertions(+), 3 deletions(-) diff --git a/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh b/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh index b05e71b03..d6fa8d5ea 100644 --- a/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh +++ b/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh @@ -19,18 +19,20 @@ #include #include #include +#include #include #include -#include +#include namespace cuspatial { namespace detail { /** * @internal - * @brief Kernel to merge segments, naive n^2 algorithm. + * @brief Kernel to merge segments. Each thread works on a pair of segment spaces. + * naive n^2 algorithm. */ template void __global__ simple_find_and_combine_segments_kernel(OffsetRange offsets, @@ -44,6 +46,9 @@ void __global__ simple_find_and_combine_segments_kernel(OffsetRange offsets, merged_flag[i] = 0; } + // For each of the segment, loop over the rest of the segment in the space and see + // if it is mergeable with the current segment. + // Note that if the current segment is already merged. Skip checking. for (auto i = offsets[pair_idx]; i < offsets[pair_idx + 1] && merged_flag[i] != 1; i++) { for (auto j = i + 1; j < offsets[pair_idx + 1]; j++) { auto res = maybe_merge_segments(segments[i], segments[j]); @@ -57,6 +62,29 @@ void __global__ simple_find_and_combine_segments_kernel(OffsetRange offsets, } } +template +struct segment_comparator { + bool __device__ operator()(thrust::tuple> const& lhs, + thrust::tuple> const& rhs) const + { + auto lhs_index = thrust::get<0>(lhs); + auto rhs_index = thrust::get<0>(rhs); + auto lhs_segment = thrust::get<1>(lhs); + auto rhs_segment = thrust::get<1>(rhs); + + // Compare space id + if (lhs_index == rhs_index) { + // Compare slope + if (lhs_segment.collinear(rhs_segment)) { + // Sort by the lower left point of the segment. + return lhs_segment.lower_left() < rhs_segment.lower_left(); + } + return lhs_segment.slope() < rhs_segment.slope(); + } + return lhs_index < rhs_index; + } +}; + /** * @internal * @brief For each pair of mergeable segment, overwrites the first segment with merged result, @@ -68,9 +96,21 @@ void find_and_combine_segment(OffsetRange offsets, OutputIt merged_flag, rmm::cuda_stream_view stream) { + using index_t = typename OffsetRange::value_type; + using T = typename SegmentRange::value_type::value_type; auto num_spaces = offsets.size() - 1; if (num_spaces == 0) return; + // Construct a key iterator using the offsets of the segment and the segment itself. + auto space_id_iter = make_geometry_id_iterator(offsets.begin(), offsets.end()); + auto space_id_segment_iter = thrust::make_zip_iterator(space_id_iter, segments.begin()); + + thrust::sort_by_key(rmm::exec_policy(stream), + space_id_segment_iter, + space_id_segment_iter + segments.size(), + segments.begin(), + segment_comparator{}); + auto [threads_per_block, num_blocks] = grid_1d(num_spaces); simple_find_and_combine_segments_kernel<<>>( offsets, segments, merged_flag); diff --git a/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh b/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh index 796365bde..c6c6e60ee 100644 --- a/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh +++ b/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh @@ -237,13 +237,25 @@ linestring_intersection_result pairwise_linestring_intersection( points.remove_if(range(point_flags.begin(), point_flags.end()), stream); if (segments.num_geoms() > 0) { + segments.debug_print(); // Merge mergeable segments rmm::device_uvector segment_flags(num_segments, stream); detail::find_and_combine_segment( segments.offset_range(), segments.geom_range(), segment_flags.begin(), stream); + std::cout << "\n\n"; + rmm::device_uvector segment_flags_int(num_segments, stream); + thrust::copy(rmm::exec_policy(stream), + segment_flags.begin(), + segment_flags.end(), + segment_flags_int.begin()); + test::print_device_vector(segment_flags_int, "segment_flags"); + std::cout << "\n\n"; + segments.remove_if(range(segment_flags.begin(), segment_flags.end()), stream); + segments.debug_print(); + // Reusing `point_flags` for merge point on segment primitive. // Notice that `point_flags` contains leftovers from previous `find_duplicate_points` call. // Here it does not need to be zero initialized because find_points_on_segments will overwrite diff --git a/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh b/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh index 8e35dfdca..2af03db64 100644 --- a/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh +++ b/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh @@ -16,6 +16,8 @@ #pragma once +#include + #include #include #include @@ -361,6 +363,16 @@ struct linestring_intersection_intermediates { /// Return the number of geometries in the intermediates auto num_geoms() { return geoms->size(); } + + void debug_print() + { + test::print_device_vector(*offsets, "offsets: "); + test::print_device_vector(*geoms, "geoms: "); + test::print_device_vector(*lhs_linestring_ids, "lhs_linestring_ids: "); + test::print_device_vector(*lhs_segment_ids, "lhs_segment_ids: "); + test::print_device_vector(*rhs_linestring_ids, "rhs_linestring_ids: "); + test::print_device_vector(*rhs_segment_ids, "rhs_segment_ids: "); + } }; /** diff --git a/cpp/include/cuspatial/geometry/segment.cuh b/cpp/include/cuspatial/geometry/segment.cuh index 9cfaf3b8e..9788d789a 100644 --- a/cpp/include/cuspatial/geometry/segment.cuh +++ b/cpp/include/cuspatial/geometry/segment.cuh @@ -55,6 +55,17 @@ class alignas(sizeof(Vertex)) segment { /// Return the length squared of segment. T CUSPATIAL_HOST_DEVICE length2() const { return dot(v2 - v1, v2 - v1); } + /// Return slope of segment. + T CUSPATIAL_HOST_DEVICE slope() { return (v2.y - v1.y) / (v2.x - v1.x); } + + /// Return the lower left vertex of segment. + Vertex CUSPATIAL_HOST_DEVICE lower_left() { return v1 < v2 ? v1 : v2; } + + bool CUSPATIAL_HOST_DEVICE collinear(segment const& other) + { + return (v1.x - v1.y) * (other.v2.x - other.v2.y) == (v2.x - v2.y) * (other.v1.x - other.v1.y); + } + private: friend std::ostream& operator<<(std::ostream& os, segment const& seg) { diff --git a/cpp/tests/intersection/linestring_intersection_large_test.cu b/cpp/tests/intersection/linestring_intersection_large_test.cu index 3f7ba4fcd..85c221678 100644 --- a/cpp/tests/intersection/linestring_intersection_large_test.cu +++ b/cpp/tests/intersection/linestring_intersection_large_test.cu @@ -2031,7 +2031,7 @@ template struct coordinate_functor { cuspatial::vec_2d __device__ operator()(std::size_t i) { - return cuspatial::vec_2d{i / T{2.0}, i / T{2.0}}; + return cuspatial::vec_2d{static_cast(i), static_cast(i)}; } }; From 04d551e855ee54e7f14b0be777d5b2afd4fb89d6 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 29 Jun 2023 11:58:58 -0700 Subject: [PATCH 03/18] add twospaces find and combine test --- .../find/find_and_combine_segments_test.cu | 22 +++++++++++++++++++ 1 file changed, 22 insertions(+) diff --git a/cpp/tests/find/find_and_combine_segments_test.cu b/cpp/tests/find/find_and_combine_segments_test.cu index ea01d7e53..e388bae73 100644 --- a/cpp/tests/find/find_and_combine_segments_test.cu +++ b/cpp/tests/find/find_and_combine_segments_test.cu @@ -281,3 +281,25 @@ TYPED_TEST(FindAndCombineSegmentsTest, nooverlap3) {0, 0}, {S{P{0.0, 0.0}, P{1.0, 1.0}}, S{P{0.0, 1.0}, P{1.0, 0.0}}}); } + +TYPED_TEST(FindAndCombineSegmentsTest, twospaces) +{ + using T = TypeParam; + using index_t = std::size_t; + using P = vec_2d; + using S = segment; + + auto segments = make_segment_array({0, 2, 4}, + {S{P{0.0, 0.0}, P{1.0, 1.0}}, + S{P{1.0, 1.0}, P{2.0, 2.0}}, + S{P{1.0, 1.0}, P{0.0, 0.0}}, + S{P{2.0, 2.0}, P{1.0, 1.0}}}); + + CUSPATIAL_RUN_TEST(this->run_single_test, + segments, + {0, 1, 0, 1}, + {S{P{0.0, 0.0}, P{2.0, 2.0}}, + S{P{1.0, 1.0}, P{2.0, 2.0}}, + S{P{0.0, 0.0}, P{2.0, 2.0}}, + S{P{2.0, 2.0}, P{1.0, 1.0}}}); +} From bbc2dce2d502edc650697be7e551852f202a1ae5 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 29 Jun 2023 12:18:23 -0700 Subject: [PATCH 04/18] add twospaces, non-contiguous input --- .../detail/find/find_and_combine_segment.cuh | 3 +++ .../find/find_and_combine_segments_test.cu | 22 +++++++++++++++++++ 2 files changed, 25 insertions(+) diff --git a/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh b/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh index d6fa8d5ea..1f436dbf0 100644 --- a/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh +++ b/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh @@ -16,6 +16,7 @@ #pragma once +#include "cuspatial_test/test_util.cuh" #include #include #include @@ -24,6 +25,8 @@ #include #include +#include + #include namespace cuspatial { diff --git a/cpp/tests/find/find_and_combine_segments_test.cu b/cpp/tests/find/find_and_combine_segments_test.cu index e388bae73..9905b53dd 100644 --- a/cpp/tests/find/find_and_combine_segments_test.cu +++ b/cpp/tests/find/find_and_combine_segments_test.cu @@ -303,3 +303,25 @@ TYPED_TEST(FindAndCombineSegmentsTest, twospaces) S{P{0.0, 0.0}, P{2.0, 2.0}}, S{P{2.0, 2.0}, P{1.0, 1.0}}}); } + +TYPED_TEST(FindAndCombineSegmentsTest, twospaces_non_contiguous_segments_with_empty) +{ + using T = TypeParam; + using index_t = std::size_t; + using P = vec_2d; + using S = segment; + + auto segments = make_segment_array({0, 4, 4}, + {S{P{1.0, 1.0}, P{2.0, 2.0}}, + S{P{3.0, 3.0}, P{4.0, 4.0}}, + S{P{0.0, 0.0}, P{1.0, 1.0}}, + S{P{2.0, 2.0}, P{3.0, 3.0}}}); + + CUSPATIAL_RUN_TEST(this->run_single_test, + segments, + {0, 1, 1, 1}, + {S{P{0.0, 0.0}, P{4.0, 4.0}}, + S{P{1.0, 1.0}, P{2.0, 2.0}}, + S{P{2.0, 2.0}, P{3.0, 3.0}}, + S{P{3.0, 3.0}, P{4.0, 4.0}}}); +} From 83d19af4d2a2d1dcf1624d29dc2302e6f84a664a Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 29 Jun 2023 12:19:07 -0700 Subject: [PATCH 05/18] Remove stale includes --- cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh | 3 --- 1 file changed, 3 deletions(-) diff --git a/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh b/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh index 1f436dbf0..d6fa8d5ea 100644 --- a/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh +++ b/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh @@ -16,7 +16,6 @@ #pragma once -#include "cuspatial_test/test_util.cuh" #include #include #include @@ -25,8 +24,6 @@ #include #include -#include - #include namespace cuspatial { From 4a822e53f56479764c3ce001c8f675b6a2726853 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 29 Jun 2023 12:27:58 -0700 Subject: [PATCH 06/18] add documentation --- .../detail/find/find_and_combine_segment.cuh | 11 +++++++++++ 1 file changed, 11 insertions(+) diff --git a/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh b/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh index d6fa8d5ea..6a0c27512 100644 --- a/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh +++ b/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh @@ -62,6 +62,14 @@ void __global__ simple_find_and_combine_segments_kernel(OffsetRange offsets, } } +/** + * @brief Comparator for sorting the segment range. + * + * This comparator makes sure that the segment range are sorted by the following keys: + * 1. Segments with the same space id are grouped together. + * 2. Segments within the same space are grouped by their slope. + * 3. Within each slope group, segments are sorted by their lower left point. + */ template struct segment_comparator { bool __device__ operator()(thrust::tuple> const& lhs, @@ -89,6 +97,9 @@ struct segment_comparator { * @internal * @brief For each pair of mergeable segment, overwrites the first segment with merged result, * sets the flag for the second segment as 1. + * + * @note This function will alter the input segment range by rearranging the order of the segments + * within each space so that merging kernel can take place. */ template void find_and_combine_segment(OffsetRange offsets, From 19fc646693f033fb9c0ea31f1ec14ec8c4dddfe4 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 29 Jun 2023 15:03:17 -0700 Subject: [PATCH 07/18] add test for overlapping and non-contiguous inputs --- .../find/find_and_combine_segments_test.cu | 18 ++++++++++++++++++ 1 file changed, 18 insertions(+) diff --git a/cpp/tests/find/find_and_combine_segments_test.cu b/cpp/tests/find/find_and_combine_segments_test.cu index 9905b53dd..764f0e678 100644 --- a/cpp/tests/find/find_and_combine_segments_test.cu +++ b/cpp/tests/find/find_and_combine_segments_test.cu @@ -325,3 +325,21 @@ TYPED_TEST(FindAndCombineSegmentsTest, twospaces_non_contiguous_segments_with_em S{P{2.0, 2.0}, P{3.0, 3.0}}, S{P{3.0, 3.0}, P{4.0, 4.0}}}); } + +TYPED_TEST(FindAndCombineSegmentsTest, onespace_non_contiguous_segments_overlaps) +{ + using T = TypeParam; + using index_t = std::size_t; + using P = vec_2d; + using S = segment; + + auto segments = make_segment_array( + {0, 3}, + {S{P{1.0, 1.0}, P{2.0, 2.0}}, S{P{4.0, 4.0}, P{5.0, 5.0}}, S{P{-1.0, -1.0}, P{4.0, 4.0}}}); + + CUSPATIAL_RUN_TEST( + this->run_single_test, + segments, + {0, 1, 1}, + {S{P{-1.0, -1.0}, P{5.0, 5.0}}, S{P{1.0, 1.0}, P{2.0, 2.0}}, S{P{4.0, 4.0}, P{5.0, 5.0}}}); +} From f0000076fe6abfe926c081260152270bff7d6538 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 29 Jun 2023 17:52:50 -0700 Subject: [PATCH 08/18] Remove stale debug prints --- .../detail/intersection/linestring_intersection.cuh | 12 ------------ .../linestring_intersection_with_duplicates.cuh | 12 ------------ 2 files changed, 24 deletions(-) diff --git a/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh b/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh index c6c6e60ee..796365bde 100644 --- a/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh +++ b/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh @@ -237,25 +237,13 @@ linestring_intersection_result pairwise_linestring_intersection( points.remove_if(range(point_flags.begin(), point_flags.end()), stream); if (segments.num_geoms() > 0) { - segments.debug_print(); // Merge mergeable segments rmm::device_uvector segment_flags(num_segments, stream); detail::find_and_combine_segment( segments.offset_range(), segments.geom_range(), segment_flags.begin(), stream); - std::cout << "\n\n"; - rmm::device_uvector segment_flags_int(num_segments, stream); - thrust::copy(rmm::exec_policy(stream), - segment_flags.begin(), - segment_flags.end(), - segment_flags_int.begin()); - test::print_device_vector(segment_flags_int, "segment_flags"); - std::cout << "\n\n"; - segments.remove_if(range(segment_flags.begin(), segment_flags.end()), stream); - segments.debug_print(); - // Reusing `point_flags` for merge point on segment primitive. // Notice that `point_flags` contains leftovers from previous `find_duplicate_points` call. // Here it does not need to be zero initialized because find_points_on_segments will overwrite diff --git a/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh b/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh index 2af03db64..8e35dfdca 100644 --- a/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh +++ b/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh @@ -16,8 +16,6 @@ #pragma once -#include - #include #include #include @@ -363,16 +361,6 @@ struct linestring_intersection_intermediates { /// Return the number of geometries in the intermediates auto num_geoms() { return geoms->size(); } - - void debug_print() - { - test::print_device_vector(*offsets, "offsets: "); - test::print_device_vector(*geoms, "geoms: "); - test::print_device_vector(*lhs_linestring_ids, "lhs_linestring_ids: "); - test::print_device_vector(*lhs_segment_ids, "lhs_segment_ids: "); - test::print_device_vector(*rhs_linestring_ids, "rhs_linestring_ids: "); - test::print_device_vector(*rhs_segment_ids, "rhs_segment_ids: "); - } }; /** From 3ec56347ab6b3f1979d308027feec10a2ce3a964 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 29 Jun 2023 17:53:02 -0700 Subject: [PATCH 09/18] Update documentation --- .../cuspatial/detail/find/find_and_combine_segment.cuh | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh b/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh index 6a0c27512..b5ca12480 100644 --- a/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh +++ b/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh @@ -31,8 +31,8 @@ namespace detail { /** * @internal - * @brief Kernel to merge segments. Each thread works on a pair of segment spaces. - * naive n^2 algorithm. + * @brief Kernel to merge segments. Each thread works on one segment space, + * with a naive n^2 algorithm. */ template void __global__ simple_find_and_combine_segments_kernel(OffsetRange offsets, From 8713c49b60ef3446feda50a3d4d553df46d5ca74 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 29 Jun 2023 17:55:36 -0700 Subject: [PATCH 10/18] remove large intersection test --- .../linestring_intersection_large_test.cu | 29 ------------------- 1 file changed, 29 deletions(-) diff --git a/cpp/tests/intersection/linestring_intersection_large_test.cu b/cpp/tests/intersection/linestring_intersection_large_test.cu index 85c221678..e67debb45 100644 --- a/cpp/tests/intersection/linestring_intersection_large_test.cu +++ b/cpp/tests/intersection/linestring_intersection_large_test.cu @@ -2026,32 +2026,3 @@ TYPED_TEST(LinestringIntersectionLargeTest, LongInput) CUSPATIAL_RUN_TEST( this->template verify_legal_result, multilinestrings1.range(), multilinestrings2.range()); } - -template -struct coordinate_functor { - cuspatial::vec_2d __device__ operator()(std::size_t i) - { - return cuspatial::vec_2d{static_cast(i), static_cast(i)}; - } -}; - -TYPED_TEST(LinestringIntersectionLargeTest, LongInput_2) -{ - using P = cuspatial::vec_2d; - auto geometry_offset = cuspatial::test::make_device_vector({0, 1}); - auto part_offset = cuspatial::test::make_device_vector({0, 130}); - auto coordinates = rmm::device_uvector

(260, this->stream()); - - thrust::tabulate(rmm::exec_policy(this->stream()), - coordinates.begin(), - thrust::next(coordinates.begin(), 128), - coordinate_functor{}); - - coordinates.set_element(128, P{127.0, 0.0}, this->stream()); - coordinates.set_element(129, P{0.0, 0.0}, this->stream()); - - auto rng = cuspatial::make_multilinestring_range( - 1, geometry_offset.begin(), 1, part_offset.begin(), 130, coordinates.begin()); - - CUSPATIAL_RUN_TEST(this->template verify_legal_result, rng, rng); -} From bac1a362f40579af4e0065d885386722364d92a8 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 13 Jul 2023 21:57:24 -0700 Subject: [PATCH 11/18] applied review comments --- .../detail/algorithm/is_point_in_polygon.cuh | 4 ++-- .../detail/find/find_and_combine_segment.cuh | 2 +- cpp/include/cuspatial/detail/utility/linestring.cuh | 12 ++++++------ .../tests/binpreds/test_contains_properly.py | 2 +- 4 files changed, 10 insertions(+), 10 deletions(-) diff --git a/cpp/include/cuspatial/detail/algorithm/is_point_in_polygon.cuh b/cpp/include/cuspatial/detail/algorithm/is_point_in_polygon.cuh index 54ad6b499..9e8d6583c 100644 --- a/cpp/include/cuspatial/detail/algorithm/is_point_in_polygon.cuh +++ b/cpp/include/cuspatial/detail/algorithm/is_point_in_polygon.cuh @@ -68,8 +68,8 @@ __device__ inline bool is_point_in_polygon(vec_2d const& test_point, PolygonR T run_to_point = test_point.x - a.x; // point-on-edge test - bool is_colinear = float_equal(run * rise_to_point, run_to_point * rise); - if (is_colinear) { + bool is_collinear = float_equal(run * rise_to_point, run_to_point * rise); + if (is_collinear) { T minx = a.x; T maxx = b.x; if (minx > maxx) thrust::swap(minx, maxx); diff --git a/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh b/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh index b5ca12480..66ab4c497 100644 --- a/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh +++ b/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh @@ -65,7 +65,7 @@ void __global__ simple_find_and_combine_segments_kernel(OffsetRange offsets, /** * @brief Comparator for sorting the segment range. * - * This comparator makes sure that the segment range are sorted by the following keys: + * This comparator makes sure that the segment range is sorted such that: * 1. Segments with the same space id are grouped together. * 2. Segments within the same space are grouped by their slope. * 3. Within each slope group, segments are sorted by their lower left point. diff --git a/cpp/include/cuspatial/detail/utility/linestring.cuh b/cpp/include/cuspatial/detail/utility/linestring.cuh index 94180a5df..239a100f8 100644 --- a/cpp/include/cuspatial/detail/utility/linestring.cuh +++ b/cpp/include/cuspatial/detail/utility/linestring.cuh @@ -93,10 +93,10 @@ __forceinline__ T __device__ point_to_segment_distance_squared(vec_2d const& * @brief Computes shortest distance between two segments (ab and cd) that don't intersect. */ template -__forceinline__ T __device__ segment_distance_no_intersect_or_colinear(vec_2d const& a, - vec_2d const& b, - vec_2d const& c, - vec_2d const& d) +__forceinline__ T __device__ segment_distance_no_intersect_or_collinear(vec_2d const& a, + vec_2d const& b, + vec_2d const& c, + vec_2d const& d) { auto dist_sqr = min( min(point_to_segment_distance_squared(a, c, d), point_to_segment_distance_squared(b, c, d)), @@ -123,7 +123,7 @@ __forceinline__ T __device__ squared_segment_distance(vec_2d const& a, if (float_equal(denom, T{0})) { // Segments parallel or collinear - return segment_distance_no_intersect_or_colinear(a, b, c, d); + return segment_distance_no_intersect_or_collinear(a, b, c, d); } auto ac = c - a; @@ -132,7 +132,7 @@ __forceinline__ T __device__ squared_segment_distance(vec_2d const& a, auto r = r_numer * denom_reciprocal; auto s = det(ac, ab) * denom_reciprocal; if (r >= 0 and r <= 1 and s >= 0 and s <= 1) { return 0.0; } - return segment_distance_no_intersect_or_colinear(a, b, c, d); + return segment_distance_no_intersect_or_collinear(a, b, c, d); } /** diff --git a/python/cuspatial/cuspatial/tests/binpreds/test_contains_properly.py b/python/cuspatial/cuspatial/tests/binpreds/test_contains_properly.py index e3a67df6c..a50094737 100644 --- a/python/cuspatial/cuspatial/tests/binpreds/test_contains_properly.py +++ b/python/cuspatial/cuspatial/tests/binpreds/test_contains_properly.py @@ -131,7 +131,7 @@ def test_float_precision_limits(point, polygon, expects): Unique success cases identified by @mharris. These go in a pair with test_float_precision_limits_failures because these are inconsistent results, where 0.6 fails above (as True, within the - polygon) and 0.66 below succeeds, though they are colinear. + polygon) and 0.66 below succeeds, though they are collinear. """ gpdpoint = gpd.GeoSeries(point) gpdpolygon = gpd.GeoSeries(polygon) From 03b83387fc051e7788c34a141d7e552e9b729f90 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Fri, 14 Jul 2023 06:47:29 +0000 Subject: [PATCH 12/18] Add streams to allocate_like call --- cpp/src/distance/haversine.cu | 2 +- cpp/src/trajectory/derive_trajectories.cu | 8 ++++---- 2 files changed, 5 insertions(+), 5 deletions(-) diff --git a/cpp/src/distance/haversine.cu b/cpp/src/distance/haversine.cu index 45c736c2f..3c4e631da 100644 --- a/cpp/src/distance/haversine.cu +++ b/cpp/src/distance/haversine.cu @@ -54,7 +54,7 @@ struct haversine_functor { if (a_lon.is_empty()) { return cudf::empty_like(a_lon); } auto mask_policy = cudf::mask_allocation_policy::NEVER; - auto result = cudf::allocate_like(a_lon, a_lon.size(), mask_policy, mr); + auto result = cudf::allocate_like(a_lon, a_lon.size(), mask_policy, stream, mr); auto lonlat_a = cuspatial::make_vec_2d_iterator(a_lon.begin(), a_lat.begin()); auto lonlat_b = cuspatial::make_vec_2d_iterator(b_lon.begin(), b_lat.begin()); diff --git a/cpp/src/trajectory/derive_trajectories.cu b/cpp/src/trajectory/derive_trajectories.cu index 0a5db39ed..8356369ef 100644 --- a/cpp/src/trajectory/derive_trajectories.cu +++ b/cpp/src/trajectory/derive_trajectories.cu @@ -47,10 +47,10 @@ struct derive_trajectories_dispatch { { auto cols = std::vector>{}; cols.reserve(4); - cols.push_back(cudf::allocate_like(object_id, cudf::mask_allocation_policy::NEVER, mr)); - cols.push_back(cudf::allocate_like(x, cudf::mask_allocation_policy::NEVER, mr)); - cols.push_back(cudf::allocate_like(y, cudf::mask_allocation_policy::NEVER, mr)); - cols.push_back(cudf::allocate_like(timestamp, cudf::mask_allocation_policy::NEVER, mr)); + cols.push_back(cudf::allocate_like(object_id, cudf::mask_allocation_policy::NEVER, stream, mr)); + cols.push_back(cudf::allocate_like(x, cudf::mask_allocation_policy::NEVER, stream, mr)); + cols.push_back(cudf::allocate_like(y, cudf::mask_allocation_policy::NEVER, stream, mr)); + cols.push_back(cudf::allocate_like(timestamp, cudf::mask_allocation_policy::NEVER, stream, mr)); auto points_begin = thrust::make_zip_iterator(x.begin(), y.begin()); auto points_out_begin = thrust::make_zip_iterator(cols[1]->mutable_view().begin(), From 571318c12894dfcc912728971960613be714961b Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Fri, 14 Jul 2023 06:47:29 +0000 Subject: [PATCH 13/18] Add streams to allocate_like call --- cpp/src/distance/haversine.cu | 2 +- cpp/src/trajectory/derive_trajectories.cu | 8 ++++---- 2 files changed, 5 insertions(+), 5 deletions(-) diff --git a/cpp/src/distance/haversine.cu b/cpp/src/distance/haversine.cu index 45c736c2f..3c4e631da 100644 --- a/cpp/src/distance/haversine.cu +++ b/cpp/src/distance/haversine.cu @@ -54,7 +54,7 @@ struct haversine_functor { if (a_lon.is_empty()) { return cudf::empty_like(a_lon); } auto mask_policy = cudf::mask_allocation_policy::NEVER; - auto result = cudf::allocate_like(a_lon, a_lon.size(), mask_policy, mr); + auto result = cudf::allocate_like(a_lon, a_lon.size(), mask_policy, stream, mr); auto lonlat_a = cuspatial::make_vec_2d_iterator(a_lon.begin(), a_lat.begin()); auto lonlat_b = cuspatial::make_vec_2d_iterator(b_lon.begin(), b_lat.begin()); diff --git a/cpp/src/trajectory/derive_trajectories.cu b/cpp/src/trajectory/derive_trajectories.cu index 0a5db39ed..8356369ef 100644 --- a/cpp/src/trajectory/derive_trajectories.cu +++ b/cpp/src/trajectory/derive_trajectories.cu @@ -47,10 +47,10 @@ struct derive_trajectories_dispatch { { auto cols = std::vector>{}; cols.reserve(4); - cols.push_back(cudf::allocate_like(object_id, cudf::mask_allocation_policy::NEVER, mr)); - cols.push_back(cudf::allocate_like(x, cudf::mask_allocation_policy::NEVER, mr)); - cols.push_back(cudf::allocate_like(y, cudf::mask_allocation_policy::NEVER, mr)); - cols.push_back(cudf::allocate_like(timestamp, cudf::mask_allocation_policy::NEVER, mr)); + cols.push_back(cudf::allocate_like(object_id, cudf::mask_allocation_policy::NEVER, stream, mr)); + cols.push_back(cudf::allocate_like(x, cudf::mask_allocation_policy::NEVER, stream, mr)); + cols.push_back(cudf::allocate_like(y, cudf::mask_allocation_policy::NEVER, stream, mr)); + cols.push_back(cudf::allocate_like(timestamp, cudf::mask_allocation_policy::NEVER, stream, mr)); auto points_begin = thrust::make_zip_iterator(x.begin(), y.begin()); auto points_out_begin = thrust::make_zip_iterator(cols[1]->mutable_view().begin(), From 76d066ea47ecdecba11bb4dc68594bdf2da87d8e Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Fri, 14 Jul 2023 07:23:55 +0000 Subject: [PATCH 14/18] address review changes --- cpp/tests/intersection/linestring_intersection_large_test.cu | 1 + 1 file changed, 1 insertion(+) diff --git a/cpp/tests/intersection/linestring_intersection_large_test.cu b/cpp/tests/intersection/linestring_intersection_large_test.cu index e67debb45..9c5833aa2 100644 --- a/cpp/tests/intersection/linestring_intersection_large_test.cu +++ b/cpp/tests/intersection/linestring_intersection_large_test.cu @@ -28,6 +28,7 @@ #include #include + #include template From 5478f37e55b0f28ace010124e85efde90af1e3d8 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Sun, 16 Jul 2023 23:45:37 -0700 Subject: [PATCH 15/18] fix collinear test bug --- cpp/include/cuspatial/geometry/segment.cuh | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/cpp/include/cuspatial/geometry/segment.cuh b/cpp/include/cuspatial/geometry/segment.cuh index 9788d789a..28974b5d5 100644 --- a/cpp/include/cuspatial/geometry/segment.cuh +++ b/cpp/include/cuspatial/geometry/segment.cuh @@ -61,9 +61,11 @@ class alignas(sizeof(Vertex)) segment { /// Return the lower left vertex of segment. Vertex CUSPATIAL_HOST_DEVICE lower_left() { return v1 < v2 ? v1 : v2; } + /// Returns true if two segments are on the same line + /// Test if the determinant of the matrix, of which the column vector is constructed from the segments is 0. bool CUSPATIAL_HOST_DEVICE collinear(segment const& other) { - return (v1.x - v1.y) * (other.v2.x - other.v2.y) == (v2.x - v2.y) * (other.v1.x - other.v1.y); + return (v1.x - v2.x) * (other.v1.y - other.v2.y) == (v1.y - v2.y) * (other.v1.x - other.v2.x); } private: From 37d380816b6e7ee2f84b2fd405b283a6a6226038 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Tue, 25 Jul 2023 09:30:38 +0000 Subject: [PATCH 16/18] style --- cpp/include/cuspatial/geometry/segment.cuh | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/cpp/include/cuspatial/geometry/segment.cuh b/cpp/include/cuspatial/geometry/segment.cuh index 28974b5d5..bc2ce86e8 100644 --- a/cpp/include/cuspatial/geometry/segment.cuh +++ b/cpp/include/cuspatial/geometry/segment.cuh @@ -62,7 +62,8 @@ class alignas(sizeof(Vertex)) segment { Vertex CUSPATIAL_HOST_DEVICE lower_left() { return v1 < v2 ? v1 : v2; } /// Returns true if two segments are on the same line - /// Test if the determinant of the matrix, of which the column vector is constructed from the segments is 0. + /// Test if the determinant of the matrix, of which the column vector is constructed from the + /// segments is 0. bool CUSPATIAL_HOST_DEVICE collinear(segment const& other) { return (v1.x - v2.x) * (other.v1.y - other.v2.y) == (v1.y - v2.y) * (other.v1.x - other.v2.x); From 5f3f9e1a4f7245e73a0a7b12804bd966acdf637c Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Tue, 25 Jul 2023 10:00:11 +0000 Subject: [PATCH 17/18] Add comments to explain algorithm complexity --- .../detail/find/find_and_combine_segment.cuh | 14 ++++++++++++-- 1 file changed, 12 insertions(+), 2 deletions(-) diff --git a/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh b/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh index 5e7bcd24b..5694088ad 100644 --- a/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh +++ b/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh @@ -34,8 +34,18 @@ namespace detail { /** * @internal - * @brief Kernel to merge segments. Each thread works on one segment space, - * with a naive n^2 algorithm. + * @pre All segments in range @p segments , it is presorted using `segment_comparator`. + * + * @brief Kernel to merge segments. Each thread works on one segment space, within each space, + * `segment_comparator` guarantees that segments with same slope are grouped together. We call + * each of such group is a "mergeable group". Within each mergeable group, the first segment is + * the "leading segment". The algorithm behave as follows: + * + * 1. For each mergeable group, loop over the rest of the segments in the group and see if it is + * mergeable with the leading segment. If it is, overwrite the leading segment with the merged + * result. Then mark the segment as merged by setting the flag to 1. This makes sure the inner loop + * for each merged segment is not run again. + * 2. Repeat 1 until all mergeable group is processed. */ template void __global__ simple_find_and_combine_segments_kernel(OffsetRange offsets, From 6f13294a8d350430621d121a56d4847faa1390b3 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Tue, 25 Jul 2023 15:35:52 +0000 Subject: [PATCH 18/18] Change test case result order due to sorting --- cpp/tests/find/find_and_combine_segments_test.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/tests/find/find_and_combine_segments_test.cu b/cpp/tests/find/find_and_combine_segments_test.cu index 764f0e678..f7a584f51 100644 --- a/cpp/tests/find/find_and_combine_segments_test.cu +++ b/cpp/tests/find/find_and_combine_segments_test.cu @@ -279,7 +279,7 @@ TYPED_TEST(FindAndCombineSegmentsTest, nooverlap3) CUSPATIAL_RUN_TEST(this->run_single_test, segments, {0, 0}, - {S{P{0.0, 0.0}, P{1.0, 1.0}}, S{P{0.0, 1.0}, P{1.0, 0.0}}}); + {S{P{0.0, 1.0}, P{1.0, 0.0}}, S{P{0.0, 0.0}, P{1.0, 1.0}}}); } TYPED_TEST(FindAndCombineSegmentsTest, twospaces)