diff --git a/.github/workflows/pr.yaml b/.github/workflows/pr.yaml index 6923ee31f..ecf4206c7 100644 --- 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/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 diff --git a/.github/workflows/test.yaml b/.github/workflows/test.yaml index 05323103b..493eca446 100644 --- 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/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 diff --git a/cpp/include/cuspatial/cuda_utils.hpp b/cpp/include/cuspatial/cuda_utils.hpp index f46fddfcb..28353ad2e 100644 --- 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 index aa955a7d6..aa3484fa7 100644 --- 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 -__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::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 index 2cdbfd425..ca86096be 100644 --- 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 -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 index de18bb927..f78424ebe 100644 --- 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 -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 index b164fa729..19009c06c 100644 --- 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 -__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 index f1a365555..5d2165437 100644 --- 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 -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 index 9786681d9..e5eaf7259 100644 --- 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 -__global__ void linestring_distance(MultiLinestringRange1 multilinestrings1, - MultiLinestringRange2 multilinestrings2, - thrust::optional intersects, - OutputIt distances_first) +CUSPATIAL_KERNEL void linestring_distance(MultiLinestringRange1 multilinestrings1, + MultiLinestringRange2 multilinestrings2, + thrust::optional 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 -__global__ void point_linestring_distance(MultiPointRange multipoints, - MultiLinestringRange multilinestrings, - thrust::optional intersects, - OutputIterator distances) +CUSPATIAL_KERNEL void point_linestring_distance(MultiPointRange multipoints, + MultiLinestringRange multilinestrings, + thrust::optional 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 index 141a85d61..157c32485 100644 --- 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 -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 index f085204eb..13270a7cc 100644 --- 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 -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; using IndexType = iterator_value_type; diff --git a/cpp/include/cuspatial_test/geometry_generator.cuh b/cpp/include/cuspatial_test/geometry_generator.cuh index 8c17ff291..fbf6f02d6 100644 --- 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 centroid, std::size_t ring_local_idx, T rad * @param params Parameters to generate the mulitpolygons */ template -void __global__ generate_multipolygon_array_coordinates(MultipolygonRange multipolygons, - multipolygon_generator_parameter params) +CUSPATIAL_KERNEL void generate_multipolygon_array_coordinates( + MultipolygonRange multipolygons, multipolygon_generator_parameter 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 index 7216f276a..2fa44cb43 100644 --- 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 __device__ order_end_points(segment const& seg) } template -__global__ void compute_intersection(segment ab, - segment cd, - Point point_out, - Segment segment_out) +CUSPATIAL_KERNEL void compute_intersection(segment ab, + segment 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 index 4938e0df4..b4d1fe6de 100644 --- 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 -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 index 501ddd6cf..3e7546bc5 100644 --- 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 -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 -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 index 1e93ef136..f39ffb7e7 100644 --- 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 -__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); }