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.