diff --git a/conda/environments/all_cuda-118_arch-x86_64.yaml b/conda/environments/all_cuda-118_arch-x86_64.yaml index b2c984d5e..324fbe1f6 100644 --- a/conda/environments/all_cuda-118_arch-x86_64.yaml +++ b/conda/environments/all_cuda-118_arch-x86_64.yaml @@ -20,8 +20,6 @@ dependencies: - gtest=1.10.0 - ipython - libcudf=23.04 -- libcusparse-dev=11.7.5.86 -- libcusparse=11.7.5.86 - librmm=23.04 - myst-parser - nbsphinx diff --git a/conda/recipes/libcuspatial/conda_build_config.yaml b/conda/recipes/libcuspatial/conda_build_config.yaml index 5cc0c3ea7..940595460 100644 --- a/conda/recipes/libcuspatial/conda_build_config.yaml +++ b/conda/recipes/libcuspatial/conda_build_config.yaml @@ -18,13 +18,3 @@ gtest_version: sysroot_version: - "2.17" - -# The CTK libraries below are missing from the conda-forge::cudatoolkit -# package. The "*_host_*" version specifiers correspond to `11.8` packages and the -# "*_run_*" version specifiers correspond to `11.x` packages. - -libcusparse_host_version: - - "=11.7.5.86" - -libcusparse_run_version: - - ">=11.6.0.43,<12" diff --git a/conda/recipes/libcuspatial/meta.yaml b/conda/recipes/libcuspatial/meta.yaml index d9e363e52..f94172e66 100644 --- a/conda/recipes/libcuspatial/meta.yaml +++ b/conda/recipes/libcuspatial/meta.yaml @@ -42,8 +42,6 @@ requirements: - gmock {{ gtest_version }} - gtest {{ gtest_version }} - libcudf ={{ minor_version }} - - libcusparse {{ libcusparse_host_version }} - - libcusparse-dev {{ libcusparse_host_version }} - librmm ={{ minor_version }} outputs: @@ -64,8 +62,6 @@ outputs: run: - cudatoolkit {{ cuda_spec }} - libcudf ={{ minor_version }} - - libcusparse {{ libcusparse_run_version }} - - libcusparse-dev {{ libcusparse_run_version }} - librmm ={{ minor_version }} test: commands: diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 0f6ad6ff9..8bd8958bc 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -111,7 +111,6 @@ include(cmake/thirdparty/CUSPATIAL_GetCUDF.cmake) add_library(cuspatial src/column/geometry_column_view.cpp src/indexing/construction/point_quadtree.cu - src/interpolate/cubic_spline.cu src/join/quadtree_point_in_polygon.cu src/join/quadtree_point_to_nearest_linestring.cu src/join/quadtree_bbox_filtering.cu @@ -188,7 +187,7 @@ endif() target_compile_definitions(cuspatial PUBLIC "SPDLOG_ACTIVE_LEVEL=SPDLOG_LEVEL_${RMM_LOGGING_LEVEL}") # Specify the target module library dependencies -target_link_libraries(cuspatial PUBLIC cudf::cudf CUDA::cusparse${_ctk_static_suffix}) +target_link_libraries(cuspatial PUBLIC cudf::cudf) add_library(cuspatial::cuspatial ALIAS cuspatial) diff --git a/cpp/include/cuspatial/cubic_spline.hpp b/cpp/include/cuspatial/cubic_spline.hpp deleted file mode 100644 index 28ee62ccf..000000000 --- a/cpp/include/cuspatial/cubic_spline.hpp +++ /dev/null @@ -1,98 +0,0 @@ -/* - * 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. - * 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 -#include -#include - -#include - -#include - -namespace cuspatial { - -/** - * @addtogroup cubic_spline - * @{ - */ - -/** - * @brief Create a table of cubic spline coefficients from columns of coordinates. - * - * Computes coefficients for a natural cubic spline similar to the method - * found on http://mathworld.wolfram.com/CubicSpline.html . - * - * The input data arrays `t` and `y` contain the vertices of many concatenated - * splines. - * - * Currently, all input splines must be the same length. The minimum supported - * length is 5. - * - * @note Ids should be prefixed with a 0, even when only a single spline - * is fit, ids will be {0, 0} - * - * @param[in] t column_view of independent coordinates for fitting splines - * @param[in] y column_view of dependent variables to be fit along t axis - * @param[in] ids of incoming coordinate sets - * @param[in] offsets the exclusive scan of the spline sizes, prefixed by - * 0. For example, for 3 splines of 5 vertices each, the offsets input array - * is {0, 5, 10, 15}. - * @param[in] mr The memory resource to use for allocating output - * - * @return cudf::table_view of coefficients for spline interpolation. The size - * of the table is ((M-n), 4) where M is `t.size()` and and n is - * `ids.size()-1`. - **/ -std::unique_ptr cubicspline_interpolate( - cudf::column_view const& query_points, - cudf::column_view const& spline_ids, - cudf::column_view const& offsets, - cudf::column_view const& source_points, - cudf::table_view const& coefficients, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); - -/** - * @brief Compute cubic interpolations of a set of points based on their - * ids and a coefficient matrix. - * - * @param[in] query_points column of coordinate values to be interpolated. - * @param[in] spline_ids ids that identift the spline to interpolate each - * coordinate into. - * @param[in] offsets int32 column of offset of the source_points. - * This is used to calculate which values from the coefficients are - * used for each interpolation. - * @param[in] source_points column of the original `t` values used - * to compute the coefficients matrix. These source points are used to - * identify which specific spline a given query_point is interpolated with. - * @param[in] coefficients table of spline coefficients produced by - * cubicspline_coefficients. - * @param[in] mr The memory resource to use for allocating output - * - * @return cudf::column `y` coordinates interpolated from `x` and `coefs`. - **/ -std::unique_ptr cubicspline_coefficients( - cudf::column_view const& t, - cudf::column_view const& y, - cudf::column_view const& ids, - cudf::column_view const& offsets, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); - -/** - * @} // end of doxygen group - */ -} // namespace cuspatial diff --git a/cpp/include/cuspatial/cusparse_error.hpp b/cpp/include/cuspatial/cusparse_error.hpp deleted file mode 100644 index b0ffa4dfb..000000000 --- a/cpp/include/cuspatial/cusparse_error.hpp +++ /dev/null @@ -1,65 +0,0 @@ -/* - * 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. - * 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 - -#include - -namespace cuspatial { - -/**---------------------------------------------------------------------------* - * @brief Exception thrown when logical precondition is violated. - * - * This exception should not be thrown directly and is instead thrown by the - * CUSPARSE_TRY macro. - * - *---------------------------------------------------------------------------**/ -struct cusparse_error : public std::runtime_error { - cusparse_error(std::string const& message) : std::runtime_error(message) {} -}; - -namespace detail { - -inline void throw_cusparse_error(cusparseStatus_t error, const char* file, unsigned int line) -{ - // would be nice to include `cusparseGetErrorName(error)` and - // `cusparseGetErrorString(error)`, but those aren't introduced until - // cuda 10.1 (and not in the initial release). - throw cuspatial::cusparse_error( - std::string{"CUSPARSE error encountered at: " + std::string{file} + ":" + std::to_string(line) + - ": " + std::to_string(error)}); -} - -} // namespace detail -} // namespace cuspatial - -/**---------------------------------------------------------------------------* - * @brief Error checking macro for cuSPARSE runtime API functions. - * - * Invokes a cuSPARSE runtime API function call, if the call does not return - * CUSPARSE_STATUS_SUCCESS, throws an exception detailing the cuSPARSE error - * that occurred. - * - *---------------------------------------------------------------------------**/ -#define CUSPARSE_TRY(call) \ - do { \ - cusparseStatus_t status = (call); \ - if (CUSPARSE_STATUS_SUCCESS != status) { \ - cuspatial::detail::throw_cusparse_error(status, __FILE__, __LINE__); \ - } \ - } while (0); diff --git a/cpp/include/cuspatial/detail/cubic_spline.hpp b/cpp/include/cuspatial/detail/cubic_spline.hpp deleted file mode 100644 index 88a9cecfd..000000000 --- a/cpp/include/cuspatial/detail/cubic_spline.hpp +++ /dev/null @@ -1,76 +0,0 @@ -/* - * 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. - * 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 -#include -#include - -#include - -#include - -namespace cuspatial { - -namespace detail { - -/** - * @brief For each query point, finds the index of the last coordinate in source_points, - * that is smaller than the query_point, grouped by prefixes and curve_ids. - * - * This implementation of cubic curve has four coefficients for each pair of control points. For - * a curve then with n control points there will be n-1 sets of coefficients. This function - * finds which set of coefficients to use for a given query_point. Each `query_point[i]` is passed - * with a corresponding `curve_ids[i], identifying which offset `j` to use from `prefixes` into - * `source_points`. For example, given two sets of `source_points` = [0, 1, 2, 3, 4, 0, 2, 5, 10, - *20] with `prefixes = [0, 5, 10], and a single `query_point = 6` with `curve_ids = 1`, the - *coefficient position `6` is returned. - * - * The first curve, specified by `curve_ids = 0` uses the first four coefficient indices 0...3, and - * the second curve uses the next four indices. `6 > 5` specifying the third pair (ordinal 2) - * of `source_points` (also known as control points). - * - * Below is a simple diagram of the cofficient indices that correspond with the `source_points`. - * - * [0, 1, 2, 3, 4, 0, 2, 5, 10, 20] - * 0 1 2 3 4 5 6 7 - * - * @param query_points column of coordinate values to be interpolated. - * @param spline_ids ids that identify the spline to interpolate each - * coordinate into. - * @param offsets int32 column of offsets of the source_points. - * This is used to calculate which values from the coefficients are - * used for each interpolation. - * @param source_points column of the original `t` values used - * to compute the coefficients matrix. - * @param mr the optional caller specified RMM memory resource - * @param stream the optional caller specified cudaStream - * - * @return cudf::column of size equal to query points, one index position - * of the first source_point mapped by offsets that is smaller than each - * query point. - **/ -std::unique_ptr find_coefficient_indices(cudf::column_view const& query_points, - cudf::column_view const& curve_ids, - cudf::column_view const& prefixes, - cudf::column_view const& source_points, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr); - -} // namespace detail - -} // namespace cuspatial diff --git a/cpp/include/doxygen_groups.h b/cpp/include/doxygen_groups.h index a22ee832f..a80fbe9ab 100644 --- a/cpp/include/doxygen_groups.h +++ b/cpp/include/doxygen_groups.h @@ -77,11 +77,6 @@ * @file point_linestring_nearest_points.hpp * @file point_linestring_nearest_points.cuh * @} - * @defgroup cubic_spline Cubic Spline - * @{ - * @brief APIs related to cubic splines - * @file cubic_spline.hpp - * @} * @} * @defgroup trajectory_api Trajectory APIs * @{ diff --git a/cpp/src/interpolate/cubic_spline.cu b/cpp/src/interpolate/cubic_spline.cu deleted file mode 100644 index 5f3f411eb..000000000 --- a/cpp/src/interpolate/cubic_spline.cu +++ /dev/null @@ -1,504 +0,0 @@ -/* - * Copyright (c) 2020-2021, 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 -#include -#include -#include - -#include -#include -#include - -#include -#include - -#include - -#include -#include - -namespace { // anonymous - -// This functor performs one linear search for each input point in query_coords -struct parallel_search { - template - std::enable_if_t::value, std::unique_ptr> operator()( - cudf::column_view const& search_coords, - cudf::column_view const& curve_ids, - cudf::column_view const& prefixes, - cudf::column_view const& query_coords, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) - { - const T SEARCH_OFFSET{0.0001}; - const T QUERY_OFFSET{0.00001}; - - const T* p_search_coords = search_coords.data(); - const int32_t* p_curve_ids = curve_ids.data(); - const int32_t* p_prefixes = prefixes.data(); - const T* p_query_coords = query_coords.data(); - auto result = cudf::make_numeric_column( - curve_ids.type(), search_coords.size(), cudf::mask_state::UNALLOCATED, stream, mr); - int32_t* p_result = result->mutable_view().data(); - thrust::for_each( - rmm::exec_policy(stream), - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(search_coords.size()), - [=] __device__(int32_t index) { - int32_t curve = p_curve_ids[index]; - int32_t len = p_prefixes[curve + 1] - p_prefixes[curve]; - int32_t query_coord_offset = p_prefixes[curve]; - int32_t coefficient_table_offset = p_prefixes[curve] - curve; - // O(n) search, can do log(n) easily - const T search_coord = p_search_coords[index] + SEARCH_OFFSET; - for (int32_t i = 1; i < len; ++i) { - if ((search_coord < p_query_coords[query_coord_offset + i] + QUERY_OFFSET)) { - p_result[index] = coefficient_table_offset + i - 1; - return; - } - } - // NOTE: Important failure case: - // This will use the final set of coefficients - // for t_ values that are outside of the original - // interpolation range. - p_result[index] = coefficient_table_offset + len - 2; - }); - return result; - } - - template - std::enable_if_t::value, std::unique_ptr> operator()( - Args&&...) - { - CUSPATIAL_FAIL("Non-floating point operation is not supported."); - } -}; - -// This functor simply computes the interpolation of each coordinate `t[i]` -// using the coefficients from row `coef_indices[i]`. -struct interpolate { - template - std::enable_if_t::value, std::unique_ptr> operator()( - cudf::column_view const& t, - cudf::column_view const& coef_indices, - cudf::table_view const& coefficients, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) - { - const T* p_t = t.data(); - const int32_t* p_coef_indices = coef_indices.data(); - const T* p_d3 = coefficients.column(3).data(); - const T* p_d2 = coefficients.column(2).data(); - const T* p_d1 = coefficients.column(1).data(); - const T* p_d0 = coefficients.column(0).data(); - auto result = - cudf::make_numeric_column(t.type(), t.size(), cudf::mask_state::UNALLOCATED, stream, mr); - T* p_result = result->mutable_view().data(); - thrust::for_each(rmm::exec_policy(stream), - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(t.size()), - [=] __device__(int32_t index) { - int32_t h = p_coef_indices[index]; - p_result[index] = - p_d3[h] + - p_t[index] * (p_d2[h] + p_t[index] * (p_d1[h] + (p_t[index] * p_d0[h]))); - }); - return result; - }; - template - std::enable_if_t::value, std::unique_ptr> operator()( - Args&&...) - { - CUSPATIAL_FAIL("Non-floating point operation is not supported."); - } -}; - -// This functor computes the coefficients table for the cubic hermite spline -// specified by the inputs `t` and `y`. -struct coefficients_compute { - template - std::enable_if_t::value, void> operator()( - cudf::column_view const& t, - cudf::column_view const& y, - cudf::column_view const& prefixes, - cudf::mutable_column_view const& h, - cudf::mutable_column_view const& i, - cudf::mutable_column_view const& z, - cudf::mutable_column_view const& d3, - cudf::mutable_column_view const& d2, - cudf::mutable_column_view const& d1, - cudf::mutable_column_view const& d0, - rmm::cuda_stream_view stream) - { - const T* p_t = t.data(); - const T* p_y = y.data(); - const int32_t* p_prefixes = prefixes.data(); - T* p_h = h.data(); - T* p_i = i.data(); - T* p_z = z.data(); - T* p_d3 = d3.data(); - T* p_d2 = d2.data(); - T* p_d1 = d1.data(); - T* p_d0 = d0.data(); - thrust::for_each( - rmm::exec_policy(stream), - thrust::make_counting_iterator(1), - thrust::make_counting_iterator(prefixes.size()), - [p_t, p_y, p_prefixes, p_h, p_i, p_z, p_d3, p_d2, p_d1, p_d0] __device__(int32_t index) { - int32_t n = p_prefixes[index] - p_prefixes[index - 1]; - int32_t h = p_prefixes[index - 1]; - int32_t dh = p_prefixes[index - 1] - (index - 1); - int32_t ci = 0; - for (ci = 0; ci < n - 1; ++ci) { - T a = p_y[h + ci]; - T b = p_i[h + ci] - p_h[h + ci] * (p_z[h + ci + 1] + 2 * p_z[h + ci]) / 6; - T c = p_z[h + ci] / 2.0; - T d = (p_z[h + ci + 1] - p_z[h + ci]) / 6 * p_h[h + ci]; - T t = p_t[h + ci]; - p_d3[dh + ci] = d; - p_d2[dh + ci] = c - 3 * d * t; - p_d1[dh + ci] = b - t * (2 * c - t * (3 * d)); - p_d0[dh + ci] = a - t * (b - t * (c - t * d)); // horners - } - }); - } - - template - std::enable_if_t::value, void> operator()(Args&&...) - { - CUSPATIAL_FAIL("Non-floating point operation is not supported."); - } -}; - -// Computes the diagonal `D` of a large sparse matrix, and also the upper and -// lower diagonals `Dlu`, which in this case are equal. -struct compute_spline_tridiagonals { - template - std::enable_if_t::value, void> operator()( - cudf::column_view const& t, - cudf::column_view const& y, - cudf::column_view const& prefixes, - cudf::mutable_column_view const& D, - cudf::mutable_column_view const& Dlu, - cudf::mutable_column_view const& u, - cudf::mutable_column_view const& h, - cudf::mutable_column_view const& i, - rmm::cuda_stream_view stream) - { - const T* p_t = t.data(); - const T* p_y = y.data(); - const int32_t* p_prefixes = prefixes.data(); - T* p_d = D.data(); - T* p_dlu = Dlu.data(); - T* p_u = u.data(); - T* p_h = h.data(); - T* p_i = i.data(); - thrust::for_each(rmm::exec_policy(stream), - thrust::make_counting_iterator(1), - thrust::make_counting_iterator(prefixes.size()), - [p_t, p_y, p_prefixes, p_d, p_dlu, p_u, p_h, p_i] __device__(int32_t index) { - int32_t n = p_prefixes[index] - p_prefixes[index - 1]; - int32_t h = p_prefixes[index - 1]; - int32_t ci = 0; - for (ci = 0; ci < n - 1; ++ci) { - p_h[h + ci] = p_t[h + ci + 1] - p_t[h + ci]; - p_i[h + ci] = (p_y[h + ci + 1] - p_y[h + ci]) / p_h[h + ci]; - } - for (ci = 0; ci < n - 2; ++ci) { - p_d[h + ci + 1] = (p_h[h + ci + 1] + p_h[h + (n - 2) - ci]) * 2; - p_u[h + ci + 1] = (p_i[h + ci + 1] - p_i[h + (n - 2) - ci]) * 6; - } - for (ci = 0; ci < n - 3; ++ci) { - p_dlu[h + ci + 1] = p_i[h + ci + 1]; - } - }); - } - - template - std::enable_if_t::value, void> operator()(Args&&...) - { - CUSPATIAL_FAIL("Non-floating point operation is not supported."); - } -}; - -} // anonymous namespace - -namespace cuspatial { - -namespace detail { - -/** - * @brief Finds the lower interpolant position of query_points from a set of - * interpolation independent variables. - * - * @param[in] query_points column of coordinate values to be interpolated. - * @param[in] spline_ids ids that identify the spline to interpolate each - * coordinate into. - * @param[in] offsets int32 column of offsets of the source_points. - * This is used to calculate which values from the coefficients are - * used for each interpolation. - * @param[in] source_points column of the original `t` values used - * to compute the coefficients matrix. These source points are used to - * identify which specific spline a given query_point is interpolated with. - * cubicspline_coefficients. - * @param[in] mr the optional caller specified RMM memory resource - * @param[in] stream the optional caller specified cudaStream - * - * @return cudf::column of size equal to query points, one index position - * of the first source_point mapped by offsets that is smaller than each - * query point. - **/ -std::unique_ptr find_coefficient_indices(cudf::column_view const& query_points, - cudf::column_view const& curve_ids, - cudf::column_view const& prefixes, - cudf::column_view const& source_points, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - auto coefficient_indices = cudf::type_dispatcher(query_points.type(), - parallel_search{}, - query_points, - curve_ids, - prefixes, - source_points, - stream, - mr); - return coefficient_indices; -} - -/** - * @brief Compute cubic interpolations of a set of points based on their - * ids and a coefficient matrix. - * - * @param[in] query_points column of coordinate values to be interpolated. - * @param[in] spline_ids ids that identift the spline to interpolate each - * coordinate into. - * @param[in] offsets int32 column of offset of the source_points. - * This is used to calculate which values from the coefficients are - * used for each interpolation. - * @param[in] source_points column of the original `t` values used - * to compute the coefficients matrix. These source points are used to - * identify which specific spline a given query_point is interpolated with. - * @param[in] coefficients table of spline coefficients produced by - * cubicspline_coefficients. - * @param[in] mr the optional caller specified RMM memory resource - * @param[in] stream the optional caller specified cudaStream - * - * @return cudf::column `y` coordinates interpolated from `x` and `coefs`. - **/ -std::unique_ptr cubicspline_interpolate(cudf::column_view const& query_points, - cudf::column_view const& curve_ids, - cudf::column_view const& prefixes, - cudf::column_view const& source_points, - cudf::table_view const& coefficients, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - auto coefficient_indices = cudf::type_dispatcher(query_points.type(), - parallel_search{}, - query_points, - curve_ids, - prefixes, - source_points, - stream, - mr); - - auto result = cudf::type_dispatcher(query_points.type(), - interpolate{}, - query_points, - coefficient_indices->view(), - coefficients, - stream, - mr); - - return result; -} - -/** - * @brief Create a table of cubic spline coefficients from columns of coordinates. - * - * Computes coefficients for a natural cubic spline similar to the method - * found on http://mathworld.wolfram.com/CubicSpline.html . - * - * The input data arrays `t` and `y` contain the vertices of many concatenated - * splines. - * - * Currently, all input splines must be the same length. The minimum supported - * length is 5. - * - * @note Ids should be prefixed with a 0, even when only a single spline - * is fit, ids will be {0, 0} - * - * @param[in] t column_view of independent coordinates for fitting splines - * @param[in] y column_view of dependent variables to be fit along t axis - * @param[in] ids of incoming coordinate sets - * @param[in] offsets the exclusive scan of the spline sizes, prefixed by - * 0. For example, for 3 splines of 5 vertices each, the offsets input array - * is {0, 5, 10, 15}. - * @param[in] mr the optional caller specified RMM memory resource - * @param[in] stream the optional caller specified cudaStream - * - * @return cudf::table_view of coefficients for spline interpolation. The size - * of the table is ((M-n), 4) where M is `t.size()` and and n is - * `ids.size()-1`. - **/ -std::unique_ptr cubicspline_coefficients(cudf::column_view const& t, - cudf::column_view const& y, - cudf::column_view const& ids, - cudf::column_view const& prefixes, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - int64_t n = y.size(); - auto h_col = make_numeric_column(y.type(), n, cudf::mask_state::UNALLOCATED, stream, mr); - auto i_col = make_numeric_column(y.type(), n, cudf::mask_state::UNALLOCATED, stream, mr); - auto D_col = make_numeric_column(y.type(), n, cudf::mask_state::UNALLOCATED, stream, mr); - auto Dlu_col = make_numeric_column(y.type(), n, cudf::mask_state::UNALLOCATED, stream, mr); - auto Dll_col = make_numeric_column(y.type(), n, cudf::mask_state::UNALLOCATED, stream, mr); - auto u_col = make_numeric_column(y.type(), n, cudf::mask_state::UNALLOCATED, stream, mr); - auto h_buffer = h_col->mutable_view(); - auto i_buffer = i_col->mutable_view(); - auto D_buffer = D_col->mutable_view(); - auto Dlu_buffer = Dll_col->mutable_view(); - auto Dll_buffer = Dll_col->mutable_view(); - auto u_buffer = u_col->mutable_view(); - - auto zero = cudf::numeric_scalar(0.0); - auto one = cudf::numeric_scalar(1.0); - cudf::fill_in_place(h_buffer, 0, h_col->size(), zero); - cudf::fill_in_place(i_buffer, 0, i_col->size(), zero); - cudf::fill_in_place(D_buffer, 0, D_col->size(), one); - cudf::fill_in_place(Dlu_buffer, 0, Dlu_col->size(), zero); - cudf::fill_in_place(u_buffer, 0, u_col->size(), zero); - - cudf::type_dispatcher(y.type(), - compute_spline_tridiagonals{}, - t, - y, - prefixes, - D_buffer, - Dlu_buffer, - u_buffer, - h_buffer, - i_buffer, - stream); - - // cusparse solve n length m tridiagonal systems - // 4. call cusparsegtsv2() to solve - // 4.1 Get cuSparse library context - // compute inputs: - // handle: the cuSparse library context - // m: size - // n: number of columns of solution matrix B - // dl, d, du: vectors of the diagonal - // B: (ldb, n) dimensional dense matrix to be solved for - // ldb: leading dimension of B - // pBuffer: get size of thisu by gtsv2_bufferSizeExt - cusparseHandle_t handle; - - CUSPATIAL_CUDA_TRY(cudaMalloc(&handle, sizeof(cusparseHandle_t))); - CUSPARSE_TRY(cusparseCreate(&handle)); - - size_t pBufferSize; - int32_t batchStride = y.size() / (prefixes.size() - 1); - int32_t batchSize = batchStride; - - CUSPARSE_TRY(cusparseSgtsv2StridedBatch_bufferSizeExt(handle, - batchSize, - Dll_buffer.data(), - D_buffer.data(), - Dlu_buffer.data(), - u_buffer.data(), - prefixes.size() - 1, - batchStride, - &pBufferSize)); - - rmm::device_vector pBuffer(pBufferSize); - - CUSPARSE_TRY(cusparseSgtsv2StridedBatch(handle, - batchSize, - Dll_buffer.data(), - D_buffer.data(), - Dlu_buffer.data(), - u_buffer.data(), - prefixes.size() - 1, - batchStride, - pBuffer.data().get())); - - CUSPARSE_TRY(cusparseDestroy(handle)); - - int32_t dn = n - (prefixes.size() - 1); - // Finally, compute coefficients via Horner's scheme - auto d3_col = make_numeric_column(y.type(), dn, cudf::mask_state::UNALLOCATED, stream, mr); - auto d2_col = make_numeric_column(y.type(), dn, cudf::mask_state::UNALLOCATED, stream, mr); - auto d1_col = make_numeric_column(y.type(), dn, cudf::mask_state::UNALLOCATED, stream, mr); - auto d0_col = make_numeric_column(y.type(), dn, cudf::mask_state::UNALLOCATED, stream, mr); - auto d3 = d3_col->mutable_view(); - auto d2 = d2_col->mutable_view(); - auto d1 = d1_col->mutable_view(); - auto d0 = d0_col->mutable_view(); - - cudf::type_dispatcher(y.type(), - coefficients_compute{}, - t, - y, - prefixes, - h_buffer, - i_buffer, - u_buffer, - d3, - d2, - d1, - d0, - stream); - - // Place d3..0 into a table and return - std::vector> table; - table.push_back(std::move(d3_col)); - table.push_back(std::move(d2_col)); - table.push_back(std::move(d1_col)); - table.push_back(std::move(d0_col)); - std::unique_ptr result = std::make_unique(move(table)); - return result; -} - -} // namespace detail - -// Calls the interpolate function using default memory resources. -std::unique_ptr cubicspline_interpolate(cudf::column_view const& query_points, - cudf::column_view const& curve_ids, - cudf::column_view const& prefixes, - cudf::column_view const& source_points, - cudf::table_view const& coefficients, - rmm::mr::device_memory_resource* mr) -{ - return cuspatial::detail::cubicspline_interpolate( - query_points, curve_ids, prefixes, source_points, coefficients, rmm::cuda_stream_default, mr); -} - -// Calls the coeffiecients function using default memory resources. -std::unique_ptr cubicspline_coefficients(cudf::column_view const& t, - cudf::column_view const& y, - cudf::column_view const& ids, - cudf::column_view const& prefixes, - rmm::mr::device_memory_resource* mr) -{ - return cuspatial::detail::cubicspline_coefficients( - t, y, ids, prefixes, rmm::cuda_stream_default, mr); -} - -} // namespace cuspatial diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 16cb67531..6182e1238 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -52,9 +52,6 @@ endfunction(ConfigureTest) ### test sources ################################################################################## ################################################################################################### -ConfigureTest(CUBIC_SPLINE_TEST - interpolate/cubic_spline_test.cpp) - ConfigureTest(SINUSOIDAL_PROJECTION_TEST spatial/sinusoidal_projection_test.cu) diff --git a/cpp/tests/interpolate/cubic_spline_test.cpp b/cpp/tests/interpolate/cubic_spline_test.cpp deleted file mode 100644 index cb0299c0e..000000000 --- a/cpp/tests/interpolate/cubic_spline_test.cpp +++ /dev/null @@ -1,273 +0,0 @@ -/* - * Copyright (c) 2020-2021, 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 -#include -#include -#include - -#include -#include -#include -#include - -#include -#include -#include - -#include -#include -#include - -struct CubicSplineTest : public cudf::test::BaseFixture { -}; - -TEST_F(CubicSplineTest, test_coefficients_single) -{ - cudf::test::fixed_width_column_wrapper t_column{{0, 1, 2, 3, 4}}; - cudf::test::fixed_width_column_wrapper y_column{{3, 2, 3, 4, 3}}; - cudf::test::fixed_width_column_wrapper ids_column{{0, 0}}; - cudf::test::fixed_width_column_wrapper prefix_column{{0, 5}}; - - auto splines = cuspatial::cubicspline_coefficients(t_column, y_column, ids_column, prefix_column); - - cudf::test::fixed_width_column_wrapper detail3_expected{{0.5, -0.5, -0.5, 0.5}}; - cudf::test::fixed_width_column_wrapper detail2_expected{{0.0, 3.0, 3.0, -6.0}}; - cudf::test::fixed_width_column_wrapper detail1_expected{{-1.5, -4.5, -4.5, 22.5}}; - cudf::test::fixed_width_column_wrapper detail0_expected{{3.0, 4.0, 4.0, -23.0}}; - - auto expected = - cudf::table_view{{detail3_expected, detail2_expected, detail1_expected, detail0_expected}}; - - CUDF_TEST_EXPECT_TABLES_EQUIVALENT(*splines, expected); -} - -TEST_F(CubicSplineTest, test_coefficients_full) -{ - cudf::test::fixed_width_column_wrapper t_column{ - {0, 1, 2, 3, 4, 0, 1, 2, 3, 4, 0, 1, 2, 3, 4}}; - cudf::test::fixed_width_column_wrapper y_column{ - {3, 2, 3, 4, 3, 3, 2, 3, 4, 3, 3, 2, 3, 4, 3}}; - cudf::test::fixed_width_column_wrapper ids_column{{0, 0, 1, 2}}; - cudf::test::fixed_width_column_wrapper prefix_column{{0, 5, 10, 15}}; - - auto splines = cuspatial::cubicspline_coefficients(t_column, y_column, ids_column, prefix_column); - - cudf::test::fixed_width_column_wrapper detail3_expected{ - {0.5, -0.5, -0.5, 0.5, 0.5, -0.5, -0.5, 0.5, 0.5, -0.5, -0.5, 0.5}}; - cudf::test::fixed_width_column_wrapper detail2_expected{ - {0.0, 3.0, 3.0, -6.0, 0.0, 3.0, 3.0, -6.0, 0.0, 3.0, 3.0, -6.0}}; - cudf::test::fixed_width_column_wrapper detail1_expected{ - {-1.5, -4.5, -4.5, 22.5, -1.5, -4.5, -4.5, 22.5, -1.5, -4.5, -4.5, 22.5}}; - cudf::test::fixed_width_column_wrapper detail0_expected{ - {3.0, 4.0, 4.0, -23.0, 3.0, 4.0, 4.0, -23.0, 3.0, 4.0, 4.0, -23.0}}; - - auto expected = - cudf::table_view{{detail3_expected, detail2_expected, detail1_expected, detail0_expected}}; - - CUDF_TEST_EXPECT_TABLES_EQUIVALENT(*splines, expected); -} - -TEST_F(CubicSplineTest, test_interpolate_between_control_points) -{ - cudf::test::fixed_width_column_wrapper t_column{ - {0, 1, 2, 3, 4, 0, 1, 2, 3, 4, 0, 1, 2, 3, 4}}; - cudf::test::fixed_width_column_wrapper new_column{ - {0.0, 0.5, 1.0, 1.5, 2.0, 2.5, 3.0, 3.5, 4.0, 0.0, 0.5, 1.0, 1.5, 2.0, - 2.5, 3.0, 3.5, 4.0, 0.0, 0.5, 1.0, 1.5, 2.0, 2.5, 3.0, 3.5, 4.0}}; - cudf::test::fixed_width_column_wrapper x_column{ - {3, 2, 3, 4, 3, 3, 2, 3, 4, 3, 3, 2, 3, 4, 3}}; - cudf::test::fixed_width_column_wrapper ids_column{{0, 0, 1, 2}}; - cudf::test::fixed_width_column_wrapper prefix_column{{0, 5, 10, 15}}; - cudf::test::fixed_width_column_wrapper old_point_ids_column{ - {0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 2, 2, 2, 2, 2}}; - cudf::test::fixed_width_column_wrapper new_point_ids_column{ - {0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 2, 2, 2, 2, 2, 2, 2, 2, 2}}; - - auto splines = cuspatial::cubicspline_coefficients(t_column, x_column, ids_column, prefix_column); - - auto interpolants_new = cuspatial::cubicspline_interpolate( - new_column, new_point_ids_column, prefix_column, t_column, *splines); - - auto gather_map = cudf::test::fixed_width_column_wrapper{ - 0, 2, 4, 6, 8, 9, 11, 13, 15, 17, 18, 20, 22, 24, 26}; - auto interpolants_gather = - cudf::gather(cudf::table_view{std::vector{interpolants_new->view()}}, - gather_map) - ->get_column(0); - - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( - interpolants_gather, - cudf::test::fixed_width_column_wrapper{{3, 2, 3, 4, 3, 3, 2, 3, 4, 3, 3, 2, 3, 4, 3}}); -} - -TEST_F(CubicSplineTest, test_interpolate_single) -{ - cudf::test::fixed_width_column_wrapper t_column{{0, 1, 2, 3, 4}}; - cudf::test::fixed_width_column_wrapper x_column{{3, 2, 3, 4, 3}}; - cudf::test::fixed_width_column_wrapper ids_column{{0, 0}}; - cudf::test::fixed_width_column_wrapper prefix_column{{0, 5}}; - cudf::test::fixed_width_column_wrapper point_ids_column{{0, 0, 0, 0, 0}}; - - auto splines = cuspatial::cubicspline_coefficients(t_column, x_column, ids_column, prefix_column); - - auto interpolants = cuspatial::cubicspline_interpolate( - t_column, point_ids_column, prefix_column, t_column, *splines); - - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( - *interpolants, cudf::test::fixed_width_column_wrapper{{3, 2, 3, 4, 3}}); -} - -TEST_F(CubicSplineTest, test_interpolate_at_control_points_full) -{ - cudf::test::fixed_width_column_wrapper t_column{ - {0, 1, 2, 3, 4, 0, 1, 2, 3, 4, 0, 1, 2, 3, 4}}; - cudf::test::fixed_width_column_wrapper x_column{ - {3, 2, 3, 4, 3, 3, 2, 3, 4, 3, 3, 2, 3, 4, 3}}; - cudf::test::fixed_width_column_wrapper ids_column{{0, 0, 1, 2}}; - cudf::test::fixed_width_column_wrapper prefix_column{{0, 5, 10, 15}}; - cudf::test::fixed_width_column_wrapper point_ids_column{ - {0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 2, 2, 2, 2, 2}}; - - auto splines = cuspatial::cubicspline_coefficients(t_column, x_column, ids_column, prefix_column); - - auto interpolants = cuspatial::cubicspline_interpolate( - t_column, point_ids_column, prefix_column, t_column, *splines); - - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( - *interpolants, - cudf::test::fixed_width_column_wrapper{{3, 2, 3, 4, 3, 3, 2, 3, 4, 3, 3, 2, 3, 4, 3}}); -} - -TEST_F(CubicSplineTest, test_parallel_search_single) -{ - cudf::test::fixed_width_column_wrapper short_single{{0, 1, 2, 3, 4}}; - cudf::test::fixed_width_column_wrapper point_ids_column{{0, 0, 0, 0, 0}}; - cudf::test::fixed_width_column_wrapper prefix_column{{0, 5}}; - - auto indexes = cuspatial::detail::find_coefficient_indices(short_single, - point_ids_column, - prefix_column, - short_single, - rmm::cuda_stream_default, - this->mr()); - - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*indexes, - cudf::test::fixed_width_column_wrapper{{0, 1, 2, 3, 3}}); -} - -TEST_F(CubicSplineTest, test_parallel_search_triple) -{ - cudf::test::fixed_width_column_wrapper short_triple{ - {0, 1, 2, 3, 4, 0, 1, 2, 3, 4, 0, 1, 2, 3, 4}}; - cudf::test::fixed_width_column_wrapper point_ids_column{ - {0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 2, 2, 2, 2, 2}}; - cudf::test::fixed_width_column_wrapper prefix_column{{0, 5, 10, 15}}; - - auto indexes = cuspatial::detail::find_coefficient_indices(short_triple, - point_ids_column, - prefix_column, - short_triple, - rmm::cuda_stream_default, - this->mr()); - - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( - *indexes, - cudf::test::fixed_width_column_wrapper{{0, 1, 2, 3, 3, 4, 5, 6, 7, 7, 8, 9, 10, 11, 11}}); -} - -TEST_F(CubicSplineTest, test_parallel_search_middle_single) -{ - cudf::test::fixed_width_column_wrapper short_single{{0, 1, 2, 3, 4}}; - cudf::test::fixed_width_column_wrapper long_single{ - {0.0, 0.5, 1.0, 1.5, 2.0, 2.5, 3.0, 3.5, 4.0}}; - cudf::test::fixed_width_column_wrapper point_ids_column{{0, 0, 0, 0, 0, 0, 0, 0, 0}}; - cudf::test::fixed_width_column_wrapper prefix_column{{0, 5}}; - - auto indexes = cuspatial::detail::find_coefficient_indices(long_single, - point_ids_column, - prefix_column, - short_single, - rmm::cuda_stream_default, - this->mr()); - - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( - *indexes, cudf::test::fixed_width_column_wrapper{{0, 0, 1, 1, 2, 2, 3, 3, 3}}); -} - -TEST_F(CubicSplineTest, test_parallel_search_middle_triple) -{ - cudf::test::fixed_width_column_wrapper short_triple{ - {0, 1, 2, 3, 4, 0, 1, 2, 3, 4, 0, 1, 2, 3, 4}}; - cudf::test::fixed_width_column_wrapper long_triple{ - {0.0, 0.5, 1.0, 1.5, 2.0, 2.5, 3.0, 3.5, 4.0, 0.0, 0.5, 1.0, 1.5, 2.0, - 2.5, 3.0, 3.5, 4.0, 0.0, 0.5, 1.0, 1.5, 2.0, 2.5, 3.0, 3.5, 4.0}}; - cudf::test::fixed_width_column_wrapper point_ids_column{ - {0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 2, 2, 2, 2, 2, 2, 2, 2, 2}}; - cudf::test::fixed_width_column_wrapper prefix_column{{0, 5, 10, 15}}; - - auto indexes = cuspatial::detail::find_coefficient_indices(long_triple, - point_ids_column, - prefix_column, - short_triple, - rmm::cuda_stream_default, - this->mr()); - - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( - *indexes, - cudf::test::fixed_width_column_wrapper{ - {0, 0, 1, 1, 2, 2, 3, 3, 3, 4, 4, 5, 5, 6, 6, 7, 7, 7, 8, 8, 9, 9, 10, 10, 11, 11, 11}}); -} - -TEST_F(CubicSplineTest, test_parallel_search_single_end_values) -{ - cudf::test::fixed_width_column_wrapper short_triple{{0, 1, 2, 3, 4}}; - cudf::test::fixed_width_column_wrapper long_triple{{4.0, 4.1, 4.5, 5.0, 10000.0}}; - cudf::test::fixed_width_column_wrapper point_ids_column{{0, 0, 0, 0, 0}}; - cudf::test::fixed_width_column_wrapper prefix_column{{0, 5}}; - - auto indexes = cuspatial::detail::find_coefficient_indices(long_triple, - point_ids_column, - prefix_column, - short_triple, - rmm::cuda_stream_default, - this->mr()); - - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*indexes, - cudf::test::fixed_width_column_wrapper{{3, 3, 3, 3, 3}}); -} - -TEST_F(CubicSplineTest, test_parallel_search_triple_end_values) -{ - cudf::test::fixed_width_column_wrapper short_triple{ - {0, 1, 2, 3, 4, 0, 1, 2, 3, 4, 0, 1, 2, 3, 4}}; - cudf::test::fixed_width_column_wrapper long_triple{ - {4.0, 4.1, 4.5, 5.0, 10000.0, 4.0, 4.1, 4.5, 5.0, 10000.0, 4.0, 4.1, 4.5, 5.0, 10000.0}}; - cudf::test::fixed_width_column_wrapper point_ids_column{ - {0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 2, 2, 2, 2, 2}}; - cudf::test::fixed_width_column_wrapper prefix_column{{0, 5}}; - - auto indexes = cuspatial::detail::find_coefficient_indices(long_triple, - point_ids_column, - prefix_column, - short_triple, - rmm::cuda_stream_default, - this->mr()); - - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*indexes, - cudf::test::fixed_width_column_wrapper{ - {3, 3, 3, 3, 3, 7, 7, 7, 7, 7, 11, 11, 11, 11, 11}}); -} diff --git a/dependencies.yaml b/dependencies.yaml index da4634061..316b96354 100644 --- a/dependencies.yaml +++ b/dependencies.yaml @@ -114,28 +114,18 @@ dependencies: cuda: "11.8" packages: - cudatoolkit=11.8 - - libcusparse-dev=11.7.5.86 - - libcusparse=11.7.5.86 - matrix: cuda: "11.5" packages: - cudatoolkit=11.5 - - libcusparse-dev>=11.7.0.31,<=11.7.0.107 - - libcusparse>=11.7.0.31,<=11.7.0.107 - matrix: cuda: "11.4" packages: - cudatoolkit=11.4 - - &libcusparse_dev114 libcusparse-dev>=11.6.0.43,<=11.6.0.120 - - &libcusparse114 libcusparse>=11.6.0.43,<=11.6.0.120 - matrix: cuda: "11.2" packages: - cudatoolkit=11.2 - # The NVIDIA channel doesn't publish pkgs older than 11.4 for these libs, - # so 11.2 uses 11.4 packages (the oldest available). - - *libcusparse_dev114 - - *libcusparse114 develop: common: - output_types: [conda, requirements] diff --git a/docs/source/api_docs/index.rst b/docs/source/api_docs/index.rst index b124e678e..15ff53c6f 100644 --- a/docs/source/api_docs/index.rst +++ b/docs/source/api_docs/index.rst @@ -10,6 +10,5 @@ This page provides a list of all publicly accessible modules, methods and classe spatial trajectory - interpolation geopandas_compatibility io diff --git a/docs/source/api_docs/interpolation.rst b/docs/source/api_docs/interpolation.rst deleted file mode 100644 index d65a3a421..000000000 --- a/docs/source/api_docs/interpolation.rst +++ /dev/null @@ -1,10 +0,0 @@ -Interpolation -------------- - -Functions to interpolate curves and sample new points. - -.. currentmodule:: cuspatial - -.. autoclass:: CubicSpline -.. automethod:: CubicSpline.__init__ -.. automethod:: CubicSpline.__call__ diff --git a/python/cuspatial/cuspatial/__init__.py b/python/cuspatial/cuspatial/__init__.py index 492cb73c0..bb9a237c9 100644 --- a/python/cuspatial/cuspatial/__init__.py +++ b/python/cuspatial/cuspatial/__init__.py @@ -1,8 +1,6 @@ from ._version import get_versions -from .core import interpolate from .core.geodataframe import GeoDataFrame from .core.geoseries import GeoSeries -from .core.interpolate import CubicSpline from .core.spatial import ( directed_hausdorff_distance, haversine_distance, diff --git a/python/cuspatial/cuspatial/_lib/CMakeLists.txt b/python/cuspatial/cuspatial/_lib/CMakeLists.txt index 6e521ccb5..4ffbfc7dc 100644 --- a/python/cuspatial/cuspatial/_lib/CMakeLists.txt +++ b/python/cuspatial/cuspatial/_lib/CMakeLists.txt @@ -15,7 +15,6 @@ set(cython_sources distance.pyx hausdorff.pyx - interpolate.pyx intersection.pyx nearest_points.pyx point_in_polygon.pyx diff --git a/python/cuspatial/cuspatial/_lib/cpp/interpolate.pxd b/python/cuspatial/cuspatial/_lib/cpp/interpolate.pxd deleted file mode 100644 index 7280c7859..000000000 --- a/python/cuspatial/cuspatial/_lib/cpp/interpolate.pxd +++ /dev/null @@ -1,25 +0,0 @@ -# Copyright (c) 2020, NVIDIA CORPORATION. - -from libcpp.memory cimport unique_ptr - -from cudf._lib.column cimport column, column_view -from cudf._lib.cpp.table.table cimport table, table_view - - -cdef extern from "cuspatial/cubic_spline.hpp" namespace "cuspatial" nogil: - cdef unique_ptr[table] cubicspline_coefficients \ - "cuspatial::cubicspline_coefficients" ( - const column_view & t, - const column_view & x, - const column_view & ids, - const column_view & prefix_sums - ) except + - - cdef unique_ptr[column] cubicspline_interpolate \ - "cuspatial::cubicspline_interpolate" ( - const column_view & p, - const column_view & ids, - const column_view & prefix_sums, - const column_view & old_t, - const table_view & coefficients - ) except + diff --git a/python/cuspatial/cuspatial/_lib/interpolate.pyx b/python/cuspatial/cuspatial/_lib/interpolate.pyx deleted file mode 100644 index 43392e10a..000000000 --- a/python/cuspatial/cuspatial/_lib/interpolate.pyx +++ /dev/null @@ -1,62 +0,0 @@ -# Copyright (c) 2019-2020, NVIDIA CORPORATION. - -from libcpp.memory cimport unique_ptr -from libcpp.utility cimport move - -from cudf._lib.column cimport Column, column -from cudf._lib.cpp.table.table cimport table -from cudf._lib.utils cimport data_from_unique_ptr, table_view_from_table - -from cuspatial._lib.cpp.interpolate cimport ( - cubicspline_coefficients as cpp_cubicspline_coefficients, - cubicspline_interpolate as cpp_cubicspline_interpolate, -) - - -cpdef cubicspline_coefficients( - Column t, - Column x, - Column ids, - Column prefixes -): - t_v = t.view() - x_v = x.view() - ids_v = ids.view() - prefixes_v = prefixes.view() - cdef unique_ptr[table] c_result - with nogil: - c_result = move( - cpp_cubicspline_coefficients( - t_v, - x_v, - ids_v, - prefixes_v - ) - ) - return data_from_unique_ptr(move(c_result), ["d3", "d2", "d1", "d0"]) - -cpdef cubicspline_interpolate( - Column points, - Column ids, - Column prefixes, - Column original_t, - object coefficients -): - p_v = points.view() - ids_v = ids.view() - prefixes_v = prefixes.view() - original_t_v = original_t.view() - coefs_v = table_view_from_table(coefficients, ignore_index=True) - cdef unique_ptr[column] c_result - with nogil: - c_result = move( - cpp_cubicspline_interpolate( - p_v, - ids_v, - prefixes_v, - original_t_v, - coefs_v - ) - ) - result = Column.from_unique_ptr(move(c_result)) - return result diff --git a/python/cuspatial/cuspatial/core/interpolate.py b/python/cuspatial/cuspatial/core/interpolate.py deleted file mode 100644 index 72a67d3c0..000000000 --- a/python/cuspatial/cuspatial/core/interpolate.py +++ /dev/null @@ -1,203 +0,0 @@ -# Copyright (c) 2020, NVIDIA CORPORATION. -import warnings - -import cupy as cp -import numpy as np - -from cudf import DataFrame, Series - -from cuspatial._lib.interpolate import ( - cubicspline_coefficients, - cubicspline_interpolate, -) - - -def _cubic_spline_coefficients(x, y, ids, prefix_sums): - x_c = x._column - y_c = y._column - ids_c = ids._column - prefix_c = prefix_sums._column - return DataFrame._from_data( - *cubicspline_coefficients(x_c, y_c, ids_c, prefix_c) - ) - - -def _cubic_spline_fit(points, points_ids, prefixes, original_t, c): - points_c = points._column - points_ids_c = points_ids._column - prefixes_c = prefixes._column - original_t_c = original_t._column - result_column = cubicspline_interpolate( - points_c, points_ids_c, prefixes_c, original_t_c, c - ) - return result_column - - -class CubicSpline: - """ - Fits each column of the input Series `y` to a hermetic cubic spline. - - ``cuspatial.CubicSpline`` supports basic usage identical to - scipy.interpolate.CubicSpline:: - - curve = cuspatial.CubicSpline(x, y) - new_points = curve(np.linspace(x.min, x.max, 50)) - - Parameters - ---------- - x : cudf.Series - 1-D array containing values of the independent variable. - Values must be real, finite and in strictly increasing order. - y : cudf.Series - Array containing values of the dependent variable. - ids (Optional) : cudf.Series - ids of each spline - size (Optional) : cudf.Series - fixed size of each spline - offset (Optional) : cudf.Series - alternative to `size`, allows splines of varying - length. Not yet fully supported. - - Returns - ------- - CubicSpline : callable `o` - ``o.c`` contains the coefficients that can be used to compute new - points along the spline fitting the original ``t`` data. ``o(n)`` - interpolates the spline coordinates along new input values ``n``. - - Note - ---- - cuSpatial will outperform scipy when many splines are - fit simultaneously. Data must be arranged in a structure of arrays (SoA) - format, and the exclusive `offset` of the separate curves must also be - passed to the function. - - Example - ------- - # The following example only serves to demonstrate the - # cuspatial.CubicSpline API. cuSpatial does not out perform scipy when - # fitting a small number of curves as shown. - >>> import cuspatial, cudf - >>> import numpy as np - >>> def f(x): - ... return x**3+4*x**2-7*x+1 - ... - >>> x = np.array([0, 1, 2, 3, 4] + [10, 11, 12, 13, 14]) - >>> y = map(f, x) - >>> x, y = cudf.Series(x, dtype='f4'), cudf.Series(y, dtype='f4') - >>> offset = cudf.Series([0, 5, 10], dtype='i4') - >>> curve = cuspatial.CubicSpline(x, y, offset=offset) - UserWarning: fitting a small number of curves on device may suffer from - kernel launch overheads. - >>> x_sample1 = [*np.arange(-1, 5, 0.3)] - >>> x_sample2 = [*np.arange(11, 12, 0.1)] - >>> curve_ids = cudf.Series( - ... [0]*len(x_sample1) + [1]*len(x_sample2), dtype='i4' - ... ) - >>> x_sample = cudf.Series(x_sample1 + x_sample2, dtype='f4') - >>> y_sampled = curve(x_sample, curve_ids) - >>> y_sampled1 = y_sampled[0:len(x_sample1)] - >>> y_sampled2 = y_sampled[len(x_sample1):] - """ - - def __init__(self, x, y, ids=None, size=None, offset=None): - # error protections: - if len(x) < 5: - raise ValueError( - "Use of GPU cubic spline requires splines of length > 4" - ) - if not isinstance(x, Series): - raise TypeError( - "Error: input independent vars must be cudf Series" - ) - if not isinstance(y, (Series, DataFrame)): - raise TypeError( - "Error: input dependent vars must be cudf Series or DataFrame" - ) - if not len(x) == len(y): - raise TypeError( - "Error: dependent and independent vars have different length" - ) - if ids is None: - self.ids = Series([0, 0]).astype("int32") - else: - if not isinstance(ids, Series): - raise TypeError("cuspatial.CubicSpline requires a cudf.Series") - if not ids.dtype == np.int32: - raise TypeError("Error: int32 only supported at this time.") - self.ids = ids - self.size = size if size is not None else len(x) - if not isinstance(self.size, int): - raise TypeError("Error: size must be an integer") - if not ((len(x) % self.size) == 0): - raise ValueError( - "Error: length of input is not a multiple of size" - ) - if not isinstance(x, Series): - raise TypeError("cuspatial.CubicSpline requires a cudf.Series") - if not x.dtype == np.float32: - raise TypeError("Error: float32 only supported at this time.") - if not isinstance(y, Series): - raise TypeError("cuspatial.CubicSpline requires a cudf.Series") - if not y.dtype == np.float32: - raise TypeError("Error: float32 only supported at this time.") - self.x = x - self.y = y - if offset is None: - self.offset = Series( - cp.arange((len(x) / self.size) + 1) * self.size - ).astype("int32") - else: - if not isinstance(offset, Series): - raise TypeError("cuspatial.CubicSpline requires a cudf.Series") - if not offset.dtype == np.int32: - raise TypeError("Error: int32 only supported at this time.") - self.offset = offset - - if self.offset.size < 15: - warnings.warn( - "Fitting a small number of curves on " - "device may suffer from kernel launch overheads." - ) - - self.c = self._compute_coefficients() - - def _compute_coefficients(self): - """ - Utility method used by __init__ once members have been initialized. - """ - if isinstance(self.y, Series): - return _cubic_spline_coefficients( - self.x, self.y, self.ids, self.offset - ) - else: - c = {} - for col in self.y.columns: - c[col] = _cubic_spline_coefficients( - self.x, self.y, self.ids, self.offset - ) - return c - - def __call__(self, coordinates, groups=None): - """ - Interpolates new input values `coordinates` using the `.c` DataFrame - or map of DataFrames. - """ - if isinstance(self.y, Series): - if groups is not None: - self.groups = groups.astype("int32") - else: - self.groups = Series( - cp.repeat(cp.array(0), len(coordinates)) - ).astype("int32") - result = _cubic_spline_fit( - coordinates, self.groups, self.offset, self.x, self.c - ) - return Series(result) - else: - result = DataFrame() - for col in self.y.columns: - result[col] = Series( - _cubic_spline_fit(self.c[col], coordinates) - ) - return result diff --git a/python/cuspatial/cuspatial/tests/test_interpolate.py b/python/cuspatial/cuspatial/tests/test_interpolate.py deleted file mode 100644 index 77f0d7238..000000000 --- a/python/cuspatial/cuspatial/tests/test_interpolate.py +++ /dev/null @@ -1,250 +0,0 @@ -# Copyright (c) 2020, NVIDIA CORPORATION. - -import cupy as cp -import numpy as np -import pytest - -import cudf - -import cuspatial - - -def test_errors(): - # t and y must have the same length - with pytest.raises(TypeError): - cuspatial.interpolate.CubicSpline( - cudf.Series([0, 0, 0, 0, 0]).astype("float32"), - cudf.Series([0]).astype("float32"), - cudf.Series([0]).astype("int32"), - cudf.Series([0, 1]).astype("int32"), - ) - # length must not be zero - with pytest.raises(ZeroDivisionError): - cuspatial.interpolate.CubicSpline( - cudf.Series([0, 0, 0, 0, 0]).astype("float32"), - cudf.Series([0, 0, 0, 0, 0]).astype("float32"), - cudf.Series([0]).astype("int32"), - 0, - ) - # Length must be greater than 4 - with pytest.raises(ValueError): - cuspatial.interpolate.CubicSpline( - cudf.Series([0]).astype("float32"), - cudf.Series([0]).astype("float32"), - cudf.Series([0]).astype("int32"), - 1, - ) - - -def test_class_coefs(): - t = cudf.Series([0, 1, 2, 3, 4]).astype("float32") - x = cudf.Series([3, 2, 3, 4, 3]).astype("float32") - g = cuspatial.interpolate.CubicSpline(t, x) - cudf.testing.assert_frame_equal( - g.c, - cudf.DataFrame( - { - "d3": [0.5, -0.5, -0.5, 0.5], - "d2": [0, 3, 3, -6], - "d1": [-1.5, -4.5, -4.5, 22.5], - "d0": [3, 4, 4, -23], - } - ), - check_dtype=False, - ) - - -def test_min(): - result = cuspatial.CubicSpline( - cudf.Series([0, 1, 2, 3, 4]).astype("float32"), - cudf.Series([3, 2, 3, 4, 3]).astype("float32"), - cudf.Series([0, 5]).astype("int32"), - ) - cudf.testing.assert_frame_equal( - result.c, - cudf.DataFrame( - { - "d3": [0.5, -0.5, -0.5, 0.5], - "d2": [0, 3, 3, -6], - "d1": [-1.5, -4.5, -4.5, 22.5], - "d0": [3, 4, 4, -23], - } - ), - check_dtype=False, - ) - - -def test_cusparse(): - result = cuspatial.CubicSpline( - cudf.Series([0, 1, 2, 3, 4, 0, 1, 2, 3, 4, 0, 1, 2, 3, 4]).astype( - "float32" - ), - cudf.Series([3, 2, 3, 4, 3, 3, 2, 3, 4, 3, 3, 2, 3, 4, 3]).astype( - "float32" - ), - offset=cudf.Series([0, 5, 10, 15]).astype("int32"), - ) - cudf.testing.assert_frame_equal( - result.c, - cudf.DataFrame( - { - "d3": [ - 0.5, - -0.5, - -0.5, - 0.5, - 0.5, - -0.5, - -0.5, - 0.5, - 0.5, - -0.5, - -0.5, - 0.5, - ], - "d2": [0, 3, 3, -6, 0, 3, 3, -6, 0, 3, 3, -6], - "d1": [ - -1.5, - -4.5, - -4.5, - 22.5, - -1.5, - -4.5, - -4.5, - 22.5, - -1.5, - -4.5, - -4.5, - 22.5, - ], - "d0": [3, 4, 4, -23, 3, 4, 4, -23, 3, 4, 4, -23], - } - ), - check_dtype=False, - ) - - -def test_class_interpolation_length_five(): - t = cudf.Series([0, 1, 2, 3, 4]).astype("float32") - x = cudf.Series([3, 2, 3, 4, 3]).astype("float32") - g = cuspatial.interpolate.CubicSpline(t, x) - cudf.testing.assert_series_equal(g(t), x) - - -def test_class_interpolation_length_six(): - t = cudf.Series([0, 1, 2, 3, 4, 5]).astype("float32") - x = cudf.Series([3, 2, 3, 4, 3, 4]).astype("float32") - g = cuspatial.interpolate.CubicSpline(t, x) - cudf.testing.assert_series_equal(g(t), x) - - -def test_class_interpolation_length_six_splits(): - t = cudf.Series([0, 1, 2, 3, 4, 5]).astype("float32") - x = cudf.Series([3, 2, 3, 4, 3, 4]).astype("float32") - g = cuspatial.interpolate.CubicSpline(t, x) - split_t = cudf.Series(np.linspace(0, 5, 11), dtype="float32") - cudf.testing.assert_series_equal( - g(split_t)[t * 2].reset_index(drop=True), x - ) - - -def test_class_triple(): - t = cudf.Series([0, 1, 2, 3, 4, 0, 1, 2, 3, 4, 0, 1, 2, 3, 4]).astype( - "float32" - ) - x = cudf.Series([3, 2, 3, 4, 3, 3, 2, 3, 4, 3, 3, 2, 3, 4, 3]).astype( - "float32" - ) - prefixes = cudf.Series([0, 5, 10, 15]).astype("int32") - g = cuspatial.interpolate.CubicSpline(t, x, offset=prefixes) - groups = cudf.Series( - np.ravel(np.array([np.repeat(0, 5), np.repeat(1, 5), np.repeat(2, 5)])) - ) - cudf.testing.assert_series_equal(g(t, groups=groups), x) - - -def test_class_triple_six(): - t = cudf.Series( - [0, 1, 2, 3, 4, 5, 0, 1, 2, 3, 4, 5, 0, 1, 2, 3, 4, 5] - ).astype("float32") - x = cudf.Series( - [3, 2, 3, 4, 3, 1, 3, 2, 3, 4, 3, 1, 3, 2, 3, 4, 3, 1] - ).astype("float32") - prefixes = cudf.Series([0, 6, 12, 18]).astype("int32") - g = cuspatial.interpolate.CubicSpline(t, x, offset=prefixes) - groups = cudf.Series( - np.ravel(np.array([np.repeat(0, 6), np.repeat(1, 6), np.repeat(2, 6)])) - ) - cudf.testing.assert_series_equal(g(t, groups=groups), x) - - -def test_class_triple_six_splits(): - t = cudf.Series( - [0, 1, 2, 3, 4, 5, 0, 1, 2, 3, 4, 5, 0, 1, 2, 3, 4, 5] - ).astype("float32") - x = cudf.Series( - [3, 2, 3, 4, 3, 1, 3, 2, 3, 4, 3, 1, 3, 2, 3, 4, 3, 1] - ).astype("float32") - prefixes = cudf.Series([0, 6, 12, 18]).astype("int32") - g = cuspatial.interpolate.CubicSpline(t, x, offset=prefixes) - groups = cudf.Series( - np.ravel( - np.array([np.repeat(0, 12), np.repeat(1, 12), np.repeat(2, 12)]) - ) - ) - split_t = cudf.Series( - np.ravel( - ( - np.linspace(0, 5, 11), - np.linspace(0, 5, 11), - np.linspace(0, 5, 11), - ) - ), - dtype="float32", - ) - split_t_ind = [ - 0, - 2, - 4, - 6, - 8, - 10, - 11, - 13, - 15, - 17, - 19, - 21, - 22, - 24, - 26, - 28, - 30, - 32, - ] - cudf.testing.assert_series_equal( - g(split_t, groups=groups)[split_t_ind].reset_index(drop=True), x - ) - - -def test_class_new_interpolation(): - t = cudf.Series(np.hstack((np.arange(5),) * 3)).astype("float32") - y = cudf.Series([3, 2, 3, 4, 3, 3, 2, 3, 4, 3, 3, 2, 3, 4, 3]).astype( - "float32" - ) - prefix_sum = cudf.Series(cp.arange(4) * 5).astype("int32") - new_samples = cudf.Series(np.hstack((np.linspace(0, 4, 9),) * 3)).astype( - "float32" - ) - curve = cuspatial.CubicSpline(t, y, offset=prefix_sum) - new_x = cudf.Series(np.repeat(np.arange(0, 3), 9)).astype("int32") - old_x = cudf.Series(np.repeat(np.arange(0, 3), 5)).astype("int32") - new_points = curve(new_samples, groups=new_x) - old_points = curve(t, groups=old_x) - new_points_at_control_points = new_points[ - 0, 2, 4, 6, 8, 9, 11, 13, 15, 17, 18, 20, 22, 24, 26 - ] - new_points_at_control_points.index = cudf.RangeIndex( - 0, len(new_points_at_control_points) - ) - cudf.testing.assert_series_equal(new_points_at_control_points, old_points)