From cf436bcb98230d9db31779ab222f5820fe948620 Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 6 Jul 2023 18:03:58 +0200 Subject: [PATCH 1/8] Reuse dpctl.tensor.sqrt for dpnp.sqrt --- dpnp/dpnp_algo/dpnp_algo.pxd | 3 -- dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi | 5 --- dpnp/dpnp_algo/dpnp_elementwise_common.py | 36 +++++++++++++++++ dpnp/dpnp_iface_trigonometric.py | 45 +++++++++++----------- tests/test_umath.py | 4 +- 5 files changed, 61 insertions(+), 32 deletions(-) diff --git a/dpnp/dpnp_algo/dpnp_algo.pxd b/dpnp/dpnp_algo/dpnp_algo.pxd index b04f01f989e3..8e7721326166 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pxd +++ b/dpnp/dpnp_algo/dpnp_algo.pxd @@ -299,8 +299,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_SINH_EXT DPNP_FN_SORT DPNP_FN_SORT_EXT - DPNP_FN_SQRT - DPNP_FN_SQRT_EXT DPNP_FN_SQUARE DPNP_FN_SQUARE_EXT DPNP_FN_STD @@ -559,7 +557,6 @@ cpdef dpnp_descriptor dpnp_radians(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_recip(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_sin(dpnp_descriptor array1, dpnp_descriptor out) cpdef dpnp_descriptor dpnp_sinh(dpnp_descriptor array1) -cpdef dpnp_descriptor dpnp_sqrt(dpnp_descriptor array1, dpnp_descriptor out) cpdef dpnp_descriptor dpnp_square(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_tan(dpnp_descriptor array1, dpnp_descriptor out) cpdef dpnp_descriptor dpnp_tanh(dpnp_descriptor array1) diff --git a/dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi b/dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi index e8e2e5b16a54..4538461ca558 100644 --- a/dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi +++ b/dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi @@ -56,7 +56,6 @@ __all__ += [ 'dpnp_recip', 'dpnp_sin', 'dpnp_sinh', - 'dpnp_sqrt', 'dpnp_square', 'dpnp_tan', 'dpnp_tanh', @@ -144,10 +143,6 @@ cpdef utils.dpnp_descriptor dpnp_sinh(utils.dpnp_descriptor x1): return call_fptr_1in_1out_strides(DPNP_FN_SINH_EXT, x1) -cpdef utils.dpnp_descriptor dpnp_sqrt(utils.dpnp_descriptor x1, utils.dpnp_descriptor out): - return call_fptr_1in_1out_strides(DPNP_FN_SQRT_EXT, x1, dtype=None, out=out, where=True, func_name='sqrt') - - cpdef utils.dpnp_descriptor dpnp_square(utils.dpnp_descriptor x1): return call_fptr_1in_1out_strides(DPNP_FN_SQUARE_EXT, x1) diff --git a/dpnp/dpnp_algo/dpnp_elementwise_common.py b/dpnp/dpnp_algo/dpnp_elementwise_common.py index 8465454b7aed..ab0d99551e07 100644 --- a/dpnp/dpnp_algo/dpnp_elementwise_common.py +++ b/dpnp/dpnp_algo/dpnp_elementwise_common.py @@ -53,6 +53,7 @@ "dpnp_multiply", "dpnp_not_equal", "dpnp_subtract", + "dpnp_sqrt", ] @@ -643,3 +644,38 @@ def dpnp_subtract(x1, x2, out=None, order="K"): ) res_usm = func(x1_usm_or_scalar, x2_usm_or_scalar, out=out_usm, order=order) return dpnp_array._create_from_usm_ndarray(res_usm) + + +_sqrt_docstring_ = """ +sqrt(x, out=None, order='K') +Computes the non-negative square-root for each element `x_i` for input array `x`. +Args: + x (dpnp.ndarray): + Input array. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + usm_ndarray: + An array containing the element-wise square-root results. +""" + + +def dpnp_sqrt(x, out=None, order="K"): + """Invokes sqrt() from dpctl.tensor implementation for sqrt() function.""" + + # dpctl.tensor only works with usm_ndarray or scalar + x_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + func = UnaryElementwiseFunc( + "sqrt", + ti._sqrt_result_type, + ti._sqrt, + _sqrt_docstring_, + ) + res_usm = func(x_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) diff --git a/dpnp/dpnp_iface_trigonometric.py b/dpnp/dpnp_iface_trigonometric.py index a2be4327fa11..6e1f878fecd7 100644 --- a/dpnp/dpnp_iface_trigonometric.py +++ b/dpnp/dpnp_iface_trigonometric.py @@ -40,7 +40,6 @@ """ -import dpctl.tensor as dpt import numpy import dpnp @@ -50,6 +49,7 @@ from .dpnp_algo.dpnp_elementwise_common import ( check_nd_call_func, dpnp_log, + dpnp_sqrt, ) __all__ = [ @@ -1019,7 +1019,17 @@ def sinh(x1): return call_origin(numpy.sinh, x1, **kwargs) -def sqrt(x1, /, out=None, **kwargs): +def sqrt( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ Return the positive square-root of an array, element-wise. @@ -1030,8 +1040,8 @@ def sqrt(x1, /, out=None, **kwargs): Input array is supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. Parameter `out` is supported as class:`dpnp.ndarray`, class:`dpctl.tensor.usm_ndarray` or with default value ``None``. + Parameters `where`, `dtype` and `subok` are supported with their default values. Otherwise the function will be executed sequentially on CPU. - Keyword arguments ``kwargs`` are currently unsupported. Input array data types are limited by supported DPNP :ref:`Data types`. Examples @@ -1044,26 +1054,17 @@ def sqrt(x1, /, out=None, **kwargs): """ - x1_desc = ( - dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False - ) - if not kwargs - else None + return check_nd_call_func( + numpy.sqrt, + dpnp_sqrt, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - if out is not None: - if not isinstance(out, (dpnp.ndarray, dpt.usm_ndarray)): - raise TypeError("return array must be of supported array type") - out_desc = ( - dpnp.get_dpnp_descriptor(out, copy_when_nondefault_queue=False) - or None - ) - else: - out_desc = None - return dpnp_sqrt(x1_desc, out=out_desc).get_pyobj() - - return call_origin(numpy.sqrt, x1, out=out, **kwargs) def square(x1): diff --git a/tests/test_umath.py b/tests/test_umath.py index e2086e77c4bd..33c612abd342 100644 --- a/tests/test_umath.py +++ b/tests/test_umath.py @@ -478,7 +478,7 @@ def test_invalid_dtype(self, dtype): dp_array = dpnp.arange(10, dtype=dpnp.float32) dp_out = dpnp.empty(10, dtype=dtype) - with pytest.raises(ValueError): + with pytest.raises(TypeError): dpnp.sqrt(dp_array, out=dp_out) @pytest.mark.parametrize( @@ -488,7 +488,7 @@ def test_invalid_shape(self, shape): dp_array = dpnp.arange(10, dtype=dpnp.float32) dp_out = dpnp.empty(shape, dtype=dpnp.float32) - with pytest.raises(ValueError): + with pytest.raises(TypeError): dpnp.sqrt(dp_array, out=dp_out) @pytest.mark.parametrize( From b5b3cd9a01a0fc9bf271de57b2fb338dcd0770ae Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Fri, 7 Jul 2023 00:29:59 +0200 Subject: [PATCH 2/8] Update tests and docstrings for dpnp.sqrt --- dpnp/dpnp_algo/dpnp_elementwise_common.py | 72 +++++++++++------------ dpnp/dpnp_iface_bitwise.py | 2 +- dpnp/dpnp_iface_trigonometric.py | 14 ++++- tests/test_umath.py | 18 +++++- 4 files changed, 64 insertions(+), 42 deletions(-) diff --git a/dpnp/dpnp_algo/dpnp_elementwise_common.py b/dpnp/dpnp_algo/dpnp_elementwise_common.py index ab0d99551e07..cbabd0b0d69d 100644 --- a/dpnp/dpnp_algo/dpnp_elementwise_common.py +++ b/dpnp/dpnp_algo/dpnp_elementwise_common.py @@ -52,8 +52,8 @@ "dpnp_log", "dpnp_multiply", "dpnp_not_equal", - "dpnp_subtract", "dpnp_sqrt", + "dpnp_subtract", ] @@ -586,6 +586,41 @@ def dpnp_not_equal(x1, x2, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_sqrt_docstring_ = """ +sqrt(x, out=None, order='K') +Computes the non-negative square-root for each element `x_i` for input array `x`. +Args: + x (dpnp.ndarray): + Input array. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise square-root results. +""" + + +def dpnp_sqrt(x, out=None, order="K"): + """Invokes sqrt() from dpctl.tensor implementation for sqrt() function.""" + + # dpctl.tensor only works with usm_ndarray or scalar + x_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + func = UnaryElementwiseFunc( + "sqrt", + ti._sqrt_result_type, + ti._sqrt, + _sqrt_docstring_, + ) + res_usm = func(x_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _subtract_docstring_ = """ subtract(x1, x2, out=None, order="K") @@ -644,38 +679,3 @@ def dpnp_subtract(x1, x2, out=None, order="K"): ) res_usm = func(x1_usm_or_scalar, x2_usm_or_scalar, out=out_usm, order=order) return dpnp_array._create_from_usm_ndarray(res_usm) - - -_sqrt_docstring_ = """ -sqrt(x, out=None, order='K') -Computes the non-negative square-root for each element `x_i` for input array `x`. -Args: - x (dpnp.ndarray): - Input array. - out ({None, dpnp.ndarray}, optional): - Output array to populate. Array must have the correct - shape and the expected data type. - order ("C","F","A","K", optional): memory layout of the new - output array, if parameter `out` is `None`. - Default: "K". -Return: - usm_ndarray: - An array containing the element-wise square-root results. -""" - - -def dpnp_sqrt(x, out=None, order="K"): - """Invokes sqrt() from dpctl.tensor implementation for sqrt() function.""" - - # dpctl.tensor only works with usm_ndarray or scalar - x_usm = dpnp.get_usm_ndarray(x) - out_usm = None if out is None else dpnp.get_usm_ndarray(out) - - func = UnaryElementwiseFunc( - "sqrt", - ti._sqrt_result_type, - ti._sqrt, - _sqrt_docstring_, - ) - res_usm = func(x_usm, out=out_usm, order=order) - return dpnp_array._create_from_usm_ndarray(res_usm) diff --git a/dpnp/dpnp_iface_bitwise.py b/dpnp/dpnp_iface_bitwise.py index ecaa68ec2f75..4bffb71417fb 100644 --- a/dpnp/dpnp_iface_bitwise.py +++ b/dpnp/dpnp_iface_bitwise.py @@ -133,7 +133,7 @@ def bitwise_and(x1, x2, dtype=None, out=None, where=True, **kwargs): Returns ------- y : dpnp.ndarray - An array containing the element-wise results. + An array containing the element-wise results of positive square root. Limitations ----------- diff --git a/dpnp/dpnp_iface_trigonometric.py b/dpnp/dpnp_iface_trigonometric.py index 6e1f878fecd7..1a4db5c1e191 100644 --- a/dpnp/dpnp_iface_trigonometric.py +++ b/dpnp/dpnp_iface_trigonometric.py @@ -1035,6 +1035,11 @@ def sqrt( For full documentation refer to :obj:`numpy.sqrt`. + Returns + ------- + y : dpnp.ndarray + An array containing the element-wise results of positive square root. + Limitations ----------- Input array is supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. @@ -1048,9 +1053,12 @@ def sqrt( -------- >>> import dpnp as np >>> x = np.array([1, 4, 9]) - >>> out = np.sqrt(x) - >>> [i for i in out] - [1.0, 2.0, 3.0] + >>> np.sqrt(x) + array([1., 2., 3.]) + + >>> x2 = np.array([4, -1, np.inf]) + >>> np.sqrt(x2) + array([ 2., nan, inf]) """ diff --git a/tests/test_umath.py b/tests/test_umath.py index 33c612abd342..c4bcab397703 100644 --- a/tests/test_umath.py +++ b/tests/test_umath.py @@ -7,7 +7,7 @@ from .helper import ( get_all_dtypes, get_complex_dtypes, - get_float_dtypes, + get_float_complex_dtypes, has_support_aspect16, has_support_aspect64, ) @@ -454,9 +454,23 @@ def test_invalid_shape(self, shape): class TestSqrt: - @pytest.mark.parametrize("dtype", get_float_dtypes()) + @pytest.mark.parametrize("dtype", get_all_dtypes()) def test_sqrt_ordinary(self, dtype): array_data = numpy.arange(10) + + # DPNP + dp_array = dpnp.array(array_data, dtype=dtype) + result = dpnp.sqrt(dp_array) + + # original + np_array = numpy.array(array_data, dtype=dtype) + expected = numpy.sqrt(np_array) + + numpy.testing.assert_allclose(expected, result) + + @pytest.mark.parametrize("dtype", get_float_complex_dtypes()) + def test_sqrt_out(self, dtype): + array_data = numpy.arange(10) out = numpy.empty(10, dtype=dtype) # DPNP From bcc20ff1aa33b4af38286ed79d55557a994ce909 Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Fri, 7 Jul 2023 14:10:53 +0200 Subject: [PATCH 3/8] Add sqrt call from OneMKL by pybind11 extension --- dpnp/backend/extensions/vm/sqrt.hpp | 78 +++++++++++++++++++++ dpnp/backend/extensions/vm/types_matrix.hpp | 19 +++++ dpnp/backend/extensions/vm/vm_py.cpp | 32 +++++++++ dpnp/dpnp_algo/dpnp_elementwise_common.py | 20 +++++- 4 files changed, 147 insertions(+), 2 deletions(-) create mode 100644 dpnp/backend/extensions/vm/sqrt.hpp diff --git a/dpnp/backend/extensions/vm/sqrt.hpp b/dpnp/backend/extensions/vm/sqrt.hpp new file mode 100644 index 000000000000..df20ee4b0509 --- /dev/null +++ b/dpnp/backend/extensions/vm/sqrt.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, 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 + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event sqrt_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::sqrt(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct SqrtContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::SqrtOutputType::value_type, void>) + { + return nullptr; + } + else { + return sqrt_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/types_matrix.hpp b/dpnp/backend/extensions/vm/types_matrix.hpp index 26235b838721..28a9d4d87f6f 100644 --- a/dpnp/backend/extensions/vm/types_matrix.hpp +++ b/dpnp/backend/extensions/vm/types_matrix.hpp @@ -86,6 +86,25 @@ struct LnOutputType dpctl_td_ns::TypeMapResultEntry, dpctl_td_ns::DefaultResultEntry>::result_type; }; + +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::sqrt function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct SqrtOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns:: + TypeMapResultEntry, std::complex>, + dpctl_td_ns:: + TypeMapResultEntry, std::complex>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; } // namespace types } // namespace vm } // namespace ext diff --git a/dpnp/backend/extensions/vm/vm_py.cpp b/dpnp/backend/extensions/vm/vm_py.cpp index 59e6612b9ed8..2306d0583cdb 100644 --- a/dpnp/backend/extensions/vm/vm_py.cpp +++ b/dpnp/backend/extensions/vm/vm_py.cpp @@ -33,6 +33,7 @@ #include "common.hpp" #include "div.hpp" #include "ln.hpp" +#include "sqrt.hpp" #include "types_matrix.hpp" namespace py = pybind11; @@ -44,6 +45,7 @@ using vm_ext::unary_impl_fn_ptr_t; static binary_impl_fn_ptr_t div_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t ln_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t sqrt_dispatch_vector[dpctl_td_ns::num_types]; PYBIND11_MODULE(_vm_impl, m) { @@ -107,4 +109,34 @@ PYBIND11_MODULE(_vm_impl, m) "OneMKL VM library can be used", py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); } + + // UnaryUfunc: ==== Sqrt(x) ==== + { + vm_ext::init_ufunc_dispatch_vector( + sqrt_dispatch_vector); + + auto sqrt_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + sqrt_dispatch_vector); + }; + m.def( + "_sqrt", sqrt_pyapi, + "Call `sqrt` from OneMKL VM library to performs element by element " + "operation of extracting the square root " + "of vector `src` to resulting vector `dst`", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto sqrt_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + sqrt_dispatch_vector); + }; + m.def("_mkl_sqrt_to_call", sqrt_need_to_call_pyapi, + "Check input arguments to answer if `sqrt` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } } diff --git a/dpnp/dpnp_algo/dpnp_elementwise_common.py b/dpnp/dpnp_algo/dpnp_elementwise_common.py index cbabd0b0d69d..a0dd7ad336b6 100644 --- a/dpnp/dpnp_algo/dpnp_elementwise_common.py +++ b/dpnp/dpnp_algo/dpnp_elementwise_common.py @@ -605,7 +605,23 @@ def dpnp_not_equal(x1, x2, out=None, order="K"): def dpnp_sqrt(x, out=None, order="K"): - """Invokes sqrt() from dpctl.tensor implementation for sqrt() function.""" + """ + Invokes sqrt() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for sqrt() function. + + """ + + def _call_sqrt(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_sqrt_to_call(sycl_queue, src, dst): + # call pybind11 extension for sqrt() function from OneMKL VM + return vmi._sqrt(sycl_queue, src, dst, depends) + return ti._sqrt(src, dst, sycl_queue, depends) # dpctl.tensor only works with usm_ndarray or scalar x_usm = dpnp.get_usm_ndarray(x) @@ -614,7 +630,7 @@ def dpnp_sqrt(x, out=None, order="K"): func = UnaryElementwiseFunc( "sqrt", ti._sqrt_result_type, - ti._sqrt, + _call_sqrt, _sqrt_docstring_, ) res_usm = func(x_usm, out=out_usm, order=order) From cc774a7f349fa9ca46903eca29a20cdeda8a32fe Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Fri, 7 Jul 2023 14:53:59 +0200 Subject: [PATCH 4/8] Update test_umath for dpnp.sqrt --- dpnp/linalg/dpnp_algo_linalg.pyx | 2 +- tests/test_umath.py | 62 ++++++++++++++++++++++---------- 2 files changed, 44 insertions(+), 20 deletions(-) diff --git a/dpnp/linalg/dpnp_algo_linalg.pyx b/dpnp/linalg/dpnp_algo_linalg.pyx index 56b18935ce42..69a3efa223f7 100644 --- a/dpnp/linalg/dpnp_algo_linalg.pyx +++ b/dpnp/linalg/dpnp_algo_linalg.pyx @@ -366,7 +366,7 @@ cpdef object dpnp_norm(object input, ord=None, axis=None): input = dpnp.ravel(input, order='K') sqnorm = dpnp.dot(input, input) - ret = dpnp.sqrt([sqnorm]) + ret = dpnp.sqrt(sqnorm) return dpnp.array(ret.reshape(1, *ret.shape), dtype=res_type) len_axis = 1 if axis is None else len(axis_) diff --git a/tests/test_umath.py b/tests/test_umath.py index c4bcab397703..c27af9ca4594 100644 --- a/tests/test_umath.py +++ b/tests/test_umath.py @@ -454,36 +454,60 @@ def test_invalid_shape(self, shape): class TestSqrt: - @pytest.mark.parametrize("dtype", get_all_dtypes()) - def test_sqrt_ordinary(self, dtype): - array_data = numpy.arange(10) + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_bool=True, no_complex=True) + ) + def test_sqrt_int_float(self, dtype): + np_array = numpy.arange(10, dtype=dtype) + np_out = numpy.empty(10, dtype=numpy.float64) # DPNP - dp_array = dpnp.array(array_data, dtype=dtype) - result = dpnp.sqrt(dp_array) + dp_out_dtype = dpnp.float32 + if has_support_aspect64() and dtype != dpnp.float32: + dp_out_dtype = dpnp.float64 - # original - np_array = numpy.array(array_data, dtype=dtype) - expected = numpy.sqrt(np_array) + dp_out = dpnp.array(np_out, dtype=dp_out_dtype) + dp_array = dpnp.array(np_array, dtype=dtype) + result = dpnp.sqrt(dp_array, out=dp_out) - numpy.testing.assert_allclose(expected, result) + # original + expected = numpy.sqrt(np_array, out=np_out) + assert_allclose(expected, result) - @pytest.mark.parametrize("dtype", get_float_complex_dtypes()) - def test_sqrt_out(self, dtype): - array_data = numpy.arange(10) - out = numpy.empty(10, dtype=dtype) + @pytest.mark.parametrize("dtype", get_complex_dtypes()) + def test_sqrt_complex(self, dtype): + np_array = numpy.arange(10, 20, dtype=dtype) + np_out = numpy.empty(10, dtype=numpy.complex128) # DPNP - dp_array = dpnp.array(array_data, dtype=dtype) - dp_out = dpnp.array(out, dtype=dtype) + dp_out_dtype = dpnp.complex64 + if has_support_aspect64() and dtype != dpnp.complex64: + dp_out_dtype = dpnp.complex128 + + dp_out = dpnp.array(np_out, dtype=dp_out_dtype) + dp_array = dpnp.array(np_array, dtype=dtype) result = dpnp.sqrt(dp_array, out=dp_out) # original - np_array = numpy.array(array_data, dtype=dtype) - expected = numpy.sqrt(np_array, out=out) + expected = numpy.sqrt(np_array, out=np_out) + assert_allclose(expected, result) + + @pytest.mark.usefixtures("suppress_divide_numpy_warnings") + @pytest.mark.skipif( + not has_support_aspect16(), reason="No fp16 support by device" + ) + def test_sqrt_bool(self): + np_array = numpy.arange(2, dtype=numpy.bool_) + np_out = numpy.empty(2, dtype=numpy.float16) + + # DPNP + dp_array = dpnp.array(np_array, dtype=np_array.dtype) + dp_out = dpnp.array(np_out, dtype=np_out.dtype) + result = dpnp.sqrt(dp_array, out=dp_out) - numpy.testing.assert_allclose(expected, result) - numpy.testing.assert_allclose(out, dp_out) + # original + expected = numpy.sqrt(np_array, out=np_out) + assert_allclose(expected, result) @pytest.mark.parametrize( "dtype", [numpy.int64, numpy.int32], ids=["numpy.int64", "numpy.int32"] From cef96177ce0e699b640a4f7413ec83e1473ddf85 Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Fri, 7 Jul 2023 18:38:39 +0200 Subject: [PATCH 5/8] Remove DPNP_FN_SQRT_EXT and update docstrings --- dpnp/backend/include/dpnp_iface_fptr.hpp | 2 -- dpnp/backend/kernels/dpnp_krnl_elemwise.cpp | 9 --------- dpnp/dpnp_iface_trigonometric.py | 8 ++++++-- tests/test_umath.py | 1 - 4 files changed, 6 insertions(+), 14 deletions(-) diff --git a/dpnp/backend/include/dpnp_iface_fptr.hpp b/dpnp/backend/include/dpnp_iface_fptr.hpp index 335c34bad583..f68361a49cb2 100644 --- a/dpnp/backend/include/dpnp_iface_fptr.hpp +++ b/dpnp/backend/include/dpnp_iface_fptr.hpp @@ -483,8 +483,6 @@ enum class DPNPFuncName : size_t DPNP_FN_SORT_EXT, /**< Used in numpy.sort() impl, requires extra parameters */ DPNP_FN_SQRT, /**< Used in numpy.sqrt() impl */ - DPNP_FN_SQRT_EXT, /**< Used in numpy.sqrt() impl, requires extra parameters - */ DPNP_FN_SQUARE, /**< Used in numpy.square() impl */ DPNP_FN_SQUARE_EXT, /**< Used in numpy.square() impl, requires extra parameters */ diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index 74141e4bd5f8..88008b55b354 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -747,15 +747,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_SQRT][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_sqrt_c_default}; - fmap[DPNPFuncName::DPNP_FN_SQRT_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_sqrt_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SQRT_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_sqrt_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SQRT_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_sqrt_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SQRT_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_sqrt_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TAN][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_tan_c_default}; fmap[DPNPFuncName::DPNP_FN_TAN][eft_LNG][eft_LNG] = { diff --git a/dpnp/dpnp_iface_trigonometric.py b/dpnp/dpnp_iface_trigonometric.py index 1a4db5c1e191..44a0285dd4c4 100644 --- a/dpnp/dpnp_iface_trigonometric.py +++ b/dpnp/dpnp_iface_trigonometric.py @@ -1031,14 +1031,18 @@ def sqrt( **kwargs, ): """ - Return the positive square-root of an array, element-wise. + Return the non-negative square-root of an array, element-wise. For full documentation refer to :obj:`numpy.sqrt`. Returns ------- y : dpnp.ndarray - An array containing the element-wise results of positive square root. + An array of the same shape as `x`, containing the positive + square-root of each element in `x`. If any element in `x` is + complex, a complex array is returned (and the square-roots of + negative reals are calculated). If all of the elements in `x` + are real, so is `y`, with negative elements returning ``nan``. Limitations ----------- diff --git a/tests/test_umath.py b/tests/test_umath.py index c27af9ca4594..f1a5a738a734 100644 --- a/tests/test_umath.py +++ b/tests/test_umath.py @@ -7,7 +7,6 @@ from .helper import ( get_all_dtypes, get_complex_dtypes, - get_float_complex_dtypes, has_support_aspect16, has_support_aspect64, ) From 93bf4d22caa639ca2d537f4bdf7fd0de03d4355e Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Fri, 7 Jul 2023 21:17:12 +0200 Subject: [PATCH 6/8] Return deleted DPNP_FN_SQRT_EXT --- dpnp/backend/include/dpnp_iface_fptr.hpp | 2 ++ dpnp/backend/kernels/dpnp_krnl_elemwise.cpp | 6 ++++++ 2 files changed, 8 insertions(+) diff --git a/dpnp/backend/include/dpnp_iface_fptr.hpp b/dpnp/backend/include/dpnp_iface_fptr.hpp index f68361a49cb2..335c34bad583 100644 --- a/dpnp/backend/include/dpnp_iface_fptr.hpp +++ b/dpnp/backend/include/dpnp_iface_fptr.hpp @@ -483,6 +483,8 @@ enum class DPNPFuncName : size_t DPNP_FN_SORT_EXT, /**< Used in numpy.sort() impl, requires extra parameters */ DPNP_FN_SQRT, /**< Used in numpy.sqrt() impl */ + DPNP_FN_SQRT_EXT, /**< Used in numpy.sqrt() impl, requires extra parameters + */ DPNP_FN_SQUARE, /**< Used in numpy.square() impl */ DPNP_FN_SQUARE_EXT, /**< Used in numpy.square() impl, requires extra parameters */ diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index 88008b55b354..a46d3f710b5f 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -747,6 +747,12 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_SQRT][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_sqrt_c_default}; + // Used in dpnp_std_c + fmap[DPNPFuncName::DPNP_FN_SQRT_EXT][eft_FLT][eft_FLT] = { + eft_FLT, (void *)dpnp_sqrt_c_ext}; + fmap[DPNPFuncName::DPNP_FN_SQRT_EXT][eft_DBL][eft_DBL] = { + eft_DBL, (void *)dpnp_sqrt_c_ext}; + fmap[DPNPFuncName::DPNP_FN_TAN][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_tan_c_default}; fmap[DPNPFuncName::DPNP_FN_TAN][eft_LNG][eft_LNG] = { From 08fe523ad84266460b6cc48af647e65df0b9f23a Mon Sep 17 00:00:00 2001 From: Anton <100830759+antonwolfy@users.noreply.github.com> Date: Fri, 7 Jul 2023 23:09:46 +0200 Subject: [PATCH 7/8] Update dpnp/backend/extensions/vm/vm_py.cpp --- dpnp/backend/extensions/vm/vm_py.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dpnp/backend/extensions/vm/vm_py.cpp b/dpnp/backend/extensions/vm/vm_py.cpp index 88fee3480ee1..6fc255fceff3 100644 --- a/dpnp/backend/extensions/vm/vm_py.cpp +++ b/dpnp/backend/extensions/vm/vm_py.cpp @@ -156,7 +156,7 @@ PYBIND11_MODULE(_vm_impl, m) }; m.def("_sin", sin_pyapi, "Call `sin` function from OneMKL VM library to compute " - "natural logarithm of vector elements", + "sine of vector elements", py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), py::arg("depends") = py::list()); From d84ba86ffe9572d15be94c1fd4ca4f806468e5f8 Mon Sep 17 00:00:00 2001 From: Anton <100830759+antonwolfy@users.noreply.github.com> Date: Fri, 7 Jul 2023 23:12:49 +0200 Subject: [PATCH 8/8] Update dpnp/backend/extensions/vm/vm_py.cpp --- dpnp/backend/extensions/vm/vm_py.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/dpnp/backend/extensions/vm/vm_py.cpp b/dpnp/backend/extensions/vm/vm_py.cpp index 6fc255fceff3..aeb12ac51094 100644 --- a/dpnp/backend/extensions/vm/vm_py.cpp +++ b/dpnp/backend/extensions/vm/vm_py.cpp @@ -50,7 +50,6 @@ static unary_impl_fn_ptr_t cos_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t ln_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t sin_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t sqrt_dispatch_vector[dpctl_td_ns::num_types]; -static unary_impl_fn_ptr_t sin_dispatch_vector[dpctl_td_ns::num_types]; PYBIND11_MODULE(_vm_impl, m) {