From 1d7aab212de859f54b9efee73df2d2ccd144e939 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Thu, 8 Aug 2024 11:56:47 -0400 Subject: [PATCH 1/5] cuspatial CUDA kernels are now all marked with internal linkage --- cpp/include/cuspatial/cuda_utils.hpp | 4 +++- .../cuspatial/detail/distance/hausdorff.cuh | 2 +- .../detail/find/find_and_combine_segment.cuh | 6 ++--- .../detail/find/find_duplicate_points.cuh | 4 ++-- .../linestring_intersection_count.cuh | 8 +++---- ...inestring_intersection_with_duplicates.cuh | 22 +++++++++---------- .../detail/kernel/pairwise_distance.cuh | 16 +++++++------- .../pairwise_multipoint_equals_count.cuh | 6 ++--- .../point_linestring_nearest_points.cuh | 2 +- .../cuspatial_test/geometry_generator.cuh | 4 ++-- cpp/tests/operators/linestrings_test.cu | 10 ++++----- cpp/tests/range/multilinestring_range_test.cu | 2 +- cpp/tests/range/multipoint_range_test.cu | 6 ++--- cpp/tests/range/multipolygon_range_test.cu | 2 +- 14 files changed, 48 insertions(+), 46 deletions(-) diff --git a/cpp/include/cuspatial/cuda_utils.hpp b/cpp/include/cuspatial/cuda_utils.hpp index f46fddfcb..27039e9b9 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..dd45bb763 100644 --- a/cpp/include/cuspatial/detail/intersection/linestring_intersection_count.cuh +++ b/cpp/include/cuspatial/detail/intersection/linestring_intersection_count.cuh @@ -34,10 +34,10 @@ 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..7cc597848 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,17 @@ 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..f62c44e05 100644 --- a/cpp/include/cuspatial/detail/point_linestring_nearest_points.cuh +++ b/cpp/include/cuspatial/detail/point_linestring_nearest_points.cuh @@ -54,7 +54,7 @@ template -void __global__ +void CUSPATIAL_KERNEL pairwise_point_linestring_nearest_points_kernel(OffsetIteratorA points_geometry_offsets_first, OffsetIteratorA points_geometry_offsets_last, Vec2dItA points_first, diff --git a/cpp/include/cuspatial_test/geometry_generator.cuh b/cpp/include/cuspatial_test/geometry_generator.cuh index 8c17ff291..7321bdf98 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..d84c82f11 100644 --- a/cpp/tests/range/multilinestring_range_test.cu +++ b/cpp/tests/range/multilinestring_range_test.cu @@ -38,7 +38,7 @@ 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..e84dd1b6e 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,7 +39,7 @@ using namespace cuspatial; using namespace cuspatial::test; template -void __global__ array_access_tester(MultiPointRange multipoints, +CUSPATIAL_KERNEL void array_access_tester(MultiPointRange multipoints, std::size_t i, OutputIt output_points) { @@ -48,7 +48,7 @@ void __global__ array_access_tester(MultiPointRange multipoints, } 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..388675b83 100644 --- a/cpp/tests/range/multipolygon_range_test.cu +++ b/cpp/tests/range/multipolygon_range_test.cu @@ -648,7 +648,7 @@ 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); } From ceb58852db96ea75058f359d72e9494aca0ea93b Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Thu, 8 Aug 2024 11:59:25 -0400 Subject: [PATCH 2/5] Add CI checks for external CUDA kernels --- .github/workflows/pr.yaml | 8 ++++++++ .github/workflows/test.yaml | 9 +++++++++ 2 files changed, 17 insertions(+) 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 From dd658e0b699ba0bc0726da055cf1022099f3cec4 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Thu, 8 Aug 2024 13:05:49 -0400 Subject: [PATCH 3/5] Correct issues found by CI --- cpp/include/cuspatial/cuda_utils.hpp | 2 +- .../linestring_intersection_count.cuh | 9 ++++---- ...inestring_intersection_with_duplicates.cuh | 23 ++++++++++--------- .../cuspatial_test/geometry_generator.cuh | 4 ++-- cpp/tests/range/multilinestring_range_test.cu | 4 +++- cpp/tests/range/multipoint_range_test.cu | 8 ++++--- cpp/tests/range/multipolygon_range_test.cu | 4 +++- 7 files changed, 31 insertions(+), 23 deletions(-) diff --git a/cpp/include/cuspatial/cuda_utils.hpp b/cpp/include/cuspatial/cuda_utils.hpp index 27039e9b9..28353ad2e 100644 --- a/cpp/include/cuspatial/cuda_utils.hpp +++ b/cpp/include/cuspatial/cuda_utils.hpp @@ -19,7 +19,7 @@ #ifdef __CUDACC__ #define CUSPATIAL_HOST_DEVICE __host__ __device__ -#define CUSPATIAL_KERNEL __global__ static +#define CUSPATIAL_KERNEL __global__ static #else #define CUSPATIAL_HOST_DEVICE #define CUSPATIAL_KERNEL diff --git a/cpp/include/cuspatial/detail/intersection/linestring_intersection_count.cuh b/cpp/include/cuspatial/detail/intersection/linestring_intersection_count.cuh index dd45bb763..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 -CUSPATIAL_KERNEL 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 7cc597848..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 -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) +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_test/geometry_generator.cuh b/cpp/include/cuspatial_test/geometry_generator.cuh index 7321bdf98..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 -CUSPATIAL_KERNEL void 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/range/multilinestring_range_test.cu b/cpp/tests/range/multilinestring_range_test.cu index d84c82f11..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 -CUSPATIAL_KERNEL void 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 e84dd1b6e..3e7546bc5 100644 --- a/cpp/tests/range/multipoint_range_test.cu +++ b/cpp/tests/range/multipoint_range_test.cu @@ -40,15 +40,17 @@ using namespace cuspatial::test; template CUSPATIAL_KERNEL void array_access_tester(MultiPointRange multipoints, - std::size_t i, - OutputIt output_points) + 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 -CUSPATIAL_KERNEL void 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 388675b83..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 -CUSPATIAL_KERNEL 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); } From d488e587994138f08567719ef0fb909534dae30e Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Mon, 12 Aug 2024 11:29:33 -0400 Subject: [PATCH 4/5] Consistent ordering of CUSPATIAL_KERNEL and void --- .../cuspatial/detail/point_linestring_nearest_points.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/include/cuspatial/detail/point_linestring_nearest_points.cuh b/cpp/include/cuspatial/detail/point_linestring_nearest_points.cuh index f62c44e05..339678865 100644 --- a/cpp/include/cuspatial/detail/point_linestring_nearest_points.cuh +++ b/cpp/include/cuspatial/detail/point_linestring_nearest_points.cuh @@ -54,7 +54,7 @@ template -void CUSPATIAL_KERNEL +CUSPATIAL_KERNEL void pairwise_point_linestring_nearest_points_kernel(OffsetIteratorA points_geometry_offsets_first, OffsetIteratorA points_geometry_offsets_last, Vec2dItA points_first, From e01e2754c185a6d69c40327039cbf69c379f0eba Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Mon, 12 Aug 2024 15:57:01 -0400 Subject: [PATCH 5/5] Correct issues found by CI --- .../point_linestring_nearest_points.cuh | 24 +++++++++---------- 1 file changed, 12 insertions(+), 12 deletions(-) diff --git a/cpp/include/cuspatial/detail/point_linestring_nearest_points.cuh b/cpp/include/cuspatial/detail/point_linestring_nearest_points.cuh index 339678865..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 -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) +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;