Implement dpnp.interp() (#2417)

This PR suggests adding an implementation of `dpnp.interp()` with
sycl_kernel including support for all function arguments.

Author: vlad-perevezentsev

CHANGELOG.md

@@ -16,6 +16,7 @@ This release achieves 100% compliance with Python Array API specification (revis
 * Added implementation of `dpnp.bartlett` [#2366](https://github.com/IntelPython/dpnp/pull/2366)
 * Added implementation of `dpnp.convolve` [#2205](https://github.com/IntelPython/dpnp/pull/2205)
 * Added implementation of `dpnp.kaiser` [#2387](https://github.com/IntelPython/dpnp/pull/2387)
+* Added implementation of `dpnp.interp` [#2417](https://github.com/IntelPython/dpnp/pull/2417)
 
 ### Changed
 

dpnp/backend/extensions/common/ext/common.hpp

@@ -106,6 +106,27 @@ struct IsNan
     }
 };
 
+template <typename T, bool hasValueType>
+struct value_type_of_impl;
+
+template <typename T>
+struct value_type_of_impl<T, false>
+{
+    using type = T;
+};
+
+template <typename T>
+struct value_type_of_impl<T, true>
+{
+    using type = typename T::value_type;
+};
+
+template <typename T>
+using value_type_of = value_type_of_impl<T, type_utils::is_complex_v<T>>;
+
+template <typename T>
+using value_type_of_t = typename value_type_of<T>::type;
+
 size_t get_max_local_size(const sycl::device &device);
 size_t get_max_local_size(const sycl::device &device,
                           int cpu_local_size_limit,

dpnp/backend/extensions/common/ext/details/validation_utils_internal.hpp

@@ -23,9 +23,17 @@
 // THE POSSIBILITY OF SUCH DAMAGE.
 //*****************************************************************************
 
+#include <pybind11/numpy.h>
+#include <pybind11/pybind11.h>
+
+#include "ext/common.hpp"
+
 #include "ext/validation_utils.hpp"
 #include "utils/memory_overlap.hpp"
 
+namespace td_ns = dpctl::tensor::type_dispatch;
+namespace common = ext::common;
+
 namespace ext::validation
 {
 inline sycl::queue get_queue(const std::vector<array_ptr> &inputs,
@@ -137,6 +145,15 @@ inline void check_num_dims(const array_ptr &arr,
     }
 }
 
+inline void check_num_dims(const std::vector<array_ptr> &arrays,
+                           const size_t ndim,
+                           const array_names &names)
+{
+    for (const auto &arr : arrays) {
+        check_num_dims(arr, ndim, names);
+    }
+}
+
 inline void check_max_dims(const array_ptr &arr,
                            const size_t max_ndim,
                            const array_names &names)
@@ -163,6 +180,103 @@ inline void check_size_at_least(const array_ptr &arr,
     }
 }
 
+inline void check_has_dtype(const array_ptr &arr,
+                            const typenum_t dtype,
+                            const array_names &names)
+{
+    if (arr == nullptr) {
+        return;
+    }
+
+    auto array_types = td_ns::usm_ndarray_types();
+    int array_type_id = array_types.typenum_to_lookup_id(arr->get_typenum());
+    int expected_type_id = static_cast<int>(dtype);
+
+    if (array_type_id != expected_type_id) {
+        py::dtype actual_dtype = common::dtype_from_typenum(array_type_id);
+        py::dtype dtype_py = common::dtype_from_typenum(expected_type_id);
+
+        std::string msg = "Array " + name_of(arr, names) + " must have dtype " +
+                          std::string(py::str(dtype_py)) + ", but got " +
+                          std::string(py::str(actual_dtype));
+
+        throw py::value_error(msg);
+    }
+}
+
+inline void check_same_dtype(const array_ptr &arr1,
+                             const array_ptr &arr2,
+                             const array_names &names)
+{
+    if (arr1 == nullptr || arr2 == nullptr) {
+        return;
+    }
+
+    auto array_types = td_ns::usm_ndarray_types();
+    int first_type_id = array_types.typenum_to_lookup_id(arr1->get_typenum());
+    int second_type_id = array_types.typenum_to_lookup_id(arr2->get_typenum());
+
+    if (first_type_id != second_type_id) {
+        py::dtype first_dtype = common::dtype_from_typenum(first_type_id);
+        py::dtype second_dtype = common::dtype_from_typenum(second_type_id);
+
+        std::string msg = "Arrays " + name_of(arr1, names) + " and " +
+                          name_of(arr2, names) +
+                          " must have the same dtype, but got " +
+                          std::string(py::str(first_dtype)) + " and " +
+                          std::string(py::str(second_dtype));
+
+        throw py::value_error(msg);
+    }
+}
+
+inline void check_same_dtype(const std::vector<array_ptr> &arrays,
+                             const array_names &names)
+{
+    if (arrays.empty()) {
+        return;
+    }
+
+    const auto *first = arrays[0];
+    for (size_t i = 1; i < arrays.size(); ++i) {
+        check_same_dtype(first, arrays[i], names);
+    }
+}
+
+inline void check_same_size(const array_ptr &arr1,
+                            const array_ptr &arr2,
+                            const array_names &names)
+{
+    if (arr1 == nullptr || arr2 == nullptr) {
+        return;
+    }
+
+    auto size1 = arr1->get_size();
+    auto size2 = arr2->get_size();
+
+    if (size1 != size2) {
+        std::string msg =
+            "Arrays " + name_of(arr1, names) + " and " + name_of(arr2, names) +
+            " must have the same size, but got " + std::to_string(size1) +
+            " and " + std::to_string(size2);
+
+        throw py::value_error(msg);
+    }
+}
+
+inline void check_same_size(const std::vector<array_ptr> &arrays,
+                            const array_names &names)
+{
+    if (arrays.empty()) {
+        return;
+    }
+
+    auto first = arrays[0];
+    for (size_t i = 1; i < arrays.size(); ++i) {
+        check_same_size(first, arrays[i], names);
+    }
+}
+
 inline void common_checks(const std::vector<array_ptr> &inputs,
                           const std::vector<array_ptr> &outputs,
                           const array_names &names)

dpnp/backend/extensions/common/ext/validation_utils.hpp

@@ -35,6 +35,7 @@ namespace ext::validation
 {
 using array_ptr = const dpctl::tensor::usm_ndarray *;
 using array_names = std::unordered_map<array_ptr, std::string>;
+using dpctl::tensor::type_dispatch::typenum_t;
 
 std::string name_of(const array_ptr &arr, const array_names &names);
 
@@ -56,6 +57,9 @@ void check_no_overlap(const std::vector<array_ptr> &inputs,
 void check_num_dims(const array_ptr &arr,
                     const size_t ndim,
                     const array_names &names);
+void check_num_dims(const std::vector<array_ptr> &arrays,
+                    const size_t ndim,
+                    const array_names &names);
 void check_max_dims(const array_ptr &arr,
                     const size_t max_ndim,
                     const array_names &names);
@@ -64,6 +68,20 @@ void check_size_at_least(const array_ptr &arr,
                          const size_t size,
                          const array_names &names);
 
+void check_has_dtype(const array_ptr &arr,
+                     const typenum_t dtype,
+                     const array_names &names);
+
+void check_same_dtype(const array_ptr &arr1,
+                      const array_ptr &arr2,
+                      const array_names &names);
+
+void check_same_size(const array_ptr &arr1,
+                     const array_ptr &arr2,
+                     const array_names &names);
+void check_same_size(const std::vector<array_ptr> &arrays,
+                     const array_names &names);
+
 void common_checks(const std::vector<array_ptr> &inputs,
                    const std::vector<array_ptr> &outputs,
                    const array_names &names);

dpnp/backend/extensions/ufunc/CMakeLists.txt

@@ -36,6 +36,7 @@ set(_elementwise_sources
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/gcd.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/heaviside.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/i0.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/interpolate.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/lcm.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/ldexp.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/logaddexp2.cpp
@@ -69,6 +70,7 @@ endif()
 set_target_properties(${python_module_name} PROPERTIES CMAKE_POSITION_INDEPENDENT_CODE ON)
 
 target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../)
+target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../common)
 
 target_include_directories(${python_module_name} PUBLIC ${Dpctl_INCLUDE_DIR})
 target_include_directories(${python_module_name} PUBLIC ${Dpctl_TENSOR_INCLUDE_DIR})

dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp

@@ -36,6 +36,7 @@
 #include "gcd.hpp"
 #include "heaviside.hpp"
 #include "i0.hpp"
+#include "interpolate.hpp"
 #include "lcm.hpp"
 #include "ldexp.hpp"
 #include "logaddexp2.hpp"
@@ -64,6 +65,7 @@ void init_elementwise_functions(py::module_ m)
     init_gcd(m);
     init_heaviside(m);
     init_i0(m);
+    init_interpolate(m);
     init_lcm(m);
     init_ldexp(m);
     init_logaddexp2(m);