Rework implementation of dpnp.deg2rad and dpnp.radians functions

doc/reference/ufunc.rst

@@ -54,6 +54,8 @@ Math operations
 
 Trigonometric functions
 ~~~~~~~~~~~~~~~~~~~~~~~
+All trigonometric functions use radians when an angle is called for.
+The ratio of degrees to radians is :math:`180^{\circ}/\pi.`
 
 .. autosummary::
    :toctree: generated/
@@ -73,6 +75,8 @@ Trigonometric functions
    dpnp.arcsinh
    dpnp.arccosh
    dpnp.arctanh
+   dpnp.degrees
+   dpnp.radians
    dpnp.deg2rad
    dpnp.rad2deg
 

dpnp/backend/extensions/ufunc/CMakeLists.txt

@@ -29,6 +29,7 @@ set(_elementwise_sources
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmax.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmin.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmod.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/radians.cpp
 )
 
 set(python_module_name _ufunc_impl)

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

@@ -29,6 +29,7 @@
 #include "fmax.hpp"
 #include "fmin.hpp"
 #include "fmod.hpp"
+#include "radians.hpp"
 
 namespace py = pybind11;
 
@@ -43,5 +44,6 @@ void init_elementwise_functions(py::module_ m)
     init_fmax(m);
     init_fmin(m);
     init_fmod(m);
+    init_radians(m);
 }
 } // namespace dpnp::extensions::ufunc

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

@@ -40,19 +40,16 @@
 #include "kernels/elementwise_functions/common.hpp"
 #include "utils/type_dispatch.hpp"
 
-namespace py = pybind11;
-
 namespace dpnp::extensions::ufunc
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
+namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
-namespace td_ns = dpctl::tensor::type_dispatch;
-
-using ew_cmn_ns::unary_contig_impl_fn_ptr_t;
-using ew_cmn_ns::unary_strided_impl_fn_ptr_t;
 
 namespace impl
 {
+namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
+namespace td_ns = dpctl::tensor::type_dispatch;
+
 /**
  * @brief A factory to define pairs of supported types for which
  * sycl::fabs<T> function is available.

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

@@ -41,20 +41,17 @@
 #include "kernels/elementwise_functions/maximum.hpp"
 #include "utils/type_dispatch.hpp"
 
-namespace py = pybind11;
-
 namespace dpnp::extensions::ufunc
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
-namespace max_ns = dpctl::tensor::kernels::maximum;
+namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
 
-using ew_cmn_ns::unary_contig_impl_fn_ptr_t;
-using ew_cmn_ns::unary_strided_impl_fn_ptr_t;
-
 namespace impl
 {
+namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
+namespace max_ns = dpctl::tensor::kernels::maximum;
+
 // Supports the same types table as for maximum function in dpctl
 template <typename T1, typename T2>
 using OutputType = max_ns::MaximumOutputType<T1, T2>;

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

@@ -41,20 +41,17 @@
 #include "kernels/elementwise_functions/minimum.hpp"
 #include "utils/type_dispatch.hpp"
 
-namespace py = pybind11;
-
 namespace dpnp::extensions::ufunc
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
-namespace min_ns = dpctl::tensor::kernels::minimum;
+namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
 
-using ew_cmn_ns::unary_contig_impl_fn_ptr_t;
-using ew_cmn_ns::unary_strided_impl_fn_ptr_t;
-
 namespace impl
 {
+namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
+namespace min_ns = dpctl::tensor::kernels::minimum;
+
 // Supports the same types table as for minimum function in dpctl
 template <typename T1, typename T2>
 using OutputType = min_ns::MinimumOutputType<T1, T2>;

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

@@ -40,19 +40,16 @@
 #include "kernels/elementwise_functions/common.hpp"
 #include "utils/type_dispatch.hpp"
 
-namespace py = pybind11;
-
 namespace dpnp::extensions::ufunc
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
+namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
-namespace td_ns = dpctl::tensor::type_dispatch;
-
-using ew_cmn_ns::unary_contig_impl_fn_ptr_t;
-using ew_cmn_ns::unary_strided_impl_fn_ptr_t;
 
 namespace impl
 {
+namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
+namespace td_ns = dpctl::tensor::type_dispatch;
+
 /**
  * @brief A factory to define pairs of supported types for which
  * sycl::fmod<T> function is available.

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

@@ -0,0 +1,127 @@
+//*****************************************************************************
+// Copyright (c) 2024, Intel Corporation
+// All rights reserved.
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+// - Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// - Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
+// THE POSSIBILITY OF SUCH DAMAGE.
+//*****************************************************************************
+
+#include <sycl/sycl.hpp>
+
+#include "dpctl4pybind11.hpp"
+
+#include "kernels/elementwise_functions/radians.hpp"
+#include "populate.hpp"
+#include "radians.hpp"
+
+// include a local copy of elementwise common header from dpctl tensor:
+// dpctl/tensor/libtensor/source/elementwise_functions/elementwise_functions.hpp
+// TODO: replace by including dpctl header once available
+#include "../../elementwise_functions/elementwise_functions.hpp"
+
+// dpctl tensor headers
+#include "kernels/elementwise_functions/common.hpp"
+#include "utils/type_dispatch.hpp"
+
+namespace dpnp::extensions::ufunc
+{
+namespace py = pybind11;
+namespace py_int = dpnp::extensions::py_internal;
+
+namespace impl
+{
+namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
+namespace td_ns = dpctl::tensor::type_dispatch;
+
+/**
+ * @brief A factory to define pairs of supported types for which
+ * sycl::radians<T> function is available.
+ *
+ * @tparam T Type of input vector `a` and of result vector `y`.
+ */
+template <typename T>
+struct OutputType
+{
+    using value_type =
+        typename std::disjunction<td_ns::TypeMapResultEntry<T, sycl::half>,
+                                  td_ns::TypeMapResultEntry<T, float>,
+                                  td_ns::TypeMapResultEntry<T, double>,
+                                  td_ns::DefaultResultEntry<void>>::result_type;
+};
+
+using dpnp::kernels::radians::RadiansFunctor;
+
+template <typename argT,
+          typename resT = argT,
+          unsigned int vec_sz = 4,
+          unsigned int n_vecs = 2,
+          bool enable_sg_loadstore = true>
+using ContigFunctor = ew_cmn_ns::UnaryContigFunctor<argT,
+                                                    resT,
+                                                    RadiansFunctor<argT, resT>,
+                                                    vec_sz,
+                                                    n_vecs,
+                                                    enable_sg_loadstore>;
+
+template <typename argTy, typename resTy, typename IndexerT>
+using StridedFunctor = ew_cmn_ns::
+    UnaryStridedFunctor<argTy, resTy, IndexerT, RadiansFunctor<argTy, resTy>>;
+
+using ew_cmn_ns::unary_contig_impl_fn_ptr_t;
+using ew_cmn_ns::unary_strided_impl_fn_ptr_t;
+
+static unary_contig_impl_fn_ptr_t
+    radians_contig_dispatch_vector[td_ns::num_types];
+static int radians_output_typeid_vector[td_ns::num_types];
+static unary_strided_impl_fn_ptr_t
+    radians_strided_dispatch_vector[td_ns::num_types];
+
+MACRO_POPULATE_DISPATCH_VECTORS(radians);
+} // namespace impl
+
+void init_radians(py::module_ m)
+{
+    using arrayT = dpctl::tensor::usm_ndarray;
+    using event_vecT = std::vector<sycl::event>;
+    {
+        impl::populate_radians_dispatch_vectors();
+        using impl::radians_contig_dispatch_vector;
+        using impl::radians_output_typeid_vector;
+        using impl::radians_strided_dispatch_vector;
+
+        auto radians_pyapi = [&](const arrayT &src, const arrayT &dst,
+                                 sycl::queue &exec_q,
+                                 const event_vecT &depends = {}) {
+            return py_int::py_unary_ufunc(src, dst, exec_q, depends,
+                                          radians_output_typeid_vector,
+                                          radians_contig_dispatch_vector,
+                                          radians_strided_dispatch_vector);
+        };
+        m.def("_radians", radians_pyapi, "", py::arg("src"), py::arg("dst"),
+              py::arg("sycl_queue"), py::arg("depends") = py::list());
+
+        auto radians_result_type_pyapi = [&](const py::dtype &dtype) {
+            return py_int::py_unary_ufunc_result_type(
+                dtype, radians_output_typeid_vector);
+        };
+        m.def("_radians_result_type", radians_result_type_pyapi);
+    }
+}
+} // namespace dpnp::extensions::ufunc

dpnp/backend/extensions/ufunc/elementwise_functions/radians.hpp

@@ -0,0 +1,35 @@
+//*****************************************************************************
+// Copyright (c) 2024, Intel Corporation
+// All rights reserved.
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+// - Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// - Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
+// THE POSSIBILITY OF SUCH DAMAGE.
+//*****************************************************************************
+
+#pragma once
+
+#include <pybind11/pybind11.h>
+
+namespace py = pybind11;
+
+namespace dpnp::extensions::ufunc
+{
+void init_radians(py::module_ m);
+} // namespace dpnp::extensions::ufunc

dpnp/backend/include/dpnp_gen_1arg_2type_tbl.hpp

@@ -89,9 +89,6 @@ MACRO_1ARG_2TYPES_OP(dpnp_copyto_c, input_elem, q.submit(kernel_func))
 MACRO_1ARG_2TYPES_OP(dpnp_degrees_c,
                      sycl::degrees(input_elem),
                      q.submit(kernel_func))
-MACRO_1ARG_2TYPES_OP(dpnp_radians_c,
-                     sycl::radians(input_elem),
-                     q.submit(kernel_func))
 MACRO_1ARG_2TYPES_OP(dpnp_sqrt_c,
                      sycl::sqrt(input_elem),
                      oneapi::mkl::vm::sqrt(q, input1_size, input1_data, result))

dpnp/backend/include/dpnp_iface_fptr.hpp

@@ -116,9 +116,6 @@ enum class DPNPFuncName : size_t
     DPNP_FN_PARTITION_EXT, /**< Used in numpy.partition() impl, requires extra
                               parameters */
     DPNP_FN_PROD,          /**< Used in numpy.prod() impl  */
-    DPNP_FN_RADIANS,       /**< Used in numpy.radians() impl  */
-    DPNP_FN_RADIANS_EXT,   /**< Used in numpy.radians() impl, requires extra
-                              parameters */
     DPNP_FN_RNG_BETA,      /**< Used in numpy.random.beta() impl  */
     DPNP_FN_RNG_BETA_EXT,  /**< Used in numpy.random.beta() impl, requires extra
                               parameters */

dpnp/backend/kernels/dpnp_krnl_elemwise.cpp

@@ -337,36 +337,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap)
     fmap[DPNPFuncName::DPNP_FN_DEGREES_EXT][eft_DBL][eft_DBL] = {
         eft_DBL, (void *)dpnp_degrees_c_ext<double, double>};
 
-    fmap[DPNPFuncName::DPNP_FN_RADIANS][eft_INT][eft_INT] = {
-        eft_DBL, (void *)dpnp_radians_c_default<int32_t, double>};
-    fmap[DPNPFuncName::DPNP_FN_RADIANS][eft_LNG][eft_LNG] = {
-        eft_DBL, (void *)dpnp_radians_c_default<int64_t, double>};
-    fmap[DPNPFuncName::DPNP_FN_RADIANS][eft_FLT][eft_FLT] = {
-        eft_FLT, (void *)dpnp_radians_c_default<float, float>};
-    fmap[DPNPFuncName::DPNP_FN_RADIANS][eft_DBL][eft_DBL] = {
-        eft_DBL, (void *)dpnp_radians_c_default<double, double>};
-
-    fmap[DPNPFuncName::DPNP_FN_RADIANS_EXT][eft_INT][eft_INT] = {
-        get_default_floating_type(),
-        (void *)dpnp_radians_c_ext<
-            int32_t, func_type_map_t::find_type<get_default_floating_type()>>,
-        get_default_floating_type<std::false_type>(),
-        (void *)dpnp_radians_c_ext<
-            int32_t, func_type_map_t::find_type<
-                         get_default_floating_type<std::false_type>()>>};
-    fmap[DPNPFuncName::DPNP_FN_RADIANS_EXT][eft_LNG][eft_LNG] = {
-        get_default_floating_type(),
-        (void *)dpnp_radians_c_ext<
-            int64_t, func_type_map_t::find_type<get_default_floating_type()>>,
-        get_default_floating_type<std::false_type>(),
-        (void *)dpnp_radians_c_ext<
-            int64_t, func_type_map_t::find_type<
-                         get_default_floating_type<std::false_type>()>>};
-    fmap[DPNPFuncName::DPNP_FN_RADIANS_EXT][eft_FLT][eft_FLT] = {
-        eft_FLT, (void *)dpnp_radians_c_ext<float, float>};
-    fmap[DPNPFuncName::DPNP_FN_RADIANS_EXT][eft_DBL][eft_DBL] = {
-        eft_DBL, (void *)dpnp_radians_c_ext<double, double>};
-
     fmap[DPNPFuncName::DPNP_FN_SQRT][eft_INT][eft_INT] = {
         eft_DBL, (void *)dpnp_sqrt_c_default<int32_t, double>};
     fmap[DPNPFuncName::DPNP_FN_SQRT][eft_LNG][eft_LNG] = {