diff --git a/conda/environments/all_cuda-118_arch-x86_64.yaml b/conda/environments/all_cuda-118_arch-x86_64.yaml index 805292e2e..40a0a0ed4 100644 --- a/conda/environments/all_cuda-118_arch-x86_64.yaml +++ b/conda/environments/all_cuda-118_arch-x86_64.yaml @@ -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 diff --git a/conda/recipes/libcuspatial/conda_build_config.yaml b/conda/recipes/libcuspatial/conda_build_config.yaml index 0692c0651..a50e5c3ff 100644 --- a/conda/recipes/libcuspatial/conda_build_config.yaml +++ b/conda/recipes/libcuspatial/conda_build_config.yaml @@ -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" diff --git a/conda/recipes/libcuspatial/meta.yaml b/conda/recipes/libcuspatial/meta.yaml index 9b2a8079a..f9e62e61d 100644 --- a/conda/recipes/libcuspatial/meta.yaml +++ b/conda/recipes/libcuspatial/meta.yaml @@ -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 }} @@ -88,6 +87,5 @@ outputs: run: - {{ pin_subpackage('libcuspatial', exact=True) }} - cudatoolkit {{ cuda_spec }} - - gdal {{ gdal_version }} - gmock {{ gtest_version }} - gtest {{ gtest_version }} diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index d9cd05e6b..20154fed1 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -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} @@ -38,7 +30,7 @@ function(ConfigureTest CMAKE_TEST_NAME) PROPERTIES RUNTIME_OUTPUT_DIRECTORY "$" 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} diff --git a/cpp/tests/join/point_in_polygon_test_large.cpp b/cpp/tests/join/point_in_polygon_test_large.cpp index 38874f6d7..9afc54506 100644 --- a/cpp/tests/join/point_in_polygon_test_large.cpp +++ b/cpp/tests/join/point_in_polygon_test_large.cpp @@ -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. @@ -15,6 +15,7 @@ */ #include +#include #include #include #include @@ -25,31 +26,18 @@ #include #include -#include - #include #include #include #include #include -#include -#include -#include -#include #include -#include - -// #include - /* - * 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}; @@ -78,57 +66,6 @@ inline auto generate_points(std::vector> const& quads, uint32_t p return std::make_pair(std::move(point_x), std::move(point_y)); } -template -inline auto make_polygons_geometry(thrust::host_vector const& poly_offsets, - thrust::host_vector const& ring_offsets, - thrust::host_vector const& poly_x, - thrust::host_vector const& poly_y) -{ - std::vector polygons{}; - for (uint32_t poly_idx = 0, poly_end = poly_offsets.size() - 1; poly_idx < poly_end; ++poly_idx) { - auto ring_idx = static_cast(poly_offsets[poly_idx]); - auto ring_end = static_cast(poly_offsets[poly_idx + 1]); - auto polygon = static_cast(OGRGeometryFactory::createGeometry(wkbPolygon)); - for (; ring_idx < ring_end; ++ring_idx) { - auto seg_idx = static_cast(ring_offsets[ring_idx]); - auto seg_end = static_cast( - ring_idx < ring_offsets.size() - 1 ? ring_offsets[ring_idx + 1] : poly_x.size()); - auto ring = static_cast(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 -auto geometry_to_poly_and_point_indices(std::vector const& polygons, - thrust::host_vector const& x, - thrust::host_vector const& y) -{ - std::vector poly_indices{}; - std::vector point_lengths{}; - std::vector point_indices{}; - - for (uint32_t i = 0, n = x.size(); i < n; i++) { - OGRPoint point(x[i], y[i]); - std::vector 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; @@ -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(poly_offsets).first, - cudf::test::to_host(ring_offsets).first, - cudf::test::to_host(poly_x).first, - cudf::test::to_host(poly_y).first); - - auto host_poly_and_point_indices = - geometry_to_poly_and_point_indices(h_poly, - cudf::test::to_host(points->get_column(0)).first, - cudf::test::to_host(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(poly_idx).first; auto actual_point_indices = cudf::test::to_host(point_idx).first; - auto actual_point_lengths = thrust::host_vector(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(hits->view()).first; - auto poly_a = fixed_width_column_wrapper(expected_poly_indices.begin(), - expected_poly_indices.end()); - auto poly_b = - fixed_width_column_wrapper(actual_poly_indices.begin(), actual_poly_indices.end()); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(poly_a, poly_b, verbosity); + std::vector expected_poly_indices; + std::vector expected_point_indices; - auto point_a = fixed_width_column_wrapper(expected_point_indices.begin(), - expected_point_indices.end()); - auto point_b = - fixed_width_column_wrapper(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(expected_point_lengths.begin(), - expected_point_lengths.end()); - auto lengths_b = - fixed_width_column_wrapper(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(expected_poly_indices.begin(), + expected_poly_indices.end()); + auto poly_b = + fixed_width_column_wrapper(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(expected_point_indices.begin(), + expected_point_indices.end()); + auto point_b = fixed_width_column_wrapper(actual_point_indices.begin(), + actual_point_indices.end()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(point_a, point_b, verbosity); + } } diff --git a/dependencies.yaml b/dependencies.yaml index bf561c893..6346cae7a 100644 --- a/dependencies.yaml +++ b/dependencies.yaml @@ -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