Skip to content
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
merged 31 commits into from
Apr 28, 2022
Merged
Show file tree
Hide file tree
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 Jan 19, 2022
0736356
Merge branch 'branch-22.04' into fea-header-only-haversine
harrism Jan 19, 2022
5830b66
Apply suggestions from code review
harrism Jan 20, 2022
4c16cd6
require RandomAccessIterator
harrism Jan 20, 2022
7580661
Merge branch 'fea-header-only-haversine' of github.com:harrism/cuspat…
harrism Jan 20, 2022
073e2d7
Convert haversine API to use AOS inputs.
harrism Feb 16, 2022
e37d61e
Revert cosmetic changes to top-level haversine.hpp
harrism Feb 16, 2022
9d8e3eb
Align location_2d and remove unused location_3d and coord_2d.
harrism Feb 16, 2022
7f2bbad
__device__ only
harrism Feb 16, 2022
2448677
"" --> <>
harrism Feb 16, 2022
0d954e0
Remove unused macro.
harrism Feb 16, 2022
7d100dc
Add refactoring guide.
harrism Mar 30, 2022
daa82a8
Add refactoring guide.
harrism Mar 30, 2022
c4ba1f7
Merge branch 'branch-22.04' into fea-header-only-haversine
harrism Mar 31, 2022
454d967
Add fancy iterator test
harrism Mar 31, 2022
a5dab4a
Merge branch 'branch-22.06' into fea-header-only-haversine
harrism Mar 31, 2022
08dfe95
.hpp->.cuh
harrism Mar 31, 2022
e373065
Add note about not making tests depend on libcudf_test
harrism Mar 31, 2022
3ae133e
gitignore
harrism Apr 5, 2022
f8947eb
Add missing include and @
harrism Apr 5, 2022
564cc4c
Don't hide the stream parameter in the detail layer.
harrism Apr 5, 2022
4a6976e
header cleanup
harrism Apr 5, 2022
c208144
Simplify coordinate types to a single vec_2d
harrism Apr 7, 2022
21db279
Clean up haversine_test.cu includes
harrism Apr 7, 2022
be1226c
type-safe vectors
harrism Apr 26, 2022
8134f12
Fix typo
harrism Apr 26, 2022
42f909e
vec_2d --> lonlat_2d in docs
harrism Apr 26, 2022
e5cb703
Merge branch 'branch-22.06' into fea-header-only-haversine
harrism Apr 26, 2022
0b78d87
Review suggestions
harrism Apr 27, 2022
17569c6
Clarified documentation / refactoring guide.
harrism Apr 27, 2022
d5ef3b7
style
harrism Apr 28, 2022
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 4 additions & 0 deletions CHANGELOG.md
Original file line number Diff line number Diff line change
@@ -1,3 +1,7 @@
# cuSpatial 22.04.00 (Date TBD)

Please see https://github.com/rapidsai/cuspatial/releases/tag/v22.04.00a for the latest changes to this development branch.

# cuSpatial 22.02.00 (Date TBD)

Please see https://github.com/rapidsai/cuspatial/releases/tag/v22.02.00a for the latest changes to this development branch.
Expand Down
2 changes: 1 addition & 1 deletion conda/environments/cuspatial_dev_cuda11.0.yml
Original file line number Diff line number Diff line change
Expand Up @@ -7,7 +7,7 @@ channels:
dependencies:
- clang=11.1.0
- clang-tools=11.1.0
- cudf=22.02.*
- cudf=22.04.*
- cudatoolkit=11.0
- gdal>=3.0.2
- geopandas>=0.9.0,<0.10.0a0
Expand Down
2 changes: 1 addition & 1 deletion conda/environments/cuspatial_dev_cuda11.1.yml
Original file line number Diff line number Diff line change
Expand Up @@ -7,7 +7,7 @@ channels:
dependencies:
- clang=11.1.0
- clang-tools=11.1.0
- cudf=22.02.*
- cudf=22.04.*
- cudatoolkit=11.1
- gdal>=3.0.2
- geopandas>=0.9.0,<0.10.0a0
Expand Down
2 changes: 1 addition & 1 deletion conda/environments/cuspatial_dev_cuda11.2.yml
Original file line number Diff line number Diff line change
Expand Up @@ -7,7 +7,7 @@ channels:
dependencies:
- clang=11.1.0
- clang-tools=11.1.0
- cudf=22.02.*
- cudf=22.04.*
- cudatoolkit=11.2
- gdal>=3.0.2
- geopandas>=0.9.0,<0.10.0a0
Expand Down
2 changes: 1 addition & 1 deletion cpp/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -28,7 +28,7 @@ elseif(CMAKE_CUDA_ARCHITECTURES STREQUAL "")
set(CUSPATIAL_BUILD_FOR_DETECTED_ARCHS TRUE)
endif()

project(CUSPATIAL VERSION 22.02.00 LANGUAGES C CXX)
project(CUSPATIAL VERSION 22.04.00 LANGUAGES C CXX)

# Needed because GoogleBenchmark changes the state of FindThreads.cmake,
# causing subsequent runs to have different values for the `Threads::Threads` target.
Expand Down
125 changes: 125 additions & 0 deletions cpp/include/cuspatial/experimental/detail/haversine.hpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,125 @@
/*
Copy link
Contributor

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?

Copy link
Member Author

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.

Copy link
Member Author

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.

Copy link
Contributor

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?

Copy link
Member Author

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).

Copy link
Contributor

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.

Copy link
Member Author

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.

* 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
75 changes: 75 additions & 0 deletions cpp/include/cuspatial/experimental/haversine.hpp
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>
3 changes: 1 addition & 2 deletions cpp/include/cuspatial/haversine.hpp
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2020, NVIDIA CORPORATION.
* Copyright (c) 2020-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.
Expand Down Expand Up @@ -44,5 +44,4 @@ std::unique_ptr<cudf::column> haversine_distance(
cudf::column_view const& b_lat,
double const radius = EARTH_RADIUS_KM,
rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());

} // namespace cuspatial
53 changes: 12 additions & 41 deletions cpp/src/spatial/haversine.cu
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2020, NVIDIA CORPORATION.
* Copyright (c) 2020-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.
Expand All @@ -16,41 +16,22 @@

#include <cuspatial/constants.hpp>
#include <cuspatial/error.hpp>
#include <cuspatial/experimental/haversine.hpp>

#include <cudf/column/column.hpp>
#include <cudf/column/column_factories.hpp>
#include <cudf/column/column_view.hpp>
#include <cudf/copying.hpp>
#include <cudf/types.hpp>
#include <cudf/utilities/type_dispatcher.hpp>

#include <rmm/cuda_stream_view.hpp>
#include <rmm/exec_policy.hpp>
#include <rmm/mr/device/device_memory_resource.hpp>

#include <memory>
#include <type_traits>

namespace {

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));
};

struct haversine_functor {
template <typename T, typename... Args>
std::enable_if_t<not std::is_floating_point<T>::value, std::unique_ptr<cudf::column>> operator()(
Expand All @@ -65,7 +46,7 @@ struct haversine_functor {
cudf::column_view const& a_lat,
cudf::column_view const& b_lon,
cudf::column_view const& b_lat,
double radius,
T radius,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
Expand All @@ -74,25 +55,15 @@ struct haversine_functor {
auto mask_policy = cudf::mask_allocation_policy::NEVER;
auto result = cudf::allocate_like(a_lon, a_lon.size(), mask_policy, mr);

auto input_tuple = thrust::make_tuple(thrust::make_constant_iterator(static_cast<T>(radius)),
a_lon.begin<T>(),
a_lat.begin<T>(),
b_lon.begin<T>(),
b_lat.begin<T>());

auto input_iter = thrust::make_zip_iterator(input_tuple);

thrust::transform(rmm::exec_policy(stream),
input_iter,
input_iter + result->size(),
result->mutable_view().begin<T>(),
[] __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));
});
cuspatial::detail::haversine_distance(
a_lon.begin<T>(),
a_lon.end<T>(),
a_lat.begin<T>(),
b_lon.begin<T>(),
b_lat.begin<T>(),
static_cast<cudf::mutable_column_view>(*result).begin<T>(),
T{radius},
stream);

return result;
}
Expand Down
6 changes: 5 additions & 1 deletion cpp/tests/CMakeLists.txt
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
#=============================================================================
# Copyright (c) 2019-2021, NVIDIA CORPORATION.
# Copyright (c) 2019-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.
Expand Down Expand Up @@ -88,3 +88,7 @@ ConfigureTest(TRAJECTORY_BOUNDING_BOXES_TEST

ConfigureTest(SPATIAL_WINDOW_POINT_TEST
spatial_window/spatial_window_test.cpp)

# Experimental API
ConfigureTest(HAVERSINE_TEST_EXP
experimental/spatial/haversine_test.cu)
Loading