Internal linkage kernels

.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/conda-cpp-build.yaml@branch-24.10
     with:
       build_type: pull-request
+  conda-cpp-checks:
+    needs: conda-cpp-build
+    secrets: inherit
+    uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-post-build-checks.yaml@branch-24.10
+    with:
+      build_type: pull-request
+      enable_check_symbols: true
   conda-cpp-tests:
     needs: conda-cpp-build
     secrets: inherit

.github/workflows/test.yaml

@@ -14,6 +14,15 @@ on:
         type: string
 
 jobs:
+  conda-cpp-checks:
+    secrets: inherit
+    uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-post-build-checks.yaml@branch-24.10
+    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/conda-cpp-tests.yaml@branch-24.10

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
 
 /**

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;

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.

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];

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())) {

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;

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;
 

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;
 

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>;

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);

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;

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);
 }

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);

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);
 }