From 50146a8bec088069eb30463cee902ca50a0bf240 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 12 May 2022 16:35:55 -0700 Subject: [PATCH 01/45] Initial conversion, passes compilation and test --- .../detail/linestring_distance.cuh | 346 ++++++++++++++++++ .../experimental/linestring_distance.cuh | 41 +++ cpp/src/spatial/linestring_distance.cu | 242 ++---------- 3 files changed, 411 insertions(+), 218 deletions(-) create mode 100644 cpp/include/cuspatial/experimental/detail/linestring_distance.cuh create mode 100644 cpp/include/cuspatial/experimental/linestring_distance.cuh diff --git a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh new file mode 100644 index 000000000..9a1f55f2a --- /dev/null +++ b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh @@ -0,0 +1,346 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include +#include + +#include + +#include +#include +#include + +#include +#include +#include + +#include +#include +#include + +namespace cuspatial { +namespace detail { + +template +constexpr bool is_same() +{ + return std::conjunction_v...>; +} + +template +constexpr bool is_floating_point() +{ + return std::conjunction_v...>; +} + +template +double __device__ point_to_segment_distance(vec_2d const& C, + vec_2d const& A, + vec_2d const& B) +{ + // Subject 1.02 of https://www.inf.pucrs.br/~pinho/CG/faq.html + // Project the point to the segment, if it lands on the segment, + // the distance is the length of proejction, otherwise it's the + // length to one of the end points. + + double L_squared = (A.x - B.x) * (A.x - B.x) + (A.y - B.y) * (A.y - B.y); + if (L_squared == 0) { return hypot(C.x - A.x, C.y - A.y); } + double r = ((C.x - A.x) * (B.x - A.x) + (C.y - A.y) * (B.y - A.y)) / L_squared; + if (r <= 0 or r >= 1) { + return std::min(hypot(C.x - A.x, C.y - A.y), hypot(C.x - B.x, C.y - B.y)); + } + double Px = A.x + r * (B.x - A.x); + double Py = A.y + r * (B.y - A.y); + return hypot(C.x - Px, C.y - Py); +} + +template +double __device__ segment_distance_no_intersect(vec_2d const& A, + vec_2d const& B, + vec_2d const& C, + vec_2d const& D) +{ + return std::min(std::min(point_to_segment_distance(A, C, D), point_to_segment_distance(B, C, D)), + std::min(point_to_segment_distance(C, A, B), point_to_segment_distance(D, A, B))); +} + +/** + * @brief Computes shortest distance between two segments. + * + * If two segment intersects, distance is 0. + */ +template +double __device__ +segment_distance(vec_2d const& A, vec_2d const& B, vec_2d const& C, vec_2d const& D) +{ + // Subject 1.03 of https://www.inf.pucrs.br/~pinho/CG/faq.html + // Construct a parametrized ray of AB and CD, solve for the parameters. + // If both parameters are within [0, 1], the intersection exists. + + double r_denom = (B.x - A.x) * (D.y - C.y) - (B.y - A.y) * (D.x - C.x); + double r_numer = (A.y - C.y) * (D.x - C.x) - (A.x - C.x) * (D.y - C.y); + if (r_denom == 0) { + if (r_numer == 0) { return 0.0; } // Segments coincides + // Segments parallel + return segment_distance_no_intersect(A, B, C, D); + } + double r = r_numer / r_denom; + double s = ((A.y - C.y) * (B.x - A.x) - (A.x - C.x) * (B.y - A.x)) / + ((B.x - A.x) * (D.y - C.y) - (B.y - A.y) * (D.x - C.x)); + if (r >= 0 and r <= 1 and s >= 0 and s <= 1) { return 0.0; } + return segment_distance_no_intersect(A, B, C, D); +} + +/** @brief Get the index that is one-past the end point of linestring at @p linestring_idx + * + * @note The last endpoint of the linestring is not included in the offset array, thus + * @p num_points is returned. + */ +template +inline SizeType __device__ +endpoint_index_of_linestring(SizeType const& linestring_idx, + OffsetIterator const& linestring_offsets_begin, + SizeType const& num_linestrings, + SizeType const& num_points) +{ + return (linestring_idx == (num_linestrings - 1) + ? (num_points) + : *(linestring_offsets_begin + linestring_idx + 1)) - + 1; +} + +/** + * @brief Computes shortest distance between @p c and segment ab + */ +template +T __device__ point_to_segment_distance_squared(vec_2d const& c, + vec_2d const& a, + vec_2d const& b) +{ + auto ab = b - a; + auto ac = c - a; + auto l_squared = dot(ab, ab); + if (l_squared == 0) { return dot(ac, ac); } + auto r = dot(ac, ab); + auto bc = c - b; + // If the projection of `c` is outside of segment `ab`, compute point-point distance. + if (r <= 0 or r >= l_squared) { return std::min(dot(ac, ac), dot(bc, bc)); } + auto p = a + (r / l_squared) * ab; + auto pc = c - p; + return dot(pc, pc); +} + +/** + * @brief Computes shortest distance between two segments (ab and cd) that + * doesn't intersect. + */ +template +T __device__ segment_distance_no_intersect_or_colinear(vec_2d const& a, + vec_2d const& b, + vec_2d const& c, + vec_2d const& d) +{ + auto dist_sqr = std::min(std::min(point_to_segment_distance_squared(a, c, d), + point_to_segment_distance_squared(b, c, d)), + std::min(point_to_segment_distance_squared(c, a, b), + point_to_segment_distance_squared(d, a, b))); + return dist_sqr; +} + +/** + * @brief Computes shortest distance between two segments. + * + * If two segments intersect, the distance is 0. Otherwise compute the shortest point + * to segment distance. + */ +template +T __device__ squared_segment_distance(vec_2d const& a, + vec_2d const& b, + vec_2d const& c, + vec_2d const& d) +{ + auto ab = b - a; + auto cd = d - c; + auto denom = det(ab, cd); + + if (denom == 0) { + // Segments parallel or collinear + return segment_distance_no_intersect_or_colinear(a, b, c, d); + } + + auto ac = c - a; + auto r_numer = det(ac, cd); + auto denom_reciprocal = 1 / denom; + auto r = r_numer * denom_reciprocal; + auto s = det(ac, ab) * denom_reciprocal; + if (r >= 0 and r <= 1 and s >= 0 and s <= 1) { return 0.0; } + return segment_distance_no_intersect_or_colinear(a, b, c, d); +} + +/** + * @brief The kernel to compute point to linestring distance + * + * Each thread of the kernel computes the distance between a segment in a linestring in pair 1 + * to a linestring in pair 2. For a segment in pair 1, the linestring index is looked up from + * the offset array and mapped to the linestring in the pair 2. The segment is then computed + * with all segments in the corresponding linestringin pair 2. This forms a local minima of the + * shortest distance, which is then combined with other segment results via an atomic operation + * to form the globally minimum distance between the linestrings. + * + * @tparam Cart2dItA Iterator to 2d cartesian coordinates. Must meet requirements of + * [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. + * @tparam Cart2dItB Iterator to 2d cartesian coordinates. Must meet requirements of + * [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. + * @tparam OffsetIterator Iterator to linestring offsets. Must meet requirements of + * [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. + * @tparam OutputIterator Iterator to output distances. Must meet requirements of + * [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. + * + * @param[in] linestring1_offsets_begin Iterator to the begin of the range of linestring offsets + * in pair 1. + * @param[in] linestring1_offsets_end Iterator to the end of the range of linestring offsets + * in pair 1. + * @param[in] linestring1_points_xs_begin Iterator to the begin of the range of x coordinates of + * points in pair 1. + * @param[in] linestring1_points_xs_end Iterator to the end of the range of x coordiantes of points + * in pair 1. + * @param[in] linestring2_offsets_begin Iterator to the begin of the range of linestring offsets + * in pair 2. + * @param[in] linestring2_points_xs_begin Iterator to the begin of the range of x coordinates of + * points in pair 2. + * @param[in] linestring2_points_xs_end Iterator to the end of the range of x coordiantes of points + * in pair 2. + * @param[out] distances Iterator to the output range of shortest distances between pairs. + * + * [LinkLRAI]: https://en.cppreference.com/w/cpp/named_req/RandomAccessIterator + * "LegacyRandomAccessIterator" + */ +template +void __global__ pairwise_linestring_distance_kernel(OffsetIterator linestring1_offsets_begin, + OffsetIterator linestring1_offsets_end, + Cart2dItA linestring1_points_begin, + Cart2dItA linestring1_points_end, + OffsetIterator linestring2_offsets_begin, + Cart2dItB linestring2_points_begin, + Cart2dItB linestring2_points_end, + OutputIterator distances) +{ + using T = typename std::iterator_traits::value_type::value_type; + + auto const p1_idx = threadIdx.x + blockIdx.x * blockDim.x; + auto const num_linestrings = thrust::distance(linestring1_offsets_begin, linestring1_offsets_end); + auto const linestring1_num_points = + thrust::distance(linestring1_points_begin, linestring1_points_end); + auto const linestring2_num_points = + thrust::distance(linestring2_points_begin, linestring2_points_end); + + if (p1_idx >= linestring1_num_points) { return; } + + auto const linestring_idx = + thrust::distance(linestring1_offsets_begin, + thrust::upper_bound( + thrust::seq, linestring1_offsets_begin, linestring1_offsets_end, p1_idx)) - + 1; + + auto ls1_end = endpoint_index_of_linestring( + linestring_idx, linestring1_offsets_begin, num_linestrings, linestring1_num_points); + + if (p1_idx == ls1_end) { + // Current point is the end point of the line string. + return; + } + + auto ls2_start = *(linestring2_offsets_begin + linestring_idx); + auto ls2_end = endpoint_index_of_linestring( + linestring_idx, linestring2_offsets_begin, num_linestrings, linestring2_num_points); + + auto const& A = linestring1_points_begin[p1_idx]; + auto const& B = linestring1_points_begin[p1_idx + 1]; + + auto min_squared_distance = std::numeric_limits::max(); + for (auto p2_idx = ls2_start; p2_idx < ls2_end; p2_idx++) { + auto C = linestring2_points_begin[p2_idx]; + auto D = linestring2_points_begin[p2_idx + 1]; + min_squared_distance = std::min(min_squared_distance, squared_segment_distance(A, B, C, D)); + } + atomicMin(distances + linestring_idx, static_cast(std::sqrt(min_squared_distance))); +} + +} // namespace detail + +template +void pairwise_linestring_distance(OffsetIterator linestring1_offsets_first, + OffsetIterator linestring1_offsets_last, + Cart2dItA linestring1_points_first, + Cart2dItA linestring1_points_last, + OffsetIterator linestring2_offsets_first, + Cart2dItB linestring2_points_first, + Cart2dItB linestring2_points_last, + OutputIt distances_first, + rmm::cuda_stream_view stream) +{ + using T = typename Cart2dA::value_type; + + static_assert(detail::is_floating_point::value_type>(), + "Inputs and output must be floating point types."); + + static_assert(detail::is_same::value_type>(), + "Inputs and output must be the same types."); + + static_assert(detail::is_same, Cart2dA, Cart2dB>(), + "Inputs must be cuspatial::cartesian_2d"); + + auto const num_string_pairs = + thrust::distance(linestring1_offsets_first, linestring1_offsets_last); + auto const num_linestring1_points = + thrust::distance(linestring1_points_first, linestring1_points_last); + auto const num_linestring2_points = + thrust::distance(linestring2_points_first, linestring2_points_last); + + thrust::fill(rmm::exec_policy(stream), + distances_first, + distances_first + num_string_pairs, + std::numeric_limits::max()); + + std::size_t constexpr threads_per_block = 64; + std::size_t const num_blocks = + (num_linestring1_points + threads_per_block - 1) / threads_per_block; + + detail::pairwise_linestring_distance_kernel<<>>( + linestring1_offsets_first, + linestring1_offsets_last, + linestring1_points_first, + linestring1_points_last, + linestring2_offsets_first, + linestring2_points_first, + linestring2_points_last, + distances_first); + + CUSPATIAL_CUDA_TRY(cudaGetLastError()); +} + +} // namespace cuspatial diff --git a/cpp/include/cuspatial/experimental/linestring_distance.cuh b/cpp/include/cuspatial/experimental/linestring_distance.cuh new file mode 100644 index 000000000..bb9f14ef5 --- /dev/null +++ b/cpp/include/cuspatial/experimental/linestring_distance.cuh @@ -0,0 +1,41 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include + +namespace cuspatial { + +template ::value_type, + class Cart2dB = typename std::iterator_traits::value_type> +void pairwise_linestring_distance(OffsetIterator linestring1_offsets_first, + OffsetIterator linestring1_offsets_last, + Cart2dItA linestring1_points_first, + Cart2dItA linestring1_points_last, + OffsetIterator linestring2_offsets_first, + Cart2dItB linestring2_points_first, + Cart2dItB linestring2_points_last, + OutputIt distances_first, + rmm::cuda_stream_view stream); + +} + +#include diff --git a/cpp/src/spatial/linestring_distance.cu b/cpp/src/spatial/linestring_distance.cu index 036086ea2..66a7cd81b 100644 --- a/cpp/src/spatial/linestring_distance.cu +++ b/cpp/src/spatial/linestring_distance.cu @@ -15,7 +15,8 @@ */ #include -#include +#include +#include #include #include @@ -37,195 +38,7 @@ #include namespace cuspatial { -namespace { - -/** @brief Get the index that is one-past the end point of linestring at @p linestring_idx - * - * @note The last endpoint of the linestring is not included in the offset array, thus - * @p num_points is returned. - */ -template -inline SizeType __device__ -endpoint_index_of_linestring(SizeType const& linestring_idx, - OffsetIterator const& linestring_offsets_begin, - SizeType const& num_linestrings, - SizeType const& num_points) -{ - return (linestring_idx == (num_linestrings - 1) - ? (num_points) - : *(linestring_offsets_begin + linestring_idx + 1)) - - 1; -} - -/** - * @brief Computes shortest distance between @p c and segment ab - */ -template -T __device__ point_to_segment_distance_squared(vec_2d const& c, - vec_2d const& a, - vec_2d const& b) -{ - auto ab = b - a; - auto ac = c - a; - auto l_squared = dot(ab, ab); - if (l_squared == 0) { return dot(ac, ac); } - auto r = dot(ac, ab); - auto bc = c - b; - // If the projection of `c` is outside of segment `ab`, compute point-point distance. - if (r <= 0 or r >= l_squared) { return std::min(dot(ac, ac), dot(bc, bc)); } - auto p = a + (r / l_squared) * ab; - auto pc = c - p; - return dot(pc, pc); -} - -/** - * @brief Computes shortest distance between two segments (ab and cd) that - * doesn't intersect. - */ -template -T __device__ segment_distance_no_intersect_or_colinear(vec_2d const& a, - vec_2d const& b, - vec_2d const& c, - vec_2d const& d) -{ - auto dist_sqr = std::min(std::min(point_to_segment_distance_squared(a, c, d), - point_to_segment_distance_squared(b, c, d)), - std::min(point_to_segment_distance_squared(c, a, b), - point_to_segment_distance_squared(d, a, b))); - return dist_sqr; -} - -/** - * @brief Computes shortest distance between two segments. - * - * If two segments intersect, the distance is 0. Otherwise compute the shortest point - * to segment distance. - */ -template -T __device__ squared_segment_distance(vec_2d const& a, - vec_2d const& b, - vec_2d const& c, - vec_2d const& d) -{ - auto ab = b - a; - auto cd = d - c; - auto denom = det(ab, cd); - - if (denom == 0) { - // Segments parallel or collinear - return segment_distance_no_intersect_or_colinear(a, b, c, d); - } - - auto ac = c - a; - auto r_numer = det(ac, cd); - auto denom_reciprocal = 1 / denom; - auto r = r_numer * denom_reciprocal; - auto s = det(ac, ab) * denom_reciprocal; - if (r >= 0 and r <= 1 and s >= 0 and s <= 1) { return 0.0; } - return segment_distance_no_intersect_or_colinear(a, b, c, d); -} - -/** - * @brief The kernel to compute point to linestring distance - * - * Each thread of the kernel computes the distance between a segment in a linestring in pair 1 - * to a linestring in pair 2. For a segment in pair 1, the linestring index is looked up from - * the offset array and mapped to the linestring in the pair 2. The segment is then computed - * with all segments in the corresponding linestringin pair 2. This forms a local minima of the - * shortest distance, which is then combined with other segment results via an atomic operation - * to form the globally minimum distance between the linestrings. - * - * @tparam CoordinateIterator Iterator to coordinates. Must meet requirements of - * [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. - * @tparam OffsetIterator Iterator to linestring offsets. Must meet requirements of - * [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. - * @tparam OutputIterator Iterator to output distances. Must meet requirements of - * [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. - * - * @param[in] linestring1_offsets_begin Iterator to the begin of the range of linestring offsets - * in pair 1. - * @param[in] linestring1_offsets_end Iterator to the end of the range of linestring offsets - * in pair 1. - * @param[in] linestring1_points_xs_begin Iterator to the begin of the range of x coordinates of - * points in pair 1. - * @param[in] linestring1_points_xs_end Iterator to the end of the range of x coordiantes of points - * in pair 1. - * @param[in] linestring1_points_ys_begin Iterator to the begin of the range of y coordinates of - * points in pair 1. - * @param[in] linestring2_offsets_begin Iterator to the begin of the range of linestring offsets - * in pair 2. - * @param[in] linestring2_points_xs_begin Iterator to the begin of the range of x coordinates of - * points in pair 2. - * @param[in] linestring2_points_xs_end Iterator to the end of the range of x coordiantes of points - * in pair 2. - * @param[in] linestring2_points_ys_begin Iterator to the begin of the range of y coordinates of - * points in pair 2. - * @param[out] distances Iterator to the output range of shortest distances between pairs. - * - * [LinkLRAI]: https://en.cppreference.com/w/cpp/named_req/RandomAccessIterator - * "LegacyRandomAccessIterator" - */ -template -void __global__ pairwise_linestring_distance_kernel(OffsetIterator linestring1_offsets_begin, - OffsetIterator linestring1_offsets_end, - CoordinateIterator linestring1_points_xs_begin, - CoordinateIterator linestring1_points_xs_end, - CoordinateIterator linestring1_points_ys_begin, - OffsetIterator linestring2_offsets_begin, - CoordinateIterator linestring2_points_xs_begin, - CoordinateIterator linestring2_points_xs_end, - CoordinateIterator linestring2_points_ys_begin, - OutputIterator distances) -{ - using T = typename std::iterator_traits::value_type; - - auto const p1_idx = threadIdx.x + blockIdx.x * blockDim.x; - cudf::size_type const num_linestrings = - thrust::distance(linestring1_offsets_begin, linestring1_offsets_end); - cudf::size_type const linestring1_num_points = - thrust::distance(linestring1_points_xs_begin, linestring1_points_xs_end); - cudf::size_type const linestring2_num_points = - thrust::distance(linestring2_points_xs_begin, linestring2_points_xs_end); - - if (p1_idx >= linestring1_num_points) { return; } - - cudf::size_type const linestring_idx = - thrust::distance(linestring1_offsets_begin, - thrust::upper_bound( - thrust::seq, linestring1_offsets_begin, linestring1_offsets_end, p1_idx)) - - 1; - - cudf::size_type ls1_end = endpoint_index_of_linestring( - linestring_idx, linestring1_offsets_begin, num_linestrings, linestring1_num_points); - - if (p1_idx == ls1_end) { - // Current point is the end point of the line string. - return; - } - - cudf::size_type ls2_start = *(linestring2_offsets_begin + linestring_idx); - cudf::size_type ls2_end = endpoint_index_of_linestring( - linestring_idx, linestring2_offsets_begin, num_linestrings, linestring2_num_points); - - vec_2d A{linestring1_points_xs_begin[p1_idx], linestring1_points_ys_begin[p1_idx]}; - vec_2d B{linestring1_points_xs_begin[p1_idx + 1], linestring1_points_ys_begin[p1_idx + 1]}; - - T min_squared_distance = std::numeric_limits::max(); - for (cudf::size_type p2_idx = ls2_start; p2_idx < ls2_end; p2_idx++) { - vec_2d C{linestring2_points_xs_begin[p2_idx], linestring2_points_ys_begin[p2_idx]}; - vec_2d D{linestring2_points_xs_begin[p2_idx + 1], linestring2_points_ys_begin[p2_idx + 1]}; - min_squared_distance = std::min(min_squared_distance, squared_segment_distance(A, B, C, D)); - } - atomicMin(distances + linestring_idx, static_cast(std::sqrt(min_squared_distance))); -} - -} // anonymous namespace - namespace detail { - -/** - * @brief Functor that launches the kernel to compute pairwise linestring distances. - */ struct pairwise_linestring_distance_functor { template std::enable_if_t::value, std::unique_ptr> operator()( @@ -238,35 +51,28 @@ struct pairwise_linestring_distance_functor { rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - using namespace cudf; - - auto const num_string_pairs = static_cast(linestring1_offsets.size()); - - auto distances = - make_numeric_column(data_type{type_to_id()}, num_string_pairs, mask_state::UNALLOCATED); - - thrust::fill(rmm::exec_policy(stream), - distances->mutable_view().begin(), - distances->mutable_view().end(), - std::numeric_limits::max()); - - std::size_t constexpr threads_per_block = 64; - std::size_t const num_blocks = - (linestring1_points_x.size() + threads_per_block - 1) / threads_per_block; - - pairwise_linestring_distance_kernel<<>>( - linestring1_offsets.begin(), - linestring1_offsets.end(), - linestring1_points_x.begin(), - linestring1_points_x.end(), - linestring1_points_y.begin(), - linestring2_offsets.begin(), - linestring2_points_x.begin(), - linestring2_points_x.end(), - linestring2_points_y.begin(), - distances->mutable_view().begin()); - - CUSPATIAL_CUDA_TRY(cudaGetLastError()); + auto const num_string_pairs = static_cast(linestring1_offsets.size()); + + auto distances = cudf::make_numeric_column(cudf::data_type{cudf::type_to_id()}, + num_string_pairs, + cudf::mask_state::UNALLOCATED, + stream, + mr); + + auto linestring1_coords_it = + make_cartesian_2d_iterator(linestring1_points_x.begin(), linestring1_points_y.begin()); + auto linestring2_coords_it = + make_cartesian_2d_iterator(linestring2_points_x.begin(), linestring2_points_y.begin()); + + pairwise_linestring_distance(linestring1_offsets.begin(), + linestring1_offsets.end(), + linestring1_coords_it, + linestring1_coords_it + linestring1_points_x.size(), + linestring2_offsets.begin(), + linestring2_coords_it, + linestring2_coords_it + linestring2_points_x.size(), + distances->mutable_view().begin(), + stream); return distances; } From 1861dc986da944cfc60014568ac3bbe1f3ebdc1c Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 12 May 2022 17:16:06 -0700 Subject: [PATCH 02/45] Add docstring --- .../experimental/linestring_distance.cuh | 34 ++++++++++++++++++- 1 file changed, 33 insertions(+), 1 deletion(-) diff --git a/cpp/include/cuspatial/experimental/linestring_distance.cuh b/cpp/include/cuspatial/experimental/linestring_distance.cuh index bb9f14ef5..fde989436 100644 --- a/cpp/include/cuspatial/experimental/linestring_distance.cuh +++ b/cpp/include/cuspatial/experimental/linestring_distance.cuh @@ -20,6 +20,38 @@ namespace cuspatial { +/** + * @copybrief cuspatial::pairwise_linestring_distance + * + * The shortest distance between two linestrings is defined as the shortest distance + * between all pairs of segments of the two linestrings. If any of the segments intersect, + * the distance is 0. + * + * @tparam Cart2dItA iterator type for point array of the first linestring of each pair + * @tparam Cart2dItB iterator type for point array of the second linestring of each pair + * @tparam OffsetIterator iterator type for offset array + * @tparam OutputIt iterator type for output array + * @tparam std::iterator_traits::value_type value type of `Cart2dItA`, must be + * `cuspatial::cartesian_2d` + * @tparam std::iterator_traits::value_type value type of `Cart2dItB`, must be + * `cuspatial::cartesian_2d` + * + * @param linestring1_offsets_first begin of range of the offsets to the first linestring of each + * pair + * @param linestring1_offsets_last end of range of the offsets to the first linestring of each pair + * @param linestring1_points_first begin of range of the point of the first linestring of each pair + * @param linestring1_points_last end of range of the point of the first linestring of each pair + * @param linestring2_offsets_first begin of range of the offsets to the second linestring of each + * pair + * @param linestring2_points_first begin of range of the point of the second linestring of each pair + * @param linestring2_points_last end of range of the point of the second linestring of each pair + * @param distances_first begin to output array + * @param stream Used for device memory operations and kernel launches. + * + * @pre all input iterators for coordinates must have `cuspatial::cartesian_2d` type. + * @pre all scalar types must be floating point types, and must be the same type for all input + * iterators and output iterators. + */ template From 0d5dc9408f94081eefd8c4ba0c0978419e0ab9a6 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 12 May 2022 17:20:56 -0700 Subject: [PATCH 03/45] Add RAI specification --- .../experimental/linestring_distance.cuh | 15 +++++++++++---- 1 file changed, 11 insertions(+), 4 deletions(-) diff --git a/cpp/include/cuspatial/experimental/linestring_distance.cuh b/cpp/include/cuspatial/experimental/linestring_distance.cuh index fde989436..6c01ea88e 100644 --- a/cpp/include/cuspatial/experimental/linestring_distance.cuh +++ b/cpp/include/cuspatial/experimental/linestring_distance.cuh @@ -27,10 +27,14 @@ namespace cuspatial { * between all pairs of segments of the two linestrings. If any of the segments intersect, * the distance is 0. * - * @tparam Cart2dItA iterator type for point array of the first linestring of each pair - * @tparam Cart2dItB iterator type for point array of the second linestring of each pair - * @tparam OffsetIterator iterator type for offset array - * @tparam OutputIt iterator type for output array + * @tparam Cart2dItA iterator type for point array of the first linestring of each pair. Must meet + * the requirements of [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. + * @tparam Cart2dItB iterator type for point array of the second linestring of each pair. Must meet + * the requirements of [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. + * @tparam OffsetIterator iterator type for offset array. Must meet the requirements of + * [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. + * @tparam OutputIt iterator type for output array. Must meet the requirements of + * [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. * @tparam std::iterator_traits::value_type value type of `Cart2dItA`, must be * `cuspatial::cartesian_2d` * @tparam std::iterator_traits::value_type value type of `Cart2dItB`, must be @@ -51,6 +55,9 @@ namespace cuspatial { * @pre all input iterators for coordinates must have `cuspatial::cartesian_2d` type. * @pre all scalar types must be floating point types, and must be the same type for all input * iterators and output iterators. + * + * [LinkLRAI]: https://en.cppreference.com/w/cpp/named_req/RandomAccessIterator + * "LegacyRandomAccessIterator" */ template Date: Thu, 12 May 2022 17:33:57 -0700 Subject: [PATCH 04/45] add default stream parameter --- cpp/include/cuspatial/experimental/linestring_distance.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/include/cuspatial/experimental/linestring_distance.cuh b/cpp/include/cuspatial/experimental/linestring_distance.cuh index 6c01ea88e..0fbd9ef2c 100644 --- a/cpp/include/cuspatial/experimental/linestring_distance.cuh +++ b/cpp/include/cuspatial/experimental/linestring_distance.cuh @@ -73,7 +73,7 @@ void pairwise_linestring_distance(OffsetIterator linestring1_offsets_first, Cart2dItB linestring2_points_first, Cart2dItB linestring2_points_last, OutputIt distances_first, - rmm::cuda_stream_view stream); + rmm::cuda_stream_view stream = rmm::cuda_stream_default); } // namespace cuspatial From a427af6eabceaceacd7b2e87cf95eb1d29550c23 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 12 May 2022 20:27:07 -0700 Subject: [PATCH 05/45] Add first test and cast references around. --- .../detail/linestring_distance.cuh | 20 +++--- cpp/tests/CMakeLists.txt | 3 + .../spatial/linestring_distance_test.cu | 70 +++++++++++++++++++ 3 files changed, 83 insertions(+), 10 deletions(-) create mode 100644 cpp/tests/experimental/spatial/linestring_distance_test.cu diff --git a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh index 9a1f55f2a..3c93bcdb6 100644 --- a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh +++ b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh @@ -168,11 +168,9 @@ T __device__ segment_distance_no_intersect_or_colinear(vec_2d const& a, * If two segments intersect, the distance is 0. Otherwise compute the shortest point * to segment distance. */ -template -T __device__ squared_segment_distance(vec_2d const& a, - vec_2d const& b, - vec_2d const& c, - vec_2d const& d) +template +T __device__ +squared_segment_distance(Vec2d const& a, Vec2d const& b, Vec2d const& c, Vec2d const& d) { auto ab = b - a; auto cd = d - c; @@ -269,16 +267,18 @@ void __global__ pairwise_linestring_distance_kernel(OffsetIterator linestring1_o auto ls2_end = endpoint_index_of_linestring( linestring_idx, linestring2_offsets_begin, num_linestrings, linestring2_num_points); - auto const& A = linestring1_points_begin[p1_idx]; - auto const& B = linestring1_points_begin[p1_idx + 1]; + auto const& A = thrust::raw_reference_cast(linestring1_points_begin[p1_idx]); + auto const& B = thrust::raw_reference_cast(linestring1_points_begin[p1_idx + 1]); auto min_squared_distance = std::numeric_limits::max(); for (auto p2_idx = ls2_start; p2_idx < ls2_end; p2_idx++) { - auto C = linestring2_points_begin[p2_idx]; - auto D = linestring2_points_begin[p2_idx + 1]; + auto const& C = thrust::raw_reference_cast(linestring2_points_begin[p2_idx]); + auto const& D = thrust::raw_reference_cast(linestring2_points_begin[p2_idx + 1]); min_squared_distance = std::min(min_squared_distance, squared_segment_distance(A, B, C, D)); } - atomicMin(distances + linestring_idx, static_cast(std::sqrt(min_squared_distance))); + + atomicMin(&thrust::raw_reference_cast(*(distances + linestring_idx)), + static_cast(std::sqrt(min_squared_distance))); } } // namespace detail diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index b6857003c..a9560fb5e 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -104,3 +104,6 @@ ConfigureTest(SPATIAL_WINDOW_POINT_TEST # Experimental API ConfigureTest(HAVERSINE_TEST_EXP experimental/spatial/haversine_test.cu) + +ConfigureTest(LINESTRING_DISTANCE_EXP + experimental/spatial/linestring_distance_test.cu) diff --git a/cpp/tests/experimental/spatial/linestring_distance_test.cu b/cpp/tests/experimental/spatial/linestring_distance_test.cu new file mode 100644 index 000000000..c3d4be133 --- /dev/null +++ b/cpp/tests/experimental/spatial/linestring_distance_test.cu @@ -0,0 +1,70 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include +#include + +#include +#include + +#include +#include +#include +#include + +namespace cuspatial { +namespace test { + +using namespace cudf; +using namespace cudf::test; + +template +struct PairwiseLinestringDistanceTest : public BaseFixture { +}; + +// float and double are logically the same but would require seperate tests due to precision. +using TestTypes = ::testing::Types; +TYPED_TEST_CASE(PairwiseLinestringDistanceTest, TestTypes); + +TYPED_TEST(PairwiseLinestringDistanceTest, FromSeparateArrayInputs) +{ + using T = TypeParam; + using CartVec = std::vector>; + + auto a_cart2d = rmm::device_vector>{ + CartVec({{0.0f, 0.0f}, {1.0f, 0.0f}, {2.0f, 0.0f}, {3.0f, 0.0f}, {4.0f, 0.0f}})}; + auto b_cart2d = rmm::device_vector>{ + CartVec({{0.0f, 1.0f}, {1.0f, 1.0f}, {2.0f, 1.0f}, {3.0f, 1.0f}, {4.0f, 1.0f}})}; + auto offset = rmm::device_vector{0}; + + auto distance = rmm::device_vector{std::vector{0.0}}; + auto expected = rmm::device_vector{std::vector{1.0}}; + + pairwise_linestring_distance(offset.begin(), + offset.end(), + a_cart2d.begin(), + a_cart2d.end(), + offset.begin(), + b_cart2d.begin(), + b_cart2d.end(), + distance.begin()); + + EXPECT_EQ(distance, expected); +} + +} // namespace test +} // namespace cuspatial From 2c6ec8533db4d4954f20530e5f12a0bc6f7ca1fe Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 12 May 2022 20:48:45 -0700 Subject: [PATCH 06/45] Add more tests --- .../spatial/linestring_distance_test.cu | 83 +++++++++++++++++++ 1 file changed, 83 insertions(+) diff --git a/cpp/tests/experimental/spatial/linestring_distance_test.cu b/cpp/tests/experimental/spatial/linestring_distance_test.cu index c3d4be133..fc6d7c5bf 100644 --- a/cpp/tests/experimental/spatial/linestring_distance_test.cu +++ b/cpp/tests/experimental/spatial/linestring_distance_test.cu @@ -16,6 +16,7 @@ #include #include +#include #include #include @@ -66,5 +67,87 @@ TYPED_TEST(PairwiseLinestringDistanceTest, FromSeparateArrayInputs) EXPECT_EQ(distance, expected); } +TYPED_TEST(PairwiseLinestringDistanceTest, FromSamePointArrayInput) +{ + using T = TypeParam; + using CartVec = std::vector>; + + auto cart2ds = rmm::device_vector>{ + CartVec({{0.0f, 0.0f}, {1.0f, 0.0f}, {2.0f, 0.0f}, {3.0f, 0.0f}, {4.0f, 0.0f}})}; + auto offset = rmm::device_vector{0}; + + auto a_begin = cart2ds.begin(); + auto a_end = cart2ds.begin() + 3; + auto b_begin = cart2ds.begin() + 1; + auto b_end = cart2ds.end(); + + auto distance = rmm::device_vector{std::vector{0.0}}; + auto expected = rmm::device_vector{std::vector{0.0}}; + + pairwise_linestring_distance( + offset.begin(), offset.end(), a_begin, a_end, offset.begin(), b_begin, b_end, distance.begin()); + + EXPECT_EQ(distance, expected); +} + +TYPED_TEST(PairwiseLinestringDistanceTest, FromTransformIterator) +{ + using T = TypeParam; + using CartVec = std::vector>; + + auto a_cart2d_x = rmm::device_vector{std::vector{0.0, 1.0, 2.0, 3.0, 4.0}}; + auto a_cart2d_y = rmm::device_vector(5, 0.0); + + auto a_begin = make_cartesian_2d_iterator(a_cart2d_x.begin(), a_cart2d_y.begin()); + auto a_end = a_begin + a_cart2d_x.size(); + + auto b_cart2d_x = rmm::device_vector{std::vector{0.0, 1.0, 2.0, 3.0, 4.0}}; + auto b_cart2d_y = rmm::device_vector(5, 1.0); + + auto b_begin = make_cartesian_2d_iterator(b_cart2d_x.begin(), b_cart2d_y.begin()); + auto b_end = b_begin + b_cart2d_x.size(); + + auto offset = rmm::device_vector{0}; + + auto distance = rmm::device_vector{std::vector{0.0}}; + auto expected = rmm::device_vector{std::vector{1.0}}; + + pairwise_linestring_distance( + offset.begin(), offset.end(), a_begin, a_end, offset.begin(), b_begin, b_end, distance.begin()); + + EXPECT_EQ(distance, expected); +} + +TYPED_TEST(PairwiseLinestringDistanceTest, FromMixedIterator) +{ + using T = TypeParam; + using CartVec = std::vector>; + + auto a_cart2d = rmm::device_vector>{ + CartVec({{0.0f, 0.0f}, {1.0f, 0.0f}, {2.0f, 0.0f}, {3.0f, 0.0f}, {4.0f, 0.0f}})}; + + auto b_cart2d_x = rmm::device_vector{std::vector{0.0, 1.0, 2.0, 3.0, 4.0}}; + auto b_cart2d_y = rmm::device_vector(5, 1.0); + + auto b_begin = make_cartesian_2d_iterator(b_cart2d_x.begin(), b_cart2d_y.begin()); + auto b_end = b_begin + b_cart2d_x.size(); + + auto offset = rmm::device_vector{0}; + + auto distance = rmm::device_vector{std::vector{0.0}}; + auto expected = rmm::device_vector{std::vector{1.0}}; + + pairwise_linestring_distance(offset.begin(), + offset.end(), + a_cart2d.begin(), + a_cart2d.end(), + offset.begin(), + b_begin, + b_end, + distance.begin()); + + EXPECT_EQ(distance, expected); +} + } // namespace test } // namespace cuspatial From 9958f4844e50e7227e43959fd359612e46f0425e Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Mon, 16 May 2022 14:19:41 -0700 Subject: [PATCH 07/45] fix offset arrays --- .../experimental/spatial/linestring_distance_test.cu | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/tests/experimental/spatial/linestring_distance_test.cu b/cpp/tests/experimental/spatial/linestring_distance_test.cu index fc6d7c5bf..cf05fedf6 100644 --- a/cpp/tests/experimental/spatial/linestring_distance_test.cu +++ b/cpp/tests/experimental/spatial/linestring_distance_test.cu @@ -50,7 +50,7 @@ TYPED_TEST(PairwiseLinestringDistanceTest, FromSeparateArrayInputs) CartVec({{0.0f, 0.0f}, {1.0f, 0.0f}, {2.0f, 0.0f}, {3.0f, 0.0f}, {4.0f, 0.0f}})}; auto b_cart2d = rmm::device_vector>{ CartVec({{0.0f, 1.0f}, {1.0f, 1.0f}, {2.0f, 1.0f}, {3.0f, 1.0f}, {4.0f, 1.0f}})}; - auto offset = rmm::device_vector{0}; + auto offset = rmm::device_vector{std::vector{std::vector{0}}}; auto distance = rmm::device_vector{std::vector{0.0}}; auto expected = rmm::device_vector{std::vector{1.0}}; @@ -74,7 +74,7 @@ TYPED_TEST(PairwiseLinestringDistanceTest, FromSamePointArrayInput) auto cart2ds = rmm::device_vector>{ CartVec({{0.0f, 0.0f}, {1.0f, 0.0f}, {2.0f, 0.0f}, {3.0f, 0.0f}, {4.0f, 0.0f}})}; - auto offset = rmm::device_vector{0}; + auto offset = rmm::device_vector{std::vector{0}}; auto a_begin = cart2ds.begin(); auto a_end = cart2ds.begin() + 3; @@ -107,7 +107,7 @@ TYPED_TEST(PairwiseLinestringDistanceTest, FromTransformIterator) auto b_begin = make_cartesian_2d_iterator(b_cart2d_x.begin(), b_cart2d_y.begin()); auto b_end = b_begin + b_cart2d_x.size(); - auto offset = rmm::device_vector{0}; + auto offset = rmm::device_vector{std::vector{0}}; auto distance = rmm::device_vector{std::vector{0.0}}; auto expected = rmm::device_vector{std::vector{1.0}}; @@ -132,7 +132,7 @@ TYPED_TEST(PairwiseLinestringDistanceTest, FromMixedIterator) auto b_begin = make_cartesian_2d_iterator(b_cart2d_x.begin(), b_cart2d_y.begin()); auto b_end = b_begin + b_cart2d_x.size(); - auto offset = rmm::device_vector{0}; + auto offset = rmm::device_vector{std::vector{0}}; auto distance = rmm::device_vector{std::vector{0.0}}; auto expected = rmm::device_vector{std::vector{1.0}}; From dc944c572e3ac2762b5f94f8e3fac99debef7c7f Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Mon, 16 May 2022 14:20:14 -0700 Subject: [PATCH 08/45] fix wrong gtest binary name --- cpp/tests/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index a9560fb5e..2dee99013 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -105,5 +105,5 @@ ConfigureTest(SPATIAL_WINDOW_POINT_TEST ConfigureTest(HAVERSINE_TEST_EXP experimental/spatial/haversine_test.cu) -ConfigureTest(LINESTRING_DISTANCE_EXP +ConfigureTest(LINESTRING_DISTANCE_TEST_EXP experimental/spatial/linestring_distance_test.cu) From e36186a9b17054756bc2fb92204ea30f2b2029c8 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Mon, 16 May 2022 14:36:41 -0700 Subject: [PATCH 09/45] Add precommit hooks and script for cmake format/lint --- .pre-commit-config.yaml | 18 +++++++++ cpp/scripts/run-cmake-format.sh | 72 +++++++++++++++++++++++++++++++++ 2 files changed, 90 insertions(+) create mode 100755 cpp/scripts/run-cmake-format.sh diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index 70b4bde47..e988539c3 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -40,5 +40,23 @@ repos: language: system files: \.(cu|cuh|h|hpp|cpp|inl)$ args: ['-fallback-style=none'] + - id: cmake-format + name: cmake-format + entry: ./cpp/scripts/run-cmake-format.sh cmake-format + language: python + types: [cmake] + # Note that pre-commit autoupdate does not update the versions + # of dependencies, so we'll have to update this manually. + additional_dependencies: + - cmakelang==0.6.13 + - id: cmake-lint + name: cmake-lint + entry: ./cpp/scripts/run-cmake-format.sh cmake-lint + language: python + types: [cmake] + # Note that pre-commit autoupdate does not update the versions + # of dependencies, so we'll have to update this manually. + additional_dependencies: + - cmakelang==0.6.13 default_language_version: python: python3 diff --git a/cpp/scripts/run-cmake-format.sh b/cpp/scripts/run-cmake-format.sh new file mode 100755 index 000000000..5c47704e3 --- /dev/null +++ b/cpp/scripts/run-cmake-format.sh @@ -0,0 +1,72 @@ +#!/bin/bash + +# This script is a wrapper for cmakelang that may be used with pre-commit. The +# wrapping is necessary because RAPIDS libraries split configuration for +# cmakelang linters between a local config file and a second config file that's +# shared across all of RAPIDS via rapids-cmake. In order to keep it up to date +# this file is only maintained in one place (the rapids-cmake repo) and +# pulled down during builds. We need a way to invoke CMake linting commands +# without causing pre-commit failures (which could block local commits or CI), +# while also being sufficiently flexible to allow users to maintain the config +# file independently of a build directory. +# +# This script provides the minimal functionality to enable those use cases. It +# searches in a number of predefined locations for the rapids-cmake config file +# and exits gracefully if the file is not found. If a user wishes to specify a +# config file at a nonstandard location, they may do so by setting the +# environment variable RAPIDS_CMAKE_FORMAT_FILE. +# +# This script can be invoked directly anywhere within the project repository. +# Alternatively, it may be invoked as a pre-commit hook via +# `pre-commit run (cmake-format)|(cmake-lint)`. +# +# Usage: +# bash run-cmake-format.sh {cmake-format,cmake-lint} infile [infile ...] + +status=0 +if [ -z ${CUSPATIAL_ROOT:+PLACEHOLDER} ]; then + CUSPATIAL_BUILD_DIR=$(git rev-parse --show-toplevel 2>&1)/cpp/build + status=$? +else + CUSPATIAL_BUILD_DIR=${CUSPATIAL_ROOT} +fi + +if ! [ ${status} -eq 0 ]; then + if [[ ${CUSPATIAL_BUILD_DIR} == *"not a git repository"* ]]; then + echo "This script must be run inside the cuspatial repository, or the CUSPATIAL_ROOT environment variable must be set." + else + echo "Script failed with unknown error attempting to determine project root:" + echo ${CUSPATIAL_BUILD_DIR} + fi + exit 1 +fi + +DEFAULT_FORMAT_FILE_LOCATIONS=( + "${CUSPATIAL_BUILD_DIR:-${HOME}}/_deps/rapids-cmake-src/cmake-format-rapids-cmake.json" +) + +if [ -z ${RAPIDS_CMAKE_FORMAT_FILE:+PLACEHOLDER} ]; then + for file_path in ${DEFAULT_FORMAT_FILE_LOCATIONS[@]}; do + if [ -f ${file_path} ]; then + RAPIDS_CMAKE_FORMAT_FILE=${file_path} + break + fi + done +fi + +if [ -z ${RAPIDS_CMAKE_FORMAT_FILE:+PLACEHOLDER} ]; then + echo "The rapids-cmake cmake-format configuration file was not found at any of the default search locations: " + echo "" + ( IFS=$'\n'; echo "${DEFAULT_FORMAT_FILE_LOCATIONS[*]}" ) + echo "" + echo "Try setting the environment variable RAPIDS_CMAKE_FORMAT_FILE to the path to the config file." + exit 0 +else + echo "Using format file ${RAPIDS_CMAKE_FORMAT_FILE}" +fi + +if [[ $1 == "cmake-format" ]]; then + cmake-format -i --config-files cpp/cmake/config.json ${RAPIDS_CMAKE_FORMAT_FILE} -- ${@:2} +elif [[ $1 == "cmake-lint" ]]; then + cmake-lint --config-files cpp/cmake/config.json ${RAPIDS_CMAKE_FORMAT_FILE} -- ${@:2} +fi From 62aa1ad4ed4444d1cd31f076f4b2c6a325d1618c Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Tue, 17 May 2022 15:08:45 -0700 Subject: [PATCH 10/45] update with optimized code --- .../detail/linestring_distance.cuh | 66 ++----------------- 1 file changed, 5 insertions(+), 61 deletions(-) diff --git a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh index 3c93bcdb6..26b1b4a46 100644 --- a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh +++ b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh @@ -48,64 +48,6 @@ constexpr bool is_floating_point() return std::conjunction_v...>; } -template -double __device__ point_to_segment_distance(vec_2d const& C, - vec_2d const& A, - vec_2d const& B) -{ - // Subject 1.02 of https://www.inf.pucrs.br/~pinho/CG/faq.html - // Project the point to the segment, if it lands on the segment, - // the distance is the length of proejction, otherwise it's the - // length to one of the end points. - - double L_squared = (A.x - B.x) * (A.x - B.x) + (A.y - B.y) * (A.y - B.y); - if (L_squared == 0) { return hypot(C.x - A.x, C.y - A.y); } - double r = ((C.x - A.x) * (B.x - A.x) + (C.y - A.y) * (B.y - A.y)) / L_squared; - if (r <= 0 or r >= 1) { - return std::min(hypot(C.x - A.x, C.y - A.y), hypot(C.x - B.x, C.y - B.y)); - } - double Px = A.x + r * (B.x - A.x); - double Py = A.y + r * (B.y - A.y); - return hypot(C.x - Px, C.y - Py); -} - -template -double __device__ segment_distance_no_intersect(vec_2d const& A, - vec_2d const& B, - vec_2d const& C, - vec_2d const& D) -{ - return std::min(std::min(point_to_segment_distance(A, C, D), point_to_segment_distance(B, C, D)), - std::min(point_to_segment_distance(C, A, B), point_to_segment_distance(D, A, B))); -} - -/** - * @brief Computes shortest distance between two segments. - * - * If two segment intersects, distance is 0. - */ -template -double __device__ -segment_distance(vec_2d const& A, vec_2d const& B, vec_2d const& C, vec_2d const& D) -{ - // Subject 1.03 of https://www.inf.pucrs.br/~pinho/CG/faq.html - // Construct a parametrized ray of AB and CD, solve for the parameters. - // If both parameters are within [0, 1], the intersection exists. - - double r_denom = (B.x - A.x) * (D.y - C.y) - (B.y - A.y) * (D.x - C.x); - double r_numer = (A.y - C.y) * (D.x - C.x) - (A.x - C.x) * (D.y - C.y); - if (r_denom == 0) { - if (r_numer == 0) { return 0.0; } // Segments coincides - // Segments parallel - return segment_distance_no_intersect(A, B, C, D); - } - double r = r_numer / r_denom; - double s = ((A.y - C.y) * (B.x - A.x) - (A.x - C.x) * (B.y - A.x)) / - ((B.x - A.x) * (D.y - C.y) - (B.y - A.y) * (D.x - C.x)); - if (r >= 0 and r <= 1 and s >= 0 and s <= 1) { return 0.0; } - return segment_distance_no_intersect(A, B, C, D); -} - /** @brief Get the index that is one-past the end point of linestring at @p linestring_idx * * @note The last endpoint of the linestring is not included in the offset array, thus @@ -168,9 +110,11 @@ T __device__ segment_distance_no_intersect_or_colinear(vec_2d const& a, * If two segments intersect, the distance is 0. Otherwise compute the shortest point * to segment distance. */ -template -T __device__ -squared_segment_distance(Vec2d const& a, Vec2d const& b, Vec2d const& c, Vec2d const& d) +template +T __device__ squared_segment_distance(vec_2d const& a, + vec_2d const& b, + vec_2d const& c, + vec_2d const& d) { auto ab = b - a; auto cd = d - c; From af386525113a35f87ae2c23f82f953da34f8633c Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Tue, 17 May 2022 16:30:51 -0700 Subject: [PATCH 11/45] remove dependency on cudf atomics --- .../detail/linestring_distance.cuh | 3 +- .../cuspatial/utility/device_atomics.cuh | 53 +++++++++++++++++++ 2 files changed, 54 insertions(+), 2 deletions(-) create mode 100644 cpp/include/cuspatial/utility/device_atomics.cuh diff --git a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh index 26b1b4a46..88ab73636 100644 --- a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh +++ b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh @@ -17,10 +17,9 @@ #pragma once #include +#include #include -#include - #include #include #include diff --git a/cpp/include/cuspatial/utility/device_atomics.cuh b/cpp/include/cuspatial/utility/device_atomics.cuh new file mode 100644 index 000000000..92d3a6894 --- /dev/null +++ b/cpp/include/cuspatial/utility/device_atomics.cuh @@ -0,0 +1,53 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +namespace cuspatial { +namespace detail { + +__device__ double atomicMin(double* addr, double val) +{ + unsigned long long int* address_as_ull = (unsigned long long int*)addr; + unsigned long long int old = *address_as_ull, assumed; + + do { + assumed = old; + old = atomicCAS( + address_as_ull, assumed, __double_as_longlong(std::min(val, __longlong_as_double(assumed)))); + + // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN) + } while (assumed != old); + + return __longlong_as_double(old); +} + +__device__ float atomicMin(float* addr, float val) +{ + unsigned int* address_as_ull = (unsigned int*)addr; + unsigned int old = *address_as_ull, assumed; + + do { + assumed = old; + old = + atomicCAS(address_as_ull, assumed, __float_as_uint(std::min(val, __uint_as_float(assumed)))); + + // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN) + } while (assumed != old); + + return __uint_as_float(old); +} + +} // namespace detail +} // namespace cuspatial From ae114435c6f41f52732d184ead1f3e0e40642d2f Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Tue, 17 May 2022 16:32:34 -0700 Subject: [PATCH 12/45] regroup includes --- .../cuspatial/experimental/detail/linestring_distance.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh index 88ab73636..5922f7cb1 100644 --- a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh +++ b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh @@ -20,7 +20,6 @@ #include #include -#include #include #include @@ -28,6 +27,7 @@ #include #include +#include #include #include #include From 79d57aa017efcd673a9950ba29bfaef7a1172f4c Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Tue, 17 May 2022 17:15:37 -0700 Subject: [PATCH 13/45] Use size_t as index type. --- .../detail/linestring_distance.cuh | 12 ++--- .../spatial/linestring_distance_test.cu | 45 +++++++++++++++++-- 2 files changed, 48 insertions(+), 9 deletions(-) diff --git a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh index 5922f7cb1..1aff1ed50 100644 --- a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh +++ b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh @@ -183,16 +183,18 @@ void __global__ pairwise_linestring_distance_kernel(OffsetIterator linestring1_o { using T = typename std::iterator_traits::value_type::value_type; - auto const p1_idx = threadIdx.x + blockIdx.x * blockDim.x; - auto const num_linestrings = thrust::distance(linestring1_offsets_begin, linestring1_offsets_end); - auto const linestring1_num_points = + auto const p1_idx = threadIdx.x + blockIdx.x * blockDim.x; + size_t const num_linestrings = + thrust::distance(linestring1_offsets_begin, linestring1_offsets_end); + + size_t const linestring1_num_points = thrust::distance(linestring1_points_begin, linestring1_points_end); - auto const linestring2_num_points = + size_t const linestring2_num_points = thrust::distance(linestring2_points_begin, linestring2_points_end); if (p1_idx >= linestring1_num_points) { return; } - auto const linestring_idx = + size_t const linestring_idx = thrust::distance(linestring1_offsets_begin, thrust::upper_bound( thrust::seq, linestring1_offsets_begin, linestring1_offsets_end, p1_idx)) - diff --git a/cpp/tests/experimental/spatial/linestring_distance_test.cu b/cpp/tests/experimental/spatial/linestring_distance_test.cu index cf05fedf6..49d76c967 100644 --- a/cpp/tests/experimental/spatial/linestring_distance_test.cu +++ b/cpp/tests/experimental/spatial/linestring_distance_test.cu @@ -27,6 +27,9 @@ #include #include +#include +#include + namespace cuspatial { namespace test { @@ -52,7 +55,7 @@ TYPED_TEST(PairwiseLinestringDistanceTest, FromSeparateArrayInputs) CartVec({{0.0f, 1.0f}, {1.0f, 1.0f}, {2.0f, 1.0f}, {3.0f, 1.0f}, {4.0f, 1.0f}})}; auto offset = rmm::device_vector{std::vector{std::vector{0}}}; - auto distance = rmm::device_vector{std::vector{0.0}}; + auto distance = rmm::device_vector{1}; auto expected = rmm::device_vector{std::vector{1.0}}; pairwise_linestring_distance(offset.begin(), @@ -81,7 +84,7 @@ TYPED_TEST(PairwiseLinestringDistanceTest, FromSamePointArrayInput) auto b_begin = cart2ds.begin() + 1; auto b_end = cart2ds.end(); - auto distance = rmm::device_vector{std::vector{0.0}}; + auto distance = rmm::device_vector{1}; auto expected = rmm::device_vector{std::vector{0.0}}; pairwise_linestring_distance( @@ -109,7 +112,7 @@ TYPED_TEST(PairwiseLinestringDistanceTest, FromTransformIterator) auto offset = rmm::device_vector{std::vector{0}}; - auto distance = rmm::device_vector{std::vector{0.0}}; + auto distance = rmm::device_vector{1}; auto expected = rmm::device_vector{std::vector{1.0}}; pairwise_linestring_distance( @@ -134,7 +137,7 @@ TYPED_TEST(PairwiseLinestringDistanceTest, FromMixedIterator) auto offset = rmm::device_vector{std::vector{0}}; - auto distance = rmm::device_vector{std::vector{0.0}}; + auto distance = rmm::device_vector{1}; auto expected = rmm::device_vector{std::vector{1.0}}; pairwise_linestring_distance(offset.begin(), @@ -149,5 +152,39 @@ TYPED_TEST(PairwiseLinestringDistanceTest, FromMixedIterator) EXPECT_EQ(distance, expected); } +TYPED_TEST(PairwiseLinestringDistanceTest, FromLongInputs) +{ + using T = TypeParam; + using CartVec = std::vector>; + + auto num_points = 1000; + + auto a_cart2d_x_begin = thrust::make_constant_iterator(T{0.0}); + auto a_cart2d_y_begin = thrust::make_counting_iterator(T{0.0}); + auto a_cart2d_begin = make_cartesian_2d_iterator(a_cart2d_x_begin, a_cart2d_y_begin); + auto a_cart2d_end = a_cart2d_begin + num_points; + + auto b_cart2d_x_begin = thrust::make_constant_iterator(T{42.0}); + auto b_cart2d_y_begin = thrust::make_counting_iterator(T{0.0}); + auto b_cart2d_begin = make_cartesian_2d_iterator(b_cart2d_x_begin, b_cart2d_y_begin); + auto b_cart2d_end = b_cart2d_begin + num_points; + + auto offset = rmm::device_vector{std::vector{0, 100, 200, 300, 400}}; + + auto distance = rmm::device_vector{1}; + auto expected = rmm::device_vector{std::vector{42.0, 42.0, 42.0, 42.0, 42.0}}; + + pairwise_linestring_distance(offset.begin(), + offset.end(), + a_cart2d_begin, + a_cart2d_end, + offset.begin(), + b_cart2d_begin, + b_cart2d_end, + distance.begin()); + + EXPECT_EQ(distance, expected); +} + } // namespace test } // namespace cuspatial From e2a3db723e4e6071ec2b10bb2a65eff2b6238795 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Tue, 17 May 2022 17:41:13 -0700 Subject: [PATCH 14/45] some fixes on tests --- cpp/tests/experimental/spatial/linestring_distance_test.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/tests/experimental/spatial/linestring_distance_test.cu b/cpp/tests/experimental/spatial/linestring_distance_test.cu index 49d76c967..adaae771c 100644 --- a/cpp/tests/experimental/spatial/linestring_distance_test.cu +++ b/cpp/tests/experimental/spatial/linestring_distance_test.cu @@ -171,7 +171,7 @@ TYPED_TEST(PairwiseLinestringDistanceTest, FromLongInputs) auto offset = rmm::device_vector{std::vector{0, 100, 200, 300, 400}}; - auto distance = rmm::device_vector{1}; + auto distance = rmm::device_vector{5}; auto expected = rmm::device_vector{std::vector{42.0, 42.0, 42.0, 42.0, 42.0}}; pairwise_linestring_distance(offset.begin(), From a1512d85a8923f65fa3c31b1e7edd9c4478df436 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 18 May 2022 12:01:23 -0700 Subject: [PATCH 15/45] Revert cmake-format and precommit hooks --- .pre-commit-config.yaml | 18 --------- cpp/scripts/run-cmake-format.sh | 72 --------------------------------- 2 files changed, 90 deletions(-) delete mode 100755 cpp/scripts/run-cmake-format.sh diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index e988539c3..70b4bde47 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -40,23 +40,5 @@ repos: language: system files: \.(cu|cuh|h|hpp|cpp|inl)$ args: ['-fallback-style=none'] - - id: cmake-format - name: cmake-format - entry: ./cpp/scripts/run-cmake-format.sh cmake-format - language: python - types: [cmake] - # Note that pre-commit autoupdate does not update the versions - # of dependencies, so we'll have to update this manually. - additional_dependencies: - - cmakelang==0.6.13 - - id: cmake-lint - name: cmake-lint - entry: ./cpp/scripts/run-cmake-format.sh cmake-lint - language: python - types: [cmake] - # Note that pre-commit autoupdate does not update the versions - # of dependencies, so we'll have to update this manually. - additional_dependencies: - - cmakelang==0.6.13 default_language_version: python: python3 diff --git a/cpp/scripts/run-cmake-format.sh b/cpp/scripts/run-cmake-format.sh deleted file mode 100755 index 5c47704e3..000000000 --- a/cpp/scripts/run-cmake-format.sh +++ /dev/null @@ -1,72 +0,0 @@ -#!/bin/bash - -# This script is a wrapper for cmakelang that may be used with pre-commit. The -# wrapping is necessary because RAPIDS libraries split configuration for -# cmakelang linters between a local config file and a second config file that's -# shared across all of RAPIDS via rapids-cmake. In order to keep it up to date -# this file is only maintained in one place (the rapids-cmake repo) and -# pulled down during builds. We need a way to invoke CMake linting commands -# without causing pre-commit failures (which could block local commits or CI), -# while also being sufficiently flexible to allow users to maintain the config -# file independently of a build directory. -# -# This script provides the minimal functionality to enable those use cases. It -# searches in a number of predefined locations for the rapids-cmake config file -# and exits gracefully if the file is not found. If a user wishes to specify a -# config file at a nonstandard location, they may do so by setting the -# environment variable RAPIDS_CMAKE_FORMAT_FILE. -# -# This script can be invoked directly anywhere within the project repository. -# Alternatively, it may be invoked as a pre-commit hook via -# `pre-commit run (cmake-format)|(cmake-lint)`. -# -# Usage: -# bash run-cmake-format.sh {cmake-format,cmake-lint} infile [infile ...] - -status=0 -if [ -z ${CUSPATIAL_ROOT:+PLACEHOLDER} ]; then - CUSPATIAL_BUILD_DIR=$(git rev-parse --show-toplevel 2>&1)/cpp/build - status=$? -else - CUSPATIAL_BUILD_DIR=${CUSPATIAL_ROOT} -fi - -if ! [ ${status} -eq 0 ]; then - if [[ ${CUSPATIAL_BUILD_DIR} == *"not a git repository"* ]]; then - echo "This script must be run inside the cuspatial repository, or the CUSPATIAL_ROOT environment variable must be set." - else - echo "Script failed with unknown error attempting to determine project root:" - echo ${CUSPATIAL_BUILD_DIR} - fi - exit 1 -fi - -DEFAULT_FORMAT_FILE_LOCATIONS=( - "${CUSPATIAL_BUILD_DIR:-${HOME}}/_deps/rapids-cmake-src/cmake-format-rapids-cmake.json" -) - -if [ -z ${RAPIDS_CMAKE_FORMAT_FILE:+PLACEHOLDER} ]; then - for file_path in ${DEFAULT_FORMAT_FILE_LOCATIONS[@]}; do - if [ -f ${file_path} ]; then - RAPIDS_CMAKE_FORMAT_FILE=${file_path} - break - fi - done -fi - -if [ -z ${RAPIDS_CMAKE_FORMAT_FILE:+PLACEHOLDER} ]; then - echo "The rapids-cmake cmake-format configuration file was not found at any of the default search locations: " - echo "" - ( IFS=$'\n'; echo "${DEFAULT_FORMAT_FILE_LOCATIONS[*]}" ) - echo "" - echo "Try setting the environment variable RAPIDS_CMAKE_FORMAT_FILE to the path to the config file." - exit 0 -else - echo "Using format file ${RAPIDS_CMAKE_FORMAT_FILE}" -fi - -if [[ $1 == "cmake-format" ]]; then - cmake-format -i --config-files cpp/cmake/config.json ${RAPIDS_CMAKE_FORMAT_FILE} -- ${@:2} -elif [[ $1 == "cmake-lint" ]]; then - cmake-lint --config-files cpp/cmake/config.json ${RAPIDS_CMAKE_FORMAT_FILE} -- ${@:2} -fi From 9b1687c19fc58b3bb567c9ff848e769cf9d3e7f4 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 18 May 2022 12:17:42 -0700 Subject: [PATCH 16/45] Remove `Cart2dA` and `Cart2dB` --- .../experimental/detail/linestring_distance.cuh | 13 +++++-------- .../cuspatial/experimental/linestring_distance.cuh | 11 +---------- 2 files changed, 6 insertions(+), 18 deletions(-) diff --git a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh index 1aff1ed50..cb9aef02f 100644 --- a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh +++ b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh @@ -228,12 +228,7 @@ void __global__ pairwise_linestring_distance_kernel(OffsetIterator linestring1_o } // namespace detail -template +template void pairwise_linestring_distance(OffsetIterator linestring1_offsets_first, OffsetIterator linestring1_offsets_last, Cart2dItA linestring1_points_first, @@ -244,7 +239,7 @@ void pairwise_linestring_distance(OffsetIterator linestring1_offsets_first, OutputIt distances_first, rmm::cuda_stream_view stream) { - using T = typename Cart2dA::value_type; + using T = typename std::iterator_traits::value_type::value_type; static_assert(detail::is_floating_point::value_type>(), "Inputs and output must be the same types."); - static_assert(detail::is_same, Cart2dA, Cart2dB>(), + static_assert(detail::is_same, + std::iterator_traits::value_type, + std::iterator_traits::value_type>(), "Inputs must be cuspatial::cartesian_2d"); auto const num_string_pairs = diff --git a/cpp/include/cuspatial/experimental/linestring_distance.cuh b/cpp/include/cuspatial/experimental/linestring_distance.cuh index 0fbd9ef2c..6209380f4 100644 --- a/cpp/include/cuspatial/experimental/linestring_distance.cuh +++ b/cpp/include/cuspatial/experimental/linestring_distance.cuh @@ -35,10 +35,6 @@ namespace cuspatial { * [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. * @tparam OutputIt iterator type for output array. Must meet the requirements of * [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. - * @tparam std::iterator_traits::value_type value type of `Cart2dItA`, must be - * `cuspatial::cartesian_2d` - * @tparam std::iterator_traits::value_type value type of `Cart2dItB`, must be - * `cuspatial::cartesian_2d` * * @param linestring1_offsets_first begin of range of the offsets to the first linestring of each * pair @@ -59,12 +55,7 @@ namespace cuspatial { * [LinkLRAI]: https://en.cppreference.com/w/cpp/named_req/RandomAccessIterator * "LegacyRandomAccessIterator" */ -template ::value_type, - class Cart2dB = typename std::iterator_traits::value_type> +template void pairwise_linestring_distance(OffsetIterator linestring1_offsets_first, OffsetIterator linestring1_offsets_last, Cart2dItA linestring1_points_first, From dea41f129ffb58b16fc250167f78f31edef4271b Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 18 May 2022 12:43:55 -0700 Subject: [PATCH 17/45] Documentation update Co-authored-by: Mark Harris --- .../cuspatial/experimental/linestring_distance.cuh | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/cpp/include/cuspatial/experimental/linestring_distance.cuh b/cpp/include/cuspatial/experimental/linestring_distance.cuh index 6209380f4..e40038ac9 100644 --- a/cpp/include/cuspatial/experimental/linestring_distance.cuh +++ b/cpp/include/cuspatial/experimental/linestring_distance.cuh @@ -36,17 +36,17 @@ namespace cuspatial { * @tparam OutputIt iterator type for output array. Must meet the requirements of * [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. * - * @param linestring1_offsets_first begin of range of the offsets to the first linestring of each + * @param linestring1_offsets_first beginning of range of the offsets to the first linestring of each * pair * @param linestring1_offsets_last end of range of the offsets to the first linestring of each pair - * @param linestring1_points_first begin of range of the point of the first linestring of each pair + * @param linestring1_points_first beginning of range of the point of the first linestring of each pair * @param linestring1_points_last end of range of the point of the first linestring of each pair - * @param linestring2_offsets_first begin of range of the offsets to the second linestring of each + * @param linestring2_offsets_first beginning of range of the offsets to the second linestring of each * pair - * @param linestring2_points_first begin of range of the point of the second linestring of each pair + * @param linestring2_points_first beginning of range of the point of the second linestring of each pair * @param linestring2_points_last end of range of the point of the second linestring of each pair - * @param distances_first begin to output array - * @param stream Used for device memory operations and kernel launches. + * @param distances_first beginning iterator to output + * @param stream The CUDA stream to use for device memory operations and kernel launches. * * @pre all input iterators for coordinates must have `cuspatial::cartesian_2d` type. * @pre all scalar types must be floating point types, and must be the same type for all input From 270cfe21b343a6ce374a350956a42931f0af53b8 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 18 May 2022 13:44:04 -0700 Subject: [PATCH 18/45] Style fix --- .../cuspatial/experimental/linestring_distance.cuh | 12 +++++++----- 1 file changed, 7 insertions(+), 5 deletions(-) diff --git a/cpp/include/cuspatial/experimental/linestring_distance.cuh b/cpp/include/cuspatial/experimental/linestring_distance.cuh index e40038ac9..9000e090b 100644 --- a/cpp/include/cuspatial/experimental/linestring_distance.cuh +++ b/cpp/include/cuspatial/experimental/linestring_distance.cuh @@ -36,14 +36,16 @@ namespace cuspatial { * @tparam OutputIt iterator type for output array. Must meet the requirements of * [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. * - * @param linestring1_offsets_first beginning of range of the offsets to the first linestring of each - * pair + * @param linestring1_offsets_first beginning of range of the offsets to the first linestring of + * each pair * @param linestring1_offsets_last end of range of the offsets to the first linestring of each pair - * @param linestring1_points_first beginning of range of the point of the first linestring of each pair + * @param linestring1_points_first beginning of range of the point of the first linestring of each + * pair * @param linestring1_points_last end of range of the point of the first linestring of each pair - * @param linestring2_offsets_first beginning of range of the offsets to the second linestring of each + * @param linestring2_offsets_first beginning of range of the offsets to the second linestring of + * each pair + * @param linestring2_points_first beginning of range of the point of the second linestring of each * pair - * @param linestring2_points_first beginning of range of the point of the second linestring of each pair * @param linestring2_points_last end of range of the point of the second linestring of each pair * @param distances_first beginning iterator to output * @param stream The CUDA stream to use for device memory operations and kernel launches. From cb5d301350b37aa0b6752eca5405f0e6788d358c Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 18 May 2022 20:24:14 -0700 Subject: [PATCH 19/45] fix broken compile --- .../experimental/detail/linestring_distance.cuh | 15 ++++++++------- 1 file changed, 8 insertions(+), 7 deletions(-) diff --git a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh index cb9aef02f..f8d19233b 100644 --- a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh +++ b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh @@ -241,19 +241,20 @@ void pairwise_linestring_distance(OffsetIterator linestring1_offsets_first, { using T = typename std::iterator_traits::value_type::value_type; - static_assert(detail::is_floating_point::value_type>(), - "Inputs and output must be floating point types."); + static_assert( + detail::is_floating_point::value_type::value_type, + typename std::iterator_traits::value_type>(), + "Inputs and output must be floating point types."); static_assert(detail::is_same::value_type::value_type, typename std::iterator_traits::value_type>(), "Inputs and output must be the same types."); static_assert(detail::is_same, - std::iterator_traits::value_type, - std::iterator_traits::value_type>(), + typename std::iterator_traits::value_type, + typename std::iterator_traits::value_type>(), "Inputs must be cuspatial::cartesian_2d"); auto const num_string_pairs = From 183c10cadde592366aff2531f23c3c92da8e17a8 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 18 May 2022 21:59:10 -0700 Subject: [PATCH 20/45] Removes `device_atomics` usage --- .../detail/linestring_distance.cuh | 7 +-- .../cuspatial/utility/device_atomics.cuh | 53 ------------------- 2 files changed, 4 insertions(+), 56 deletions(-) delete mode 100644 cpp/include/cuspatial/utility/device_atomics.cuh diff --git a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh index f8d19233b..763a0ce16 100644 --- a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh +++ b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh @@ -17,7 +17,6 @@ #pragma once #include -#include #include #include @@ -27,6 +26,8 @@ #include #include +#include + #include #include #include @@ -222,8 +223,8 @@ void __global__ pairwise_linestring_distance_kernel(OffsetIterator linestring1_o min_squared_distance = std::min(min_squared_distance, squared_segment_distance(A, B, C, D)); } - atomicMin(&thrust::raw_reference_cast(*(distances + linestring_idx)), - static_cast(std::sqrt(min_squared_distance))); + cuda::atomic_ref ref{thrust::raw_reference_cast(*(distances + linestring_idx))}; + ref.fetch_min(static_cast(std::sqrt(min_squared_distance)), cuda::memory_order_relaxed); } } // namespace detail diff --git a/cpp/include/cuspatial/utility/device_atomics.cuh b/cpp/include/cuspatial/utility/device_atomics.cuh deleted file mode 100644 index 92d3a6894..000000000 --- a/cpp/include/cuspatial/utility/device_atomics.cuh +++ /dev/null @@ -1,53 +0,0 @@ -/* - * Copyright (c) 2022, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -namespace cuspatial { -namespace detail { - -__device__ double atomicMin(double* addr, double val) -{ - unsigned long long int* address_as_ull = (unsigned long long int*)addr; - unsigned long long int old = *address_as_ull, assumed; - - do { - assumed = old; - old = atomicCAS( - address_as_ull, assumed, __double_as_longlong(std::min(val, __longlong_as_double(assumed)))); - - // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN) - } while (assumed != old); - - return __longlong_as_double(old); -} - -__device__ float atomicMin(float* addr, float val) -{ - unsigned int* address_as_ull = (unsigned int*)addr; - unsigned int old = *address_as_ull, assumed; - - do { - assumed = old; - old = - atomicCAS(address_as_ull, assumed, __float_as_uint(std::min(val, __uint_as_float(assumed)))); - - // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN) - } while (assumed != old); - - return __uint_as_float(old); -} - -} // namespace detail -} // namespace cuspatial From b2fbaad10195cfef5b4207e003b9dd6c167ec86a Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 18 May 2022 22:08:21 -0700 Subject: [PATCH 21/45] Add `internal` marker to internal docstrings. --- .../experimental/detail/linestring_distance.cuh | 11 +++++++++-- 1 file changed, 9 insertions(+), 2 deletions(-) diff --git a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh index 763a0ce16..913e68242 100644 --- a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh +++ b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh @@ -48,7 +48,9 @@ constexpr bool is_floating_point() return std::conjunction_v...>; } -/** @brief Get the index that is one-past the end point of linestring at @p linestring_idx +/** + * @internal + * @brief Get the index that is one-past the end point of linestring at @p linestring_idx * * @note The last endpoint of the linestring is not included in the offset array, thus * @p num_points is returned. @@ -67,6 +69,7 @@ endpoint_index_of_linestring(SizeType const& linestring_idx, } /** + * @internal * @brief Computes shortest distance between @p c and segment ab */ template @@ -88,6 +91,7 @@ T __device__ point_to_segment_distance_squared(vec_2d const& c, } /** + * @internal * @brief Computes shortest distance between two segments (ab and cd) that * doesn't intersect. */ @@ -105,6 +109,7 @@ T __device__ segment_distance_no_intersect_or_colinear(vec_2d const& a, } /** + * @internal * @brief Computes shortest distance between two segments. * * If two segments intersect, the distance is 0. Otherwise compute the shortest point @@ -135,6 +140,7 @@ T __device__ squared_segment_distance(vec_2d const& a, } /** + * @internal * @brief The kernel to compute point to linestring distance * * Each thread of the kernel computes the distance between a segment in a linestring in pair 1 @@ -223,7 +229,8 @@ void __global__ pairwise_linestring_distance_kernel(OffsetIterator linestring1_o min_squared_distance = std::min(min_squared_distance, squared_segment_distance(A, B, C, D)); } - cuda::atomic_ref ref{thrust::raw_reference_cast(*(distances + linestring_idx))}; + cuda::atomic_ref ref{ + thrust::raw_reference_cast(*(distances + linestring_idx))}; ref.fetch_min(static_cast(std::sqrt(min_squared_distance)), cuda::memory_order_relaxed); } From 1752485b38123fb049fd8603fa640a74a90b40b2 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 18 May 2022 22:59:53 -0700 Subject: [PATCH 22/45] Move derived traits to traits.hpp --- .../detail/linestring_distance.cuh | 16 ++--------- cpp/include/cuspatial/utility/traits.hpp | 27 +++++++++++++++++++ 2 files changed, 29 insertions(+), 14 deletions(-) create mode 100644 cpp/include/cuspatial/utility/traits.hpp diff --git a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh index 913e68242..d5ba7a5c2 100644 --- a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh +++ b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh @@ -17,6 +17,7 @@ #pragma once #include +#include #include #include @@ -36,18 +37,6 @@ namespace cuspatial { namespace detail { -template -constexpr bool is_same() -{ - return std::conjunction_v...>; -} - -template -constexpr bool is_floating_point() -{ - return std::conjunction_v...>; -} - /** * @internal * @brief Get the index that is one-past the end point of linestring at @p linestring_idx @@ -229,8 +218,7 @@ void __global__ pairwise_linestring_distance_kernel(OffsetIterator linestring1_o min_squared_distance = std::min(min_squared_distance, squared_segment_distance(A, B, C, D)); } - cuda::atomic_ref ref{ - thrust::raw_reference_cast(*(distances + linestring_idx))}; + cuda::atomic_ref ref{*(distances + linestring_idx)}; ref.fetch_min(static_cast(std::sqrt(min_squared_distance)), cuda::memory_order_relaxed); } diff --git a/cpp/include/cuspatial/utility/traits.hpp b/cpp/include/cuspatial/utility/traits.hpp new file mode 100644 index 000000000..e4464a272 --- /dev/null +++ b/cpp/include/cuspatial/utility/traits.hpp @@ -0,0 +1,27 @@ +#include + +namespace cuspatial { +namespace detail { + +/** + * @internal + * @brief returns true if all types are the same. + */ +template +constexpr bool is_same() +{ + return std::conjunction_v...>; +} + +/** + * @internal + * @brief returns true if all types are floating point types. + */ +template +constexpr bool is_floating_point() +{ + return std::conjunction_v...>; +} + +} // namespace detail +} // namespace cuspatial From fe83ea2acef701672db8b2985bd603f7026534e7 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 19 May 2022 11:08:01 -0700 Subject: [PATCH 23/45] add back raw_reference_cast --- .../cuspatial/experimental/detail/linestring_distance.cuh | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh index d5ba7a5c2..dd1b297de 100644 --- a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh +++ b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh @@ -218,7 +218,8 @@ void __global__ pairwise_linestring_distance_kernel(OffsetIterator linestring1_o min_squared_distance = std::min(min_squared_distance, squared_segment_distance(A, B, C, D)); } - cuda::atomic_ref ref{*(distances + linestring_idx)}; + cuda::atomic_ref ref{ + thrust::raw_reference_cast(*(distances + linestring_idx))}; ref.fetch_min(static_cast(std::sqrt(min_squared_distance)), cuda::memory_order_relaxed); } From 38dfffd4a7fc0068bd070529e9fed4490f1b51eb Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Tue, 24 May 2022 21:07:33 -0700 Subject: [PATCH 24/45] initial of point distance --- cpp/CMakeLists.txt | 1 + .../cuspatial/distances/point_distance.hpp | 34 ++++++ .../experimental/detail/point_distance.cuh | 66 +++++++++++ .../cuspatial/experimental/point_distance.cuh | 60 ++++++++++ cpp/src/spatial/point_distance.cu | 110 ++++++++++++++++++ 5 files changed, 271 insertions(+) create mode 100644 cpp/include/cuspatial/distances/point_distance.hpp create mode 100644 cpp/include/cuspatial/experimental/detail/point_distance.cuh create mode 100644 cpp/include/cuspatial/experimental/point_distance.cuh create mode 100644 cpp/src/spatial/point_distance.cu diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 375403f1e..3e49d3342 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -120,6 +120,7 @@ add_library(cuspatial src/spatial/haversine.cu src/spatial/hausdorff.cu src/spatial/linestring_distance.cu + src/spatial/point_distance.cu src/spatial/lonlat_to_cartesian.cu src/trajectory/derive_trajectories.cu src/trajectory/trajectory_bounding_boxes.cu diff --git a/cpp/include/cuspatial/distances/point_distance.hpp b/cpp/include/cuspatial/distances/point_distance.hpp new file mode 100644 index 000000000..47b947a50 --- /dev/null +++ b/cpp/include/cuspatial/distances/point_distance.hpp @@ -0,0 +1,34 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include +#include + +namespace cuspatial { + +/** + * @brief Compute pairwise point to point cartesian distance + */ +std::unique_ptr pairwise_point_distance( + cudf::column_view const& points1_x, + cudf::column_view const& points1_y, + cudf::column_view const& points2_x, + cudf::column_view const& points2_y, + rmm::cuda_stream_view stream = rmm::cuda_stream_default); + +} // namespace cuspatial diff --git a/cpp/include/cuspatial/experimental/detail/point_distance.cuh b/cpp/include/cuspatial/experimental/detail/point_distance.cuh new file mode 100644 index 000000000..8027b070a --- /dev/null +++ b/cpp/include/cuspatial/experimental/detail/point_distance.cuh @@ -0,0 +1,66 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include +#include +#include + +#include +#include + +#include + +namespace cuspatial { + +template +OutputIt pairwise_point_distance(Cart2dItA points1_first, + Cart2dItA points1_last, + Cart2dItB points2_first, + OutputIt distances_first, + rmm::cuda_stream_view stream) +{ + using T = typename std::iterator_traits::value_type::value_type; + + static_assert( + detail::is_floating_point::value_type::value_type, + typename std::iterator_traits::value_type>(), + "Inputs and output must be floating point types."); + + static_assert(detail::is_same::value_type::value_type, + typename std::iterator_traits::value_type>(), + "Inputs and output must be the same types."); + + static_assert(detail::is_same, + typename std::iterator_traits::value_type, + typename std::iterator_traits::value_type>(), + "Inputs must be cuspatial::cartesian_2d"); + + return thrust::transform(rmm::exec_policy(stream), + points1_first, + points1_last, + points2_first, + distances_first, + [] __device__(auto p1, auto p2) { + auto v = p1 - p2; + return std::sqrt(dot(v, v)); + }); +} + +} // namespace cuspatial diff --git a/cpp/include/cuspatial/experimental/point_distance.cuh b/cpp/include/cuspatial/experimental/point_distance.cuh new file mode 100644 index 000000000..79d7298b5 --- /dev/null +++ b/cpp/include/cuspatial/experimental/point_distance.cuh @@ -0,0 +1,60 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include + +namespace cuspatial { + +/** + * @copybrief cuspatial::pairwise_point_distance + * + * Computes pairwise distance between points. + * + * @tparam Cart2dItA iterator type for point array of the first point of each pair. Must meet + * the requirements of [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. + * @tparam Cart2dItB iterator type for point array of the second point of each pair. Must meet + * the requirements of [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. + * @tparam OutputIt iterator type for output array. Must meet the requirements of + * [LegacyRandomAccessIterator][LinkLRAI], be mutable and be device-accessible. + * + * @param points1_first beginning of range of the point of the first point of each + * pair + * @param points1_last end of range of the point of the first point of each pair + * @param points2_first beginning of range of the point of the second point of each + * pair + * @param distances_first beginning iterator to output + * @param stream The CUDA stream to use for device memory operations and kernel launches + * @return Output iterator to one past the last element in the output range + * + * @pre all input iterators for coordinates must have `cuspatial::cartesian_2d` type. + * @pre all scalar types must be floating point types, and must be the same type for all input + * iterators and output iterators. + * + * [LinkLRAI]: https://en.cppreference.com/w/cpp/named_req/RandomAccessIterator + * "LegacyRandomAccessIterator" + */ +template +OutputIt pairwise_point_distance(Cart2dItA points1_first, + Cart2dItA points1_last, + Cart2dItB points2_first, + OutputIt distances_first, + rmm::cuda_stream_view stream = rmm::cuda_stream_default); + +} // namespace cuspatial + +#include diff --git a/cpp/src/spatial/point_distance.cu b/cpp/src/spatial/point_distance.cu new file mode 100644 index 000000000..b0ea5a1fe --- /dev/null +++ b/cpp/src/spatial/point_distance.cu @@ -0,0 +1,110 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include +#include +#include + +#include +#include +#include +#include + +#include +#include + +#include +#include + +namespace cuspatial { +namespace detail { +struct pairwise_point_distance_functor { + template + std::enable_if_t::value, std::unique_ptr> operator()( + cudf::column_view const& points1_x, + cudf::column_view const& points1_y, + cudf::column_view const& points2_x, + cudf::column_view const& points2_y, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) + { + auto distances = cudf::make_numeric_column(cudf::data_type{cudf::type_to_id()}, + points1_x.size(), + cudf::mask_state::UNALLOCATED, + stream, + mr); + + auto points1_it = make_cartesian_2d_iterator(points1_x.begin(), points1_y.begin()); + auto points2_it = make_cartesian_2d_iterator(points2_x.begin(), points2_y.begin()); + + pairwise_point_distance(points1_it, + points1_it + points1_x.size(), + points2_it, + distances->mutable_view().begin(), + stream); + + return distances; + } + + template + std::enable_if_t::value, std::unique_ptr> operator()( + Args&&...) + { + CUSPATIAL_FAIL("Point distances only supports floating point coordinates."); + } +}; + +std::unique_ptr pairwise_point_distance(cudf::column_view const& points1_x, + cudf::column_view const& points1_y, + cudf::column_view const& points2_x, + cudf::column_view const& points2_y, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + CUSPATIAL_EXPECTS(points1_x.size() == points1_y.size() and + points2_x.size() == points2_y.size() and points1_x.size() == points2_x.size(), + "Mismatch number of coordinate or number of points."); + + CUSPATIAL_EXPECTS(points1_x.type() == points1_y.type() and + points2_x.type() == points2_y.type() and points1_x.type() == points2_x.type(), + "The types of point coordinates arrays mismatch."); + + if (points1_x.size() == 0) { return cudf::empty_like(points1_x); } + + return cudf::type_dispatcher(points1_x.type(), + pairwise_point_distance_functor{}, + points1_x, + points1_y, + points2_x, + points2_y, + stream, + mr); +} + +} // namespace detail + +std::unique_ptr pairwise_point_distance(cudf::column_view const& points1_x, + cudf::column_view const& points1_y, + cudf::column_view const& points2_x, + cudf::column_view const& points2_y, + rmm::mr::device_memory_resource* mr) +{ + return detail::pairwise_point_distance( + points1_x, points1_y, points2_x, points2_y, rmm::cuda_stream_default, mr); +} + +} // namespace cuspatial From 6089a791ca76801989e4bdab3e26b4bcba2c4c85 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 8 Jun 2022 14:36:25 -0700 Subject: [PATCH 25/45] add header only api test --- cpp/tests/CMakeLists.txt | 3 + .../spatial/point_distance_test.cu | 276 ++++++++++++++++++ 2 files changed, 279 insertions(+) create mode 100644 cpp/tests/experimental/spatial/point_distance_test.cu diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index a2c400788..186ab660c 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -105,6 +105,9 @@ ConfigureTest(SPATIAL_WINDOW_POINT_TEST ConfigureTest(HAVERSINE_TEST_EXP experimental/spatial/haversine_test.cu) +ConfigureTest(POINT_DISTANCE_TEST_EXP + experimental/spatial/point_distance_test.cu) + ConfigureTest(LINESTRING_DISTANCE_TEST_EXP experimental/spatial/linestring_distance_test.cu) diff --git a/cpp/tests/experimental/spatial/point_distance_test.cu b/cpp/tests/experimental/spatial/point_distance_test.cu new file mode 100644 index 000000000..3bd5557d4 --- /dev/null +++ b/cpp/tests/experimental/spatial/point_distance_test.cu @@ -0,0 +1,276 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include +#include +#include + +#include +#include + +#include +#include +#include +#include +#include +#include + +#include +#include + +namespace cuspatial { + +template +struct PairwisePointDistanceTest : public ::testing::Test { +}; + +template +auto compute_point_distance_host(Cart2DVec const& point1, Cart2DVec const& point2) +{ + using Cart2D = typename Cart2DVec::value_type; + using T = typename Cart2D::value_type; + thrust::host_vector h_point1(point1); + thrust::host_vector h_point2(point2); + auto pair_iter = + thrust::make_zip_iterator(thrust::make_tuple(h_point1.begin(), h_point2.begin())); + auto result_iter = thrust::make_transform_iterator(pair_iter, [](auto p) { + auto p0 = thrust::get<0>(p); + auto p1 = thrust::get<1>(p); + return std::sqrt((p0.x - p1.x) * (p0.x - p1.x) + (p0.y - p1.y) * (p0.y - p1.y)); + }); + + return thrust::host_vector(result_iter, result_iter + point1.size()); +} + +using TestTypes = ::testing::Types; + +TYPED_TEST_CASE(PairwisePointDistanceTest, TestTypes); + +TYPED_TEST(PairwisePointDistanceTest, Empty) +{ + using T = TypeParam; + using Cart2D = cartesian_2d; + using Cart2DVec = std::vector; + + rmm::device_vector points1{}; + rmm::device_vector points2{}; + + rmm::device_vector expected{}; + rmm::device_vector got(points1.size()); + + pairwise_point_distance(points1.begin(), points1.end(), points2.begin(), got.begin()); + + EXPECT_EQ(expected, got); +} + +TYPED_TEST(PairwisePointDistanceTest, OnePair) +{ + using T = TypeParam; + using Cart2D = cartesian_2d; + using Cart2DVec = std::vector; + + rmm::device_vector points1{Cart2DVec{{1.0, 1.0}}}; + rmm::device_vector points2{Cart2DVec{{0.0, 0.0}}}; + + rmm::device_vector expected{std::vector{std::sqrt(T{2.0})}}; + rmm::device_vector got(points1.size()); + + pairwise_point_distance(points1.begin(), points1.end(), points2.begin(), got.begin()); + + EXPECT_EQ(expected, got); +} + +template +struct RandomPointGenerator { + using Cart2D = cartesian_2d; + thrust::minstd_rand rng{}; + thrust::random::normal_distribution norm_dist{}; + + Cart2D __device__ operator()(size_t const& i) + { + rng.discard(i); + return Cart2D{norm_dist(rng), norm_dist(rng)}; + } +}; + +TYPED_TEST(PairwisePointDistanceTest, ManyRandom) +{ + using T = TypeParam; + using Cart2D = cartesian_2d; + using Cart2DVec = std::vector; + + std::size_t constexpr num_points = 1000; + + rmm::device_vector points1(num_points); + rmm::device_vector points2(num_points); + + auto counting_iter1 = thrust::make_counting_iterator(0); + auto counting_iter2 = thrust::make_counting_iterator(num_points); + + thrust::transform( + counting_iter1, counting_iter1 + num_points, points1.begin(), RandomPointGenerator{}); + thrust::transform( + counting_iter2, counting_iter2 + num_points, points2.begin(), RandomPointGenerator{}); + + auto expected = compute_point_distance_host(points1, points2); + rmm::device_vector got(points1.size()); + + pairwise_point_distance(points1.begin(), points1.end(), points2.begin(), got.begin()); + thrust::host_vector hgot(got); + + if constexpr (std::is_same_v) { + EXPECT_THAT(expected, ::testing::Pointwise(::testing::FloatEq(), hgot)); + } else { + EXPECT_THAT(expected, ::testing::Pointwise(::testing::DoubleEq(), hgot)); + } +} + +TYPED_TEST(PairwisePointDistanceTest, CompareWithShapely) +{ + using T = TypeParam; + using Cart2D = cartesian_2d; + using Cart2DVec = std::vector; + + std::size_t constexpr num_points = 100; + + std::vector x1{ + -12.30983106, -7.92705956, -49.9570584, -1.05124645, -89.39777526, -32.46014839, + -20.64749623, 74.88373211, -3.56663354, -91.43203925, 1.68283845, 30.90993924, + 2.52087164, -47.13990143, -89.6038701, 15.79930126, -22.88872897, 81.6430658, + 28.3240726, -43.32017928, 31.15072851, -90.9256331, -17.07797375, -88.54243713, + -83.67679143, -78.86701539, 60.11416346, 38.38679261, 86.29202144, 90.51425714, + -72.13954337, -29.90930958, -72.27943372, 49.18231191, -84.50393601, -94.33250534, + -9.93256832, 36.99556838, -24.20862704, -50.44204232, -59.14098804, 30.67322574, + 48.6740379, -63.20731556, 29.52859942, 26.1739035, 47.79243984, -99.38850933, + -83.31147453, 5.84133312, -47.87029605, 95.82254404, -55.52829901, 74.87973108, + -84.05457105, -95.877361, -6.48011262, -78.09036923, 62.14707651, -43.34499838, + 77.42752654, 12.53076343, 97.98997833, -51.38957136, 59.66745814, 65.98475889, + 30.40020778, -49.59550931, 9.93012318, -19.28373689, -78.06236248, 63.68142859, + 79.4625226, 54.24426312, 30.45840289, 70.78206731, -15.30635468, 91.01665772, + -32.76589235, -72.46623074, 58.8632721, -41.35480445, -61.06943341, 81.15104129, + -77.69660768, 95.47462924, -97.46155919, -81.54704047, 84.92285342, -16.08257532, + 52.50986409, 63.78396724, 13.60523945, -63.70301611, -63.10763374, -61.10864955, + 57.26635791, -46.96569014, -43.63636501, -29.30674629}; + std::vector y1{ + -18.05187594, -72.61500308, -23.91931729, 74.04449323, 27.00365642, 5.1319236, + 14.38149555, -44.99859038, 66.15308743, 31.82686363, 60.19621369, 36.0210066, + -18.16429723, 23.06381469, -34.39959102, -80.65093615, -50.66614352, 30.69653939, + -62.06159839, -55.67574679, 2.25702849, 49.26091313, -69.7029038, -14.16800789, + 87.743588, -88.40683092, -78.23312583, 18.95008158, -13.00117829, -88.72165573, + 29.1323603, 0.96433644, -58.14148269, 98.23977047, 87.65596264, -68.42627074, + -61.49539737, 95.22412232, -71.36634131, -87.93028627, -63.70741872, 1.83023166, + -44.18487939, -29.21226692, 36.63070499, 90.55120946, 35.40957933, -53.4846641, + 85.05271776, 80.18938384, -21.31383215, -64.49346601, -72.18007668, 50.73463806, + 7.31981159, -56.54419098, -80.58912509, 6.9148441, -22.67913193, 75.95466325, + 69.60650343, 27.61785095, -17.79886571, -78.36406108, 6.59132839, 64.32222104, + 55.24725933, -53.49018276, -71.57964472, -9.67121623, -29.99957675, -54.15829041, + 29.2535217, 57.8310291, 76.77316186, -54.7557032, 58.71741302, -89.00648352, + -62.5722641, 55.11808159, -72.80219812, 56.12298346, -9.07364408, 87.38574222, + 16.65929972, -91.77505634, -99.47758027, 6.65748231, 19.82536216, -22.91831102, + 30.17048427, 83.6666866, -91.70882742, 78.70726479, 86.04667134, -83.58460595, + 84.27888265, 6.37422824, 62.58260785, -87.64421056}; + + std::vector x2{ + -69.89840831561355, 78.8460456024616, 39.85341596822734, -24.391223974913235, + 13.303395979112231, -12.113621295331923, 65.76955972393912, 32.88000233887396, + 75.15679902070009, 70.42968479275325, -70.48373074669782, -67.41906709787041, + 24.0317463752441, 15.6825064869063, 22.786346338534358, -20.418849974209763, + 34.82105661248487, 38.24867453316148, -25.835471974453984, -99.8181927392706, + 89.84785718125181, 92.62449528299297, -15.692938009982782, 42.32594734729251, + -60.14762773795758, 74.97034158301297, 49.83345296858048, -8.799811548418369, + 35.12809596314472, 93.18344995215058, -94.67426883200939, 52.863378156989384, + 80.55592370229223, -9.708518300250157, 58.19902373613033, 94.71328595396487, + -41.956496383879006, -99.23900353260521, -96.8820547539014, -61.540850851797046, + 10.60351610840815, -86.06663137958869, -19.76183018904282, -52.98140516951296, + -60.77170988936312, -67.64765557651907, 45.61193823583003, 56.92515530750559, + -33.35973933318071, -51.94527984432248, -14.582250347543601, -96.83073470861669, + -47.25698648583708, 48.904375839188006, 14.554162511314495, 38.237373081363344, + -32.7325518620032, 57.537241341535015, -70.50257367880944, -83.11435173667108, + 1.3843207970826832, -61.35647094743536, 43.70708320820875, -81.93488230360825, + -53.098660448910465, 70.16656087048054, 0.7197864636628637, 92.59459361315123, + -77.37226816319428, -32.66885376463454, 34.32370196646004, 71.72963476414482, + 1.5234779242439433, 3.0626652169396085, -1.600973288116736, -1.875116500268692, + 24.115900341387686, -6.818007491235834, -37.57206985691543, 46.48919986671669, + 99.81587509298548, 26.961573147884856, -57.411420876126954, -78.90146907605978, + 37.2322492476274, 67.99231943510561, 64.95985406157519, -21.195261701977287, + 78.89518238318205, -95.50952525706322, 76.75637507677297, -63.30961059551444, + 88.07294705390709, 12.963110252847354, -59.3400766172247, 18.016669829562915, + 0.024732013514316975, -47.68463698812436, -16.12846919710843, 57.85570255646779, + }; + + std::vector y2{ + 96.98573446, -58.67543342, -15.58533008, -14.69764415, 85.96236693, 38.92770099, + 19.79169398, 27.48346165, 53.91447893, 75.83100042, 17.73746514, 51.50105094, + 33.83904611, -9.5980519, 27.56740206, 33.72816966, 48.98821931, -14.86179498, + 0.13287706, 35.05682116, 88.14369171, -20.65562107, -36.15962607, 23.46390886, + 95.9320668, 10.93618875, -76.64604957, -44.27118733, -17.066191, 51.82799017, + -55.47233099, 82.31391458, -99.25207116, -8.96223612, -14.76459615, 35.51101965, + -7.51521537, -12.73466947, -76.18168201, -58.82174033, -64.55998759, -66.29491005, + 96.9048821, -42.97997452, -31.86598156, -96.36343702, -84.20827194, 26.79428452, + 62.9120389, -87.22767369, 11.29343689, -65.44214692, 85.68799019, 61.94678236, + 83.46238187, 21.33376867, 61.87186017, -35.70805035, 68.43167378, -18.40025139, + 25.27768848, -74.94714348, 2.39102813, -78.06742778, 73.16329192, -5.42551355, + -17.11543473, -21.57167168, 60.95981138, -87.3077912, 46.07464277, -26.73518669, + 77.34113841, -10.89097658, -7.48300521, -24.16332469, 66.03877278, 46.51467863, + 86.52324723, 23.88758094, 32.7046036, 47.38730439, -40.72743971, 96.60257607, + -93.12849376, -70.2629721, 94.52718105, 68.27804048, -74.27404657, -21.16650115, + -34.93847764, 66.55335171, -88.44856488, -23.53818607, -29.02780535, -29.34648183, + 74.28318391, -38.37789666, 56.28623834, -81.09317815}; + + std::vector expect{ + 128.64717656028176, 87.88562670763609, 90.19632281028372, 91.76013021796666, + 118.4215357030851, 39.44788631062081, 86.58624490836462, 83.77327247860025, + 79.6690804001798, 167.7366440763836, 83.73027552297903, 99.54006861093508, + 56.276686562837135, 70.80573751073386, 128.34122090714868, 119.97639069191793, + 115.15820154183437, 62.91768450568626, 82.47065566268454, 106.88509910638807, + 104.02822613477268, 196.4153033352887, 33.57186030542483, 136.17156536458378, + 24.91330426477482, 183.12555244130633, 10.402491013960068, 78.8891909881514, + 51.325155608916646, 140.57498906651185, 87.55436962189877, 116.056329846112, + 158.26789618636312, 122.3127143880106, 175.65336769339257, 215.7342613973661, + 62.764576137605516, 173.82450721651924, 72.83278521088664, 31.152704497923047, + 69.74971493014701, 135.16371248533227, 156.8113160405468, 17.14989841904493, + 113.33993969348232, 209.1400727201678, 119.63772368951071, 175.72328059917774, + 54.63868144260578, 177.1094683844075, 46.59751045319419, 192.6556145241176, + 158.08460123131488, 28.29189388239395, 124.5848038188644, 155.08622923365692, + 144.85966618486546, 142.16736573734084, 160.92578604203126, 102.39361006963875, + 88.02052587541618, 126.40767969785988, 57.91601260573419, 30.546751247397513, + 130.95046326396607, 69.87298439379285, 78.21308650467851, 145.7285720718672, + 158.70858501146054, 78.78197209323828, 135.71261679147338, 28.579717281106852, + 91.58009372078648, 85.68704702725647, 90.14934991196503, 78.83501748640491, + 40.09634022929284, 167.14546691120552, 149.17295612658907, 122.98674172809133, + 113.17597316247064, 68.87263271597341, 31.86446035391618, 160.31767244865998, + 158.94024585517036, 34.900531808173085, 253.01889830507722, 86.25213267010419, + 94.2922665997649, 79.44626620532313, 69.47712008431841, 128.24056985459816, + 74.53904203761351, 127.79603731531678, 115.13613538697125, 95.93013225849919, + 58.10781125509778, 44.75789949605465, 28.21929483784659, 87.40828630126103}; + + rmm::device_vector dx1(x1), dy1(y1), dx2(x2), dy2(y2), dexpect(expect); + rmm::device_vector got(dx1.size()); + + auto p1_begin = make_cartesian_2d_iterator(dx1.begin(), dy1.begin()); + auto p2_begin = make_cartesian_2d_iterator(dx2.begin(), dy2.begin()); + + pairwise_point_distance(p1_begin, p1_begin + dx1.size(), p2_begin, got.begin()); + + if constexpr (std::is_same_v) { + EXPECT_THAT(expected, ::testing::Pointwise(::testing::FloatEq(), hgot)); + } else { + EXPECT_THAT(expected, ::testing::Pointwise(::testing::DoubleEq(), hgot)); + } +} + +} // namespace cuspatial From 1757853e2971d220e66d38ebd17c97350b3ebd6d Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 8 Jun 2022 15:16:54 -0700 Subject: [PATCH 26/45] Supply more precision to hard coded test, minor fixes --- .../spatial/point_distance_test.cu | 143 ++++++++++-------- 1 file changed, 83 insertions(+), 60 deletions(-) diff --git a/cpp/tests/experimental/spatial/point_distance_test.cu b/cpp/tests/experimental/spatial/point_distance_test.cu index 3bd5557d4..1872cb830 100644 --- a/cpp/tests/experimental/spatial/point_distance_test.cu +++ b/cpp/tests/experimental/spatial/point_distance_test.cu @@ -145,44 +145,58 @@ TYPED_TEST(PairwisePointDistanceTest, CompareWithShapely) using Cart2D = cartesian_2d; using Cart2DVec = std::vector; - std::size_t constexpr num_points = 100; - std::vector x1{ - -12.30983106, -7.92705956, -49.9570584, -1.05124645, -89.39777526, -32.46014839, - -20.64749623, 74.88373211, -3.56663354, -91.43203925, 1.68283845, 30.90993924, - 2.52087164, -47.13990143, -89.6038701, 15.79930126, -22.88872897, 81.6430658, - 28.3240726, -43.32017928, 31.15072851, -90.9256331, -17.07797375, -88.54243713, - -83.67679143, -78.86701539, 60.11416346, 38.38679261, 86.29202144, 90.51425714, - -72.13954337, -29.90930958, -72.27943372, 49.18231191, -84.50393601, -94.33250534, - -9.93256832, 36.99556838, -24.20862704, -50.44204232, -59.14098804, 30.67322574, - 48.6740379, -63.20731556, 29.52859942, 26.1739035, 47.79243984, -99.38850933, - -83.31147453, 5.84133312, -47.87029605, 95.82254404, -55.52829901, 74.87973108, - -84.05457105, -95.877361, -6.48011262, -78.09036923, 62.14707651, -43.34499838, - 77.42752654, 12.53076343, 97.98997833, -51.38957136, 59.66745814, 65.98475889, - 30.40020778, -49.59550931, 9.93012318, -19.28373689, -78.06236248, 63.68142859, - 79.4625226, 54.24426312, 30.45840289, 70.78206731, -15.30635468, 91.01665772, - -32.76589235, -72.46623074, 58.8632721, -41.35480445, -61.06943341, 81.15104129, - -77.69660768, 95.47462924, -97.46155919, -81.54704047, 84.92285342, -16.08257532, - 52.50986409, 63.78396724, 13.60523945, -63.70301611, -63.10763374, -61.10864955, - 57.26635791, -46.96569014, -43.63636501, -29.30674629}; + -12.309831056315302, -7.927059559371418, -49.95705839647165, -1.0512464476733485, + -89.39777525663895, -32.460148393873666, -20.64749623324501, 74.88373211296442, + -3.566633537053898, -91.4320392524529, 1.68283845329249, 30.90993923507801, + 2.5208716416609267, -47.13990142514067, -89.60387010381702, 15.799301259524867, + -22.8887289692815, 81.6430657985936, 28.324072604115624, -43.3201792789866, + 31.15072850958005, -90.9256331022774, -17.077973750390452, -88.54243712973691, + -83.67679143413889, -78.86701538797912, 60.11416346218348, 38.38679261335849, + 86.29202143733288, 90.51425714428673, -72.13954336543273, -29.909309579787713, + -72.27943372189681, 49.182311914851205, -84.50393600760954, -94.33250533960667, + -9.932568319346647, 36.99556837875937, -24.20862704113279, -50.442042319693705, + -59.14098804172897, 30.673225738449172, 48.67403790478693, -63.207315558126545, + 29.52859942242645, 26.173903500998197, 47.79243983907904, -99.38850933058964, + -83.31147453301942, 5.8413331217636255, -47.87029604603307, 95.82254403897923, + -55.52829900834991, 74.87973107553039, -84.05457104705182, -95.87736100367613, + -6.480112617573386, -78.09036923042659, 62.14707651291427, -43.34499838125344, + 77.42752654240155, 12.530763429172254, 97.98997832862835, -51.389571363066, + 59.66745813871337, 65.98475889051292, 30.40020778235388, -49.595509308751595, + 9.930123176564942, -19.283736893878867, -78.06236247946624, 63.68142858698178, + 79.46252260195803, 54.24426311960122, 30.458402886352822, 70.7820673095687, + -15.306354680748024, 91.01665772140062, -32.765892351019666, -72.46623073604916, + 58.863272100444334, -41.35480445335994, -61.06943341086172, 81.15104128608479, + -77.69660768219927, 95.47462923834442, -97.46155919360085, -81.54704046899158, + 84.9228534190681, -16.082575320922533, 52.509864091786355, 63.78396723518307, + 13.605239448412032, -63.70301611378514, -63.10763374202178, -61.108649551804895, + 57.266357913172385, -46.96569013769979, -43.636365011489566, -29.306746287827558}; std::vector y1{ - -18.05187594, -72.61500308, -23.91931729, 74.04449323, 27.00365642, 5.1319236, - 14.38149555, -44.99859038, 66.15308743, 31.82686363, 60.19621369, 36.0210066, - -18.16429723, 23.06381469, -34.39959102, -80.65093615, -50.66614352, 30.69653939, - -62.06159839, -55.67574679, 2.25702849, 49.26091313, -69.7029038, -14.16800789, - 87.743588, -88.40683092, -78.23312583, 18.95008158, -13.00117829, -88.72165573, - 29.1323603, 0.96433644, -58.14148269, 98.23977047, 87.65596264, -68.42627074, - -61.49539737, 95.22412232, -71.36634131, -87.93028627, -63.70741872, 1.83023166, - -44.18487939, -29.21226692, 36.63070499, 90.55120946, 35.40957933, -53.4846641, - 85.05271776, 80.18938384, -21.31383215, -64.49346601, -72.18007668, 50.73463806, - 7.31981159, -56.54419098, -80.58912509, 6.9148441, -22.67913193, 75.95466325, - 69.60650343, 27.61785095, -17.79886571, -78.36406108, 6.59132839, 64.32222104, - 55.24725933, -53.49018276, -71.57964472, -9.67121623, -29.99957675, -54.15829041, - 29.2535217, 57.8310291, 76.77316186, -54.7557032, 58.71741302, -89.00648352, - -62.5722641, 55.11808159, -72.80219812, 56.12298346, -9.07364408, 87.38574222, - 16.65929972, -91.77505634, -99.47758027, 6.65748231, 19.82536216, -22.91831102, - 30.17048427, 83.6666866, -91.70882742, 78.70726479, 86.04667134, -83.58460595, - 84.27888265, 6.37422824, 62.58260785, -87.64421056}; + -18.051875936208862, -72.61500308351708, -23.919317289360777, 74.04449323147637, + 27.003656419402276, 5.131923603252009, 14.381495553187262, -44.998590378882795, + 66.15308743061799, 31.82686362809011, 60.19621369406618, 36.02100660419922, + -18.164297228344505, 23.06381468579426, -34.39959102364766, -80.65093614662105, + -50.66614351982265, 30.696539385917852, -62.06159838829518, -55.67574678891346, + 2.2570284921564987, 49.260913129155036, -69.70290379728544, -14.168007892316037, + 87.743587998508, -88.40683092249026, -78.23312582934655, 18.950081576813904, + -13.001178290210335, -88.72165572783072, 29.13236030074242, 0.9643364439866353, + -58.14148269328302, 98.23977047259831, 87.65596263514071, -68.42627074347195, + -61.49539737381592, 95.22412232012014, -71.3663413078797, -87.93028627383005, + -63.70741871892348, 1.83023166369769, -44.184879390345245, -29.212266921068498, + 36.63070498793903, 90.55120945758097, 35.40957933073132, -53.484664102448285, + 85.05271776288717, 80.18938384135001, -21.313832146230382, -64.49346600820266, + -72.18007667511924, 50.73463806168728, 7.319811593578507, -56.54419097667299, + -80.58912509276239, 6.9148441008914485, -22.67913193215382, 75.95466324740005, + 69.60650343179027, 27.61785095385285, -17.798865714702472, -78.36406107867042, + 6.59132839160077, 64.32222103875719, 55.24725933014744, -53.49018275541756, + -71.57964472201111, -9.671216230543001, -29.999576747551593, -54.15829040618368, + 29.253521698849028, 57.83102910157538, 76.77316185511351, -54.755703196886174, + 58.71741301597688, -89.00648352439477, -62.572264098389354, 55.118081589496626, + -72.80219811987917, 56.12298345685937, -9.073644079329679, 87.3857422229443, + 16.65929971566098, -91.77505633845232, -99.4775802747735, 6.657482305470497, + 19.82536215719839, -22.918311016363912, 30.170484267010387, 83.6666865961853, + -91.70882742463144, 78.70726479431833, 86.04667133973348, -83.58460594914955, + 84.27888264842167, 6.374228239422575, 62.58260784755962, -87.64421055779096}; std::vector x2{ -69.89840831561355, 78.8460456024616, 39.85341596822734, -24.391223974913235, @@ -209,27 +223,34 @@ TYPED_TEST(PairwisePointDistanceTest, CompareWithShapely) 37.2322492476274, 67.99231943510561, 64.95985406157519, -21.195261701977287, 78.89518238318205, -95.50952525706322, 76.75637507677297, -63.30961059551444, 88.07294705390709, 12.963110252847354, -59.3400766172247, 18.016669829562915, - 0.024732013514316975, -47.68463698812436, -16.12846919710843, 57.85570255646779, - }; + 0.024732013514316975, -47.68463698812436, -16.12846919710843, 57.85570255646779}; std::vector y2{ - 96.98573446, -58.67543342, -15.58533008, -14.69764415, 85.96236693, 38.92770099, - 19.79169398, 27.48346165, 53.91447893, 75.83100042, 17.73746514, 51.50105094, - 33.83904611, -9.5980519, 27.56740206, 33.72816966, 48.98821931, -14.86179498, - 0.13287706, 35.05682116, 88.14369171, -20.65562107, -36.15962607, 23.46390886, - 95.9320668, 10.93618875, -76.64604957, -44.27118733, -17.066191, 51.82799017, - -55.47233099, 82.31391458, -99.25207116, -8.96223612, -14.76459615, 35.51101965, - -7.51521537, -12.73466947, -76.18168201, -58.82174033, -64.55998759, -66.29491005, - 96.9048821, -42.97997452, -31.86598156, -96.36343702, -84.20827194, 26.79428452, - 62.9120389, -87.22767369, 11.29343689, -65.44214692, 85.68799019, 61.94678236, - 83.46238187, 21.33376867, 61.87186017, -35.70805035, 68.43167378, -18.40025139, - 25.27768848, -74.94714348, 2.39102813, -78.06742778, 73.16329192, -5.42551355, - -17.11543473, -21.57167168, 60.95981138, -87.3077912, 46.07464277, -26.73518669, - 77.34113841, -10.89097658, -7.48300521, -24.16332469, 66.03877278, 46.51467863, - 86.52324723, 23.88758094, 32.7046036, 47.38730439, -40.72743971, 96.60257607, - -93.12849376, -70.2629721, 94.52718105, 68.27804048, -74.27404657, -21.16650115, - -34.93847764, 66.55335171, -88.44856488, -23.53818607, -29.02780535, -29.34648183, - 74.28318391, -38.37789666, 56.28623834, -81.09317815}; + 96.98573446222625, -58.675433421313485, -15.58533007526851, -14.697644147821276, + 85.96236693008059, 38.92770099339309, 19.791693980620906, 27.483461653596166, + 53.91447892576453, 75.83100042363395, 17.73746513670771, 51.50105094020323, + 33.83904611309756, -9.59805189545494, 27.567402061211244, 33.72816965802343, + 48.98821930718205, -14.861794980690213, 0.13287706149869294, 35.05682115680253, + 88.14369170856402, -20.655621067301244, -36.15962607484525, 23.463908856814932, + 95.93206680397306, 10.936188747304243, -76.64604957338365, -44.27118733203363, + -17.066191002518682, 51.827990165726675, -55.472330987826744, 82.31391457552668, + -99.25207116240846, -8.9622361202783, -14.764596152666753, 35.51101965248979, + -7.515215371057382, -12.734669471901016, -76.18168200736743, -58.82174033449078, + -64.55998759489724, -66.29491004534883, 96.90488209719925, -42.97997451919843, + -31.865981559056365, -96.36343702487376, -84.20827193890962, 26.79428452012931, + 62.912038904465774, -87.227673692568, 11.2934368901489, -65.442146916886, + 85.68799018964843, 61.94678236143925, 83.46238187197174, 21.333768673112008, + 61.8718601660381, -35.70805034839669, 68.43167377857928, -18.400251392936294, + 25.277688476279536, -74.94714347783905, 2.391028130810602, -78.06742777647494, + 73.16329191776757, -5.425513550228256, -17.11543472509981, -21.571671681683625, + 60.95981137578463, -87.30779120172515, 46.07464276698177, -26.735186694206213, + 77.34113840661823, -10.89097657623882, -7.483005212073712, -24.163324686785494, + 66.03877277717585, 46.514678630068175, 86.52324722682492, 23.88758093704468, + 32.70460360118328, 47.3873043949026, -40.72743971179719, 96.60257606822059, + -93.1284937647867, -70.26297209791194, 94.52718104748459, 68.27804048047095, + -74.27404656785302, -21.16650114972075, -34.93847763736745, 66.55335171298651, + -88.44856487882186, -23.53818606503958, -29.02780534888051, -29.346481830318815, + 74.28318391238213, -38.37789665677865, 56.28623833724116, -81.09317815145866}; std::vector expect{ 128.64717656028176, 87.88562670763609, 90.19632281028372, 91.76013021796666, @@ -258,18 +279,20 @@ TYPED_TEST(PairwisePointDistanceTest, CompareWithShapely) 74.53904203761351, 127.79603731531678, 115.13613538697125, 95.93013225849919, 58.10781125509778, 44.75789949605465, 28.21929483784659, 87.40828630126103}; - rmm::device_vector dx1(x1), dy1(y1), dx2(x2), dy2(y2), dexpect(expect); + rmm::device_vector dx1(x1), dy1(y1), dx2(x2), dy2(y2); rmm::device_vector got(dx1.size()); - auto p1_begin = make_cartesian_2d_iterator(dx1.begin(), dy1.begin()); - auto p2_begin = make_cartesian_2d_iterator(dx2.begin(), dy2.begin()); + auto p1_begin = make_cartesian_2d_iterator(dx1.begin(), dy1.begin()); + auto p2_begin = make_cartesian_2d_iterator(dx2.begin(), dy2.begin()); pairwise_point_distance(p1_begin, p1_begin + dx1.size(), p2_begin, got.begin()); + thrust::host_vector hgot(got); + if constexpr (std::is_same_v) { - EXPECT_THAT(expected, ::testing::Pointwise(::testing::FloatEq(), hgot)); + EXPECT_THAT(expect, ::testing::Pointwise(::testing::FloatEq(), hgot)); } else { - EXPECT_THAT(expected, ::testing::Pointwise(::testing::DoubleEq(), hgot)); + EXPECT_THAT(expect, ::testing::Pointwise(::testing::DoubleEq(), hgot)); } } From 2f4ac2d39bfd5993ec735d2b66a41bee2bf1d2d3 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 8 Jun 2022 15:21:28 -0700 Subject: [PATCH 27/45] Add cudf dependent API test --- cpp/tests/spatial/point_distance_test.cpp | 71 +++++++++++++++++++++++ 1 file changed, 71 insertions(+) create mode 100644 cpp/tests/spatial/point_distance_test.cpp diff --git a/cpp/tests/spatial/point_distance_test.cpp b/cpp/tests/spatial/point_distance_test.cpp new file mode 100644 index 000000000..aca9e6810 --- /dev/null +++ b/cpp/tests/spatial/point_distance_test.cpp @@ -0,0 +1,71 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include + +#include +#include + +#include +#include + +namespace cuspatial { + +using namespace cudf; +using namespace cudf::test; + +template +struct PairwisePointDistanceTest : public ::testing::Test { +}; + +using TestTypes = ::testing::Types; + +TYPED_TEST_CASE(PairwisePointDistanceTest, TestTypes); + +TYPED_TEST(PairwisePointDistanceTest, Empty) +{ + using T = TypeParam; + + auto x1 = fixed_width_column_wrapper{}; + auto y1 = fixed_width_column_wrapper{}; + auto x2 = fixed_width_column_wrapper{}; + auto y2 = fixed_width_column_wrapper{}; + + auto expect = fixed_width_column_wrapper{}; + + auto got = pairwise_point_distance(x1, y1, x2, y2); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expect, *got); +} + +TYPED_TEST(PairwisePointDistanceTest, OnePair) +{ + using T = TypeParam; + + auto x1 = fixed_width_column_wrapper{1.0}; + auto y1 = fixed_width_column_wrapper{1.0}; + auto x2 = fixed_width_column_wrapper{0.0}; + auto y2 = fixed_width_column_wrapper{0.0}; + + auto expect = fixed_width_column_wrapper{std::sqrt(2)}; + + auto got = pairwise_point_distance(x1, y1, x2, y2); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expect, *got); +} + +} // namespace cuspatial From e77a3bf069382fcb9a03779f872ebafd250a06f5 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 8 Jun 2022 15:39:31 -0700 Subject: [PATCH 28/45] documentation updates --- cpp/include/cuspatial/distances/point_distance.hpp | 10 ++++++++++ .../cuspatial/experimental/linestring_distance.cuh | 1 + cpp/include/cuspatial/experimental/point_distance.cuh | 3 ++- cpp/include/doxygen_groups.h | 1 + 4 files changed, 14 insertions(+), 1 deletion(-) diff --git a/cpp/include/cuspatial/distances/point_distance.hpp b/cpp/include/cuspatial/distances/point_distance.hpp index 47b947a50..9b150179d 100644 --- a/cpp/include/cuspatial/distances/point_distance.hpp +++ b/cpp/include/cuspatial/distances/point_distance.hpp @@ -22,7 +22,17 @@ namespace cuspatial { /** + * @ingroup distance * @brief Compute pairwise point to point cartesian distance + * + * Computes cartesian distances between points. + * + * @param points1_x Column of x coordinates to the first point in pair + * @param points1_y Column of y coordinates to the first point in pair + * @param points2_x Column of x coordinates to the second point in pair + * @param points2_y Column of y coordinates to the second point in pair + * @param stream The CUDA stream to use for device memory operations and kernel launches + * @return Column of distances */ std::unique_ptr pairwise_point_distance( cudf::column_view const& points1_x, diff --git a/cpp/include/cuspatial/experimental/linestring_distance.cuh b/cpp/include/cuspatial/experimental/linestring_distance.cuh index 9000e090b..563a0a771 100644 --- a/cpp/include/cuspatial/experimental/linestring_distance.cuh +++ b/cpp/include/cuspatial/experimental/linestring_distance.cuh @@ -21,6 +21,7 @@ namespace cuspatial { /** + * @ingroup distance * @copybrief cuspatial::pairwise_linestring_distance * * The shortest distance between two linestrings is defined as the shortest distance diff --git a/cpp/include/cuspatial/experimental/point_distance.cuh b/cpp/include/cuspatial/experimental/point_distance.cuh index 79d7298b5..cd9a29a18 100644 --- a/cpp/include/cuspatial/experimental/point_distance.cuh +++ b/cpp/include/cuspatial/experimental/point_distance.cuh @@ -21,9 +21,10 @@ namespace cuspatial { /** + * @ingroup distance * @copybrief cuspatial::pairwise_point_distance * - * Computes pairwise distance between points. + * Computes cartesian distances between points. * * @tparam Cart2dItA iterator type for point array of the first point of each pair. Must meet * the requirements of [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. diff --git a/cpp/include/doxygen_groups.h b/cpp/include/doxygen_groups.h index 3672c5bd8..f38e60b27 100644 --- a/cpp/include/doxygen_groups.h +++ b/cpp/include/doxygen_groups.h @@ -47,6 +47,7 @@ * @{ * @brief Distance computation APIs * + * @file point_distance.hpp * @file linestring_distance.hpp * @file hausdorff.hpp * @file haversine.hpp From bbd28555cbacc46e96abde1a40a1c6166ec059c1 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 15 Jun 2022 17:05:02 -0700 Subject: [PATCH 29/45] update docs --- cpp/include/cuspatial/distances/point_distance.hpp | 12 +++++------- .../cuspatial/experimental/point_distance.cuh | 12 ++++-------- 2 files changed, 9 insertions(+), 15 deletions(-) diff --git a/cpp/include/cuspatial/distances/point_distance.hpp b/cpp/include/cuspatial/distances/point_distance.hpp index 9b150179d..15c322bba 100644 --- a/cpp/include/cuspatial/distances/point_distance.hpp +++ b/cpp/include/cuspatial/distances/point_distance.hpp @@ -25,14 +25,12 @@ namespace cuspatial { * @ingroup distance * @brief Compute pairwise point to point cartesian distance * - * Computes cartesian distances between points. - * - * @param points1_x Column of x coordinates to the first point in pair - * @param points1_y Column of y coordinates to the first point in pair - * @param points2_x Column of x coordinates to the second point in pair - * @param points2_y Column of y coordinates to the second point in pair + * @param points1_x Column of x coordinates to the first point in each pair + * @param points1_y Column of y coordinates to the first point in each pair + * @param points2_x Column of x coordinates to the second point in each pair + * @param points2_y Column of y coordinates to the second point in each pair * @param stream The CUDA stream to use for device memory operations and kernel launches - * @return Column of distances + * @return Column of distances between each pair of input points */ std::unique_ptr pairwise_point_distance( cudf::column_view const& points1_x, diff --git a/cpp/include/cuspatial/experimental/point_distance.cuh b/cpp/include/cuspatial/experimental/point_distance.cuh index cd9a29a18..51260df8c 100644 --- a/cpp/include/cuspatial/experimental/point_distance.cuh +++ b/cpp/include/cuspatial/experimental/point_distance.cuh @@ -24,8 +24,6 @@ namespace cuspatial { * @ingroup distance * @copybrief cuspatial::pairwise_point_distance * - * Computes cartesian distances between points. - * * @tparam Cart2dItA iterator type for point array of the first point of each pair. Must meet * the requirements of [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. * @tparam Cart2dItB iterator type for point array of the second point of each pair. Must meet @@ -33,16 +31,14 @@ namespace cuspatial { * @tparam OutputIt iterator type for output array. Must meet the requirements of * [LegacyRandomAccessIterator][LinkLRAI], be mutable and be device-accessible. * - * @param points1_first beginning of range of the point of the first point of each - * pair - * @param points1_last end of range of the point of the first point of each pair - * @param points2_first beginning of range of the point of the second point of each - * pair + * @param points1_first beginning of range of the first point of each pair + * @param points1_last end of range of the first point of each pair + * @param points2_first beginning of range of the second point of each pair * @param distances_first beginning iterator to output * @param stream The CUDA stream to use for device memory operations and kernel launches * @return Output iterator to one past the last element in the output range * - * @pre all input iterators for coordinates must have `cuspatial::cartesian_2d` type. + * @pre all input iterators for coordinates must have a `value_type` of `cuspatial::cartesian_2d`. * @pre all scalar types must be floating point types, and must be the same type for all input * iterators and output iterators. * From 456b34275fff1335fe54b64c1d30bb490a9087c1 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 15 Jun 2022 17:06:33 -0700 Subject: [PATCH 30/45] Remove redundant static assert check --- .../cuspatial/experimental/detail/linestring_distance.cuh | 6 ++---- .../cuspatial/experimental/detail/point_distance.cuh | 6 ++---- 2 files changed, 4 insertions(+), 8 deletions(-) diff --git a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh index 1299e9c88..761201681 100644 --- a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh +++ b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh @@ -241,15 +241,13 @@ void pairwise_linestring_distance(OffsetIterator linestring1_offsets_first, typename std::iterator_traits::value_type>(), "Inputs and output must have floating point value type."); - static_assert(detail::is_same::value_type::value_type, - typename std::iterator_traits::value_type>(), + static_assert(detail::is_same::value_type>(), "Inputs and output must have the same value type."); static_assert(detail::is_same, typename std::iterator_traits::value_type, typename std::iterator_traits::value_type>(), - "Inputs must be cuspatial::cartesian_2d"); + "Input types mismatches or input types are not cuspatial::cartesian_2d"); auto const num_linestring_pairs = thrust::distance(linestring1_offsets_first, linestring1_offsets_last); diff --git a/cpp/include/cuspatial/experimental/detail/point_distance.cuh b/cpp/include/cuspatial/experimental/detail/point_distance.cuh index 8027b070a..cc4d11496 100644 --- a/cpp/include/cuspatial/experimental/detail/point_distance.cuh +++ b/cpp/include/cuspatial/experimental/detail/point_distance.cuh @@ -42,15 +42,13 @@ OutputIt pairwise_point_distance(Cart2dItA points1_first, typename std::iterator_traits::value_type>(), "Inputs and output must be floating point types."); - static_assert(detail::is_same::value_type::value_type, - typename std::iterator_traits::value_type>(), + static_assert(detail::is_same::value_type>(), "Inputs and output must be the same types."); static_assert(detail::is_same, typename std::iterator_traits::value_type, typename std::iterator_traits::value_type>(), - "Inputs must be cuspatial::cartesian_2d"); + "Input types mismatches or input types are not cuspatial::cartesian_2d"); return thrust::transform(rmm::exec_policy(stream), points1_first, From 3853a1d363a473ddab957d2223c221bf5344667b Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 15 Jun 2022 17:07:06 -0700 Subject: [PATCH 31/45] Remove redundant test --- cpp/tests/spatial/point_distance_test.cpp | 16 ---------------- 1 file changed, 16 deletions(-) diff --git a/cpp/tests/spatial/point_distance_test.cpp b/cpp/tests/spatial/point_distance_test.cpp index aca9e6810..dc704d737 100644 --- a/cpp/tests/spatial/point_distance_test.cpp +++ b/cpp/tests/spatial/point_distance_test.cpp @@ -52,20 +52,4 @@ TYPED_TEST(PairwisePointDistanceTest, Empty) CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expect, *got); } -TYPED_TEST(PairwisePointDistanceTest, OnePair) -{ - using T = TypeParam; - - auto x1 = fixed_width_column_wrapper{1.0}; - auto y1 = fixed_width_column_wrapper{1.0}; - auto x2 = fixed_width_column_wrapper{0.0}; - auto y2 = fixed_width_column_wrapper{0.0}; - - auto expect = fixed_width_column_wrapper{std::sqrt(2)}; - - auto got = pairwise_point_distance(x1, y1, x2, y2); - - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expect, *got); -} - } // namespace cuspatial From 4ffc18bf354f3957d674bed63d903284ea32b68e Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 15 Jun 2022 17:08:30 -0700 Subject: [PATCH 32/45] Add column null tests --- cpp/src/spatial/point_distance.cu | 3 +++ 1 file changed, 3 insertions(+) diff --git a/cpp/src/spatial/point_distance.cu b/cpp/src/spatial/point_distance.cu index b0ea5a1fe..224157169 100644 --- a/cpp/src/spatial/point_distance.cu +++ b/cpp/src/spatial/point_distance.cu @@ -82,6 +82,9 @@ std::unique_ptr pairwise_point_distance(cudf::column_view const& p CUSPATIAL_EXPECTS(points1_x.type() == points1_y.type() and points2_x.type() == points2_y.type() and points1_x.type() == points2_x.type(), "The types of point coordinates arrays mismatch."); + CUSPATIAL_EXPECTS(not points1_x.has_nulls() and not points1_y.has_nulls() and + not points2_x.has_nulls() and not points2_y.has_nulls(), + "The coordinate columns cannot have nulls."); if (points1_x.size() == 0) { return cudf::empty_like(points1_x); } From bebede63d4d87e5783f69a889c588f601e5948c7 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Tue, 21 Jun 2022 17:32:54 -0700 Subject: [PATCH 33/45] test returned iterator --- .../spatial/point_distance_test.cu | 21 ++++++++++++------- 1 file changed, 14 insertions(+), 7 deletions(-) diff --git a/cpp/tests/experimental/spatial/point_distance_test.cu b/cpp/tests/experimental/spatial/point_distance_test.cu index 1872cb830..4fd5e61b9 100644 --- a/cpp/tests/experimental/spatial/point_distance_test.cu +++ b/cpp/tests/experimental/spatial/point_distance_test.cu @@ -72,9 +72,11 @@ TYPED_TEST(PairwisePointDistanceTest, Empty) rmm::device_vector expected{}; rmm::device_vector got(points1.size()); - pairwise_point_distance(points1.begin(), points1.end(), points2.begin(), got.begin()); + auto ret_it = + pairwise_point_distance(points1.begin(), points1.end(), points2.begin(), got.begin()); EXPECT_EQ(expected, got); + EXPECT_EQ(expected.size(), std::distance(got.begin(), ret_it)); } TYPED_TEST(PairwisePointDistanceTest, OnePair) @@ -89,9 +91,11 @@ TYPED_TEST(PairwisePointDistanceTest, OnePair) rmm::device_vector expected{std::vector{std::sqrt(T{2.0})}}; rmm::device_vector got(points1.size()); - pairwise_point_distance(points1.begin(), points1.end(), points2.begin(), got.begin()); + auto ret_it = + pairwise_point_distance(points1.begin(), points1.end(), points2.begin(), got.begin()); EXPECT_EQ(expected, got); + EXPECT_EQ(expected.size(), std::distance(got.begin(), ret_it)); } template @@ -129,7 +133,8 @@ TYPED_TEST(PairwisePointDistanceTest, ManyRandom) auto expected = compute_point_distance_host(points1, points2); rmm::device_vector got(points1.size()); - pairwise_point_distance(points1.begin(), points1.end(), points2.begin(), got.begin()); + auto ret_it = + pairwise_point_distance(points1.begin(), points1.end(), points2.begin(), got.begin()); thrust::host_vector hgot(got); if constexpr (std::is_same_v) { @@ -137,6 +142,7 @@ TYPED_TEST(PairwisePointDistanceTest, ManyRandom) } else { EXPECT_THAT(expected, ::testing::Pointwise(::testing::DoubleEq(), hgot)); } + EXPECT_EQ(expected.size(), std::distance(got.begin(), ret_it)); } TYPED_TEST(PairwisePointDistanceTest, CompareWithShapely) @@ -252,7 +258,7 @@ TYPED_TEST(PairwisePointDistanceTest, CompareWithShapely) -88.44856487882186, -23.53818606503958, -29.02780534888051, -29.346481830318815, 74.28318391238213, -38.37789665677865, 56.28623833724116, -81.09317815145866}; - std::vector expect{ + std::vector expected{ 128.64717656028176, 87.88562670763609, 90.19632281028372, 91.76013021796666, 118.4215357030851, 39.44788631062081, 86.58624490836462, 83.77327247860025, 79.6690804001798, 167.7366440763836, 83.73027552297903, 99.54006861093508, @@ -285,15 +291,16 @@ TYPED_TEST(PairwisePointDistanceTest, CompareWithShapely) auto p1_begin = make_cartesian_2d_iterator(dx1.begin(), dy1.begin()); auto p2_begin = make_cartesian_2d_iterator(dx2.begin(), dy2.begin()); - pairwise_point_distance(p1_begin, p1_begin + dx1.size(), p2_begin, got.begin()); + auto ret_it = pairwise_point_distance(p1_begin, p1_begin + dx1.size(), p2_begin, got.begin()); thrust::host_vector hgot(got); if constexpr (std::is_same_v) { - EXPECT_THAT(expect, ::testing::Pointwise(::testing::FloatEq(), hgot)); + EXPECT_THAT(expected, ::testing::Pointwise(::testing::FloatEq(), hgot)); } else { - EXPECT_THAT(expect, ::testing::Pointwise(::testing::DoubleEq(), hgot)); + EXPECT_THAT(expected, ::testing::Pointwise(::testing::DoubleEq(), hgot)); } + EXPECT_EQ(expected.size(), std::distance(got.begin(), ret_it)); } } // namespace cuspatial From 38286470b8d645c476280821a336822a291a02a4 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 23 Jun 2022 15:12:01 -0700 Subject: [PATCH 34/45] updating utility locations --- cpp/include/cuspatial/experimental/detail/point_distance.cuh | 4 ++-- cpp/src/spatial/point_distance.cu | 2 +- cpp/tests/experimental/spatial/point_distance_test.cu | 2 +- 3 files changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/include/cuspatial/experimental/detail/point_distance.cuh b/cpp/include/cuspatial/experimental/detail/point_distance.cuh index cc4d11496..9b5050498 100644 --- a/cpp/include/cuspatial/experimental/detail/point_distance.cuh +++ b/cpp/include/cuspatial/experimental/detail/point_distance.cuh @@ -16,9 +16,9 @@ #pragma once +#include #include -#include -#include +#include #include #include diff --git a/cpp/src/spatial/point_distance.cu b/cpp/src/spatial/point_distance.cu index 224157169..b42d22ae5 100644 --- a/cpp/src/spatial/point_distance.cu +++ b/cpp/src/spatial/point_distance.cu @@ -17,7 +17,7 @@ #include #include #include -#include +#include #include #include diff --git a/cpp/tests/experimental/spatial/point_distance_test.cu b/cpp/tests/experimental/spatial/point_distance_test.cu index 4fd5e61b9..38680d06c 100644 --- a/cpp/tests/experimental/spatial/point_distance_test.cu +++ b/cpp/tests/experimental/spatial/point_distance_test.cu @@ -17,7 +17,7 @@ #include #include #include -#include +#include #include #include From 7009a94415130909691b32f68ef211e6dfd49163 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 28 Jul 2022 08:49:45 -0700 Subject: [PATCH 35/45] Apply docs review comments Co-authored-by: Mark Harris --- cpp/include/cuspatial/distances/point_distance.hpp | 2 +- .../cuspatial/experimental/detail/linestring_distance.cuh | 2 +- cpp/include/cuspatial/experimental/detail/point_distance.cuh | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/cpp/include/cuspatial/distances/point_distance.hpp b/cpp/include/cuspatial/distances/point_distance.hpp index 15c322bba..0d903cfa6 100644 --- a/cpp/include/cuspatial/distances/point_distance.hpp +++ b/cpp/include/cuspatial/distances/point_distance.hpp @@ -23,7 +23,7 @@ namespace cuspatial { /** * @ingroup distance - * @brief Compute pairwise point to point cartesian distance + * @brief Compute pairwise point-to-point Cartesian distance * * @param points1_x Column of x coordinates to the first point in each pair * @param points1_y Column of y coordinates to the first point in each pair diff --git a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh index 4c5e71860..b3882419d 100644 --- a/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh +++ b/cpp/include/cuspatial/experimental/detail/linestring_distance.cuh @@ -158,7 +158,7 @@ void pairwise_linestring_distance(OffsetIterator linestring1_offsets_first, static_assert(detail::is_same, typename std::iterator_traits::value_type, typename std::iterator_traits::value_type>(), - "Input types mismatches or input types are not cuspatial::cartesian_2d"); + "All input types must be cuspatial::cartesian_2d with the same value type"); auto const num_linestring_pairs = thrust::distance(linestring1_offsets_first, linestring1_offsets_last); diff --git a/cpp/include/cuspatial/experimental/detail/point_distance.cuh b/cpp/include/cuspatial/experimental/detail/point_distance.cuh index 9b5050498..ce9d03819 100644 --- a/cpp/include/cuspatial/experimental/detail/point_distance.cuh +++ b/cpp/include/cuspatial/experimental/detail/point_distance.cuh @@ -48,7 +48,7 @@ OutputIt pairwise_point_distance(Cart2dItA points1_first, static_assert(detail::is_same, typename std::iterator_traits::value_type, typename std::iterator_traits::value_type>(), - "Input types mismatches or input types are not cuspatial::cartesian_2d"); + "All Input types must be cuspatial::cartesian_2d with the same value type"); return thrust::transform(rmm::exec_policy(stream), points1_first, From 412526bea940b53f344d5387ef73f211333a3e6d Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 28 Jul 2022 08:54:11 -0700 Subject: [PATCH 36/45] Renamed `distances` -> `distance` --- .../cuspatial/{distances => distance}/linestring_distance.hpp | 0 cpp/include/cuspatial/{distances => distance}/point_distance.hpp | 0 2 files changed, 0 insertions(+), 0 deletions(-) rename cpp/include/cuspatial/{distances => distance}/linestring_distance.hpp (100%) rename cpp/include/cuspatial/{distances => distance}/point_distance.hpp (100%) diff --git a/cpp/include/cuspatial/distances/linestring_distance.hpp b/cpp/include/cuspatial/distance/linestring_distance.hpp similarity index 100% rename from cpp/include/cuspatial/distances/linestring_distance.hpp rename to cpp/include/cuspatial/distance/linestring_distance.hpp diff --git a/cpp/include/cuspatial/distances/point_distance.hpp b/cpp/include/cuspatial/distance/point_distance.hpp similarity index 100% rename from cpp/include/cuspatial/distances/point_distance.hpp rename to cpp/include/cuspatial/distance/point_distance.hpp From 8f974d61b10548cd7208476421a26de44954432e Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 28 Jul 2022 10:51:35 -0700 Subject: [PATCH 37/45] Fix include paths after `distance` rename --- cpp/tests/spatial/linestring_distance_test.cpp | 2 +- cpp/tests/spatial/linestring_distance_test_medium.cpp | 2 +- cpp/tests/spatial/point_distance_test.cpp | 4 ++-- 3 files changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/tests/spatial/linestring_distance_test.cpp b/cpp/tests/spatial/linestring_distance_test.cpp index 87d77e6e9..80edf1b6e 100644 --- a/cpp/tests/spatial/linestring_distance_test.cpp +++ b/cpp/tests/spatial/linestring_distance_test.cpp @@ -14,7 +14,7 @@ * limitations under the License. */ -#include +#include #include #include diff --git a/cpp/tests/spatial/linestring_distance_test_medium.cpp b/cpp/tests/spatial/linestring_distance_test_medium.cpp index d99433f1d..0f0752d3d 100644 --- a/cpp/tests/spatial/linestring_distance_test_medium.cpp +++ b/cpp/tests/spatial/linestring_distance_test_medium.cpp @@ -14,7 +14,7 @@ * limitations under the License. */ -#include +#include #include #include diff --git a/cpp/tests/spatial/point_distance_test.cpp b/cpp/tests/spatial/point_distance_test.cpp index dc704d737..f79af2638 100644 --- a/cpp/tests/spatial/point_distance_test.cpp +++ b/cpp/tests/spatial/point_distance_test.cpp @@ -14,8 +14,8 @@ * limitations under the License. */ -#include -#include +#include +#include #include #include From 7e1afdcd177b9498c67da6267be6bff70bcd17b7 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 28 Jul 2022 10:51:55 -0700 Subject: [PATCH 38/45] Include cudf API test --- cpp/tests/CMakeLists.txt | 3 +++ 1 file changed, 3 insertions(+) diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 6f28ca0c3..dab1ee9b8 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -77,6 +77,9 @@ ConfigureTest(POLYLINE_BBOX_TEST ConfigureTest(POLYGON_BBOX_TEST spatial/polygon_bbox_test.cu) +ConfigureTest(POINT_DISTANCE_TEST + spatial/point_distance_test.cpp) + ConfigureTest(LINESTRING_DISTANCE_TEST spatial/linestring_distance_test.cpp spatial/linestring_distance_test_medium.cpp) From e8e74978db57a4b6eb7afeed92d2b6bc19674dbc Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 28 Jul 2022 10:52:10 -0700 Subject: [PATCH 39/45] Fix function signature --- cpp/include/cuspatial/distance/point_distance.hpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/cpp/include/cuspatial/distance/point_distance.hpp b/cpp/include/cuspatial/distance/point_distance.hpp index 0d903cfa6..51d506d62 100644 --- a/cpp/include/cuspatial/distance/point_distance.hpp +++ b/cpp/include/cuspatial/distance/point_distance.hpp @@ -17,7 +17,6 @@ #pragma once #include -#include namespace cuspatial { @@ -37,6 +36,6 @@ std::unique_ptr pairwise_point_distance( cudf::column_view const& points1_y, cudf::column_view const& points2_x, cudf::column_view const& points2_y, - rmm::cuda_stream_view stream = rmm::cuda_stream_default); + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); } // namespace cuspatial From f5ca35cbf9fc9ea4c6f53e51a698ea705c3822d6 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 28 Jul 2022 10:55:46 -0700 Subject: [PATCH 40/45] relocate haversine and hausdorff distance --- cpp/benchmarks/hausdorff_benchmark.cpp | 2 +- cpp/include/cuspatial/{ => distance}/hausdorff.hpp | 0 cpp/include/cuspatial/{ => distance}/haversine.hpp | 0 cpp/tests/spatial/hausdorff_test.cpp | 2 +- cpp/tests/spatial/haversine_test.cpp | 2 +- 5 files changed, 3 insertions(+), 3 deletions(-) rename cpp/include/cuspatial/{ => distance}/hausdorff.hpp (100%) rename cpp/include/cuspatial/{ => distance}/haversine.hpp (100%) diff --git a/cpp/benchmarks/hausdorff_benchmark.cpp b/cpp/benchmarks/hausdorff_benchmark.cpp index 6f59f042d..8ff80c331 100644 --- a/cpp/benchmarks/hausdorff_benchmark.cpp +++ b/cpp/benchmarks/hausdorff_benchmark.cpp @@ -14,7 +14,7 @@ * limitations under the License. */ -#include +#include #include #include diff --git a/cpp/include/cuspatial/hausdorff.hpp b/cpp/include/cuspatial/distance/hausdorff.hpp similarity index 100% rename from cpp/include/cuspatial/hausdorff.hpp rename to cpp/include/cuspatial/distance/hausdorff.hpp diff --git a/cpp/include/cuspatial/haversine.hpp b/cpp/include/cuspatial/distance/haversine.hpp similarity index 100% rename from cpp/include/cuspatial/haversine.hpp rename to cpp/include/cuspatial/distance/haversine.hpp diff --git a/cpp/tests/spatial/hausdorff_test.cpp b/cpp/tests/spatial/hausdorff_test.cpp index 665619285..e0e782726 100644 --- a/cpp/tests/spatial/hausdorff_test.cpp +++ b/cpp/tests/spatial/hausdorff_test.cpp @@ -16,7 +16,7 @@ #include #include -#include +#include #include #include diff --git a/cpp/tests/spatial/haversine_test.cpp b/cpp/tests/spatial/haversine_test.cpp index 80fa38d1b..9405a528a 100644 --- a/cpp/tests/spatial/haversine_test.cpp +++ b/cpp/tests/spatial/haversine_test.cpp @@ -15,7 +15,7 @@ */ #include -#include +#include #include #include From 050092412f357835a7ef807d61e3abd0f8a22119 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 28 Jul 2022 11:20:02 -0700 Subject: [PATCH 41/45] style --- cpp/tests/spatial/hausdorff_test.cpp | 2 +- cpp/tests/spatial/haversine_test.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/tests/spatial/hausdorff_test.cpp b/cpp/tests/spatial/hausdorff_test.cpp index e0e782726..9a165a652 100644 --- a/cpp/tests/spatial/hausdorff_test.cpp +++ b/cpp/tests/spatial/hausdorff_test.cpp @@ -15,8 +15,8 @@ */ #include -#include #include +#include #include #include diff --git a/cpp/tests/spatial/haversine_test.cpp b/cpp/tests/spatial/haversine_test.cpp index 9405a528a..373dea484 100644 --- a/cpp/tests/spatial/haversine_test.cpp +++ b/cpp/tests/spatial/haversine_test.cpp @@ -14,8 +14,8 @@ * limitations under the License. */ -#include #include +#include #include #include From 285f35a6f8280060a1c0c6827e668fd1e1868eaf Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 28 Jul 2022 15:24:48 -0700 Subject: [PATCH 42/45] Update docs --- cpp/include/cuspatial/distance/point_distance.hpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/include/cuspatial/distance/point_distance.hpp b/cpp/include/cuspatial/distance/point_distance.hpp index 51d506d62..82c6fc520 100644 --- a/cpp/include/cuspatial/distance/point_distance.hpp +++ b/cpp/include/cuspatial/distance/point_distance.hpp @@ -24,10 +24,10 @@ namespace cuspatial { * @ingroup distance * @brief Compute pairwise point-to-point Cartesian distance * - * @param points1_x Column of x coordinates to the first point in each pair - * @param points1_y Column of y coordinates to the first point in each pair - * @param points2_x Column of x coordinates to the second point in each pair - * @param points2_y Column of y coordinates to the second point in each pair + * @param points1_x Column of x-coordinates of the first point in each pair + * @param points1_y Column of y-coordinates of the first point in each pair + * @param points2_x Column of x-coordinates of the second point in each pair + * @param points2_y Column of y-coordinates of the second point in each pair * @param stream The CUDA stream to use for device memory operations and kernel launches * @return Column of distances between each pair of input points */ From 3c6b2af91965a85482bdf414bfb5520f87afeb20 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 28 Jul 2022 17:59:01 -0700 Subject: [PATCH 43/45] Add throw tests --- cpp/tests/spatial/point_distance_test.cpp | 64 +++++++++++++++++++++++ 1 file changed, 64 insertions(+) diff --git a/cpp/tests/spatial/point_distance_test.cpp b/cpp/tests/spatial/point_distance_test.cpp index f79af2638..76588a4a1 100644 --- a/cpp/tests/spatial/point_distance_test.cpp +++ b/cpp/tests/spatial/point_distance_test.cpp @@ -15,6 +15,7 @@ */ #include +#include #include #include @@ -52,4 +53,67 @@ TYPED_TEST(PairwisePointDistanceTest, Empty) CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expect, *got); } +struct PairwisePointDistanceTestThrow : public ::testing::Test { +}; + +TEST_F(PairwisePointDistanceTestThrow, SizeMismatch) +{ + auto x1 = fixed_width_column_wrapper{1, 2, 3}; + auto y1 = fixed_width_column_wrapper{1, 2, 3}; + auto x2 = fixed_width_column_wrapper{}; + auto y2 = fixed_width_column_wrapper{}; + + auto expect = fixed_width_column_wrapper{}; + + EXPECT_THROW(pairwise_point_distance(x1, y1, x2, y2), cuspatial::logic_error); +} + +TEST_F(PairwisePointDistanceTestThrow, SizeMismatch2) +{ + auto x1 = fixed_width_column_wrapper{}; + auto y1 = fixed_width_column_wrapper{1, 2, 3}; + auto x2 = fixed_width_column_wrapper{1, 2, 3}; + auto y2 = fixed_width_column_wrapper{1, 2, 3}; + + auto expect = fixed_width_column_wrapper{}; + + EXPECT_THROW(pairwise_point_distance(x1, y1, x2, y2), cuspatial::logic_error); +} + +TEST_F(PairwisePointDistanceTestThrow, TypeMismatch) +{ + auto x1 = fixed_width_column_wrapper{1, 2, 3}; + auto y1 = fixed_width_column_wrapper{1, 2, 3}; + auto x2 = fixed_width_column_wrapper{1, 2, 3}; + auto y2 = fixed_width_column_wrapper{1, 2, 3}; + + auto expect = fixed_width_column_wrapper{}; + + EXPECT_THROW(pairwise_point_distance(x1, y1, x2, y2), cuspatial::logic_error); +} + +TEST_F(PairwisePointDistanceTestThrow, TypeMismatch2) +{ + auto x1 = fixed_width_column_wrapper{1, 2, 3}; + auto y1 = fixed_width_column_wrapper{1, 2, 3}; + auto x2 = fixed_width_column_wrapper{1, 2, 3}; + auto y2 = fixed_width_column_wrapper{1, 2, 3}; + + auto expect = fixed_width_column_wrapper{}; + + EXPECT_THROW(pairwise_point_distance(x1, y1, x2, y2), cuspatial::logic_error); +} + +TEST_F(PairwisePointDistanceTestThrow, ContainsNull) +{ + auto x1 = fixed_width_column_wrapper{{1, 2, 3}, {1, 0, 1}}; + auto y1 = fixed_width_column_wrapper{1, 2, 3}; + auto x2 = fixed_width_column_wrapper{1, 2, 3}; + auto y2 = fixed_width_column_wrapper{1, 2, 3}; + + auto expect = fixed_width_column_wrapper{}; + + EXPECT_THROW(pairwise_point_distance(x1, y1, x2, y2), cuspatial::logic_error); +} + } // namespace cuspatial From 31638c79ca0b8a9b6d2d01a8880106a3b40ed754 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 28 Jul 2022 18:07:57 -0700 Subject: [PATCH 44/45] Fix cython builds --- python/cuspatial/cuspatial/_lib/cpp/distance/__init__.pxd | 0 python/cuspatial/cuspatial/_lib/cpp/distance/__init__.pyx | 0 .../cuspatial/cuspatial/_lib/cpp/{ => distance}/hausdorff.pxd | 3 ++- .../cuspatial/_lib/cpp/{spatial.pxd => distance/haversine.pxd} | 3 ++- .../cuspatial/_lib/cpp/{ => distance}/linestring_distance.pxd | 2 +- python/cuspatial/cuspatial/_lib/hausdorff.pyx | 2 +- python/cuspatial/cuspatial/_lib/linestring_distance.pyx | 2 +- python/cuspatial/cuspatial/_lib/spatial.pyx | 2 +- 8 files changed, 8 insertions(+), 6 deletions(-) create mode 100644 python/cuspatial/cuspatial/_lib/cpp/distance/__init__.pxd create mode 100644 python/cuspatial/cuspatial/_lib/cpp/distance/__init__.pyx rename python/cuspatial/cuspatial/_lib/cpp/{ => distance}/hausdorff.pxd (80%) rename python/cuspatial/cuspatial/_lib/cpp/{spatial.pxd => distance/haversine.pxd} (78%) rename python/cuspatial/cuspatial/_lib/cpp/{ => distance}/linestring_distance.pxd (91%) diff --git a/python/cuspatial/cuspatial/_lib/cpp/distance/__init__.pxd b/python/cuspatial/cuspatial/_lib/cpp/distance/__init__.pxd new file mode 100644 index 000000000..e69de29bb diff --git a/python/cuspatial/cuspatial/_lib/cpp/distance/__init__.pyx b/python/cuspatial/cuspatial/_lib/cpp/distance/__init__.pyx new file mode 100644 index 000000000..e69de29bb diff --git a/python/cuspatial/cuspatial/_lib/cpp/hausdorff.pxd b/python/cuspatial/cuspatial/_lib/cpp/distance/hausdorff.pxd similarity index 80% rename from python/cuspatial/cuspatial/_lib/cpp/hausdorff.pxd rename to python/cuspatial/cuspatial/_lib/cpp/distance/hausdorff.pxd index 90fd80da1..6f5fc69e3 100644 --- a/python/cuspatial/cuspatial/_lib/cpp/hausdorff.pxd +++ b/python/cuspatial/cuspatial/_lib/cpp/distance/hausdorff.pxd @@ -6,7 +6,8 @@ from cudf._lib.cpp.column.column cimport column from cudf._lib.cpp.column.column_view cimport column_view -cdef extern from "cuspatial/hausdorff.hpp" namespace "cuspatial" nogil: +cdef extern from "cuspatial/distance/hausdorff.hpp" \ + namespace "cuspatial" nogil: cdef unique_ptr[column] directed_hausdorff_distance( const column_view& xs, diff --git a/python/cuspatial/cuspatial/_lib/cpp/spatial.pxd b/python/cuspatial/cuspatial/_lib/cpp/distance/haversine.pxd similarity index 78% rename from python/cuspatial/cuspatial/_lib/cpp/spatial.pxd rename to python/cuspatial/cuspatial/_lib/cpp/distance/haversine.pxd index 764458f05..4e39c6ad5 100644 --- a/python/cuspatial/cuspatial/_lib/cpp/spatial.pxd +++ b/python/cuspatial/cuspatial/_lib/cpp/distance/haversine.pxd @@ -5,7 +5,8 @@ from libcpp.memory cimport unique_ptr from cudf._lib.column cimport column, column_view -cdef extern from "cuspatial/haversine.hpp" namespace "cuspatial" nogil: +cdef extern from "cuspatial/distance/haversine.hpp" \ + namespace "cuspatial" nogil: cdef unique_ptr[column] haversine_distance( const column_view& a_lon, const column_view& a_lat, diff --git a/python/cuspatial/cuspatial/_lib/cpp/linestring_distance.pxd b/python/cuspatial/cuspatial/_lib/cpp/distance/linestring_distance.pxd similarity index 91% rename from python/cuspatial/cuspatial/_lib/cpp/linestring_distance.pxd rename to python/cuspatial/cuspatial/_lib/cpp/distance/linestring_distance.pxd index 4523cf045..0a0b1a6b9 100644 --- a/python/cuspatial/cuspatial/_lib/cpp/linestring_distance.pxd +++ b/python/cuspatial/cuspatial/_lib/cpp/distance/linestring_distance.pxd @@ -8,7 +8,7 @@ from cudf._lib.cpp.column.column_view cimport column_view from cudf._lib.cpp.types cimport size_type -cdef extern from "cuspatial/distances/linestring_distance.hpp" \ +cdef extern from "cuspatial/distance/linestring_distance.hpp" \ namespace "cuspatial" nogil: cdef unique_ptr[column] pairwise_linestring_distance( const column_view linestring1_offsets, diff --git a/python/cuspatial/cuspatial/_lib/hausdorff.pyx b/python/cuspatial/cuspatial/_lib/hausdorff.pyx index ca9497813..7465efc7a 100644 --- a/python/cuspatial/cuspatial/_lib/hausdorff.pyx +++ b/python/cuspatial/cuspatial/_lib/hausdorff.pyx @@ -5,7 +5,7 @@ from libcpp.utility cimport move from cudf._lib.column cimport Column, column, column_view -from cuspatial._lib.cpp.hausdorff cimport ( +from cuspatial._lib.cpp.distance.hausdorff cimport ( directed_hausdorff_distance as directed_cpp_hausdorff_distance, ) diff --git a/python/cuspatial/cuspatial/_lib/linestring_distance.pyx b/python/cuspatial/cuspatial/_lib/linestring_distance.pyx index 36589810b..49ed7beeb 100644 --- a/python/cuspatial/cuspatial/_lib/linestring_distance.pyx +++ b/python/cuspatial/cuspatial/_lib/linestring_distance.pyx @@ -5,7 +5,7 @@ from cudf._lib.column cimport Column from cudf._lib.cpp.column.column cimport column from cudf._lib.cpp.column.column_view cimport column_view -from cuspatial._lib.cpp.linestring_distance cimport ( +from cuspatial._lib.cpp.distance.linestring_distance cimport ( pairwise_linestring_distance as cpp_pairwise_linestring_distance, ) diff --git a/python/cuspatial/cuspatial/_lib/spatial.pyx b/python/cuspatial/cuspatial/_lib/spatial.pyx index f18eaadd1..06a64014c 100644 --- a/python/cuspatial/cuspatial/_lib/spatial.pyx +++ b/python/cuspatial/cuspatial/_lib/spatial.pyx @@ -13,7 +13,7 @@ from cudf._lib.cpp.column.column_view cimport column_view from cuspatial._lib.cpp.coordinate_transform cimport ( lonlat_to_cartesian as cpp_lonlat_to_cartesian, ) -from cuspatial._lib.cpp.spatial cimport ( +from cuspatial._lib.cpp.distance.haversine cimport ( haversine_distance as cpp_haversine_distance, ) From 544d70861d661484205c8437525960e2f729edac Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 28 Jul 2022 18:11:48 -0700 Subject: [PATCH 45/45] Apply review comments --- cpp/tests/experimental/spatial/point_distance_test.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/tests/experimental/spatial/point_distance_test.cu b/cpp/tests/experimental/spatial/point_distance_test.cu index 38680d06c..c80efcfcc 100644 --- a/cpp/tests/experimental/spatial/point_distance_test.cu +++ b/cpp/tests/experimental/spatial/point_distance_test.cu @@ -50,7 +50,7 @@ auto compute_point_distance_host(Cart2DVec const& point1, Cart2DVec const& point auto result_iter = thrust::make_transform_iterator(pair_iter, [](auto p) { auto p0 = thrust::get<0>(p); auto p1 = thrust::get<1>(p); - return std::sqrt((p0.x - p1.x) * (p0.x - p1.x) + (p0.y - p1.y) * (p0.y - p1.y)); + return std::sqrt(dot(p0 - p1, p0 - p1)); }); return thrust::host_vector(result_iter, result_iter + point1.size());