-
Notifications
You must be signed in to change notification settings - Fork 157
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Refactor haversine_distance to a header-only API #477
Changes from 8 commits
314580d
0736356
5830b66
4c16cd6
7580661
073e2d7
e37d61e
9d8e3eb
7f2bbad
2448677
0d954e0
7d100dc
daa82a8
c4ba1f7
454d967
a5dab4a
08dfe95
e373065
3ae133e
f8947eb
564cc4c
4a6976e
c208144
21db279
be1226c
8134f12
42f909e
e5cb703
0b78d87
17569c6
d5ef3b7
File filter
Filter by extension
Conversations
Jump to
Diff view
Diff view
There are no files selected for viewing
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,111 @@ | ||
/* | ||
* 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 "cuspatial/types.hpp" | ||
harrism marked this conversation as resolved.
Show resolved
Hide resolved
|
||
#include <cuspatial/constants.hpp> | ||
#include <cuspatial/error.hpp> | ||
|
||
#include <rmm/cuda_stream_view.hpp> | ||
#include <rmm/exec_policy.hpp> | ||
#include <rmm/mr/device/device_memory_resource.hpp> | ||
#include <rmm/mr/device/per_device_resource.hpp> | ||
|
||
#include <thrust/iterator/constant_iterator.h> | ||
#include <thrust/iterator/zip_iterator.h> | ||
#include <thrust/transform.h> | ||
#include <thrust/tuple.h> | ||
|
||
#include <type_traits> | ||
|
||
namespace cuspatial { | ||
|
||
namespace detail { | ||
|
||
template <typename T> | ||
struct haversine_distance_functor { | ||
haversine_distance_functor(T radius) : radius_(radius) {} | ||
|
||
__device__ T operator()(location_2d<T> a_lonlat, location_2d<T> b_lonlat) | ||
{ | ||
auto ax = a_lonlat.longitude * DEGREE_TO_RADIAN; | ||
auto ay = a_lonlat.latitude * DEGREE_TO_RADIAN; | ||
auto bx = b_lonlat.longitude * DEGREE_TO_RADIAN; | ||
auto by = b_lonlat.latitude * DEGREE_TO_RADIAN; | ||
|
||
// haversine formula | ||
auto x = (bx - ax) / 2; | ||
auto y = (by - ay) / 2; | ||
auto sinysqrd = sin(y) * sin(y); | ||
auto sinxsqrd = sin(x) * sin(x); | ||
auto scale = cos(ay) * cos(by); | ||
|
||
return 2 * radius_ * asin(sqrt(sinysqrd + sinxsqrd * scale)); | ||
} | ||
|
||
T radius_{}; | ||
}; | ||
|
||
template <class LonLatItA, | ||
class LonLatItB, | ||
class OutputIt, | ||
class Location = typename std::iterator_traits<LonLatItA>::value_type, | ||
class T = typename Location::value_type> | ||
OutputIt haversine_distance(LonLatItA a_lonlat_first, | ||
LonLatItA a_lonlat_last, | ||
LonLatItB b_lonlat_first, | ||
OutputIt distance_first, | ||
T const radius = EARTH_RADIUS_KM, | ||
rmm::cuda_stream_view stream = rmm::cuda_stream_default) | ||
{ | ||
using LocationB = typename std::iterator_traits<LonLatItB>::value_type; | ||
static_assert(std::conjunction_v<std::is_same<location_2d<T>, Location>, | ||
std::is_same<location_2d<T>, LocationB>>, | ||
"Inputs must be cuspatial::location_2d"); | ||
static_assert( | ||
std::conjunction_v<std::is_floating_point<T>, | ||
std::is_floating_point<typename LocationB::value_type>, | ||
std::is_floating_point<typename std::iterator_traits<OutputIt>::value_type>>, | ||
"Haversine distance supports only floating-point coordinates."); | ||
|
||
CUSPATIAL_EXPECTS(radius > 0, "radius must be positive."); | ||
|
||
return thrust::transform(rmm::exec_policy(stream), | ||
a_lonlat_first, | ||
a_lonlat_last, | ||
b_lonlat_first, | ||
distance_first, | ||
haversine_distance_functor<T>(radius)); | ||
} | ||
} // namespace detail | ||
|
||
template <class LonLatItA, class LonLatItB, class OutputIt, class Location, class T> | ||
OutputIt haversine_distance(LonLatItA a_lonlat_first, | ||
LonLatItA a_lonlat_last, | ||
LonLatItB b_lonlat_first, | ||
OutputIt distance_first, | ||
T const radius) | ||
{ | ||
return detail::haversine_distance(a_lonlat_first, | ||
a_lonlat_last, | ||
b_lonlat_first, | ||
distance_first, | ||
radius, | ||
rmm::cuda_stream_default); | ||
} | ||
|
||
} // namespace cuspatial |
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,74 @@ | ||
/* | ||
* 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 <cuspatial/constants.hpp> | ||
|
||
#include <rmm/mr/device/device_memory_resource.hpp> | ||
#include <rmm/mr/device/per_device_resource.hpp> | ||
|
||
namespace cuspatial { | ||
|
||
/** | ||
* brief Compute haversine distances between points in set A to the corresponding points in set B. | ||
* | ||
* Computes N haversine distances, where N is `std::distance(a_lon_first, a_lon_last)`. | ||
* The distance for each `(a_lon[i], a_lat[i])` and `(b_lon[i], b_lat[i])` point pair is assigned to | ||
* `distance_first[i]`. `distance_first` must be an iterator to output storage allocated for N | ||
* distances. | ||
* | ||
* https://en.wikipedia.org/wiki/Haversine_formula | ||
* | ||
* @param[in] a_lonlat_first: beginning of range of (longitude, latitude) locations in set A | ||
* @param[in] a_lonlat_last: end of range of (longitude, latitude) locations in set A | ||
* @param[in] b_lonlat_first: beginning of range of (longitude, latitude) locations in set B | ||
* @param[out] distance_first: beginning of output range of haversine distances | ||
* @param[in] radius: radius of the sphere on which the points reside. default: 6371.0 | ||
* (approximate radius of Earth in km) | ||
* | ||
* All iterators must have the same floating-point `value_type`. | ||
* | ||
* Computed distances will have the same units as `radius`. | ||
* | ||
* @tparam LonLatItA Iterator to input location set A. Must meet the requirements of | ||
* [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. | ||
* @tparam LonLatItB Iterator to input location set B. Must meet the requirements of | ||
* [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. | ||
* @tparam OutputIt Output iterator. Must meet the requirements of | ||
* [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible. | ||
* @tparam Location The `value_type` of `LonLatItA` and `LonLatItB`. Must be | ||
* `cuspatial::location_2d<T>`. | ||
* @tparam T The underlying coordinate type. Must be a floating-point type. | ||
* | ||
* @return Output iterator to the element past the last distance computed. | ||
harrism marked this conversation as resolved.
Show resolved
Hide resolved
|
||
* | ||
* [LinkLRAI]: https://en.cppreference.com/w/cpp/named_req/RandomAccessIterator | ||
* "LegacyRandomAccessIterator" | ||
*/ | ||
template <class LonLatItA, | ||
class LonLatItB, | ||
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. The top level API seems to have asserted that A and B coordinates must have the same type. The transform iterator generated from these two arrays should have the same type as well? There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. I'm not sure I follow your question. Can you elaborate? There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. There's not necessarily a transform iterator. And we do assert that they have the same type. See the implementation file, where it has
It uses There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Thanks, I think I understand this a bit better now. We don't enforce that the iterator on coordinate A to be the same type as coordinate B. As long as they are all iterating on |
||
class OutputIt, | ||
class Location = typename std::iterator_traits<LonLatItA>::value_type, | ||
class T = typename Location::value_type> | ||
OutputIt haversine_distance(LonLatItA a_lonlat_first, | ||
LonLatItA a_lonlat_last, | ||
LonLatItB b_lonlat_first, | ||
OutputIt distance_first, | ||
T const radius = EARTH_RADIUS_KM); | ||
} // namespace cuspatial | ||
|
||
#include <cuspatial/experimental/detail/haversine.hpp> |
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,49 @@ | ||
/* | ||
* 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 <cuspatial/types.hpp> | ||
|
||
#include <thrust/iterator/transform_iterator.h> | ||
#include <thrust/iterator/zip_iterator.h> | ||
|
||
namespace cuspatial { | ||
|
||
namespace { | ||
harrism marked this conversation as resolved.
Show resolved
Hide resolved
|
||
template <typename T> | ||
struct tuple_to_location_2d { | ||
// using result_type = cuspatial::location_2d<T>; | ||
|
||
CUSPATIAL_HOST_DEVICE cuspatial::location_2d<T> operator()(thrust::tuple<T, T> lonlat) | ||
{ | ||
return cuspatial::location_2d<T>{thrust::get<0>(lonlat), thrust::get<1>(lonlat)}; | ||
} | ||
}; | ||
|
||
} // anonymous namespace | ||
|
||
// convert two iterators to an iterator<location_2d<T> | ||
template <typename LonIter, typename LatIter> | ||
auto to_location_2d(LonIter lon, LatIter lat) | ||
{ | ||
using T = typename std::iterator_traits<LonIter>::value_type; | ||
static_assert(std::is_same_v<T, typename std::iterator_traits<LatIter>::value_type>, | ||
"Iterator value_type mismatch"); | ||
|
||
auto zipped = thrust::make_zip_iterator(thrust::make_tuple(lon, lat)); | ||
return thrust::make_transform_iterator(zipped, tuple_to_location_2d<T>()); | ||
} | ||
|
||
} // namespace cuspatial |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Shouldn't this be named
haversine.cuh
?There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Depends on your style. Thrust doesn't follow that convention.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Actually, in the future we may want to enable both host and device execution of these algorithms, a la Thrust.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I thought anything with CUDA C++ in it was supposed to either be
.cu
or.cuh
? If someone includes this from a.cpp
file, GCC won't know what to do with the__device__
functions?There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
True. Thrust uses .h for all headers, and makes sure to protect everything so that it can be portable to different backends. I don't know if we want to go that direction. I just have a particular aversion to ".cuh" used in non-detail headers (not sure why, it's just ugly I guess).
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Generally my rule is that
.h/.hpp
files should be able to included in a host-only TU without issue.I'd go with
.cuh
unless we decide to make cuSpatial work on both CPU/GPU via Thrust.There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Done. It would be nice to support CPU and GPU backends but I'm not sure how feasible that is with the more complicated algorithms. Also it would have a large impact on tests, which currently use device vectors. So this is definitely future work.