Merge branch 'branch-24.10' into my_new_branch

rapidsai · Aug 21, 2024 · e6fc59e · e6fc59e
2 parents 6a5eaaa + f678ad5
commit e6fc59e
Show file tree

Hide file tree

Showing 52 changed files with 176 additions and 124 deletions.
diff --git a/.github/workflows/pr.yaml b/.github/workflows/pr.yaml
@@ -14,6 +14,7 @@ jobs:
     needs:
       - checks
       - conda-cpp-build
+      - conda-cpp-checks
       - conda-cpp-tests
       - conda-python-build
       - conda-python-tests
@@ -37,6 +38,13 @@ jobs:
     uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: pull-request
+  conda-cpp-checks:
+    needs: conda-cpp-build
+    secrets: inherit
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
+    with:
+      build_type: pull-request
+      enable_check_symbols: true
   conda-cpp-tests:
     needs: conda-cpp-build
     secrets: inherit

diff --git a/.github/workflows/test.yaml b/.github/workflows/test.yaml
@@ -14,6 +14,15 @@ on:
         type: string
 
 jobs:
+  conda-cpp-checks:
+    secrets: inherit
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
+    with:
+      build_type: nightly
+      branch: ${{ inputs.branch }}
+      date: ${{ inputs.date }}
+      sha: ${{ inputs.sha }}
+      enable_check_symbols: true
   conda-cpp-tests:
     secrets: inherit
     uses: rapidsai/shared-workflows/.github/workflows/[email protected]

diff --git a/ci/release/update-version.sh b/ci/release/update-version.sh
@@ -43,6 +43,7 @@ done
 sed_runner "s/RAPIDS_VERSION_NUMBER=\".*/RAPIDS_VERSION_NUMBER=\"${NEXT_SHORT_TAG}\"/g" ci/build_docs.sh
 
 DEPENDENCIES=(
+  pylibcudf
   cudf
   cuml
   cuspatial

diff --git a/conda/environments/all_cuda-118_arch-x86_64.yaml b/conda/environments/all_cuda-118_arch-x86_64.yaml
@@ -39,6 +39,7 @@ dependencies:
 - proj
 - pydata-sphinx-theme!=0.14.2
 - pydeck
+- pylibcudf==24.10.*,>=0.0.0a0
 - pytest
 - pytest-cov
 - pytest-xdist

diff --git a/conda/environments/all_cuda-125_arch-x86_64.yaml b/conda/environments/all_cuda-125_arch-x86_64.yaml
@@ -41,6 +41,7 @@ dependencies:
 - proj
 - pydata-sphinx-theme!=0.14.2
 - pydeck
+- pylibcudf==24.10.*,>=0.0.0a0
 - pytest
 - pytest-cov
 - pytest-xdist

diff --git a/conda/recipes/cuspatial/meta.yaml b/conda/recipes/cuspatial/meta.yaml
@@ -60,6 +60,7 @@ requirements:
     {% endif %}
     - cuda-version ={{ cuda_version }}
     - cudf ={{ minor_version }}
+    - pylibcudf ={{ minor_version }}
     - cython >=3.0.0
     - libcuspatial ={{ version }}
     - python
@@ -74,6 +75,7 @@ requirements:
     {% endif %}
     - {{ pin_compatible('cuda-version', max_pin='x', min_pin='x') }}
     - cudf ={{ minor_version }}
+    - pylibcudf ={{ minor_version }}
     - geopandas >=0.11.0
     - numpy >=1.23,<3.0a0
     - python

diff --git a/cpp/include/cuspatial/cuda_utils.hpp b/cpp/include/cuspatial/cuda_utils.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -19,8 +19,10 @@
 
 #ifdef __CUDACC__
 #define CUSPATIAL_HOST_DEVICE __host__ __device__
+#define CUSPATIAL_KERNEL      __global__ static
 #else
 #define CUSPATIAL_HOST_DEVICE
+#define CUSPATIAL_KERNEL
 #endif
 
 /**

diff --git a/cpp/include/cuspatial/detail/distance/hausdorff.cuh b/cpp/include/cuspatial/detail/distance/hausdorff.cuh
@@ -82,7 +82,7 @@ constexpr auto magnitude_squared(T a, T b)
  * @param results directed Hausdorff distances computed by kernel
  */
 template <typename T, typename Index, typename PointIt, typename OffsetIt, typename OutputIt>
-__global__ void kernel_hausdorff(
+CUSPATIAL_KERNEL void kernel_hausdorff(
   Index num_points, PointIt points, Index num_spaces, OffsetIt space_offsets, OutputIt results)
 {
   using Point = typename std::iterator_traits<PointIt>::value_type;

diff --git a/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh b/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh
@@ -48,9 +48,9 @@ namespace detail {
  * 2. Repeat 1 until all mergeable group is processed.
  */
 template <typename OffsetRange, typename SegmentRange, typename OutputIt>
-void __global__ simple_find_and_combine_segments_kernel(OffsetRange offsets,
-                                                        SegmentRange segments,
-                                                        OutputIt merged_flag)
+CUSPATIAL_KERNEL void simple_find_and_combine_segments_kernel(OffsetRange offsets,
+                                                              SegmentRange segments,
+                                                              OutputIt merged_flag)
 {
   for (auto pair_idx : ranger::grid_stride_range(offsets.size() - 1)) {
     // Zero-initialize flags for all segments in current space.

diff --git a/cpp/include/cuspatial/detail/find/find_duplicate_points.cuh b/cpp/include/cuspatial/detail/find/find_duplicate_points.cuh
@@ -34,8 +34,8 @@ namespace detail {
  * @brief Kernel to compute duplicate points in each multipoint. Naive N^2 algorithm.
  */
 template <typename MultiPointRange, typename OutputIt>
-void __global__ find_duplicate_points_kernel_simple(MultiPointRange multipoints,
-                                                    OutputIt duplicate_flags)
+CUSPATIAL_KERNEL void find_duplicate_points_kernel_simple(MultiPointRange multipoints,
+                                                          OutputIt duplicate_flags)
 {
   for (auto idx : ranger::grid_stride_range(multipoints.size())) {
     auto multipoint    = multipoints[idx];

diff --git a/cpp/include/cuspatial/detail/intersection/linestring_intersection_count.cuh b/cpp/include/cuspatial/detail/intersection/linestring_intersection_count.cuh
@@ -34,10 +34,11 @@ template <typename MultiLinestringRange1,
           typename MultiLinestringRange2,
           typename OutputIt1,
           typename OutputIt2>
-__global__ void count_intersection_and_overlaps_simple(MultiLinestringRange1 multilinestrings1,
-                                                       MultiLinestringRange2 multilinestrings2,
-                                                       OutputIt1 point_count_it,
-                                                       OutputIt2 segment_count_it)
+CUSPATIAL_KERNEL void count_intersection_and_overlaps_simple(
+  MultiLinestringRange1 multilinestrings1,
+  MultiLinestringRange2 multilinestrings2,
+  OutputIt1 point_count_it,
+  OutputIt2 segment_count_it)
 {
   using T = typename MultiLinestringRange1::element_t;
   for (auto idx : ranger::grid_stride_range(multilinestrings1.num_points())) {

diff --git a/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh b/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh
@@ -394,17 +394,18 @@ template <typename MultiLinestringRange1,
           typename IdRanges,
           typename OutputIt1,
           typename OutputIt2>
-void __global__ pairwise_linestring_intersection_simple(MultiLinestringRange1 multilinestrings1,
-                                                        MultiLinestringRange2 multilinestrings2,
-                                                        TempIt1 n_points_stored,
-                                                        TempIt2 n_segments_stored,
-                                                        Offsets1 num_points_offsets_first,
-                                                        Offsets2 num_segments_offsets_first,
-                                                        Offsets3 num_points_per_pair_first,
-                                                        IdRanges point_ids_range,
-                                                        IdRanges segment_ids_range,
-                                                        OutputIt1 points_first,
-                                                        OutputIt2 segments_first)
+CUSPATIAL_KERNEL void pairwise_linestring_intersection_simple(
+  MultiLinestringRange1 multilinestrings1,
+  MultiLinestringRange2 multilinestrings2,
+  TempIt1 n_points_stored,
+  TempIt2 n_segments_stored,
+  Offsets1 num_points_offsets_first,
+  Offsets2 num_segments_offsets_first,
+  Offsets3 num_points_per_pair_first,
+  IdRanges point_ids_range,
+  IdRanges segment_ids_range,
+  OutputIt1 points_first,
+  OutputIt2 segments_first)
 {
   using T       = typename MultiLinestringRange1::element_t;
   using types_t = uint8_t;

diff --git a/cpp/include/cuspatial/detail/kernel/pairwise_distance.cuh b/cpp/include/cuspatial/detail/kernel/pairwise_distance.cuh
@@ -48,10 +48,10 @@ namespace detail {
  * @note This kernel does not compute pairs that contains empty geometry.
  */
 template <class MultiLinestringRange1, class MultiLinestringRange2, class OutputIt>
-__global__ void linestring_distance(MultiLinestringRange1 multilinestrings1,
-                                    MultiLinestringRange2 multilinestrings2,
-                                    thrust::optional<uint8_t*> intersects,
-                                    OutputIt distances_first)
+CUSPATIAL_KERNEL void linestring_distance(MultiLinestringRange1 multilinestrings1,
+                                          MultiLinestringRange2 multilinestrings2,
+                                          thrust::optional<uint8_t*> intersects,
+                                          OutputIt distances_first)
 {
   using T = typename MultiLinestringRange1::element_t;
 
@@ -91,10 +91,10 @@ __global__ void linestring_distance(MultiLinestringRange1 multilinestrings1,
  * set to nullopt, no distance computation will be bypassed.
  */
 template <class MultiPointRange, class MultiLinestringRange, class OutputIterator>
-__global__ void point_linestring_distance(MultiPointRange multipoints,
-                                          MultiLinestringRange multilinestrings,
-                                          thrust::optional<uint8_t*> intersects,
-                                          OutputIterator distances)
+CUSPATIAL_KERNEL void point_linestring_distance(MultiPointRange multipoints,
+                                                MultiLinestringRange multilinestrings,
+                                                thrust::optional<uint8_t*> intersects,
+                                                OutputIterator distances)
 {
   using T = typename MultiPointRange::element_t;
 

diff --git a/cpp/include/cuspatial/detail/pairwise_multipoint_equals_count.cuh b/cpp/include/cuspatial/detail/pairwise_multipoint_equals_count.cuh
@@ -43,9 +43,9 @@ namespace cuspatial {
 namespace detail {
 
 template <class MultiPointRangeA, class MultiPointRangeB, class OutputIt>
-void __global__ pairwise_multipoint_equals_count_kernel(MultiPointRangeA lhs,
-                                                        MultiPointRangeB rhs,
-                                                        OutputIt output)
+CUSPATIAL_KERNEL void pairwise_multipoint_equals_count_kernel(MultiPointRangeA lhs,
+                                                              MultiPointRangeB rhs,
+                                                              OutputIt output)
 {
   using T = typename MultiPointRangeA::point_t::value_type;
 

diff --git a/cpp/include/cuspatial/detail/point_linestring_nearest_points.cuh b/cpp/include/cuspatial/detail/point_linestring_nearest_points.cuh
@@ -54,18 +54,18 @@ template <class Vec2dItA,
           class OffsetIteratorB,
           class OffsetIteratorC,
           class OutputIt>
-void __global__
-pairwise_point_linestring_nearest_points_kernel(OffsetIteratorA points_geometry_offsets_first,
-                                                OffsetIteratorA points_geometry_offsets_last,
-                                                Vec2dItA points_first,
-                                                Vec2dItA points_last,
-                                                OffsetIteratorB linestring_geometry_offsets_first,
-                                                OffsetIteratorB linestring_geometry_offsets_last,
-                                                OffsetIteratorC linestring_part_offsets_first,
-                                                OffsetIteratorC linestring_part_offsets_last,
-                                                Vec2dItB linestring_points_first,
-                                                Vec2dItB linestring_points_last,
-                                                OutputIt output_first)
+CUSPATIAL_KERNEL void pairwise_point_linestring_nearest_points_kernel(
+  OffsetIteratorA points_geometry_offsets_first,
+  OffsetIteratorA points_geometry_offsets_last,
+  Vec2dItA points_first,
+  Vec2dItA points_last,
+  OffsetIteratorB linestring_geometry_offsets_first,
+  OffsetIteratorB linestring_geometry_offsets_last,
+  OffsetIteratorC linestring_part_offsets_first,
+  OffsetIteratorC linestring_part_offsets_last,
+  Vec2dItB linestring_points_first,
+  Vec2dItB linestring_points_last,
+  OutputIt output_first)
 {
   using T         = iterator_vec_base_type<Vec2dItA>;
   using IndexType = iterator_value_type<OffsetIteratorA>;

diff --git a/cpp/include/cuspatial_test/geometry_generator.cuh b/cpp/include/cuspatial_test/geometry_generator.cuh
@@ -194,8 +194,8 @@ ring_centroid_displacement(vec_2d<T> centroid, std::size_t ring_local_idx, T rad
  * @param params Parameters to generate the mulitpolygons
  */
 template <typename T, typename MultipolygonRange>
-void __global__ generate_multipolygon_array_coordinates(MultipolygonRange multipolygons,
-                                                        multipolygon_generator_parameter<T> params)
+CUSPATIAL_KERNEL void generate_multipolygon_array_coordinates(
+  MultipolygonRange multipolygons, multipolygon_generator_parameter<T> params)
 {
   for (auto idx : ranger::grid_stride_range(multipolygons.num_points())) {
     auto ring_idx     = multipolygons.ring_idx_from_point_idx(idx);

diff --git a/cpp/tests/operators/linestrings_test.cu b/cpp/tests/operators/linestrings_test.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -69,10 +69,10 @@ segment<T> __device__ order_end_points(segment<T> const& seg)
 }
 
 template <typename T, typename Point, typename Segment>
-__global__ void compute_intersection(segment<T> ab,
-                                     segment<T> cd,
-                                     Point point_out,
-                                     Segment segment_out)
+CUSPATIAL_KERNEL void compute_intersection(segment<T> ab,
+                                           segment<T> cd,
+                                           Point point_out,
+                                           Segment segment_out)
 {
   auto [p, s]    = detail::segment_intersection(ab, cd);
   point_out[0]   = p;

diff --git a/cpp/tests/range/multilinestring_range_test.cu b/cpp/tests/range/multilinestring_range_test.cu
@@ -38,7 +38,9 @@ using namespace cuspatial;
 using namespace cuspatial::test;
 
 template <typename MultiLineStringRange, typename OutputIt>
-void __global__ array_access_tester(MultiLineStringRange mls, std::size_t i, OutputIt output_points)
+CUSPATIAL_KERNEL void array_access_tester(MultiLineStringRange mls,
+                                          std::size_t i,
+                                          OutputIt output_points)
 {
   thrust::copy(thrust::seq, mls[i].point_begin(), mls[i].point_end(), output_points);
 }

diff --git a/cpp/tests/range/multipoint_range_test.cu b/cpp/tests/range/multipoint_range_test.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -39,16 +39,18 @@ using namespace cuspatial;
 using namespace cuspatial::test;
 
 template <typename MultiPointRange, typename OutputIt>
-void __global__ array_access_tester(MultiPointRange multipoints,
-                                    std::size_t i,
-                                    OutputIt output_points)
+CUSPATIAL_KERNEL void array_access_tester(MultiPointRange multipoints,
+                                          std::size_t i,
+                                          OutputIt output_points)
 {
   using T = typename MultiPointRange::element_t;
   thrust::copy(thrust::seq, multipoints[i].begin(), multipoints[i].end(), output_points);
 }
 
 template <typename MultiPointRange, typename OutputIt>
-void __global__ point_accessor_tester(MultiPointRange multipoints, std::size_t i, OutputIt point)
+CUSPATIAL_KERNEL void point_accessor_tester(MultiPointRange multipoints,
+                                            std::size_t i,
+                                            OutputIt point)
 {
   using T  = typename MultiPointRange::element_t;
   point[0] = multipoints.point(i);

diff --git a/cpp/tests/range/multipolygon_range_test.cu b/cpp/tests/range/multipolygon_range_test.cu
@@ -648,7 +648,9 @@ TYPED_TEST(MultipolygonRangeTest, MultipolygonAsMultiPoint3)
 }
 
 template <typename MultiPolygonRange, typename PointOutputIt>
-__global__ void array_access_tester(MultiPolygonRange rng, std::size_t i, PointOutputIt output)
+CUSPATIAL_KERNEL 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);
 }

diff --git a/dependencies.yaml b/dependencies.yaml
@@ -424,6 +424,7 @@ dependencies:
       - output_types: conda
         packages:
           - &cudf_unsuffixed cudf==24.10.*,>=0.0.0a0
+          - &pylibcudf_unsuffixed pylibcudf==24.10.*,>=0.0.0a0
       - output_types: requirements
         packages:
           # pip recognizes the index as a global option for the requirements.txt file
@@ -437,11 +438,13 @@ dependencies:
               cuda_suffixed: "true"
             packages:
               - cudf-cu12==24.10.*,>=0.0.0a0
+              - pylibcudf-cu12==24.10.*,>=0.0.0a0
           - matrix:
               cuda: "11.*"
               cuda_suffixed: "true"
             packages:
               - cudf-cu11==24.10.*,>=0.0.0a0
+              - pylibcudf-cu11==24.10.*,>=0.0.0a0
           - {matrix: null, packages: [*cudf_unsuffixed]}
 
   depends_on_cuml:

diff --git a/python/cuproj/pyproject.toml b/python/cuproj/pyproject.toml
@@ -65,7 +65,7 @@ force_grid_wrap = 0
 combine_as_imports = true
 order_by_type = true
 known_dask = ["dask", "distributed", "dask_cuda"]
-known_rapids = ["rmm", "cudf"]
+known_rapids = ["rmm", "cudf", "pylibcudf"]
 known_first_party = ["cuspatial", "cuproj"]
 default_section = "THIRDPARTY"
 sections = [

diff --git a/python/cuspatial/cuspatial/_lib/cpp/column/geometry_column_view.pxd b/python/cuspatial/cuspatial/_lib/cpp/column/geometry_column_view.pxd
@@ -1,8 +1,9 @@
-# Copyright (c) 2023, NVIDIA CORPORATION.
+# Copyright (c) 2023-2024, NVIDIA CORPORATION.
 
 from libcpp.memory cimport unique_ptr
 
-from cudf._lib.column cimport column, column_view
+from pylibcudf.libcudf.column.column cimport column
+from pylibcudf.libcudf.column.column_view cimport column_view
 
 from cuspatial._lib.cpp.types cimport collection_type_id, geometry_type_id
 

diff --git a/python/cuspatial/cuspatial/_lib/cpp/distance.pxd b/python/cuspatial/cuspatial/_lib/cpp/distance.pxd
@@ -3,9 +3,9 @@
 from libcpp.memory cimport unique_ptr
 from libcpp.utility cimport pair
 
-from cudf._lib.pylibcudf.libcudf.column.column cimport column
-from cudf._lib.pylibcudf.libcudf.column.column_view cimport column_view
-from cudf._lib.pylibcudf.libcudf.table.table_view cimport table_view
+from pylibcudf.libcudf.column.column cimport column
+from pylibcudf.libcudf.column.column_view cimport column_view
+from pylibcudf.libcudf.table.table_view cimport table_view
 
 from cuspatial._lib.cpp.column.geometry_column_view cimport (
     geometry_column_view,

diff --git a/python/cuspatial/cuspatial/_lib/cpp/distance/polygon_distance.pxd b/python/cuspatial/cuspatial/_lib/cpp/distance/polygon_distance.pxd
@@ -2,7 +2,7 @@
 
 from libcpp.memory cimport unique_ptr
 
-from cudf._lib.pylibcudf.libcudf.column.column cimport column
+from pylibcudf.libcudf.column.column cimport column
 
 from cuspatial._lib.cpp.column.geometry_column_view cimport (
     geometry_column_view,