-
Notifications
You must be signed in to change notification settings - Fork 154
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
Merged
rapids-bot
merged 31 commits into
rapidsai:branch-22.06
from
harrism:fea-header-only-haversine
Apr 28, 2022
Merged
Changes from 3 commits
Commits
Show all changes
31 commits
Select commit
Hold shift + click to select a range
314580d
Create header-only refactoring of cuspatial::haversine_distance
harrism 0736356
Merge branch 'branch-22.04' into fea-header-only-haversine
harrism 5830b66
Apply suggestions from code review
harrism 4c16cd6
require RandomAccessIterator
harrism 7580661
Merge branch 'fea-header-only-haversine' of github.com:harrism/cuspat…
harrism 073e2d7
Convert haversine API to use AOS inputs.
harrism e37d61e
Revert cosmetic changes to top-level haversine.hpp
harrism 9d8e3eb
Align location_2d and remove unused location_3d and coord_2d.
harrism 7f2bbad
__device__ only
harrism 2448677
"" --> <>
harrism 0d954e0
Remove unused macro.
harrism 7d100dc
Add refactoring guide.
harrism daa82a8
Add refactoring guide.
harrism c4ba1f7
Merge branch 'branch-22.04' into fea-header-only-haversine
harrism 454d967
Add fancy iterator test
harrism a5dab4a
Merge branch 'branch-22.06' into fea-header-only-haversine
harrism 08dfe95
.hpp->.cuh
harrism e373065
Add note about not making tests depend on libcudf_test
harrism 3ae133e
gitignore
harrism f8947eb
Add missing include and @
harrism 564cc4c
Don't hide the stream parameter in the detail layer.
harrism 4a6976e
header cleanup
harrism c208144
Simplify coordinate types to a single vec_2d
harrism 21db279
Clean up haversine_test.cu includes
harrism be1226c
type-safe vectors
harrism 8134f12
Fix typo
harrism 42f909e
vec_2d --> lonlat_2d in docs
harrism e5cb703
Merge branch 'branch-22.06' into fea-header-only-haversine
harrism 0b78d87
Review suggestions
harrism 17569c6
Clarified documentation / refactoring guide.
harrism d5ef3b7
style
harrism File filter
Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
125 changes: 125 additions & 0 deletions
125
cpp/include/cuspatial/experimental/detail/haversine.hpp
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,125 @@ | ||
/* | ||
* 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 <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> | ||
__device__ T calculate_haversine_distance(T radius, T a_lon, T a_lat, T b_lon, T b_lat) | ||
{ | ||
auto ax = a_lon * DEGREE_TO_RADIAN; | ||
auto ay = a_lat * DEGREE_TO_RADIAN; | ||
auto bx = b_lon * DEGREE_TO_RADIAN; | ||
auto by = b_lat * 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)); | ||
}; | ||
|
||
template <class LonItA, | ||
class LatItA, | ||
class LonItB, | ||
class LatItB, | ||
class OutputIt, | ||
class T = typename std::iterator_traits<LonItA>::value_type> | ||
OutputIt haversine_distance(LonItA a_lon_first, | ||
LonItA a_lon_last, | ||
LatItA a_lat_first, | ||
LonItB b_lon_first, | ||
LatItB b_lat_first, | ||
OutputIt distance_first, | ||
T const radius = EARTH_RADIUS_KM, | ||
rmm::cuda_stream_view stream = rmm::cuda_stream_default) | ||
{ | ||
static_assert( | ||
std::conjunction_v<std::is_floating_point<T>, | ||
std::is_floating_point<typename std::iterator_traits<LonItA>::value_type>, | ||
std::is_floating_point<typename std::iterator_traits<LatItA>::value_type>, | ||
std::is_floating_point<typename std::iterator_traits<LonItB>::value_type>, | ||
std::is_floating_point<typename std::iterator_traits<LatItB>::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."); | ||
|
||
auto input_tuple = thrust::make_tuple(thrust::make_constant_iterator(static_cast<T>(radius)), | ||
a_lon_first, | ||
a_lat_first, | ||
b_lon_first, | ||
b_lat_first); | ||
|
||
auto input_iter = thrust::make_zip_iterator(input_tuple); | ||
|
||
auto output_size = std::distance(a_lon_first, a_lon_last); | ||
|
||
return thrust::transform(rmm::exec_policy(stream), | ||
input_iter, | ||
input_iter + output_size, | ||
distance_first, | ||
[] __device__(auto inputs) { | ||
return calculate_haversine_distance(thrust::get<0>(inputs), | ||
thrust::get<1>(inputs), | ||
thrust::get<2>(inputs), | ||
thrust::get<3>(inputs), | ||
thrust::get<4>(inputs)); | ||
}); | ||
} | ||
} // namespace detail | ||
|
||
template <class LonItA, class LatItA, class LonItB, class LatItB, class OutputIt, class T> | ||
OutputIt haversine_distance(LonItA a_lon_first, | ||
LonItA a_lon_last, | ||
LatItA a_lat_first, | ||
LonItB b_lon_first, | ||
LatItB b_lat_first, | ||
OutputIt distance_first, | ||
T const radius) | ||
{ | ||
return detail::haversine_distance(a_lon_first, | ||
a_lon_last, | ||
a_lat_first, | ||
b_lon_first, | ||
b_lat_first, | ||
distance_first, | ||
radius, | ||
rmm::cuda_stream_default); | ||
} | ||
|
||
} // namespace cuspatial |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,75 @@ | ||
/* | ||
* 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)`. One distance | ||
* is computed for each (a_lon[i], a_lat[i]) and (b_lon[i], b_lat[i]) point pair. `distance_first` | ||
* must be an iterator to output storage allocated for N distances. | ||
harrism marked this conversation as resolved.
Show resolved
Hide resolved
|
||
* | ||
* All iterators must have the same floating-point `value_type`. | ||
* | ||
* https://en.wikipedia.org/wiki/Haversine_formula | ||
* | ||
* @param[in] a_lon_first: beginning of range of longitude of points in set A | ||
* @param[in] a_lon_last: end of the range of longitude of points in set A | ||
* @param[in] b_lon_first: beginning of range of longitude of points in set B | ||
* @param[in] b_lat_first: beginning of range of latitude of points 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) | ||
* | ||
* @tparam LonItA must meet the requirements of [LegacyForwardIterator][LinkLFI] and be | ||
* device-accessible. | ||
* @tparam LatItA must meet the requirements of [LegacyForwardIterator][LinkLFI] and be | ||
* device-accessible. | ||
* @tparam LonItB must meet the requirements of [LegacyForwardIterator][LinkLFI] and be | ||
* device-accessible. | ||
* @tparam LatItB must meet the requirements of [LegacyForwardIterator][LinkLFI] and be | ||
* device-accessible. | ||
* @tparam OutputIt must meet the requirements of [LegacyForwardIterator][LinkLFI] and be | ||
* device-accessible. | ||
harrism marked this conversation as resolved.
Show resolved
Hide resolved
|
||
* | ||
* @return Output iterator to the element past the last distance computed. | ||
harrism marked this conversation as resolved.
Show resolved
Hide resolved
|
||
* | ||
* [LinkLFI]: https://en.cppreference.com/w/cpp/named_req/ForwardIterator "LegacyForwardIterator" | ||
*/ | ||
template <class LonItA, | ||
class LatItA, | ||
class LonItB, | ||
class LatItB, | ||
class OutputIt, | ||
class T = typename std::iterator_traits<LonItA>::value_type> | ||
OutputIt haversine_distance(LonItA a_lon_first, | ||
LonItA a_lon_last, | ||
LatItA a_lat_first, | ||
LonItB b_lon_first, | ||
LatItB b_lat_first, | ||
OutputIt distance_first, | ||
T const radius = EARTH_RADIUS_KM); | ||
} // namespace cuspatial | ||
|
||
#include <cuspatial/experimental/detail/haversine.hpp> |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Oops, something went wrong.
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
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.