From b4a2421d8a660d8dc47567a99675032aa5bc73f1 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Mon, 12 Jun 2023 13:24:07 -0700 Subject: [PATCH 01/19] add multipoint_tests --- .devcontainer/devcontainer.json | 6 +- cpp/CMakeLists.txt | 6 +- cpp/tests/CMakeLists.txt | 1 + cpp/tests/range/multipoint_range_test.cu | 267 +++++++++++++++++++++++ 4 files changed, 275 insertions(+), 5 deletions(-) create mode 100644 cpp/tests/range/multipoint_range_test.cu diff --git a/.devcontainer/devcontainer.json b/.devcontainer/devcontainer.json index 902276b71..5bcea0a1f 100644 --- a/.devcontainer/devcontainer.json +++ b/.devcontainer/devcontainer.json @@ -40,14 +40,16 @@ "source=${localWorkspaceFolder}/../.cache,target=/home/coder/.cache,type=bind,consistency=consistent", "source=${localWorkspaceFolder}/../.config,target=/home/coder/.config,type=bind,consistency=consistent", "source=${localWorkspaceFolder}/../.conda/pkgs,target=/home/coder/.conda/pkgs,type=bind,consistency=consistent", - "source=${localWorkspaceFolder}/../.conda/${localWorkspaceFolderBasename}/single,target=/home/coder/.conda/envs,type=bind,consistency=consistent" + "source=${localWorkspaceFolder}/../.conda/${localWorkspaceFolderBasename}/single,target=/home/coder/.conda/envs,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../scratch,target=/home/coder/scratch,type=bind,consistency=consistent" ], "customizations": { "vscode": { "extensions": [ "mutantdino.resourcemonitor", - "tamasfe.even-better-toml" + "tamasfe.even-better-toml", + "ms-toolsai.jupyter" ], "settings": { "files.trimFinalNewlines": true, diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index c3477d10d..ecaae8302 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -47,7 +47,7 @@ option(BUILD_BENCHMARKS "Configure CMake to build (google) benchmarks" OFF) option(PER_THREAD_DEFAULT_STREAM "Build with per-thread default stream" OFF) option(DISABLE_DEPRECATION_WARNING "Disable warnings generated from deprecated declarations." OFF) # Option to enable line info in CUDA device compilation to allow introspection when profiling / memchecking -option(CUDA_ENABLE_LINEINFO "Enable the -lineinfo option for nvcc (useful for cuda-memcheck / profiler" OFF) +option(CUDA_ENABLE_LINEINFO "Enable the -lineinfo option for nvcc (useful for cuda-memcheck / profiler" ON) # cudart can be statically linked or dynamically linked. The python ecosystem wants dynamic linking option(CUDA_STATIC_RUNTIME "Statically link the CUDA toolkit runtime and libraries" OFF) @@ -63,12 +63,12 @@ message(STATUS "CUSPATIAL: Enable the -lineinfo option for nvcc (useful for cuda message(STATUS "CUSPATIAL: Statically link the CUDA toolkit runtime and libraries: ${CUDA_STATIC_RUNTIME}") # Set a default build type if none was specified -rapids_cmake_build_type("Release") +rapids_cmake_build_type("Debug") set(CUSPATIAL_BUILD_TESTS ${BUILD_TESTS}) set(CUSPATIAL_BUILD_BENCHMARKS ${BUILD_BENCHMARKS}) set(CUSPATIAL_CXX_FLAGS "") -set(CUSPATIAL_CUDA_FLAGS "") +set(CUSPATIAL_CUDA_FLAGS "-ftemplate-backtrace-limit=0") set(CUSPATIAL_CXX_DEFINITIONS "") set(CUSPATIAL_CUDA_DEFINITIONS "") diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 8f344e534..b732d91fa 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/multipoint_range_test.cu b/cpp/tests/range/multipoint_range_test.cu new file mode 100644 index 000000000..fa80d625d --- /dev/null +++ b/cpp/tests/range/multipoint_range_test.cu @@ -0,0 +1,267 @@ +/* + * 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 + +using namespace cuspatial; +using namespace cuspatial::test; + +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; + + 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(); + + test_is_single_point_range(); + } + + virtual void test_num_multipoints() = 0; + + virtual void test_num_points() = 0; + + virtual void test_size() = 0; + + virtual void test_multipoint_it() = 0; + + virtual void test_begin() = 0; + + virtual void test_end() = 0; + + 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() = 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.multipoin_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.offset_begin(), rng.offset_end(), offsets.begin()); + return offsets; + }; + + rmm::device_uvector copy_geoemtry_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_uvector> copy_ith_multipoint(std::size_t i) + { + rmm::device_scalar num_points(this->stream()); + + } + + multipoint_array test_multipoints; +}; + +template +class EmptyMultiPointRangeTest : public MultipointRangeTest { + public: + void make_test_multipoints() { this->test_multipoints = make_multipoint_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_size() { EXPECT_EQ(this->range().size(), 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_begin() { EXPECT_EQ(this->range().begin(), this->range().multipoint_begin()); } + + void test_end() { EXPECT_EQ(this->range().end(), this->range().multipoint_end()); } + + 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_all_points(); + auto expected = make_device_vector({0}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(offsets, expected); + } + + void test_geometry_idx_from_point_idx() + { + auto geometry_indices = rmm::device_uvector(0, this->stream()); + auto expected = rmm::device_uvector(0, this->stream()) ; + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_indices, expected); + } + + void test_subscript_operator() + { + + } + +}; + +template +class LengthOneMultiPointRangeTest : public MultipointRangeTest { + public: + void make_test_multipoints() + { + this->test_multipoints = make_multipoint_array({{{1.0, 1.0}}}); + } +}; + +template +class LengthFiveMultiPointRangeTest : public MultipointRangeTest { + public: + void make_test_multipoints() + { + this->test_multipoints = 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}}, + {{}}}); + } +}; + +template +class LengthOneThousandRangeTest : public MultipointRangeTest { + public: + void make_test_multipoints() + { + std::size_t constexpr num_multipoints = 1000; + std::size_t constexpr num_point_per_multipoint = 3; + + rmm::device_uvector geometry_offsets(num_multipoints + 1, this->stream()); + rmm::device_uvector> coordinates(num_multipoints * num_point_per_multipoint, + this->stream()); + + thrust::sequence( + rmm::exec_policy(this->stream()), geometry_offsets.begin(), geometry_offsets.end(), 0, 3); + + thrust::generate_n(rmm::exec_policy(this->stream()), + geometry_offsets.begin(), + geometry_offsets.end(), + [] __device__() { + return vec_2d{0.0, 10.0}; + }); + + this->test_multipoints = + make_multipoint_array(std::move(geometry_offsets), std::move(coordinates)); + } +}; From a924ed7a4f8402f93561f96b19bc9eb8bf379de7 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Mon, 12 Jun 2023 20:28:49 +0000 Subject: [PATCH 02/19] Revert "devcontainer and cmake changes" This reverts commit b4a2421d8a660d8dc47567a99675032aa5bc73f1. --- .devcontainer/devcontainer.json | 6 ++---- cpp/CMakeLists.txt | 6 +++--- 2 files changed, 5 insertions(+), 7 deletions(-) diff --git a/.devcontainer/devcontainer.json b/.devcontainer/devcontainer.json index 5bcea0a1f..902276b71 100644 --- a/.devcontainer/devcontainer.json +++ b/.devcontainer/devcontainer.json @@ -40,16 +40,14 @@ "source=${localWorkspaceFolder}/../.cache,target=/home/coder/.cache,type=bind,consistency=consistent", "source=${localWorkspaceFolder}/../.config,target=/home/coder/.config,type=bind,consistency=consistent", "source=${localWorkspaceFolder}/../.conda/pkgs,target=/home/coder/.conda/pkgs,type=bind,consistency=consistent", - "source=${localWorkspaceFolder}/../.conda/${localWorkspaceFolderBasename}/single,target=/home/coder/.conda/envs,type=bind,consistency=consistent", - "source=${localWorkspaceFolder}/../scratch,target=/home/coder/scratch,type=bind,consistency=consistent" + "source=${localWorkspaceFolder}/../.conda/${localWorkspaceFolderBasename}/single,target=/home/coder/.conda/envs,type=bind,consistency=consistent" ], "customizations": { "vscode": { "extensions": [ "mutantdino.resourcemonitor", - "tamasfe.even-better-toml", - "ms-toolsai.jupyter" + "tamasfe.even-better-toml" ], "settings": { "files.trimFinalNewlines": true, diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index ecaae8302..c3477d10d 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -47,7 +47,7 @@ option(BUILD_BENCHMARKS "Configure CMake to build (google) benchmarks" OFF) option(PER_THREAD_DEFAULT_STREAM "Build with per-thread default stream" OFF) option(DISABLE_DEPRECATION_WARNING "Disable warnings generated from deprecated declarations." OFF) # Option to enable line info in CUDA device compilation to allow introspection when profiling / memchecking -option(CUDA_ENABLE_LINEINFO "Enable the -lineinfo option for nvcc (useful for cuda-memcheck / profiler" ON) +option(CUDA_ENABLE_LINEINFO "Enable the -lineinfo option for nvcc (useful for cuda-memcheck / profiler" OFF) # cudart can be statically linked or dynamically linked. The python ecosystem wants dynamic linking option(CUDA_STATIC_RUNTIME "Statically link the CUDA toolkit runtime and libraries" OFF) @@ -63,12 +63,12 @@ message(STATUS "CUSPATIAL: Enable the -lineinfo option for nvcc (useful for cuda message(STATUS "CUSPATIAL: Statically link the CUDA toolkit runtime and libraries: ${CUDA_STATIC_RUNTIME}") # Set a default build type if none was specified -rapids_cmake_build_type("Debug") +rapids_cmake_build_type("Release") set(CUSPATIAL_BUILD_TESTS ${BUILD_TESTS}) set(CUSPATIAL_BUILD_BENCHMARKS ${BUILD_BENCHMARKS}) set(CUSPATIAL_CXX_FLAGS "") -set(CUSPATIAL_CUDA_FLAGS "-ftemplate-backtrace-limit=0") +set(CUSPATIAL_CUDA_FLAGS "") set(CUSPATIAL_CXX_DEFINITIONS "") set(CUSPATIAL_CUDA_DEFINITIONS "") From ebfa6999012b522c3d27bc8f0e74a4c91a6bfcb7 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Mon, 12 Jun 2023 20:29:56 +0000 Subject: [PATCH 03/19] [skip ci] block ci --- cpp/tests/range/multipoint_range_test.cu | 1 + 1 file changed, 1 insertion(+) diff --git a/cpp/tests/range/multipoint_range_test.cu b/cpp/tests/range/multipoint_range_test.cu index fa80d625d..04ea42182 100644 --- a/cpp/tests/range/multipoint_range_test.cu +++ b/cpp/tests/range/multipoint_range_test.cu @@ -213,6 +213,7 @@ class EmptyMultiPointRangeTest : public MultipointRangeTest { void test_subscript_operator() { + } }; From 0733e336c97fe5b0ed264389b2a494e00d03fddc Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Mon, 12 Jun 2023 20:43:35 -0700 Subject: [PATCH 04/19] implemented multipoint tests --- cpp/include/cuspatial/iterator_factory.cuh | 8 ++ .../cuspatial/range/multipoint_range.cuh | 1 + cpp/tests/range/multipoint_range_test.cu | 83 ++++++++++++++----- 3 files changed, 69 insertions(+), 23 deletions(-) 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/multipoint_range.cuh b/cpp/include/cuspatial/range/multipoint_range.cuh index b41edbd7e..d937b14b0 100644 --- a/cpp/include/cuspatial/range/multipoint_range.cuh +++ b/cpp/include/cuspatial/range/multipoint_range.cuh @@ -148,6 +148,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/tests/range/multipoint_range_test.cu b/cpp/tests/range/multipoint_range_test.cu index 04ea42182..eac90b206 100644 --- a/cpp/tests/range/multipoint_range_test.cu +++ b/cpp/tests/range/multipoint_range_test.cu @@ -14,15 +14,18 @@ * limitations under the License. */ -#include #include #include #include +#include +#include #include #include +#include #include +#include #include #include #include @@ -34,30 +37,39 @@ 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; + for (vec_2d point : multipoints[i]) + output_points[i] = point; +} + template class MultipointRangeTest : public BaseFixture { public: struct copy_leading_point_per_multipoint { - template + template vec_2d __device__ operator()(MultiPointRef multipoint) { return multipoint.size() > 0 ? multipoint[0] : vec_2d{-1, -1}; } }; - template + template struct point_idx_to_geometry_idx { - MultiPointRange rng; std::size_t __device__ operator()(std::size_t pidx) { - return rng.geometry_idx_from_point_idx(pidx); + return rng.geometry_idx_from_point_idx(pidx); } }; void SetUp() { make_test_multipoints(); } - auto range() { return test_multipoints.range(); } + auto range() { return test_multipoints->range(); } virtual void make_test_multipoints() = 0; @@ -83,7 +95,7 @@ class MultipointRangeTest : public BaseFixture { test_subscript_operator(); - test_point(); + test_point_accessor(); test_is_single_point_range(); } @@ -108,7 +120,7 @@ class MultipointRangeTest : public BaseFixture { virtual void test_subscript_operator() = 0; - virtual void test_point() = 0; + virtual void test_point_accessor() = 0; virtual void test_is_single_point_range() = 0; @@ -119,7 +131,7 @@ class MultipointRangeTest : public BaseFixture { rmm::device_uvector> leading_points(rng.num_multipoints(), this->stream()); thrust::transform(rmm::exec_policy(this->stream()), rng.multipoint_begin(), - rng.multipoin_end(), + rng.multipoint_end(), leading_points.begin(), copy_leading_point_per_multipoint{}); @@ -135,12 +147,12 @@ class MultipointRangeTest : public BaseFixture { return points; }; - rmm::device_uvector> copy_offsets() + rmm::device_uvector copy_offsets() { auto rng = this->range(); - rmm::device_uvector offsets(rng.num_multipoints()+1, this->stream()); + rmm::device_uvector offsets(rng.num_multipoints() + 1, this->stream()); thrust::copy( - rmm::exec_policy(this->stream()), rng.offset_begin(), rng.offset_end(), offsets.begin()); + rmm::exec_policy(this->stream()), rng.offsets_begin(), rng.offsets_end(), offsets.begin()); return offsets; }; @@ -150,27 +162,37 @@ class MultipointRangeTest : public BaseFixture { 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} - ); + rmm::exec_policy(this->stream()), idx.begin(), idx.end(), point_idx_to_geometry_idx{rng}); return idx; } 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().geometry_offset_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; } - multipoint_array test_multipoints; + std::unique_ptr, rmm::device_vector>>> + test_multipoints; }; template class EmptyMultiPointRangeTest : public MultipointRangeTest { public: - void make_test_multipoints() { this->test_multipoints = make_multipoint_array({}); } + 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); } @@ -196,8 +218,9 @@ class EmptyMultiPointRangeTest : public MultipointRangeTest { CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(points, expected); } - void test_offsets_it() { - auto offsets = this->copy_all_points(); + void test_offsets_it() + { + auto offsets = this->copy_offsets(); auto expected = make_device_vector({0}); CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(offsets, expected); } @@ -205,19 +228,33 @@ class EmptyMultiPointRangeTest : public MultipointRangeTest { void test_geometry_idx_from_point_idx() { auto geometry_indices = rmm::device_uvector(0, this->stream()); - auto expected = rmm::device_uvector(0, this->stream()) ; + 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: From 4f1a3301e1fb4407ae3bfc8b40cc8312712f385a Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Mon, 12 Jun 2023 22:23:56 -0700 Subject: [PATCH 05/19] add length one test --- cpp/tests/range/multipoint_range_test.cu | 89 ++++++++++++++++++++---- 1 file changed, 77 insertions(+), 12 deletions(-) diff --git a/cpp/tests/range/multipoint_range_test.cu b/cpp/tests/range/multipoint_range_test.cu index eac90b206..c39d10d6f 100644 --- a/cpp/tests/range/multipoint_range_test.cu +++ b/cpp/tests/range/multipoint_range_test.cu @@ -47,6 +47,13 @@ void __global__ array_access_tester(MultiPointRange multipoints, output_points[i] = point; } +template +void __global__ point_accessor_tester(MultiPointRange multipoints, std::size_t i, OutputIt point) +{ + using T = typename MultiPointRange::element_t; + point = multipoints.point(i); +} + template class MultipointRangeTest : public BaseFixture { public: @@ -62,6 +69,8 @@ class MultipointRangeTest : public BaseFixture { 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); @@ -104,13 +113,13 @@ class MultipointRangeTest : public BaseFixture { virtual void test_num_points() = 0; - virtual void test_size() = 0; + void test_size() { EXPECT_EQ(this->range().size(), this->range().num_multipoints()); } virtual void test_multipoint_it() = 0; - virtual void test_begin() = 0; + void test_begin() { EXPECT_EQ(this->range().begin(), this->range().multipoint_begin()); } - virtual void test_end() = 0; + void test_end() { EXPECT_EQ(this->range().end(), this->range().multipoint_end()); } virtual void test_point_it() = 0; @@ -156,7 +165,7 @@ class MultipointRangeTest : public BaseFixture { return offsets; }; - rmm::device_uvector copy_geoemtry_idx() + rmm::device_uvector copy_geometry_idx() { auto rng = this->range(); rmm::device_uvector idx(rng.num_points(), this->stream()); @@ -166,6 +175,17 @@ class MultipointRangeTest : public BaseFixture { return idx; } + rmm::device_scalar> copy_ith_point(std::size_t i) + { + auto rng = this->range(); + + rmm::device_scalar> point(this->stream()); + array_access_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(); @@ -198,8 +218,6 @@ class EmptyMultiPointRangeTest : public MultipointRangeTest { void test_num_points() { EXPECT_EQ(this->range().num_points(), 0); } - void test_size() { EXPECT_EQ(this->range().size(), 0); } - void test_multipoint_it() { auto leading_points = this->copy_leading_points(); @@ -207,10 +225,6 @@ class EmptyMultiPointRangeTest : public MultipointRangeTest { CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(leading_points, expected); } - void test_begin() { EXPECT_EQ(this->range().begin(), this->range().multipoint_begin()); } - - void test_end() { EXPECT_EQ(this->range().end(), this->range().multipoint_end()); } - void test_point_it() { auto points = this->copy_all_points(); @@ -227,7 +241,7 @@ class EmptyMultiPointRangeTest : public MultipointRangeTest { void test_geometry_idx_from_point_idx() { - auto geometry_indices = rmm::device_uvector(0, this->stream()); + auto geometry_indices = this->copy_geometry_idx(); auto expected = rmm::device_uvector(0, this->stream()); CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_indices, expected); @@ -260,8 +274,59 @@ class LengthOneMultiPointRangeTest : public MultipointRangeTest { public: void make_test_multipoints() { - this->test_multipoints = make_multipoint_array({{{1.0, 1.0}}}); + 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_offset_it() + { + auto offsets = this->copy_offsets(); + auto expected = make_device_vector({0, 1}); + 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}); + + 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()); } }; template From 6c24e2b5f6e507c5473cb1073f023fb31364d4f0 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Tue, 13 Jun 2023 10:37:59 -0700 Subject: [PATCH 06/19] add length 5 multipoint range test --- cpp/tests/range/multipoint_range_test.cu | 75 ++++++++++++++++++++++-- 1 file changed, 69 insertions(+), 6 deletions(-) diff --git a/cpp/tests/range/multipoint_range_test.cu b/cpp/tests/range/multipoint_range_test.cu index c39d10d6f..dd6b83886 100644 --- a/cpp/tests/range/multipoint_range_test.cu +++ b/cpp/tests/range/multipoint_range_test.cu @@ -14,6 +14,7 @@ * limitations under the License. */ +#include "gtest/gtest.h" #include #include #include @@ -190,8 +191,7 @@ class MultipointRangeTest : public BaseFixture { { auto rng = this->range(); rmm::device_scalar num_points(this->stream()); - auto count_iterator = - make_count_iterator_from_offset_iterator(this->range().geometry_offset_begin()); + 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()); @@ -274,7 +274,7 @@ class LengthOneMultiPointRangeTest : public MultipointRangeTest { public: void make_test_multipoints() { - auto array = make_multipoint_array({{{1.0, 1.0}}, {10.0, 10.0}}); + auto array = make_multipoint_array({{{1.0, 1.0}, {10.0, 10.0}}}); this->test_multipoints = std::make_unique(std::move(array)); } @@ -285,7 +285,7 @@ class LengthOneMultiPointRangeTest : public MultipointRangeTest { void test_multipoint_it() { auto leading_points = this->copy_leading_points(); - auto expected = make_device_vector>({1.0, 1.0}); + auto expected = make_device_vector>({{1.0, 1.0}}); CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(leading_points, expected); } @@ -296,7 +296,7 @@ class LengthOneMultiPointRangeTest : public MultipointRangeTest { CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(points, expected); } - void test_offset_it() + void test_offsets_it() { auto offsets = this->copy_offsets(); auto expected = make_device_vector({0, 1}); @@ -329,19 +329,82 @@ class LengthOneMultiPointRangeTest : public MultipointRangeTest { 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() { - this->test_multipoints = make_multipoint_array({{{0.0, 0.0}, {1.0, 1.0}}, + 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: From 08ebf2db295aeb96f5d5603e530a0ca703be97f3 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Tue, 13 Jun 2023 10:50:05 -0700 Subject: [PATCH 07/19] fix test case error --- cpp/tests/range/multipoint_range_test.cu | 29 ++++++++++++------------ 1 file changed, 14 insertions(+), 15 deletions(-) diff --git a/cpp/tests/range/multipoint_range_test.cu b/cpp/tests/range/multipoint_range_test.cu index dd6b83886..9a9a6e996 100644 --- a/cpp/tests/range/multipoint_range_test.cu +++ b/cpp/tests/range/multipoint_range_test.cu @@ -44,8 +44,7 @@ void __global__ array_access_tester(MultiPointRange multipoints, OutputIt output_points) { using T = typename MultiPointRange::element_t; - for (vec_2d point : multipoints[i]) - output_points[i] = point; + thrust::copy(thrust::seq, multipoints[i].begin(), multipoints[i].end(), output_points); } template @@ -85,29 +84,29 @@ class MultipointRangeTest : public BaseFixture { void run_test() { - test_num_multipoints(); + // test_num_multipoints(); - test_num_points(); + // test_num_points(); - test_size(); + // test_size(); - test_multipoint_it(); + // test_multipoint_it(); - test_begin(); + // test_begin(); - test_end(); + // test_end(); - test_point_it(); + // test_point_it(); - test_offsets_it(); + // test_offsets_it(); - test_geometry_idx_from_point_idx(); + // test_geometry_idx_from_point_idx(); test_subscript_operator(); - test_point_accessor(); + // test_point_accessor(); - test_is_single_point_range(); + // test_is_single_point_range(); } virtual void test_num_multipoints() = 0; @@ -299,14 +298,14 @@ class LengthOneMultiPointRangeTest : public MultipointRangeTest { void test_offsets_it() { auto offsets = this->copy_offsets(); - auto expected = make_device_vector({0, 1}); + 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}); + auto expected = make_device_vector({0, 0}); CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_indices, expected); } From 613ed2bfd56247839c20841ea71a6ead6b8b7479 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Tue, 13 Jun 2023 11:21:48 -0700 Subject: [PATCH 08/19] fix length five test --- cpp/tests/range/multipoint_range_test.cu | 35 ++++++++++++------------ 1 file changed, 18 insertions(+), 17 deletions(-) diff --git a/cpp/tests/range/multipoint_range_test.cu b/cpp/tests/range/multipoint_range_test.cu index 9a9a6e996..9388d7c3b 100644 --- a/cpp/tests/range/multipoint_range_test.cu +++ b/cpp/tests/range/multipoint_range_test.cu @@ -50,8 +50,8 @@ void __global__ array_access_tester(MultiPointRange multipoints, template void __global__ point_accessor_tester(MultiPointRange multipoints, std::size_t i, OutputIt point) { - using T = typename MultiPointRange::element_t; - point = multipoints.point(i); + using T = typename MultiPointRange::element_t; + point[0] = multipoints.point(i); } template @@ -84,29 +84,29 @@ class MultipointRangeTest : public BaseFixture { void run_test() { - // test_num_multipoints(); + test_num_multipoints(); - // test_num_points(); + test_num_points(); - // test_size(); + test_size(); - // test_multipoint_it(); + test_multipoint_it(); - // test_begin(); + test_begin(); - // test_end(); + test_end(); - // test_point_it(); + test_point_it(); - // test_offsets_it(); + test_offsets_it(); - // test_geometry_idx_from_point_idx(); + test_geometry_idx_from_point_idx(); test_subscript_operator(); - // test_point_accessor(); + test_point_accessor(); - // test_is_single_point_range(); + test_is_single_point_range(); } virtual void test_num_multipoints() = 0; @@ -180,7 +180,7 @@ class MultipointRangeTest : public BaseFixture { auto rng = this->range(); rmm::device_scalar> point(this->stream()); - array_access_tester<<<1, 1, 0, this->stream()>>>(rng, i, point.data()); + point_accessor_tester<<<1, 1, 0, this->stream()>>>(rng, i, point.data()); CUSPATIAL_CHECK_CUDA(this->stream()); return point; @@ -340,7 +340,7 @@ class LengthFiveMultiPointRangeTest : public MultipointRangeTest { {{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)); } @@ -426,7 +426,8 @@ class LengthOneThousandRangeTest : public MultipointRangeTest { return vec_2d{0.0, 10.0}; }); - this->test_multipoints = - make_multipoint_array(std::move(geometry_offsets), std::move(coordinates)); + auto array = make_multipoint_array(std::move(geometry_offsets), std::move(coordinates)); + + this->test_multipoints = std::make_unique(std::move(array)); } }; From e271ff343d02d4ab6cd195decb3ed8c3e96aa25d Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Tue, 13 Jun 2023 12:21:21 -0700 Subject: [PATCH 09/19] add length one thousand test --- .../cuspatial_test/vector_factories.cuh | 12 ++ cpp/tests/range/multipoint_range_test.cu | 112 +++++++++++++++--- 2 files changed, 109 insertions(+), 15 deletions(-) diff --git a/cpp/include/cuspatial_test/vector_factories.cuh b/cpp/include/cuspatial_test/vector_factories.cuh index f2572c17e..0409adf2e 100644 --- a/cpp/include/cuspatial_test/vector_factories.cuh +++ b/cpp/include/cuspatial_test/vector_factories.cuh @@ -406,5 +406,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/range/multipoint_range_test.cu b/cpp/tests/range/multipoint_range_test.cu index 9388d7c3b..bd62994c2 100644 --- a/cpp/tests/range/multipoint_range_test.cu +++ b/cpp/tests/range/multipoint_range_test.cu @@ -31,6 +31,7 @@ #include #include +#include #include #include @@ -407,27 +408,108 @@ 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() { - std::size_t constexpr num_multipoints = 1000; - std::size_t constexpr num_point_per_multipoint = 3; + rmm::device_vector geometry_offsets(num_multipoints + 1); + rmm::device_vector> coordinates(num_points); - rmm::device_uvector geometry_offsets(num_multipoints + 1, this->stream()); - rmm::device_uvector> coordinates(num_multipoints * num_point_per_multipoint, - this->stream()); + thrust::sequence(rmm::exec_policy(this->stream()), + geometry_offsets.begin(), + geometry_offsets.end(), + 0ul, + num_point_per_multipoint); - thrust::sequence( - rmm::exec_policy(this->stream()), geometry_offsets.begin(), geometry_offsets.end(), 0, 3); + thrust::tabulate(rmm::exec_policy(this->stream()), + coordinates.begin(), + coordinates.end(), + [] __device__(auto i) { + return vec_2d{static_cast(i), 10.0}; + }); - thrust::generate_n(rmm::exec_policy(this->stream()), - geometry_offsets.begin(), - geometry_offsets.end(), - [] __device__() { - return vec_2d{0.0, 10.0}; - }); - - auto array = make_multipoint_array(std::move(geometry_offsets), std::move(coordinates)); + 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(); } From 9f79718fc288cfcae70eefcd0b3d6859090f9013 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Tue, 13 Jun 2023 12:31:12 -0700 Subject: [PATCH 10/19] [skip ci] block ci --- cpp/tests/range/multipoint_range_test.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/tests/range/multipoint_range_test.cu b/cpp/tests/range/multipoint_range_test.cu index bd62994c2..303998dcf 100644 --- a/cpp/tests/range/multipoint_range_test.cu +++ b/cpp/tests/range/multipoint_range_test.cu @@ -14,7 +14,6 @@ * limitations under the License. */ -#include "gtest/gtest.h" #include #include #include @@ -157,6 +156,7 @@ class MultipointRangeTest : public BaseFixture { return points; }; + rmm::device_uvector copy_offsets() { auto rng = this->range(); From 75c5872f996e39bdc56f146142ed7b95f22be4ae Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Tue, 18 Jul 2023 19:43:29 -0700 Subject: [PATCH 11/19] define multilinestring test interface --- cpp/tests/range/multilinestring_range_test.cu | 108 ++++++++++++++++++ 1 file changed, 108 insertions(+) diff --git a/cpp/tests/range/multilinestring_range_test.cu b/cpp/tests/range/multilinestring_range_test.cu index 250e88355..2ac96e4c2 100644 --- a/cpp/tests/range/multilinestring_range_test.cu +++ b/cpp/tests/range/multilinestring_range_test.cu @@ -533,3 +533,111 @@ 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 MultilinestringRangeTest2 { + 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_part_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_offsets_it(); + test_part_offsets_it(); + + test_as_multipoint_range(); + } + + void test_size() { EXPECT_EQ(this->range().size(), this->range().num_muiltilinestrings()); } + + virtual void test_num_multilinestrings() = 0; + + virtual void test_num_linestrings() = 0; + + virtual void test_num_points() = 0; + + virtual void test_multilinestring_it(); + + 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_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_part_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; + + virtual void test_geometry_offsets_it() = 0; + virtual void test_part_offsets_it() = 0; + + virtual void test_as_multipoint_range() = 0; + + private: + std::unique_ptr, + rmm::device_vector, + rmm::device_vector>>> + test_multilinestring; +}; From 882b67a8eae0a0dcf83bd262fbd075e8c09937e6 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Sun, 23 Jul 2023 19:07:43 -0700 Subject: [PATCH 12/19] add empty test for multilinestring_range --- .../detail/range/multilinestring_range.cuh | 7 +- .../cuspatial/range/multilinestring_range.cuh | 3 +- cpp/tests/range/multilinestring_range_test.cu | 402 +++++++++++++++++- 3 files changed, 393 insertions(+), 19 deletions(-) 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/range/multilinestring_range.cuh b/cpp/include/cuspatial/range/multilinestring_range.cuh index b4bb6eaf5..6cf6dec49 100644 --- a/cpp/include/cuspatial/range/multilinestring_range.cuh +++ b/cpp/include/cuspatial/range/multilinestring_range.cuh @@ -144,8 +144,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(); diff --git a/cpp/tests/range/multilinestring_range_test.cu b/cpp/tests/range/multilinestring_range_test.cu index 2ac96e4c2..1f28f1872 100644 --- a/cpp/tests/range/multilinestring_range_test.cu +++ b/cpp/tests/range/multilinestring_range_test.cu @@ -22,14 +22,26 @@ #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) +{ + using T = typename MultiLineStringRange::element_t; + 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, @@ -535,12 +547,90 @@ TYPED_TEST(MultilinestringRangeTest, MultilinestringAsMultipointTest6) } template -class MultilinestringRangeTest2 { +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_index_functor { + MultiLineStringRange _rng; + std::size_t __device__ operator()(std::size_t segment_idx) + { + return _rng.geometry_idx_from_segment_idx(segment_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) { return _rng.segment(i); } + }; + void SetUp() { make_test_multilinestring(); } virtual void make_test_multilinestring() = 0; - auto range() { return test_multilinestring.range(); } + auto range() { return test_multilinestring->range(); } void run_test() { @@ -566,8 +656,6 @@ class MultilinestringRangeTest2 { test_part_idx_from_segment_idx(); - test_geometry_idx_from_part_idx(); - test_geometry_idx_from_point_idx(); test_intra_part_idx(); @@ -585,12 +673,11 @@ class MultilinestringRangeTest2 { test_array_access_operator(); test_geometry_offsets_it(); - test_part_offsets_it(); - test_as_multipoint_range(); + test_part_offsets_it(); } - void test_size() { EXPECT_EQ(this->range().size(), this->range().num_muiltilinestrings()); } + void test_size() { EXPECT_EQ(this->range().size(), this->range().num_multilinestrings()); } virtual void test_num_multilinestrings() = 0; @@ -598,7 +685,7 @@ class MultilinestringRangeTest2 { virtual void test_num_points() = 0; - virtual void test_multilinestring_it(); + virtual void test_multilinestring_it() = 0; void test_begin() { EXPECT_EQ(this->range().begin(), this->range().multilinestring_begin()); } @@ -612,8 +699,6 @@ class MultilinestringRangeTest2 { virtual void test_part_idx_from_segment_idx() = 0; - virtual void test_geometry_idx_from_part_idx() = 0; - virtual void test_geometry_idx_from_point_idx() = 0; virtual void test_intra_part_idx() = 0; @@ -631,13 +716,304 @@ class MultilinestringRangeTest2 { virtual void test_array_access_operator() = 0; virtual void test_geometry_offsets_it() = 0; - virtual void test_part_offsets_it() = 0; - virtual void test_as_multipoint_range() = 0; + virtual void test_part_offsets_it() = 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_part_offset() + { + auto rng = this->range(); + auto d_part_offset = rmm::device_uvector(rng.num_multilinestrings() + 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_points(), 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_multipoint(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; + } + + 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_offsets_begin(), + rng.geometry_offsets_end(), + d_geometry_offsets.begin()); + return d_geometry_offsets; + } + + rmm::device_uvector copy_part_offsets() + { + auto rng = this->range(); + auto d_part_offsets = rmm::device_uvector(rng.num_linestrings() + 1, stream()); + thrust::copy(rmm::exec_policy(stream()), + rng.part_offsets_begin(), + rng.part_offsets_end(), + d_part_offsets.begin()); + return d_part_offsets; + } - private: + 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_offsets_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(); } From 5e326c91dbfb6960fd9db928cc8aaceb7e994e1e Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Sun, 23 Jul 2023 20:19:30 -0700 Subject: [PATCH 13/19] add size one multilinestring range test --- .../range/multilinestring_segment_range.cuh | 2 +- .../cuspatial/range/multilinestring_range.cuh | 13 +- cpp/tests/range/multilinestring_range_test.cu | 193 ++++++++++++++---- 3 files changed, 165 insertions(+), 43 deletions(-) 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/range/multilinestring_range.cuh b/cpp/include/cuspatial/range/multilinestring_range.cuh index 6cf6dec49..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; } @@ -167,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/tests/range/multilinestring_range_test.cu b/cpp/tests/range/multilinestring_range_test.cu index 1f28f1872..246fedb2a 100644 --- a/cpp/tests/range/multilinestring_range_test.cu +++ b/cpp/tests/range/multilinestring_range_test.cu @@ -21,6 +21,7 @@ #include +#include #include #include #include @@ -623,7 +624,13 @@ class MultilinestringRangeTestBase : public BaseFixture { template struct segment_functor { MultiLineStringRange _rng; - segment __device__ operator()(std::size_t i) { return _rng.segment(i); } + 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(); } @@ -672,9 +679,7 @@ class MultilinestringRangeTestBase : public BaseFixture { test_array_access_operator(); - test_geometry_offsets_it(); - - test_part_offsets_it(); + test_geometry_offset_it(); } void test_size() { EXPECT_EQ(this->range().size(), this->range().num_multilinestrings()); } @@ -693,6 +698,8 @@ class MultilinestringRangeTestBase : public BaseFixture { 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; @@ -715,10 +722,6 @@ class MultilinestringRangeTestBase : public BaseFixture { virtual void test_array_access_operator() = 0; - virtual void test_geometry_offsets_it() = 0; - - virtual void test_part_offsets_it() = 0; - // Helper functions to be used by all subclass (test cases). rmm::device_uvector> copy_leading_points() { @@ -741,10 +744,22 @@ class MultilinestringRangeTestBase : public BaseFixture { 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_multilinestrings() + 1, stream()); + 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(), @@ -785,7 +800,7 @@ class MultilinestringRangeTestBase : public BaseFixture { rmm::device_uvector copy_intra_part_idx() { auto rng = this->range(); - auto d_intra_part_idx = rmm::device_uvector(rng.num_points(), stream()); + 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(), @@ -850,7 +865,7 @@ class MultilinestringRangeTestBase : public BaseFixture { return d_multilinestring_linestring_count; } - rmm::device_uvector> copy_all_points_of_ith_multipoint(std::size_t i) + rmm::device_uvector> copy_all_points_of_ith_multilinestring(std::size_t i) { auto rng = this->range(); rmm::device_scalar num_points(stream()); @@ -866,29 +881,6 @@ class MultilinestringRangeTestBase : public BaseFixture { 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_offsets_begin(), - rng.geometry_offsets_end(), - d_geometry_offsets.begin()); - return d_geometry_offsets; - } - - rmm::device_uvector copy_part_offsets() - { - auto rng = this->range(); - auto d_part_offsets = rmm::device_uvector(rng.num_linestrings() + 1, stream()); - thrust::copy(rmm::exec_policy(stream()), - rng.part_offsets_begin(), - rng.part_offsets_end(), - d_part_offsets.begin()); - return d_part_offsets; - } - protected: std::unique_ptr, rmm::device_vector, @@ -1000,7 +992,7 @@ class MultilinestringRangeEmptyTest : public MultilinestringRangeTestBase { SUCCEED(); } - void test_geometry_offsets_it() + void test_geometry_offset_it() { auto geometry_offsets = this->copy_geometry_offsets(); auto expected = make_device_vector({0}); @@ -1017,3 +1009,134 @@ class MultilinestringRangeEmptyTest : public MultilinestringRangeTestBase { 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(); } From b00606ce6b61ab8e3d6ee5a2f7d922a434ad3630 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Mon, 24 Jul 2023 08:37:36 +0000 Subject: [PATCH 14/19] Add one thousand multilinestring test --- .../multilinestring_ref.cuh | 6 +- .../cuspatial/detail/utility/zero_data.cuh | 2 + .../cuspatial_test/vector_factories.cuh | 8 +- cpp/tests/range/multilinestring_range_test.cu | 223 +++++++++++++++++- 4 files changed, 231 insertions(+), 8 deletions(-) diff --git a/cpp/include/cuspatial/detail/geometry_collection/multilinestring_ref.cuh b/cpp/include/cuspatial/detail/geometry_collection/multilinestring_ref.cuh index 7a76f92b8..ba2950a4c 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 is the one-past the last part index to the point 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/utility/zero_data.cuh b/cpp/include/cuspatial/detail/utility/zero_data.cuh index c6196a4c2..1a4f72a54 100644 --- a/cpp/include/cuspatial/detail/utility/zero_data.cuh +++ b/cpp/include/cuspatial/detail/utility/zero_data.cuh @@ -19,6 +19,8 @@ #include +#include + namespace cuspatial { namespace detail { diff --git a/cpp/include/cuspatial_test/vector_factories.cuh b/cpp/include/cuspatial_test/vector_factories.cuh index 73f167ae6..5b359fce2 100644 --- a/cpp/include/cuspatial_test/vector_factories.cuh +++ b/cpp/include/cuspatial_test/vector_factories.cuh @@ -280,13 +280,11 @@ auto make_multilinestring_array(IndexRangeA geometry_inl, CoordRange coord_inl) { using CoordType = typename CoordRange::value_type; - using DeviceIndexVector = thrust::device_vector; - using DeviceCoordVector = thrust::device_vector; + using DeviceIndexVector = rmm::device_vector; + using DeviceCoordVector = rmm::device_vector; return multilinestring_array( - make_device_vector(std::move(geometry_inl)), - make_device_vector(std::move(part_inl)), - make_device_vector(std::move(coord_inl))); + std::move(geometry_inl), std::move(part_inl), std::move(coord_inl)); } /** diff --git a/cpp/tests/range/multilinestring_range_test.cu b/cpp/tests/range/multilinestring_range_test.cu index 246fedb2a..7bce1de42 100644 --- a/cpp/tests/range/multilinestring_range_test.cu +++ b/cpp/tests/range/multilinestring_range_test.cu @@ -19,9 +19,9 @@ #include #include +#include #include -#include #include #include #include @@ -29,9 +29,11 @@ #include #include +#include #include #include +#include using namespace cuspatial; using namespace cuspatial::test; @@ -1140,3 +1142,222 @@ class MultilinestringRangeOneTest : public MultilinestringRangeTestBase { 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(); } From e72ec4bfdc9e4c90362d33bfffa5ee89f8ee54d9 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Mon, 24 Jul 2023 20:33:58 +0000 Subject: [PATCH 15/19] Add empty multipolygon range test --- .../detail/range/multipolygon_range.cuh | 41 -- .../cuspatial/range/multipolygon_range.cuh | 23 +- cpp/tests/range/multilinestring_range_test.cu | 10 - cpp/tests/range/multipolygon_range_test.cu | 391 +++++++++++++++++- 4 files changed, 386 insertions(+), 79 deletions(-) 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 ; 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/tests/range/multilinestring_range_test.cu b/cpp/tests/range/multilinestring_range_test.cu index 7bce1de42..11c17e5ac 100644 --- a/cpp/tests/range/multilinestring_range_test.cu +++ b/cpp/tests/range/multilinestring_range_test.cu @@ -41,7 +41,6 @@ using namespace cuspatial::test; template void __global__ array_access_tester(MultiLineStringRange mls, std::size_t i, OutputIt output_points) { - using T = typename MultiLineStringRange::element_t; thrust::copy(thrust::seq, mls[i].point_begin(), mls[i].point_end(), output_points); } @@ -592,15 +591,6 @@ class MultilinestringRangeTestBase : public BaseFixture { } }; - template - struct intra_index_functor { - MultiLineStringRange _rng; - std::size_t __device__ operator()(std::size_t segment_idx) - { - return _rng.geometry_idx_from_segment_idx(segment_idx); - } - }; - template struct intra_part_idx_functor { MultiLineStringRange _rng; diff --git a/cpp/tests/range/multipolygon_range_test.cu b/cpp/tests/range/multipolygon_range_test.cu index 2fe133945..a23c68fcc 100644 --- a/cpp/tests/range/multipolygon_range_test.cu +++ b/cpp/tests/range/multipolygon_range_test.cu @@ -22,10 +22,13 @@ #include #include +#include #include #include #include +#include + #include using namespace cuspatial; @@ -121,10 +124,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 +489,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 +501,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 +535,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 +645,373 @@ 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[0].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::copy(rmm::exec_policy(stream()), rng.point_begin(), rng.point_end(), d_points.begin()); + 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(); } From 9c655e804689ad1135756a1ac26dc452e028f786 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Mon, 24 Jul 2023 19:35:22 -0700 Subject: [PATCH 16/19] add one multipolygon test --- cpp/tests/range/multipolygon_range_test.cu | 123 +++++++++++++++++++++ 1 file changed, 123 insertions(+) diff --git a/cpp/tests/range/multipolygon_range_test.cu b/cpp/tests/range/multipolygon_range_test.cu index a23c68fcc..4714dba6a 100644 --- a/cpp/tests/range/multipolygon_range_test.cu +++ b/cpp/tests/range/multipolygon_range_test.cu @@ -1015,3 +1015,126 @@ class MultipolygonRangeEmptyTest : public MultipolygonRangeTestBase { 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(); } From 75a8657d2b3a1ad6c55da13379b7fa5811015a88 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Mon, 24 Jul 2023 22:59:50 -0700 Subject: [PATCH 17/19] add one thousands test for multipolygon range, and fix polygon/multipolygon_ref --- .../cuspatial/detail/geometry/polygon_ref.cuh | 4 +- .../geometry_collection/multipolygon_ref.cuh | 4 +- .../cuspatial_test/geometry_generator.cuh | 2 +- .../cuspatial_test/vector_factories.cuh | 53 +++-- cpp/tests/range/multipolygon_range_test.cu | 186 +++++++++++++++++- 5 files changed, 228 insertions(+), 21 deletions(-) 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/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_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 5b359fce2..65c1c7140 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)) { } @@ -271,19 +282,33 @@ class multilinestring_array { * @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 = rmm::device_vector; - using DeviceCoordVector = rmm::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( +/** + * @brief Construct an owning object of a multilinestring array from ranges + * + * @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)); } diff --git a/cpp/tests/range/multipolygon_range_test.cu b/cpp/tests/range/multipolygon_range_test.cu index 4714dba6a..1ba2d9935 100644 --- a/cpp/tests/range/multipolygon_range_test.cu +++ b/cpp/tests/range/multipolygon_range_test.cu @@ -27,6 +27,7 @@ #include #include +#include #include #include @@ -659,7 +660,7 @@ class MultipolygonRangeTestBase : public BaseFixture { template __device__ vec_2d operator()(MultiPolygonRef mpolygon) { - return mpolygon.size() > 0 ? mpolygon[0].point_begin()[0] : vec_2d{-1, -1}; + return mpolygon.size() > 0 ? mpolygon.point_begin()[0] : vec_2d{-1, -1}; } }; @@ -776,7 +777,11 @@ class MultipolygonRangeTestBase : public BaseFixture { { auto rng = range(); auto d_points = rmm::device_uvector>(rng.num_multipolygons(), stream()); - thrust::copy(rmm::exec_policy(stream()), rng.point_begin(), rng.point_end(), d_points.begin()); + thrust::transform(rmm::exec_policy(stream()), + rng.multipolygon_begin(), + rng.multipolygon_end(), + d_points.begin(), + copy_leading_point_functor{}); return d_points; } @@ -1138,3 +1143,180 @@ class MultipolygonRangeOneTest : public MultipolygonRangeTestBase { 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(); } From 293e4c30b5b43ea57225fe4e52a55e18f4e99e3a Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Tue, 25 Jul 2023 06:00:44 +0000 Subject: [PATCH 18/19] style fix --- cpp/tests/range/multipoint_range_test.cu | 1 - 1 file changed, 1 deletion(-) diff --git a/cpp/tests/range/multipoint_range_test.cu b/cpp/tests/range/multipoint_range_test.cu index 303998dcf..501ddd6cf 100644 --- a/cpp/tests/range/multipoint_range_test.cu +++ b/cpp/tests/range/multipoint_range_test.cu @@ -156,7 +156,6 @@ class MultipointRangeTest : public BaseFixture { return points; }; - rmm::device_uvector copy_offsets() { auto rng = this->range(); From 7dfc481dcbed35d05e31c258bc6765b388714936 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 27 Jul 2023 03:04:04 +0000 Subject: [PATCH 19/19] Address review comments --- .../detail/geometry_collection/multilinestring_ref.cuh | 2 +- cpp/include/cuspatial_test/vector_factories.cuh | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/cpp/include/cuspatial/detail/geometry_collection/multilinestring_ref.cuh b/cpp/include/cuspatial/detail/geometry_collection/multilinestring_ref.cuh index ba2950a4c..15472485f 100644 --- a/cpp/include/cuspatial/detail/geometry_collection/multilinestring_ref.cuh +++ b/cpp/include/cuspatial/detail/geometry_collection/multilinestring_ref.cuh @@ -80,7 +80,7 @@ CUSPATIAL_HOST_DEVICE auto multilinestring_ref::point template CUSPATIAL_HOST_DEVICE auto multilinestring_ref::point_end() const { - // _part_end is the one-past the last part index to the point of this multilinestring. + // _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)); } diff --git a/cpp/include/cuspatial_test/vector_factories.cuh b/cpp/include/cuspatial_test/vector_factories.cuh index 65c1c7140..cbf1d8f56 100644 --- a/cpp/include/cuspatial_test/vector_factories.cuh +++ b/cpp/include/cuspatial_test/vector_factories.cuh @@ -275,7 +275,7 @@ 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 @@ -294,7 +294,7 @@ auto make_multilinestring_array(rmm::device_uvector&& geometry_inl, } /** - * @brief Construct an owning object of a multilinestring array from ranges + * @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