diff --git a/cpp/include/cuspatial/detail/geometry/polygon_ref.cuh b/cpp/include/cuspatial/detail/geometry/polygon_ref.cuh index 882ae22ae..66ad89003 100644 --- a/cpp/include/cuspatial/detail/geometry/polygon_ref.cuh +++ b/cpp/include/cuspatial/detail/geometry/polygon_ref.cuh @@ -59,13 +59,13 @@ CUSPATIAL_HOST_DEVICE auto polygon_ref::ring_end() co template CUSPATIAL_HOST_DEVICE auto polygon_ref::point_begin() const { - return _point_begin; + return thrust::next(_point_begin, *_ring_begin); } template CUSPATIAL_HOST_DEVICE auto polygon_ref::point_end() const { - return _point_end; + return thrust::next(_point_begin, *thrust::prev(_ring_end)); } template diff --git a/cpp/include/cuspatial/detail/geometry_collection/multilinestring_ref.cuh b/cpp/include/cuspatial/detail/geometry_collection/multilinestring_ref.cuh index 7a76f92b8..15472485f 100644 --- a/cpp/include/cuspatial/detail/geometry_collection/multilinestring_ref.cuh +++ b/cpp/include/cuspatial/detail/geometry_collection/multilinestring_ref.cuh @@ -74,13 +74,15 @@ CUSPATIAL_HOST_DEVICE auto multilinestring_ref::part_ template CUSPATIAL_HOST_DEVICE auto multilinestring_ref::point_begin() const { - return _point_begin; + return thrust::next(_point_begin, *_part_begin); } template CUSPATIAL_HOST_DEVICE auto multilinestring_ref::point_end() const { - return _point_end; + // _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)); } template diff --git a/cpp/include/cuspatial/detail/geometry_collection/multipolygon_ref.cuh b/cpp/include/cuspatial/detail/geometry_collection/multipolygon_ref.cuh index a663a2dde..e13b786ba 100644 --- a/cpp/include/cuspatial/detail/geometry_collection/multipolygon_ref.cuh +++ b/cpp/include/cuspatial/detail/geometry_collection/multipolygon_ref.cuh @@ -112,14 +112,14 @@ template CUSPATIAL_HOST_DEVICE auto multipolygon_ref::point_begin() const { - return _point_begin; + return thrust::next(_point_begin, *thrust::next(_ring_begin, *_part_begin)); } template CUSPATIAL_HOST_DEVICE auto multipolygon_ref::point_end() const { - return _point_end; + return thrust::next(_point_begin, *thrust::next(_ring_begin, *thrust::prev(_part_end))); } template diff --git a/cpp/include/cuspatial/detail/range/multilinestring_range.cuh b/cpp/include/cuspatial/detail/range/multilinestring_range.cuh index a73d423ac..b9b53bfc0 100644 --- a/cpp/include/cuspatial/detail/range/multilinestring_range.cuh +++ b/cpp/include/cuspatial/detail/range/multilinestring_range.cuh @@ -200,12 +200,11 @@ multilinestring_range::is_valid_seg template template -CUSPATIAL_HOST_DEVICE thrust::pair< - vec_2d::element_t>, - vec_2d::element_t>> +CUSPATIAL_HOST_DEVICE auto multilinestring_range::segment(IndexType segment_idx) { - return thrust::make_pair(_point_begin[segment_idx], _point_begin[segment_idx + 1]); + using T = iterator_vec_base_type; + return cuspatial::segment{_point_begin[segment_idx], _point_begin[segment_idx + 1]}; } template diff --git a/cpp/include/cuspatial/detail/range/multilinestring_segment_range.cuh b/cpp/include/cuspatial/detail/range/multilinestring_segment_range.cuh index 028756347..b684146d3 100644 --- a/cpp/include/cuspatial/detail/range/multilinestring_segment_range.cuh +++ b/cpp/include/cuspatial/detail/range/multilinestring_segment_range.cuh @@ -200,7 +200,7 @@ class multilinestring_segment_range { CUSPATIAL_HOST_DEVICE auto multigeometry_offset_begin() { return thrust::make_permutation_iterator(_per_linestring_offset_begin(), - _parent.geometry_offsets_begin()); + _parent.geometry_offset_begin()); } /// Returns end iterator to the range of the starting segment index per multilinestring diff --git a/cpp/include/cuspatial/detail/range/multipolygon_range.cuh b/cpp/include/cuspatial/detail/range/multipolygon_range.cuh index a76803229..1e126fb89 100644 --- a/cpp/include/cuspatial/detail/range/multipolygon_range.cuh +++ b/cpp/include/cuspatial/detail/range/multipolygon_range.cuh @@ -241,23 +241,6 @@ multipolygon_range:: thrust::prev(thrust::upper_bound(thrust::seq, _geometry_begin, _geometry_end, part_idx))); } -template -template -CUSPATIAL_HOST_DEVICE auto -multipolygon_range:: - geometry_idx_from_segment_idx(IndexType segment_idx) -{ - auto ring_idx = ring_idx_from_point_idx(segment_idx); - if (!is_valid_segment_id(segment_idx, ring_idx)) - return multipolygon_range:: - INVALID_INDEX; - - return geometry_idx_from_part_idx(part_idx_from_ring_idx(ring_idx)); -} - template ::o return multipolygon_begin()[multipolygon_idx]; } -template -template -CUSPATIAL_HOST_DEVICE auto -multipolygon_range::get_segment( - IndexType segment_idx) -{ - return segment{_point_begin[segment_idx], _point_begin[segment_idx + 1]}; -} - template -template -CUSPATIAL_HOST_DEVICE bool -multipolygon_range:: - is_first_point_of_multipolygon(IndexType1 point_idx, IndexType2 geometry_idx) -{ - return point_idx == _ring_begin[_part_begin[_geometry_begin[geometry_idx]]]; -} - template +#include + namespace cuspatial { namespace detail { diff --git a/cpp/include/cuspatial/iterator_factory.cuh b/cpp/include/cuspatial/iterator_factory.cuh index 1f026512c..9b22ba9bc 100644 --- a/cpp/include/cuspatial/iterator_factory.cuh +++ b/cpp/include/cuspatial/iterator_factory.cuh @@ -16,6 +16,7 @@ #pragma once +#include #include #include #include @@ -424,6 +425,13 @@ auto make_geometry_id_iterator(GeometryIter geometry_offsets_begin, std::distance(geometry_offsets_begin, geometry_offsets_end))); } +template +auto make_count_iterator_from_offset_iterator(OffsetIterator it) +{ + auto zipped_offsets_it = thrust::make_zip_iterator(it, thrust::next(it)); + return thrust::make_transform_iterator(zipped_offsets_it, detail::offset_pair_to_count_functor{}); +} + /** * @} // end of doxygen group */ diff --git a/cpp/include/cuspatial/range/multilinestring_range.cuh b/cpp/include/cuspatial/range/multilinestring_range.cuh index b4bb6eaf5..b4b396ee5 100644 --- a/cpp/include/cuspatial/range/multilinestring_range.cuh +++ b/cpp/include/cuspatial/range/multilinestring_range.cuh @@ -101,6 +101,12 @@ class multilinestring_range { /// Return the iterator to the one past the last point in the range. CUSPATIAL_HOST_DEVICE auto point_end() { return _point_end; } + /// Return the iterator to the first geometry offset in the range. + CUSPATIAL_HOST_DEVICE auto geometry_offset_begin() { return _geometry_begin; } + + /// Return the iterator to the one past the last geometry offset in the range. + CUSPATIAL_HOST_DEVICE auto geometry_offset_end() { return _geometry_end; } + /// Return the iterator to the first part offset in the range. CUSPATIAL_HOST_DEVICE auto part_offset_begin() { return _part_begin; } @@ -144,8 +150,7 @@ class multilinestring_range { /// Returns the segment given a segment index. template - CUSPATIAL_HOST_DEVICE thrust::pair, vec_2d> segment( - IndexType segment_idx); + CUSPATIAL_HOST_DEVICE auto segment(IndexType segment_idx); /// Returns an iterator to the counts of points per multilinestring CUSPATIAL_HOST_DEVICE auto multilinestring_point_count_begin(); @@ -168,13 +173,6 @@ class multilinestring_range { template CUSPATIAL_HOST_DEVICE auto operator[](IndexType multilinestring_idx); - /// Raw offsets iterator - - CUSPATIAL_HOST_DEVICE auto geometry_offsets_begin() { return _geometry_begin; } - CUSPATIAL_HOST_DEVICE auto geometry_offsets_end() { return _geometry_end; } - CUSPATIAL_HOST_DEVICE auto part_offsets_begin() { return _part_begin; } - CUSPATIAL_HOST_DEVICE auto part_offsets_end() { return _part_end; } - /// Range Casts /// Casts the multilinestring range into a multipoint range. diff --git a/cpp/include/cuspatial/range/multipoint_range.cuh b/cpp/include/cuspatial/range/multipoint_range.cuh index 92f7bff76..caee16ad4 100644 --- a/cpp/include/cuspatial/range/multipoint_range.cuh +++ b/cpp/include/cuspatial/range/multipoint_range.cuh @@ -149,6 +149,7 @@ class multipoint_range { /** * @brief Returns `true` if the range contains only single points + * Undefined behavior if the range is an empty range. */ CUSPATIAL_HOST_DEVICE bool is_single_point_range(); diff --git a/cpp/include/cuspatial/range/multipolygon_range.cuh b/cpp/include/cuspatial/range/multipolygon_range.cuh index 2af70e066..1aae07346 100644 --- a/cpp/include/cuspatial/range/multipolygon_range.cuh +++ b/cpp/include/cuspatial/range/multipolygon_range.cuh @@ -72,8 +72,6 @@ class multipolygon_range { using index_t = iterator_value_type; using element_t = iterator_vec_base_type; - int64_t static constexpr INVALID_INDEX = -1; - multipolygon_range(GeometryIterator geometry_begin, GeometryIterator geometry_end, PartIterator part_begin, @@ -117,10 +115,10 @@ class multipolygon_range { CUSPATIAL_HOST_DEVICE auto point_end(); /// Return the iterator to the first geometry offset in the range. - CUSPATIAL_HOST_DEVICE auto geometry_offsets_begin() { return _part_begin; } + CUSPATIAL_HOST_DEVICE auto geometry_offset_begin() { return _part_begin; } /// Return the iterator to the one past the last geometry offset in the range. - CUSPATIAL_HOST_DEVICE auto geometry_offsets_end() { return _part_end; } + CUSPATIAL_HOST_DEVICE auto geometry_offset_end() { return _part_end; } /// Return the iterator to the first part offset in the range. CUSPATIAL_HOST_DEVICE auto part_offset_begin() { return _part_begin; } @@ -134,13 +132,6 @@ class multipolygon_range { /// Return the iterator to the one past the last ring offset in the range. CUSPATIAL_HOST_DEVICE auto ring_offset_end() { return _ring_end; } - /// Given the index of a segment, return the index of the geometry (multipolygon) that contains - /// the segment. Segment index is the index to the starting point of the segment. If the index is - /// the last point of the ring, then it is not a valid index. This function returns - /// multipolygon_range::INVALID_INDEX if the index is invalid. - template - CUSPATIAL_HOST_DEVICE auto geometry_idx_from_segment_idx(IndexType segment_idx); - /// Given the index of a point, return the index of the ring that contains the point. template CUSPATIAL_HOST_DEVICE auto ring_idx_from_point_idx(IndexType point_idx); @@ -158,16 +149,6 @@ class multipolygon_range { template CUSPATIAL_HOST_DEVICE auto operator[](IndexType multipolygon_idx); - /// Returns the `segment_idx`th segment in the multipolygon range. - template - CUSPATIAL_HOST_DEVICE auto get_segment(IndexType segment_idx); - - /// Returns `true` if `point_idx`th point is the first point of `geometry_idx`th - /// multipolygon - template - CUSPATIAL_HOST_DEVICE bool is_first_point_of_multipolygon(IndexType1 point_idx, - IndexType2 geometry_idx); - /// Returns an iterator to the number of points of the first multipolygon /// @note The count includes the duplicate first and last point of the ring. CUSPATIAL_HOST_DEVICE auto multipolygon_point_count_begin(); diff --git a/cpp/include/cuspatial_test/geometry_generator.cuh b/cpp/include/cuspatial_test/geometry_generator.cuh index 2fd3be28c..f966664fb 100644 --- a/cpp/include/cuspatial_test/geometry_generator.cuh +++ b/cpp/include/cuspatial_test/geometry_generator.cuh @@ -354,7 +354,7 @@ auto generate_multilinestring_array(multilinestring_generator_parameter param params.origin, detail::random_walk_functor{params.segment_length}); - return make_multilinestring_array( + return make_multilinestring_array( std::move(geometry_offset), std::move(part_offset), std::move(points)); } diff --git a/cpp/include/cuspatial_test/vector_factories.cuh b/cpp/include/cuspatial_test/vector_factories.cuh index 91d2c04d7..cbf1d8f56 100644 --- a/cpp/include/cuspatial_test/vector_factories.cuh +++ b/cpp/include/cuspatial_test/vector_factories.cuh @@ -101,6 +101,17 @@ class multipolygon_array { { } + multipolygon_array(rmm::device_vector&& geometry_offsets_array, + rmm::device_vector&& part_offsets_array, + rmm::device_vector&& ring_offsets_array, + rmm::device_vector&& coordinates_array) + : _geometry_offsets_array(std::move(geometry_offsets_array)), + _part_offsets_array(std::move(part_offsets_array)), + _ring_offsets_array(std::move(ring_offsets_array)), + _coordinates_array(std::move(coordinates_array)) + { + } + multipolygon_array(rmm::device_uvector&& geometry_offsets_array, rmm::device_uvector&& part_offsets_array, rmm::device_uvector&& ring_offsets_array, @@ -230,9 +241,9 @@ class multilinestring_array { multilinestring_array(GeometryArray geometry_offsets_array, PartArray part_offsets_array, CoordinateArray coordinate_array) - : _geometry_offset_array(geometry_offsets_array), - _part_offset_array(part_offsets_array), - _coordinate_array(coordinate_array) + : _geometry_offset_array(std::move(geometry_offsets_array)), + _part_offset_array(std::move(part_offsets_array)), + _coordinate_array(std::move(coordinate_array)) { } @@ -264,29 +275,41 @@ class multilinestring_array { }; /** - * @brief Construct an owning object of a multilinestring array from ranges + * @brief Construct an owning object of a multilinestring array from `device_uvectors` * * @param geometry_inl Range of geometry offsets * @param part_inl Range of part offsets * @param coord_inl Ramge of coordinate * @return multilinestring array object */ -template -auto make_multilinestring_array(IndexRangeA geometry_inl, - IndexRangeB part_inl, - CoordRange coord_inl) +template +auto make_multilinestring_array(rmm::device_uvector&& geometry_inl, + rmm::device_uvector&& part_inl, + rmm::device_uvector>&& coord_inl) { - using CoordType = typename CoordRange::value_type; - using DeviceIndexVector = thrust::device_vector; - using DeviceCoordVector = thrust::device_vector; + return multilinestring_array, + rmm::device_uvector, + rmm::device_uvector>>( + std::move(geometry_inl), std::move(part_inl), std::move(coord_inl)); +} - return multilinestring_array( - make_device_vector(std::move(geometry_inl)), - make_device_vector(std::move(part_inl)), - make_device_vector(std::move(coord_inl))); +/** + * @brief Construct an owning object of a multilinestring array from `device_vectors` + * + * @param geometry_inl Range of geometry offsets + * @param part_inl Range of part offsets + * @param coord_inl Ramge of coordinate + * @return multilinestring array object + */ +template +auto make_multilinestring_array(rmm::device_vector&& geometry_inl, + rmm::device_vector&& part_inl, + rmm::device_vector>&& coord_inl) +{ + return multilinestring_array, + rmm::device_vector, + rmm::device_vector>>( + std::move(geometry_inl), std::move(part_inl), std::move(coord_inl)); } /** @@ -414,5 +437,17 @@ auto make_multipoint_array(rmm::device_uvector geometry_offsets, std::move(geometry_offsets), std::move(coords)}; } +/** + * @brief Factory method to construct multipoint array by moving the offsets and coordinates from + * `rmm::device_vector`. + */ +template +auto make_multipoint_array(rmm::device_vector geometry_offsets, + rmm::device_vector> coords) +{ + return multipoint_array, rmm::device_vector>>{ + std::move(geometry_offsets), std::move(coords)}; +} + } // namespace test } // namespace cuspatial diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index b67b214ce..7f7cb112a 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -171,6 +171,7 @@ ConfigureTest(SINUSOIDAL_PROJECTION_TEST_EXP # range ConfigureTest(RANGE_TEST_EXP + range/multipoint_range_test.cu range/multilinestring_range_test.cu range/multipolygon_range_test.cu) diff --git a/cpp/tests/range/multilinestring_range_test.cu b/cpp/tests/range/multilinestring_range_test.cu index 250e88355..11c17e5ac 100644 --- a/cpp/tests/range/multilinestring_range_test.cu +++ b/cpp/tests/range/multilinestring_range_test.cu @@ -19,17 +19,31 @@ #include #include +#include #include #include +#include +#include +#include #include #include +#include +#include + #include +#include using namespace cuspatial; using namespace cuspatial::test; +template +void __global__ array_access_tester(MultiLineStringRange mls, std::size_t i, OutputIt output_points) +{ + thrust::copy(thrust::seq, mls[i].point_begin(), mls[i].point_end(), output_points); +} + template struct MultilinestringRangeTest : public BaseFixture { void run_segment_test_single(std::initializer_list geometry_offset, @@ -533,3 +547,807 @@ TYPED_TEST(MultilinestringRangeTest, MultilinestringAsMultipointTest6) {P{1, 1}, P{0, 0}, P{6, 6}, P{6, 7}}, {{P{1, 1}, P{0, 0}, P{6, 6}, P{6, 7}}}); } + +template +class MultilinestringRangeTestBase : public BaseFixture { + public: + struct copy_leading_point_per_multilinestring { + template + vec_2d __device__ operator()(MultiLineStringRef m) + { + return m.size() > 0 ? m[0].point_begin()[0] : vec_2d{-1, -1}; + } + }; + + template + struct part_idx_from_point_idx_functor { + MultiLineStringRange _rng; + std::size_t __device__ operator()(std::size_t point_idx) + { + return _rng.part_idx_from_point_idx(point_idx); + } + }; + + template + struct part_idx_from_segment_idx_functor { + MultiLineStringRange _rng; + std::size_t __device__ operator()(std::size_t segment_idx) + { + auto opt = _rng.part_idx_from_segment_idx(segment_idx); + if (opt.has_value()) { + return opt.value(); + } else { + return std::numeric_limits::max(); + } + } + }; + + template + struct geometry_idx_from_point_idx_functor { + MultiLineStringRange _rng; + std::size_t __device__ operator()(std::size_t point_idx) + { + return _rng.geometry_idx_from_point_idx(point_idx); + } + }; + + template + struct intra_part_idx_functor { + MultiLineStringRange _rng; + std::size_t __device__ operator()(std::size_t i) { return _rng.intra_part_idx(i); } + }; + + template + struct intra_point_idx_functor { + MultiLineStringRange _rng; + std::size_t __device__ operator()(std::size_t i) { return _rng.intra_point_idx(i); } + }; + + template + struct is_valid_segment_id_functor { + MultiLineStringRange _rng; + bool __device__ operator()(std::size_t i) + { + auto part_idx = _rng.part_idx_from_point_idx(i); + return _rng.is_valid_segment_id(i, part_idx); + } + }; + + template + struct segment_functor { + MultiLineStringRange _rng; + segment __device__ operator()(std::size_t i) + { + auto part_idx = _rng.part_idx_from_point_idx(i); + return _rng.is_valid_segment_id(i, part_idx) + ? _rng.segment(i) + : segment{vec_2d{-1, -1}, vec_2d{-1, -1}}; + } + }; + + void SetUp() { make_test_multilinestring(); } + + virtual void make_test_multilinestring() = 0; + + auto range() { return test_multilinestring->range(); } + + void run_test() + { + test_size(); + + test_num_multilinestrings(); + + test_num_linestrings(); + + test_num_points(); + + test_multilinestring_it(); + + test_begin(); + + test_end(); + + test_point_it(); + + test_part_offset_it(); + + test_part_idx_from_point_idx(); + + test_part_idx_from_segment_idx(); + + test_geometry_idx_from_point_idx(); + + test_intra_part_idx(); + + test_intra_point_idx(); + + test_is_valid_segment_id(); + + test_segment(); + + test_multilinestring_point_count_it(); + + test_multilinestring_linestring_count_it(); + + test_array_access_operator(); + + test_geometry_offset_it(); + } + + void test_size() { EXPECT_EQ(this->range().size(), this->range().num_multilinestrings()); } + + virtual void test_num_multilinestrings() = 0; + + virtual void test_num_linestrings() = 0; + + virtual void test_num_points() = 0; + + virtual void test_multilinestring_it() = 0; + + void test_begin() { EXPECT_EQ(this->range().begin(), this->range().multilinestring_begin()); } + + void test_end() { EXPECT_EQ(this->range().end(), this->range().multilinestring_end()); } + + virtual void test_point_it() = 0; + + virtual void test_geometry_offset_it() = 0; + + virtual void test_part_offset_it() = 0; + + virtual void test_part_idx_from_point_idx() = 0; + + virtual void test_part_idx_from_segment_idx() = 0; + + virtual void test_geometry_idx_from_point_idx() = 0; + + virtual void test_intra_part_idx() = 0; + + virtual void test_intra_point_idx() = 0; + + virtual void test_is_valid_segment_id() = 0; + + virtual void test_segment() = 0; + + virtual void test_multilinestring_point_count_it() = 0; + + virtual void test_multilinestring_linestring_count_it() = 0; + + virtual void test_array_access_operator() = 0; + + // Helper functions to be used by all subclass (test cases). + rmm::device_uvector> copy_leading_points() + { + auto rng = this->range(); + auto d_leading_point = rmm::device_uvector>(rng.num_multilinestrings(), stream()); + thrust::transform(rmm::exec_policy(stream()), + rng.begin(), + rng.end(), + d_leading_point.begin(), + copy_leading_point_per_multilinestring()); + return d_leading_point; + } + + rmm::device_uvector> copy_all_points() + { + auto rng = this->range(); + auto d_all_points = rmm::device_uvector>(rng.num_points(), stream()); + thrust::copy( + rmm::exec_policy(stream()), rng.point_begin(), rng.point_end(), d_all_points.begin()); + return d_all_points; + } + + rmm::device_uvector copy_geometry_offsets() + { + auto rng = this->range(); + auto d_geometry_offsets = + rmm::device_uvector(rng.num_multilinestrings() + 1, stream()); + thrust::copy(rmm::exec_policy(stream()), + rng.geometry_offset_begin(), + rng.geometry_offset_end(), + d_geometry_offsets.begin()); + return d_geometry_offsets; + } + + rmm::device_uvector copy_part_offset() + { + auto rng = this->range(); + auto d_part_offset = rmm::device_uvector(rng.num_linestrings() + 1, stream()); + thrust::copy(rmm::exec_policy(stream()), + rng.part_offset_begin(), + rng.part_offset_end(), + d_part_offset.begin()); + return d_part_offset; + } + + rmm::device_uvector copy_part_idx_from_point_idx() + { + auto rng = this->range(); + auto d_part_idx = rmm::device_uvector(rng.num_points(), stream()); + auto f = part_idx_from_point_idx_functor{rng}; + thrust::tabulate(rmm::exec_policy(stream()), d_part_idx.begin(), d_part_idx.end(), f); + return d_part_idx; + } + + rmm::device_uvector copy_part_idx_from_segment_idx() + { + auto rng = this->range(); + auto d_part_idx = rmm::device_uvector(rng.num_points(), stream()); + auto f = part_idx_from_segment_idx_functor{rng}; + + thrust::tabulate(rmm::exec_policy(stream()), d_part_idx.begin(), d_part_idx.end(), f); + return d_part_idx; + } + + rmm::device_uvector copy_geometry_idx_from_point_idx() + { + auto rng = this->range(); + auto d_geometry_idx = rmm::device_uvector(rng.num_points(), stream()); + thrust::tabulate(rmm::exec_policy(stream()), + d_geometry_idx.begin(), + d_geometry_idx.end(), + geometry_idx_from_point_idx_functor{rng}); + return d_geometry_idx; + } + + rmm::device_uvector copy_intra_part_idx() + { + auto rng = this->range(); + auto d_intra_part_idx = rmm::device_uvector(rng.num_linestrings(), stream()); + thrust::tabulate(rmm::exec_policy(stream()), + d_intra_part_idx.begin(), + d_intra_part_idx.end(), + intra_part_idx_functor{rng}); + return d_intra_part_idx; + } + + rmm::device_uvector copy_intra_point_idx() + { + auto rng = this->range(); + auto d_intra_point_idx = rmm::device_uvector(rng.num_points(), stream()); + thrust::tabulate(rmm::exec_policy(stream()), + d_intra_point_idx.begin(), + d_intra_point_idx.end(), + intra_point_idx_functor{rng}); + return d_intra_point_idx; + } + + rmm::device_uvector copy_is_valid_segment_id() + { + auto rng = this->range(); + auto d_is_valid_segment_id = rmm::device_uvector(rng.num_points(), stream()); + thrust::tabulate(rmm::exec_policy(stream()), + d_is_valid_segment_id.begin(), + d_is_valid_segment_id.end(), + is_valid_segment_id_functor{rng}); + return d_is_valid_segment_id; + } + + rmm::device_uvector> copy_segments() + { + auto rng = this->range(); + auto d_segments = rmm::device_uvector>(rng.num_points(), stream()); + thrust::tabulate(rmm::exec_policy(stream()), + d_segments.begin(), + d_segments.end(), + segment_functor{rng}); + return d_segments; + } + + rmm::device_uvector copy_multilinestring_point_count() + { + auto rng = this->range(); + auto d_multilinestring_point_count = + rmm::device_uvector(rng.num_multilinestrings(), stream()); + thrust::copy(rmm::exec_policy(stream()), + rng.multilinestring_point_count_begin(), + rng.multilinestring_point_count_end(), + d_multilinestring_point_count.begin()); + return d_multilinestring_point_count; + } + + rmm::device_uvector copy_multilinestring_linestring_count() + { + auto rng = this->range(); + auto d_multilinestring_linestring_count = + rmm::device_uvector(rng.num_multilinestrings(), stream()); + thrust::copy(rmm::exec_policy(stream()), + rng.multilinestring_linestring_count_begin(), + rng.multilinestring_linestring_count_end(), + d_multilinestring_linestring_count.begin()); + return d_multilinestring_linestring_count; + } + + rmm::device_uvector> copy_all_points_of_ith_multilinestring(std::size_t i) + { + auto rng = this->range(); + rmm::device_scalar num_points(stream()); + + thrust::copy_n(rmm::exec_policy(stream()), + rng.multilinestring_point_count_begin() + i, + 1, + num_points.data()); + + auto d_all_points = rmm::device_uvector>(num_points.value(stream()), stream()); + + array_access_tester<<<1, 1, 0, stream()>>>(rng, i, d_all_points.data()); + return d_all_points; + } + + protected: + std::unique_ptr, + rmm::device_vector, + rmm::device_vector>>> + test_multilinestring; +}; + +template +class MultilinestringRangeEmptyTest : public MultilinestringRangeTestBase { + void make_test_multilinestring() + { + auto array = make_multilinestring_array({0}, {0}, {}); + this->test_multilinestring = std::make_unique(std::move(array)); + } + + void test_num_multilinestrings() { EXPECT_EQ(this->range().num_multilinestrings(), 0); } + + void test_num_linestrings() { EXPECT_EQ(this->range().num_linestrings(), 0); } + + void test_num_points() { EXPECT_EQ(this->range().num_points(), 0); } + + void test_multilinestring_it() + { + auto leading_points = this->copy_leading_points(); + auto expected = rmm::device_uvector>(0, this->stream()); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(leading_points, expected); + } + + void test_point_it() + { + auto all_points = this->copy_all_points(); + auto expected = rmm::device_uvector>(0, this->stream()); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(all_points, expected); + } + + void test_part_offset_it() + { + auto part_offset = this->copy_part_offset(); + auto expected = make_device_vector({0}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_offset, expected); + } + + void test_part_idx_from_point_idx() + { + auto part_idx = this->copy_part_idx_from_point_idx(); + auto expected = rmm::device_vector(0); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_idx, expected); + } + + void test_part_idx_from_segment_idx() + { + auto part_idx = this->copy_part_idx_from_segment_idx(); + auto expected = rmm::device_vector(0); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_idx, expected); + } + + void test_geometry_idx_from_point_idx() + { + auto geometry_idx = this->copy_geometry_idx_from_point_idx(); + auto expected = rmm::device_vector(0); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_idx, expected); + } + + void test_intra_part_idx() + { + auto intra_part_idx = this->copy_intra_part_idx(); + auto expected = rmm::device_vector(0); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(intra_part_idx, expected); + } + + void test_intra_point_idx() + { + auto intra_point_idx = this->copy_intra_point_idx(); + auto expected = rmm::device_vector(0); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(intra_point_idx, expected); + } + + void test_is_valid_segment_id() + { + auto is_valid_segment_id = this->copy_is_valid_segment_id(); + auto expected = rmm::device_vector(0); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(is_valid_segment_id, expected); + } + + void test_segment() + { + auto segments = this->copy_segments(); + auto expected = rmm::device_vector>(0); + CUSPATIAL_EXPECT_VEC2D_PAIRS_EQUIVALENT(segments, expected); + } + + void test_multilinestring_point_count_it() + { + auto multilinestring_point_count = this->copy_multilinestring_point_count(); + auto expected = rmm::device_vector(0); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(multilinestring_point_count, expected); + } + + void test_multilinestring_linestring_count_it() + { + auto multilinestring_linestring_count = this->copy_multilinestring_linestring_count(); + auto expected = rmm::device_vector(0); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(multilinestring_linestring_count, expected); + } + + void test_array_access_operator() + { + // Nothing to access + SUCCEED(); + } + + void test_geometry_offset_it() + { + auto geometry_offsets = this->copy_geometry_offsets(); + auto expected = make_device_vector({0}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_offsets, expected); + } + + void test_part_offsets_it() + { + auto part_offsets = this->copy_part_offsets(); + auto expected = make_device_vector({0}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_offsets, expected); + } +}; + +TYPED_TEST_CASE(MultilinestringRangeEmptyTest, FloatingPointTypes); +TYPED_TEST(MultilinestringRangeEmptyTest, EmptyTest) { this->run_test(); } + +template +class MultilinestringRangeOneTest : public MultilinestringRangeTestBase { + void make_test_multilinestring() + { + auto array = make_multilinestring_array( + {0, 2}, {0, 2, 5}, {{10, 10}, {20, 20}, {100, 100}, {200, 200}, {300, 300}}); + this->test_multilinestring = std::make_unique(std::move(array)); + } + + void test_num_multilinestrings() { EXPECT_EQ(this->range().num_multilinestrings(), 1); } + + void test_num_linestrings() { EXPECT_EQ(this->range().num_linestrings(), 2); } + + void test_num_points() { EXPECT_EQ(this->range().num_points(), 5); } + + void test_multilinestring_it() + { + auto leading_points = this->copy_leading_points(); + auto expected = make_device_vector>({{10, 10}}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(leading_points, expected); + } + + void test_point_it() + { + auto all_points = this->copy_all_points(); + auto expected = + make_device_vector>({{10, 10}, {20, 20}, {100, 100}, {200, 200}, {300, 300}}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(all_points, expected); + } + + void test_part_offset_it() + { + auto part_offset = this->copy_part_offset(); + auto expected = make_device_vector({0, 2, 5}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_offset, expected); + } + + void test_part_idx_from_point_idx() + { + auto part_idx = this->copy_part_idx_from_point_idx(); + auto expected = make_device_vector({0, 0, 1, 1, 1}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_idx, expected); + } + + void test_part_idx_from_segment_idx() + { + auto part_idx = this->copy_part_idx_from_segment_idx(); + auto expected = make_device_vector( + {0, std::numeric_limits::max(), 1, 1, std::numeric_limits::max()}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_idx, expected); + } + + void test_geometry_idx_from_point_idx() + { + auto geometry_idx = this->copy_geometry_idx_from_point_idx(); + auto expected = make_device_vector({0, 0, 0, 0, 0}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_idx, expected); + } + + void test_intra_part_idx() + { + auto intra_part_idx = this->copy_intra_part_idx(); + auto expected = make_device_vector({0, 1}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(intra_part_idx, expected); + } + + void test_intra_point_idx() + { + auto intra_point_idx = this->copy_intra_point_idx(); + auto expected = make_device_vector({0, 1, 0, 1, 2}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(intra_point_idx, expected); + } + + void test_is_valid_segment_id() + { + auto is_valid_segment_id = this->copy_is_valid_segment_id(); + auto expected = make_device_vector({1, 0, 1, 1, 0}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(is_valid_segment_id, expected); + } + + void test_segment() + { + auto segments = this->copy_segments(); + auto expected = make_device_vector>({{{10, 10}, {20, 20}}, + {{-1, -1}, {-1, -1}}, + {{100, 100}, {200, 200}}, + {{200, 200}, {300, 300}}, + {{-1, -1}, {-1, -1}}}); + CUSPATIAL_EXPECT_VEC2D_PAIRS_EQUIVALENT(segments, expected); + } + + void test_multilinestring_point_count_it() + { + auto multilinestring_point_count = this->copy_multilinestring_point_count(); + auto expected = make_device_vector({5}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(multilinestring_point_count, expected); + } + + void test_multilinestring_linestring_count_it() + { + auto multilinestring_linestring_count = this->copy_multilinestring_linestring_count(); + auto expected = make_device_vector({2}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(multilinestring_linestring_count, expected); + } + + void test_array_access_operator() + { + auto all_points = this->copy_all_points_of_ith_multilinestring(0); + auto expected = + make_device_vector>({{10, 10}, {20, 20}, {100, 100}, {200, 200}, {300, 300}}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(all_points, expected); + } + + void test_geometry_offset_it() + { + auto geometry_offsets = this->copy_geometry_offsets(); + auto expected = make_device_vector({0, 2}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_offsets, expected); + } + + void test_part_offsets_it() + { + auto part_offsets = this->copy_part_offsets(); + auto expected = make_device_vector({0, 2, 5}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_offsets, expected); + } +}; + +TYPED_TEST_CASE(MultilinestringRangeOneTest, FloatingPointTypes); +TYPED_TEST(MultilinestringRangeOneTest, OneTest) { this->run_test(); } + +template +class MultilinestringRangeOneThousandTest : public MultilinestringRangeTestBase { + public: + struct make_points_functor { + vec_2d __device__ operator()(std::size_t i) + { + auto part_idx = i / 2; + auto intra_point_idx = i % 2; + return vec_2d{static_cast(part_idx * 10 + intra_point_idx), + static_cast(part_idx * 10 + intra_point_idx)}; + } + }; + + void make_test_multilinestring() + { + rmm::device_vector geometry_offset(1001); + rmm::device_vector part_offset(1001); + rmm::device_vector> points(2000); + + thrust::sequence( + rmm::exec_policy(this->stream()), geometry_offset.begin(), geometry_offset.end()); + + thrust::sequence( + rmm::exec_policy(this->stream()), part_offset.begin(), part_offset.end(), 0, 2); + + thrust::tabulate( + rmm::exec_policy(this->stream()), points.begin(), points.end(), make_points_functor{}); + + auto array = make_multilinestring_array( + std::move(geometry_offset), std::move(part_offset), std::move(points)); + + this->test_multilinestring = std::make_unique(std::move(array)); + } + + void test_num_multilinestrings() { EXPECT_EQ(this->range().num_multilinestrings(), 1000); } + + void test_num_linestrings() { EXPECT_EQ(this->range().num_linestrings(), 1000); } + + void test_num_points() { EXPECT_EQ(this->range().num_points(), 2000); } + + void test_multilinestring_it() + { + auto leading_points = this->copy_leading_points(); + auto expected = rmm::device_uvector>(1000, this->stream()); + + thrust::tabulate(rmm::exec_policy(this->stream()), + expected.begin(), + expected.end(), + [] __device__(std::size_t i) { + return vec_2d{i * T{10.}, i * T{10.}}; + }); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(leading_points, expected); + } + + void test_point_it() + { + auto all_points = this->copy_all_points(); + auto expected = rmm::device_uvector>(2000, this->stream()); + + thrust::tabulate( + rmm::exec_policy(this->stream()), expected.begin(), expected.end(), make_points_functor{}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(all_points, expected); + } + + void test_part_offset_it() + { + auto part_offset = this->copy_part_offset(); + auto expected = rmm::device_uvector(1001, this->stream()); + + thrust::sequence(rmm::exec_policy(this->stream()), expected.begin(), expected.end(), 0, 2); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_offset, expected); + } + + void test_part_idx_from_point_idx() + { + auto part_idx = this->copy_part_idx_from_point_idx(); + auto expected = rmm::device_uvector(2000, this->stream()); + + thrust::tabulate(rmm::exec_policy(this->stream()), + expected.begin(), + expected.end(), + [] __device__(std::size_t i) { return i / 2; }); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_idx, expected); + } + + void test_part_idx_from_segment_idx() + { + auto part_idx = this->copy_part_idx_from_segment_idx(); + auto expected = rmm::device_uvector(2000, this->stream()); + + thrust::tabulate(rmm::exec_policy(this->stream()), + expected.begin(), + expected.end(), + [] __device__(std::size_t i) { + return i % 2 == 0 ? i / 2 : std::numeric_limits::max(); + }); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_idx, expected); + } + + void test_geometry_idx_from_point_idx() + { + auto geometry_idx = this->copy_geometry_idx_from_point_idx(); + auto expected = rmm::device_uvector(2000, this->stream()); + + thrust::tabulate(rmm::exec_policy(this->stream()), + expected.begin(), + expected.end(), + [] __device__(std::size_t i) { return i / 2; }); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_idx, expected); + } + + void test_intra_part_idx() + { + auto intra_part_idx = this->copy_intra_part_idx(); + auto expected = rmm::device_uvector(1000, this->stream()); + + detail::zero_data_async(expected.begin(), expected.end(), this->stream()); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(intra_part_idx, expected); + } + + void test_intra_point_idx() + { + auto intra_point_idx = this->copy_intra_point_idx(); + auto expected = rmm::device_uvector(2000, this->stream()); + + thrust::tabulate(rmm::exec_policy(this->stream()), + expected.begin(), + expected.end(), + [] __device__(std::size_t i) { return i % 2; }); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(intra_point_idx, expected); + } + + void test_is_valid_segment_id() + { + auto is_valid_segment_id = this->copy_is_valid_segment_id(); + auto expected = rmm::device_uvector(2000, this->stream()); + + thrust::tabulate(rmm::exec_policy(this->stream()), + expected.begin(), + expected.end(), + [] __device__(std::size_t i) { return (i + 1) % 2; }); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(is_valid_segment_id, expected); + } + + void test_segment() + { + auto segments = this->copy_segments(); + auto expected = rmm::device_uvector>(2000, this->stream()); + + thrust::tabulate( + rmm::exec_policy(this->stream()), + expected.begin(), + expected.end(), + [] __device__(std::size_t i) { + auto part_idx = i / 2; + auto intra_point_idx = i % 2; + return i % 2 == 0 + ? segment{vec_2d{static_cast(part_idx * 10 + intra_point_idx), + static_cast(part_idx * 10 + intra_point_idx)}, + vec_2d{static_cast(part_idx * 10 + intra_point_idx + 1), + static_cast(part_idx * 10 + intra_point_idx + 1)}} + : segment{vec_2d{-1, -1}, vec_2d{-1, -1}}; + }); + CUSPATIAL_EXPECT_VEC2D_PAIRS_EQUIVALENT(segments, expected); + } + + void test_multilinestring_point_count_it() + { + auto multilinestring_point_count = this->copy_multilinestring_point_count(); + auto expected = rmm::device_uvector(1000, this->stream()); + + thrust::fill(rmm::exec_policy(this->stream()), expected.begin(), expected.end(), 2); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(multilinestring_point_count, expected); + } + + void test_multilinestring_linestring_count_it() + { + auto multilinestring_linestring_count = this->copy_multilinestring_linestring_count(); + auto expected = rmm::device_uvector(1000, this->stream()); + + thrust::fill(rmm::exec_policy(this->stream()), expected.begin(), expected.end(), 1); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(multilinestring_linestring_count, expected); + } + + void test_array_access_operator() + { + auto all_points = this->copy_all_points_of_ith_multilinestring(513); + auto expected = make_device_vector>({{5130, 5130}, {5131, 5131}}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(all_points, expected); + } + + void test_geometry_offset_it() + { + auto geometry_offsets = this->copy_geometry_offsets(); + auto expected = rmm::device_uvector(1001, this->stream()); + + thrust::sequence(rmm::exec_policy(this->stream()), expected.begin(), expected.end()); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_offsets, expected); + } + + void test_part_offsets_it() + { + auto part_offsets = this->copy_part_offsets(); + auto expected = rmm::device_uvector(1001, this->stream()); + + thrust::sequence(rmm::exec_policy(this->stream()), expected.begin(), expected.end(), 0, 2); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_offsets, expected); + } +}; + +TYPED_TEST_CASE(MultilinestringRangeOneThousandTest, FloatingPointTypes); +TYPED_TEST(MultilinestringRangeOneThousandTest, OneThousandTest) { this->run_test(); } diff --git a/cpp/tests/range/multipoint_range_test.cu b/cpp/tests/range/multipoint_range_test.cu new file mode 100644 index 000000000..501ddd6cf --- /dev/null +++ b/cpp/tests/range/multipoint_range_test.cu @@ -0,0 +1,514 @@ +/* + * Copyright (c) 2023, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include +#include + +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include + +#include +#include + +#include + +using namespace cuspatial; +using namespace cuspatial::test; + +template +void __global__ array_access_tester(MultiPointRange multipoints, + std::size_t i, + OutputIt output_points) +{ + using T = typename MultiPointRange::element_t; + thrust::copy(thrust::seq, multipoints[i].begin(), multipoints[i].end(), output_points); +} + +template +void __global__ point_accessor_tester(MultiPointRange multipoints, std::size_t i, OutputIt point) +{ + using T = typename MultiPointRange::element_t; + point[0] = multipoints.point(i); +} + +template +class MultipointRangeTest : public BaseFixture { + public: + struct copy_leading_point_per_multipoint { + template + vec_2d __device__ operator()(MultiPointRef multipoint) + { + return multipoint.size() > 0 ? multipoint[0] : vec_2d{-1, -1}; + } + }; + + template + struct point_idx_to_geometry_idx { + MultiPointRange rng; + + point_idx_to_geometry_idx(MultiPointRange r) : rng(r) {} + + std::size_t __device__ operator()(std::size_t pidx) + { + return rng.geometry_idx_from_point_idx(pidx); + } + }; + + void SetUp() { make_test_multipoints(); } + auto range() { return test_multipoints->range(); } + + virtual void make_test_multipoints() = 0; + + void run_test() + { + test_num_multipoints(); + + test_num_points(); + + test_size(); + + test_multipoint_it(); + + test_begin(); + + test_end(); + + test_point_it(); + + test_offsets_it(); + + test_geometry_idx_from_point_idx(); + + test_subscript_operator(); + + test_point_accessor(); + + test_is_single_point_range(); + } + + virtual void test_num_multipoints() = 0; + + virtual void test_num_points() = 0; + + void test_size() { EXPECT_EQ(this->range().size(), this->range().num_multipoints()); } + + virtual void test_multipoint_it() = 0; + + void test_begin() { EXPECT_EQ(this->range().begin(), this->range().multipoint_begin()); } + + void test_end() { EXPECT_EQ(this->range().end(), this->range().multipoint_end()); } + + virtual void test_point_it() = 0; + + virtual void test_offsets_it() = 0; + + virtual void test_geometry_idx_from_point_idx() = 0; + + virtual void test_subscript_operator() = 0; + + virtual void test_point_accessor() = 0; + + virtual void test_is_single_point_range() = 0; + + protected: + rmm::device_uvector> copy_leading_points() + { + auto rng = this->range(); + rmm::device_uvector> leading_points(rng.num_multipoints(), this->stream()); + thrust::transform(rmm::exec_policy(this->stream()), + rng.multipoint_begin(), + rng.multipoint_end(), + leading_points.begin(), + copy_leading_point_per_multipoint{}); + + return leading_points; + } + + rmm::device_uvector> copy_all_points() + { + auto rng = this->range(); + rmm::device_uvector> points(rng.num_points(), this->stream()); + thrust::copy( + rmm::exec_policy(this->stream()), rng.point_begin(), rng.point_end(), points.begin()); + return points; + }; + + rmm::device_uvector copy_offsets() + { + auto rng = this->range(); + rmm::device_uvector offsets(rng.num_multipoints() + 1, this->stream()); + thrust::copy( + rmm::exec_policy(this->stream()), rng.offsets_begin(), rng.offsets_end(), offsets.begin()); + return offsets; + }; + + rmm::device_uvector copy_geometry_idx() + { + auto rng = this->range(); + rmm::device_uvector idx(rng.num_points(), this->stream()); + + thrust::tabulate( + rmm::exec_policy(this->stream()), idx.begin(), idx.end(), point_idx_to_geometry_idx{rng}); + return idx; + } + + rmm::device_scalar> copy_ith_point(std::size_t i) + { + auto rng = this->range(); + + rmm::device_scalar> point(this->stream()); + point_accessor_tester<<<1, 1, 0, this->stream()>>>(rng, i, point.data()); + CUSPATIAL_CHECK_CUDA(this->stream()); + + return point; + } + + rmm::device_uvector> copy_ith_multipoint(std::size_t i) + { + auto rng = this->range(); + rmm::device_scalar num_points(this->stream()); + auto count_iterator = make_count_iterator_from_offset_iterator(this->range().offsets_begin()); + thrust::copy_n(rmm::exec_policy(this->stream()), count_iterator, 1, num_points.data()); + + rmm::device_uvector> multipoint(num_points.value(this->stream()), this->stream()); + array_access_tester<<<1, 1, 0, this->stream()>>>(rng, i, multipoint.begin()); + CUSPATIAL_CHECK_CUDA(this->stream()); + + return multipoint; + } + + std::unique_ptr, rmm::device_vector>>> + test_multipoints; +}; + +template +class EmptyMultiPointRangeTest : public MultipointRangeTest { + public: + void make_test_multipoints() + { + auto array = make_multipoint_array({}); + this->test_multipoints = std::make_unique(std::move(array)); + } + + void test_num_multipoints() { EXPECT_EQ(this->range().num_multipoints(), 0); } + + void test_num_points() { EXPECT_EQ(this->range().num_points(), 0); } + + void test_multipoint_it() + { + auto leading_points = this->copy_leading_points(); + auto expected = rmm::device_uvector>(0, this->stream()); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(leading_points, expected); + } + + void test_point_it() + { + auto points = this->copy_all_points(); + auto expected = rmm::device_uvector>(0, this->stream()); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(points, expected); + } + + void test_offsets_it() + { + auto offsets = this->copy_offsets(); + auto expected = make_device_vector({0}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(offsets, expected); + } + + void test_geometry_idx_from_point_idx() + { + auto geometry_indices = this->copy_geometry_idx(); + auto expected = rmm::device_uvector(0, this->stream()); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_indices, expected); + } + + void test_subscript_operator() + { + // Range is empty, nothing to test. + SUCCEED(); + } + + void test_point_accessor() + { + // Range is empty, nothing to test. + SUCCEED(); + } + + void test_is_single_point_range() + { + // Range is empty, undefined behavior. + SUCCEED(); + } +}; + +TYPED_TEST_CASE(EmptyMultiPointRangeTest, FloatingPointTypes); +TYPED_TEST(EmptyMultiPointRangeTest, Test) { this->run_test(); } + +template +class LengthOneMultiPointRangeTest : public MultipointRangeTest { + public: + void make_test_multipoints() + { + auto array = make_multipoint_array({{{1.0, 1.0}, {10.0, 10.0}}}); + this->test_multipoints = std::make_unique(std::move(array)); + } + + void test_num_multipoints() { EXPECT_EQ(this->range().num_multipoints(), 1); } + + void test_num_points() { EXPECT_EQ(this->range().num_points(), 2); } + + void test_multipoint_it() + { + auto leading_points = this->copy_leading_points(); + auto expected = make_device_vector>({{1.0, 1.0}}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(leading_points, expected); + } + + void test_point_it() + { + auto points = this->copy_all_points(); + auto expected = make_device_vector>({{1.0, 1.0}, {10.0, 10.0}}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(points, expected); + } + + void test_offsets_it() + { + auto offsets = this->copy_offsets(); + auto expected = make_device_vector({0, 2}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(offsets, expected); + } + + void test_geometry_idx_from_point_idx() + { + auto geometry_indices = this->copy_geometry_idx(); + auto expected = make_device_vector({0, 0}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_indices, expected); + } + + void test_subscript_operator() + { + auto multipoint = this->copy_ith_multipoint(0); + auto expected = make_device_vector>({{1.0, 1.0}, {10.0, 10.0}}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(multipoint, expected); + } + + void test_point_accessor() + { + auto point = this->copy_ith_point(1); + auto expected = vec_2d{10.0, 10.0}; + EXPECT_EQ(point.value(this->stream()), expected); + } + + void test_is_single_point_range() { EXPECT_FALSE(this->range().is_single_point_range()); } +}; + +TYPED_TEST_CASE(LengthOneMultiPointRangeTest, FloatingPointTypes); +TYPED_TEST(LengthOneMultiPointRangeTest, Test) { this->run_test(); } + +template +class LengthFiveMultiPointRangeTest : public MultipointRangeTest { + public: + void make_test_multipoints() + { + auto array = make_multipoint_array({{{0.0, 0.0}, {1.0, 1.0}}, + {{10.0, 10.0}}, + {{20.0, 21.0}, {22.0, 23.0}}, + {{30.0, 31.0}, {32.0, 33.0}, {34.0, 35.0}}, + {}}); + this->test_multipoints = std::make_unique(std::move(array)); + } + + void test_num_multipoints() { EXPECT_EQ(this->range().num_multipoints(), 5); } + + void test_num_points() { EXPECT_EQ(this->range().num_points(), 8); } + + void test_multipoint_it() + { + auto leading_points = this->copy_leading_points(); + auto expected = make_device_vector>( + {{0.0, 0.0}, {10.0, 10.0}, {20.0, 21.0}, {30.0, 31.0}, {-1, -1}}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(leading_points, expected); + } + + void test_point_it() + { + auto points = this->copy_all_points(); + auto expected = make_device_vector>({{0.0, 0.0}, + {1.0, 1.0}, + {10.0, 10.0}, + {20.0, 21.0}, + {22.0, 23.0}, + {30.0, 31.0}, + {32.0, 33.0}, + {34.0, 35.0}}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(points, expected); + } + + void test_offsets_it() + { + auto offsets = this->copy_offsets(); + auto expected = make_device_vector({0, 2, 3, 5, 8, 8}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(offsets, expected); + } + + void test_geometry_idx_from_point_idx() + { + auto geometry_indices = this->copy_geometry_idx(); + auto expected = make_device_vector({0, 0, 1, 2, 2, 3, 3, 3}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_indices, expected); + } + + void test_subscript_operator() + { + auto second_multipoint = this->copy_ith_multipoint(2); + auto expected = make_device_vector>({{20.0, 21.0}, {22.0, 23.0}}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(second_multipoint, expected); + } + + void test_point_accessor() + { + auto third_point = this->copy_ith_point(3); + auto expected = vec_2d{20.0, 21.0}; + EXPECT_EQ(third_point.value(this->stream()), expected); + } + + void test_is_single_point_range() { EXPECT_FALSE(this->range().is_single_point_range()); } +}; + +TYPED_TEST_CASE(LengthFiveMultiPointRangeTest, FloatingPointTypes); +TYPED_TEST(LengthFiveMultiPointRangeTest, Test) { this->run_test(); } + +template +class LengthOneThousandRangeTest : public MultipointRangeTest { + public: + std::size_t static constexpr num_multipoints = 1000; + std::size_t static constexpr num_point_per_multipoint = 3; + std::size_t static constexpr num_points = num_multipoints * num_point_per_multipoint; + void make_test_multipoints() + { + rmm::device_vector geometry_offsets(num_multipoints + 1); + rmm::device_vector> coordinates(num_points); + + thrust::sequence(rmm::exec_policy(this->stream()), + geometry_offsets.begin(), + geometry_offsets.end(), + 0ul, + num_point_per_multipoint); + + thrust::tabulate(rmm::exec_policy(this->stream()), + coordinates.begin(), + coordinates.end(), + [] __device__(auto i) { + return vec_2d{static_cast(i), 10.0}; + }); + + auto array = + make_multipoint_array(std::move(geometry_offsets), std::move(coordinates)); + + this->test_multipoints = std::make_unique(std::move(array)); + } + + void test_num_multipoints() { EXPECT_EQ(this->range().num_multipoints(), num_multipoints); } + + void test_num_points() { EXPECT_EQ(this->range().num_points(), num_points); } + + void test_multipoint_it() + { + auto leading_points = this->copy_leading_points(); + auto expect = rmm::device_uvector>(num_multipoints, this->stream()); + + thrust::tabulate( + rmm::exec_policy(this->stream()), expect.begin(), expect.end(), [] __device__(auto i) { + return vec_2d{static_cast(i) * num_point_per_multipoint, 10.0}; + }); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(leading_points, expect); + } + + void test_point_it() + { + auto all_points = this->copy_all_points(); + auto expect = rmm::device_uvector>(num_points, this->stream()); + + thrust::tabulate( + rmm::exec_policy(this->stream()), expect.begin(), expect.end(), [] __device__(auto i) { + return vec_2d{static_cast(i), 10.0}; + }); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(all_points, expect); + } + + void test_offsets_it() + { + auto offsets = this->copy_offsets(); + auto expect = rmm::device_uvector(num_multipoints + 1, this->stream()); + thrust::sequence(rmm::exec_policy(this->stream()), + expect.begin(), + expect.end(), + 0ul, + num_point_per_multipoint); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(offsets, expect); + } + + void test_geometry_idx_from_point_idx() + { + auto indices = this->copy_geometry_idx(); + auto expect = rmm::device_uvector(3000, this->stream()); + thrust::tabulate( + rmm::exec_policy(this->stream()), expect.begin(), expect.end(), [] __device__(auto i) { + return i / num_point_per_multipoint; + }); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(indices, expect); + } + + void test_subscript_operator() + { + auto multipoint_five_hundred_thirty_third = this->copy_ith_multipoint(533); + auto expect = make_device_vector>({{533 * num_point_per_multipoint, 10.0}, + {533 * num_point_per_multipoint + 1, 10.0}, + {533 * num_point_per_multipoint + 2, 10.0}}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(multipoint_five_hundred_thirty_third, expect); + } + + void test_point_accessor() + { + auto point_seventeen_hundred_seventy_six = this->copy_ith_point(1776); + auto expect = vec_2d{1776, 10.0}; + + EXPECT_EQ(point_seventeen_hundred_seventy_six.value(this->stream()), expect); + } + + void test_is_single_point_range() { EXPECT_FALSE(this->range().is_single_point_range()); } +}; + +TYPED_TEST_CASE(LengthOneThousandRangeTest, FloatingPointTypes); +TYPED_TEST(LengthOneThousandRangeTest, Test) { this->run_test(); } diff --git a/cpp/tests/range/multipolygon_range_test.cu b/cpp/tests/range/multipolygon_range_test.cu index 2fe133945..1ba2d9935 100644 --- a/cpp/tests/range/multipolygon_range_test.cu +++ b/cpp/tests/range/multipolygon_range_test.cu @@ -22,10 +22,14 @@ #include #include +#include #include #include #include +#include +#include + #include using namespace cuspatial; @@ -121,10 +125,14 @@ struct MultipolygonRangeTest : public BaseFixture { multipolygon_coordinates); auto rng = multipolygon_array.range().as_multilinestring_range(); - auto got = - make_multilinestring_array(range(rng.geometry_offsets_begin(), rng.geometry_offsets_end()), - range(rng.part_offsets_begin(), rng.part_offsets_end()), - range(rng.point_begin(), rng.point_end())); + auto geometry_offsets = + rmm::device_vector(rng.geometry_offset_begin(), rng.geometry_offset_end()); + auto part_offsets = + rmm::device_vector(rng.part_offset_begin(), rng.part_offset_end()); + auto points = rmm::device_vector>(rng.point_begin(), rng.point_end()); + + auto got = make_multilinestring_array( + std::move(geometry_offsets), std::move(part_offsets), std::move(points)); auto expected = make_multilinestring_array( multilinestring_geometry_offset, multilinestring_part_offset, multilinestring_coordinates); @@ -482,7 +490,7 @@ TYPED_TEST(MultipolygonRangeTest, MultipolygonSegmentCount_ContainsEmptyPart) {6, 3}); } -TYPED_TEST(MultipolygonRangeTest, MultipolygonAsMultilinestring1) +TYPED_TEST(MultipolygonRangeTest, MultipolygonAsMultipolygon1) { CUSPATIAL_RUN_TEST(this->test_multipolygon_as_multilinestring, {0, 1, 2}, @@ -494,7 +502,7 @@ TYPED_TEST(MultipolygonRangeTest, MultipolygonAsMultilinestring1) {{0, 0}, {1, 0}, {1, 1}, {0, 0}, {10, 10}, {11, 10}, {11, 11}, {10, 10}}); } -TYPED_TEST(MultipolygonRangeTest, MultipolygonAsMultilinestring2) +TYPED_TEST(MultipolygonRangeTest, MultipolygonAsMultipolygon2) { CUSPATIAL_RUN_TEST(this->test_multipolygon_as_multilinestring, {0, 1, 2}, @@ -528,7 +536,7 @@ TYPED_TEST(MultipolygonRangeTest, MultipolygonAsMultilinestring2) {20, 20}}); } -TYPED_TEST(MultipolygonRangeTest, MultipolygonAsMultilinestring3) +TYPED_TEST(MultipolygonRangeTest, MultipolygonAsMultipolygon3) { CUSPATIAL_RUN_TEST(this->test_multipolygon_as_multilinestring, {0, 1, 2}, @@ -638,3 +646,677 @@ TYPED_TEST(MultipolygonRangeTest, MultipolygonAsMultiPoint3) {21, 21}, {20, 20}}); } + +template +__global__ void array_access_tester(MultiPolygonRange rng, std::size_t i, PointOutputIt output) +{ + thrust::copy(thrust::seq, rng[i].point_begin(), rng[i].point_end(), output); +} + +template +class MultipolygonRangeTestBase : public BaseFixture { + public: + struct copy_leading_point_functor { + template + __device__ vec_2d operator()(MultiPolygonRef mpolygon) + { + return mpolygon.size() > 0 ? mpolygon.point_begin()[0] : vec_2d{-1, -1}; + } + }; + + template + struct ring_idx_from_point_idx_functor { + MultiPolygonRange mpolygons; + __device__ std::size_t operator()(std::size_t point_idx) + { + return mpolygons.ring_idx_from_point_idx(point_idx); + } + }; + + template + struct part_idx_from_ring_idx_functor { + MultiPolygonRange mpolygons; + __device__ std::size_t operator()(std::size_t ring_idx) + { + return mpolygons.part_idx_from_ring_idx(ring_idx); + } + }; + + template + struct geometry_idx_from_part_idx_functor { + MultiPolygonRange mpolygons; + __device__ std::size_t operator()(std::size_t part_idx) + { + return mpolygons.geometry_idx_from_part_idx(part_idx); + } + }; + + void SetUp() { make_test_multipolygon(); } + + virtual void make_test_multipolygon() = 0; + + auto range() { return test_multipolygon->range(); } + + void run_test() + { + test_size(); + + test_num_multipolygons(); + + test_num_polygons(); + + test_num_rings(); + + test_num_points(); + + test_multipolygon_it(); + + test_begin(); + + test_end(); + + test_point_it(); + + test_geometry_offsets_it(); + + test_part_offset_it(); + + test_ring_offset_it(); + + test_ring_idx_from_point_idx(); + + test_part_idx_from_ring_idx(); + + test_geometry_idx_from_part_idx(); + + test_array_access_operator(); + + test_multipolygon_point_count_it(); + + test_multipolygon_ring_count_it(); + } + + void test_size() { EXPECT_EQ(range().size(), range().num_multipolygons()); }; + + virtual void test_num_multipolygons() = 0; + + virtual void test_num_polygons() = 0; + + virtual void test_num_rings() = 0; + + virtual void test_num_points() = 0; + + virtual void test_multipolygon_it() = 0; + + void test_begin() { EXPECT_EQ(range().begin(), range().multipolygon_begin()); } + + void test_end() { EXPECT_EQ(range().end(), range().multipolygon_end()); } + + virtual void test_point_it() = 0; + + virtual void test_geometry_offsets_it() = 0; + + virtual void test_part_offset_it() = 0; + + virtual void test_ring_offset_it() = 0; + + virtual void test_ring_idx_from_point_idx() = 0; + + virtual void test_part_idx_from_ring_idx() = 0; + + virtual void test_geometry_idx_from_part_idx() = 0; + + virtual void test_array_access_operator() = 0; + + virtual void test_multipolygon_point_count_it() = 0; + + virtual void test_multipolygon_ring_count_it() = 0; + + // helper method to access multipolygon range + rmm::device_uvector> copy_leading_point_multipolygon() + { + auto rng = range(); + auto d_points = rmm::device_uvector>(rng.num_multipolygons(), stream()); + thrust::transform(rmm::exec_policy(stream()), + rng.multipolygon_begin(), + rng.multipolygon_end(), + d_points.begin(), + copy_leading_point_functor{}); + return d_points; + } + + rmm::device_uvector> copy_all_points() + { + auto rng = range(); + auto d_points = rmm::device_uvector>(rng.num_points(), stream()); + thrust::copy(rmm::exec_policy(stream()), rng.point_begin(), rng.point_end(), d_points.begin()); + return d_points; + } + + rmm::device_uvector copy_geometry_offsets() + { + auto rng = range(); + auto d_offsets = rmm::device_uvector(rng.num_multipolygons() + 1, stream()); + thrust::copy(rmm::exec_policy(stream()), + rng.geometry_offset_begin(), + rng.geometry_offset_end(), + d_offsets.begin()); + return d_offsets; + } + + rmm::device_uvector copy_part_offsets() + { + auto rng = range(); + auto d_offsets = rmm::device_uvector(rng.num_polygons() + 1, stream()); + thrust::copy(rmm::exec_policy(stream()), + rng.part_offset_begin(), + rng.part_offset_end(), + d_offsets.begin()); + return d_offsets; + } + + rmm::device_uvector copy_ring_offsets() + { + auto rng = range(); + auto d_offsets = rmm::device_uvector(rng.num_rings() + 1, stream()); + thrust::copy(rmm::exec_policy(stream()), + rng.ring_offset_begin(), + rng.ring_offset_end(), + d_offsets.begin()); + return d_offsets; + } + + rmm::device_uvector copy_ring_idx_from_point_idx() + { + auto rng = range(); + auto d_ring_idx = rmm::device_uvector(rng.num_points(), stream()); + thrust::tabulate(rmm::exec_policy(stream()), + d_ring_idx.begin(), + d_ring_idx.end(), + ring_idx_from_point_idx_functor{rng}); + return d_ring_idx; + } + + rmm::device_uvector copy_part_idx_from_ring_idx() + { + auto rng = range(); + auto d_part_idx = rmm::device_uvector(rng.num_rings(), stream()); + thrust::tabulate(rmm::exec_policy(stream()), + d_part_idx.begin(), + d_part_idx.end(), + part_idx_from_ring_idx_functor{rng}); + return d_part_idx; + } + + rmm::device_uvector copy_geometry_idx_from_part_idx() + { + auto rng = range(); + auto d_geometry_idx = rmm::device_uvector(rng.num_polygons(), stream()); + thrust::tabulate(rmm::exec_policy(stream()), + d_geometry_idx.begin(), + d_geometry_idx.end(), + geometry_idx_from_part_idx_functor{rng}); + return d_geometry_idx; + } + + rmm::device_uvector> copy_all_points_of_ith_multipolygon(std::size_t i) + { + auto rng = this->range(); + rmm::device_scalar num_points(stream()); + + thrust::copy_n( + rmm::exec_policy(stream()), rng.multipolygon_point_count_begin() + i, 1, num_points.data()); + + auto d_all_points = rmm::device_uvector>(num_points.value(stream()), stream()); + + array_access_tester<<<1, 1, 0, stream()>>>(rng, i, d_all_points.data()); + return d_all_points; + } + + rmm::device_uvector copy_multipolygon_point_count() + { + auto rng = this->range(); + auto d_point_count = rmm::device_uvector(rng.num_multipolygons(), stream()); + thrust::copy(rmm::exec_policy(stream()), + rng.multipolygon_point_count_begin(), + rng.multipolygon_point_count_end(), + d_point_count.begin()); + return d_point_count; + } + + rmm::device_uvector copy_multipolygon_ring_count() + { + auto rng = this->range(); + auto d_ring_count = rmm::device_uvector(rng.num_multipolygons(), stream()); + thrust::copy(rmm::exec_policy(stream()), + rng.multipolygon_ring_count_begin(), + rng.multipolygon_ring_count_end(), + d_ring_count.begin()); + return d_ring_count; + } + + protected: + std::unique_ptr, + rmm::device_vector, + rmm::device_vector, + rmm::device_vector>>> + test_multipolygon; +}; + +template +class MultipolygonRangeEmptyTest : public MultipolygonRangeTestBase { + void make_test_multipolygon() + { + auto geometry_offsets = make_device_vector({0}); + auto part_offsets = make_device_vector({0}); + auto ring_offsets = make_device_vector({0}); + auto coordinates = make_device_vector>({}); + + this->test_multipolygon = std::make_unique, + rmm::device_vector, + rmm::device_vector, + rmm::device_vector>>>( + std::move(geometry_offsets), + std::move(part_offsets), + std::move(ring_offsets), + std::move(coordinates)); + } + + void test_num_multipolygons() { EXPECT_EQ(this->range().num_multipolygons(), 0); } + + void test_num_polygons() { EXPECT_EQ(this->range().num_polygons(), 0); } + + void test_num_rings() { EXPECT_EQ(this->range().num_rings(), 0); } + + void test_num_points() { EXPECT_EQ(this->range().num_points(), 0); } + + void test_multipolygon_it() + { + rmm::device_uvector> d_points = this->copy_leading_point_multipolygon(); + rmm::device_uvector> expected(0, this->stream()); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_points, expected); + } + + void test_point_it() + { + rmm::device_uvector> d_points = this->copy_all_points(); + rmm::device_uvector> expected(0, this->stream()); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_points, expected); + } + + void test_geometry_offsets_it() + { + rmm::device_uvector d_offsets = this->copy_geometry_offsets(); + auto expected = make_device_vector({0}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_offsets, expected); + } + + void test_part_offset_it() + { + rmm::device_uvector d_offsets = this->copy_part_offsets(); + auto expected = make_device_vector({0}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_offsets, expected); + } + + void test_ring_offset_it() + { + rmm::device_uvector d_offsets = this->copy_ring_offsets(); + auto expected = make_device_vector({0}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_offsets, expected); + } + + void test_ring_idx_from_point_idx() + { + rmm::device_uvector d_ring_idx = this->copy_ring_idx_from_point_idx(); + auto expected = make_device_vector({}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_ring_idx, expected); + } + + void test_part_idx_from_ring_idx() + { + rmm::device_uvector d_part_idx = this->copy_part_idx_from_ring_idx(); + auto expected = make_device_vector({}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_part_idx, expected); + } + + void test_geometry_idx_from_part_idx() + { + rmm::device_uvector d_geometry_idx = this->copy_geometry_idx_from_part_idx(); + auto expected = make_device_vector({}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_geometry_idx, expected); + } + + void test_array_access_operator() + { + // Nothing to access + SUCCEED(); + } + + void test_multipolygon_point_count_it() + { + rmm::device_uvector d_point_count = this->copy_multipolygon_point_count(); + rmm::device_uvector expected(0, this->stream()); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_point_count, expected); + } + + void test_multipolygon_ring_count_it() + { + rmm::device_uvector d_ring_count = this->copy_multipolygon_ring_count(); + rmm::device_uvector expected(0, this->stream()); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_ring_count, expected); + } +}; + +TYPED_TEST_CASE(MultipolygonRangeEmptyTest, FloatingPointTypes); + +TYPED_TEST(MultipolygonRangeEmptyTest, EmptyMultipolygonRange) { this->run_test(); } + +template +class MultipolygonRangeOneTest : public MultipolygonRangeTestBase { + void make_test_multipolygon() + { + auto geometry_offsets = make_device_vector({0, 2}); + auto part_offsets = make_device_vector({0, 1, 2}); + auto ring_offsets = make_device_vector({0, 4, 8}); + auto coordinates = make_device_vector>( + {{0, 0}, {1, 0}, {1, 1}, {0, 0}, {10, 10}, {11, 10}, {11, 11}, {10, 10}}); + + this->test_multipolygon = std::make_unique, + rmm::device_vector, + rmm::device_vector, + rmm::device_vector>>>( + std::move(geometry_offsets), + std::move(part_offsets), + std::move(ring_offsets), + std::move(coordinates)); + } + + void test_num_multipolygons() { EXPECT_EQ(this->range().num_multipolygons(), 1); } + + void test_num_polygons() { EXPECT_EQ(this->range().num_polygons(), 2); } + + void test_num_rings() { EXPECT_EQ(this->range().num_rings(), 2); } + + void test_num_points() { EXPECT_EQ(this->range().num_points(), 8); } + + void test_multipolygon_it() + { + rmm::device_uvector> d_points = this->copy_leading_point_multipolygon(); + auto expected = make_device_vector>({{0, 0}}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_points, expected); + } + + void test_point_it() + { + rmm::device_uvector> d_points = this->copy_all_points(); + auto expected = make_device_vector>( + {{0, 0}, {1, 0}, {1, 1}, {0, 0}, {10, 10}, {11, 10}, {11, 11}, {10, 10}}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_points, expected); + } + + void test_geometry_offsets_it() + { + rmm::device_uvector d_offsets = this->copy_geometry_offsets(); + auto expected = make_device_vector({0, 1}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_offsets, expected); + } + + void test_part_offset_it() + { + rmm::device_uvector d_offsets = this->copy_part_offsets(); + auto expected = make_device_vector({0, 1, 2}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_offsets, expected); + } + + void test_ring_offset_it() + { + rmm::device_uvector d_offsets = this->copy_ring_offsets(); + auto expected = make_device_vector({0, 4, 8}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_offsets, expected); + } + + void test_ring_idx_from_point_idx() + { + rmm::device_uvector d_ring_idx = this->copy_ring_idx_from_point_idx(); + auto expected = make_device_vector({0, 0, 0, 0, 1, 1, 1, 1}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_ring_idx, expected); + } + + void test_part_idx_from_ring_idx() + { + rmm::device_uvector d_part_idx = this->copy_part_idx_from_ring_idx(); + auto expected = make_device_vector({0, 1}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_part_idx, expected); + } + + void test_geometry_idx_from_part_idx() + { + rmm::device_uvector d_geometry_idx = this->copy_geometry_idx_from_part_idx(); + auto expected = make_device_vector({0, 0}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_geometry_idx, expected); + } + + void test_array_access_operator() + { + auto all_points = this->copy_all_points_of_ith_multipolygon(0); + auto expected = make_device_vector>( + {{0, 0}, {1, 0}, {1, 1}, {0, 0}, {10, 10}, {11, 10}, {11, 11}, {10, 10}}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(all_points, expected); + } + + void test_multipolygon_point_count_it() + { + rmm::device_uvector d_point_count = this->copy_multipolygon_point_count(); + auto expected = make_device_vector({8}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_point_count, expected); + } + + void test_multipolygon_ring_count_it() + { + rmm::device_uvector d_ring_count = this->copy_multipolygon_ring_count(); + auto expected = make_device_vector({2}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_ring_count, expected); + } +}; + +TYPED_TEST_CASE(MultipolygonRangeOneTest, FloatingPointTypes); + +TYPED_TEST(MultipolygonRangeOneTest, OneMultipolygonRange) { this->run_test(); } + +template +class MultipolygonRangeOneThousandTest : public MultipolygonRangeTestBase { + public: + struct make_points_functor { + __device__ auto operator()(std::size_t i) + { + auto geometry_idx = i / 4; + auto intra_point_idx = i % 4; + return vec_2d{geometry_idx * T{10.} + intra_point_idx, + geometry_idx * T{10.} + intra_point_idx}; + } + }; + + void make_test_multipolygon() + { + auto geometry_offsets = rmm::device_vector(1001); + auto part_offsets = rmm::device_vector(1001); + auto ring_offsets = rmm::device_vector(1001); + auto coordinates = rmm::device_vector>(4000); + + thrust::sequence( + rmm::exec_policy(this->stream()), geometry_offsets.begin(), geometry_offsets.end()); + + thrust::sequence(rmm::exec_policy(this->stream()), part_offsets.begin(), part_offsets.end()); + + thrust::sequence( + rmm::exec_policy(this->stream()), ring_offsets.begin(), ring_offsets.end(), 0, 4); + + thrust::tabulate(rmm::exec_policy(this->stream()), + coordinates.begin(), + coordinates.end(), + make_points_functor{}); + + this->test_multipolygon = std::make_unique, + rmm::device_vector, + rmm::device_vector, + rmm::device_vector>>>( + std::move(geometry_offsets), + std::move(part_offsets), + std::move(ring_offsets), + std::move(coordinates)); + } + + void test_num_multipolygons() { EXPECT_EQ(this->range().num_multipolygons(), 1000); } + + void test_num_polygons() { EXPECT_EQ(this->range().num_polygons(), 1000); } + + void test_num_rings() { EXPECT_EQ(this->range().num_rings(), 1000); } + + void test_num_points() { EXPECT_EQ(this->range().num_points(), 4000); } + + void test_multipolygon_it() + { + rmm::device_uvector> d_points = this->copy_leading_point_multipolygon(); + rmm::device_uvector> expected(1000, this->stream()); + thrust::tabulate(rmm::exec_policy(this->stream()), + expected.begin(), + expected.end(), + [] __device__(std::size_t i) { + return vec_2d{i * T{10.}, i * T{10.}}; + }); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_points, expected); + } + + void test_point_it() + { + rmm::device_uvector> d_points = this->copy_all_points(); + rmm::device_uvector> expected(4000, this->stream()); + + thrust::tabulate( + rmm::exec_policy(this->stream()), expected.begin(), expected.end(), make_points_functor{}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_points, expected); + } + + void test_geometry_offsets_it() + { + rmm::device_uvector d_offsets = this->copy_geometry_offsets(); + auto expected = rmm::device_uvector(1001, this->stream()); + + thrust::sequence(rmm::exec_policy(this->stream()), expected.begin(), expected.end()); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_offsets, expected); + } + + void test_part_offset_it() + { + rmm::device_uvector d_offsets = this->copy_part_offsets(); + auto expected = rmm::device_uvector(1001, this->stream()); + + thrust::sequence(rmm::exec_policy(this->stream()), expected.begin(), expected.end()); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_offsets, expected); + } + + void test_ring_offset_it() + { + rmm::device_uvector d_offsets = this->copy_ring_offsets(); + auto expected = rmm::device_uvector(1001, this->stream()); + + thrust::sequence(rmm::exec_policy(this->stream()), expected.begin(), expected.end(), 0, 4); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_offsets, expected); + } + + void test_ring_idx_from_point_idx() + { + rmm::device_uvector d_ring_idx = this->copy_ring_idx_from_point_idx(); + auto expected = rmm::device_uvector(4000, this->stream()); + + thrust::tabulate(rmm::exec_policy(this->stream()), + expected.begin(), + expected.end(), + [] __device__(std::size_t i) { return i / 4; }); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_ring_idx, expected); + } + + void test_part_idx_from_ring_idx() + { + rmm::device_uvector d_part_idx = this->copy_part_idx_from_ring_idx(); + auto expected = rmm::device_uvector(1000, this->stream()); + + thrust::sequence(rmm::exec_policy(this->stream()), expected.begin(), expected.end()); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_part_idx, expected); + } + + void test_geometry_idx_from_part_idx() + { + rmm::device_uvector d_geometry_idx = this->copy_geometry_idx_from_part_idx(); + auto expected = rmm::device_uvector(1000, this->stream()); + + thrust::sequence(rmm::exec_policy(this->stream()), expected.begin(), expected.end()); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_geometry_idx, expected); + } + + void test_array_access_operator() + { + auto all_points = this->copy_all_points_of_ith_multipolygon(777); + auto expected = make_device_vector>({ + {7770, 7770}, + {7771, 7771}, + {7772, 7772}, + {7773, 7773}, + }); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(all_points, expected); + } + + void test_multipolygon_point_count_it() + { + rmm::device_uvector d_point_count = this->copy_multipolygon_point_count(); + auto expected = rmm::device_uvector(1000, this->stream()); + + thrust::fill(rmm::exec_policy(this->stream()), expected.begin(), expected.end(), 4); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_point_count, expected); + } + + void test_multipolygon_ring_count_it() + { + rmm::device_uvector d_ring_count = this->copy_multipolygon_ring_count(); + auto expected = rmm::device_uvector(1000, this->stream()); + + thrust::fill(rmm::exec_policy(this->stream()), expected.begin(), expected.end(), 1); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_ring_count, expected); + } +}; + +TYPED_TEST_CASE(MultipolygonRangeOneThousandTest, FloatingPointTypes); + +TYPED_TEST(MultipolygonRangeOneThousandTest, OneThousandMultipolygonRange) { this->run_test(); }