Refactor lonlat_to_cartesian to header-only API (#514)

Following #477, implements the header-only API for `cuspatial::lonlat_to_cartesian` and implements existing C++ API on top of it. Follows the refactoring guide introduced in #477. Note this branch is based on the branch for #477 so diff includes all changes from that PR as well. Once #477 is merged this PR will simplify a lot. Authors: - Mark Harris (https://github.com/harrism) Approvers: - Vyas Ramasubramani (https://github.com/vyasr) - Michael Wang (https://github.com/isVoid) URL: #514
rapidsai · Jun 1, 2022 · 01702a7 · 01702a7
1 parent c149e17
commit 01702a7
Showing 10 changed files with 432 additions and 52 deletions.
diff --git a/.gitignore b/.gitignore
@@ -67,6 +67,9 @@ build/
 cpp/build/
 cpp/thirdparty/googletest/
 
+## CI build directories
+build_rapidsai*/
+
 ## Eclipse IDE
 .project
 .cproject

diff --git a/cpp/include/cuspatial/constants.hpp b/cpp/include/cuspatial/constants.hpp
@@ -18,8 +18,12 @@
 
 #pragma once
 
+namespace cuspatial {
+
 constexpr double DEGREE_TO_RADIAN = M_PI / 180.0;
 constexpr double RADIAN_TO_DEGREE = 180.0 / M_PI;
 
 constexpr double EARTH_RADIUS_KM                = 6371.0;
 constexpr double EARTH_CIRCUMFERENCE_EQUATOR_KM = 40000.0;
+
+}  // namespace cuspatial
diff --git a/cpp/include/cuspatial/experimental/coordinate_transform.cuh b/cpp/include/cuspatial/experimental/coordinate_transform.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 <cuspatial/utility/vec_2d.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+
+#include <iterator>
+
+namespace cuspatial {
+
+/**
+ * @brief Translate longitude/latitude relative to origin to cartesian (x/y) coordinates in km.
+ *
+ * @param[in]  lon_lat_first beginning of range of input longitude/latitude coordinates.
+ * @param[in]  lon_lat_last end of range of input longitude/latitude coordinates.
+ * @param[in]  origin: longitude and latitude of origin.
+ * @param[out] xy_first: beginning of range of output x/y coordinates.
+ * @param[in]  stream: The CUDA stream on which to perform computations and allocate memory.
+ *
+ * All input iterators must have a `value_type` of `cuspatial::lonlat_2d<T>` (Lat/Lon coordinates),
+ * and the output iterator must be able to accept for storage values of type
+ * `cuspatial::cartesian_2d<T>` (Cartesian coordinates).
+ *
+ * @tparam InputIt Iterator over longitude/latitude locations. Must meet the requirements of
+ * [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible.
+ * @tparam OutputIt Iterator over Cartesian output points. Must meet the requirements of
+ * [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible and mutable.
+ * @tparam T the floating-point coordinate value type of input longitude/latitude coordinates.
+ *
+ * @return Output iterator to the element past the last x/y coordinate computed.
+ *
+ * [LinkLRAI]: https://en.cppreference.com/w/cpp/named_req/RandomAccessIterator
+ * "LegacyRandomAccessIterator"
+ */
+template <class InputIt, class OutputIt, class T>
+OutputIt lonlat_to_cartesian(InputIt lon_lat_first,
+                             InputIt lon_lat_last,
+                             OutputIt xy_first,
+                             lonlat_2d<T> origin,
+                             rmm::cuda_stream_view stream = rmm::cuda_stream_default);
+
+}  // namespace cuspatial
+
+#include <cuspatial/experimental/detail/coordinate_transform.cuh>
diff --git a/cpp/include/cuspatial/experimental/detail/coordinate_transform.cuh b/cpp/include/cuspatial/experimental/detail/coordinate_transform.cuh
@@ -0,0 +1,90 @@
+/*
+ * 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/types.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
+
+#include <thrust/transform.h>
+
+#include <iterator>
+#include <type_traits>
+
+namespace cuspatial {
+
+namespace detail {
+
+constexpr double EARTH_CIRCUMFERENCE_KM_PER_DEGREE = EARTH_CIRCUMFERENCE_EQUATOR_KM / 360.0;
+
+template <typename T>
+__device__ inline T midpoint(T a, T b)
+{
+  return (a + b) / 2;
+}
+
+template <typename T>
+__device__ inline T lon_to_x(T lon, T lat)
+{
+  return lon * EARTH_CIRCUMFERENCE_KM_PER_DEGREE * cos(lat * DEGREE_TO_RADIAN);
+};
+
+template <typename T>
+__device__ inline T lat_to_y(T lat)
+{
+  return lat * EARTH_CIRCUMFERENCE_KM_PER_DEGREE;
+};
+
+template <typename T>
+struct to_cartesian_functor {
+  to_cartesian_functor(lonlat_2d<T> origin) : _origin(origin) {}
+
+  cartesian_2d<T> __device__ operator()(lonlat_2d<T> loc)
+  {
+    return cartesian_2d<T>{lon_to_x(_origin.x - loc.x, midpoint(loc.y, _origin.y)),
+                           lat_to_y(_origin.y - loc.y)};
+  }
+
+ private:
+  lonlat_2d<T> _origin{};
+};
+
+}  // namespace detail
+
+template <class InputIt, class OutputIt, class T>
+OutputIt lonlat_to_cartesian(InputIt lon_lat_first,
+                             InputIt lon_lat_last,
+                             OutputIt xy_first,
+                             lonlat_2d<T> origin,
+                             rmm::cuda_stream_view stream)
+{
+  static_assert(std::is_floating_point_v<T>,
+                "lonlat_to_cartesian supports only floating-point coordinates.");
+
+  CUSPATIAL_EXPECTS(origin.x >= -180 && origin.x <= 180 && origin.y >= -90 && origin.y <= 90,
+                    "origin must have valid longitude [-180, 180] and latitude [-90, 90]");
+
+  return thrust::transform(rmm::exec_policy(stream),
+                           lon_lat_first,
+                           lon_lat_last,
+                           xy_first,
+                           detail::to_cartesian_functor{origin});
+}
+
+}  // namespace cuspatial
diff --git a/cpp/include/cuspatial/experimental/haversine.cuh b/cpp/include/cuspatial/experimental/haversine.cuh
@@ -49,7 +49,7 @@ namespace cuspatial {
  * @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.
+ * [LegacyRandomAccessIterator][LinkLRAI] and be device-accessible and mutable.
  * @tparam Location The `value_type` of `LonLatItA` and `LonLatItB`. Must be
  * `cuspatial::lonlat_2d<T>`.
  * @tparam T The underlying coordinate type. Must be a floating-point type.

diff --git a/cpp/include/cuspatial/experimental/type_utils.hpp b/cpp/include/cuspatial/experimental/type_utils.hpp
@@ -20,9 +20,12 @@
 #include <thrust/iterator/transform_output_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
 
+#include <type_traits>
+
 namespace cuspatial {
 
 namespace detail {
+
 template <typename T, typename VectorType>
 struct tuple_to_vec_2d {
   __device__ VectorType operator()(thrust::tuple<T, T> const& pos)

diff --git a/cpp/include/cuspatial/utility/vec_2d.hpp b/cpp/include/cuspatial/utility/vec_2d.hpp
@@ -1,3 +1,19 @@
+/*
+ * 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>
 
@@ -6,9 +22,9 @@ namespace cuspatial {
 /**
  * @brief A 2D vector
  *
- * Used in cuspatial for both Longitude/Latitude (LonLat) coordinate pairs and Cartesian (X/Y)
- * coordinate pairs. For LonLat pairs, the `x` member represents Longitude, and `y` represents
- * Latitude.
+ * This is the base type used in cuspatial for both Longitude/Latitude (LonLat) coordinate pairs and
+ * Cartesian (X/Y) coordinate pairs. For LonLat pairs, the `x` member represents Longitude, and `y`
+ * represents Latitude.
  *
  * @tparam T the base type for the coordinates
  */
@@ -19,16 +35,37 @@ struct alignas(2 * sizeof(T)) vec_2d {
   value_type y;
 };
 
+/**
+ * @brief A strongly typed geographical longitude/latitude coordinate pair.
+ *
+ * `x` is the longitude coordinate, `y` is the latitude coordinate.
+ *
+ * @tparam T the base type for the coordinates.
+ */
 template <typename T>
 struct alignas(2 * sizeof(T)) lonlat_2d : vec_2d<T> {
 };
 
+/**
+ * @brief A strongly typed Cartesian x/y coordinate pair.
+ *
+ * @tparam T the base type for the coordinates.
+ */
 template <typename T>
 struct alignas(2 * sizeof(T)) cartesian_2d : vec_2d<T> {
 };
 
 /**
- * @brief Element-wise add of two 2d vectors.
+ * @brief Compare two 2D vectors for equality.
+ */
+template <typename T>
+bool operator==(vec_2d<T> const& lhs, vec_2d<T> const& rhs)
+{
+  return (lhs.x == rhs.x) && (lhs.y == rhs.y);
+}
+
+/**
+ * @brief Element-wise addition of two 2D vectors.
  */
 template <typename T>
 vec_2d<T> CUSPATIAL_HOST_DEVICE operator+(vec_2d<T> const& a, vec_2d<T> const& b)
@@ -37,7 +74,7 @@ vec_2d<T> CUSPATIAL_HOST_DEVICE operator+(vec_2d<T> const& a, vec_2d<T> const& b
 }
 
 /**
- * @brief Element-wise subtract of two 2d vectors.
+ * @brief Element-wise subtraction of two 2D vectors.
  */
 template <typename T>
 vec_2d<T> CUSPATIAL_HOST_DEVICE operator-(vec_2d<T> const& a, vec_2d<T> const& b)
@@ -46,7 +83,7 @@ vec_2d<T> CUSPATIAL_HOST_DEVICE operator-(vec_2d<T> const& a, vec_2d<T> const& b
 }
 
 /**
- * @brief Scale a 2d vector by ratio @p r.
+ * @brief Scale a 2D vector by a factor @p r.
  */
 template <typename T>
 vec_2d<T> CUSPATIAL_HOST_DEVICE operator*(vec_2d<T> vec, T const& r)
@@ -64,7 +101,7 @@ vec_2d<T> CUSPATIAL_HOST_DEVICE operator*(T const& r, vec_2d<T> vec)
 }
 
 /**
- * @brief Compute dot product of two 2d vectors.
+ * @brief Compute dot product of two 2D vectors.
  */
 template <typename T>
 T CUSPATIAL_HOST_DEVICE dot(vec_2d<T> const& a, vec_2d<T> const& b)
@@ -73,7 +110,7 @@ T CUSPATIAL_HOST_DEVICE dot(vec_2d<T> const& a, vec_2d<T> const& b)
 }
 
 /**
- * @brief Compute 2d determinant of a 2x2 matrix with column vectors @p a and @p b.
+ * @brief Compute 2D determinant of a 2x2 matrix with column vectors @p a and @p b.
  */
 template <typename T>
 T CUSPATIAL_HOST_DEVICE det(vec_2d<T> const& a, vec_2d<T> const& b)

diff --git a/cpp/src/spatial/lonlat_to_cartesian.cu b/cpp/src/spatial/lonlat_to_cartesian.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, 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.
@@ -15,6 +15,9 @@
  */
 
 #include <cuspatial/error.hpp>
+#include <cuspatial/experimental/coordinate_transform.cuh>
+#include <cuspatial/experimental/type_utils.hpp>
+#include <cuspatial/utility/vec_2d.hpp>
 
 #include <cudf/column/column.hpp>
 #include <cudf/column/column_device_view.cuh>
@@ -31,19 +34,6 @@ namespace {
 
 using pair_of_columns = std::pair<std::unique_ptr<cudf::column>, std::unique_ptr<cudf::column>>;
 
-constexpr double earth_circumference_km            = 40000.0;
-constexpr double earth_circumference_km_per_degree = earth_circumference_km / 360.0;
-constexpr double deg_to_rad                        = M_PI / 180;
-
-__device__ inline double midpoint(double a, double b) { return (a + b) / 2; }
-
-__device__ inline double lon_to_x(double lon, double lat)
-{
-  return lon * earth_circumference_km_per_degree * cos(lat * deg_to_rad);
-};
-
-__device__ inline double lat_to_y(double lat) { return lat * earth_circumference_km_per_degree; };
-
 struct lonlat_to_cartesian_functor {
   template <typename T, typename... Args>
   std::enable_if_t<not std::is_floating_point<T>::value, pair_of_columns> operator()(Args&&...)
@@ -53,38 +43,30 @@ struct lonlat_to_cartesian_functor {
 
   template <typename T>
   std::enable_if_t<std::is_floating_point<T>::value, pair_of_columns> operator()(
-    double origin_lon,
-    double origin_lat,
+    T origin_lon,
+    T origin_lat,
     cudf::column_view const& input_lon,
     cudf::column_view const& input_lat,
     rmm::cuda_stream_view stream,
     rmm::mr::device_memory_resource* mr)
   {
     auto size = input_lon.size();
-    auto tid  = cudf::type_to_id<T>();
-    auto type = cudf::data_type{tid};
+    auto type = cudf::data_type{cudf::type_to_id<T>()};
+
     auto output_x =
       cudf::make_fixed_width_column(type, size, cudf::mask_state::UNALLOCATED, stream, mr);
     auto output_y =
       cudf::make_fixed_width_column(type, size, cudf::mask_state::UNALLOCATED, stream, mr);
 
-    auto input_iters = thrust::make_tuple(input_lon.begin<T>(), input_lat.begin<T>());
-
-    auto output_iters =
-      thrust::make_tuple(output_x->mutable_view().begin<T>(), output_y->mutable_view().begin<T>());
+    auto lonlat_begin = cuspatial::make_lonlat_iterator(input_lon.begin<T>(), input_lat.begin<T>());
 
-    auto input_zip  = thrust::make_zip_iterator(input_iters);
-    auto output_zip = thrust::make_zip_iterator(output_iters);
+    auto output_zip = cuspatial::make_zipped_cartesian_2d_output_iterator(
+      output_x->mutable_view().begin<T>(), output_y->mutable_view().begin<T>());
 
-    auto to_cartesian = [=] __device__(auto lonlat) {
-      auto lon = thrust::get<0>(lonlat);
-      auto lat = thrust::get<1>(lonlat);
-      return thrust::make_pair(lon_to_x(origin_lon - lon, midpoint(lat, origin_lat)),
-                               lat_to_y(origin_lat - lat));
-    };
+    auto origin = cuspatial::lonlat_2d<T>{origin_lon, origin_lat};
 
-    thrust::transform(
-      rmm::exec_policy(stream), input_zip, input_zip + input_lon.size(), output_zip, to_cartesian);
+    cuspatial::lonlat_to_cartesian(
+      lonlat_begin, lonlat_begin + input_lon.size(), output_zip, origin, stream);
 
     return std::make_pair(std::move(output_x), std::move(output_y));
   }
@@ -102,6 +84,17 @@ pair_of_columns lonlat_to_cartesian(double origin_lon,
                                     rmm::cuda_stream_view stream,
                                     rmm::mr::device_memory_resource* mr)
 {
+  CUSPATIAL_EXPECTS(
+    origin_lon >= -180 && origin_lon <= 180 && origin_lat >= -90 && origin_lat <= 90,
+    "origin must have valid longitude [-180, 180] and latitude [-90, 90]");
+
+  CUSPATIAL_EXPECTS(input_lon.size() == input_lat.size(), "inputs must have the same length");
+
+  CUSPATIAL_EXPECTS(input_lon.type() == input_lat.type(), "inputs must have the same type");
+
+  CUSPATIAL_EXPECTS(not input_lon.has_nulls() && not input_lat.has_nulls(),
+                    "input cannot contain nulls");
+
   return cudf::type_dispatcher(input_lon.type(),
                                lonlat_to_cartesian_functor(),
                                origin_lon,
@@ -120,17 +113,6 @@ pair_of_columns lonlat_to_cartesian(double origin_lon,
                                     cudf::column_view const& input_lat,
                                     rmm::mr::device_memory_resource* mr)
 {
-  CUSPATIAL_EXPECTS(
-    origin_lon >= -180 && origin_lon <= 180 && origin_lat >= -90 && origin_lat <= 90,
-    "origin must have valid longitude [-180, 180] and latitude [-90, 90]");
-
-  CUSPATIAL_EXPECTS(input_lon.size() == input_lat.size(), "inputs must have the same length");
-
-  CUSPATIAL_EXPECTS(input_lon.type() == input_lat.type(), "inputs must have the same type");
-
-  CUSPATIAL_EXPECTS(not input_lon.has_nulls() && not input_lat.has_nulls(),
-                    "input cannot contain nulls");
-
   return detail::lonlat_to_cartesian(
     origin_lon, origin_lat, input_lon, input_lat, rmm::cuda_stream_default, mr);
 }

diff --git 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(COORDINATE_TRANSFORM_TEST_EXP
+    experimental/spatial/coordinate_transform_test.cu)
diff --git a/cpp/tests/experimental/spatial/coordinate_transform_test.cu b/cpp/tests/experimental/spatial/coordinate_transform_test.cu
@@ -0,0 +1,198 @@
+/*
+ * 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/error.hpp>
+#include <cuspatial/experimental/coordinate_transform.cuh>
+#include <cuspatial/types.hpp>
+
+#include <rmm/device_vector.hpp>
+
+#include <gtest/gtest.h>
+
+template <typename T>
+struct LonLatToCartesianTest : public ::testing::Test {
+};
+
+// float and double are logically the same but would require seperate tests due to precision.
+using TestTypes = ::testing::Types<float, double>;
+TYPED_TEST_CASE(LonLatToCartesianTest, TestTypes);
+
+TYPED_TEST(LonLatToCartesianTest, Empty)
+{
+  using T    = TypeParam;
+  using Loc  = cuspatial::lonlat_2d<T>;
+  using Cart = cuspatial::cartesian_2d<T>;
+
+  auto origin = Loc{-90.66511046, 42.49197018};
+
+  auto h_point_lonlat = std::vector<Loc>{};
+  auto h_expected     = std::vector<Cart>{};
+
+  auto point_lonlat = rmm::device_vector<Loc>{};
+  auto expected     = rmm::device_vector<Cart>{};
+
+  auto xy_output = rmm::device_vector<Cart>{};
+
+  auto xy_end = cuspatial::lonlat_to_cartesian(
+    point_lonlat.begin(), point_lonlat.end(), xy_output.begin(), origin);
+
+  EXPECT_EQ(expected, xy_output);
+  EXPECT_EQ(0, std::distance(xy_output.begin(), xy_end));
+}
+
+TYPED_TEST(LonLatToCartesianTest, Single)
+{
+  using T    = TypeParam;
+  using Loc  = cuspatial::lonlat_2d<T>;
+  using Cart = cuspatial::cartesian_2d<T>;
+
+  auto origin = Loc{-90.66511046, 42.49197018};
+
+  auto h_point_lonlat = std::vector<Loc>({{-90.664973, 42.493894}});
+  auto h_expected     = std::vector<Cart>({{-0.01126195531216838, -0.21375777777718794}});
+
+  auto point_lonlat = rmm::device_vector<Loc>{h_point_lonlat};
+  auto expected     = rmm::device_vector<Cart>{h_expected};
+
+  auto xy_output = rmm::device_vector<Cart>(1);
+
+  auto xy_end = cuspatial::lonlat_to_cartesian(
+    point_lonlat.begin(), point_lonlat.end(), xy_output.begin(), origin);
+
+  EXPECT_EQ(expected, xy_output);
+  EXPECT_EQ(1, std::distance(xy_output.begin(), xy_end));
+}
+
+TYPED_TEST(LonLatToCartesianTest, Extremes)
+{
+  using T    = TypeParam;
+  using Loc  = cuspatial::lonlat_2d<T>;
+  using Cart = cuspatial::cartesian_2d<T>;
+
+  auto origin = Loc{0, 0};
+
+  auto h_points_lonlat = std::vector<Loc>(
+    {{0.0, -90.0}, {0.0, 90.0}, {-180.0, 0.0}, {180.0, 0.0}, {45.0, 0.0}, {-180.0, -90.0}});
+  auto h_expected = std::vector<Cart>({{0.0, 10000.0},
+                                       {0.0, -10000.0},
+                                       {20000.0, 0.0},
+                                       {-20000.0, 0.0},
+                                       {-5000.0, 0.0},
+                                       {14142.13562373095192015, 10000.0}});
+
+  auto points_lonlat = rmm::device_vector<Loc>{h_points_lonlat};
+  auto expected      = rmm::device_vector<Cart>{h_expected};
+
+  auto xy_output = rmm::device_vector<Cart>(6, Cart{-1, -1});
+
+  auto xy_end = cuspatial::lonlat_to_cartesian(
+    points_lonlat.begin(), points_lonlat.end(), xy_output.begin(), origin);
+
+  EXPECT_EQ(expected, xy_output);
+  EXPECT_EQ(6, std::distance(xy_output.begin(), xy_end));
+}
+
+TYPED_TEST(LonLatToCartesianTest, Multiple)
+{
+  using T    = TypeParam;
+  using Loc  = cuspatial::lonlat_2d<T>;
+  using Cart = cuspatial::cartesian_2d<T>;
+
+  auto origin = Loc{-90.66511046, 42.49197018};
+
+  auto h_points_lonlat = std::vector<Loc>({{-90.664973, 42.493894},
+                                           {-90.665393, 42.491520},
+                                           {-90.664976, 42.491420},
+                                           {-90.664537, 42.493823}});
+  auto h_expected      = std::vector<Cart>({
+    {-0.01126195531216838, -0.21375777777718794},
+    {0.02314864865181343, 0.05002000000015667},
+    {-0.01101638630252916, 0.06113111111163663},
+    {-0.04698301003584082, -0.20586888888847929},
+  });
+
+  auto points_lonlat = rmm::device_vector<Loc>{h_points_lonlat};
+  auto expected      = rmm::device_vector<Cart>{h_expected};
+
+  auto xy_output = rmm::device_vector<Cart>(4, Cart{-1, -1});
+
+  auto xy_end = cuspatial::lonlat_to_cartesian(
+    points_lonlat.begin(), points_lonlat.end(), xy_output.begin(), origin);
+
+  EXPECT_EQ(expected, xy_output);
+  EXPECT_EQ(4, std::distance(xy_output.begin(), xy_end));
+}
+
+TYPED_TEST(LonLatToCartesianTest, OriginOutOfBounds)
+{
+  using T    = TypeParam;
+  using Loc  = cuspatial::lonlat_2d<T>;
+  using Cart = cuspatial::cartesian_2d<T>;
+
+  auto origin = Loc{-181, -91};
+
+  auto h_point_lonlat = std::vector<Loc>{};
+  auto h_expected     = std::vector<Cart>{};
+
+  auto point_lonlat = rmm::device_vector<Loc>{};
+  auto expected     = rmm::device_vector<Cart>{};
+
+  auto xy_output = rmm::device_vector<Cart>{};
+
+  EXPECT_THROW(cuspatial::lonlat_to_cartesian(
+                 point_lonlat.begin(), point_lonlat.end(), xy_output.begin(), origin),
+               cuspatial::logic_error);
+}
+
+template <typename T>
+struct identity_xform {
+  using Location = cuspatial::lonlat_2d<T>;
+  __device__ Location operator()(Location const& loc) { return loc; };
+};
+
+// This test verifies that fancy iterators can be passed by using a pass-through transform_iterator
+TYPED_TEST(LonLatToCartesianTest, TransformIterator)
+{
+  using T    = TypeParam;
+  using Loc  = cuspatial::lonlat_2d<T>;
+  using Cart = cuspatial::cartesian_2d<T>;
+
+  auto origin = Loc{-90.66511046, 42.49197018};
+
+  auto h_points_lonlat = std::vector<Loc>({{-90.664973, 42.493894},
+                                           {-90.665393, 42.491520},
+                                           {-90.664976, 42.491420},
+                                           {-90.664537, 42.493823}});
+  auto h_expected      = std::vector<Cart>({
+    {-0.01126195531216838, -0.21375777777718794},
+    {0.02314864865181343, 0.05002000000015667},
+    {-0.01101638630252916, 0.06113111111163663},
+    {-0.04698301003584082, -0.20586888888847929},
+  });
+
+  auto points_lonlat = rmm::device_vector<Loc>{h_points_lonlat};
+  auto expected      = rmm::device_vector<Cart>{h_expected};
+
+  auto xy_output = rmm::device_vector<Cart>(4, Cart{-1, -1});
+
+  auto xform_begin = thrust::make_transform_iterator(points_lonlat.begin(), identity_xform<T>{});
+  auto xform_end   = thrust::make_transform_iterator(points_lonlat.end(), identity_xform<T>{});
+
+  auto xy_end = cuspatial::lonlat_to_cartesian(xform_begin, xform_end, xy_output.begin(), origin);
+
+  EXPECT_EQ(expected, xy_output);
+  EXPECT_EQ(4, std::distance(xy_output.begin(), xy_end));
+}