Skip to content

Commit

Permalink
Remove GDAL dependency in quadtree spatial join tests. (rapidsai#974)
Browse files Browse the repository at this point in the history
Closes rapidsai#958.  Replaces usage of GDAL `contains` in the quadtree point-in-polygon tests with a call to a different cuSpatial point-in-polygon function. This allows us to remove all current dependencies on GDAL in cuSpatial.

Depends on rapidsai#973. This PR currently updates some expectations to make tests pass that are fixed in that PR.

Authors:
  - Mark Harris (https://github.com/harrism)

Approvers:
  - Michael Wang (https://github.com/isVoid)
  - H. Thomson Comer (https://github.com/thomcom)
  - Paul Taylor (https://github.com/trxcllnt)
  - Jordan Jacobelli (https://github.com/jjacobelli)

URL: rapidsai#974
  • Loading branch information
harrism authored and trxcllnt committed Mar 15, 2023
1 parent 21481e8 commit 4283625
Show file tree
Hide file tree
Showing 6 changed files with 40 additions and 127 deletions.
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 @@ -40,7 +40,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 @@ -88,6 +87,5 @@ outputs:
run:
- {{ pin_subpackage('libcuspatial', exact=True) }}
- cudatoolkit {{ cuda_spec }}
- gdal {{ gdal_version }}
- gmock {{ gtest_version }}
- gtest {{ gtest_version }}
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];
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

0 comments on commit 4283625

Please sign in to comment.