Add CUDA device-callable projection API

[harrism]

Description

Closes #1489

This PR refactors device cuspatial::detail::pipeline into the public API via a type alias cuspatial::device_projection which can be passed to a CUDA kernel and invoked to transform coordinates. cuspatial::projection::get_device_projection(direction) can be used to get a device_projection.

This required changing the direction parameter for cuspatial::detail::pipeline to a constructor parameter rather than a template parameter. I benchmarked before and after this change and saw no significant difference.

I have added tests and an example in README.txt.

Checklist

• I am familiar with the Contributing Guidelines.
• New or existing tests cover these changes.
• The documentation is up to date with these changes.

[tmartin-gh]

Compile error message in CUDA 12.6

cuspatial/cpp/cuproj/include/cuproj/operation/clamp_angular_coordinates.cuh(95): warning #20013-D: calling a constexpr __host__ function("clamp") fr
om a __host__ __device__ function("forward") is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.
      xy.y = std::clamp(xy.y, -half_pi, half_pi);

Fixed by

cuspatial$ git diff -b
diff --git a/cpp/cuproj/include/cuproj/operation/clamp_angular_coordinates.cuh b/cpp/cuproj/include/cuproj/operation/clamp_angular_coordinates.cuh
index e3c7dd7f..bfac81de 100644
--- a/cpp/cuproj/include/cuproj/operation/clamp_angular_coordinates.cuh
+++ b/cpp/cuproj/include/cuproj/operation/clamp_angular_coordinates.cuh
@@ -24,7 +24,7 @@

 #include <thrust/iterator/transform_iterator.h>

-#include <algorithm>
+#include <cuda/std/__algorithm_>

 namespace cuproj {

@@ -92,7 +92,7 @@ class clamp_angular_coordinates : operation<Coordinate> {

     /* Clamp latitude to -pi/2..pi/2 degree range */
     auto half_pi = static_cast<T>(M_PI_2);
-    xy.y         = std::clamp(xy.y, -half_pi, half_pi);
+    xy.y         = cuda::std::clamp(xy.y, -half_pi, half_pi);

[tmartin-gh]

Compiler error in CUDA 12.6

cuspatial/cpp/cuproj/include/cuproj/detail/pipeline.cuh(122): warning #940-D: missing return statement at end of non-void function "cuproj::detail::pipeline<Coordinate, T>::dispatch_op [with Coordinate=cuproj::vec_2d<double>, T=double]"

One solution

cuspatial$ git diff -b cpp/cuproj/include/cuproj/detail/pipeline.cuh
diff --git a/cpp/cuproj/include/cuproj/detail/pipeline.cuh b/cpp/cuproj/include/cuproj/detail/pipeline.cuh
index ab482f85..2f97f9bb 100644
--- a/cpp/cuproj/include/cuproj/detail/pipeline.cuh
+++ b/cpp/cuproj/include/cuproj/detail/pipeline.cuh
@@ -114,7 +114,7 @@ class pipeline {
         auto op = offset_scale_cartesian_coordinates<Coordinate>{params_};
         return op(c, dir_);
       }
-      case operation_type::TRANSVERSE_MERCATOR: {
+      default: { //case operation_type::TRANSVERSE_MERCATOR: {
         auto op = transverse_mercator<Coordinate>{params_};
         return op(c, dir_);
       }

[tmartin-gh]

Compiler error with CUDA 12.6

cuspatial/cpp/cuproj/include/cuproj/detail/pipeline.cuh(75): error: calling a constexpr __host__ function("std::reverse_iterator<const  ::cuproj::operation_type *> ::reverse_iterator(const  ::cuproj::operation_type *)") from a __device__ function("cuproj::detail::pipeline< ::cuproj::vec_2d<double> , double> ::operator () const") is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.

One solution

cuspatial$ git diff -b cpp/cuproj/include/cuproj/detail/pipeline.cuh
diff --git a/cpp/cuproj/include/cuproj/detail/pipeline.cuh b/cpp/cuproj/include/cuproj/detail/pipeline.cuh
index ab482f85..ba9da725 100644
--- a/cpp/cuproj/include/cuproj/detail/pipeline.cuh
+++ b/cpp/cuproj/include/cuproj/detail/pipeline.cuh
@@ -72,7 +72,7 @@ class pipeline {
       thrust::for_each_n(
         thrust::seq, first, num_stages, [&](auto const& op) { c_out = dispatch_op(c_out, op); });
     } else {
-      auto first = std::reverse_iterator(d_ops + num_stages);
+      auto first = cuda::std::reverse_iterator(d_ops + num_stages);
       thrust::for_each_n(
         thrust::seq, first, num_stages, [&](auto const& op) { c_out = dispatch_op(c_out, op); });
     }

[tmartin-gh]

In CUDA 12.6 there are many -Wshadow errors...

cuspatial/cpp/cuproj/include/cuproj/ellipsoid.hpp: In constructor 'constexpr cuproj::ellipsoid<T>::ellipsoid(T, T)':
/nfs/scratch.timothym_ate_1/cuspatial/cpp/cuproj/include/cuproj/ellipsoid.hpp:44:23: warning: declaration of 'a' shadows a member of 'cuproj::ellipsoid<T>' [-Wshadow]
   44 |   constexpr ellipsoid(T a, T inverse_flattening) : a(a)
      |                     ~~^
cuspatial/cpp/cuproj/include/cuproj/ellipsoid.hpp:55:3: note: shadowed declaration is here
   55 |   T a{};      // semi-major axis

Solution left to reader :)

[harrism]

Thanks @tmartin-gh ! I don't think we support CUDA 12.6 in RAPIDS yet.

<cuda/std/__algorithm> is obviously not a public header and it doesn't exist in the version of CCCL that RAPIDS currently depends on. Will certainly adopt cuda::std::clamp() once it is publicly available.

The missing default: is a good catch, it should throw. WIll fix in this PR.
Should be able to fix the shadowing warnings. The constexpr is weird, something must have changed with relaxed-constexpr.

[isVoid]

It seems like in theory, this device projection object also supports thrust transform using the device projection object as input. How about the following test?

TEST(ProjectionTest, device_readme_thrust_example)
{
  using coordinate = cuproj::vec_2d<float>;

  // Make a projection to convert WGS84 (lat, lon) coordinates to
  // UTM zone 56S (x, y) coordinates
  auto proj = cuproj::make_projection<coordinate>("EPSG:4326", "EPSG:32756");

  rmm::cuda_stream_view stream = rmm::cuda_stream_default;

  // Sydney, NSW, Australia
  coordinate sydney{-33.858700, 151.214000};
  thrust::device_vector<coordinate> d_in{1, sydney};
  thrust::device_vector<coordinate> d_out(d_in.size());

  auto d_proj            = proj->get_device_projection(cuproj::direction::FORWARD);
  thrust::transform(
    rmm::exec_policy(stream), d_in.begin(), d_in.end(), d_out.begin(), d_proj
  );
}

[harrism]

In CUDA 12.6 there are many -Wshadow errors...

cuspatial/cpp/cuproj/include/cuproj/ellipsoid.hpp: In constructor 'constexpr cuproj::ellipsoid<T>::ellipsoid(T, T)':
/nfs/scratch.timothym_ate_1/cuspatial/cpp/cuproj/include/cuproj/ellipsoid.hpp:44:23: warning: declaration of 'a' shadows a member of 'cuproj::ellipsoid<T>' [-Wshadow]
   44 |   constexpr ellipsoid(T a, T inverse_flattening) : a(a)
      |                     ~~^
cuspatial/cpp/cuproj/include/cuproj/ellipsoid.hpp:55:3: note: shadowed declaration is here
   55 |   T a{};      // semi-major axis

Solution left to reader :)

After investigation, -Wshadow is not a default warning, and is not included in -Wall. cuSpatial does not specify -Wall so I'm inclined not to fix this one.

[harrism]

It seems like in theory, this device projection object also supports thrust transform using the device projection object as input. How about the following test?

TEST(ProjectionTest, device_readme_thrust_example)
{
  using coordinate = cuproj::vec_2d<float>;

  // Make a projection to convert WGS84 (lat, lon) coordinates to
  // UTM zone 56S (x, y) coordinates
  auto proj = cuproj::make_projection<coordinate>("EPSG:4326", "EPSG:32756");

  rmm::cuda_stream_view stream = rmm::cuda_stream_default;

  // Sydney, NSW, Australia
  coordinate sydney{-33.858700, 151.214000};
  thrust::device_vector<coordinate> d_in{1, sydney};
  thrust::device_vector<coordinate> d_out(d_in.size());

  auto d_proj            = proj->get_device_projection(cuproj::direction::FORWARD);
  thrust::transform(
    rmm::exec_policy(stream), d_in.begin(), d_in.end(), d_out.begin(), d_proj
  );
}

But is this test necessary? After all this is exactly the implementation of projection::transform():

template <class InputCoordIter, class OutputCoordIter>
  void transform(InputCoordIter first,
                 InputCoordIter last,
                 OutputCoordIter result,
                 direction dir,
                 rmm::cuda_stream_view stream = rmm::cuda_stream_default) const
  {
    thrust::transform(rmm::exec_policy(stream), first, last, result, get_device_projection(dir));
  }

[tmartin-gh]

In CUDA 12.6 there are many -Wshadow errors...

cuspatial/cpp/cuproj/include/cuproj/ellipsoid.hpp: In constructor 'constexpr cuproj::ellipsoid<T>::ellipsoid(T, T)':
/nfs/scratch.timothym_ate_1/cuspatial/cpp/cuproj/include/cuproj/ellipsoid.hpp:44:23: warning: declaration of 'a' shadows a member of 'cuproj::ellipsoid<T>' [-Wshadow]
   44 |   constexpr ellipsoid(T a, T inverse_flattening) : a(a)
      |                     ~~^
cuspatial/cpp/cuproj/include/cuproj/ellipsoid.hpp:55:3: note: shadowed declaration is here
   55 |   T a{};      // semi-major axis

Solution left to reader :)

After investigation, -Wshadow is not a default warning, and is not included in -Wall. cuSpatial does not specify -Wall so I'm inclined not to fix this one.

FWIW MatX turns this warning on. https://github.com/NVIDIA/MatX/blob/0e5c634ee29f7e5a60802d34afbeadab6dac4e35/CMakeLists.txt#L145

[harrism]

@tmartin-gh I have fixed all the errors you found except the -Wshadow warnings. Please test again.

[review-notebook-app]

Check out this pull request on 

See visual diffs & provide feedback on Jupyter Notebooks.

---

Powered by ReviewNB

[isVoid]

But is this test necessary? After all this is exactly the implementation of projection::transform():

template <class InputCoordIter, class OutputCoordIter>
  void transform(InputCoordIter first,
                 InputCoordIter last,
                 OutputCoordIter result,
                 direction dir,
                 rmm::cuda_stream_view stream = rmm::cuda_stream_default) const
  {
    thrust::transform(rmm::exec_policy(stream), first, last, result, get_device_projection(dir));
  }

You are right about this. Resolved!

[tmartin-gh]

@tmartin-gh I have fixed all the errors you found except the -Wshadow warnings. Please test again.

Looks good. Approved.

[harrism]

/merge