Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Remove GDAL dependency in quadtree spatial join tests. #974

Merged
merged 4 commits into from
Mar 13, 2023
Merged
Show file tree
Hide file tree
Changes from 3 commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 0 additions & 1 deletion conda/environments/all_cuda-118_arch-x86_64.yaml
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,6 @@ dependencies:
- cython>=0.29,<0.30
- doxygen
- gcc_linux-64=11.*
- gdal>3.5.0,<3.6.0
- geopandas>=0.11.0
- gmock=1.10.0
- gtest=1.10.0
Expand Down
3 changes: 0 additions & 3 deletions conda/recipes/libcuspatial/conda_build_config.yaml
Original file line number Diff line number Diff line change
Expand Up @@ -10,9 +10,6 @@ cuda_compiler:
cmake_version:
- ">=3.23.1,!=3.25.0"

gdal_version:
- ">3.5.0,<3.6.0"

gtest_version:
- "1.10.0"

Expand Down
2 changes: 0 additions & 2 deletions conda/recipes/libcuspatial/meta.yaml
Original file line number Diff line number Diff line change
Expand Up @@ -38,7 +38,6 @@ requirements:
- sysroot_{{ target_platform }} {{ sysroot_version }}
host:
- cudatoolkit ={{ cuda_version }}
- gdal {{ gdal_version }}
- gmock {{ gtest_version }}
- gtest {{ gtest_version }}
- libcudf ={{ minor_version }}
Expand Down Expand Up @@ -86,6 +85,5 @@ outputs:
run:
- {{ pin_subpackage('libcuspatial', exact=True) }}
- cudatoolkit {{ cuda_spec }}
- gdal {{ gdal_version }}
- gmock {{ gtest_version }}
- gtest {{ gtest_version }}
Original file line number Diff line number Diff line change
Expand Up @@ -102,8 +102,8 @@ OutputIt point_in_polygon(Cart2dItA test_points_first,
using T = iterator_vec_base_type<Cart2dItA>;

auto const num_test_points = std::distance(test_points_first, test_points_last);
auto const num_polys = std::distance(polygon_offsets_first, polygon_offsets_last);
auto const num_rings = std::distance(poly_ring_offsets_first, poly_ring_offsets_last);
auto const num_polys = std::distance(polygon_offsets_first, polygon_offsets_last) - 1;
auto const num_rings = std::distance(poly_ring_offsets_first, poly_ring_offsets_last) - 1;
auto const num_poly_points = std::distance(polygon_points_first, polygon_points_last);

static_assert(is_same_floating_point<T, iterator_vec_base_type<Cart2dItB>>(),
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/spatial/point_in_polygon.cu
Original file line number Diff line number Diff line change
Expand Up @@ -143,7 +143,7 @@ std::unique_ptr<cudf::column> point_in_polygon(cudf::column_view const& test_poi
CUSPATIAL_EXPECTS(poly_ring_offsets.size() >= poly_offsets.size(),
"Each polygon must have at least one ring");

CUSPATIAL_EXPECTS(poly_points_x.size() >= poly_offsets.size() * 4,
CUSPATIAL_EXPECTS(poly_points_x.size() >= (poly_ring_offsets.size() - 1) * 4,
"Each ring must have at least four vertices");

return cuspatial::detail::point_in_polygon(test_points_x,
Expand Down
10 changes: 1 addition & 9 deletions cpp/tests/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -17,14 +17,6 @@
###################################################################################################
# - compiler function -----------------------------------------------------------------------------

# find gdal
rapids_find_package(
GDAL REQUIRED
GLOBAL_TARGETS GDAL::GDAL
BUILD_EXPORT_SET cuspatial-exports
INSTALL_EXPORT_SET cuspatial-exports
)

function(ConfigureTest CMAKE_TEST_NAME)
add_executable(${CMAKE_TEST_NAME} ${ARGN})
target_compile_options(${CMAKE_TEST_NAME}
Expand All @@ -38,7 +30,7 @@ function(ConfigureTest CMAKE_TEST_NAME)
PROPERTIES RUNTIME_OUTPUT_DIRECTORY "$<BUILD_INTERFACE:${CUSPATIAL_BINARY_DIR}/gtests>"
INSTALL_RPATH "\$ORIGIN/../../../lib"
)
target_link_libraries(${CMAKE_TEST_NAME} GDAL::GDAL GTest::gtest_main GTest::gmock_main cudf::cudftestutil cuspatial)
target_link_libraries(${CMAKE_TEST_NAME} GTest::gtest_main GTest::gmock_main cudf::cudftestutil cuspatial)
add_test(NAME ${CMAKE_TEST_NAME} COMMAND ${CMAKE_TEST_NAME})
install(
TARGETS ${CMAKE_TEST_NAME}
Expand Down
150 changes: 39 additions & 111 deletions cpp/tests/join/point_in_polygon_test_large.cpp
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2020-2021, NVIDIA CORPORATION.
* Copyright (c) 2020-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.
Expand All @@ -15,6 +15,7 @@
*/

#include <cuspatial/error.hpp>
#include <cuspatial/point_in_polygon.hpp>
#include <cuspatial/point_quadtree.hpp>
#include <cuspatial/polygon_bounding_box.hpp>
#include <cuspatial/spatial_join.hpp>
Expand All @@ -25,31 +26,18 @@
#include <cudf/table/table.hpp>
#include <cudf/table/table_view.hpp>

#include <rmm/device_uvector.hpp>

#include <cudf_test/base_fixture.hpp>
#include <cudf_test/column_utilities.hpp>
#include <cudf_test/column_wrapper.hpp>
#include <cudf_test/table_utilities.hpp>
#include <cudf_test/type_lists.hpp>

#include <thrust/distance.h>
#include <thrust/host_vector.h>
#include <thrust/iterator/constant_iterator.h>
#include <thrust/reduce.h>
#include <thrust/sort.h>

#include <ogrsf_frmts.h>

// #include <algorithm>

/*
* The test uses the same quadtree structure as in pip_refine_test_small.
* However, the numbers of randomly generated points under all quadrants (min_size) are increased
* to be more than the number of threads per-block (currently fixed to 256, but can be set between
* 32 2048 (CUDA Compute Capacity 7.0, multiples of warp size, which is 32) The test is designed to
* fully test the two kernels in the refinment code, including both warp level reduce and scan, vote
* and popc. Thrust primitives to divide quadrants into sub-blocks are also tested.
* The test uses the same quadtree structure as in pip_refine_test_small. However, the number of
* randomly generated points under all quadrants (min_size) are increased to be more than the
* number of threads per-block.
*/

constexpr cudf::test::debug_output_level verbosity{cudf::test::debug_output_level::ALL_ERRORS};
Expand Down Expand Up @@ -78,57 +66,6 @@ inline auto generate_points(std::vector<std::vector<T>> const& quads, uint32_t p
return std::make_pair(std::move(point_x), std::move(point_y));
}

template <typename T>
inline auto make_polygons_geometry(thrust::host_vector<uint32_t> const& poly_offsets,
thrust::host_vector<uint32_t> const& ring_offsets,
thrust::host_vector<T> const& poly_x,
thrust::host_vector<T> const& poly_y)
{
std::vector<OGRGeometry*> polygons{};
for (uint32_t poly_idx = 0, poly_end = poly_offsets.size() - 1; poly_idx < poly_end; ++poly_idx) {
auto ring_idx = static_cast<size_t>(poly_offsets[poly_idx]);
auto ring_end = static_cast<size_t>(poly_offsets[poly_idx + 1]);
auto polygon = static_cast<OGRPolygon*>(OGRGeometryFactory::createGeometry(wkbPolygon));
for (; ring_idx < ring_end; ++ring_idx) {
auto seg_idx = static_cast<size_t>(ring_offsets[ring_idx]);
auto seg_end = static_cast<size_t>(
ring_idx < ring_offsets.size() - 1 ? ring_offsets[ring_idx + 1] : poly_x.size());
auto ring = static_cast<OGRLineString*>(OGRGeometryFactory::createGeometry(wkbLinearRing));
for (; seg_idx < seg_end; ++seg_idx) {
ring->addPoint(poly_x[seg_idx], poly_y[seg_idx]);
}
polygon->addRing(ring);
}
polygons.push_back(polygon);
}
return polygons;
}

template <typename T>
auto geometry_to_poly_and_point_indices(std::vector<OGRGeometry*> const& polygons,
thrust::host_vector<T> const& x,
thrust::host_vector<T> const& y)
{
std::vector<uint32_t> poly_indices{};
std::vector<uint32_t> point_lengths{};
std::vector<uint32_t> point_indices{};

for (uint32_t i = 0, n = x.size(); i < n; i++) {
OGRPoint point(x[i], y[i]);
std::vector<uint32_t> found_poly_idxs{};
for (uint32_t j = 0; j < polygons.size(); j++) {
if (polygons[j]->Contains(&point)) { found_poly_idxs.push_back(j); }
}
if (found_poly_idxs.size() > 0) {
point_lengths.push_back(found_poly_idxs.size());
point_indices.push_back(i);
poly_indices.insert(poly_indices.end(), found_poly_idxs.begin(), found_poly_idxs.end());
}
}
return std::make_tuple(
std::move(poly_indices), std::move(point_indices), std::move(point_lengths));
}

TYPED_TEST(PIPRefineTestLarge, TestLarge)
{
using T = TypeParam;
Expand Down Expand Up @@ -235,57 +172,48 @@ TYPED_TEST(PIPRefineTestLarge, TestLarge)
auto poly_idx = point_in_polygon_pairs->get_column(0).view();
auto point_idx = point_in_polygon_pairs->get_column(1).view();

// verify

auto h_poly = make_polygons_geometry(cudf::test::to_host<uint32_t>(poly_offsets).first,
cudf::test::to_host<uint32_t>(ring_offsets).first,
cudf::test::to_host<T>(poly_x).first,
cudf::test::to_host<T>(poly_y).first);

auto host_poly_and_point_indices =
geometry_to_poly_and_point_indices(h_poly,
cudf::test::to_host<T>(points->get_column(0)).first,
cudf::test::to_host<T>(points->get_column(1)).first);

auto& expected_poly_indices = std::get<0>(host_poly_and_point_indices);
auto& expected_point_indices = std::get<1>(host_poly_and_point_indices);
auto& expected_point_lengths = std::get<2>(host_poly_and_point_indices);

auto actual_poly_indices = cudf::test::to_host<uint32_t>(poly_idx).first;
auto actual_point_indices = cudf::test::to_host<uint32_t>(point_idx).first;
auto actual_point_lengths = thrust::host_vector<uint32_t>(point_in_polygon_pairs->num_rows());

thrust::stable_sort_by_key(
actual_point_indices.begin(), actual_point_indices.end(), actual_poly_indices.begin());

auto num_search_points = thrust::distance(actual_point_indices.begin(),
thrust::reduce_by_key(actual_point_indices.begin(),
actual_point_indices.end(),
thrust::make_constant_iterator(1),
actual_point_indices.begin(),
actual_point_lengths.begin())
.first);
{
// verify

auto hits = cuspatial::point_in_polygon(
points->get_column(0), points->get_column(1), poly_offsets, ring_offsets, poly_x, poly_y);

actual_point_indices.resize(num_search_points);
actual_point_lengths.resize(num_search_points);
actual_point_indices.shrink_to_fit();
actual_point_lengths.shrink_to_fit();
auto hits_host = cudf::test::to_host<int32_t>(hits->view()).first;

auto poly_a = fixed_width_column_wrapper<uint32_t>(expected_poly_indices.begin(),
expected_poly_indices.end());
auto poly_b =
fixed_width_column_wrapper<uint32_t>(actual_poly_indices.begin(), actual_poly_indices.end());
CUDF_TEST_EXPECT_COLUMNS_EQUAL(poly_a, poly_b, verbosity);
std::vector<uint32_t> expected_poly_indices;
std::vector<uint32_t> expected_point_indices;

auto point_a = fixed_width_column_wrapper<uint32_t>(expected_point_indices.begin(),
expected_point_indices.end());
auto point_b =
fixed_width_column_wrapper<uint32_t>(actual_point_indices.begin(), actual_point_indices.end());
CUDF_TEST_EXPECT_COLUMNS_EQUAL(point_a, point_b, verbosity);
for (int point_index = 0; point_index < hits->size(); point_index++) {
// iterate over set bits
std::uint32_t bits = hits_host[point_index];
trxcllnt marked this conversation as resolved.
Show resolved Hide resolved
while (bits != 0) {
std::uint32_t t = bits & -bits; // get only LSB
std::uint32_t poly_index = __builtin_ctz(bits); // get index of LSB
expected_poly_indices.push_back(poly_index);
expected_point_indices.push_back(point_index);
bits ^= t; // reset LSB to zero to advance to next set bit
}
}

auto lengths_a = fixed_width_column_wrapper<uint32_t>(expected_point_lengths.begin(),
expected_point_lengths.end());
auto lengths_b =
fixed_width_column_wrapper<uint32_t>(actual_point_lengths.begin(), actual_point_lengths.end());
CUDF_TEST_EXPECT_COLUMNS_EQUAL(lengths_a, lengths_b, verbosity);
thrust::stable_sort_by_key(
expected_point_indices.begin(), expected_point_indices.end(), expected_poly_indices.begin());

auto poly_a = fixed_width_column_wrapper<uint32_t>(expected_poly_indices.begin(),
expected_poly_indices.end());
auto poly_b =
fixed_width_column_wrapper<uint32_t>(actual_poly_indices.begin(), actual_poly_indices.end());
CUDF_TEST_EXPECT_COLUMNS_EQUAL(poly_a, poly_b, verbosity);

auto point_a = fixed_width_column_wrapper<uint32_t>(expected_point_indices.begin(),
expected_point_indices.end());
auto point_b = fixed_width_column_wrapper<uint32_t>(actual_point_indices.begin(),
actual_point_indices.end());
CUDF_TEST_EXPECT_COLUMNS_EQUAL(point_a, point_b, verbosity);
}
}
1 change: 0 additions & 1 deletion dependencies.yaml
Original file line number Diff line number Diff line change
Expand Up @@ -53,7 +53,6 @@ dependencies:
- &cmake_ver cmake>=3.23.1,!=3.25.0
- c-compiler
- cxx-compiler
- gdal>3.5.0,<3.6.0
- gmock=1.10.0
- gtest=1.10.0
- libcudf=23.04
Expand Down